283

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

4

4 回答 4

346

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

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

然后，您可以使用宏包装每个 API 调用，该gpuErrchk宏将处理它包装的 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 Dynamic Parallelism时，可以并且应该将非常相似的方法应用于设备内核中 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);
   }
}
于 2012-12-26T09:35:49.810 回答
73

上面 talonmies 的回答是以assert-style 方式中止应用程序的好方法。

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

这是一个相当简洁的方法，通过抛出从std::runtime_errorusing派生的 C++ 异常thrust::system_error：

#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_erroris a 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;
}
于 2012-12-27T07:31:21.220 回答
30

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

我曾经对这个问题感到厌烦；我曾经有一个宏兼包装功能解决方案，就像在 Talonmies 和 Jared 的答案中一样，但是，老实说？它使使用 CUDA 运行时 API 变得更加丑陋和类似 C。

所以我以一种不同的、更基本的方式来解决这个问题。对于结果示例，以下是 CUDAvectorAdd示例的一部分 - 对每个运行时 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);

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

再次 - 检查所有潜在错误，如果发生错误则异常（警告：如果内核在启动后导致某些错误，它将在尝试复制结果之后被捕获，而不是之前；为确保内核成功，您将需要使用cuda::outstanding_error::ensure_none()命令检查启动和复制之间的错误）。

上面的代码使用我的

CUDA 运行时 API 库(Github)

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

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

于 2013-12-09T18:58:16.583 回答
13

这里讨论的解决方案对我来说效果很好。该解决方案使用内置的 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;
}
于 2015-12-19T15:35:20.793 回答