DAY14:阅读CUDA C runtime之错误检查和Call stack

今天这个部分讲完后,下期将开始讲解 Texture and Surface Memory

3.2.9. Error Checking

All runtime functions return an error code, but for an asynchronous function (see Asynchronous Concurrent Execution), this error code cannot possibly report any of the asynchronous errors that could occur on the device since the function returns before the device has completed the task; the error code only reports errors that occur on the host prior to executing the task, typically related to parameter validation; if an asynchronous error occurs, it will be reported by some subsequent unrelated runtime function call.

The only way to check for asynchronous errors just after some asynchronous function call is therefore to synchronize just after the call by calling cudaDeviceSynchronize() (or by using any other synchronization mechanisms described in Asynchronous Concurrent Execution) and checking the error code returned by cudaDeviceSynchronize().

The runtime maintains an error variable for each host thread that is initialized to cudaSuccess and is overwritten by the error code every time an error occurs (be it a parameter validation error or an asynchronous error). cudaPeekAtLastError() returns this variable. cudaGetLastError() returns this variable and resets it to cudaSuccess.

Kernel launches do not return any error code, so cudaPeekAtLastError() or cudaGetLastError() must be called just after the kernel launch to retrieve any pre-launch errors. To ensure that any error returned by cudaPeekAtLastError() or cudaGetLastError() does not originate from calls prior to the kernel launch, one has to make sure that the runtime error variable is set to cudaSuccess just before the kernel launch, for example, by calling cudaGetLastError() just before the kernel launch. Kernel launches are asynchronous, so to check for asynchronous errors, the application must synchronize in-between the kernel launch and the call to cudaPeekAtLastError() or cudaGetLastError().

Note that cudaErrorNotReady that may be returned by cudaStreamQuery() and cudaEventQuery() is not considered an error and is therefore not reported by cudaPeekAtLastError() or cudaGetLastError().

3.2.10. Call Stack

On devices of compute capability 2.x and higher, the size of the call stack can be queried using cudaDeviceGetLimit() and set using cudaDeviceSetLimit().

When the call stack overflows, the kernel call fails with a stack overflow error if the application is run via a CUDA debugger (cuda-gdb, Nsight) or an unspecified launch error, otherwise.

本文备注/经验分享:

Error Checking——

所有的Runtime函数都返回错误代码。但是对于异步函数来说,返回的代码不会报告设备上将来可能发生的错误。因为异步函数在设备完成任务之前,就已经返回了。该错误代码只会报告host上发生的,任务开始执行之前的错误----这一般都是参数验证方面的。 举个例子说,不仅仅是Runtime函数,kernel启动也是如此,

<<<>>>如果导致了错误,可以在立刻的cudaGetLastError()上得知。(因为<<<>>>的错误不能直接写在一行),而我们知道<<<>>>是启动了异步任务(kernel),所以立刻得到的错误,只是实质的该kernel启动所可能发生的所有任务的一部分。例如说,用户要求使用了200KB的shared memory, 或者要求使用3000个线程的block,这些是做不到的,会立刻在跟随的cudaGetLastError()出错。这就是说的,相关参数信息验证错误,因为现在所有的计算能力都没有200KB的shared memory,也不能支持这么巨大的block,如果这第一步验证通过(参数相关),则kernel会在设备上开始启动,此时将随时可能发生第二步的错误,例如kernel中途访存挂了,这个时候因为时机已经错过去了,那么只能由后续的其他无关调用(例如你分配一段显存)来报告。所以用户会看到一个奇特的现象:后面无论调用什么都是错的,哪怕检查了代码,后面的行写的看上去都没有问题。用户就会困惑。而此章节,很好的解说了这点。我们的技术论坛上有一个帖子:http://bbs.gpuworld.cn/thread-58784-1-3.html;

这就是异步错误。很多新人都不知道这点,导致无法定位错误位置(他们会定位到其他不相关的后续位置,例如这个楼主),而现在,你通过阅读本章节,已经能成功的识别这种情况了。 需要说明的是,因为异步错误发生需要一定的时间, 例如刚才说的kernel,kernel可能执行到100ms后,才会挂,此时可能在kernel启动后很久才能出现错误,例如:

<<<>>> A B C D cudaMemcpy之类的会等待的函数。则错误可能在A,B,C,D处都没事。到了很久之后的cudaMemcpy才出现错误。 此时错误就和错误的发生源(实际上的kernel)相隔很远了。需要用户认真往上找才可以。这种错误一般总是会在后面的cudaDeviceSynchronize(), cudaMemcpy()这种同步调用才会暴露,因为同步调用会等待,无论kernel是100ms后挂了,还是200ms后挂了,同步调用总会在那里等着,因此几乎总能在这些地方暴露出之前的错误。

Call Stack—— 计算能力2.0+开始(Fermi开始),能开始支持设备函数调用了。例如有:__global__ 的A()和__device__的B(),现在A能真的调用B了,像在CPU上那样的。以前只能进行Inline操作,也就是将B嵌入到A中。现在可以直接调用了。这样如果以前的CUDA代码,A调用了好多次B,以前会导致嵌入多次B,代码膨胀,性能下降。而现在只有一份B的代码存在,A每次通过call stack那样真的调用B,代码体积缩小,I-Cache命中率提高,性能上升。不过需要说明的是,如果调用的太多层,call stack可能会不够(overflow,溢出),此时可以手工增大。类似的情况还有动态并行(计算能力3.5+),如果同步深度过大,也需要手工增大的。不过是另外一个参数了。以及,类似的,printf(也是计算能力2.0+)如果输出太多,也需要手工增加。

When the call stack overflows, the kernel call fails with a stack overflow error if the application is run via a CUDA debugger (cuda-gdb, Nsight) or an unspecified launch error, otherwise.当Call stack Overflows, 如果在调试器下运行的话,能捕获到精确错误(调用栈溢出),如果没有调试器的话,则返回未知启动错误。没有调试器就是直接启动程序,有调试器就是在调试器下面启动程序(例如cuda-gdb your_program或者nsight->start cuda debugging)。

有不明白的地方,请在本文后留言

或者在我们的技术论坛bbs.gpuworld.cn上发帖

原文发布于微信公众号 - 吉浦迅科技(gpusolution)

原文发表时间:2018-05-17

本文参与腾讯云自媒体分享计划,欢迎正在阅读的你也加入,一起分享。

发表于

我来说两句

0 条评论
登录 后参与评论

相关文章

来自专栏大数据杂谈

Python 爬虫实战:股票数据定向爬虫

2274
来自专栏晨星先生的自留地

全方位绕过安全狗

2736
来自专栏程序猿

Git面试常见问题

有次推送了Git的基本配置。 Git作为团队开发的利器,在面试的时候,被问到的概率很大。 基础部分--大家平时都用什么工具? 出了Git自带的命令行工具外,作为...

3335
来自专栏前端知识分享

第130天:移动端-rem布局

当拿到设计师给的UI设计图,前端的首要任务就是布局和样式,相信这对于大部分前端工程师来说已经不是什么难题了。移动端的布局相对PC较为简单,关键在于对不同设备的适...

784
来自专栏北京马哥教育

Python 爬虫实战:股票数据定向爬虫

功能简介 目标: 获取上交所和深交所所有股票的名称和交易信息。 输出: 保存到文件中。 技术路线: requests—bs4–re 语言:python3.5 ...

36411
来自专栏我的小碗汤

【插件】解放鼠标,让浏览器更智能

首先感谢大家参与这周的送书活动。已经给小助手微信发送好友请求的,不要着急,这两天会通过验证,并拉大家进入抽奖群参与抽奖,大家肯定都可以参与的。

722
来自专栏社区的朋友们

基于 python 、js 的一个网页模块开发流程总结

这篇文章主要介绍了在功能模块中的一些关键处理思路和流程,以及一些比较典型的问题,都是比较基础的东西。其中的内容,相信各位大牛还有许多更好的处理方式。水平有限,总...

8560
来自专栏Seebug漏洞平台

TP-LINK 远程代码执行漏洞 CVE-2017-13772 趣谈

原文地址:https://www.fidusinfosec.com/tp-link-remote-code-execution-cve-2017-13772/ ...

3286
来自专栏web前端教室

Vue2.0,lifeCycle ['laɪfˌsaɪkl] -- 生命周期大白话~

生命周期,这词太屌了,头一次在前端相关文章中看到这个词的时候,我真是被唬住了。心里想,这前端还跟生命周期搞一块了,是不是还带转生投胎啊,跪着看了一半,我就站起来...

2128
来自专栏逍遥剑客的游戏开发

游戏配置序列化

2664

扫码关注云+社区