《CUDA编程》4.CUDA步伐的错误检测
在编写CUDA步伐时，有的错误在编译过程中被发现，称为编译错误，有的在运行时出现，称为运行时刻错误，本章讨论怎样排查运行时刻错误1 一个检测CUDA运行时错误的宏函数
1.1 编写错误查抄宏函数
在《CUDA编程》3.简朴CUDA步伐的基本框架 中枚举的函数，返回值是cudaError_t，只有在返回cudaSuccess时，才表现调用乐成，否则返回一个错误代码，下面新建一个CUDA头文件并编写一个错误查抄的宏函数：
①新建CUDA头文件
https://i-blog.csdnimg.cn/direct/a97868d13c7d4971ad613fe946f5e59e.png#pic_center
新建的文件是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 = a;
h_y = 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 = x + y;
}
void check(const double* z, const int N) {
bool has_error = false;
for (int n = 0; n < N; ++n) {
if (fabs(z - c) > EPS) {
has_error = true;
}
}
printf("Has error: %d\n", has_error);
}
运行结果如下：
https://i-blog.csdnimg.cn/direct/5deee0c36d6143f09310a628eb2c9f05.png
指出了错误代码的信息，包罗文件位置、行数、个数、和错误范例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 = a;
h_y = 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 = x + y;
}
void check(const double* z, const int N) {
bool has_error = false;
for (int n = 0; n < N; ++n) {
if (fabs(z - c) > EPS) {
has_error = true;
}
}
printf("Has error: %d\n", has_error);
}
输出结果如下：
https://i-blog.csdnimg.cn/direct/bcb505e19a934f259b4ceff3eea5e848.png
即表现配置错误，如果不利用该函数，则只能发现有一个错误，而不知道具体的错误信息。
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
免责声明：如果侵犯了您的权益，请联系站长，我们会及时删除侵权内容，谢谢合作！更多信息从访问主页：qidao123.com:ToB企服之家，中国第一个企服评测及商务社交产业平台。
页:
[1]