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上发帖