DGX Spark 上如何定位 CUDA 越界访问并验证修复

CUDA 内核已经启动，错误却在结果回传时才出现；同一段程序缩小输入后又能运行。这类现象首先需要检查设备访存，而不是立即调整驱动。在 DGX Spark 上，可以用一个独立的小程序建立诊断链：主动触发越界、定位源码行、修复索引，再确认错误检测没有被关闭。整个过程不需要模型权重，也不涉及线上服务。

一、固定环境并划清测试边界

TcpNode 的 DGX Spark 使用 NVIDIA GB10 Grace Blackwell、ARM64 架构与 NVIDIA DGX OS，配备 128GB LPDDR5x 统一内存。统一内存由 CPU 与 GPU 共用，不意味着任意指针都能跨越分配边界。下面使用 cudaMalloc 分配设备内存，使合法地址范围明确。

付款后在控制台查看连接凭据，经 SSH 登录。测试应放在独立工作目录，以普通用户运行，不在正在承担请求的推理进程中注入错误。

mkdir -p "$HOME/cuda-access-check"
cd "$HOME/cuda-access-check"
uname -m
nvidia-smi
command -v nvcc
command -v compute-sanitizer
nvcc --version
compute-sanitizer --version

架构应显示 aarch64。记录工具版本与驱动信息，后续复验使用同一环境。CUDA 栈预装并不等于每个工具都已进入当前 PATH；命令不存在时，先核对已安装工具链与安装路径，不要直接替换整套驱动。

越界样例只能在可丢弃的独立进程中执行。出现非法访存后，应结束该进程重新运行，不能把同一 CUDA 上下文里的后续结果当作有效证据。

二、构造只有一个变量的越界样例

保存以下内容为 repro.cu。它只分配一个整数，合法索引为零；编译时通过 TEST_INDEX 切换写入位置。这样错误版与修复版共享源码，避免同时改变输入、分配大小和执行路径。

#include <cuda_runtime.h>
#include <cstdio>
#include <cstdlib>

static void checked(cudaError_t status) {
    if (status != cudaSuccess) {
        std::fprintf(stderr, "%s\n", cudaGetErrorString(status));
        std::exit(EXIT_FAILURE);
    }
}

__global__ void write_value(int* data) {
    data[TEST_INDEX] = 7;
}

int main() {
    int* data = nullptr;
    int result = 0;
    checked(cudaMalloc(reinterpret_cast<void**>(&data), sizeof(int)));
    checked(cudaMemset(data, 0, sizeof(int)));
    write_value<<<1, 1>>>(data);
    checked(cudaGetLastError());
    checked(cudaDeviceSynchronize());
    checked(cudaMemcpy(&result, data, sizeof(int), cudaMemcpyDeviceToHost));
    checked(cudaFree(data));
    std::printf("result=%d\n", result);
    return result == 7 ? EXIT_SUCCESS : EXIT_FAILURE;
}

分别编译两个程序：

nvcc -arch=native -lineinfo -DTEST_INDEX=1 repro.cu -o repro-bad
nvcc -arch=native -lineinfo -DTEST_INDEX=0 repro.cu -o repro-good
sha256sum repro.cu repro-bad repro-good > build.sha256

-arch=native 面向当前可见 GPU 编译，适用于本机验证，不用于生成跨设备通用产物。若编译器不支持当前设备，先核对工具链兼容性，不要随意套用其他架构参数。-lineinfo 保留设备源码行信息；首次定位不必启用会明显改变执行特征的完整设备调试模式。

为什么同时检查启动与同步

cudaGetLastError 检查启动相关错误，但内核通常异步执行。cudaDeviceSynchronize 等待前面的设备工作完成，执行期间的错误才有机会在这里返回。两者必须分别检查，不能只看程序是否打印了结果。

三、让检测结果成为可判断的退出码

先运行错误版，并分别保存标准输出与检测日志：

if compute-sanitizer --tool memcheck --error-exitcode 86 \
    ./repro-bad > bad.stdout 2> bad.stderr; then
    status=0
else
    status=$?
fi
printf 'exit=%s\n' "$status"
cat bad.stdout bad.stderr

预期能看到非法全局内存写入报告，并定位到 write_value 中的写操作。报告通常还包含访问大小、线程位置和地址关系，具体排版随工具版本变化，不应把某段固定文字作为唯一判断条件。

--error-exitcode 86 用于让检测到的错误影响自动化结果；若被测程序自身已经失败，也可能得到程序返回的其他非零值。因此，验收要求是“退出非零且报告指向预期越界”，而不是强制所有故障都返回 86。

不要把命令直接接到普通日志管道后，只读取最后一个日志命令的状态。上面的重定向与条件分支可以保留真正的运行结果，也适用于启用了失败即退出策略的脚本。

四、根据地址范围修复，不掩盖症状

样例的分配长度为一个整数，写索引却为一，已经越过合法范围。真实程序需要把报告中的地址与对应分配关联，再检查索引计算、元素类型、步长及批次偏移。

常见问题包括：分配参数使用字节数，索引却按元素数解释；最后一个线程块不足整块；二维数组的行跨度与逻辑宽度不同；释放后的指针又被另一个任务使用。扩大分配可能暂时隐藏错误，却不能证明算法正确。

对一维任务，通常应在访问前检查 index < count，同时确认索引计算使用的整数类型能够容纳乘积。对共享内存协作内核，不能随意让部分线程提前返回，再让剩余线程进入块级屏障，否则可能引入新的同步问题。

若错误只出现在大型工作流，可暂时在可疑阶段后增加同步检查，逐段缩小范围。一次只增加一个检查点，保留首个错误；不要在错误发生后继续提交任务，产生大量次生报错。

五、用正反向测试完成修复验收

运行修复版：

if compute-sanitizer --tool memcheck --error-exitcode 86 \
    ./repro-good > good.stdout 2> good.stderr; then
    status=0
else
    status=$?
fi
printf 'exit=%s\n' "$status"
cat good.stdout good.stderr

本例应同时满足三个条件：退出码为零、检测报告没有访存错误、输出为 result=7。只看“没有报错”不够，内核没有执行或结果没有真正写入，也可能让测试表面通过。

随后重新运行错误版，确认它依然失败。这项反向测试能够发现检测命令被删掉、测试误用了旧二进制，或流水线吞掉失败状态等问题。真实内核还应覆盖最小有效输入、整块边界与不足整块的尾部输入；若接口允许空输入，也要明确其处理方式。

最终归档源码、编译命令、版本信息、校验和及两次运行日志。检测器会改变执行开销，检测期间的耗时不能直接作为性能结论；memcheck 通过也不代表数据竞争和同步问题已经排除。可靠的收尾是让访问合法性、计算结果与失败传播分别拥有可复查的证据，而不是仅留下一个成功截图。

常见问题