最開始十天都在 RISC-V CPU 上打轉，從 RVV intrinsic、SGLang RVV 後端走到 torch.compile 和 PyTorch Inductor RVV，那時候基本單位是 CPU core、vector register、cache、C++ kernel 和 compiler codegen。

接下來十天換到 NVIDIA GPU、CUDA、CUTLASS、cuBLAS、CuTeDSL、cuTile、MLIR，再一路追 Triton 從 TTIR、TTGIR、LLVM IR、PTX、cubin 走到 CUDA Driver API launch。名詞也全部換成 SM、warp、shared memory、Tensor Core，還有 compiler 和 runtime 怎麼把 tile-level program 送上 GPU。

今天終於輪到 Tenstorrent。問題還是資料在哪裡、怎麼搬、怎麼重用、compiler 在哪一層做 mapping，但硬體名詞會換成 Tensix tile、L1、NoC、DRAM tile，最佳化方式也更靠近 dataflow planning、broadcast 和 core placement。

接下來這十天會出現幾個新的夥伴

• TT-Metalium / TT-Metal：Tenstorrent 的低階 C++ SDK，可以直接管理 kernel、buffer、NoC、L1、device program
• TileLoom：一篇研究 paper，目標是把 Triton 或 Helion 這種 tile-based program，自動規劃到 Tenstorrent 這種 core grid + local memory + NoC 的硬體上

今天只要先看懂 Tenstorrent 硬體為什麼會讓 compiler 思考資料放哪裡、怎麼搬、哪些 core 可以共用。

本篇大綱

• 先用 Wormhole 和 Blackhole 看 Tenstorrent 架構如何延續。
• 拆開 Wormhole ASIC，確認裡面有哪些 tile。
• 再看 Tensix tile、DRAM tile、Ethernet tile、PCIe tile、ARC tile 各自的用途。
• 接著看 2D torus NoC 對資料搬移的意義。
• 用同一個 matmul 比較 CPU、NVIDIA GPU 與 Tenstorrent 的執行方式。
• 最後把硬體差異接到 compiler 的 vectorization、GPU mapping 與 dataflow planning。

先認識 Wormhole 和 Blackhole

Tenstorrent 的硬體名稱很多，這次介紹 Wormhole 和 Blackhole。
兩者都沿用 Tensix core grid、本地 L1、可程式化資料搬移與
NoC 的設計。先看懂 Wormhole 的 tile 分工，之後讀 Blackhole 或新的軟體
文件時，Tensix、DRAM、Ethernet、PCIe 和 NoC 這些名詞仍然接得起來。

Wormhole 是目前許多 TT-Metal 教材、範例和研究工作的硬體背景。一顆 ASIC
最多有 80 個 Tensix tile，實際產品可能開放 64、72 或 80 個，並搭配
12 GiB GDDR6。接下來以 Wormhole 拆解架構；它的 ISA documentation
清楚列出各種 tile 的數量與用途。

Blackhole 是更新的架構。以目前官方列出的 Blackhole p100a、p150a 與
p150b 卡為例，每顆處理器提供 120 個可用 Tensix core、180 MB SRAM，也就是
每個 Tensix core 1.5 MB，並搭配 16 個 SiFive X280「Big RISC-V」core。
不同卡提供 28 GB 或 32 GB GDDR6，host interface 也從 Wormhole 的 PCIe
4.0 x16 前進到 PCIe 5.0 x16。

這張表用來定位兩個架構，不適合直接拿來判斷整體效能。核心數、記憶體、
資料格式、時脈、軟體版本和 workload 都會影響最後結果。接下來先固定在
Wormhole，把一顆 ASIC 裡面的角色逐一看懂。

先看一顆 Wormhole ASIC 裡有什麼

Wormhole README 把一顆 ASIC 拆成幾種 tile

80x Tensix tile
18x DRAM tile
16x Ethernet tile
1x PCI Express tile
1x ARC tile
2x NoC

這裡的 tile 是 NoC grid 上的一個硬體節點。每種 tile 接到同一張
on-chip network，但內部資源與工作不同：Tensix tile 做運算，DRAM tile
提供外部記憶體入口，Ethernet tile 連其他 ASIC，PCIe tile 連 host，ARC
tile 管理晶片與板卡。兩套 NoC 再把這些節點接起來。

先看官方 ISA 文件裡的實際 NoC layout

圖片來源：Tenstorrent tt-isa-documentation 的 Wormhole NoC layout。

讀這張圖時先不要數每一格，先看顏色與字母

• T0 到 T79 是 80 個實體 Tensix tile，佔據 grid 的主要區域。
• D0 到 D5 是六組 GDDR6 資源的入口。同一個 D 編號在圖上出現三次，對應「每 2 GiB GDDR6 由三個 DRAM tile 暴露」的設計。
• E0 到 E15 集中在上下邊緣，用來和其他 ASIC 連線。
• PCIe 與 ARC 在左側，分別處理 host interface 與板卡管理。
• 白色空格不是運算核心；NoC grid 仍可能在該座標保留 router 與 NoC interface unit（NIU）。

這張圖也解釋了為什麼後面的程式會使用 (x, y) 座標。DRAM、運算核心與外部連線都在同一張 grid 上，資料從哪裡出發、送到哪一格，會影響走過的 NoC 路徑與可能遇到的壅塞。

CPU 文章裡我常問：這個欄位會影響 -march、intrinsic、cache、threading 哪裡。Tenstorrent 這裡要換成另一組問題

• 這個 tensor tile 會放在哪個 Tensix core 的 L1？
• 從 DRAM 搬到 L1 時要使用 NoC0 還是 NoC1，來源與目的座標是什麼？
• 哪些 core 可以共用同一份資料？
• 有沒有必要把資料 broadcast 到一排或一列 core？
• 最後結果要從哪個 core 寫回 DRAM？

這些問題會一路接到後面要讀的 TileLoom。這裡先把 TileLoom 當成「幫 compiler 規劃資料流的研究系統」就好。它要處理的事情包含：單一 tile program 怎麼編、很多 logical tile 要怎麼放到實體 core、資料要不要 broadcast、哪些 tile 可以重用。

Tensix tile：計算發生的地方

Tensix tile 也叫 worker tile 或 compute tile，是 Wormhole 裡執行主要
workload 的節點。ASIC 內有 80 個實體 Tensix tile；依產品組態，使用者會
看到 64、72 或 80 個，其餘 tile 會被 fused off。因此寫 benchmark 時，
應該查詢裝置實際回報的 compute grid，不能只拿 ASIC 上的實體數量來算。

每個 Tensix tile 可以先拆成幾類資源

• 1464 KiB 本地 SRAM，也就是 L1，保存輸入、輸出與中間 tile；
• 五個稱為 Baby RISC-V 的 RV32IM core，負責監督資料搬移與 Tensix coprocessor；
• 兩個 unpacker，把資料從 L1 搬進 Tensix coprocessor；
• Matrix Unit（FPU），執行低精度矩陣乘加等運算；
• Vector Unit（SFPU），執行 32-lane SIMD 運算；
• Scalar Unit（ThCon），執行整數 scalar 與部分 L1 memory operation；
• 四個 packer，把運算結果送回 L1；
• 兩個 NoC connection，讓本地 core 存取其他 tile，也讓其他 tile 存取這個 tile。

圖片來源：Tenstorrent TT Architecture and Metalium Guide。藍色箭頭表示指令派送，橘色箭頭表示資料傳輸。

這裡最容易看錯的是藍色方塊。五個 Baby RISC-V core 主要負責控制與派送：兩個 Data Movement core 操作 NoC interface，另外三個 core 分別控制 Unpack、Math、Pack。大量矩陣與向量運算由橘色的 FPU／SFPU 等單元完成。各單元經由中央綠色的 SRAM 交換資料；這塊 SRAM 就是文件常說的本地 L1。方塊圖以 1.5 MB 概括容量，Wormhole ISA 文件列出的精確數字則是 1464 KiB。

這些部件讓一個 Tensix tile 同時具有 local memory、data movement 和
compute 能力。官方文件給出的自然分工是：兩個 RISC-V core 監督 NoC，
另外三個監督 Tensix coprocessor。RISC-V core 負責控制這些硬體部件，
主要吞吐量則來自 NoC 與 Tensix coprocessor。

資料不會因為程式讀了某個 address 就自動出現在本地 L1；reader kernel
需要經由 NoC 搬入資料，compute kernel 才能在本地運算。Day22 的 reader、
compute、writer 分工就是從這個硬體結構長出來的。

這和一般 GPU 的想像不太一樣。GPU 開發時常常先想 thread block、warp、shared memory。Tenstorrent 這邊更容易先想

一個 tile program
-> 搬資料進本地 L1
-> 在本地 compute unit 上算
-> 把結果留在 L1、傳給別的 tile，或寫回 DRAM

所以後面會一直看到 tile 這個詞。不過 Tenstorrent 文件裡有兩種 tile，要先分開

一顆 Tensix hardware tile 會反覆接收、處理許多 tensor tiles。兩者名稱相同，卻不是「一顆核心只能放一個 32×32 矩陣」。

Tensor tile：為什麼常看到 32×32

Tenstorrent 的主要運算與資料搬移以 tile 為基本單位。官方文件目前把 32×32 稱為最常見、支援最完整的形狀；其他 tile size 也存在，但支援範圍較有限。資料型別會決定一個 tile 佔多少空間，例如 FP32 的 32×32 tile 是 4 KiB，BF16 則是 2 KiB。

若矩陣最後兩個維度不是 32 的倍數，就要補齊到 tile boundary。以下官方例子把 48×1024 矩陣的 48 rows 補成 64 rows

圖片來源：Tenstorrent TT Architecture and Metalium Guide。更完整的 tile 與 face layout 可參考 TT-Metalium Tiles 文件。

這個例子沿著 column 方向有 1024 ÷ 32 = 32 個 tile，沿著 row 方向補齊後有 64 ÷ 32 = 2 個 tile，合計 64 個 tensor tiles。做矩陣相乘時，padding 通常要填 0，避免補上的值改變計算結果。

等到 Day22 看 circular buffer，buffer 的容量也常用「可以容納幾個 tensor tiles」來描述。硬體 grid 決定資料送到哪顆 Tensix，32×32 tensor tile 則是送進去的資料單位。

DRAM tile：資料還是要付搬移成本

Wormhole 的 18 個 DRAM tile 合計暴露 12 GiB GDDR6。這裡不能把「18 個
tile」直接理解成 18 份獨立記憶體容量：官方文件說明，每 2 GiB GDDR6 會
以相同方式暴露在 3 個 DRAM tile 上。DRAM tile 就是 NoC grid 上
通往外部記憶體的入口，18 個入口共同對應 12 GiB 容量。

多個入口讓 Tensix tile 可以從不同位置進出 DRAM，也使 placement、
NoC0／NoC1 選擇與傳輸端點影響實際流量。資料從 GDDR6 進入 compute tile 的 L1 時要穿過 NoC；
若大量 core 都反覆讀取同一份資料，DRAM 頻寬和 NoC link 都可能成為限制。
compiler 或手寫 TT-Metal 程式因此要安排資料從哪個 DRAM bank 讀取、放到
哪個 core，以及能否在 core 之間重用。

如果一個 matmul 的 A tile 被很多 output tile 重複使用，那 compiler 可以問

• 是否先把 A tile 搬到某一群 core 附近？
• 是否 broadcast 給一列 core？
• 是否讓多個 output tile 在時間上重用同一份 A tile？

這些選擇對 CPU cache 來說多半交給硬體和 cache hierarchy。Tenstorrent 這類 spatial dataflow accelerator 讓軟體有更多控制權，也讓 compiler 有更多責任。

Ethernet tile 和多 ASIC

16 個 Ethernet tile 各自提供雙向 100 GbE，設計用途是把多顆 Wormhole
ASIC 串在一起。連線形式可能是 QSFP-DD cable、Warp 100 bridge，或直接做在
電路板上的 static link。資料從某個 Tensix tile 前往另一顆 ASIC 時，會先
走晶片內 NoC 到 Ethernet tile，再經過外部 link。

有些產品不會把全部 Ethernet tile 接到外部。未接線的 tile 仍保有自己的
RISC-V 與 L1，可以被使用，但不會收到 Ethernet packet，送出的 packet
也沒有外部目的地。閱讀硬體圖時要分清「ASIC 具備哪些 tile」與
「板卡實際接出哪些 link」。

等到模型或 workload 跨多顆 ASIC，問題會從「一顆晶片內怎麼安排 core」再擴大成「不同晶片之間怎麼分資料、排通訊」。後面讀 TileLoom paper 時，會再回到這個問題：compiler 需要一份硬體描述，才知道 core、memory、NoC 和多晶片連線大概長什麼樣。

PCI Express tile：host 進入 ASIC 的入口

Wormhole 有一個 PCI Express tile，提供 PCIe 4.0 x16 host interface。Host
經由它設定裝置、提交 command，也能進行 host 與 device 之間的資料傳輸。
從資料路徑看，PCIe tile 先把 host traffic 帶進晶片，再由 NoC 送往 DRAM
tile 或 worker tile。

在多 ASIC 產品中，不一定每顆 ASIC 的 PCIe tile 都直接接到 host。沒有
直接 PCIe 連線的 ASIC 仍可經由其他晶片與 Ethernet fabric 被存取。因此
PCIe 解決 host-to-device 入口，Ethernet 則繼續處理 device-to-device
連線，兩者的角色要分開看。

ARC tile：管理晶片，不執行使用者 kernel

ARC tile 處理 chip 與 board management，例如初始化、監控及管理相關工作。
官方 Wormhole 文件特別指出，它不執行 customer workload，也不參與
customer workload 的 dispatch。學習 TT-Metal kernel 時通常不會直接操作
ARC tile，但系統能否正常啟動、管理時脈與維持板卡狀態，仍需要這個管理
節點。

NoC：資料流架構的重點

Wormhole README 說晶片裡有兩個 NoC，NoC 把不同 tile 接成 2D torus。2D torus 可以先想成一個會左右、上下繞回去的 2D grid。

兩套 NoC 都能在同一顆 ASIC 內傳送資料，並以相反方向繞行 torus。程式可以
讓不同資料搬移使用不同 NoC，減少所有 traffic 擠在同一條路徑，也能讓
read、write 與 compute 更容易重疊。NoC 上搬的內容不只包括 tensor data，
還包含不同 tile 之間的訊號與同步需求。

圖片來源：Tenstorrent TT Architecture and Metalium Guide。綠色與藍色代表彼此獨立、方向相反的兩套 NoC。

可以把它想成同一個路口疊了兩組單向道路。典型安排會讓 NoC0 讀入資料、NoC1 送出結果，因此兩個方向可以同時工作；這只是常見配置，程式也能為了特定資料搬移同時使用兩個 NoC。軟體選擇 NoC、來源與目的座標，硬體 routing 規則再決定封包如何逐跳前進。

這個結構會直接影響資料搬移策略。假設一個矩陣乘法被切成 64 個 output tile，放在 8x8 的 compute grid 上。同一塊 A tile 可供橫列中的多個 output tile 使用；同一塊 B tile 則可供直欄中的多個 output tile 使用。合理的 mapping 能沿著 grid 重用資料，減少每個 core 各自回 DRAM 抓資料的次數。

可以用這張簡化圖想

      B tile broadcast
        ↓ ↓ ↓
A -> [C][C][C]
A -> [C][C][C]
A -> [C][C][C]

這張圖只是幫我們抓方向：NoC 讓「資料怎麼走」變成 compiler 可以規劃的東西。

比較三種硬體

前面已經走過 RISC-V CPU 與 NVIDIA GPU。Tenstorrent 同樣能執行矩陣乘法，但三種硬體提供給 compiler 的控制介面不同。

這張表描述本系列採用的程式模型，不表示 CPU 完全沒有多核資料組態，也不表示 GPU 的資料搬移都由硬體自動完成。在 Ampere 上可用 cp.async，在 Hopper 上可用 TMA，兩者都能安排 global memory 到 shared memory 的傳輸。Tenstorrent 再把控制範圍擴大到多顆 Tensix core 的 L1 與 NoC。

用同一個 matmul 比較三條資料路徑

假設要計算

C[M, N] = A[M, K] × B[K, N]

CPU：把 loop 切成向量長度

RVV kernel 會沿著某個矩陣維度做 strip-mining

for M/N/K blocks:
    vl = vsetvl(remaining)
    vector load A/B
    vector multiply-accumulate
    vector store C

Compiler 會考慮 loop order、vector length、LMUL、register pressure 與 cache blocking。資料仍經由一般 load/store 存取共享位址空間，cache 是否命中會影響效能。

GPU：把 output tile 分給 CTA 與 warp

GPU matmul 常把一塊 C 分給一個 CTA，再把更小的 fragment 分給 warp

HBM/global memory
  → A/B tile 搬入 shared memory
  → warp 讀取 fragment
  → Tensor Core MMA
  → C tile 寫回 global memory

Triton 的 TTGIR layout 會描述 tensor element 如何分給 CTA、warp 與 lane。Compiler 還要安排 shared-memory layout、async copy、barrier 與 MMA instruction。GPU 硬體用大量 resident warp 隱藏 latency，occupancy 與 register/shared-memory 用量會互相牽制。

Tenstorrent：把 output tile 放到實體 Tensix 座標

Tenstorrent 的資料路徑可以寫成

DRAM tile
  → reader 經 NoC 把 A/B tile 搬到各 core 的 L1
  → Tensix Matrix Unit 計算
  → writer 經 NoC 寫回 DRAM

若同一塊 A 被一列 core 使用，程式可以把它沿著 NoC broadcast；B 則可能沿另一個方向傳送。Compiler 因此要知道 logical M/N tile 對應哪個 (x, y) 座標，才能判斷資料能否共用。

GPU 和 Tenstorrent 都是 tile-based，差別在哪裡

兩者都會把 tensor 切成 tile，也都要使用鄰近運算單元的快速記憶體。差異主要落在硬體向軟體暴露的排程方式。

NVIDIA GPU 的程式先提交 grid。硬體把 CTA 排到可用 SM，warp scheduler 再從 ready warp 中發射指令。Kernel 可以控制 block 內的 shared memory、同步與 pipeline，但一般不會把某個 CTA 永久綁定到指定 SM 座標，再規劃 SM 之間傳輸。

Tenstorrent 程式會直接面對 2D Tensix grid。資料位於各 core 的 L1，跨 core 傳輸走 NoC，因此 core coordinate、來源與目的區域會進入程式和 IR。Day28 看到的：

area : [1, 8]
region : (UL : [...], LR : [...])

就是這種 spatial data movement 的表示。

可以把兩種 compiler 問題並排

GPU:
tile → CTA/warp/lane layout
     → shared-memory layout
     → async pipeline
     → MMA

Tenstorrent:
tile → Tensix (x, y) placement
     → per-core L1
     → NoC0／NoC1、傳輸端點與 broadcast 規劃
     → Matrix Unit

這裡的 compiler 決策是選擇 NoC0／NoC1、來源與目的座標、unicast／broadcast 區域及傳輸時機。封包經過哪些 router，則由 NoC topology 與 routing 規則決定，不能把它理解成 compiler 任意指定每一跳。

所以 TileLoom 關注的 dataflow planning，會比前面 Triton NVIDIA flow 多出跨 core placement 與顯式 broadcast。GPU compiler 同樣重視 data movement，只是硬體抽象與可控制的空間不同。

三種硬體對 AI Compiler 提出的問題

把三條路線放在一起，compiler optimization 的重心會跟著硬體改變

CPU:
  哪個 loop 要向量化？
  RVV 的 vl、SEW、LMUL 怎麼選？
  tile 能否留在 cache？

GPU:
  tensor tile 如何分給 CTA、warp、lane？
  global → shared 的搬移如何和 compute 重疊？
  哪種 layout 能餵給 Tensor Core？

Tenstorrent:
  logical tile 要放到哪個 Tensix core？
  資料應放在哪顆 core 的 L1？
  哪些資料值得經 NoC broadcast？

後面講 Triton on Tenstorrent 時，這些問題會進入具體 IR。Triton kernel 先描述 tile-level program，TileLoom 再嘗試決定 tile instance 的硬體位置、資料搬移與 reuse。

今天先走到這裡

今天先記住幾個 Tenstorrent 的關鍵名詞

• Wormhole 和 Blackhole 延續 Tensix grid、本地 L1 與顯式資料搬移的設計。
• Wormhole 是由多種 tile 組成的 ASIC，Tensix tile 負責主要 compute。
• DRAM tile 是 GDDR6 的 NoC 入口；Ethernet tile 處理多 ASIC 連線；PCIe tile 連 host；ARC tile 負責管理。
• 兩個 NoC 把 tile 接成 2D torus，讓資料搬移成為 compiler / runtime 必須面對的問題。
• GPU 與 Tenstorrent 都採用大量平行運算和 tile-based compute；GPU 著重 CTA/warp、shared memory 與 MMA，Tenstorrent 額外把 core coordinate、L1 placement 與 NoC movement 暴露給軟體。
• 後面看 TT-Metal 時，會看到 reader、compute、writer kernel 如何在 Tensix core 上協作。
• 後面看 TileLoom 時，會看到 compiler 怎麼把 Triton / Helion 的 tile program 放到這種 2D grid 上，到那時候再細講 TileLoom 的 pipeline。

明天來看 TT-Metal，今天看的是硬體地圖，明天看程式要怎麼描述搬資料、算、寫回去。

參考資料

• Tenstorrent tt-isa-documentation: WormholeB0 README
  https://github.com/tenstorrent/tt-isa-documentation/blob/main/WormholeB0/README.md
• Tenstorrent tt-isa-documentation: WormholeB0 NoC
  https://github.com/tenstorrent/tt-isa-documentation/blob/main/WormholeB0/NoC/README.md
• Tenstorrent Blackhole PCIe Cards
  https://docs.tenstorrent.com/aibs/blackhole/
• Tenstorrent TT-Metalium Tiles
  https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/tt_metal/advanced_topics/tiles.html
• Tenstorrent Hardware Overview
  https://docs.tenstorrent.com/docs-test/pybuda/latest/hardware.html
• Tenstorrent TT-Metalium Guide
  https://github.com/tenstorrent/tt-metal/blob/main/METALIUM_GUIDE.md
• Tenstorrent YouTube
  https://www.youtube.com/@Tenstorrent
• NVIDIA CUDA C++ Programming Guide
  https://docs.nvidia.com/cuda/cuda-c-programming-guide/