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 ) 返回上一次cudaruntime调用时的错误码，并将cuda错误码设置为cudaSuccess
• __host__ __device__ cudaError_t cudaPeekAtLastError( void ) 与上面功能一样，但是不会重置cuda的错误信息。

核函数被调用或者cuda API被调用时，都会返回一个错误码，如果返回的错误码为0，表示该调用没有出错，即cudaSuccess（错误名称），错误码一共有82个。

对于类似cudaMemcpy的API调用函数，可以使用cudaGetErrorString来进行错误检查； 对于核函数，其调用之后，必然会更新cuda Runtime的错误值，可以在核函数调用之后调用cudaGetLastError 进行核函数调用的错误检查。

由于CPU与GPU程序是异步执行的，所以为了避免获取到的核函数调用时的错误信息有误，应该在调用cudaGetLastError 函数之前使用cudaDeviceSynchronize() 对CPU进行阻塞，以同步GPU和CPU的线程，保证核函数已经调用完成。

以下是可以参考的错误处理核函数：

//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);
   }
}

example:

//统计用于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));

参考：cuda编程中的错误处理