H100 顯示很忙，LLM 卻可能沒吃滿：研究拆解 Hopper GPU「假性高利用率」

當監控工具告訴你 GPU utilization 已經來到 70%、90%，甚至接近滿載時，直覺通常會認為：這張 GPU 已經被榨乾了。

但對大型語言模型推理而言，事情可能完全不是這樣。

KTH Royal Institute of Technology 研究團隊在 2026 年 9 月公布的論文 Dissecting GPU Utilization for LLM Inference on Nvidia Hopper，直接挑戰了這個非常常見的效能判讀方式。研究團隊指出，NVIDIA Nsight Compute 中經常被視為 SM utilization 的 sm__throughput.avg.pct_of_peak_sustained_elapsed，本身並沒有錯，但如果把它當成「GPU 有多少比例正在做有效 AI 運算」，就很容易得到錯誤結論。

研究真正想回答的是：GPU 顯示很忙時，到底有多少運算是在處理真正的 token？

答案在某些 LLM decode 情境下，相當驚人。

一個 utilization 數字，藏了太多不同問題

現代 GPU 的 utilization 並不像汽車引擎轉速那麼單純。

同樣一個看似很高的百分比，背後可能包含完全不同的原因，包括：

• Tensor Core 正在有效執行矩陣乘法
• Warp 正在等待 HBM 或快取資料
• Shared Memory 或 Register 限制了 occupancy
• 執行網格沒有剛好填滿所有 SM
• GPU 正在執行包含大量無效 row 的矩陣 fragment
• Kernel 被 barrier 或 synchronization 卡住

因此研究團隊沒有使用單一 utilization 指標，而是建立 8 種互補的 GPU utilization view，利用 NVIDIA Nsight Compute 的底層 counter 去區分 occupancy、memory stall、wave efficiency、fragment fill 與 kernel selection 等不同因素。

測試平台是一張 NVIDIA H100 NVL 94GB HBM3，包含 132 個 Streaming Multiprocessor，也就是 SM，研究使用 BF16 推理，並搭配 vLLM、FlashAttention-3 與 cuBLASLt。

測試模型則包括：

• Meta Llama 3 8B
• Qwen3-14B
• Qwen3-32B
• Qwen3-30B-A3B MoE

研究進一步把 LLM inference 拆成三種典型工作狀態：cold prefill、warm prefill 與 decode。

而真正有趣的問題，就出現在 decode。

Prefill 與 Decode，其實是兩種非常不一樣的 GPU 工作

使用者送出一大段 prompt 時，模型首先進入 prefill。

假設輸入有數千個 token，GPU 可以一次處理大量 token，因此矩陣乘法中的 M dimension 很大，很容易形成尺寸漂亮的大型 GEMM。這種 workload 非常符合 Tensor Core 喜歡的形狀，大量計算可以一起塞進 GPU。

但模型開始回答後，事情改變了。

Autoregressive LLM 每次 decode iteration，對每一個 request 通常只生成 一個新 token。

如果 batch size 很小，例如只有一名或少數使用者正在生成內容，那麼矩陣乘法突然從大型矩陣變成非常矮的矩陣。

問題在於：Hopper Tensor Core 並不會跟著你的 token 數量無限縮小工作單位。

Hopper 的 GMMA 有一個固定的 64-row 地板

Hopper 架構使用 BF16 的 wgmma.mma_async 指令執行 Warp-Group Matrix Multiply-Accumulate，也就是 GMMA。

研究指出，它在 M 軸上的基本 fragment 是固定 64 rows。

可以把它想成一台一次至少要放 64 個位置的運算機器。

如果你一次真的有 64 個 token rows：

64 / 64 = 100%

非常漂亮。

但如果 decode batch size 只有 1：

1 / 64 ≈ 1.56%

也就是說，即使真正需要的資料只有一列，硬體仍然按照 64-row fragment 執行，其餘位置必須被 padding。

研究團隊把這個指標稱作 GMMA M-fragment fill。

在小 batch decode 中，他們得到：

• Batch 1：1.56%
• Batch 2：3.12%
• Batch 4：6.25%
• Batch 8：12.5%

在整個 decode 測試樣本中，平均 fragment fill 只有 5.86%。

相較之下，cold prefill 因為同時處理大量 token，平均 fragment fill 為 100%；warm prefill 則約為 70.2%。

這就是為什麼「GPU 正在執行 Tensor Core 指令」跟「Tensor Core 裡面裝的是不是有效 token」必須分開討論。

最極端案例：帳面 6.9%，有效矩陣運算只剩約 0.11%

研究甚至提供了一個非常具體的案例。

在 Llama 3 8B、sequence length 2048、batch size 1 的 decode QKV projection 中，研究觀察到 active SM 上的 compute throughput 約為 9.5%。

該 kernel 觸及 132 個 SM 中的 96 個，因此換算整張 GPU，device-wide compute utilization 約為：

9.5% × 96 / 132 ≈ 6.9%

如果只看到這裡，工程師可能會認為 GPU 還有大量 compute headroom，甚至開始嘗試重新調整 cuBLASLt tile。

但 batch size 1 的 GMMA fragment fill 只有 1.56%。

再把有效 row 比例算進去：

9.5% × 1.56% × 96 / 132 ≈ 0.11%

研究估算，對照 H100 約 835 TFLOP/s 的 BF16 dense Tensor Core peak，真正對有效 token row 做出的矩陣運算約只有 0.9 TFLOP/s。

這並不代表剩下 99.89% 的 H100 算力只要重新調校就能全部拿回來。

這點非常重要。

研究稱這是 measurement-attribution gap，也就是測量指標與真正有效工作之間的歸因落差，而不是單純的可回收效能空間。

原因是 Hopper BF16 GMMA 的 m64 fragment 本身就是指令集與硬體架構限制。

所以換一個 cuBLAS tile 就好了嗎？答案是不行

研究團隊進一步掃描 cuBLASLt 的 tile space，希望確認問題是不是 NVIDIA library 沒挑到最好的 tile。

結果不是。

研究針對 Llama 3 8B decode QKV，掃描 17 個 BF16 NT algorithm ID × 12 種 tile shape，共 204 種候選組合。

部分 algorithm 對該 workload 沒有可用 tile，最後留下 59 組有效配置。

結果沒有任何一組能完全解決問題。

原因很簡單：BLOCKM 無法低於硬體 GMMA 的 m64 基本限制。

所以這不是把 64×128 tile 換成另一個神奇 tile 就會消失的軟體問題。

研究認為，真正值得改變的是有效的 M dimension。

換句話說，不是改變「64 格的箱子」，而是想辦法一次多放幾個 token 進箱子。

Batch 越大，GPU 突然開始正常了

這也解釋了另一個非常重要的實驗結果。

當 Qwen3-14B 從 batch size 1 提升到 batch size 32 時，dense decode GEMM 的 SM-busy 從約 12% 上升到 71.9%。

Qwen3-32B 在 batch size 32、sequence length 2048 的測試中也達到 72.6% SM-busy，每 token wall-clock time 降到約 0.98 ms。

這不是 GPU 突然變快了。

而是大量 request 被一起 batching 後，矩陣的 M dimension 變大，硬體原本浪費掉的 fragment 空間開始被真正 token 填滿。

這也是大型 AI inference provider 為什麼如此重視 batching 與 request scheduling。

對雲端服務來說，一張昂貴 H100 最怕的可能不是沒有 request，而是 request 太零碎。

還有第二種浪費：SM 有了，工作卻沒平均分配

Fragment fill 並不是唯一問題。

研究還分析了 wave quantization。

H100 NVL 擁有 132 個 SM，因此理想狀況下，kernel launch 最好可以把工作均勻分給這 132 個 SM。

但實際 CTA 數量不一定剛好是 132 的倍數。

例如 Llama 3 8B 的某個 decode QKV kernel 只 launch 96 個 CTA。

結果就是：

132 個 SM − 96 個有工作的 SM = 36 個 SM 暫時沒有工作

該案例 wave efficiency 約為 72.7%。

整體測試中，研究觀察到 cold prefill 的平均 wave efficiency 高達 99.8%，warm prefill 為 85.5%，decode 則降至 80.5%。

也就是 decode 除了每個 GMMA fragment 裡可能裝不滿 token，整顆 GPU 的 SM 空間也可能沒有完全鋪滿。

這兩種浪費會疊加。

記憶體等待依然是 Decode 的另一隻大象

即使忽略 fragment fill，decode 本身還面臨 memory-bound 問題。

研究使用 Nsight Compute 的 long_scoreboard stall 分析發現，在 dense decode GEMM 中，約 70%～84% 的 stall budget 與長延遲 memory dependency 有關。

warm prefill 的情況也非常接近，約為 72%～83%。

反之，cold prefill 因為大型矩陣可以把資料載入成本攤到更多計算上，long_scoreboard 比例下降至約 19%～36%。

因此，小 batch decode 實際上同時遭遇幾種問題：

矩陣太矮 → GMMA fragment 裝不滿

Kernel grid 太小 → SM wave 裝不滿

Arithmetic intensity 太低 → 更容易受記憶體頻寬與 latency 限制

這就是為什麼單一 GPU utilization 百分比幾乎不可能完整描述 LLM inference。

FlashAttention 又是完全不同的瓶頸

另一個有趣發現是，FlashAttention-3 並沒有單純複製 dense GEMM 的問題。

FA3 在 H100 上使用 persistent kernel，通常會建立一個能覆蓋全部 132 個 SM 的 launch grid，然後在 kernel 內部持續處理 attention tile。

因此它在 wave utilization 上相對漂亮。

但 FA3 decode 的 stall signature 又變成以 barrier 與 wait 為主，而不是 dense GEMM 常見的 long_scoreboard memory stall。

也就是：

Dense GEMM 慢，不代表 FlashAttention 也因為同樣原因慢。

這正是研究團隊強調的核心觀點：不同 kernel 的 utilization 損失具有完全不同的物理原因，因此需要不同的解法。

那未來 LLM inference 應該怎麼優化？

對 Hopper 小 batch decode 而言，論文認為只重新調整 cuBLASLt tile 的幫助有限。

因為 GMMA m64 是硬體層面的最低 fragment。

比較合理的方向包括：

1. Cross-request token packing

把多個不同使用者正在生成的 token 收集起來，一起形成更大的矩陣 M dimension。

這其實就是更積極、更聰明的 batching。

2. Persistent decode kernel

不要讓每一個超小 GEMM 都按照傳統方式獨立 launch，而是設計更適合 autoregressive decode 工作型態的 persistent kernel，在 GPU 內持續接收與處理 token work。

3. Serving scheduler 與硬體共同最佳化

LLM serving 最佳化不再只是「選最快 GPU」。

Scheduler 決定哪些 request 同時執行，就等於間接決定 Tensor Core 的矩陣形狀。

這讓 scheduling 從後端軟體問題，變成實際影響硬體利用率的 architecture 問題。

對 AI 資料中心來說，這件事為什麼重要？

一張 H100 本身極其昂貴，而 hyperscale AI inference service 動輒部署數千甚至數萬張 GPU。

假如容量規劃只看一個 utilization 指標，就可能產生兩種完全相反的錯誤判斷。

第一種是看到 SM utilization 很高，認為 GPU 已經沒有餘裕，因此繼續加卡。

第二種則是看到 compute utilization 偏低，認為換 tile 或 kernel 就能拿回大量算力。

實際上，有些浪費來自 memory stalls，有些來自 occupancy，有些來自 wave quantization，還有一些像 GMMA m64 fragment floor，根本就是現有硬體執行單位的基本粒度。

它們不能用同一種最佳化方法處理。

這也意味著未來 NVIDIA Blackwell、Rubin，以及其他 AI accelerator 的競爭，不能只比較 TFLOPS、HBM bandwidth 與 VRAM 容量。

硬體能否有效處理大量細碎、小矩陣、低 batch 的 autoregressive workload，同樣可能決定真實 inference 效率。

對 AI GPU 而言，真正重要的問題逐漸從「一秒可以算多少 FLOPs」，變成「有多少 FLOPs 真的是使用者需要的那一部分」。

---

Tech English｜本文技術名詞

GPU Utilization（GPU 利用率）：描述 GPU 某類運算資源忙碌程度的指標，但不同 counter 的分母與定義不同，因此不能直接等同有效運算比例。

GMMA — Warp-Group Matrix Multiply-Accumulate（Warp 群組矩陣乘加）：NVIDIA Hopper Tensor Core 使用的重要矩陣運算機制。BF16 GMMA 在 M 軸具有固定 64-row fragment，是本文小 batch decode 效率問題的核心。

Fragment Fill（矩陣片段填充率）：真正有效資料占硬體固定運算 fragment 的比例。例如只使用 m64 fragment 中 1 row 時，fill 約為 1.56%。

Wave Quantization（執行波量化效應）：Kernel 的 CTA 數量無法剛好填滿所有 SM 時，最後一個 execution wave 留下一部分閒置 SM 的現象。

Prefill / Decode（提示詞預填充／逐 Token 解碼）：Prefill 一次處理大量輸入 token，通常形成大型矩陣；decode 每輪只替每個 request 產生一個 token，因此容易形成小 M GEMM。

Persistent Kernel（常駐核心）：Kernel 啟動後長時間留在 GPU 上持續處理工作，而不是每個小任務重新 launch，可減少啟動成本並改善細碎 workload 的硬體利用率。