在写CUDA程序时，经常会遇到各种各样的错误，这些错误大致可分为两类：

• 编译阶段错误：例如语法错误等
• 运行阶段错误：例如函数返回非预期结果等

编译阶段的错误编译器会检查出来，本文主要讨论CUDA运行阶段函数的错误检查，主要分为带返回码函数的错误检查和不带返回码函数的错误检查，相关内容可参考官方API文档的Error Handling章节。

CUDA运行时API可参考官方文档：https://docs.nvidia.com/cuda/pdf/CUDA_Runtime_API.pdf

1、 带返回码函数的错误检查

CUDA运行时API的很多函数都会返回一个cudaError_t类型的错误码，这和常规的c/c++风格类似，可以通过该返回码来判断操作是否成功。cudaError_t是一个枚举类型的数据，常见的枚举值如下，更多枚举值可查看官方API文档enum cudaError章节。

• cudaSuccess = 0
• cudaErrorInvalidValue = 1
• cudaErrorMemoryAllocation = 2
• cudaErrorInitializationError = 3
• cudaErrorCudartUnloading = 4
• …

对于这类带返回码的CUDA函数，可以判断返回码是否等于cudaSuccess来判断操作是否成功。但是当返回码不等于cudaSuccess时，记录错误日志的时候只记录一个枚举数值的返回码对于分析问题是不够直观的，拿到错误码还需要比对官方文档才能知道具体是什么错误。

为了在日志中提供更直观的日志信息，可以借助CUDA提供的两个函数：

• cudaGetErrorName

函数定义：__host____device__const char *cudaGetErrorName(cudaError_t error)

函数功能：传入一个cudaError_t类型的错误码，返回该错误码的错误名称，如果传入的错误码不在CUDA错误码枚举列表里，则会返回“unrecognized error code”。

• cudaGetErrorString ：

函数定义：__host____device__const char *cudaGetErrorString(cudaError_t error)

函数功能：传入一个cudaError_t类型的错误码，返回该错误码的描述信息，如果传入的错误码不在CUDA错误码枚举列表里，则会返回“unrecognized error code”。

因此可以封装一个错误检查函数，当错误码为cudaSuccess时直接返回，否则记录错误码信息：

// 可以通过__LINE__获取当前代码行数
// 考虑到项目多文件，还可以加上__FILE__获取文件名
cudaError_t error_check(cudaError_t error_code, int line)
{
        if (error_code != cudaSuccess)
        {
                printf("line: %d, error_code: %d, error_name: %s, error_description: %s\n",
                                line, error_code, cudaGetErrorName(error_code), cudaGetErrorString(error_code));
        }
        return error_code;
}

结合一个实际例子error.cu（为了方便说明，加入了行号）：

  1 #include<stdio.h>
  2
  3 cudaError_t error_check(cudaError_t error_code, int line)
  4 {
  5         if (error_code != cudaSuccess)
  6         {
  7                 printf("line: %d, error_code: %d, error_name: %s, error_description: %s\n",
  8                                 line, error_code, cudaGetErrorName(error_code), cudaGetErrorString(error_code));
  9         }
 10         return error_code;
 11 }
 12
 13 int main()
 14 {
 15         // host上申请内存空间
 16         float *p_host = (float *)malloc(4);
 17         memset(p_host, 0, 4);
 18
 19         // device上申请相同大小空间
 20         float *p_device;
 21         error_check(cudaMalloc((float**)&p_device, 4), __LINE__);
 22         error_check(cudaMemset(p_device, 0, 4), __LINE__);
 23
 24         // host数据拷贝到device
 25         // 【注意这里是想把p_host数据拷贝到p_device，但是类型用的是cudaMemcpyDeviceToHost，即设备到主机，明显是错误的】
 26         error_check(cudaMemcpy(p_device, p_host, 4, cudaMemcpyDeviceToHost), __LINE__);
 27
 28         free(p_host); // 释放host内存
 29         error_check(cudaFree(p_device), __LINE__); // 释放device内存
 30
 31         return 0;
 32 }

注意第26行故意写了错误代码，看看编译运行结果：

$ nvcc error.cu -o error
$ ./error
line: 26, error_code: 1, error_name: cudaErrorInvalidValue, error_description: invalid argument

2、 不带返回码函数的错误检查

除了上述带cudaError_t返回码的函数，CUDA也有一些不带返回码的函数，例如核函数，必须返回void，这类函数该怎么左错误检查呢？

如果熟悉linux的读者，应该对$?会有印象，echo $?就是表示查看上一次执行命令的错误码，CUDA也有一个类似功能的函数cudaGetLastError来获取上次函数调用的“错误码”，我们可以用这个函数来检查不带返回码的函数。先来看看cudaGetLastError的定义和功能：

• cudaGetLastError

函数定义：__host____device__cudaError_t cudaGetLastError(void)

函数功能：官方文档中写的是Returns the last error from a runtime call，翻译过来就是返回上一次运行时调用的错误。

结合error_check函数和前面的自定义核函数，来看看cudaGetLastError的示例last_error.cu：

  1 #include<stdio.h>
  2
  3 __global__ void hello()
  4 {
  5         printf("hello from GPU");
  6 }
  7
  8 cudaError_t error_check(cudaError_t error_code, int line)
  9 {
 10         if (error_code != cudaSuccess)
 11         {
 12                 printf("line: %d, error_code: %d, error_name: %s, error_description: %s\n",
 13                                 line, error_code, cudaGetErrorName(error_code), cudaGetErrorString(error_code));
 14         }
 15         return error_code;
 16 }
 17
 18 int main()
 19 {
 20         // 查看线程块的最大线程数
 21         cudaDeviceProp prop;
 22         cudaGetDeviceProperties(&prop, 0);
 23         int maxThreadsPerBlock = prop.maxThreadsDim[0];
 24         printf("maxThreadsPerBlock: %d\n", maxThreadsPerBlock);
 25
 26         // 注意这里用的是一维线程模型，block_size配置的是2048，即一个线程块里有2048个线程
 27         // 但是我电脑GPU当前最多允许一个线程块1024个线程，所以这里会报错
 28         hello<<<1, 2048>>>();
 29         error_check(cudaGetLastError(), __LINE__);
 30
 31         return 0;
 32 }

编译运行：

$ nvcc last_error.cu -o last_error
$ ./last_error
maxThreadsPerBlock: 1024
line: 29, error_code: 9, error_name: cudaErrorInvalidConfiguration, error_description: invalid configuration argument

微信公众号卡巴斯同步发布，欢迎大家关注。