在编写CUDA程序时，有的错误在编译过程中被发现，称为编译错误，有的在运行时出现，称为运行时刻错误，本章讨论如何排查运行时刻错误
1 一个检测CUDA运行时错误的宏函数
1.1 编写错误检查宏函数
在《CUDA编程》3.简单CUDA程序的基本框架 中列举的函数，返回值是cudaError_t，只有在返回cudaSuccess时，才表示调用成功，否则返回一个错误代码，下面新建一个CUDA头文件并编写一个错误检查的宏函数：
①新建CUDA头文件
新建的文件是error_check.cuh，注意后缀变化。
②编写错误检查代码
在定义宏时，如果一行写不下，需要在行末写 \，表示续行*，错误检查代码如下：
#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 message: %s\n", cudaGetErrorString(error_code)); \
exit(1); \
} \
} while (0)
该段代码会检查返回值是否为cudaSuccess，如果不是，则返回错误代码的位置
1.2 把检查函数添加到CUDA程序中
这里以《CUDA编程》3.简单CUDA程序的基本框架中的代码例子为例，注意在头文件中添加#include "error_check.cuh"，并为代码中分配内存的函数进行检查。
然后手动将39行的代码修改为 CHECK(cudaMemcpy(d_x, h_x, M, cudaMemcpyDeviceToHost));，修改后是错误代码，原本应该是cudaMemcpyHostToDevice
#include <cuda.h>
#include <cuda_runtime.h>
#include <math.h>
#include <stdio.h>
#include "error_check.cuh"
const double EPS = 1.0e-15;
const double a = 1.23;
const double b = 2.34;
const double c = 3.57;
// 希望 add 函数在 GPU 上执行
__global__ void add(const double* x, const double* y, double* z);
void check(const double* z, const int N);
int main(void) {
const int N = 100000000; // 定义数组的长度为 10 的 8 次方
const int M = sizeof(double) * N; // 每个数组所需的字节数
// 分配host内存
double* h_x = (double*)malloc(M);
double* h_y = (double*)malloc(M);
double* h_z = (double*)malloc(M);
for (int n = 0; n < N; ++n) {
h_x[n] = a;
h_y[n] = b;
}
//分配device内存
double* d_x, * d_y, * d_z;
CHECK(cudaMalloc((void**)&d_x, M));
CHECK(cudaMalloc((void**)&d_y, M));
CHECK(cudaMalloc((void**)&d_z, M));
// 将数据从主机复制到设备上
CHECK(cudaMemcpy(d_x, h_x, M, cudaMemcpyDeviceToHost));
CHECK(cudaMemcpy(d_y, h_y, M, cudaMemcpyHostToDevice));
const int block_size = 128;
// 计算网格尺寸，确保所有元素都能被处理
const int grid_size = (N + block_size - 1) / block_size;
// 调用内核函数在设备中进行计算
add << <grid_size, block_size >> > (d_x, d_y, d_z);
// 将计算结果从设备复制回主机
CHECK(cudaMemcpy(h_z, d_z, M, cudaMemcpyDeviceToHost));
check(h_z, N);
// 释放内存
free(h_x);
free(h_y);
free(h_z);
CHECK(cudaFree(d_x));
CHECK(cudaFree(d_y));
CHECK(cudaFree(d_z));
return 0;
}
__global__ void add(const double* x, const double* y, double* z) {
const int n = blockIdx.x * blockDim.x + threadIdx.x;
z[n] = x[n] + y[n];
}
void check(const double* z, const int N) {
bool has_error = false;
for (int n = 0; n < N; ++n) {
if (fabs(z[n] - c) > EPS) {
has_error = true;
}
}
printf("Has error: %d\n", has_error);
}
运行结果如下：
指出了错误代码的信息，包括文件位置、行数、个数、和错误类型invalid argument，及代表该行函数出现了非法参数，正是由于我们手动修改导致的错误
PS: 大部分代码都可以使用该宏函数，除了cudaEventQuery()，因为它可能返回cudaErrorNotReady，但并不是代码出错了
1.3 使用该宏函数检查核函数错误
使用上述方法并不能捕捉核函数的错误，因为核函数不返回任何值，所以若想捕捉和函数的错误，应该在调用核函数之后使用如下语句：
CHECK(cudaDeviceSynchronize());
CHECK(cudaGetLastError());
- 第一个语句是同步主机和设备，因为核函数的调用是异步的，使用该函数可以确保之前的CUDA操作全部完成，以便检查这些操作是否成功
- 返回自上次调用
cudaGetLastError()或者自程序开始以来最后一个 CUDA API 调用的错误代码。
依旧以上面的函数作为例子，手动的将block_size修改为1280，但我们知道该参数不能超过1024，所以会报错，代码如下：
#include <cuda.h>
#include <cuda_runtime.h>
#include <math.h>
#include <stdio.h>
#include "error_check.cuh"
const double EPS = 1.0e-15;
const double a = 1.23;
const double b = 2.34;
const double c = 3.57;
// 希望 add 函数在 GPU 上执行
__global__ void add(const double* x, const double* y, double* z);
void check(const double* z, const int N);
int main(void) {
const int N = 100000000; // 定义数组的长度为 10 的 8 次方
const int M = sizeof(double) * N; // 每个数组所需的字节数
// 分配host内存
double* h_x = (double*)malloc(M);
double* h_y = (double*)malloc(M);
double* h_z = (double*)malloc(M);
for (int n = 0; n < N; ++n) {
h_x[n] = a;
h_y[n] = b;
}
//分配device内存
double* d_x, * d_y, * d_z;
CHECK(cudaMalloc((void**)&d_x, M));
CHECK(cudaMalloc((void**)&d_y, M));
CHECK(cudaMalloc((void**)&d_z, M));
// 将数据从主机复制到设备上
CHECK(cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice));
CHECK(cudaMemcpy(d_y, h_y, M, cudaMemcpyHostToDevice));
const int block_size = 1280;
// 计算网格尺寸，确保所有元素都能被处理
const int grid_size = (N + block_size - 1) / block_size;
// 调用内核函数在设备中进行计算
add << <grid_size, block_size >> > (d_x, d_y, d_z);
CHECK(cudaDeviceSynchronize());
CHECK(cudaGetLastError());
// 将计算结果从设备复制回主机
CHECK(cudaMemcpy(h_z, d_z, M, cudaMemcpyDeviceToHost));
check(h_z, N);
// 释放内存
free(h_x);
free(h_y);
free(h_z);
CHECK(cudaFree(d_x));
CHECK(cudaFree(d_y));
CHECK(cudaFree(d_z));
return 0;
}
__global__ void add(const double* x, const double* y, double* z) {
const int n = blockIdx.x * blockDim.x + threadIdx.x;
z[n] = x[n] + y[n];
}
void check(const double* z, const int N) {
bool has_error = false;
for (int n = 0; n < N; ++n) {
if (fabs(z[n] - c) > EPS) {
has_error = true;
}
}
printf("Has error: %d\n", has_error);
}
输出结果如下：
即表示配置错误，如果不使用该函数，则只能发现有一个错误，而不知道具体的错误信息。
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 ，命令如下：
cuda-memcheck ./my_cuda_program.out
其他三个不可以简化
使用racecheck工具：
cuda-memcheck --tool racecheck ./my_cuda_program.out
使用synccheck工具
cuda-memcheck --tool synccheck ./my_cuda_program.out
使用initcheck工具
cuda-memcheck --tool initcheck ./my_cuda_program.out
