昨天已經用 ptxas 把 PTX 組成 SM121 binary。今天剩下最後兩個問題：cubin 裡到底放了什麼？Triton runtime 又怎麼把它載入並送到 GB10？這次會一起用 ELF 工具、cuobjdump、nvdisasm、Triton 本機原始程式碼與 Nsight Systems 把證據補齊。

想看完整檔案可以看 matmul_kernel.cubin、cuobjdump --dump-sass 輸出 與 nvdisasm -g 輸出。可以用位址或 instruction mnemonic 搜尋這篇節錄的 SASS。

本篇大綱

• 對照 PTX 與 SASS，看 ptxas 具體改了什麼。
• 對照 cubin、CUmodule、CUfunction 與 launch configuration。
• 把 cubin 當作 ELF device binary 檢視。
• 從 resource metadata 分清 register、static shared 與 dynamic shared。
• 對照 PTX 與 SASS machine instruction。
• 追 cuModuleLoadData、cuModuleGetFunction、cuLaunchKernelEx。
• 用 Nsight Systems 驗證實際 grid、block 與 shared memory。

.ptx → .cubin / SASS：ptxas

昨天的 make_cubin() 呼叫 ptxas。這段呼叫可在 Triton 3.6.0 NVIDIA backend compiler.py 第 461～535 行直接看到，其中第 489～494 行組出並執行 ptxas_cmd。輸入 PTX 是 virtual ISA，輸出 cubin 則包含 SM121 可執行的 machine code 與 resource metadata。

這裡會做什麼?

• 配置 physical register，把 %r163 這類 virtual register 改成 R52 等實體編號。
• 選擇對應 SM121 的 machine instruction，例如 mma.sync 變成 HMMA。
• 安排 instruction 並編碼 dependency control，所以 PTX 順序不一定原樣保留。
• 產生 .text、.nv.info、constant/shared-memory section 與 symbol table。
• 若 register 不足，決定 spill 與 local-memory access；本次結果是 80 registers、0 spill。

用本次 artifact 可做三組直接對照

表中的對應是 instruction family 與資料流對應，不代表每一行 PTX 都只生成一行 SASS。

Compiler pipeline 在 cubin 結束，runtime 從這裡開始

cuModuleLoadData、cuModuleGetFunction 和 cuLaunchKernelEx 都是 runtime Driver API call，不是 compiler pass。Triton 3.6.0 的 loadBinary() 在 driver.c 第 102～163 行，其中第 121～130 行處理 CUDA context、cuModuleLoadData() 與 cuModuleGetFunction()。實際路徑可拆成

cubin bytes
  → driver.c: loadBinary()
  → cuCtxGetCurrent()
  → 必要時 cuDevicePrimaryCtxRetain() / cuCtxSetCurrent()
  → cuModuleLoadData()
  → cuModuleGetFunction()
  → cuFuncGetAttribute()
  → CUmodule + CUfunction + resource metadata

呼叫 kernel 時，Python driver 產生的 launcher stub 會填好 CUlaunchConfig。Triton 3.6.0 的launcher C stub 在 driver.py 第 295～376 行；CUlaunchConfig 的 grid、block、shared memory 與 stream 設定在第 325～334 行，第 376 行呼叫 cuLaunchKernelExHandle()：

gridDim = grid
blockDimX = 32 * num_warps
sharedMemBytes = metadata["shared"]
hStream = current CUDA stream
attrs = cluster / cooperative / PDL attributes（需要時）

最後呼叫

cuLaunchKernelExHandle(&config, function, params, 0);

這段對照也說明了 artifact 的界線，SASS 能告訴我們 GPU 會執行哪些 machine instruction，卻不包含這次 launch 的 grid、stream 或 dynamic shared-memory bytes，這些是 runtime 每次呼叫時才填入的資訊。

cubin 是 container，SASS 是其中的 machine code

用 file 與 readelf 檢查後，可以確認 cubin 是

ELF 64-bit LSB executable, NVIDIA CUDA architecture

和 kernel 最直接相關的 section 包括

.text.matmul_kernel
.nv.info.matmul_kernel
.nv.constant0.matmul_kernel
.nv.shared.matmul_kernel

Symbol table 中也有全域 FUNC symbol

matmul_kernel

CUDA Driver 後面會用相同 symbol name 取得 CUfunction

Resource usage：三種 shared-memory 數字

用 cuobjdump --dump-resource-usage 實際輸出重點如下

REG:80
STACK:0
SHARED:1024
LOCAL:0
CONSTANT[0]:936

REG:80和昨天的 ptxas -v 相符，SHARED:1024 卻和 Triton metadata 的 8192 不同，原因是它們不是同一種配置

• cubin resource report 的 1024 bytes 對應 .nv.shared.matmul_kernel，是這個 kernel 的 static shared-memory 配置
• binary 另有 64-byte .nv.shared.reserved.0 section，不包含在前一項 1024 bytes 的解讀裡
• Triton metadata 的 8192 bytes 是 launch 時傳入的 dynamic shared memory
• Nsight 的 kernel record 會看到 dynamic shared memory = 8192 bytes

因此看到兩個 shared-memory 數字時，不應直接判定工具互相矛盾，先確認它們報的是 static、dynamic，還是總量。

反組譯 SASS

cuobjdump --dump-sass kernel.cubin
nvdisasm kernel.cubin

Header 顯示目標為 sm_121a，並帶有

EF_CUDA_SM121
EF_CUDA_ACCELERATORS

例如第一條 SASS 前面的 /*0010*/，代表這條 machine instruction 位於 code section offset 0x10。後面的 R7、R73 已是 physical register。先看 kernel 如何取得 block/thread id

/*0010*/ S2R R73, SR_TID.X;
/*0050*/ S2R R7,  SR_CTAID.X;

S2R 是 special-register to register。PTX 的 %tid.x、%ctaid.x 到這裡已經成為實際機器指令。

再看資料搬移與 matrix load

/*0360*/ LDGSTS.E.BYPASS.128
           [R68], desc[UR6][R8.64], !P0;
/*0530*/ LDSM.16.M88.4 R12, [R2+UR4];
/*0540*/ LDSM.16.MT88.4 R32, [R0+UR4+0x1000];

LDGSTS 對應 PTX cp.async，把 global memory 資料送到 shared memory。.128 表示 128-bit，也就是 16-byte transfer，!P0 是 predicate，延續前面的邊界條件。LDSM 對應 ldmatrix，MT88 中的 T 表示 transpose 形式。

最後是 Tensor Core machine instruction

/*05b0*/ HMMA.16816.F32 R52, R12, R32, RZ;
/*0630*/ HMMA.16816.F32 R52, R44, R34, R52;

第一條使用 RZ（zero register）作為 accumulator 輸入，相當於從 0 開始；第二條把 R52 當成舊 accumulator，再把新結果寫回 R52。16816 對應 PTX 的 m16n8k16，.F32 表示 accumulator/output 為 FP32。

從 SASS 可以建立以下對照：

這才是 GB10 最終執行的 machine instruction。TTGIR 的 tt.dot、LLIR/PTX 的 mma.sync，到 cubin 中成為 HMMA.16816.F32。

PTX 與 SASS 通常不是一對一逐行翻譯。ptxas 會配置 physical register、選擇 machine instruction、插入 dependency control 並重新排程。閱讀時應對照「instruction family 和資料流」，不要期待 virtual register 編號或指令順序原樣保留。

這個不會看到 TCGEN05，因為 SM121 target 不提供 tcgen05 / TMEM 指令。這和 Day16 的 nvidia_mma versionMajor=2、tmem_size=0 一致。到目前為止，四層證據指向同一條 lowering

TTIR:   tt.dot
TTGIR:  nvidia_mma v2，無 TMEM
LLIR:   mma.sync.m16n8k16
PTX:    mma.sync.m16n8k16
SASS:   HMMA.16816.F32

Runtime 如何載入 cubin

目前環境的 Triton NVIDIA driver.c 會先確認 current CUDA context。若還沒有，會取得 device primary context 並設成 current。載入 binary 的核心流程是

cuCtxGetCurrent(&ctx);
cuModuleLoadData(&mod, data);
cuModuleGetFunction(&fun, mod, name);
cuFuncGetAttribute(..., fun);

各物件的角色為

cubin bytes
  ↓ cuModuleLoadData
CUmodule
  ↓ cuModuleGetFunction("matmul_kernel")
CUfunction

Runtime 還會查詢 function 的 register、local memory 與最大 thread 數，並依 metadata 處理 dynamic shared memory opt-in。

這一段發生在 host CPU，cuModuleLoadData 接收的是 cubin bytes，不會執行 kernel，cuModuleGetFunction 只取得可供 launch 的 function handle。後面的 launch API 才會把工作排入 CUDA stream。

Launch configuration 怎麼形成

Triton 3.6.0 的 NVIDIA launcher 使用 cuLaunchKernelEx。這個例子的主要欄位是

gridDimX      = 16
gridDimY/Z    = 1
blockDimX     = 32 × num_warps = 128
sharedMemBytes= 8192
stream        = PyTorch current CUDA stream
function      = matmul_kernel 的 CUfunction

num_ctas=1，所以不需要 cluster launch attribute。A、B、C tensor 會轉成 device pointer；已 specialization 的 M/N/K、stride 與 block size 不必全部當作一般 runtime scalar 傳進 kernel。

把 Python launch 和 Driver 欄位對起來

這張表也說明哪些資訊不在 SASS 裡：grid size、stream 與 dynamic shared-memory bytes 都由每次 launch 決定。

用 Nsight Systems 看實際 launch

這裡使用 trace 來驗證 runtime 路徑

nsys profile --trace=cuda,nvtx,osrt -o matmul-trace \
  python matmul.py

這次 matmul_kernel 的 launch record 是

grid:                  16 × 1 × 1
block:                128 × 1 × 1
registers per thread:  80
dynamic shared memory: 8192 bytes
kernel duration:       30,688 ns

Grid、block、register 與 dynamic shared memory 分別和 Day14 metadata、昨天ptxas log 對上。Runtime trace 為 compiler 原始程式碼的推論補上執行證據。

這些欄位來自當次 Nsight Systems 實驗紀錄。Repository 目前沒有保存 .nsys-rep、SQLite 或匯出文字，因此 grid、block、register 與 shared-memory 數字可以用 metadata / cubin 交叉核對，30,688 ns 則無法由現有 artifact 獨立重驗。

30,688 ns 是 profiler 中這一次小矩陣 launch 的觀測值，包含 profiler 擾動與 cold-state 影響，不適合拿來宣稱 GB10 matmul 的代表效能。

Launch 返回不表示 GPU 已完成

cuLaunchKernelEx 一般只把工作排入 stream。Day14 程式隨後呼叫

torch.cuda.synchronize()

所以 correctness check 前能保證 kernel 完成。若用 CPU timer 只包住 launch call，量到的多半是 host enqueue latency；非同步執行錯誤也可能延後到 synchronization 才回報。

複習一下

我們這幾天用同一個 matmul 程式取得了可以互相驗證的證據

Python tl.dot
  → TTIR：64×32 與 32×64 的 tile-level dot
  → TTGIR：4 warps、8 KiB shared、NVIDIA MMA v2
  → LLIR：NVVM cp.async / ldmatrix / mma.sync
  → PTX：sm_121a、128 threads、mma.sync.m16n8k16
  → ptxas：80 registers、0 spill
  → cubin：SM121 ELF + HMMA SASS
  → Driver：CUmodule → CUfunction → cuLaunchKernelEx
  → Nsight：grid 16、block 128、dynamic shared 8192

閱讀 compiler pipeline 時，每一層都要能在下一層找到對應證據。若某個環節對不上，就能從那一層開始追 correctness 或 performance 問題。

明天我們會回顧這幾天 GPU 硬體、MLIR 和整個 Triton 編譯流程。

參考資料

• CUDA Binary Utilities
  https://docs.nvidia.com/cuda/cuda-binary-utilities/
• CUDA Driver API：Module Management
  https://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MODULE.html
• CUDA Driver API：Execution Control
  https://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__EXEC.html
• Triton NVIDIA driver
  https://github.com/triton-lang/triton/blob/v3.6.0/third_party/nvidia/backend/driver.c
• Triton NVIDIA Python driver
  https://github.com/triton-lang/triton/blob/v3.6.0/third_party/nvidia/backend/driver.py
• Nsight Systems User Guide
  https://docs.nvidia.com/nsight-systems/UserGuide/