使用CUDA运行时API检查错误的规范方法是什么？
通过对CUDA问题和CUDA标记wiki的回答和评论，我发现通常会build议每个API调用的返回状态都应该检查错误。 API文档包含像cudaGetLastError ， cudaPeekAtLastError和cudaGetErrorString这样的函数，但是将这些函数放在一起可以可靠地捕获和报告错误而不需要大量额外代码的最佳方法是什么？
可能检查运行时API代码中错误的最好方法是定义一个断言样式处理函数和包装macros，如下所示：
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); } inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true) { if (code != cudaSuccess) { fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); if (abort) exit(code); } }
然后，您可以用gpuErrchkmacros来包装每个API调用，它将处理它所包装的API调用的返回状态，例如：
gpuErrchk( cudaMalloc(&a_d, size*sizeof(int)) );
如果在调用中发生错误，描述错误的文本消息以及发生错误的代码中的文件和行将被发送到stderr ，应用程序将退出。 你可以想象修改gpuAssert来引发exception，而不是在更复杂的应用程序中调用exit()如果需要的话）。
第二个相关的问题是如何检查内核启动中的错误，这些错误不能像标准运行时API调用那样直接包装在macros调用中。 对于内核来说，就像这样：
kernel<<<1,1>>>(a); gpuErrchk( cudaPeekAtLastError() ); gpuErrchk( cudaDeviceSynchronize() );
将首先检查无效的启动参数，然后强制主机等待，直到内核停止并检查执行错误。 如果您有如下的阻塞API调用，则可以消除同步：
kernel<<<1,1>>>(a_d); gpuErrchk( cudaPeekAtLastError() ); gpuErrchk( cudaMemcpy(a_h, a_d, size * sizeof(int), cudaMemcpyDeviceToHost) );
在这种情况下， cudaMemcpy调用可以返回内核执行期间发生的错误或者内存拷贝自身发生的错误。 这可能会让初学者感到困惑，我build议在debugging过程中在内核启动后使用显式同步，以便更容易理解可能出现问题的位置。
上面的talonmies的答案是以assert风格的方式中止应用程序的好方法。
偶尔，我们可能希望在C ++上下文中作为更大的应用程序的一部分报告并从错误状态中恢复。
这里有一个相当简洁的方法，通过使用thrust::system_error抛出一个从std::runtime_error派生的C ++exception：
#include <thrust/system_error.h> #include <thrust/system/cuda/error.h> #include <sstream> void throw_on_cuda_error(cudaError_t code, const char *file, int line) { if(code != cudaSuccess) { std::stringstream ss; ss << file << "(" << line << ")"; std::string file_and_line; ss >> file_and_line; throw thrust::system_error(code, thrust::cuda_category(), file_and_line); } }
这会将文件名，行号和cudaError_t的英文描述合并到抛出的exception的.what()成员中：
#include <iostream> int main() { try { // do something crazy throw_on_cuda_error(cudaSetDevice(-1), __FILE__, __LINE__); } catch(thrust::system_error &e) { std::cerr << "CUDA error after cudaSetDevice: " << e.what() << std::endl; // oops, recover cudaSetDevice(0); } return 0; }
输出：
$ nvcc exception.cu -run CUDA error after cudaSetDevice: exception.cu(23): invalid device ordinal
some_function的客户端可以根据需要将CUDA错误与其他types的错误区分开来：
try { // call some_function which may throw something some_function(); } catch(thrust::system_error &e) { std::cerr << "CUDA error during some_function: " << e.what() << std::endl; } catch(std::bad_alloc &e) { std::cerr << "Bad memory allocation during some_function: " << e.what() << std::endl; } catch(std::runtime_error &e) { std::cerr << "Runtime error during some_function: " << e.what() << std::endl; } catch(...) { std::cerr << "Some other kind of error during some_function" << std::endl; // no idea what to do, so just rethrow the exception throw; }
由于thrust::system_error是一个std::runtime_error ，如果我们不需要前面例子的精确度，我们可以用一个广泛的类错误来处理它：
try { // call some_function which may throw something some_function(); } catch(std::runtime_error &e) { std::cerr << "Runtime error during some_function: " << e.what() << std::endl; }
C ++ – 规范的方式：不要检查错误！
…只要使用现代的C ++绑定，在错误时抛出exception
我曾经被这个问题所厌恶， 我曾经有过一个macros观兼容的函数解决scheme，就像Talonmies和Jared的答案一样，但是，老实说呢？ 它使得使用CUDA运行时API更加丑陋和类C。
所以我以一种不同的，更根本的方式来解决这个问题。 对于结果示例，下面是CUDA vectorAdd示例的一部分 – 对每个运行时API调用进行完整的错误检查：
// (... prepare host-side buffers here ...) auto current_device = cuda::device::current::get(); auto d_A = cuda::memory::device::make_unique<float[]>(current_device, numElements); auto d_B = cuda::memory::device::make_unique<float[]>(current_device, numElements); auto d_C = cuda::memory::device::make_unique<float[]>(current_device, numElements); cuda::memory::copy(d_A.get(), h_A.get(), size); cuda::memory::copy(d_B.get(), h_B.get(), size); // (... prepare a launch configuration here... ) cuda::launch( vectorAdd, launch_config, d_A.get(), d_B.get(), d_C.get(), numElements ); cuda::memory::copy(h_C.get(), d_C.get(), size); // (...check results here...)
再次 – 检查所有潜在的错误是通过抛出的exception报告。 这段代码使用我的
用于CUDA运行时API库 （Github）的Thin Modern-C ++包装器
请注意，exception在调用失败后携带string解释和CUDA运行时API状态码
有几个CUDA错误如何自动检查这些包装的链接：
- 一个testing程序引发和捕捉一堆例外
- 有关错误相关function的文档
这里讨论的解决scheme对我来说效果很好。 该解决scheme使用内置的cudafunction，实现起来非常简单。
相关代码复制如下：
#include <stdio.h> #include <stdlib.h> __global__ void foo(int *ptr) { *ptr = 7; } int main(void) { foo<<<1,1>>>(0); // make the host block until the device is finished with foo cudaDeviceSynchronize(); // check for error cudaError_t error = cudaGetLastError(); if(error != cudaSuccess) { // print the CUDA error message and exit printf("CUDA error: %s\n", cudaGetErrorString(error)); exit(-1); } return 0; }
我通常这样做：
#define CHK_ERROR if (erro != cudaSuccess) goto Error; ... erro = cudaMalloc((void**)&d_image,sizeof(unsigned char)*nBlocks); CHK_ERROR ... erro = cudaDeviceSynchronize(); CHK_ERROR ... Error: std::cerr << "Error on CUDA: " << cudaGetErrorString(erro); cudaFree(d_image);
我认为这更可读
