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的过程中,如有错误请批评指正!