blueyi's notes

Follow Excellence,Success will chase you!

0%

CUDA编程中的错误处理

Posted on In
Symbols count in article: 3.2k Reading time ≈ 3 mins.

CUDA程序往往很难高度，特别是核函数出错时，甚至有些情况下核函数中内容访问错误时，运行程序之后系统返回的报错信息是free() error之类的错误，会严重误导错误查找过程。所以编写CUDA程序一定记得进行错误错误处理，多写几行代码，有可能省出很多的时间。假如自己的CUDA程序出现了各种奇葩、难以理解的错误，可以考虑加上错误处理，也许就柳岸花明了。
在CUDA学习笔记最后的总结中刻意提到编写CUDA程序一定要做错误处理，但其中一直未提及如何进行错误处理，这里给补上。

CUDA编程中的错误处理实际上非常简单，错误处理相关的函数一共也只有4个：

  • __host__ __device__ const char* cudaGetErrorName( cudaError_t error ) 接受一个错误码，返回错误名称
  • __host__ __device__ const char* cudaGetErrorString( cudaError_t error ) 接受一个错误码，返回错误描述
  • __host__ __device__ cudaError_t cudaGetLastError( void ) 返回上一次CUDA runtime调用时的错误码，并将CUDA错误码置为cudaSuccess
  • __host__ __device__ cudaError_t cudaPeekAtLastError( void ) 与上面功能一样，但不会重置CUDA的错误信息
    当核函数被调用或者CUDA API被调用时，通常都会返回一个错误码，如果返回的错误码为0，则表示该调用没有出错，即cudaSuccess（即错误名称），CUDA的错误码一共有82个（截止到2017.3.29），所有的错误码信息可以在这里查看cudaError，其实就是个enum类型的数据。

根据官方提代的4个错误处理函数，明显可以看出对于cudaMemcpy这类API调用函数，可以直接传递给cudaGetErrorString来进行错误检查，而核函数调用之后必然会更新CUDA Runtime的错误值，那么就可以在核函数调用之后调用cudaGetLastError()来进行核函数调用的错误检查。由于CPU与GPU程序时异步执行，所以为了避免获取到的核函数调用错误信息有误，应该在调用cudaGetLastError()这类函数之前使用cudaDeviceSynchronize()对CPU进行阻塞，以同步GPU与CPU的线程，保证核函数已经调用完成。
为了便于错误处理，通常会将错误处理写成函数，或者进行宏定义，不然每次获取到错误码之后，都还要去调用cudaGetErrorString来查看错误描述。
下面是我常用的一个错误处理宏定义，兼容C与C++，可以参考：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
//API调用错误处理，可以接受CUDA的API函数调用作为参数
#define CHECK_ERROR(error) checkCudaError(error, __FILE__, __LINE__)
//检查CUDA Runtime状态码，可以接受一个指定的提示信息
#define CHECK_STATE(msg) checkCudaState(msg, __FILE__, __LINE__)

inline void checkCudaError(cudaError_t error, const char *file, const int line)
{
   if (error != cudaSuccess) {
      std::cerr << "CUDA CALL FAILED:" << file << "( " << line << ")- " << cudaGetErrorString(error) << std::endl;
      exit(EXIT_FAILURE);
   }
}

inline void checkCudaState(const char *msg, const char *file, const int line)
{
   cudaError_t error = cudaGetLastError();
   if (error != cudaSuccess) {
      std::cerr << "---" << msg << " Error---" << std::endl;
      std::cerr << file << "( " << line << ")- " << cudaGetErrorString(error) << std::endl;
      exit(EXIT_FAILURE);
   }
}

调用举例，下面是从以前一个项目代码中抽出来的代码片断：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
//统计用于GPU计算的时间
cudaEvent_t start, stop;
CHECK_ERROR(cudaEventCreate(&start));
CHECK_ERROR(cudaEventCreate(&stop));
CHECK_ERROR(cudaEventRecord(start, 0));
CHECK_ERROR(cudaEventSynchronize(start));

CHECK_ERROR(cudaMemcpy(dev_grid_in, grid_in, SIZE * sizeof(double), cudaMemcpyHostToDevice));
CHECK_ERROR(cudaMemcpy(dev_grid_out, grid_out, SIZE * sizeof(double), cudaMemcpyHostToDevice));

for (int i=0; i<times; ++i) {
   kernel<<<grids, blocks>>>(dev_grid_in, dev_grid_out);
   std::swap(dev_grid_in, dev_grid_out);
}
cudaDeviceSynchronize();

CHECK_STATE("kernel call");

CHECK_ERROR(cudaMemcpy(grid_in, dev_grid_in, SIZE * sizeof(double), cudaMemcpyDeviceToHost));

//计算统计的时间
CHECK_ERROR(cudaEventRecord(stop, 0));
CHECK_ERROR(cudaEventSynchronize(stop));
float elapsedTime;
CHECK_ERROR(cudaEventElapsedTime(&elapsedTime, start, stop));

上述代码片断的完整代码链接在https://github.com/blueyi/stencil_computation_sample/blob/master/7_point_stencil_jacobi3d_Optimize.cu

参考：
1.http://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__ERROR.html#group__CUDART__ERROR
2.How to Query Device Properties and Handle Errors in CUDA C/C++
3.How to do error checking in CUDA

Welcome to my other publishing channels