CUDA捕捉调用核函数是否发生错误的方法（之一）

前言

在CUDA中，当你调用一个核函数或者使用CUDA 的API，错误可能会发生。调用CUDA 的API出错会返回cudaError_t 变量，检查错误比较方便，这里不做展开，本文主要举例一种捕捉调用核函数时，是否发生错误的方法，为了捕捉和处理这些错误，本文提供一种较为简单的检测方法：

举例（1）

核函数调用后检查错误:
核函数调用是异步的，意味着CPU代码会继续执行，而不会等待GPU上的核函数完成。要检查核函数是否出错，可以在调用之后立即使用cudaDeviceSynchronize()来等待GPU完成，然后检查是否有错误。具体代码如下：

kernelFunction<<<blocks, threads>>>(args...);//自定义核函数

cudaDeviceSynchronize();

//返回上一个在当前线程上发生的CUDA运行时API调用的错误
cudaError_t err = cudaGetLastError();

if (err != cudaSuccess) {
    printf("Error: %s\n", cudaGetErrorString(err));
}

这里需要注意的一点是，是否要加上cudaDeviceSynchronize();，

如果不加上同步的语句，当你在核函数调用之后立即使用cudaGetLastError() 实际上是检查核函数的启动是否有问题, 而不是核函数执行过程中是否有问题。也就是说，核函数执行过程中的出错很难被检测到，但能第一时间检查并打印启东时的错误信息。这个也很好理解，因为CUDA核函数的调用是异步的。

如果你在核函数调用后使用cudaDeviceSynchronize() 再使用cudaGetLastError() ，你就可以检测到与这个核函数相关的所有错误，无论是启动时的还是执行过程中的，只不过错误打印要等核函数执行完，会慢一点。

在有上面的认知前提，我们可以对这个流程再优化一下，既能捕捉到核函数启动时的错误并打印，同时，能捕捉核函数运行时的错误并打印，更规范的代码如下：

myKernelFunction<<<grid, block>>>(...);
ErrorCheck(cudaGetLastError(), __FILE__, __LINE__);
ErrorCheck(cudaDeviceSynchronize(), __FILE__, __LINE__);

ErrorCheck()是自定义的错误检查函数，这里不赘述，有很多示例。
ErrorCheck(cudaDeviceSynchronize(), __FILE__, __LINE__);会确保当前设备完成所有之前的CUDA操作，包括核函数的执行。如果核函数在执行过程中出现了错误，cudaDeviceSynchronize()会返回一个错误码。当核函数执行完成并且没有错误时，这个函数返回cudaSuccess。

但注意，这里的写法还是要确保完成同步之后，即所有GPU任务都已完成，才能打印核函数运行中的错误信息，因为这样才可以确信在此之后发生的所有操作（例如错误检查）都是基于完整的GPU执行结果的。

因此，这两个函数的组合确保了捕捉到与核函数相关的所有错误。