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

Question

阅读有关CUDA的问题的答案和评论，并在 CUDA标签wiki 中，我看到常常建议检查每个API调用的返回状态是否存在错误。 API文档包含诸如cudaGetLastError、cudaPeekAtLastError和cudaGetErrorString等函数，但最可靠的方式是如何将它们组合起来以可靠地捕获和报告错误，而不需要大量的额外代码？

- talonmies

---

NVIDIA的CUDA示例中包含一个名为helper_cuda.h的头文件，其中有称为“getLastCudaError”和“checkCudaErrors”的宏，其功能与接受的答案描述的基本相同。请参阅示例以进行演示。只需选择与工具包一起安装示例，您就可以使用它了。 - chappjc

@chappjc，我不认为这个问题和回答是原创的，如果你的意思是这样的话。但它有一个好处，就是让人们使用CUDA错误检查变得更加有教育意义。 - Vitality

@JackOLantern 不，那不是我想表达的意思。这个问答对我非常有帮助，而且肯定比在SDK中找到某些头文件要容易得多。我认为指出这也是NVIDIA处理方式的价值所在，并告诉大家可以在哪里寻找更多信息。如果可以的话，我会缓和我的评论语气。 :) - chappjc

自2012年起，允许您“逼近”错误开始的调试工具在CUDA上有了很大改进。我没有使用基于GUI的调试器，但CUDA标签wiki提到了命令行cuda-gdb。这是非常强大的工具，因为它允许您在GPU本身上逐步执行实际的线程和warp（大多数时候需要2.0+架构）。 - opetrenko

@bluefeet: 你为什么要撤回那个编辑呢？看起来在 Markdown 中什么都没有改变，但它被接受为一个编辑。是否有什么不良的企图？ - talonmies

@talonmies 我不会评论其恶劣性质，但这里有更多细节。 - Taryn

Answer 1

检查运行时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调用，可以消除同步问题，例如：

将首先检查无效的启动参数，然后强制主机等待直到内核停止并检查执行错误。如果您有后续的阻塞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

---

在退出之前，我们不应该添加cudaDeviceReset()吗？还需要加上内存释放的语句吗？ - Aurelius

对于异步CUDA运行时调用，例如cudaMemsetAsync和cudaMemcpyAsync，是否还需要通过调用gpuErrchk(cudaDeviceSynchronize())来同步GPU设备和主机线程？ - nurabha

请注意，内核启动后进行显式同步操作并不是错误的，但可能严重影响执行性能和交错语义。如果您正在使用交错技术，在调试时进行显式同步操作可能会隐藏一整类难以在发布版本中跟踪的错误。 - masterxilo

有没有办法获取内核执行的更具体错误？我得到的所有错误都只给出主机代码的行号，而不是内核的行号。 - Azmisov

请注意，与所有其他CUDA错误不同，内核_launch_错误将不会在后续的CUDA运行时API同步调用中报告。因此，仅在下一个cudaMemcpy()或cudaDeviceSynchronize()调用周围放置gpuErrchk()是不足以捕获所有可能的错误条件的。我认为，在内核启动后立即调用cudaGetLastError()而不是cudaPeekAtLastError()更好，即使它们具有相同的效果，也可以帮助不知情的读者。 - tera

Answer 2

talonmies的上面的回答是一种通过assert方式中止应用程序的好方法。

偶尔我们可能希望在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;
}

- Jared Hoberock

---

推进头似乎被重新排列了。 <thrust/system/cuda_error.h>现在实际上是<thrust/system/cuda/error.h>。 - chappjc

Jared，我认为我的包装库已经包含了你提出的解决方案 - 大部分，并且足够轻量级，可能可以替代它。（请参见我的回答） - einpoklum

Answer 3

C++的规范方式：不要检查错误；使用C++绑定，它会抛出异常。

我曾经为这个问题感到恼火；我曾经有一个类似于Talonmies和Jared答案中的宏包装函数解决方案，但是，说实话？它使得使用CUDA Runtime API更加丑陋和类C。

因此，我采用了一种不同而更基本的方法。以下是CUDA vectorAdd示例的部分结果 - 对每个运行时API调用进行了完整的错误检查：

// (... prepare host-side buffers here ...)

auto current_device = cuda::device::current::get();
namespace cm = cuda::memory;
auto d_A = cm::device::make_unique<float[]>(current_device, numElements);
auto d_B = cm::device::make_unique<float[]>(current_device, numElements);
auto d_C = cm::device::make_unique<float[]>(current_device, numElements);

cm::copy(d_A.get(), h_A.get(), size);
cm::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);    
cm::copy(h_C.get(), d_C.get(), size);

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

再次说明-所有潜在的错误都经过检查，如果出现错误则会抛出异常（注意：如果内核在启动后引起了一些错误，则将在尝试复制结果之后捕获，而不是之前；如果要确保内核成功，您需要同步设备或默认流）。

上面的代码使用了我的

CUDA Runtime API库的轻量级现代C++包装器 (Github)

请注意，这些异常除了包含字符串解释之外，还包含失败调用后的CUDA运行时API状态码。

以下是如何使用这些包装器自动检查CUDA错误的一些链接：

• 一个测试程序，抛出和捕获一堆异常
• 与错误相关的功能文档

- einpoklum

---

看起来可以通过使用C++的作用域命名空间来改进答案。 using namespace cuda::memory。 每行超过三个::真的让我远离C++。 - Dimitri Lesnoff

@DimitriLesnoff：现在怎么样？ - einpoklum

这看起来好多了！谢谢！ 我们需要为每个数组使用.get()吗？为什么？ 我们能否获取这些数组对应的设备并在执行的其余部分中使用相同的数组？ 我还将此make_unique函数别名为cuda_make_unique_float函数，以抽象掉所有这些作用域解析和模板实例化。 make_unique将作用于数组，那么为什么我们需要指定方括号[]？ 该函数的模板定义可能会得到改进。我希望语法不要太严格。 - Dimitri Lesnoff

@DimitriLesnoff：我相信我们现在已经进入了我的包装库设计领域——这超出了本页面的范围（它只关注错误处理）。请考虑在此处提交问题，或给我发送电子邮件，或开始聊天会话。 - einpoklum

Answer 4

这里讨论的解决方案对我很有效。该解决方案使用内置的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;
}

- jthomas

Answer 5

让我添加一个我最喜欢的CUDA错误检查宏。

#define CUDACHECK(err) do { cuda_check((err), __FILE__, __LINE__); } while(false)
inline void cuda_check(cudaError_t error_code, const char *file, int line)
{
    if (error_code != cudaSuccess)
    {
        fprintf(stderr, "CUDA Error %d: %s. In file '%s' on line %d\n", error_code, cudaGetErrorString(error_code), file, line);
        fflush(stderr);
        exit(error_code);
    }
}

将每个CUDA函数调用都包装在CUDACHECK()宏中，在每次内核启动后，使用CUDACHECK(cudaPeekAtLastError())。

do{...}while(false)的存在主要是为了在宏后面强制加上分号。请参阅this article以获取有关为什么使用它的详细说明。

- Jakub Homola

---

虽然本质上与已接受的答案没有区别，但这个答案至少使用do { } while正确地定义了宏，并加以解释。 - Mark Gates