cuda-3-错误

cuda

  • CUDA运行时API大多支持返回错误代码，返回值类型：cudaError_t
  • 运行时API成功执行，返回值为cudaSuccess
  • 运行时API返回的执行状态值是枚举变量，对应官方文档查看错误代码可以知道错误来源

捕捉主机函数错误：

​ CUDA代码在主机端（CPU）和设备端（GPU）的错误检测方法并不相同

  • 在主机端（CPU）执行的CUDA代码主要包括对CUDA API函数的调用、内存分配与管理、数据传输以及内核启动等操作。主机端错误检测通常涉及检查CUDA API函数的返回值，这些函数会返回一个cudaError_t类型的值
  • 而在设备端（GPU）执行的CUDA内核代码无法直接检查错误，因为且内核函数本身不能主动报告错误。这时也需要在主机端进行错误检查。对于设备端可能发生的错误，如访问越界、浮点异常、并发错误等，但是CUDA提供方式进行事后分析

获取错误代码对应名称：

1
__host__ __device const char* cudaGetErrorName(cudaError_t Error)

获取错误代码描述信息：

1
__host__ __device__ const char* cudaGetErrorString(cadaError_t Error)

错误检查函数：

1
2
3
4
5
6
7
8
9
10
11
12
13
#pragma once
#define ErrorCheck(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)

只需要把这个 CHECK函数包裹可能发生错误的函数即可

​ #define CHECK(call) 定义了一个宏，其功能是将 CUDA 函数调用放入一个 do-while 循环中进行错误检查。这段代码采用C/C++预处理器宏（Macro）来实现错误检查机制，利用了C语言的宏替换以及一些编译器提供的特殊标识符（如 __FILE__ 和 __LINE__）等特性

捕捉核函数错误：

​ 在调用核函数后，追加如下代码：

1
2
ErrorCheck(cudaGetLastError());
ErrorCheck(cudaDeviceSynchronize());
  1. ErrorCheck函数定义与之前相同
  2. 第一条语句作用是捕捉第二条同步函数之前的最后一个错误
  3. 第二条语句同步主机与设备，因为CPU和GPU是异构架构

应用实例：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
#include <stdio.h>
#include "../tools/common.cuh"

__device__ float add(const float x, const float y) {
    return x + y;
}

__global__ void addFromGPU(float *A, float *B, float *C, const int N) {
    const int bid = blockIdx.x;
    const int tid = threadIdx.x;
    const int id = tid + bid * blockDim.x; 

    if (id >= N) return;
    C[id] = add(A[id], B[id]);
    
}


void initialData(float *addr, int elemCount)
{
    for (int i = 0; i < elemCount; i++) {
        addr[i] = (float)(rand() & 0xFF) / 10.f;
    }
    return;
}


int main(void)
{
    // 1、设置GPU设备
    setGPU();

    // 2、分配主机内存和设备内存，并初始化
    int iElemCount = 4096;                     // 设置元素数量
    size_t stBytesCount = iElemCount * sizeof(float); // 字节数
    
    // （1）分配主机内存，并初始化
    float *fpHost_A, *fpHost_B, *fpHost_C;
    fpHost_A = (float *)malloc(stBytesCount);
    fpHost_B = (float *)malloc(stBytesCount);
    fpHost_C = (float *)malloc(stBytesCount);
    if (fpHost_A != NULL && fpHost_B != NULL && fpHost_C != NULL) {
        memset(fpHost_A, 0, stBytesCount);  // 主机内存初始化为0
        memset(fpHost_B, 0, stBytesCount);
        memset(fpHost_C, 0, stBytesCount);
    
    }
    else {
        printf("Fail to allocate host memory!\n");
        exit(-1);
    }


    // （2）分配设备内存，并初始化
    float *fpDevice_A, *fpDevice_B, *fpDevice_C;
    cudaMalloc((float**)&fpDevice_A, stBytesCount);
    cudaMalloc((float**)&fpDevice_B, stBytesCount);
    cudaMalloc((float**)&fpDevice_C, stBytesCount);
    if (fpDevice_A != NULL && fpDevice_B != NULL && fpDevice_C != NULL) {
        cudaMemset(fpDevice_A, 0, stBytesCount);  // 设备内存初始化为0
        cudaMemset(fpDevice_B, 0, stBytesCount);
        cudaMemset(fpDevice_C, 0, stBytesCount);
    }
    else {
        printf("fail to allocate memory\n");
        free(fpHost_A);
        free(fpHost_B);
        free(fpHost_C);
        exit(-1);
    }

    // 3、初始化主机中数据
    srand(666); // 设置随机种子
    initialData(fpHost_A, iElemCount);
    initialData(fpHost_B, iElemCount);
    
    // 4、数据从主机复制到设备
    cudaMemcpy(fpDevice_A, fpHost_A, stBytesCount, cudaMemcpyHostToDevice); 
    cudaMemcpy(fpDevice_B, fpHost_B, stBytesCount, cudaMemcpyHostToDevice); 
    cudaMemcpy(fpDevice_C, fpHost_C, stBytesCount, cudaMemcpyHostToDevice);


    // 5、调用核函数在设备中进行计算
    dim3 block(2048);
    dim3 grid((iElemCount + block.x - 1) / 2048); 

    addFromGPU<<<grid, block>>>(fpDevice_A, fpDevice_B, fpDevice_C, iElemCount);    // 调用核函数
    ErrorCheck(cudaGetLastError(), __FILE__, __LINE__);
    ErrorCheck(cudaDeviceSynchronize(), __FILE__, __LINE__);

    // 6、将计算得到的数据从设备传给主机
    cudaMemcpy(fpHost_C, fpDevice_C, stBytesCount, cudaMemcpyDeviceToHost);
  
    for (int i = 0; i < 10; i++) {
        printf("idx=%2d\tmatrix_A:%.2f\tmatrix_B:%.2f\tresult=%.2f\n", i+1, fpHost_A[i], fpHost_B[i], fpHost_C[i]);
    }

    // 7、释放主机与设备内存
    free(fpHost_A);
    free(fpHost_B);
    free(fpHost_C);
    cudaFree(fpDevice_A);
    cudaFree(fpDevice_B);
    cudaFree(fpDevice_C);

    cudaDeviceReset();
    return 0;
}

>>> output:
The count of GPUs is 1.
set GPU 0 for computing.
CUDA error:
code=9, name=cudaErrorInvalidConfiguration, description=invalid configuration argument
file=errorCheckKernel.cu, line95
idx= 1  matrix_A:0.90   matrix_B:10.90  result=0.00
idx= 2  matrix_A:19.00  matrix_B:4.00   result=0.00
idx= 3  matrix_A:15.80  matrix_B:15.00  result=0.00
idx= 4  matrix_A:5.00   matrix_B:3.30   result=0.00
idx= 5  matrix_A:11.10  matrix_B:4.20   result=0.00
idx= 6  matrix_A:23.50  matrix_B:18.60  result=0.00
idx= 7  matrix_A:20.90  matrix_B:4.50   result=0.00
idx= 8  matrix_A:23.40  matrix_B:17.70  result=0.00
idx= 9  matrix_A:16.90  matrix_B:18.40  result=0.00
idx=10  matrix_A:7.30   matrix_B:18.30  result=0.00

​ 其中"../tools/common.cuh"中包含的文件就是上面写出来的ErrorCheck函数