Cuda编程模型中常见的错误检测方法

---

1 CUDA错误检测简介

CUDA编程模型中的错误检测是确保在GPU上运行的程序能够正确执行的关键步骤。CUDA（Compute Unified Device Architecture）提供了多种错误检测机制，以帮助开发者识别和处理在CUDA程序执行过程中可能出现的错误。这些错误检测机制包括：

• API错误检测：CUDA提供了一系列API函数，每个函数调用后都会返回一个错误码，开发者可以通过检查这个错误码来确定函数调用是否成功。例如，cudaError_t err = cudaMalloc(&d_ptr, size); 后可以使用 if (err != cudaSuccess) { /* 错误处理代码 */ } 来检测内存分配是否成功。

• 内核错误检测：在CUDA中，内核函数在GPU上并行执行，内核的执行状态可以通过 cudaGetLastError() 和 cudaPeekAtLastError() 函数来检查。这些函数返回最后一个内核执行的错误状态，帮助开发者捕捉内核执行过程中发生的错误。

• 同步和异步错误检测：CUDA操作分为同步和异步操作。同步操作的错误可以立即检测到，而异步操作的错误（如内核执行）可能会延迟。使用 cudaDeviceSynchronize() 可以强制同步，确保所有先前的CUDA调用完成，从而检测任何潜在的错误。

• 调试和分析工具：CUDA提供了多种调试和分析工具，如cuda-gdb、NVIDIA Nsight等，这些工具可以帮助开发者深入分析CUDA程序的执行情况，检测和修复错误。

通过结合这些错误检测机制，开发者可以有效地检测和处理CUDA编程中的各种错误，从而提高程序的稳定性和可靠性。

---

2 直接嵌入检测函数

2.1 检测函数介绍

在CUDA编程中，错误检测是确保代码正确性和性能优化的重要部分。以下是几个常用的CUDA错误检测函数的介绍：

__host__​__device__​const char* 	cudaGetErrorName ( cudaError_t error )
Returns the string representation of an error code enum name.  

__host__​__device__​const char* 	cudaGetErrorString ( cudaError_t error )
Returns the description string for an error code.  

__host__​__device__​cudaError_t 	cudaGetLastError ( void )
Returns the last error from a runtime call.  

__host__​__device__​cudaError_t 	cudaPeekAtLastError ( void )
Returns the last error from a runtime call.  

1. cudaGetErrorName(cudaError_t error)

• 函数原型:

  __host__ __device__ const char* cudaGetErrorName(cudaError_t error);
  

• 功能:
  这个函数返回给定错误码的名称字符串。它对调试和日志记录非常有用，使开发人员能够通过错误名称快速识别问题。

• 参数:
  error: 需要获取名称的CUDA错误码。

• 返回值:
  对应错误码的名称字符串。

• 示例:

  cudaError_t err = cudaMalloc(...);
  if (err != cudaSuccess) {
      printf("CUDA Error: %s\n", cudaGetErrorName(err));
  }
  

2. cudaGetErrorString(cudaError_t error)

• 函数原型:

  __host__ __device__ const char* cudaGetErrorString(cudaError_t error);
  

• 功能:
  这个函数返回给定错误码的描述字符串。与cudaGetErrorName类似，这个函数提供了更详细的错误信息，有助于理解错误的原因。

• 参数:
  error: 需要获取描述的CUDA错误码。

• 返回值:
  对应错误码的描述字符串。

• 示例:

  cudaError_t err = cudaMalloc(...);
  if (err != cudaSuccess) {
      printf("CUDA Error: %s\n", cudaGetErrorString(err));
  }
  

3. cudaGetLastError(void)

• 函数原型:

  __host__ __device__ cudaError_t cudaGetLastError(void);
  

• 功能:
  这个函数返回并清除最近一次执行的CUDA API调用产生的错误。如果没有错误发生，则返回cudaSuccess。这个函数通常在核函数调用之后使用，以检查核函数是否成功执行。

• 参数:
  无。

• 返回值:
  最近一次CUDA错误码。

• 示例:

  kernel<<<blocks, threads>>>(...);
  cudaError_t err = cudaGetLastError();
  if (err != cudaSuccess) {
      printf("Kernel launch failed: %s\n", cudaGetErrorString(err));
  }
  

4. cudaPeekAtLastError(void)

• 函数原型:

  __host__ __device__ cudaError_t cudaPeekAtLastError(void);
  

• 功能:
  这个函数返回最近一次执行的CUDA API调用产生的错误，但不清除错误状态。这对于需要多次检查相同错误状态的情况非常有用。

• 参数:
  无。

• 返回值:
  最近一次CUDA错误码。

• 示例:

  kernel<<<blocks, threads>>>(...);
  cudaError_t err = cudaPeekAtLastError();
  if (err != cudaSuccess) {
      printf("Kernel launch status: %s\n", cudaGetErrorString(err));
  }
  

2.2 使用示例

在示例中，我们使用最近一次执行的CUDA API调用产生的错误，并将更详细的错误信息给输出，检测代码如下：

cudaError_t err = cudaGetLastError();		// Get error code
if(err != cudaSuccess)
{
	printf("CUDA Error:%s\n"，cudaGetErrorstring(err));
	exit(-1);
}

将上述检测代码嵌入到我们执行的CUDA核函数下方，如下图所示：

重新编译CUDA文件，如果代码中存在错误，将会如下图所示输出。

---

3 封装在.cuh头文件中嵌入

3.1 创建 error.cuh 头文件

我们将检测代码做一个宏定义，并将其封装在 error.cuh 文件中，用于检查CUDA函数调用的返回值是否表示成功。如果CUDA函数返回一个错误代码，这个宏将打印错误信息并退出程序。

#pragma once
#include <stdio.h>

#define CHECK(call)                                   \
do                                                    \
{                                                     \
    const cudaError_t error_code = call;              \
    if (error_code != cudaSuccess)                    \
    {                                                 \
        printf("CUDA Error:\n");                      \
        printf("    File:       %s\n", __FILE__);     \
        printf("    Line:       %d\n", __LINE__);     \
        printf("    Error code: %d\n", error_code);   \
        printf("    Error text: %s\n",                \
            cudaGetErrorString(error_code));          \
        exit(1);                                      \
    }                                                 \
} while (0)

如果发生错误，执行以下代码块：

• printf("CUDA Error:\n")
  • 打印错误信息头。
• printf(" File: %s\n", __FILE__)
  • 打印发生错误的源文件名，__FILE__是预定义的宏，表示当前文件名。
• printf(" Line: %d\n", __LINE__)
  • 打印发生错误的行号，__LINE__是预定义的宏，表示当前行号。
• printf(" Error code: %d\n", error_code)
  • 打印错误代码。
• printf(" Error text: %s\n", cudaGetErrorString(error_code))
  • 打印错误文本描述，cudaGetErrorString函数将错误代码转换为可读的字符串。
• exit(1)
  • 退出程序，并返回状态码1，表示发生错误。

3.2 在 CUDA 程序中包含 error.cuh 并调用 CHECK 宏

在你的CUDA源文件中，包括 error.cuh 头文件，然后使用 CHECK 宏来包装CUDA API调用。以下是一个完整的使用示例程序：

// main.cu
#include <cuda_runtime.h>
#include <stdio.h>
#include "error.cuh"

// 简单的CUDA核函数
__global__ void kernel()
{
    // 核函数内部代码
}

int main()
{
    // 设备属性
    cudaDeviceProp deviceProp;
    int dev = 0;

    // 获取设备属性
    CHECK(cudaGetDeviceProperties(&deviceProp, dev));

    // 设置设备
    CHECK(cudaSetDevice(dev));

    // 分配设备内存
    int* d_array;
    size_t size = 100 * sizeof(int);
    CHECK(cudaMalloc((void**)&d_array, size));

    // 启动核函数
    kernel<<<1, 1>>>();

    // 检查核函数执行情况
    CHECK(cudaGetLastError());

    // 释放设备内存
    CHECK(cudaFree(d_array));

    // 重置设备
    CHECK(cudaDeviceReset());

    printf("CUDA程序成功完成。\n");
    return 0;
}

3.3 使用示例

首先在你的CUDA源文件中，引入 error.cuh 头文件。

在需要检查的位置插入CHECK。

编译并运行，如果CUDA源文件有错误，就会如下所示：

通过这种方式，你可以在不同的CUDA源文件中复用 CHECK 宏，而无需在每个文件中重复定义。

---

本篇采用的错误代码来源于：block_size设置过大错误分析（查看CUDA设备线程块大小）

部分参考信息来源：https://mp.weixin.qq.com/s/em309Ho6AaV5f1ogWpajJw

本人是一名学生，也是正在学习Jetson的过程中，如有错误请批评指正！