前十天一直盯著 RISC-V CPU、RVV intrinsic、SGLang RVV 後端和 PyTorch Inductor RVV，滿腦子都是 CPU core、vector register、cache 和 C++ kernel，今天終於要換到 GPU 了。

今天先從工作組織、硬體執行與資料存取三個角度拆解 NVIDIA GPU。後面看到 warp、SM、shared memory、Tensor Core、TMA、WGMMA 時，才有辦法把這些硬體名詞接回 compiler lowering 和 runtime launch，不然名詞很多很容易全部混在一起。

我這次會沿著 NVIDIA CUDA Programming Guide 與官方架構資料，以 Hopper、Blackwell 與 Vera Rubin 三代比較。Vera Rubin 部分會再區分 Rubin GPU 與整體平台，最後回到 DGX Spark，我們之後實驗會在上面跑。

本篇大綱

• 先用比較 CPU 和 GPU 架構
• 接著從工作組織、硬體執行、資料存取三個角度，串起 grid、block、warp、thread 與 SM
• 再把三個角度合在同一個 SM 裡，看 scheduler、register、shared memory、Tensor Core、L2、HBM 怎麼合作
• 接著比較 Hopper、Blackwell 與 Vera Rubin 三代主線的資料搬移、Tensor Core 與記憶體設計，並看 Vera CPU 如何承接 agentic AI 的 CPU 工作
• 再看後續實驗平台 DGX Spark 的 GB10、Arm CPU、統一記憶體、sm_121

先比較 CPU 和 GPU

這是概念圖，圖中的 core、SM 和快取數量不代表某一款處理器的實際規格。

之前看 CPU 的時候，我的思考方式比較像這樣

少量強核心
-> 每個 core 有複雜控制邏輯、cache hierarchy、branch predictor
-> 用 SIMD 把單一 core 的資料平行度拉高
-> 盡量讓 hot loop 吃滿 vector register 和 cache

CPU 很擅長處理控制流程複雜、分支多、資料相依性強的工作。像 runtime control、operator dispatch、系統呼叫、資料前處理，CPU 的強項是低延遲和彈性。就算資料不完全規則，CPU 也比較能靠 branch predictor、out-of-order execution、cache 做補救。

GPU 的方向不同，GPU 把大量電晶體放在平行運算單元和高頻寬記憶體上，適合處理大量相同形狀的資料工作

大量較小的執行單元
-> 用 warp / thread block 管理平行工作
-> 用 shared memory / register 做 tile 內資料重用
-> 用 HBM、GDDR 或高頻寬統一記憶體供應資料
-> 用 Tensor Core 加速矩陣乘法

所以 CPU 和 GPU 寫 kernel 時，在意的方向不太相同

從三個角度理解 GPU kernel 執行

GPU 這裡有很多名詞容易搞混，讓我們從三個角度拆解 GPU kernel 執行

如果名詞混在一起，先問它是工作、硬體，還是資料？例如 block 是一組工作，SM 是執行工作的硬體，shared memory 是這組工作在 SM 上使用的共用資料區。

CPU 啟動 kernel
  ↓
產生一個 grid，grid 裡有很多 blocks
  ↓
每個 block 被分派到某個 SM
  ↓
block 裡的 threads 以 warp 為群組執行
  ↓
SM 的 warp scheduler 選擇 ready warp，發射下一道指令
  ↓
指令使用 register / shared memory，或經 L1 / L2 存取 global memory

block 被分派到 SM 是三個角度之間重要的接點，下面會沿著這條主線依序展開。

先看工作組織：程式如何切分 kernel

CUDA kernel 啟動時，工作會組織成 grid -> blocks -> threads，NVIDIA GPU 執行時還會把同一 block 內的 threads 組成 warps

例如這個啟動設定

kernel<<<80, 256>>>(...);

它會產生 1 個 grid，其中有 80 個 blocks，每個 block 有 256 個 threads。一個 block 的 256 個 threads 會形成 8 個 warps，因為 256 / 32 = 8。

這 80 個 blocks 不會各自綁定一個專屬 SM，GPU 會把 blocks 分派給當下有資源的 SM，前一批 blocks 完成後，同一批 SM 還可以繼續執行剩下的 blocks。

進入一個 SM 後，warp scheduler 會從已駐留的 warps 中挑選當下可執行的 warp。CUDA 的 SIMT programming model 讓程式看起來像每個 thread 各自執行 scalar code，實際發射指令時則以 warp 為群組。同一 warp 內的 threads 走不同 branch 會產生 divergence，讀取連續 global-memory 位址時，則較容易合併成有效率的記憶體交易。

再看硬體執行：這些工作被送到哪裡

前一節描述工作如何組織，現在再把它放到實體 GPU 裡

grid  → 其中的 blocks 可分散到 GPU 上多個 SM
block → 執行時駐留在一個 SM
warp  → 由 SM 內的 warp scheduler 排程
thread → 在 warp 中以一個邏輯執行個體的角度執行程式

一個 thread 不會從 kernel 開始到結束，始終獨佔某一個 CUDA core。
CUDA core、Tensor Core、load/store unit 是執行不同類型指令的硬體單元，warp scheduler 會依指令類型和資源狀態發射工作。

如果把一顆 NVIDIA GPU 從外到內拆開，可以用這張簡化圖對照

Host CPU
  |
PCIe / NVLink
  |
GPU package
  |-- HBM / GDDR memory
  |-- Memory controllers
  |-- L2 cache
  |-- GPC / TPC / SM groups
        |-- SM
             |-- warp schedulers
             |-- dispatch units
             |-- register file
             |-- shared memory / L1
             |-- CUDA cores
             |-- Tensor Cores
             |-- load/store units
             |-- special function units

不同 NVIDIA 世代的實際命名和數量會變，後面會再針對各元件仔細說明。

最後看資料存取：誰能使用和資料在哪裡

知道 block 會駐留在一個 SM 後，記憶體名詞就有了參考點：thread 執行指令需要私有資料，同一 block 的 threads 需要一塊合作區域，所有 blocks 還需要存取大容量的輸入與輸出。

因此，CUDA 討論記憶體時，要把資料存取和前兩個角度連起來

thread 私有資料 → register，或 spill 到 local memory
block 合作資料  → 該 block 駐留 SM 內的 shared memory
所有 threads 存取 → global memory，途中可能經過 L1 / L2 快取

這裡又要再分成兩個問題，看到一筆資料時可以問

1. 哪些 thread 可以用這個名稱或位址存取資料？ 這是程式模型中的可見範圍與生命週期。例如 register 和 local memory 歸單一 thread 使用，shared memory 供同一 block 的 threads 合作，global memory 則可由整個 grid 存取。
2. 這次 load / store 實際走過哪些硬體？ 這是實體位置與存取路徑。資料可能留在 SM 的 register file 或 shared memory，也可能從 HBM / GDDR 出發，途中經過 L2 和 L1 快取。

圖片來源：NVIDIA CUDA C++ Programming Guide, Figure 6: Memory Hierarchy

第一次看可以先記住 thread → register / local、block → shared、grid → global 。
下面是完整版，可以多看幾次不需要一次記住很多細節，這裡名詞太多了。

Host 連到 GPU：PCIe / NVLink

GPU 不是單獨存在，程式通常從 CPU 端啟動 kernel，資料也常常要從 host memory 搬到 device memory。桌機或一般伺服器常見是 PCIe，大型 NVIDIA 系統會看到 NVLink / NVSwitch。
這裡影響的是 host-device transfer 成本，如果資料一直在 CPU 和 GPU 之間來回搬，就算 kernel 本身很快，整體時間也會被 transfer 吃掉。
所以 benchmark 需要區分只量 kernel time，還是量 host-to-device + kernel + device-to-host。

容量、延遲與頻寬是三個不同的問題

記憶體階層我喜歡從容量、延遲與頻寬來觀察

• 容量：模型 weight、KV cache 和中間 tensor 能不能放進 GPU 記憶體。
• 延遲：單次讀取要等多久。GPU 常靠多個 resident warps 和非同步 pipeline 隱藏延遲。
• 頻寬：持續串流讀寫每秒能搬多少資料。Decode GEMV、attention 與 KV-cache 讀取常會先碰到頻寬上限。

算術密度（arithmetic intensity）是連接運算與記憶體的簡單方法

算術密度 = 運算量 / 從記憶體搬移的資料量

同一個 weight tile 如果只用一次，kernel 很容易被 HBM 頻寬限制，搬到 shared memory 或 register 後重用多次，算術密度才會上升。這也是 prefill 的大型 GEMM 較容易吃滿 Tensor Core，小 batch decode 卻常受 weight 與 KV cache 搬移限制的原因。

HBM / GDDR：GPU 的主要外部記憶體

資料進到 GPU 後，主要會放在 device global memory。資料中心 GPU 常見 HBM，消費級或部分工作站卡會看到 GDDR。HBM / GDDR 的頻寬很高，但延遲仍然比 register / shared memory 大很多。

所以 GPU kernel 常見策略是

從 global memory 讀一大塊
-> 搬到 shared memory 或 register
-> 在近的地方重用多次
-> 最後再寫回 global memory

之後會談到的 Triton matmul 裡的 BLOCK_SIZE_M/N/K，本質上就是在決定一次搬多少 A/B tile，讓資料在 register / shared memory 裡被重用。

L2 cache：跨 SM 的資料重用

L2 cache 位在 SM 和外部記憶體之間，它不是某個 thread block 私有的，會被多個 SM 共享。

這對 matmul 很重要，假設不同 program instance 會讀到相同的 A tile 或 B tile，program order 會影響 L2 hit。這也是 Triton matmul tutorial 裡 GROUP_SIZE_M 會出現的原因：把相關 output tile 的計算順序排近一點，讓剛讀過的資料比較可能還在 L2 裡。

SM：GPU kernel 的主要執行單位

Streaming Multiprocessor (SM) 是 CUDA kernel 執行時重要的硬體單位。可以先把它想成負責執行 thread block 的工作區，一個 GPU 有很多 SM，kernel launch 後，thread blocks 會被分派到不同 SM 上跑。

一個 SM 裡面通常會有

• warp schedulers：決定下一個要發射哪個 warp 的指令。
• dispatch units：把指令送到對應 execution units。
• register file：每個 thread 使用的 register 都從這裡分配。
• shared memory / L1：同一個 block 內 threads 可以合作使用的高速記憶體。
• CUDA cores：處理一般 FP32 / INT 等運算。
• Tensor Cores：處理矩陣乘加，AI kernel 很仰賴它。
• load/store units：處理 memory access。
• special function units：處理一些特殊數學函數。

Hopper H100 的官方 SM 圖可以把上面的清單對回實體區塊

圖片來源：NVIDIA Hopper Architecture In-Depth, Figure 4。圖中一個 SM 分成四個處理分區，每區有 warp scheduler、dispatch unit、register file、運算單元與第四代 Tensor Core；下方是 TMA 與共用的 L1 data cache / shared memory。

這裡直接決定 occupancy，當一個 block 使用太多 register 或 shared memory，同一個 SM 能同時駐留的 block / warp 數量就會下降。反過來，如果 block 太小或 work 太少，也可能吃不滿 SM。

Warp scheduler：GPU 怎麼藏 memory latency

GPU global memory latency 很高，GPU 常用的做法是讓很多 warp 同時駐留在 SM 上，當某個 warp 等 memory，scheduler 可以切到另一個 ready warp 執行。

所以 GPU 最佳化裡常看到 occupancy，它不是越高一定越好，但太低會讓 GPU 沒有足夠 warp 可以切換，memory latency 比較容易暴露出來。

Register file：最快，但很容易變成壓力

Register 是 thread 私有的高速儲存空間。Triton matmul 裡的 accumulator 通常會放在 register。BLOCK_SIZE_M/N 變大時，accumulator 變大，register 用量也會上升。

register 用太多會讓同一個 SM 能放的 warps 變少。大 tile 可能讓單次 tl.dot 處理更多資料，也可能讓 ptxas 配出更多 register，最後壓低 occupancy。

Shared memory / L1：block 內合作的地方

Shared memory 是同一個 thread block 內 threads 可以共同使用的高速記憶體。它常被用來放 A/B tile、做 data staging、降低 global memory 重複讀取。

對 matmul 來說，常見流程是

global memory 讀 A/B tile
-> 放到 shared memory
-> Tensor Core / CUDA cores 重複使用
-> 沿 K 維度處理下一個 tile

shared memory 也有自己的限制，例如容量、bank conflict、同步成本。Triton 和 CUTLASS 的 pipeline 設計會一直碰到這些問題。

Tensor Core：AI kernel 的主要加速器

Tensor Core 是 NVIDIA GPU 上專門處理矩陣乘加的硬體。對 AI workload 來說，GEMM、attention、convolution、MLP projection 都會盡量接到 Tensor Core。

在 Triton 裡，tl.dot 是很重要的入口。前端看起來是一個 tile-level dot，後端會依照 GPU 架構和 dtype，把它 lowering 到 MMA、WGMMA 或其他 target-specific path。後面談 NVIDIA TTGIR 時，我們會再把這條路徑拆開。

一次 kernel 執行的資料路徑

把上面合起來，可以用一個簡化版 matmul path 來看

CPU launch kernel
  ↓
thread blocks 分派到 SM
  ↓
warp scheduler 發射 warps
  ↓
load/store units 從 global memory / L2 讀 A/B tile
  ↓
shared memory / registers 暫存資料
  ↓
Tensor Cores 執行 tl.dot 對應的矩陣乘加
  ↓
accumulator 留在 registers
  ↓
結果寫回 global memory

後面讀 TTGIR、PTX 和 ptxas 資源資訊時，很多參數都可以放回這條路徑理解

• BLOCK_SIZE_M/N/K：一次處理多大的 tile。
• num_warps：一個 program instance 用多少 warp 合作。
• num_stages：資料搬移和計算 pipeline 多深。
• dtype：能不能走 Tensor Core 的快路徑。

三代架構：先看 compute capability

NVIDIA 用 compute capability 表示 GPU 支援的硬體功能和指令能力，格式是 major.minor。例如 H100 的 compute capability 是 9.0，編譯目標就常寫成 sm_90。

這裡的 SM 有兩種用法，我之前第一次讀不小心搞錯

• 硬體單位：Streaming Multiprocessor。一張 H100 有很多個 SM，thread block 會被分派到這些 SM 執行。
• 架構目標：sm_90、sm_100 這類名稱，表示 CUDA 要產生哪一版實體 GPU ISA。sm_90 不是有 90 個 SM。

CUDA 編譯時還會看到 compute_90。它是產生 PTX 時使用的虛擬架構，sm_90 則是 ptxas 產生 cubin 時使用的實體架構。fat binary 可以同時放多個 cubin 和一份 PTX，執行時由 driver 選擇相符版本，或用 PTX JIT 補上未預先編譯的目標。

# 產生 Hopper H100 的 sm_90 cubin
nvcc kernel.cu -arch=sm_90

# 查看機器上的 GPU 型號與 compute capability
nvidia-smi --query-gpu=name,compute_cap --format=csv

我對最近三代的 GPU 比較有興趣，先來看資料中心等級 GPU 官方公開的效能規格

GPU 記憶體容量和頻寬會依 SKU 與封裝而變，表格不能代替實際部署型號的資料表。NVLink 數字也是 GPU 到 scale-up fabric 的雙向峰值，不等於任意 collective 都會達到這個吞吐量。

Compute capability 主要是指令與硬體功能，實際產品規格則回答 SM 數量、register、shared memory、L2、記憶體容量與頻寬。
Rubin 的 sm_107 來自預覽版工具鏈文件，CUDA 13.4 Developer Preview 也是預先發行軟體，發布效能當作是參考。

Hopper：TMA 和 WGMMA

H100 / H200 對應 Hopper sm_90。以 H100 SXM 為例，它有 132 個 SM、每個 SM 128 個 FP32 CUDA cores 和 4 個第四代 Tensor Core，另有 50 MB L2、每個 SM 最高 228 KB shared memory，以及約 3 TB/s 的 HBM3 頻寬。H100 PCIe 啟用的 SM 較少，記憶體也改為 HBM2e，H200 沿用 Hopper SM90，主要把記憶體升級為更大容量與更高頻寬的 HBM3e。

圖片來源：NVIDIA Hopper Architecture In-Depth, Figure 3。這是完整 GH100 設計的全晶片圖，共有 8 個 GPC、72 個 TPC 與 144 個 SM；實際 H100 SXM5 啟用 132 個 SM，H100 PCIe 啟用 114 個 SM。

閱讀這張圖時可以從外往內看，晶片邊緣是 HBM controller、NVLink 與 PCIe 介面，中央藍色區塊是分區的 L2 cache，大量重複的綠色區塊則是組成 GPC / TPC 的 SM。資料從 HBM 進入 GPU 後，會經 L2 抵達 SM，再由前面的 SM 內部圖繼續拆成 shared memory、register、CUDA Core 與 Tensor Core。

第四代 Tensor Core 加入 FP8，Transformer Engine 會在 FP8 與較高精度格式之間管理 scaling。對 LLM 訓練與推論來說，Hopper 的硬體主線可以看成兩部分，TMA 負責更有效率地餵資料，WGMMA 負責更大粒度的矩陣乘加。

H100 和 H200 它們都是 Hopper SM90，H200 的主要改變是把 HBM 從 H100 SXM 的 80 GB HBM3、3.35 TB/s，提高到 141 GB HBM3e、4.8 TB/s。同一代指令架構下，只改記憶體容量和頻寬，就可能改變模型能否單卡放入、KV cache 上限與 memory-bound kernel 的吞吐量。

TMA：把多維 tensor 搬移交給專用硬體

Tensor Memory Accelerator (TMA) 是 Hopper 的非同步資料搬移單元。程式先建立 tensor map，描述 global memory tensor 的 shape、stride、tile 與 swizzle，再由少數 thread 發起一維到五維的 transfer。TMA 可以做 global memory 與 shared memory 之間的搬移，也能配合 multicast 把資料送到 Thread Block Cluster 中多個 block 的 shared memory。

TMA 的關鍵是讓專用單元處理多維位址與大塊 transfer，其他 threads 可以繼續計算。搬移完成後，TMA 會透過 asynchronous transaction barrier 通知 consumer，所以 pipeline 必須一起安排 barrier、producer 與 consumer。

Hopper 也加入 Thread Block Cluster。cluster 內的 blocks 可以保證同時排程，並透過 Distributed Shared Memory 存取彼此的 shared memory。這讓 TMA multicast 和跨 block 合作成為可用的 kernel 設計選項。

WGMMA：由 4 個 warp 合作做非同步矩陣乘加

WGMMA 的 PTX 指令是 wgmma.mma_async。一個 warp group 由 4 個連續 warp 組成，也就是 128 個 threads。它比傳統 warp-level mma.sync 使用更大的合作範圍，矩陣 operand 會依指令形式來自 registers 或 shared memory，accumulator 則分散在 warp group 的 registers。

名稱裡的 async 也會反映在同步規則。程式要用 wgmma.fence、wgmma.commit_group、wgmma.wait_group 管理非同步 MMA group；shared memory layout、swizzle 和 matrix descriptor 如果不符合指令要求，WGMMA 就不能正確取資料。

這裡還要分清楚 sm_90 和 sm_90a。H100 / H200 的 compute capability 都是 9.0；使用 WGMMA 等 architecture-accelerated features 的 CUTLASS Hopper kernel，需要用 90a target 才能啟用這些功能。一般 SM90 baseline code 仍可使用 sm_90。a 代表程式使用特定架構才有的功能，沒有一般 PTX 的跨世代 forward compatibility 保證。

一個典型 Hopper matmul pipeline 可以簡化成

producer warp：用 TMA 把下一個 A/B tile 搬進 shared memory
                         ↓ transaction barrier
consumer warp group：用 WGMMA 計算目前 tile
                         ↓ commit / wait group
accumulator：留在 registers，最後做 epilogue 並寫回

對 Triton / CUTLASS 這類 compiler 或 template library 來說，Hopper 的重點是

• copy 可以由 TMA 解讀 tensor descriptor，不一定由每個 thread 自己算 address。
• matmul 可以由 4 個 warp 合作執行 WGMMA，不只使用 warp-level MMA。
• pipeline 設計要一起考慮 TMA barrier、warp specialization、shared memory layout 和 WGMMA async group。
• 228 KB shared memory 讓較深的 pipeline 成為可能，但 tile、stage 與 occupancy 仍會互相牽制。

Hopper 對 LLM serving 的實際影響

• Prefill 有較大的 M 與 N，容易組成大型 GEMM。TMA 與 WGMMA 的 pipeline 重點是持續供應 Tensor Core，避免運算單元等資料。
• Decode 的 batch 小、M 也小，weight 與 KV cache 的讀取量相對突出。H200 的較大、較快 HBM3e 在這類工作負載更容易發揮作用。
• Kernel 整合 可減少中間 tensor 寫回 HBM 再讀回的次數。例如把 bias、activation 或量化 scaling 放進 GEMM epilogue，節省的常是記憶體流量，不只是 kernel launch。

所以後面講 TTGIR Nvidia 時，不能只停在把 tl.dot 變成 MMA，不同 NVIDIA 架構上的 lowering 會受 Tensor Core 指令形狀和資料搬移能力影響。

Blackwell：tcgen05 和 Tensor Memory

這裡把 Blackwell 聚焦在 B100、B200 與 GB200 這條資料中心路線，Blackwell 由兩個 reticle-limited dies 組成，中間用 10 TB/s NV-HBI 連接，但 CUDA 程式仍把它視為一個 GPU。B200 對應 compute capability 10.0 / sm_100，支援最高 228 KB shared memory per SM 與第五代 NVLink。B200 有 180 GB HBM3e 與最高 8 TB/s 頻寬，GB200 的 L2 容量則增加到 126 MB。

NVIDIA 官方公開的 Blackwell Ultra SM 圖很適合用來找 TMEM 的位置

圖片來源：NVIDIA Blackwell Ultra Architecture, Figure 2。這是 Blackwell Ultra 的 SM 圖，用來說明 Blackwell 世代的 TMEM、第五代 Tensor Core 與 SM 內部關係；圖中數量不應直接套到 B200 SKU。

SM100 的 Tensor Core 指令族稱為 tcgen05，它也引入幾個會改變 kernel 設計的概念

• Tensor Memory（TMEM）：accumulator 可放在晶片內的專用記憶體，降低 register file 壓力。
• CTA-pair cooperation：兩個相鄰 CTA 可合作執行 MMA，擴大一次處理的 tile。
• 低精度和 block scaling：在 FP16、BF16、TF32 與 INT8 之外，進一步支援 NVFP4 等低精度格式。
• 新的 issue model：硬體指令的發出粒度與 Hopper WGMMA 不同，compiler 和 CUTLASS 必須選擇對應的 SM100 schedule。

TMEM 需要單獨說明，它不是一塊可以用一般 load / store 隨意操作的 shared memory，也不是 L1 / L2 快取。Blackwell 的 tcgen05 指令會把矩陣乘加結果寫入 TMEM，後續指令再將結果取出做 epilogue。這是大型 accumulator 不必全部長時間佔用一般 register file。

HBM / L2
  ↓
shared memory 中的 A/B tiles
  ↓ tcgen05 MMA
Tensor Memory 中的 accumulator
  ↓ epilogue 需要時讀出
registers
  ↓
global memory output

對 compiler 來說，這多了 TMEM 的取得、使用和釋放，也多了 tcgen05 的 commit / wait 和 CTA 合作規則。如果只把 Hopper WGMMA 指令名稱替換掉，沒有重新設計 accumulator 和 pipeline，就還沒有用到 Blackwell 的主要變化。

Blackwell 對 LLM serving 的實際影響

• NVFP4 和 microscaling 可大幅減少 weight 記憶體用量與每個 token 需搬移的位元組數，但精度與 scaling recipe 必須由模型與軟體共同配合。
• TMEM 讓 GEMM 的 accumulator 生命期和一般 registers 分開，有機會改善大 tile 下的 register pressure，但 kernel 也多了專用狀態與同步規則。
• 180 GB HBM3e、8 TB/s 與 126 MB L2（GB200） 同時改變容量與資料重用條件。Decode 是否加速，仍取決於 batch、量化格式、KV-cache 流量和 kernel 是否真正使用 SM100 路徑。

從 Hopper 到 Blackwell，變化不只是 Tensor Core 峰值提高。Accumulator 的放置、MMA 指令形式、資料精度與跨 GPU 連線都改變了，因此 Hopper 的 WGMMA kernel 不會自動變成最適合 Blackwell 的 kernel。

Vera Rubin：Rubin GPU 加上 Vera CPU 的平台

這裡先分清名稱：Rubin 是 GPU 架構，Vera 是 CPU，Vera Rubin 是把兩者與 NVLink 6、ConnectX-9、BlueField-4 等元件整合在一起的平台。Vera Rubin NVL72 則是包含 72 顆 Rubin GPU 與 36 顆 Vera CPU 的機架級系統。

先看 agentic AI 為什麼重新增加 CPU 壓力

一般 LLM 推論可以簡化成 CPU 準備請求，GPU 連續產生 token。Agentic AI 的一次請求會在 GPU model call 之間插入許多 CPU 與 I/O 工作：

GPU：模型推論，決定下一個動作
  ↓
CPU：解析輸出、更新 agent 狀態、執行排程與安全檢查
  ↓
CPU / I/O：呼叫工具、查詢資料庫、檢索文件、執行 Python 或 shell sandbox
  ↓
CPU：整理工具結果、序列化資料、組合下一段 prompt
  ↓
GPU：帶著新 context 繼續推論

強化學習的 rollout 也有相似情況。GPU 產生 action 之後，大量 CPU 環境要執行 action、計算 reward、跑測試或模擬器，再把結果送回下一輪訓練。只要其中一批環境太慢，GPU 就可能等資料，或訓練迭代必須排除來不及完成的樣本。

這類工作有幾個共同特徵，單一 agent loop 內有序列相依性，Python、JavaScript、資料庫與作業系統呼叫帶來大量分支，pointer-heavy 資料結構和不規則存取不容易轉成 GPU 大型矩陣運算；同一台機器還要同時處理成百上千個 sandbox、工具呼叫和 agent 狀態，CPU 需要兼顧單一工作的完成時間，以及滿載時大量平行工作的尾端延遲。

因此，針對 agentic AI 的 CPU可以拆成四個具體問題：

1. 一條分支多、難向量化的 agent 執行路徑，能不能快速往前走？
2. 整顆 CPU 滿載時，每個 agent 還能不能取得足夠的記憶體頻寬？
3. 大量 sandbox 同時執行時，延遲會不會因資源競爭而突然拉長？
4. CPU 完成工具工作後，能不能快速把狀態與資料交回 GPU？

Vera 的設計正好沿著這四個問題展開。

Vera CPU：Olympus 核心、Spatial Multithreading 與 SCF

Vera 是 NVIDIA 自行設計 CPU 核心的資料中心處理器，一顆 CPU 有 88 個 Armv9.2 相容的 Olympus 核心。每個核心可透過 NVIDIA Spatial Multithreading 執行兩個硬體執行緒，整顆 CPU 合計 176 個硬體執行緒。

Olympus 把單一執行緒效能放在很前面，原因來自 agent loop 的序列相依性。模型要先提出工具參數，CPU 才能執行工具，工具結果回來後，runtime 才能組下一個 prompt。這段 critical path 不能單靠增加更多核心縮短。NVIDIA 公開的微架構資料包含 10-wide decode、neural branch predictor、較深的亂序執行，以及處理間接存取的 graph prefetcher，目標都是讓分支多、存取不規則的程式維持較高 IPC。

Spatial Multithreading 則處理另一面，一座 AI 工廠會同時跑大量 agent 和強化學習環境。每個 Olympus 核心可以在需要最低延遲時把較多資源留給單一硬體執行緒，也能將核心資源分給兩個硬體執行緒，提高滿載時的環境密度。

記憶體與 coherency fabric

88 個快速核心如果共用一般伺服器 DRAM 頻寬，大量 sandbox 同時掃描狀態、解壓縮資料、跑資料庫或整理檢索結果時，很容易讓核心等記憶體。Vera 使用 SOCAMM2 LPDDR5X，每個 socket 提供最高 1.5 TB 容量和 1.2 TB/s 頻寬。LPDDR5X 在這裡服務的是資料中心 CPU，SOCAMM2 採模組化設計，讓記憶體仍可維修與更換。

第二代 NVIDIA Scalable Coherency Fabric（SCF）把 88 個核心、164 MB 統一 L3、記憶體控制器、I/O 與 NVLink-C2C 接在單一 compute die 上。NVIDIA 公開的最高雙向剖分頻寬為 3.4 TB/s。單一 compute die 可避免核心與共用資料跨越多個 CPU compute chiplet 時產生的額外 hop，對 agent 服務來說，設計目標是讓滿載時的 core-to-core 與 core-to-memory 延遲比較一致。

NVLink-C2C 把 CPU 工作接回 Rubin GPU

在 Vera Rubin 系統裡，Vera CPU 與 Rubin GPU 透過第二代 NVLink-C2C 連接，提供最高 1.8 TB/s 的 cache-coherent CPU–GPU 頻寬。它的工作是把前面那段 agent loop 接起來：GPU 做模型推論，CPU 執行工具與 orchestration，接著把結果交回 GPU。

Coherent link 也讓 KV cache offload、共享資料結構與 CPU–GPU pipeline 有更大的設計空間。舉例來說，agent 等待較慢的工具回傳時，系統可能把暫時不用的 KV cache 從 GPU HBM 移出；工具完成後若要還原該對話，CPU–GPU 連線與記憶體管理會影響還原成本。這不是 NVLink-C2C 自動解決的問題，runtime 仍要決定哪些資料保留、刪除和預先搬回。

Vera 的硬體選擇可以整理成一條因果鏈

agent / RL 加入大量 CPU critical path
  ↓
Olympus 提高分支多、序列性工作的單執行緒進度
  ↓
Spatial Multithreading 提高大量 sandbox 的平行密度
  ↓
SCF + 164 MB L3 + 1.2 TB/s LPDDR5X 避免滿載時只增加等待記憶體的核心
  ↓
NVLink-C2C 將 CPU 工具工作與 Rubin GPU 模型運算接回同一條 pipeline

NVIDIA 2026 年 7 月公開的 Rubin GPU 資料列出

Rubin GPU：224 SM、896 Tensor Cores
記憶體：最高 288 GB HBM4
記憶體頻寬：最高 22 TB/s
Transformer Engine：第三代
NVLink 6：每顆 GPU 最高 3.6 TB/s 雙向 scale-up 頻寬

圖片來源：NVIDIA Rubin GPU Architecture, Figure 2。這是全 GPU 層級的圖，可看到兩個 compute dies、GPC、中央 L2、HBM controller、NV-HBI、NVLink 6 與 PCIe Gen 6；它不是單一 SM 的內部圖。

從圖中可以讀到 Rubin 的設計方向：運算單元增加後，HBM controller、L2、die-to-die link、CPU-GPU 與 GPU-GPU 連線也一起擴充。否則 224 個 SM 與 896 個 Tensor Cores 只會更快地等待資料。

這一代的 kernel 重點開始明確對準 MoE、長上下文 attention 與多 GPU agentic inference

• 加強版 TMA 支援 inline descriptor update。多個 MoE experts 若共用 layout，kernel 可在 TMA 指令中更新 pointer 與 stride，減少 descriptor 管理成本。
• Tensor Core 每個 clock 可處理更大的 K 維度，讓高 tensor parallelism 下的 context 與 decode GEMM 減少 K-loop 次數。
• Activation sparsity（啟動值稀疏性）與 adaptive compression（自適應壓縮） 可減少 attention 中間資料的儲存與後續運算，同時增加 softmax 需要的指數運算吞吐量。
• 更細粒度的 dependent-kernel triggering 讓 consumer kernel 在所需 tile 就緒後更早開始，減少 producer 與 consumer 之間的閒置空檔。
• NVLink counted writes 讓接收端用 counter 追蹤資料搬移完成，減少多 GPU producer / consumer 流程中的 barrier、acknowledgement 與 atomic flag 協調。

Rubin 對 LLM serving 的實際影響

• 長上下文 attention 會掃描大量 KV cache。288 GB HBM4 提供容量，22 TB/s 提供搬移速度，adaptive compression 則嘗試減少中間資料流量。
• MoE 會在不同 expert weight 之間切換。TMA inline descriptor update 可在 layout 相同時更新 pointer 和 stride，降低為每個 expert 管理複製 descriptor 的成本。
• 高 tensor parallelism 會把每顆 GPU 的 M / N 工作切小，K 維度仍可能很大。Rubin 每 clock 處理更大 K 維度，就是直接對應這個 GEMM 形狀。
• Agentic inference 包含較多短 kernel、工具呼叫與跨 GPU 相依工作。Dependent-kernel triggering 和 counted writes 著重減少工作之間的等待，不只是提高單一 GEMM 峰值。

DGX Spark：Grace Blackwell 桌上型系統

這個系列後續的 Triton 實驗，都會以 NVIDIA DGX Spark 為主要平台。

DGX Spark 是採用 GB10 Grace Blackwell Superchip 的完整桌上型系統。GB10 把 20 核心 Arm CPU、Blackwell GPU、記憶體控制器與其他 I/O 整合在同一個 SoC，和常見的 x86_64 主機加 PCIe 獨立顯示卡有明顯差異。

圖是依照 NVIDIA DGX Spark Hardware Overview 與 DGX Spark Porting Guide：System Architecture 重新畫的。綠色大框是 GB10 SoC，內部整合 20 核心 Arm CPU 與 Blackwell GPU；兩者透過統一記憶體控制器共用 128 GB LPDDR5x。

先把這台實驗機器的規格攤開

不需要對 1 PFLOP 有太高的期待，理想很豐滿、現實很骨感，它描述的是 FP4 而且包含稀疏性條件的峰值，後續若量 BF16 GEMM、FP16 attention 或 LLM decode，資料格式、Tensor Core 路徑、記憶體頻寬與 batch shape 都不同。

128 GB 是 CPU 與 GPU 共用的系統記憶體

這張架構圖最值得注意的是 CPU 和 GPU 沒有各自獨立的 DRAM 。DGX Spark 採用動態統一記憶體架構（Unified Memory Architecture, UMA），CPU、整合式 GPU 和作業系統共用同一組 128 GB LPDDR5x，沒有預先切一塊固定容量給 GPU。

在 x86_64 主機加獨立 GPU 的系統上，常見資料路徑是

CPU DRAM 中的 tensor
  ↓ cudaMemcpy，通常經 PCIe
GPU HBM / GDDR 中的另一份 tensor
  ↓
GPU kernel

DGX Spark 的實體記憶體路徑則較接近

128 GB LPDDR5x 實體記憶體
  ├─ Arm CPU 存取
  ├─ Blackwell GPU 存取
  ├─ 作業系統與 page cache 使用
  └─ 顯示與其他裝置使用

CPU 與 GPU 共用實體 DRAM，可減少 CPU RAM 與 GPU VRAM 各留一份資料的需求。不過，UMA 不等於每一種 CUDA 配置都自動是零複製，也不等於 CPU 和 GPU 同時存取同一筆資料沒有同步成本。程式仍要看自己使用 cudaMalloc、cudaMallocManaged、pageable host memory、pinned memory 或其他配置方式，並遵守 CUDA 的可見性與同步規則。

128 GB 也不是應用程式可以獨佔的 128 GB GPU VRAM。Linux 核心、桌面環境、其他行程、檔案 page cache 與顯示保留區都會占用它。NVIDIA 目前的 DGX Spark 軟體可在 UEFI 中把顯示保留記憶體設為 2 GB 或 4 GB；這項設定和背景行程都可能改變同一個模型在不同時間能否成功載入。

在 UMA 系統上，nvidia-smi 不會像獨立 GPU 一樣提供可直接比較的 Memory-Usage，cudaMemGetInfo 也不會把作業系統可回收的 page cache 與 swap 全部算進可分配容量。

它屬於 Blackwell，但不是桌上的 B200

DGX Spark、B200 與 GB200 都使用 Blackwell 世代技術，產品的記憶體與系統目標差很多。

這個比較會直接影響實驗結論。GB10 可以驗證 Blackwell / SM121 上的 CUDA 程式是否正確、Triton 或 CUTLASS 是否產生原生 kernel，也適合研究低功耗桌上型系統的推論行為。GB10 的 273 GB/s LPDDR5x 無法代表 B200 的 HBM3e，還有在 DGX Spark 上觀察到 memory-bound，也不能直接推論同一個 kernel 在 B200 上有相同比例的瓶頸。

今天先走到這裡

今天分享很多 GPU 相關知識

• CPU 比較適合低延遲、控制流程複雜、資料形狀不規則的工作；GPU 比較適合大量相同形狀的平行運算。
• CUDA kernel 的執行單位是 grid、block、warp、thread。
• SM 裡有 register、shared memory、CUDA cores、Tensor Cores。
• 一顆 GPU 可以從 PCIe / NVLink、HBM / GDDR、L2、SM、warp scheduler、register file、shared memory、Tensor Core 這幾個構造拆開來看。
• Global memory、L2、shared memory、register 的資料流會決定 kernel 效能。
• Compute capability 決定可用功能，compute_xy 是 PTX 虛擬架構，sm_xy 是 cubin 的實體 ISA 目標，兩者不要和 GPU 裡的 SM 數量混在一起。
• Hopper 用 TMA、transaction barrier、Thread Block Cluster 和 WGMMA 組成更大粒度的資料搬移與矩陣運算 pipeline。
• Blackwell 的 tcgen05、Tensor Memory、NVFP4 與第五代 NVLink 改變了 accumulator、MMA 與系統連線的設計。
• Rubin GPU 進一步強化 TMA、HBM4、長上下文 attention 和跨 kernel 協調，Vera CPU 則用 Olympus 核心、Spatial Multithreading、高頻寬 LPDDR5X、SCF 與 NVLink-C2C 承接 agentic AI 的工具執行、sandbox、資料處理及 orchestration。
• DGX Spark 是 GB10 Grace Blackwell 桌上型系統，它的 Arm64、SM121 與 128 GB 統一記憶體需要一起納入軟體與效能判讀。

明天會分享一下 CUDA、CUTLASS、cuBLAS。今天看硬體，明天看 kernel 和呼叫 library 的部分。

參考資料

• NVIDIA CUDA Programming Guide
  https://docs.nvidia.com/cuda/cuda-programming-guide/index.html

• NVIDIA CUDA GPU Compute Capability
  https://developer.nvidia.com/cuda/gpus

• NVIDIA CUDA Compiler Driver：GPU Compilation
  https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html#gpu-compilation

• NVIDIA CUDA Best Practices Guide
  https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html

• NVIDIA Hopper Architecture In-Depth
  https://developer.nvidia.com/blog/nvidia-hopper-architecture-in-depth/

• NVIDIA H100 Product Specifications
  https://www.nvidia.com/en-us/data-center/h100/

• NVIDIA H200 Product Specifications
  https://www.nvidia.com/en-us/data-center/h200/

• NVIDIA Hopper Tuning Guide
  https://docs.nvidia.com/cuda/hopper-tuning-guide/index.html

• NVIDIA PTX ISA：Asynchronous Warpgroup Level MMA
  https://docs.nvidia.com/cuda/parallel-thread-execution/#asynchronous-warpgroup-level-multiply-and-accumulate-operation-using-wgmma-mma-async-instruction

• NVIDIA CUTLASS：Warpgroup MMA Programming Guide
  https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/mma_docs/wgmma_programming.html

• NVIDIA Blackwell Architecture
  https://www.nvidia.com/en-us/data-center/technologies/blackwell-architecture/

• NVIDIA Blackwell Ultra Architecture In-Depth
  https://developer.nvidia.com/blog/inside-nvidia-blackwell-ultra-the-chip-powering-the-ai-factory-era/

• NVIDIA Blackwell Tuning Guide
  https://docs.nvidia.com/cuda/blackwell-tuning-guide/

• NVIDIA CUTLASS：tcgen05 MMA Programming Guide
  https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/mma_docs/tcgen05_programming.html

• NVIDIA Rubin GPU Architecture
  https://developer.nvidia.com/blog/inside-nvidia-rubin-gpu-architecture-powering-the-era-of-agentic-ai/

• NVIDIA Vera Rubin Platform
  https://nvidianews.nvidia.com/news/nvidia-vera-rubin-platform

• NVIDIA Vera CPU
  https://www.nvidia.com/en-gb/data-center/vera-cpu/

• NVIDIA Vera CPU：Agentic Workloads in AI Factories
  https://developer.nvidia.com/blog/nvidia-vera-cpu-sets-a-new-standard-for-agentic-workloads-in-ai-factories/

• NVIDIA CUDA 13.4 Developer Preview Release Notes
  https://docs.nvidia.com/cuda/developer-preview/13.4/cuda-toolkit-release-notes/index.html

• NVIDIA CUDA 13.4 NVCC：GPU Feature List
  https://docs.nvidia.com/cuda/developer-preview/13.4/cuda-compiler-driver-nvcc/index.html#gpu-feature-list

• NVIDIA DGX Spark Hardware Overview
  https://docs.nvidia.com/dgx/dgx-spark/hardware.html

• NVIDIA DGX Spark Porting Guide：System Overview
  https://docs.nvidia.com/dgx/dgx-spark-porting-guide/overview.html