DGX Spark 的 CUDA 編譯產物驗收：目標架構、SASS 與 PTX 備援

CUDA 專案在 DGX Spark 上編譯成功，啟動核心時卻出現 no kernel image is available for execution on the device，首先應檢查編譯產物涵蓋的 GPU 架構，而不是重裝整套環境。另一種更隱晦的情況是，程式一直依賴 PTX 即時編譯執行，發佈紀錄卻誤認為產物已包含原生裝置程式碼。本文透過一個獨立的小程式，將目標查詢、產物檢查與兩條執行路徑串成一套驗收流程。

一、區分主機架構與 GPU 目標

DGX Spark 採用 NVIDIA GB10 Grace Blackwell，主機架構為 ARM64，系統為 NVIDIA DGX OS，以 Ubuntu 24.04 為基礎，並預先安裝 CUDA 軟體堆疊。付款後可在控制台查看連線憑證，透過 SSH 操作即可。

ARM64 描述的是 CPU 端執行檔的架構；CUDA 的運算能力則描述 GPU 裝置程式碼的目標，兩者並不是同一個相容性條件。file 顯示 AArch64，只能說明主機端格式正確，無法證明其中嵌入了可用的核心。

SASS 是針對目標 GPU 的機器碼，PTX 則是可由驅動程式即時編譯的中介表示。發佈套件可以同時包含兩者，但必須分別檢查，不能將一般執行成功視為兩條路徑都已驗證。

以下實驗只會在新建目錄中產生原始碼與執行檔，不會修改驅動程式、系統工具鏈或業務程式。範例結果僅用於相容性驗收，不代表效能測試。

二、查詢實際運算能力，不猜測架構編號

先建立獨立的工作目錄，保存系統與工具鏈資訊。nvidia-smi 顯示的 CUDA 版本反映驅動程式的支援能力，不等於目前 nvcc 所屬的工具套件版本，因此兩份輸出都要保留。

mkdir -p "$HOME/cuda-artifact-check"
cd "$HOME/cuda-artifact-check"
uname -m > host-arch.txt
nvidia-smi > driver.txt
nvcc --version > toolkit.txt
nvcc --list-gpu-code > supported-targets.txt

將以下內容儲存為 probe.cu。程式透過 CUDA Runtime 查詢裝置零的運算能力；若查詢失敗，會回傳非零結束碼，不會輸出可能遭誤用的目標編號。

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

int main() {
    cudaDeviceProp p{};
    cudaError_t e = cudaGetDeviceProperties(&p, 0);
    if (e != cudaSuccess) {
        std::fprintf(stderr, "%s\n", cudaGetErrorString(e));
        return 1;
    }
    std::printf("%d%d\n", p.major, p.minor);
    return 0;
}

nvcc probe.cu -o probe
CC=$(./probe) || exit 1
printf '%s\n' "$CC" > compute-capability.txt
grep -Fx "sm_${CC}" supported-targets.txt || exit 1

這裡刻意不將 GB10 的目標編號寫死。若工具套件的清單不包含查詢所得的 sm_ 目標，應先確認工具套件的支援範圍，不能任意改用較舊的編號繼續發佈。若探測程式失敗，應先處理裝置存取或執行階段函式庫的問題，不要進入編譯階段。

三、明確產生並檢查 SASS 與 PTX

將以下內容儲存為 check.cu。它只會啟動一個執行緒寫入整數，再將結果複製回主機核對。每次 CUDA 呼叫都會檢查錯誤，避免核心啟動失敗後仍顯示成功。

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

#define CUDA_OK(call) do { \
    cudaError_t e = (call); \
    if (e != cudaSuccess) { \
        std::fprintf(stderr, "%s\n", cudaGetErrorString(e)); \
        return 1; \
    } \
} while (0)

__global__ void write_value(int* p) {
    *p = 7;
}

int main() {
    int* device = nullptr;
    int host = 0;
    CUDA_OK(cudaMalloc(&device, sizeof(int)));
    write_value<<<1, 1>>>(device);
    CUDA_OK(cudaGetLastError());
    CUDA_OK(cudaDeviceSynchronize());
    CUDA_OK(cudaMemcpy(&host, device, sizeof(int),
                       cudaMemcpyDeviceToHost));
    CUDA_OK(cudaFree(device));
    if (host != 7) return 2;
    std::puts("PASS");
    return 0;
}

在同一個工作目錄中執行，明確指定同時產生原生程式碼與 PTX：

CC=$(cat compute-capability.txt)
nvcc -O2 \
  -gencode "arch=compute_${CC},code=[sm_${CC},compute_${CC}]" \
  check.cu -o check

file check > binary-format.txt
cuobjdump --list-elf ./check > elf-list.txt
cuobjdump --dump-sass ./check > sass.txt
cuobjdump --dump-ptx ./check > ptx.txt

檢查 ELF 清單與 SASS 輸出中是否包含對應的目標及核心，以及 PTX 輸出是否包含 .version、.target 與核心進入點。只要出現空檔案、目標不符或缺少預期核心的情況，就不應通過驗收；僅檢查命令的結束碼並不足夠。

檢查最終檔案，而非中間目的檔

實際專案可能經過靜態函式庫連結、裝置連結或封裝步驟。即使目的檔中存在 PTX，也不能證明最終交付檔案仍保留了它。應針對最終執行檔及實際載入的共用函式庫重複檢查，並記錄檔案雜湊值。

四、分別驗證原生執行與即時編譯

先停用 PTX 即時編譯，要求目前產物依靠原生裝置程式碼執行；接著強制使用 PTX 路徑，並停用磁碟編譯快取。兩次測試都使用新行程，避免行程內已載入的模組影響判斷。

env -u CUDA_FORCE_PTX_JIT \
  CUDA_DISABLE_PTX_JIT=1 \
  ./check > native.log 2>&1
native_rc=$?

env -u CUDA_DISABLE_PTX_JIT \
  CUDA_FORCE_PTX_JIT=1 CUDA_CACHE_DISABLE=1 \
  ./check > ptx-jit.log 2>&1
ptx_rc=$?

printf 'native=%s\nptx_jit=%s\n' \
  "$native_rc" "$ptx_rc" > execution-status.txt

cat execution-status.txt native.log ptx-jit.log
test "$native_rc" -eq 0 && test "$ptx_rc" -eq 0

兩份日誌都應出現 PASS，結束碼也都應為零。測試變數僅對該次命令生效，不要寫入全域環境，讓業務行程長期被強制使用 JIT。

若原生執行失敗但 PTX 成功，應優先回查原生目標與最終連結步驟；若原生執行成功但 PTX 失敗，則檢查 PTX 是否存在，以及驅動程式是否支援其版本。若兩者都失敗，應根據具體錯誤檢查工具鏈、執行階段函式庫與裝置狀態，不要以反覆切換編譯選項取代問題定位。

五、將實驗結果納入發佈檢核

小程式通過測試，只能證明基本編譯鏈與這段核心在目前環境中可用。業務程式還可能載入擴充函式庫、依輸入選擇不同核心，或在執行階段產生程式碼，因此必須補上真實進入點的測試，涵蓋主要資料型別與分支。

建議每次發佈保留四組佐證資料：

• 主機架構、驅動程式輸出、工具套件版本與實際編譯命令。
• GPU 運算能力，以及最終檔案的 SASS 與 PTX 檢查結果。
• 兩條執行路徑的結束碼、日誌與業務結果核對紀錄。
• 原始碼版本、最終檔案的 SHA-256，以及動態相依項目清單。

sha256sum ./check 可明確識別本次驗收的對象；後續若重新連結，即使檔名不變，也需要重新驗收。若使用容器，應在實際交付的映像檔內檢查工具鏈與最終檔案，並在實際執行環境中完成執行驗證。

PTX 是相容性備援路徑，並不保證支援所有驅動程式，更不能由一次成功推論出其他裝置也能執行。發佈檢核條件應明確寫成「目前驅動程式、目前產物、已涵蓋進入點通過」，架構或工具鏈變更後就要重新產生佐證資料，而不是沿用舊日誌。

常見問題