昨天已經看到 ttg.async_copy_global_to_local、shared-memory descriptor 與 NVIDIA MMA layout。今天再往下一層看 LLVM IR，追這些 GPU 語意怎麼變成 LLVM address space、NVVM intrinsic 與 inline assembly。

LLVM IR 很長，今天講解片段會縮短 SSA 名稱、debug metadata 與 inline-assembly operand list，calling convention、address space、intrinsic 名稱與 instruction shape 保留。

想追完整 SSA 資料流可以看這裡 matmul_kernel.llir

本篇大綱

• 對照 TTGIR 與 LLIR，找出 Triton 實際使用的 lowering passes。
• 找到 LLVM kernel entry 與 GPU thread/block id。
• 讀 global、shared memory address space。
• 對照 async copy、barrier、matrix load 與 MMA。
• 說明這次 artifact 中沒有出現的 target triple 與 libdevice call。
• 建立 LLIR 到 PTX 的對照表。

.ttgir → .llir

Triton 3.6.0 的完整 make_llir() 實作在 NVIDIA backend compiler.py 第 339～433 行。其中第 347～378 行建立並執行 MLIR pass pipeline，第 398～417 行再把 MLIR LLVM dialect module 轉成 native LLVM module、附加 NVPTX target data layout 並執行 LLVM O3。

將這段原始程式碼按責任分組後，可整理成

完整 pipeline 還包含 add_combine_tensor_select_and_if、Gluon inliner、add_check_matmul_two_cta、add_proxy_fence_insertion 與 debug scope。Concurrency sanitizer 的條件判斷與 pass 呼叫可直接看第 354～356 行：只有 instrumentation_mode == "consan" 時才加入。本次 metadata 的 instrumentation_mode 為空字串，所以不執行這個條件式 pass。

中間還會執行 canonicalizer、CSE 與 symbol DCE。這些 pass 會改變 SSA 編號和局部指令形狀，因此閱讀時應追 operation family 與資料流，不要假設 TTGIR 和 LLIR 能逐行對齊。

四組前後對照

例如 TTGIR 只用 memory descriptor 表示 A/B 的兩塊 shared-memory tile

!ttg.memdesc<1x64x32xf16, ... #smem ...>
!ttg.memdesc<1x32x64xf16, ... #smem ...>

LLIR 則已分配成具體 byte offset

@global_smem = external addrspace(3) global [0 x i8], align 16
%a_smem = getelementptr i8, ptr addrspace(3) @global_smem, i32 %a_offset
%b_smem = getelementptr i8,
    ptr addrspace(3) getelementptr (i8, ptr addrspace(3) @global_smem, i32 4096),
    i32 %b_offset

TTGIR 對「一塊 tile 在 shared memory」的描述，到 LLIR 已變成「從 @global_smem + 4096 開始的 B tile」。這就是 allocation pass 將語意變成位址的具體例子。

LLVM IR 很長，先找錨點

同一個 tt.dot 到 LLVM IR 會展開成許多 scalar SSA value。逐行從頭讀，很快會迷失在 %154、%155 這類編號。本篇先找五個錨點

ptx_kernel
addrspace(1) / addrspace(3)
read.ptx.sreg
cp.async / ldmatrix
mma.sync

找到錨點後，再向上追它的 pointer，向下追它產生的值。這比試著理解每一個 and、xor、shl 更有效。

從 kernel entry 開始

Function entry 是：

define ptx_kernel void @matmul_kernel(
    ptr addrspace(1) %0,
    ptr addrspace(1) %1,
    ptr addrspace(1) %2,
    ptr addrspace(1) %3,
    ptr addrspace(1) %4) {

ptx_kernel calling convention 告訴 LLVM NVPTX backend，這個 function 之後要成為 PTX .entry kernel。前三個 addrspace(1) pointer 對應 A、B、C 的 global-memory 位址。

實際 signature 還有 compiler/runtime 產生的額外 pointer。這裡不根據位置猜用途；判斷參數角色時，應追蹤它在 function body 中的 use，或對照下一站 PTX parameter。

Thread 與 block id 變成 NVVM intrinsic

LLIR 藉由 NVVM special-register intrinsic 取得硬體座標

%ctaid = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
%tid = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()

它們下一站會分別變成 PTX %ctaid.x 與 %tid.x。原本 Triton 的 tl.program_id(0)，到這裡已經接到 NVIDIA execution model；warp/lane 的 address calculation 也被展成 integer 與 bit operation。

真實 IR 緊接著出現大量位元運算

%lane_group = lshr i32 %tid, 2
%lane_group_masked = and i32 %lane_group, 31
%lane = and i32 %tid, 7
%lane_offset = shl i32 %lane, 3
%row = or i32 %program_offset, %lane_group_masked

這些運算是 TTGIR layout 的具體化，compiler 正把 CTA、warp 與 lane 映射回 global/shared-memory 位址。它已經不會長得像 Python 的 offsets_m[:, None]，所以閱讀時要追最後餵給哪個 getelementptr 或 intrinsic。

Global 與 shared address space

LLIR 開頭宣告

@global_smem = external addrspace(3) global [0 x i8], align 16

這是動態 shared-memory base。對照關係為

Address space 會影響後端選擇 ld.global、shared-memory operation 或 address conversion。若這層標錯，除了效能可能變差，程式還可能讀取完全不同的 memory space。LLVM SSA value 不保證逐一變成實體 register；後續最佳化可能消除或合併它，ptxas 也可能重新配置或 spill。Physical register 數量要到明天的 ptxas report 才能確認。

這份 IR 也能看到 A/B 在 shared memory 的分界

%a_smem = getelementptr i8,
    ptr addrspace(3) @global_smem, i32 %a_offset
%b_smem_base = getelementptr i8,
    ptr addrspace(3) @global_smem, i32 4096
%b_smem = getelementptr i8,
    ptr addrspace(3) %b_smem_base, i32 %b_offset

4096 正好是 A tile 的大小：64×32×2 bytes。因此 B tile 從 dynamic shared-memory base 的第 4096 byte 開始。這和昨天的兩個 local_alloc 對得上。

TTGIR operation 如何 lower

這次 artifact 的關鍵對照如下

例如 async copy 在 LLIR 中已接近 PTX

call void asm sideeffect
  "cp.async.cg.shared.global
   [ $0 + 0 ], [ $1 + 0 ], 0x10, $2;", ...
call void @llvm.nvvm.cp.async.commit.group()
...
call void @llvm.nvvm.cp.async.wait.group(i32 0)
call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0)

0x10 是一次 16-byte copy。最後一個 operand 控制有效 copy bytes；mask 為 false 時可以讓這次 copy 以 0 bytes 的有效資料處理，延續 TTIR 的 boundary mask。commit 表示一組 copy 已發出，wait_group(0) 等待所有未完成 group，barrier 再讓 CTA 內 threads 在讀 shared memory 前會合。

接著用 ldmatrix 把 shared-memory tile 讀成 MMA fragment

%a_frag = call {i32, i32, i32, i32}
    @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x4.b16.p3(
      ptr addrspace(3) %a_smem)
%b_frag = call {i32, i32, i32, i32}
    @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x4.trans.b16.p3(
      ptr addrspace(3) %b_smem)

p3 表示輸入 pointer 位於 address space 3，也就是 shared memory；B 使用 trans 版本，對應 dot operand 需要的矩陣方向。回傳的四個 i32 裝著打包後的 FP16 fragment。

Matrix multiply 也明確寫出 shape 與 dtype

call ... asm sideeffect
  "mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 ...", ...

到這裡已經能確認昨天的 nvidia_mma versionMajor=2 會走 mma.sync，不必等看到 SASS 才第一次知道。

Inline asm 回傳四個 float，後面會看到

%d0 = extractvalue {float, float, float, float} %mma, 0
%d1 = extractvalue {float, float, float, float} %mma, 1
%d2 = extractvalue {float, float, float, float} %mma, 2
%d3 = extractvalue {float, float, float, float} %mma, 3

這四個值是目前 thread 持有的一部分 FP32 accumulator，不是完整的 64×64 C tile。完整 tile 分散在四個 warp 的 threads 中。

輸出 layout conversion 如何出現在 LLIR

昨天看到輸出端有 ttg.convert_layout。LLIR 中可以找到

%packed0 = bitcast <2 x half> %c_fragment0 to i32
%packed1 = bitcast <2 x half> %c_fragment1 to i32
call void @llvm.nvvm.stmatrix.sync.aligned.m8n8.x4.b16.p3(
    ptr addrspace(3) %smem_out,
    i32 %packed0, i32 %packed1, ...)
call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0)
%out = load <4 x i32>, ptr addrspace(3) %smem_out, align 16

這份 codegen 使用 shared memory 作為 layout conversion 的中繼站，先用 stmatrix 把 fragment 寫入 shared memory，barrier 後再依 global store 所需排列讀出。

為什麼 LLVM IR 還會有 inline PTX

LLVM/NVVM intrinsic 能表達許多 NVIDIA GPU operation，但 backend 支援與 Triton 想控制的細節不一定完全重合，Triton 會混用

• 一般 LLVM instruction
• NVVM intrinsic
• inline PTX assembly

所以已經進入 LLVM IR不代表所有 operation 都必須是 target-independent LLVM op。這份 LLIR 本來就是通往 NVPTX 的交接層。

這份 LLIR 沒有明列 target triple

許多 LLVM IR 教學會先找

target triple = "nvptx64-nvidia-cuda"

但 Triton 3.6.0 這次輸出的 LLVM IR 並沒有把 target triple 印在檔案開頭。Target 仍由 NVIDIA backend 的編譯流程傳入 LLVM translation。

判斷這是 NVIDIA kernel 的直接證據仍很充分：ptx_kernel calling convention、NVVM intrinsic、PTX inline asm，以及上一層 ttg.target="cuda:121"。

Libdevice：已設定，不代表這個例子有呼叫

Metadata 中記錄了 CUDA libdevice bitcode 路徑。它讓 sin、cos、exp 等 device math function 在需要時能被連結。

這個矩陣乘法只做 load、dot、cast 與 store，實際 function body 沒有 libdevice math call。因此比較精確的說法是編譯環境配置了 libdevice，而不是這個 kernel 已連入並執行 libdevice。

LLIR 適合查什麼

明天把 PTX 與 ptxas -v log 放在一起讀。會分清 PTX 的 196 個 virtual .b32 register，和 ptxas 最終配置的 80 個 physical register。

參考資料

• LLVM NVPTX backend guide
  https://llvm.org/docs/NVPTXUsage.html
• LLVM NVVM intrinsic documentation
  https://llvm.org/docs/NVPTXUsage.html#nvvm-intrinsics
• MLIR NVVM dialect
  https://mlir.llvm.org/docs/Dialects/NVVMDialect/
• Triton NVIDIA backend compiler
  https://github.com/triton-lang/triton/blob/v3.6.0/third_party/nvidia/backend/compiler.py