在 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 也不代表已排除資料競爭與同步問題。可靠的收尾方式，是讓存取合法性、計算結果與失敗狀態傳遞各自都有可供複查的證據，而不只是留下一張成功的螢幕擷圖。

常見問題