cuda-memcheck/compute-sanitizer工具使用

        cuda-memcheck是CUDA工具包中的一个功能正确性检查套件。它包含多个可以执行不同类型检查的工具。这些工具可以精确地检测和归因CUDA应用程序中的越界和未对齐的内存访问错误。具体包括 memcheck、 racecheck、initcheck、 synccheck 共 4 个工具。它们可由可执行文件 cuda-memcheck 调用：

$ cuda-memcheck --tool memcheck [options] app_name [options]
$ cuda-memcheck --tool racecheck [options] app_name [options]
$ cuda-memcheck --tool initcheck [options] app_name [options]
$ cuda-memcheck --tool synccheck [options] app_name [options]

• memcheck：用于内存访问错误和泄漏检测
• racecheck：共享内存数据访问危险检测工具
• initcheck：未初始化的设备全局内存访问检测工具
• synccheck：用于线程同步危险检测

具体参数信息可以参考官网：CUDA-MEMCHECK :: CUDA Toolkit Documentation (nvidia.com)

 注意：此工具已在 CUDA 11.6 中被弃用，并在12.0 及更高版本中被删除。Compute Sanitizer 已取代它的位置，用法差不多，但是提供了额外的功能，如改进的性能和对 Microsoft hardware-accelerated GPU scheduling 的支持，以及对内存检查之外的功能的更广泛支持。

简单使用：

        如下有问题代码编译后，使用 cuda-memcheck ./test.out 或者compute-sanitizer ./test.out执行后则会得到一大段输出以及

========= ERROR SUMMARY: 2 errors

#include <math.h>
#include <stdio.h>

const double EPSILON = 1.0e-15;
const double a = 1.23;
const double b = 2.34;
const double c = 3.57;
void __global__ add(const double *x, const double *y, double *z);
void check(const double *z, const int N);

int main(void)
{
    const int N = 100000000;
    const int M = sizeof(double) * N;
    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;
    }

    double *d_x, *d_y, *d_z;
    cudaMalloc((void **)&d_x, M);
    cudaMalloc((void **)&d_y, M);
    cudaMalloc((void **)&d_z, M);
    cudaMemcpy(d_x, h_x, M, cudaMemcpyDeviceToHost);
    cudaMemcpy(d_y, h_y, M, cudaMemcpyDeviceToHost);

    const int block_size = 128;
    const int grid_size = N / block_size;
    add<<<grid_size, block_size>>>(d_x, d_y, d_z);

    cudaMemcpy(h_z, d_z, M, cudaMemcpyDeviceToHost);
    check(h_z, N);

    free(h_x);
    free(h_y);
    free(h_z);
    cudaFree(d_x);
    cudaFree(d_y);
    cudaFree(d_z);
    return 0;
}

void __global__ add(const double *x, const double *y, double *z)
{
    // 内存泄露
    const int n = blockDim.x * blockIdx.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) > EPSILON)
        {
            has_error = true;
        }
    }
    printf("%s\n", has_error ? "Has errors" : "No errors");
}