昨天先看 Wormhole 的 tile 和 NoC，今天再往下走一層，看 Tenstorrent software stack，最後把焦點放到 TT-Metal / Metalium。我最想先搞懂兩件事，上層模型怎麼一路接到硬體，還有 TT-Metal 程式最常見的 reader、compute、writer 分工。
在看 TT-Metal 之前，可以先理解 Tenstorrent 的 software stack。
圖片來源：Tenstorrent 官方 GitHub 組織頁面的 AI SW 圖
TT-Forge 是比較上層的 compiler stack。官方文件把它描述成 end-to-end compiler stack，包含 TT-XLA、TT-Forge-ONNX、TT-MLIR、TT-Blacksmith 等元件。它負責把 PyTorch、JAX、ONNX、TensorFlow 這類模型來源接進 Tenstorrent 的 compiler pipeline。
TT-NN 是高階 neural network operation library，有 Python 和 C++ API。對熟悉 PyTorch 的開發者來說，它比較像「我想呼叫一個 tensor op 或跑一段模型，不想一開始就管理 NoC、L1、circular buffer」。TT-NN 底下仍然建立在 TT-Metalium 上。
TT-Metalium，也常被簡稱成 TT-Metal，是低階 C++ SDK。官方文件把它定位成讓開發者直接控制 Tenstorrent 硬體的 programming environment。這一層會碰到 program、kernel、buffer、command queue、NoC、L1、Tensix compute。
TT-LLK 可以先理解成更靠近 Tensix Engine 的 low-level kernel 層。Tenstorrent 的 TT-LLK 頁面提到 LLK 會控制 Tensix Engine，提供使用 Tensix instruction set 的路徑，讓低階數學操作能貼近硬體執行。
TT-Lang 則是比較新的 Python-based DSL，用來寫 Tenstorrent 上的高效 custom kernel / fused op。官方文件把 TT-Lang 放在 TT-NN 和 TT-Metalium 之間的使用情境：它讓開發者用 Python DSL 顯式描述 computation 和 data movement，也支援 simulator / compiler / hardware flow。
所以今天看 TT-Metal 時，可以先把它放在這個位置
想快速跑模型：
TT-Forge / TT-NN
想寫常見 NN op：
TT-NN
想寫 fused custom kernel，但希望用 Python DSL：
TT-Lang
想直接控制 buffer、NoC、kernel、program：
TT-Metalium
想理解 Tensix engine 最低階數學 kernel：
TT-LLK
後面讀 TileLoom 會一直碰到 TT-Metalium，所以今天先把這層看懂。可以暫時把 TileLoom 想成一個 compiler 系統：它會把高階 tile program 轉成能在 Tenstorrent 上執行的資料流計畫，最後再落到 TT-Metalium C++ API。
Tenstorrent 官方 guide 把 TT-Metalium 定位成低階 programming model。它負責讓開發者在 Tensix processor 上建立 program、配置 buffer、啟動 kernel、管理資料流。
如果只畫出這一篇要追的低階落點，可以寫成
TT-Forge 編譯模型 ───────────┐
TT-NN 呼叫高階 tensor op ────┼─→ TT-Metalium → Tensix hardware
TT-Lang 寫 fused custom kernel ┘
TT-Forge 比較像模型進 Tenstorrent compiler stack 的入口。TT-NN 讓你用高階 tensor op 跑模型或組 op。TT-Lang 讓你用 Python DSL 寫 custom / fused kernel。TT-Metalium 再往下，會直接碰到資料從 DRAM 搬到 L1、kernel 等 circular buffer、compute 做 tile operation、writer 把結果送回去。
這也是後面 TileLoom 會選 TT-Metalium 當 backend 的原因。TileLoom 做 dataflow planning 後，需要一個能表達 buffer allocation、synchronization、data movement、compute operation 的後端。
直接看 reader、compute、writer 時，很容易漏掉「誰建立這三個 kernel」。完整程式還有 host 端，兩側分工如下
Host C++ program
1. 開啟 device、取得 command queue
2. 建立 DRAM／L1 buffer 與 device program
3. 建立 circular buffer
4. 建立 reader、compute、writer kernel
5. 設定 compile-time／runtime arguments
6. 把輸入寫到 device、enqueue program、讀回結果
Device kernels on Tensix
reader → compute → writer
Host 端負責配置與啟動，device kernel 負責實際資料搬移與運算。本文聚焦 device 端，是因為這三個 kernel 最能把 Day21 的 NoC、L1、Unpack／Math／Pack 接到程式碼；若要自己跑完整範例，仍需要 host 端把 program 組起來。
Metalium guide 裡的典型資料流會用三種 kernel
Reader kernel
-> 從 DRAM 讀 input tile
-> 寫進 circular buffer
Compute kernel
-> 等 input circular buffer 有資料
-> 做 tile compute
-> 寫到 output circular buffer
Writer kernel
-> 等 output circular buffer 有資料
-> 寫回 DRAM 或指定輸出位置
官方圖把這三種 kernel 實際框在 Tensix 方塊圖上
圖片來源：Tenstorrent
TT Architecture and Metalium Guide。圖中的 Movement kernel 0／1 對應常見的 reader／writer 分工；這是典型配置，其他 operation 也能採用不同安排。
這張圖回答了「為什麼程式看起來只有三個 kernel，硬體卻有五個 Baby RISC-V core」：reader 與 writer 各由一個 Data Movement core 執行；同一份 compute kernel 會針對 Unpack、Math、Pack 編譯成三份對應的 binary，由三個 core 並行合作。開發者不需要把 compute kernel 手寫成三個 thread。
CPU 函式通常依序完成載入、運算與儲存；Tenstorrent 在這裡採用一條小型管線：資料搬移和計算可以分開排程，中間用 circular buffer 溝通。
下面用 C = A + B 看一個 tile 如何穿過三支 kernel。為了先看懂協作關係，程式碼省略 #include、TensorAccessor 建立、compile-time arguments 與錯誤處理，只保留主要控制流程。因此，這些片段是依照官方範例縮短的教學版，不能單獨複製編譯。完整 host 與 device 程式可對照官方的 Eltwise binary example。
三個 circular buffer 的角色先固定下來
| Circular buffer | 放什麼 | Producer | Consumer |
|---|---|---|---|
cb_in0（c_0） |
A 的 input tile | reader | compute |
cb_in1（c_1） |
B 的 input tile | reader | compute |
cb_out（c_16） |
C 的 output tile | compute | writer |
Reader 的迴圈可以縮成
for (uint32_t tile = 0; tile < n_tiles; ++tile) {
cb_reserve_back(cb_in0, 1);
cb_reserve_back(cb_in1, 1);
noc_async_read_page(tile, input_a, get_write_ptr(cb_in0));
noc_async_read_page(tile, input_b, get_write_ptr(cb_in1));
noc_async_read_barrier();
cb_push_back(cb_in0, 1);
cb_push_back(cb_in1, 1);
}
這裡的順序不能交換
cb_reserve_back 先確認 circular buffer 有可寫入的空位。get_write_ptr 取得這個空位在 L1 的位址。noc_async_read_page 發出 DRAM 到 L1 的非同步 NoC read。noc_async_read_barrier 等兩筆讀取完成。cb_push_back 才把 A、B 公布給 compute。如果 reader 在 barrier 前就呼叫 cb_push_back，compute 可能看到「buffer 已有資料」，實際的 NoC read 卻還沒完成。
Compute kernel 先用 compute_kernel_hw_startup 和 add_init 準備運算，再進入 tile loop
compute_kernel_hw_startup(cb_in0, cb_in1, cb_out);
add_init(cb_in0, cb_in1);
for (uint32_t tile = 0; tile < n_tiles; ++tile) {
cb_wait_front(cb_in0, 1);
cb_wait_front(cb_in1, 1);
tile_regs_acquire();
add_tiles(cb_in0, cb_in1, 0, 0, dst_reg);
tile_regs_commit();
cb_reserve_back(cb_out, 1);
tile_regs_wait();
pack_tile(dst_reg, cb_out, 0);
tile_regs_release();
cb_pop_front(cb_in0, 1);
cb_pop_front(cb_in1, 1);
cb_push_back(cb_out, 1);
}
這段可分成三個部分
cb_wait_front：等 reader 公布 A、B。add_tiles：把兩個 input tiles 加到 destination register。pack_tile：把 register 裡的結果 pack 到 cb_out 的 L1 空間。最後，cb_pop_front 表示 input tiles 已經用完；cb_push_back(cb_out, 1) 則向 writer 公布一個 output tile。
Compute kernel 雖然只有一份 source，Metalium 會替 Unpack、Math、Pack 三個 Baby RISC-V core 產生對應 binary。像 add_tiles 會同時涉及 Unpack 與 Math，pack_tile 則由 Pack 負責。
Writer 只需要等待 compute 的結果
for (uint32_t tile = 0; tile < n_tiles; ++tile) {
cb_wait_front(cb_out, 1);
noc_async_write_page(tile, output_c, get_read_ptr(cb_out));
noc_async_write_barrier();
cb_pop_front(cb_out, 1);
}
cb_wait_front 確認 cb_out 有資料；get_read_ptr 取得結果的 L1 位址；NoC write 完成後，cb_pop_front 才能釋放這個 slot，讓 compute 後續重用。
Reader、compute、writer 都是 device kernels。Host 會把它們建立在同一個 Metalium program，替三支 kernel 分別設定 runtime arguments，再 enqueue 整個 program
reader = CreateKernel(program, reader_source, core, DataMovementConfig{RISCV_0, ...});
writer = CreateKernel(program, writer_source, core, DataMovementConfig{RISCV_1, ...});
compute = CreateKernel(program, compute_source, core, ComputeConfig{...});
SetRuntimeArgs(program, reader, core, {a_addr, b_addr, n_tiles});
SetRuntimeArgs(program, compute, core, {n_tiles});
SetRuntimeArgs(program, writer, core, {c_addr, n_tiles});
EnqueueMeshWorkload(command_queue, workload, false);
三支 kernel 啟動後可以同時向前推進。Circular buffer 的 producer／consumer 狀態與 NoC barrier 共同維持資料相依：reader 可以開始搬下一個 tile，compute 處理目前的 tile，writer 同時送出上一個結果。
以第 i 個 output tile 為例，可以把三個 kernel 的協作排成以下時間順序
cb_reserve_back，先確認 input circular buffer 還有空位。noc_async_read_page。noc_async_read_barrier 確認兩次 NoC read 已完成，reader 才能 cb_push_back 公布資料。cb_wait_front 等到 A、B 都可用，再由 Unpack core 把資料送進 compute engine。add_tiles，結果先停在 destination register。pack_tile 到 L1，再 cb_push_back。cb_wait_front 看到結果後，發出 noc_async_write_page。noc_async_write_barrier 確認資料已送出，writer 才能 cb_pop_front 釋放空位。這個順序裡有兩種等待：NoC barrier 等「硬體搬完資料」，circular buffer API 等「producer／consumer 的 buffer 狀態」。兩者解決的問題不同，少了任何一種都可能在資料尚未準備好時往下執行。
Circular buffer 是 TT-Metal 程式裡很重要的結構，它除了保存 tile，也帶著同步語意。
圖片來源：Tenstorrent
TT Architecture and Metalium Guide。綠色方塊代表位於 SRAM 的 circular buffer。Circular buffer 也能讓同一個 kernel 把資料送回自己，並不只用於三段式 pipeline。
對 reader 來說
cb_reserve_back 表示我要寫一個 tile 進來，先保留空間。cb_push_back 表示資料已經準備好，可以給下一段使用。對 compute 來說
cb_wait_front 表示我要等資料真的出現。cb_pop_front 表示這筆資料我用完了，可以釋放。Vector addition 會使用至少三個 circular buffer：cb_in0 放 A、cb_in1 放 B、cb_out 放 C。Reader 是前兩個 buffer 的 producer，compute 是它們的 consumer；compute 同時又是 cb_out 的 producer，writer 則是 consumer。用 producer／consumer 角色來讀 API，比背函式名稱更容易檢查同步是否完整。
這個設計讓 reader、compute、writer 可以各自前進。只要 buffer 和 barrier 設計正確，資料搬移和計算就能形成 pipeline。
這裡要小心 race condition。guide 也強調 NoC interface 和 compute engine 是獨立 peripheral，很多操作是 asynchronous。程式要自己放 barrier 和 buffer synchronization，否則 compute 可能讀到還沒搬完的資料，writer 也可能寫出還沒算完的結果。
Metalium guide 裡提到 compute API 會被編譯成給不同小核心執行的 code section，例如 Unpack、Math、Pack。像 add_tiles 這類 API 會對應到資料 unpack、FPU 運算、pack result 這些硬體步驟。
可以先用這個簡化圖理解
input circular buffer
-> Unpack
-> Math / FPU
-> Pack
-> output circular buffer
官方資料流圖把 NoC 與 L1 也放回來
圖片來源：Tenstorrent
TT Architecture and Metalium Guide。粉紅線是這張圖示範的一條典型資料路徑。
把三張圖合起來後，可以得到一個實用的讀碼順序：先看資料從哪個 DRAM／core 進來，再看它進入哪個 circular buffer，最後才看 compute API 做什麼。TT-Metalium 的錯誤常發生在 address、tile count、buffer 容量或同步；若只盯著 add_tiles，會漏掉大部分資料流。
這和昨天的架構能接起來，Tensix core 裡的資料搬移、unpack、math、pack 由專門部件分工完成。Metalium API 再把這些步驟包成開發者可以呼叫的單位。
如果手寫 TT-Metal，開發者要決定
後面要看的 TileLoom，價值就在這裡。它希望讓使用者從 Triton / Helion 這類 tile-based program 開始，讓 compiler 自動規劃一部分 dataflow，再產生能接到 TT-Metalium 的程式。
今天先記住幾個 TT-Metal 基本觀念
明天先來看一下 TT-MLIR 的 dialect 與 pipeline 文件，看 TTIR、TTNN、D2M、TTKernel 和 TTMetal 分別保留什麼資訊。
iThome鐵人賽
看影片追技術