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

相关文章

来自专栏点滴积累

geotrellis使用(三十二)大量GeoTiff文件实时发布TMS服务

前言 在上一篇文章中我讲了如何直接将Geotiff文件发布为TMS服务,在其中只讲了单幅Geotiff的操作,其实单幅这种量级的数据对Geotrellis来说就...

3167
来自专栏zingpLiu

【实战小项目】python开发自动化运维工具--批量操作主机

有很多开源自动化运维工具都很好用如ansible/salt stack等,完全不用重复造轮子。只不过,很多运维同学学习Python之后,苦于没小项目训练,本篇演...

2522
来自专栏Petrichor的专栏

git:git commit 书写格式

  正如 git add 的作用是将文件放入暂存区, git commit 的作用是将修改提交到分支上。

4062
来自专栏信安之路

一道 CTF 题 get 到的新姿势

本文是从一个 CTF 题目中学到的一个新姿势,下面对我的学习做一个记录总结,给大家分享一下,希望大家多多参与一起分享学习。

1070
来自专栏Kirito的技术分享

深入理解 RPC 之集群篇

上一篇文章分析了服务的注册与发现,这一篇文章着重分析下 RPC 框架都会用到的集群的相关知识。 集群(Cluster)本身并不具备太多知识点,在分布式系统中,...

3729
来自专栏Android 研究

Android跨进程通信IPC之1——Linux基础

由于Android系统是基于Linux系统的,所以有必要简单的介绍下Linux的跨进程通信,对大家后续了解Android的跨进程通信是有帮助的,本篇的主要内容如...

1993
来自专栏搜云库

在 Linux 上搭建Jekyll静态博客

在CentOS,Ubuntu 按照同样步骤安装,Ruby Gems 往往都无法搭建成,每次都是依赖不对,各种奇葩原因,解决办法就是使用 RVM 安装,解决 Ru...

3958
来自专栏小灰灰

QuickTask动态脚本支持框架整体介绍篇

一个简单的动态脚本调度框架,支持运行时,实时增加,删除和修改动态脚本,可用于后端的进行接口验证、数据订正,执行定时任务或校验脚本

1212
来自专栏有趣的Python

8- vue django restful framework 打造生鲜超市 -商品类别数据展示(下)

Vue+Django REST framework实战 搭建一个前后端分离的生鲜超市网站 Django rtf 完成 商品类别页 vue展示商品列表页数...

3614
来自专栏coding

oh-my-zsh,让你的终端从未这么爽过

7.4K5

扫码关注云+社区

领取腾讯云代金券