今天來複習從 Day11 到 Day19，我們從 GPU 硬體出發，認識 CUDA 與幾種 kernel 開發介面，再沿著一個 Triton 矩陣乘法一路打開 TTIR、TTGIR、LLVM IR、PTX、cubin 與 SASS。最後，CUDA Driver 載入 binary 並把 kernel 送進 GPU。

今天要把這九天的內容接起來，從一串 artifact 整理出三個之後也能帶去其他 compiler 與 accelerator 的問題

1. 開發者用什麼抽象描述運算？
2. Compiler 在哪一站加入資料布局、硬體映射與目標指令資訊？
3. Binary 產生後，runtime 還要做哪些事，GPU 才會真的開始執行？

本篇大綱

• 先建立一個可套用到 GPU compiler 的五段式框架
• 複習這9天的內容和帶著走的知識

先用五段式框架定位

這五個位置不要求每套系統都使用相同名稱。例如 cuTile 的公開 IR 邊界是 CUDA Tile IR，Triton NVIDIA backend 可看到 TTIR、TTGIR、LLVM IR 與 PTX，換到 AMD 或其他 accelerator，最後的 ISA、binary 格式與 runtime API 也會改變。讀法仍然相同，找出每一站的輸入、輸出，以及這一站新增的資訊。

Day11：GPU 架構：從 SM、warp、memory hierarchy 看 NVIDIA GPU 系列

Day11 把 GPU kernel 拆成三個互相連動的角度

• 工作組織：grid、thread block／CTA、warp、thread／lane。
• 硬體執行：SM 接收 blocks，warp scheduler 安排可執行的 warps，再由 CUDA Cores 或 Tensor Cores 完成運算。
• 資料存取：global memory 經過 L2 進入 SM；SM 上有 shared memory／L1 與 register file。

這些硬體名詞會在後面的 compiler artifact 重新出現。TTGIR 的 layout 描述 element 如何分給 warp 與 lane，LLVM IR 的 address space 區分 global 與 shared memory；ptxas 的 register 配置影響同一個 SM 能同時容納多少工作，launch configuration 則帶入 grid、block 與 dynamic shared-memory 大小。

Day12：從 CUDA 到 Python DSL：NVIDIA GPU Kernel 的不同寫法.md

橫軸是開發者主要描述的單位，縱軸是開發者直接指定硬體細節的程度。CUDA C++ 可以明確安排 thread、memory space、同步與 inline PTX；CuTe DSL 則以 layout、copy、MMA 等結構化抽象控制硬體。兩者控制的面向不同，無法只靠一條由低到高的直線完整比較。

• cuBLAS／cuDNN：開發者提供 GEMM、attention、normalization 等既有 operation 與參數，函式庫內部實作與 heuristic 負責執行緒映射。
• CUDA C++：開發者描述每個 thread 的 scalar program、memory space、同步與 launch，並明確安排多數執行緒工作。
• CUTLASS C++：開發者組合 tile、layout、MMA、pipeline、epilogue 等元件，搭配 CUTLASS 已提供的實作。
• CuTe DSL：開發者在 Python 中操作 layout、tensor、copy、MMA、pipeline 與硬體階層，保有詳細的硬體控制。
• cuTile：開發者描述 logical block 內的 tile operations，執行緒映射主要交給 CUDA Tile compiler。
• Triton：開發者描述 program instance 內的 block of values 與 tiling strategy，Triton compiler 再映射到 threads／warps。

還有兩組容易混淆的 Python DSL

CuTe DSL：把 CUTLASS／CuTe 的硬體抽象帶進 Python

CuTe DSL 屬於 CUTLASS 的 Python DSL 家族。它保留 layout、tensor、copy atom、MMA atom 與 pipeline 等硬體導向抽象，適合需要直接設計資料搬移與 Tensor Core kernel 的工作。

官方文件把 code generation 拆成三個階段

Python AST preprocessing
→ Python interpreter 執行 meta-stage tracing
→ 產生 internal IR
→ backend lowering 與最佳化
→ PTX／SASS 與 device binary
→ runtime load / launch

Python 在 meta-stage 處理 metaprogramming、compile-time value 與控制流程，產生的 kernel 再交給 backend。這讓開發者避開大量 C++ template metaprogramming，同時保留對 layout、copy、MMA 與硬體執行階層的控制。

cuTile：先描述 tile operation

cuTile 是建立在 CUDA Tile programming model 上的 Python DSL。Kernel 在 logical grid of blocks 上執行；每個 block 操作 immutable tile values，編譯器再決定 block 內的 thread mapping、實體儲存位置與可使用的硬體能力。

Python @ct.kernel
→ 依實際參數 specialization
→ CUDA Tile IR／TileIR bytecode
→ tileiras + libNVVM + ptxas
→ cubin
→ CUDA launch

目前官方 quickstart 列出的支援範圍已包含 compute capability 8.x 到 12.x ，不再只適用最初推出時的 Blackwell。cuTile 也支援 AOT 匯出 TileIR bytecode 或指定 GPU 的 cubin。支援版本仍會隨 CUDA Toolkit 與 cuTile 套件更新，實際使用前應重新檢視 release notes。

Triton：用 program instance 描述 block of values

Triton kernel 由 grid 上的 program instances 組成。一個 program instance 通常處理一塊資料，使用 tl.program_id、tl.arange、tl.load、tl.store 與 tl.dot 表達 block-level 運算。開發者負責 grid、tiling strategy 與 meta-parameters；compiler 負責更細的資料布局與 thread mapping。

cuTile 與 Triton 都讓使用者從 tile／block 出發，但 API、IR、backend、可觀察 artifact 與可調參數不同。CuTe DSL 提供另一條路：開發者仍直接操作更接近硬體的 layout、copy 與 MMA 抽象。

Day13：MLIR ：多層 IR 從哪裡來，又想解決什麼問題

Day13 我們一起讀 MLIR paper，了解現代 compiler 為什麼需要多層 IR。高階表示適合保留 tensor、shape 與 operation，靠近硬體的表示需要 control flow、address space、register 與 target instruction。太早丟掉高階語意，compiler 很難再做 tile、fusion 或 layout transformation；停在高階表示，也無法產生可執行程式。

MLIR 提供 dialect、operation、type、attribute、pass、rewrite 與 verification 等共用機制。它允許一個 compiler 定義多個適合不同問題的表示，再用 lowering 逐步改寫。

然後設計一層 IR 需要去思考，這一層要保留哪些資訊，以及哪些 transformation 需要這些資訊。

Day14：Triton JIT 的流程

@triton.jit 會把 Python function 包成可交給 Triton compiler 的 JITFunction。執行 kernel[grid](...) 時，runtime 會綁定參數、建立 specialization key、查詢 cache，必要時才啟動 compiler。

可以先用這條通用流程理解

kernel source
+ runtime types／constexpr／compile options／target
→ specialization key
→ cache lookup
→ miss 時編譯
→ binary + metadata
→ load + launch

這也把 benchmark 分成幾種不同成本，第一次編譯、硬碟 cache 命中、process 內已載入 module，以及純 kernel execution。

Day15：TTIR 流程

Triton frontend 走訪 Python AST，建立初始 Triton IR。進入 TTIR passes 後，artifact 仍保留 tt.load、tt.store、tt.dot、scf.for、tensor shape 與 pointer 計算等語意。

TTIR 適合回答

• program_id 是否選到正確的 output tile？
• Pointer、stride 與 offsets 是否符合 tensor layout？
• 邊界 mask 是否完整？
• K loop 與 accumulator shape 是否正確？
• constexpr 與 specialization 是否產生預期常數？

如果錯誤已經出現在 shape、mask 或 pointer，繼續追 PTX 只會增加閱讀成本。先在最接近原始語意的 artifact 確認 correctness，再往硬體映射移動。

Day16：TTGIR 流程

TTGIR 開始加入 target 與 GPU layout。相同的 tensor shape 會多出 element 如何分給 CTA、warp 與 lane 的 encoding，也能看到 global-memory coalescing、shared-memory staging、layout conversion、MMA representation 與 software pipeline。

以矩陣乘法簡化後的資料路徑為例

A / B global tiles
→ shared-memory staging
→ warp / lane register fragments
→ Tensor Core MMA
→ accumulator
→ C global memory

BLOCK_SIZE_M/N/K、num_warps、num_stages 與 target architecture 都可能改變映射。Compiler pipeline 中出現某個 pass，也不保證這支 kernel 一定使用對應硬體路徑；仍要回到實際 TTGIR、metadata 與 profiler 證據確認。

Day17：LLVM IR／NVVM

從 TTGIR 進入 LLVM IR 時，lowering 會處理 Triton GPU operations，輸出 LLVM NVPTX backend 能接收的表示

• Kernel entry 與 calling convention。
• Global、shared 等 address spaces。
• Basic blocks、branch 與 phi 等 control flow。
• NVVM intrinsics。
• 必要的 inline PTX。
• Function attributes、target data layout 與 libdevice calls。

例如 TTGIR 的 shared-memory descriptor 會變成具體的 addrspace(3) 資源；GPU thread id、barrier 與部分 matrix instructions 則會以 NVVM intrinsic 或 inline assembly 表示。到這一站，tl.load 這類 Triton 語意已大幅降低，但 GPU-specific information 仍然清楚可見。

Day18：PTX

LLVM NVPTX backend 接收 LLVM module，產生 PTX assembly text。PTX 會包含 kernel entry、address space、virtual registers 與 NVIDIA GPU operations。

接著，ptxas 針對指定的 sm_XX 目標產生 cubin

LLVM IR
→ LLVM NVPTX backend
→ PTX virtual ISA
→ ptxas
→ cubin：SASS machine code + resource metadata

ptxas 會完成 physical register allocation、target-specific instruction encoding 與 binary metadata。ptxas -v 提供的 register、spill、stack frame、constant memory 與 shared-memory 用量，可以接回 Day11 的 SM 資源限制。

PTX 的 %r163 是 virtual register，不能直接當成硬體實際使用 163 個 registers。應檢視 ptxas log、cubin metadata 或 profiler，才能判斷 physical register 壓力。

Day19：cubin 和 SASS，再到 CUDA Driver launch

cubin 是 compiler 的輸出，仍需要 runtime 把它變成一次具體執行。在 Triton 3.6.0 NVIDIA launcher 中，可以簡化成

cubin bytes
→ cuModuleLoadData()
→ CUmodule
→ cuModuleGetFunction()
→ CUfunction
→ grid / block / dynamic shared memory / stream / arguments
→ cuLaunchKernelEx()
→ GPU 執行 SASS

Module load 與 function lookup 建立目前 CUDA context 可使用的 handle，launch configuration 則描述這一次要執行多少工作、使用哪個 stream、傳入哪些 arguments。

Kernel launch 對 host 通常是非同步 enqueue。API 成功回傳只代表工作已送入 queue，並不等於 GPU 已完成。讀回結果、定位非同步錯誤或量測執行時間時，需要在正確位置使用 CUDA event 或 synchronization。

Day20：把 compile path 與 runtime path 接起來

過一次整體流程

Compile path
Python AST
→ TTIR
→ TTGIR
→ LLVM IR
→ LLVM NVPTX backend
→ PTX
→ ptxas
→ cubin / SASS

Runtime path
cubin
→ CUmodule
→ CUfunction
→ launch configuration
→ CUDA stream
→ GPU execution

編譯路徑到可載入的 target binary 為止。Runtime path 從載入 binary 開始，處理 context、function handle、arguments、stream 與 launch。把兩條路分開後，cache、module load、launch latency 與 kernel execution time 也比較不容易混淆。

問題發生時，先停在哪一層

今天先走到這裡

• GPU kernel 同時受到工作階層、運算單元與記憶體階層限制。
• 函式庫、CUDA C++、CUTLASS／CuTe DSL、cuTile 與 Triton 提供不同抽象與控制界線。
• 多層 IR 讓 compiler 在適合的位置保留運算語意、加入 layout，再降低到 target ISA。
• Triton NVIDIA backend 的可觀察主線是 TTIR、TTGIR、LLVM IR、PTX 與 cubin。
• cubin 產生後，CUDA Driver 還要載入 module、取得 function、準備 launch configuration 並送入 stream。
• 除錯時先選對 artifact，效能分析則要把靜態 artifact、resource metadata 與 profiler 時間線放在一起。

明天開始進入 Tenstorrent 好耶。硬體名稱會換成 Tensix core、L1 SRAM 與 NoC，軟體堆疊會看到 TT-Metalium 與 TT-MLIR。我們可以沿用今天的五個問題：程式從哪個抽象開始、有哪些 IR、tile 如何映射到 core、資料如何移動、runtime 如何載入並啟動工作。

參考資料

• CUDA C++ Programming Guide
• CUDA Driver API：Module Management
• CUDA Binary Utilities
• CUTLASS documentation
• CuTe DSL：End-to-End Code Generation
• cuTile Python
• Triton documentation