cuda中的错误捕获

cuda提供的API分为：
● 同步调用API，即当host端调用该API函数后，会一直hold在调用的位置，直到该函数返回为止，一般同步调用API会返回一个cudaError_t 类型的结果
● 异步调用API，即当host端调用该API函数后，会直接返回，继续执行后面的代码，异步调用API则会返回void
在CUDA当中，自定义的核函数kernel的执行总是异步的

因此在进行错误处理的时候有两个场景：
● 对同步调用的错误处理，因为它会返回cudaError_t 类型的结果，因此可以直接检查返回结果
● 对异步调用的错误处理，因为它不会返回错误信息，因此需要调用错误获取接口得到错误信息

cuda 中用于获取函数执行的结果有两个接口：
● cudaPeekAtLastError()
● cudaGetLastError()
CUDA 运行时为每个主机线程维护一个错误变量，它初始化为 cudaSuccess，每当发生错误时（无论是参数验证错误还是异步错误），都会被错误码覆盖。cudaPeekAtLastError() 返回这个变量。而 cudaGetLastError() 返回这个变量，并将其重置为 cudaSuccess。

因此总结下来：
● 对于异步调用API（包括自定义的kernel函数），由于host线程不知道它什么时候结束，因此要获取它的执行结果，需要先进行同步处理，通常是调用cudaDeviceSynchronize()函数，然后再用cudaGetLastError()或者cudaPeekAtLastError()进行捕获
● 对于同步调用API，如果它会返回cudaError_t 类型的结果，则直接对返回结果进行检查，如果它不返回错误信息，则和异步调用API一样，使用错误获取函数来得到。

示例代码：

#include <cuda_runtime.h>
#include <system_error>

#define CUDA_CHECK(call)             __cudaCheck(call, __FILE__, __LINE__)
#define LAST_KERNEL_CHECK()          __kernelCheck(__FILE__, __LINE__)
#define BLOCKSIZE 16

inline static void __cudaCheck(cudaError_t err, const char* file, const int line) {
    if (err != cudaSuccess) {
        printf("ERROR: %s:%d, ", file, line);
        printf("CODE:%s, DETAIL:%s\n", cudaGetErrorName(err), cudaGetErrorString(err));
        exit(1);
    }
}

inline static void __kernelCheck(const char* file, const int line) {
    cudaError_t err = cudaPeekAtLastError();
    if (err != cudaSuccess) {
        printf("ERROR: %s:%d, ", file, line);
        printf("CODE:%s, DETAIL:%s\n", cudaGetErrorName(err), cudaGetErrorString(err));
        exit(1);
    }
}

// 使用示例

// 同步API
CUDA_CHECK(cudaMalloc(&M_device, size));

// 异步API或自定义kernel
Kernel <<<dimGrid, dimBlock>>> (M_device, N_device, P_device, width);
CUDA_CHECK(cudaDeviceSynchronize());
LAST_KERNEL_CHECK();