《CUDA编程》4.CUDA步伐的错误检测 - IT评测·应用市场-qidao123.com技术社区

#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)

复制代码

#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);
}

复制代码

CHECK(cudaDeviceSynchronize());
CHECK(cudaGetLastError());

复制代码

#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);
}

复制代码

cuda-memcheck ./my_cuda_program.out

复制代码

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

复制代码

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

复制代码

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

复制代码