DGX Spark 的 CUDA 多流依赖：事件排序与缓冲区验收

两个 CUDA 流共用同一块设备缓冲区时，程序可能连续运行多次都正确，换一个输入规模后却出现旧值。常见原因不是算子公式，而是生产流尚未写完，消费流已经开始读取。本文在 DGX Spark 上建立双流最小实验，把提交顺序、执行依赖和缓冲区复用时机分开验收，不以一次运行成功作为结论。

一、先画出数据依赖，再创建流

DGX Spark 使用 NVIDIA GB10 Grace Blackwell，提供 128GB LPDDR5x 统一内存，CPU 与 GPU 共用同一内存池。但物理上的统一内存并不会替应用自动建立跨流执行顺序。显式复制、内核启动与主机访问，仍须遵守各自的完成条件。

实验采用显式设备分配及页锁定主机内存，便于观察边界。生产流执行“输入复制到设备、每个元素乘二”，消费流执行“每个元素加一、结果复制回主机”。完整依赖是：

输入就绪 → 复制完成 → 乘二完成 → 加一开始 → 回传完成 → 主机检查。

同一流内的操作按流顺序执行；两个流之间则必须显式连接。事件应记录在生产流的乘二内核之后，消费流应先等待事件，再提交加一内核。

主机先调用生产流、后调用消费流，只说明提交顺序。跨流数据依赖必须有明确的同步依据，不能依赖负载较轻时碰巧成立的执行顺序。

二、建立不依赖默认流的最小程序

在 TcpNode 付款后在控制台查看连接凭据，经 SSH 进入 NVIDIA DGX OS。先执行 uname -m、nvidia-smi 和 nvcc --version，分别保存主机架构、设备与驱动信息、编译工具链信息。架构应为 ARM64 对应的 aarch64；编译器版本不应从驱动输出推断。

以下内容保存为 stream_contract.cu。两个流均使用非阻塞标志，避免传统默认流的隐式同步掩盖缺失依赖。事件只承担排序，不采集耗时。

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

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

__global__ void scale(float* data, int count) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < count) data[i] *= 2.0f;
}

__global__ void add(float* data, int count) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < count) data[i] += 1.0f;
}

int main() {
    constexpr int count = 4096;
    const size_t bytes = count * sizeof(float);
    float *input = nullptr, *output = nullptr, *device = nullptr;
    cudaStream_t producer, consumer;
    cudaEvent_t ready;

    check(cudaMallocHost(reinterpret_cast<void**>(&input), bytes));
    check(cudaMallocHost(reinterpret_cast<void**>(&output), bytes));
    check(cudaMalloc(reinterpret_cast<void**>(&device), bytes));
    check(cudaStreamCreateWithFlags(&producer, cudaStreamNonBlocking));
    check(cudaStreamCreateWithFlags(&consumer, cudaStreamNonBlocking));
    check(cudaEventCreateWithFlags(&ready, cudaEventDisableTiming));

    for (int i = 0; i < count; ++i) {
        input[i] = static_cast<float>(i);
        output[i] = -1.0f;
    }

    check(cudaMemcpyAsync(device, input, bytes,
                          cudaMemcpyHostToDevice, producer));
    scale<<<(count + 255) / 256, 256, 0, producer>>>(device, count);
    check(cudaGetLastError());
    check(cudaEventRecord(ready, producer));

    check(cudaStreamWaitEvent(consumer, ready, 0));
    add<<<(count + 255) / 256, 256, 0, consumer>>>(device, count);
    check(cudaGetLastError());
    check(cudaMemcpyAsync(output, device, bytes,
                          cudaMemcpyDeviceToHost, consumer));
    check(cudaStreamSynchronize(consumer));

    int failures = 0;
    for (int i = 0; i < count; ++i) {
        float expected = 2.0f * static_cast<float>(i) + 1.0f;
        if (output[i] != expected) ++failures;
    }
    std::printf("elements=%d mismatches=%d\n", count, failures);

    check(cudaEventDestroy(ready));
    check(cudaStreamDestroy(consumer));
    check(cudaStreamDestroy(producer));
    check(cudaFree(device));
    check(cudaFreeHost(output));
    check(cudaFreeHost(input));
    return failures == 0 ? 0 : 2;
}

三、用输出与退出码共同验收

在同一目录编译并执行：

nvcc -O2 stream_contract.cu -o stream_contract
./stream_contract
printf 'exit=%s\n' "$?"

合格结果是 elements=4096 mismatches=0，退出码为零。接口或同步失败返回一，数值不符返回二。仅检查进程没有崩溃，会漏掉已经完成但内容错误的计算。

示例数值均处于单精度可精确表示的整数范围，因此直接比较相等。实际浮点模型不能照搬这一判据，应依据算法规定绝对误差、相对误差及非有限值处理规则。

接着分别将元素数改为 1、257、4097，重新编译运行。这些规模覆盖单元素、线程块尾部和多个块的边界；它们是测试输入，不是性能数据。每次保存源文件、编译命令、工具链信息与运行输出，避免只保留一张成功截图。

检查失败路径是否真的生效

临时把期望式末尾的 1.0f 改为 2.0f，程序应报告不匹配并返回二，然后恢复原式。这验证结果检查没有失效。不要用“删除等待后必须报错”作为反向测试：竞态是否出现取决于调度，缺失依赖的程序也可能暂时通过。

四、把事件边界落实到内存生命周期

cudaMemcpyAsync 返回，通常只能说明调用已经提交，不能据此改写输入。示例中的输入从提交复制到消费流同步结束一直保持不变；输出也只在回传所在流同步后读取。这个策略偏保守，但边界清楚，适合作为初始正确性基线。

页锁定内存有助于满足异步传输条件，却不是完成通知。设备缓冲区同样不能在其他流仍访问时提前释放。128GB 统一内存也不意味着应把大量主机内存全部锁页，批次容量应受实际工作集约束。

生产事件只覆盖它之前的操作，并未覆盖消费流的加一与回传。因此，生产事件完成后可以确认生产阶段结束，却不能确认设备缓冲区已经可供下一批覆盖。扩展为流水线时，应在消费流回传之后再记录完成事件，以此控制整组缓冲区的回收。

每个并行批次槽位应有独立的输入、输出、设备缓冲区和完成状态。先用单槽位验收，再增加槽位；不要只增加流数量，却继续无保护地共用同一个指针。

五、排查时逐步缩小同步范围

结果偶发错误时，先核对等待是否位于消费内核之前、事件是否记录在生产内核之后，以及记录和等待使用的是否为同一个事件。事件记录位置提前一行，就可能把“数据已写完”错误地缩减成“复制已完成”。

随后检查每个内核启动后的错误，以及最终同步返回值。前者能发现部分启动问题，后者用于接收异步执行中的错误；两者不能互相替代。若只检查启动返回而不等待完成，错误可能在更晚的接口调用中才暴露。

诊断副本中可暂时在生产阶段加入流同步，判断问题是否与跨流顺序有关。若错误消失，应继续检查依赖和缓冲区所有权，而不是直接用全设备同步收尾。全局同步会改变调度，并可能隐藏原有生命周期问题。

最终保留的应是明确的生产事件、消费等待和回收事件，以及可重复执行的输出检查。只有依赖边和内存归属均能解释清楚，多流程序才具备继续扩展并发的基础。

常见问题