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 条评论
登录 后参与评论

相关文章

来自专栏java 成神之路

使用 NIO 实现 echo 服务器

4607
来自专栏一个会写诗的程序员的博客

Spring Reactor 项目核心库Reactor Core

Non-Blocking Reactive Streams Foundation for the JVM both implementing a Reactiv...

2142
来自专栏闻道于事

js登录滑动验证,不滑动无法登陆

js的判断这里是根据滑块的位置进行判断,应该是用一个flag判断 <%@ page language="java" contentType="text/html...

6768
来自专栏Golang语言社区

【Golang语言社区】GO1.9 map并发安全测试

var m sync.Map //全局 func maintest() { // 第一个 YongHuomap := make(map[st...

4708
来自专栏pangguoming

Spring Boot集成JasperReports生成PDF文档

由于工作需要,要实现后端根据模板动态填充数据生成PDF文档,通过技术选型,使用Ireport5.6来设计模板,结合JasperReports5.6工具库来调用渲...

1.2K7
来自专栏张善友的专栏

Miguel de Icaza 细说 Mix 07大会上的Silverlight和DLR

Mono之父Miguel de Icaza 详细报道微软Mix 07大会上的Silverlight和DLR ,上面还谈到了Mono and Silverligh...

2707
来自专栏C#

DotNet加密方式解析--非对称加密

    新年新气象,也希望新年可以挣大钱。不管今年年底会不会跟去年一样,满怀抱负却又壮志未酬。(不过没事,我已为各位卜上一卦,卦象显示各位都能挣钱...)...

4848
来自专栏Ceph对象存储方案

Luminous版本PG 分布调优

Luminous版本开始新增的balancer模块在PG分布优化方面效果非常明显,操作也非常简便,强烈推荐各位在集群上线之前进行这一操作,能够极大的提升整个集群...

3105
来自专栏陈仁松博客

ASP.NET Core 'Microsoft.Win32.Registry' 错误修复

今天在发布Asp.net Core应用到Azure的时候出现错误InvalidOperationException: Cannot find compilati...

4838
来自专栏魂祭心

原 canvas绘制clock

4054

扫码关注云+社区