使用CUDA运行时API检查错误的规范方法是什么?

通过对CUDA问题的回答和评论，在CUDA标签维基中，我看到它经常建议每个API调用的返回状态都应该检查错误。API文档包含cudaGetLastError、cudaPeekAtLastError和cudaGetErrorString这样的函数，但是将这些函数组合在一起以可靠地捕获和报告错误而不需要大量额外代码的最佳方法是什么呢?

138057 次浏览
小开
最佳答案

在运行时API代码中检查错误的最好方法可能是定义一个assert样式的处理函数和包装宏，如下所示:

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

然后你可以用gpuErrchk宏包装每个API调用，它将处理被包装的API调用的返回状态，例如:

gpuErrchk( cudaMalloc(&a_d, size*sizeof(int)) );

如果在调用中出现错误，则描述错误的文本消息以及发生错误的代码中的文件和行将被发送到stderr，并且应用程序将退出。如果需要，你可以修改gpuAssert来引发异常，而不是在更复杂的应用程序中调用exit()。

第二个相关问题是如何检查内核启动中的错误，这些错误不能像标准运行时API调用那样直接包装在宏调用中。对于内核，是这样的:

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调用可以返回内核执行期间发生的错误或来自内存复制本身的错误。这可能会让初学者感到困惑，我建议在调试内核启动后使用显式同步，以便更容易理解哪里可能出现问题。

注意，当使用CUDA动态并行时，非常相似的方法可以并且应该应用于设备内核中CUDA运行时API的任何用法，以及在任何设备内核启动后:

#include <assert.h>
#define cdpErrchk(ans) { cdpAssert((ans), __FILE__, __LINE__); }
__device__ void cdpAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
if (code != cudaSuccess)
{
printf("GPU kernel assert: %s %s %d\n", cudaGetErrorString(code), file, line);
if (abort) assert(0);
}
}

CUDA Fortran错误检查是类似的。参见在这里和在这里了解典型的函数错误返回语法。使用类似CUDA c++的方法来收集与内核启动相关的错误。

小开

上面talonmies的回答是一个以__abc0风格的方式中止应用程序的好方法。

有时，我们可能希望在c++上下文中报告错误并从错误中恢复，这是大型应用程序的一部分。

下面是一个相当简洁的方法，通过使用thrust::system_error抛出一个派生自std::runtime_error的c++异常:

#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的英文语言描述合并到抛出异常的.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错误和其他类型的错误:

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++绑定。

我曾经为这个问题感到烦恼;我曾经有一个宏加包装函数的解决方案，就像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);


auto launch_config = cuda::launch_config_builder()
.overall_size(numElements)
.block_size(256)
.build();


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


// (... verify results here...)

同样，所有潜在的错误都会被检查，如果发生错误则会出现异常(警告:如果内核导致一些错误后启动，它将在尝试复制结果之后被捕获，而不是在之前;为了确保内核运行成功，您需要同步设备或默认流)。

上面的代码使用my

用于CUDA运行时API库的Thin modern - c++包装器 (Github)

请注意，在调用失败后，异常同时带有字符串解释和CUDA运行时API状态代码。

下面是一些关于如何使用这些包装器自动检查CUDA错误的链接:

小开

讨论的解决方案在这里很适合我。该解决方案使用内置cuda函数，实现非常简单。

相关代码复制如下:

#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;
}

提交答案