CUDA C++ 是一种带有扩展功能的标准 C++，支持函数在 GPU 的多个并行线程上执行。它在推动广泛应用的同时，也使开发者能够实现高性能。然而，CUDA C++ 并非内存安全的语言。在开发过程中，若缺乏合适的调试工具，细微的内存错误可能难以被察觉。
NVIDIA Compute Sanitizer 是一款帮助开发者发现程序中错误的工具。 NVIDIA CUDA 13.1 引入了一个新的编译器选项，旨在增强 Compute Sanitizer 的“memcheck”功能。该编译器选项可为开发者提供更全面的错误覆盖以及更高效的执行速度。
如果您不了解 Compute Sanitizer 是什么，请继续阅读，了解它如何帮助您避免程序出错。如果您已经使用过 Compute Sanitizer，请继续阅读，了解我们的新编译器分析功能如何帮助发现代码中一些难以察觉的内存错误。
使用 Compute Sanitizer 发现错误
我们来深入理解并使用 Compute Sanitizer，以发现以下程序中的错误。该程序改编自上一篇关于 Compute Sanitizer 博客文章中的示例。你能在下面的代码中找出逻辑错误吗？提示：这是一个典型的“差一”错误，我们都曾遇到过，次数可能比我们愿意承认的还要多。在我们的代码示例中，不仅启动了过多的线程 scaleArray，还在 scaleArray 的 if 条件中使用了错误的判断条件。这样一来，线程 512 将会访问超出边界的 array[512]。
#include <assert.h>
#include <stdio.h>
__global__ void scaleArray(float* array, size_t N, float value) {
int threadGlobalID = threadIdx.x + blockIdx.x * blockDim.x;
if (threadGlobalID <= N) {
array[threadGlobalID] = array[threadGlobalID]*value;
}
}
int main() {
float* array = nullptr;
float* buffer = nullptr;
const size_t N = 512;
// Allocate N float-size elements, visible to both CPU and GPU
cudaMallocManaged(&array, N*sizeof(float));
//cudaMallocManaged(&buffer, N*sizeof(float));
for (int i=0; i<N; i++) array[i] = 1.0f; // Initialize array
printf("Before: Array 0, 1 .. N-1: %f %f %f\n", array[0], array[1], array[N-1]);
scaleArray<<<1,1024>>>(array, N, 3.0);
cudaDeviceSynchronize();
printf("After : Array 0, 1 .. N-1: %f %f %f\n", array[0], array[1], array[N-1]);
assert(array[N/2] == 3.0); // Check that it has worked
cudaFree(array);
cudaFree(buffer);
exit(0);
}
我们来看看在编译和运行程序时能否观察到此错误：
$ /usr/local/cuda-13.1/bin/nvcc -o example1 -arch sm_86 example1.cu
$ ./example1
Before: Array 0, 1 .. N-1: 1.000000 1.000000 1.000000
After : Array 0, 1 .. N-1: 3.000000 3.000000 3.000000
在我们的测试机器上，此程序运行正常。示例中的漏洞在开发过程中容易被忽视，但后续可能引发数据损坏、段错误，甚至代码安全问题。
现在，让我们通过 Compute Sanitizer 运行程序，看看是否能发现潜在问题。
$ compute-sanitizer ./example1
========= COMPUTE-SANITIZER
Before: Array 0, 1 .. N-1: 1.000000 1.000000 1.000000
========= Invalid __global__ read of size 4 bytes
========= at scaleArray(float *, unsigned long, float)+0xb0
========= by thread (512,0,0) in block (0,0,0)
========= Access to 0x7f7754000800 is out of bounds
========= and is 1 bytes after the nearest allocation at 0x7f7754000000 of size 2,048 bytes
========= Saved host backtrace up to driver entry point at kernel launch time
========= Host Frame: main [0x8deb] in example1
=========
========= Program hit cudaErrorLaunchFailure (error 719) due to "unspecified launch failure" on CUDA API call to cudaDeviceSynchronize.
========= Saved host backtrace up to driver entry point at error
========= Host Frame: main [0x8df0] in example1
=========
After : Array 0, 1 .. N-1: 1.000000 1.000000 1.000000
example1: example1.cu:27: int main(): Assertion `array[N/2] == 3.0' failed.
========= Error: process didn't terminate successfully
========= Target application returned an error
========= ERROR SUMMARY: 2 errors
好，确实如此！但是，与所有商用内存清理工具一样，Compute Sanitizer 也可能存在假阴性情况，即可能遗漏某些实际存在的内存安全错误。请取消对 cudaMallocManaged 第二次调用的注释，以修改上述示例。在测试机器上，重新使用 compute-sanitizer 编译并运行程序后，我们现在会看到以下提示信息：
$ compute-sanitizer ./example1
========= COMPUTE-SANITIZER
Before: Array 0, 1 .. N-1: 1.000000 1.000000 1.000000
After : Array 0, 1 .. N-1: 3.000000 3.000000 3.000000
========= ERROR SUMMARY: 0 errors
没有错误？发生了什么？缓冲区溢出依然存在，但工具却不再报告问题。这一微小改动为何能掩盖缓冲区溢出？
简而言之，Compute Sanitizer 的 memcheck 工具用于确保内存访问位于有效的内存地址范围内，即已分配且未被释放的地址。在我们的测试过程中，CUDA 运行时的内存分配器将 array 和 buffer 在内存中连续放置。也就是说，buffer 紧随 array 之后，因此从 array 到 buffer 的溢出仍会访问合法的内存区域。需要注意的是，GCC 和 Clang 的“address sanitizer”工具同样存在这一局限性。
我们来了解一下为什么 Compute Sanitizer 会这样运行。到目前为止，Compute Sanitizer 完全依赖二进制插桩技术，通过对应用程序中的内存和同步指令插入检测代码，使工具能够在运行时发现异常行为。虽然我们对此过程进行了简化，但可以设想，工具能够在每次内存访问前插入一个复杂的断言，用以检查该地址是否位于当前已分配的对象范围内：
__global__ void scaleArray(float* array, size_t N, float value) {
int threadGlobalID = threadIdx.x + blockIdx.x * blockDim.x;
if (threadGlobalID <= N) {
assert(isAllocated(&array[threadGlobalID])); // Inserted instrumentation.
float tmp = array[threadGlobalID]*value;
assert(isAllocated(&array[threadGlobalID])); // Inserted instrumentation.
array[threadGlobalID] = tmp;
}
}
Compute Sanitizer 插入这些断言后，会运行应用程序，检测在运行时是否有任何断言被触发。
Compute Sanitizer 不具备静态分析程序的能力，部分原因在于它仅在原始二进制文件上运行，而原始二进制文件众所周知难以分析。由于缺乏整体分析能力，它只能以窥视孔的方式独立检查每个内存引用。在我们正在运行的示例中，即使某些线程从 array 溢出到 buffer，窥视孔检测机制仍会因 buffer 已被分配而忽略该非法访问。因此，`isAllocated(&array[threadGlobalID]` 的值为 true，且不会触发断言。
Compute Sanitizer 选项 `--padding’` 可在内存分配之间添加无效区域，从而使 Compute Sanitizer 能够检测此示例中的溢出问题。但该选项默认处于禁用状态，因为它会显著增加内存使用量。此外，该选项无法对所有类型的内存分配进行填充，例如共享内存和全局变量的分配。
我们可以利用编译时插桩来实现更好的效果。
使用编译器分析提高覆盖率
从 CUDA 13.1 开始，Compute Sanitizer 能够结合编译时分析与插桩技术以提升覆盖率。这种方法可以在不增加误报（即将非错误误判为错误）的情况下，有效减少漏报（即未能检测出的实际错误）。
我们的分析目标是迫切地将 CUDA 指针转换为所谓的“胖指针”，从而将指针与其对应的基数和边界信息绑定在一起。在胖指针上进行指针运算时，仅会修改指针部分，而基数和边界信息保持不变。这样一来，即使发生溢出并波及相邻的内存分配区域，我们仍有较大可能检测到该溢出行为。这项技术无法在本篇博客中详尽阐述，但您可在此处查阅完整的相关内容。
如果忽略编译时间，清理的开销也应该更低，在某些情况下，几乎可以减少一个数量级。也就是说，编译时插桩的一个考虑因素在于，它要求我们花费时间重新编译代码。
使用编译时插桩进行清理包含两个步骤：
- 使用
-fdevice-sanitize=memcheck标志编译程序。此步骤会插入仪器代码，用于在运行时验证程序的逻辑。例如：nvcc -fdevice-sanitize=memcheck -arch sm_86 -o example1 example1.cu.。 - 在 Compute Sanitizer 运行时环境下执行程序。此步骤会在运行时检测内存安全问题。例如：
compute-sanitizer ./example1。
虽然编译时仪器会执行静态程序分析，但必须注意，清理是一个动态过程。只有当被测试的应用程序恰好执行到存在错误的代码部分时，才能发现错误。
我们来运行之前未使用 Compute Sanitizer 的示例程序，但这次是在通过编译时插桩重新编译之后。首先，编译该程序。
nvcc -o example1 -arch sm_86 -fdevice-sanitize=memcheck example1.cu
现在，让我们使用 Compute Sanitizer 运行程序。
compute-sanitizer ./example1
========= COMPUTE-SANITIZER
Before: Array 0, 1 .. N-1: 1.000000 1.000000 1.000000
========= Invalid __global__ read of size 4 bytes
========= at scaleArray(float *, unsigned long, float)+0x150
========= by thread (512,0,0) in block (0,0,0)
========= Access to 0x7b3566000800 is out of bounds
========= and is inside the nearest allocation at 0x7b3566000800 of size 2,048 bytes
========= Saved host backtrace up to driver entry point at kernel launch time
========= Host Frame: main [0x8e01] in example1
=========
========= Program hit cudaErrorLaunchFailure (error 719) due to "unspecified launch failure" on CUDA API call to cudaDeviceSynchronize.
========= Saved host backtrace up to driver entry point at error
========= Host Frame: main [0x8e06] in example1
=========
After : Array 0, 1 .. N-1: 1.000000 1.000000 1.000000
example1: example1.cu:27: int main(): Assertion `array[N/2] == 3.0' failed.
========= Error: process didn't terminate successfully
========= Target application returned an error
========= ERROR SUMMARY: 2 errors
这次我们发现了这个错误。编译器的分析限制了对指向核函数启动时所用数组对象的数组指针的所有访问，从而防止溢出到缓冲区对象。
需要预先了解的说明和锐边
使用基于编译器的插桩技术时，需要注意一些问题。
仪器将使用额外的资源
使用“ 编译的代码仅用于调试目的，不适合部署。您的内核会膨胀，占用额外的寄存器，甚至可能使用部分堆栈。寄存器的额外占用可能导致某些内核无法启动，并引发“请求启动的资源过多”的错误提示。为解决此问题，建议通过 __launch_bounds__ 装饰或 maxregcount 编译器标志 明确限制资源使用。
虽然不太可能（我们尚未观察到），但在具有深度调用堆栈的程序中，插桩代码所增加的堆栈使用量可能会导致堆栈溢出。
使用 -G 在调试模式下编译代码会加剧上述所有问题。我们建议使用 -lineinfo，以确保 Compute Sanitizer 提供准确的属性，但应避免启用其他调试选项。
未定义的行为可能导致意外的代码，以及误报和漏报
未定义行为（UB）并非编译时工具独有的问题，也是一个颇具争议的话题，因此请勿击杀信使：如果您的代码中存在 UB，编译器将不再承担生成预期代码的责任。若将未定义行为转化为结果“67”，许多编译器开发者仍能安然入睡。遗憾的是，在 CUDA 中，编写出具有未定义行为的代码非常容易。
我们已经看到，由于程序员利用未定义行为构建指针，编译器可能会跳过某些内存操作。最近出现的一个例子涉及“假阳性”问题，其根源在于指针运算先对对象进行下采样，直到解引用前才恢复至正确边界。计算越界指针的类型（即使从未对其进行解引用）同样属于未定义行为。即便您的应用程序当前看似运行正常，但在下一次编译器更新对您的代码做出“创造性”优化之前，消除未定义行为仍是明智之举。
比赛条件可能会导致 UB
比赛条件从技术上讲属于上一节，因为在 CUDA C++ 中，比赛条件属于未定义行为（UB），但这些条件在 CUDA 中十分常见，因此值得单独设立一节进行讨论。若缺乏正确的同步机制，编译器无需保持预期的线程执行顺序和内存顺序。
当我们通过添加大量额外的逻辑和内存请求来分析代码时，会改变线程调度。如果您的代码存在潜在的竞争条件（具体取决于竞争的性质），则极有可能导致应用出现意外行为。若发现程序挂起、崩溃或内存损坏，请使用 Compute Sanitizer 中的 racecheck 工具 检查您的代码。
不支持 HMM 内存分配
目前不支持 HMM 显存。我们的工具现阶段仅跟踪通过 CUDA 运行时执行的内存分配。若您的应用使用 HMM，新工具可能会产生较多误报。我们计划在后续版本中解决此问题。
总结
最后，我们想传达两个要点。其一，我们鼓励 CUDA 开发者通过 Compute Sanitizer 运行其应用。CUDA C++ 赋予开发者诸多责任。我们持续借助 Compute Sanitizer 在专业编写的代码中发现错误。即便您的代码“在您的机器上能够运行”，仍可能存在一两处错误。
其次，我们的编译时方法大幅提升了运行时 Compute Sanitizer 的错误覆盖率与性能。如果您倾向于重新编译代码库，不妨尝试一下。
