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

1 CUDA错误检测简介
2 直接嵌入检测函数
- 2.1 检测函数介绍
- 2.2 使用示例
3 封装在`.cuh`头文件中嵌入
- 3.1 创建 `error.cuh` 头文件
- 3.2 在 CUDA 程序中包含 `error.cuh` 并调用 CHECK 宏
- 3.3 使用示例

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