昨天那台空車跳表的小黃，我們把它攔了下來。錶還在跳、後座沒人。問題是，為什麼？
大部分人遇到 GPU 慢，直覺反應是"不夠力，換張更貴的卡"。可是換卡之前，你得先知道它到底卡在哪。就像車開不快，可能是引擎老了、可能是後車廂還在等裝貨、也可能是司機根本還沒發動。這三種你都換一顆更大的引擎，只有第一種有救，另外兩種錢照燒、車照樣不動。
GPU 也一樣。它慢的方式，粗分只有三種。今天我們不談工具、不跑指令，只做一件事，把慢這件事分成三格。這三格，是這一季後面所有診斷的起點。
一張跑得不夠快的 GPU，逃不出這三種狀態：
而關鍵在於，這三種的救法完全不同。發呆要去補那些空檔、搬不完要減少資料搬運、算不動則代表你其實已經在做真功，該慶祝（或換演算法）。搞錯是哪一種，你會拿錯工具、修錯地方，白忙一場。這在鐵人賽裡叫做 "認真地做錯事"。
所以之後每抓一份 profile，第一個要問的永遠是同一句話，這是哪一種死法？
還是用昨天那台小黃。它要賺的是 "載客里程"（有用功、MFU），但它可能用三種方式賺不到錢。
司機把車停在路邊，還在滑手機決定下一趟去哪。引擎沒發動、輪子沒轉、跳表……欸，跳表有時候還在跳。
這是最容易被忽略的一種，因為 nvidia-smi 有時候還會給你不錯的數字。但實際上 GPU 都在等，等 CPU 發下一個工作、等 dataloader 把資料準備好、等另外七張卡在多卡訓練裡對齊。這段時間，運算單元是空的，電費一毛沒少。
發呆長什麼樣子？在 timeline 上就是一段一段的空白。最常見的有兩種：一種是你的程式每個 iteration 都有個固定的洞（CPU 來不及餵）；另一種是成千上萬個小到不行的 kernel，每一顆都要 CPU 花幾微秒去發射，結果 CPU 發射的速度跟不上 GPU 執行的速度，GPU 就一直在等下一顆。
舉個每天都在發生的例子，訓練迴圈裡每一步都 print(loss) 或 loss.item()，看起來人畜無害，其實它會逼 CPU 停下來、等 GPU 把這一步算完、再把數字讀回來，於是每個 iteration 都被硬生生插進一個小洞。
for x, y in loader:
loss = model(x, y)
loss.backward()
optimizer.step()
print(loss.item()) # ← 這行會等 GPU 算完才拿得到數字，順手插一個洞
一行順手的除錯碼，很可能就是你 timeline 上那排整齊空白的兇手。畫成圖大概像這樣，綠色是 GPU 真的在算，中間那些灰斜線的空白就是它在發呆、在空等，而錢照樣在燒。
這種洞用眼睛看訓練速度是抓不到的，得靠 profile 才現形，這也是為什麼我們這一季要學會量測。
這種 GPU 就是昨天說的薪水小偷的極端版，連裝忙都懶得裝，直接發呆。
發呆怎麼救、怎麼量，我們留到 Day 14（idle gap） 和 Day 15（kernel launch） 專門拆。
這台車引擎其實很猛，但它每跑一小段就得停下來，等後車廂把貨裝好、卸好。引擎在那邊空轉乾等，不是它不想跑，是料還沒到。
這就是 memory-bound（記憶體頻寬受限）。GPU 的運算單元快得誇張，但它算之前得先把資料從記憶體（HBM）搬進來、算完再搬回去。當你的運算每搬一筆資料、只算幾下，那大部分時間就都花在搬運上，算力在旁邊看戲。
給你一個數量級的感覺。A100 的 tensor core 一秒可以做三百多兆（TFLOP）次浮點運算，但它的記憶體一秒只搬得動大約 2 TB 的資料。換算起來，你每從記憶體搬 1 個 byte 進來，最好餵它做上百個浮點運算，否則算力就在等記憶體。
// SAXPY：每個元素讀 x[i]、讀 y[i]、寫 y[i]，搬 12 bytes，只換 2 個 FLOP
__global__ void saxpy(int n, float a, float *x, float *y) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) y[i] = a * x[i] + y[i]; // 1 乘 + 1 加
}
這顆 kernel 的算術強度大約只有 0.17（2 個 FLOP 配 12 bytes），遠低於前面說的每 byte 上百 FLOP 甜蜜點。所以再猛的卡，跑這種 kernel 都是在等記憶體。很多常見的操作（加法、normalization、activation）都長這樣，每個元素只算一兩下就要讀寫一次，天生就是 memory-bound。
把一個 elementwise 運算拆成階段來看就更清楚了，讀進來、算一下、寫回去。memory-bound 的意思就是，讀跟寫這兩段（藍）幾乎吃光整個 cycle，真正在算的那段（綠）只是薄薄一條：
有個很反直覺的實測，BERT 這種模型裡，那些非矩陣乘法的雜項操作，只佔了大約 0.2% 的浮點運算量，卻可能吃掉可觀的時間。做的事很少，花的時間很多，這就是搬不完的簽名。
搬不完怎麼抓，我們在 Day 16（H2D／記憶體傳輸） 動手。
第三種最特別，因為它其實不是壞事。
車在爬一段很陡的坡，引擎全開、油門踩到底、馬力用好用滿，它就是這麼快了。你再怎麼優化搬貨流程也沒用，因為瓶頸不在搬貨，在引擎本身的極限。
這就是 compute-bound，你已經把 GPU 的算力吃滿，慢是因為撞到了硬體的天花板。一個乾淨、寫得好的大矩陣乘法就長這樣。
// 矩陣乘法內圈：載入一次資料，換到大量乘加，資料重複利用
for (int k = 0; k < K; ++k)
acc += A[i][k] * B[k][j]; // K 次乘加，每筆搬進來的資料都被用很多次
跟上面的 SAXPY 剛好相反，這種算得多、搬得少的樣子，算術強度高，才吃得滿算力。這時候 util 高、SM 也真的在忙、MFU 拉得起來，這才是真的在載客里程。
所以算不動在三種死法裡是最健康的一種。撞到這裡，代表你前面兩關（發呆、搬不完）都過了。接下來能做的，通常不是調參數，而是換更聰明的演算法、換更低的精度、或接受這就是物理極限。純矩陣乘法的實務天花板大概也就在峰值的 70–80%，不會到 100%。
好，框架有了，那我怎麼知道我這台卡是哪一種死法？
今天先給你直覺版的判斷順序，精準的量法留給後面幾天各自負責：
這個順序不是隨便排的，先抓發呆、再抓搬不完、最後才是算不動。因為發呆通常最好修、也最浪費；沒有人希望自己花了大錢，結果卡在第一關還不自知。
其實搬不完和算不動，是同一件事的兩端。
想像一條路，一邊是搬料的速度（記憶體頻寬），一邊是算的速度（算力）。你的每個運算落在這條路的某個位置。如果它搬得多、算得少，瓶頸在搬料那端（memory-bound）；如果它算得多、搬得少，瓶頸在算力那端（compute-bound）。中間有一個交界點，過了它，瓶頸就從記憶體換成算力。
這張圖有個正式名字叫 roofline（屋頂線），它是 GPU 效能分析最重要的一張圖之一。但它的座標軸、怎麼把你的 kernel 畫成上面一個點、那個點又怎麼隨著優化移動，這些我們留到 Day 17 正式上場，跟 MFU 合在一起講。今天你只要記得那個直覺，同一台卡，可能因為搬料、也可能因為計算而慢，而且是兩種完全不同的病。
發呆和搬不完都是紅的，錶在跳，車沒在賺里程；只有算不動是綠的，因為那才是真的把算力用在刀口上。你的目標不是消滅所有的慢，而是把前兩種紅的，逼成第三種綠的。
今天最容易被誤讀的一句話是："所以我要讓 GPU 永遠不慢？"
不是。慢本身不是罪，浪費才是。 compute-bound 的慢，是你已經把錢花在刀口上、撞到物理極限的慢，那是好的慢。真正該抓的，是那種明明可以更快、卻卡在等的慢，也就是發呆和搬不完。
換句話說，我們要做的，不是把 GPU 榨到 100%（那不存在），而是把假忙碌（發呆、空等）換成真有用功（算不動）。當你的卡終於算不動了，恭喜，你才真的在載客。
今天只要帶走一件事，GPU 慢，先分三種死法，發呆、搬不完、算不動。 分錯格，你會拿錯工具修錯地方。
不過你可能已經發現一個問題，我一直在說 util 看起來很忙、其實假高，那 nvidia-smi 那個百分比到底在量什麼、為什麼會騙人？明天 Day 3，我們就把這個每個人都在看、卻沒幾個人真的看懂的數字，拆開來看。nvidia-smi 100%，其實可能什麼都沒證明。
錶還在跳。我們明天，來拆穿那支錶。
iThome鐵人賽
看影片追技術