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