《CUDA编程》4.CUDA步伐的错误检测

在编写CUDA步伐时，有的错误在编译过程中被发现，称为编译错误，有的在运行时出现，称为运行时刻错误，本章讨论怎样排查运行时刻错误
  1 一个检测CUDA运行时错误的宏函数

1.1 编写错误查抄宏函数

在《CUDA编程》3.简朴CUDA步伐的基本框架 中枚举的函数，返回值是cudaError_t，只有在返回cudaSuccess时，才表现调用乐成，否则返回一个错误代码，下面新建一个CUDA头文件并编写一个错误查抄的宏函数：
①新建CUDA头文件

新建的文件是error_check.cuh，注意后缀变化。
②编写错误查抄代码
在界说宏时，如果一行写不下，必要在行末写 \，表现续行*，错误查抄代码如下：

1. #pragma once
2. #include <stdio.h>
4. #define CHECK(call) \
5. do { \
6.     const cudaError_t error_code = call; \
7.     if (error_code != cudaSuccess) { \
8.         printf("CUDA Error:\n"); \
9.         printf("File: %s\n", __FILE__); \
10.         printf("Line: %d\n", __LINE__); \
11.         printf("Error code: %d\n", error_code); \
12.         printf("Error message: %s\n", cudaGetErrorString(error_code)); \
13.         exit(1); \
14.     } \
15. } while (0)

复制代码

该段代码会查抄返回值是否为cudaSuccess，如果不是，则返回错误代码的位置
1.2 把查抄函数添加到CUDA步伐中

这里以《CUDA编程》3.简朴CUDA步伐的基本框架中的代码例子为例，注意在头文件中添加#include "error_check.cuh"，并为代码中分配内存的函数进行查抄。
然后手动将39行的代码修改为 CHECK(cudaMemcpy(d_x, h_x, M, cudaMemcpyDeviceToHost));，修改后是错误代码，原本应该是cudaMemcpyHostToDevice

1. #include <cuda.h>
2. #include <cuda_runtime.h>
3. #include <math.h>
4. #include <stdio.h>
5. #include "error_check.cuh"
8. const double EPS = 1.0e-15;
9. const double a = 1.23;
10. const double b = 2.34;
11. const double c = 3.57;
13. // 希望 add 函数在 GPU 上执行
14. __global__ void add(const double* x, const double* y, double* z);
15. void check(const double* z, const int N);
17. int main(void) {
18.     const int N = 100000000; // 定义数组的长度为 10 的 8 次方
19.     const int M = sizeof(double) * N; // 每个数组所需的字节数
21.     // 分配host内存
22.     double* h_x = (double*)malloc(M);
23.     double* h_y = (double*)malloc(M);
24.     double* h_z = (double*)malloc(M);
27.     for (int n = 0; n < N; ++n) {
28.         h_x[n] = a;
29.         h_y[n] = b;
30.     }
32.     //分配device内存
33.     double* d_x, * d_y, * d_z;
34.     CHECK(cudaMalloc((void**)&d_x, M));
35.     CHECK(cudaMalloc((void**)&d_y, M));
36.     CHECK(cudaMalloc((void**)&d_z, M));
37.    
38.     // 将数据从主机复制到设备上
39.     CHECK(cudaMemcpy(d_x, h_x, M, cudaMemcpyDeviceToHost));
40.     CHECK(cudaMemcpy(d_y, h_y, M, cudaMemcpyHostToDevice));
42.     const int block_size = 128;
43.     // 计算网格尺寸，确保所有元素都能被处理
44.     const int grid_size = (N + block_size - 1) / block_size;
46.     // 调用内核函数在设备中进行计算
47.     add << <grid_size, block_size >> > (d_x, d_y, d_z);
49.     // 将计算结果从设备复制回主机
50.     CHECK(cudaMemcpy(h_z, d_z, M, cudaMemcpyDeviceToHost));
51.     check(h_z, N);
53.     // 释放内存
54.     free(h_x);
55.     free(h_y);
56.     free(h_z);
57.     CHECK(cudaFree(d_x));
58.     CHECK(cudaFree(d_y));
59.     CHECK(cudaFree(d_z));
60.     return 0;
61. }
63. __global__ void add(const double* x, const double* y, double* z) {
64.     const int n = blockIdx.x * blockDim.x + threadIdx.x;
65.     z[n] = x[n] + y[n];
66. }
68. void check(const double* z, const int N) {
69.     bool has_error = false;
70.     for (int n = 0; n < N; ++n) {
71.         if (fabs(z[n] - c) > EPS) {
72.             has_error = true;
73.         }
74.     }
75.     printf("Has error: %d\n", has_error);
76. }

复制代码

运行结果如下：

指出了错误代码的信息，包罗文件位置、行数、个数、和错误范例invalid argument，及代表该行函数出现了非法参数，正是由于我们手动修改导致的错误
PS: 大部分代码都可以利用该宏函数，除了cudaEventQuery()，因为它大概返回cudaErrorNotReady，但并不是代码出错了
1.3 利用该宏函数查抄核函数错误

利用上述方法并不能捕捉核函数的错误，因为核函数不返回任何值，以是若想捕捉和函数的错误，应该在调用核函数之后利用如下语句：

1. CHECK(cudaDeviceSynchronize());
2. CHECK(cudaGetLastError());

复制代码

• 第一个语句是同步主机和设备，因为核函数的调用是异步的，利用该函数可以确保之前的CUDA操作全部完成，以便查抄这些操作是否乐成
  
• 返回自上次调用 cudaGetLastError() 或者自步伐开始以来最后一个 CUDA API 调用的错误代码。
  

仍然以上面的函数作为例子，手动的将block_size修改为1280，但我们知道该参数不能高出1024，以是会报错，代码如下：

1. #include <cuda.h>
2. #include <cuda_runtime.h>
3. #include <math.h>
4. #include <stdio.h>
5. #include "error_check.cuh"
8. const double EPS = 1.0e-15;
9. const double a = 1.23;
10. const double b = 2.34;
11. const double c = 3.57;
13. // 希望 add 函数在 GPU 上执行
14. __global__ void add(const double* x, const double* y, double* z);
15. void check(const double* z, const int N);
17. int main(void) {
18.     const int N = 100000000; // 定义数组的长度为 10 的 8 次方
19.     const int M = sizeof(double) * N; // 每个数组所需的字节数
21.     // 分配host内存
22.     double* h_x = (double*)malloc(M);
23.     double* h_y = (double*)malloc(M);
24.     double* h_z = (double*)malloc(M);
27.     for (int n = 0; n < N; ++n) {
28.         h_x[n] = a;
29.         h_y[n] = b;
30.     }
32.     //分配device内存
33.     double* d_x, * d_y, * d_z;
34.     CHECK(cudaMalloc((void**)&d_x, M));
35.     CHECK(cudaMalloc((void**)&d_y, M));
36.     CHECK(cudaMalloc((void**)&d_z, M));
37.    
38.     // 将数据从主机复制到设备上
39.     CHECK(cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice));
40.     CHECK(cudaMemcpy(d_y, h_y, M, cudaMemcpyHostToDevice));
42.     const int block_size = 1280;
43.     // 计算网格尺寸，确保所有元素都能被处理
44.     const int grid_size = (N + block_size - 1) / block_size;
46.     // 调用内核函数在设备中进行计算
47.     add << <grid_size, block_size >> > (d_x, d_y, d_z);
48.     CHECK(cudaDeviceSynchronize());
49.     CHECK(cudaGetLastError());
51.     // 将计算结果从设备复制回主机
52.     CHECK(cudaMemcpy(h_z, d_z, M, cudaMemcpyDeviceToHost));
53.     check(h_z, N);
55.     // 释放内存
56.     free(h_x);
57.     free(h_y);
58.     free(h_z);
59.     CHECK(cudaFree(d_x));
60.     CHECK(cudaFree(d_y));
61.     CHECK(cudaFree(d_z));
62.     return 0;
63. }
65. __global__ void add(const double* x, const double* y, double* z) {
66.     const int n = blockIdx.x * blockDim.x + threadIdx.x;
67.     z[n] = x[n] + y[n];
68. }
70. void check(const double* z, const int N) {
71.     bool has_error = false;
72.     for (int n = 0; n < N; ++n) {
73.         if (fabs(z[n] - c) > EPS) {
74.             has_error = true;
75.         }
76.     }
77.     printf("Has error: %d\n", has_error);
78. }

复制代码

输出结果如下：

即表现配置错误，如果不利用该函数，则只能发现有一个错误，而不知道具体的错误信息。
PS: cudaDeviceSynchronize()非常斲丧时间，以是一般不在内存循环中调用，否则会严重低沉步伐性能
2 用CUDA-MEMCHECK查抄内存错误

CUDA提供了CUDA-MEMCHECK工具集，可以帮助你发现诸如越界访问、未初始化内存访问、内存泄漏等内存错误，从而进步代码的可靠性和性能。一共包含了4个工具：

• memcheck：用于检测内存访问错误，包罗越界访问、未初始化内存访问等，常见错误范例有：
  –Global Out-of-bounds：访问超出全局内存范围。
  –Local Out-of-bounds：访问超出局部内存范围。
  –Uninitialized Access：访问未初始化的内存。
  –Invalid Device Pointer：利用无效的设备指针
  
• racecheck：用于检测数据竞争，即多个线程同时访问同一内存位置且至少有一个线程在写入，常见错误范例有：
  –Race Condition：多个线程同时访问同一内存位置且至少有一个线程在写入。
  
• synccheck：用于检测同步错误，即线程之间的同步题目，常见错误范例有：
  –Barrier Synchronization Error：线程在屏障同步点出现错误。
  –Grid Synchronization Error：线程在网格同步点出现错误。
  
• initcheck：用于检测未初始化内存的利用，常见错误范例有：
  –Uninitialized Memory Use：利用未初始化的内存。
  

以上4个工具都可由cuda-memcheck执行文件调用，此中调用memcheck时，可以简化，注意，只能对编译后的文件进行查抄，通常是.out ，命令如下：

1. cuda-memcheck ./my_cuda_program.out

复制代码

其他三个不可以简化
利用racecheck工具：

1. cuda-memcheck --tool racecheck ./my_cuda_program.out

复制代码

利用synccheck工具

1. cuda-memcheck --tool synccheck ./my_cuda_program.out

复制代码

利用initcheck工具

1. cuda-memcheck --tool initcheck ./my_cuda_program.out

复制代码

免责声明：如果侵犯了您的权益，请联系站长，我们会及时删除侵权内容，谢谢合作！更多信息从访问主页：qidao123.com:ToB企服之家，中国第一个企服评测及商务社交产业平台。