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 統一記憶體也不代表應將大量主機記憶體全部鎖頁，批次容量仍應以實際工作集為限。

生產事件只涵蓋它之前的操作，並未涵蓋消費串流的加一與回傳。因此，生產事件完成後，可以確認生產階段已結束，卻不能確認裝置緩衝區已可讓下一批資料覆寫。擴充為管線時，應在消費串流的回傳操作之後，再記錄完成事件，以此控制整組緩衝區的回收。

每個並行批次槽位都應有獨立的輸入、輸出、裝置緩衝區與完成狀態。先以單一槽位驗收，再增加槽位；不要只增加串流數量，卻繼續在沒有保護的情況下共用同一個指標。

五、排查時逐步縮小同步範圍

結果偶發錯誤時，先核對等待是否位於消費核心函式之前、事件是否記錄在生產核心函式之後，以及記錄與等待使用的是否為同一個事件。事件的記錄位置只要提早一行，就可能把「資料已寫完」錯誤地縮減為「複製已完成」。

接著檢查每次核心函式啟動後的錯誤，以及最終同步的回傳值。前者能發現部分啟動問題，後者用於接收非同步執行期間的錯誤；兩者不能互相取代。若只檢查啟動回傳而不等待完成，錯誤可能到後續的介面呼叫才出現。

在供診斷使用的程式副本中，可暫時於生產階段加入串流同步，判斷問題是否與跨串流順序有關。若錯誤消失，應繼續檢查相依關係與緩衝區所有權，而不是直接以全裝置同步收尾。全域同步會改變排程，並可能掩蓋原有的生命週期問題。

最終保留的應是明確的生產事件、消費等待與回收事件，以及可重複執行的輸出檢查。唯有相依關係與記憶體歸屬都能解釋清楚，多串流程式才具備繼續擴充並行度的基礎。

常見問題