昨天我一直說 util 假高、那個數字會騙你。今天我們就把它抓來對質。
小黃錶上那個數字跳到 100%,你以為引擎全開了。但它其實只告訴你一件事,錶在跳。它從來沒說車上有幾個人、引擎出了幾成力。你盯著它安心了一整天,它卻可能整天都在對你說謊。
先講結論。nvidia-smi 的 utilization,量的不是"你用了多少算力",而是"過去這一小段取樣時間裡,GPU 上有沒有至少一顆 kernel 在跑"。有跑,就算 100%。哪怕那顆 kernel 只用了整張卡的百分之一。
換句話說,util 100% 的意思比較接近"這張卡有打卡上班",而不是"這張卡有在做事"。這不就是 Day 1 那個薪水小偷嗎,卡有打、人在工位,正事一件沒做。
而且這個薪水小偷特別難抓,因為它手上握著一張看起來很漂亮的成績單。Day 1 我們算過,一張 H100 一小時要六到十二塊美金,你以為 util 100% 是把這筆錢花在刀口上,其實你可能只買到那張卡的一小角。錶讀數愈漂亮,你愈懶得去查,這才是它最貴的地方。真正燒錢的浪費,往往都躲在一片綠油油的 100% 底下。
nvidia-smi 每隔一小段時間去看一眼 GPU,問它一個是非題,這個瞬間有沒有 kernel 在執行?有,就記一筆"忙"。它把一段時間裡"忙"的比例算出來,就是你看到的那個百分比。
看出問題了嗎?這個問題只有"有"跟"沒有",沒有"用了多少"。一顆吃滿整張卡的大矩陣乘法,跟一顆只用一個角落的迷你 kernel,在 util 眼裡完全一樣,都叫"有在跑"。
一張 A100 上有 108 個 SM,你可以把它想成 108 條生產線。假設我故意寫一顆爛 kernel,只開一個 block、一個 thread:
// 只開 1 個 block、1 個 thread:108 條生產線,只用到 1 條
slow_kernel<<<1, 1>>>();
這顆 kernel 跑多久,nvidia-smi 的 util 就給你顯示 100% 多久。107 條生產線全空著、閒到長香菇,錶還是滿格。反過來,一個把整張卡塞好塞滿的正常 launch,長這樣:
// 開滿整張卡:每個 SM 都有活幹
big_kernel<<<108 * 4, 256>>>();
問題來了,這兩顆 kernel,nvidia-smi 顯示的 util 可以一模一樣,都是 100%。util 分不出它們,因為它從頭到尾只在問那個是非題。這不是 bug,這就是 util 的定義。
要怪,其實不能全怪 nvidia-smi。utilization 這個指標,是從更早的圖形時代留下來的。那個年代你在意的是"GPU 有沒有在畫東西、有沒有卡頓","有沒有在動"就是個很合理的問題。
可是到了拿 GPU 跑訓練跟推論的今天,我們在意的問題徹底變了。我們要的不是"它有沒有在動",而是"它那幾百條生產線,到底開了幾成、有沒有在做我付錢要它做的那些浮點運算"。同一個老指標,被拿來回答一個它從來不是為了回答而生的問題。錶沒有壞,只是你拿計程車的跳表,去問一台貨車今天送了幾噸貨。
Modal 有一篇很好的文章,把 utilization 拆成三層。你之所以會被 100% 騙到,是因為你把上面那層,當成了最下面那層。
| 層 | 它回答的問題 | 誰量 |
|---|---|---|
| 佔用率(allocation) | 這張卡有沒有被誰租走、佔住 | 排程器 / 雲平台 |
kernel util(= nvidia-smi) |
有沒有 kernel 在跑 | nvidia-smi |
| 算力(FLOP/s,= MFU) | 算力真的用了幾成 | profile 工具 |
nvidia-smi 卡在中間那層。它跟最下面那層,也就是你真正在乎的 MFU,差了十萬八千里。錶滿格只代表"有人上車",不代表"車正在往目的地開"。你想省的錢在最底層,你天天盯的數字在中間層,中間隔著兩種死法(發呆跟搬不完)在偷你的里程。
補一下最底層那個 MFU(Model FLOPs Utilization)的定義,因為跑 LLM 的人遲早要跟它打交道。它就是一個分數:
MFU = 模型"真正需要"的 FLOPs ÷ (花掉的時間 × 這張卡的峰值 FLOP/s)
分母好算,查 spec 就有(例如 A100 的 BF16 tensor core 峰值約 312 TFLOP/s)。分子才是關鍵,它是"你的模型完成這份工作、理論上非做不可的浮點運算",跟你實際上多花了多少無關。對一個 Transformer,這個數字有個很好用的估法:訓練約 6 × 參數量 × token 數,推論的 forward 約 2 × 參數量 × token 數。舉例,一個 7B 模型訓練時每個 token 大約要 6 × 7e9 ≈ 4.2e10 次浮點運算,這是它的"有效里程",剩下的通訊、bubble、非 matmul 的雜項,全是分母吃掉、分子不算的浪費。這也是為什麼 Day 1 說真實 MFU 常落在 40% 上下,完整的座標軸版本我們留到 Day 17,今天你先知道它是一個"有效 FLOPs 除以峰值"的分數就好。
那要看"引擎到底出幾成力",該看哪裡?
先講一個還在 nvidia-smi 裡、但比較少人用的:
nvidia-smi dmon -s u # 每秒印一次 util,但 sm 欄一樣是"有沒有在跑"的是非題
dmon 會每秒印一行,讓你看到 util 隨時間跳動,比一個總平均好一點。但你要有心理準備,它骨子裡還是同一個是非題,只是印得比較勤而已。真正誠實的數字叫 SM Active,它問的是"108 條生產線裡,這一刻有幾條真的在動"。這個數字得往更底層的硬體計數器拿,DCGM 的 PROF_SM_ACTIVE 就是在報這件事。今天先知道有這個東西、也知道它跟 util 是兩回事就好,怎麼實際抓、怎麼讀,期待Day 4 打開 Nsight 之後我們慢慢摸。
順手把三個最愛被當同義詞的名詞分清楚,之後整季都用得到:
nvidia-smi):有沒有 kernel 在跑。是非題。它們是一路往下的三道關卡,而且每一關都可能漏水。util 高,不保證 SM Active 高;SM Active 高,也不保證 Occupancy 高。更麻煩的是,就算三個都很漂亮,你的 MFU 還是可能很難看,因為那些 SM 也許正忙著等記憶體,也就是昨天的搬不完。所以請把這條不等式刻在腦子裡,util ≠ SM Active ≠ Occupancy ≠ 有用功。
Occupancy 特別容易被誤會,順便講清楚。A100 每個 SM 最多同時容納 2048 個 thread,也就是 64 個 warp。如果你的 kernel 寫得太吃暫存器,每個 thread 都占用一堆 register,可能只塞得下 16 個 warp,occupancy 就剩 25%。這就像一間能住 64 人的旅館,因為每個客人都帶了一卡車行李,最後只住得下 16 個人。占用率一低,SM 手上能輪替的 warp 就不夠,沒辦法在某些 warp 等記憶體的時候,抓別的 warp 上來頂著,於是延遲藏不住、算力就空在那。但反過來也成立,occupancy 高不代表快(塞滿 warp 但每個都在等資料,一樣慢),低也不代表慢(幾顆又肥又有效率的 kernel 反而可能很讚)。它只是一個線索,不是結論。
最常見的鬼故事發生在多卡訓練。八張 H100 用 NCCL 狂通訊的時候,nvidia-smi 給你 util≈100%,八張卡看起來忙翻天。但實際上 SM Active 可能只有 10–20%(nvidia-smi util 跟真實 SM 活躍度的落差,Arthur Chiao 這篇 拆得很清楚)。
為什麼?因為那些讓 util 顯示"忙"的 kernel,是在搬資料、在等其他卡對齊,也就是昨天講的搬不完加發呆,根本不是在算。util 滿格,八張卡整整齊齊地一起假忙,帳單一起跳。"我 util 都 100% 了為什麼還這麼慢",是每個做多卡的人遲早會問的問題,而答案很簡單,util 從來沒回答過快不快這件事。它甚至不知道有這個問題。
講一個很多團隊踩過的坑。有人把 nvidia-smi 的 util 接到監控面板上,設個告警,低於某個值就叫,看到滿江綠的 100% 就安心睡覺。結果月底帳單來了,訓練慢得要死,大家一臉問號,明明 util 都滿的啊。
問題就出在,他們把一個是非題當成了 KPI。util 適合拿來回答"這張卡是不是根本沒人用、可以回收去跑別的",這種調度層的問題它答得很好。但你一旦想用它回答"我的訓練跑得好不好、划不划算",它就會準時把你帶進坑裡。判斷"划不划算"的數字,是 MFU,不是 util。
不用等到後面幾天學會 profile,今天就有一個 30 秒的小實驗,讓你親眼抓到 util 說謊。
開兩個終端機,一個照常跑你的訓練,另一個下這行:
nvidia-smi dmon -s u # 會每秒串流一行,盯著 sm util 那一欄看
如果那一欄長時間貼著 100,可是你另一邊的 step/sec、tokens/sec 卻慢得讓你想哭,恭喜,你剛剛當場人贓俱獲,util 正在對你說"我很忙",但你的訓練根本沒往前走多少。這時候先別急著換卡、也別急著調 batch size,那多半是亂槍打鳥。真正該做的,是往下一層去看它到底卡在哪,而那正是明天開始的事。這個小實驗的重點不是解決問題,是讓你對 util 這個數字,永遠失去信任。這種健康的不信任,是這一季最值錢的習慣。

看那個 util,死死釘在 100% 不動。可是底下 108 個 SM,真正在閃的只有零星幾格。只要有一格在亮,錶就給你滿格。忙 ≠ 有用,而這一次,連"忙"都是裝出來的。
nvidia-smi 的百分比,是這一季的頭號騙子。它不是完全沒用,它能告訴你"這張卡有沒有被叫醒",但它永遠回答不了"這張卡在做有用的事嗎"。以後再有人拿 util 100% 跟你邀功,你可以微笑著回一句,那 SM Active 多少?
那要看真相,到底該去哪裡看?明天 Day 4,我們第一次打開 Nsight Systems 的 timeline,把"有沒有在忙"這個是非題,換成"誰在等誰"這張時間軸,用眼睛看那張卡到底在幹嘛。
錶還在跳。但從明天起,我們不看錶了,我們看時間軸。
PROF_SM_ACTIVE 等 profiling 指標:NVIDIA DCGM 文件。