核心概念

PyTorch 的 torch.profiler 是效能優化的起點。Hugging Face 用「矩陣乘法 + 偏移加法」這個最簡單的運算,系統性示範如何設定 profiler、讀懂輸出,以及從中得出可行動的洞察。

設定流程只有四步

  1. torch.profiler.record_function("name") 標記要分析的區段
  2. torch.profiler.profile() 包住執行迴圈,指定 CPU 與 CUDA activities
  3. 設定 schedule(wait=1, warmup=1, active=3) 過濾冷啟動雜訊
  4. 輸出 profiler table(.txt,統計摘要)和 trace(.json,Perfetto 視覺化)

兩個輸出對應兩個不同問題:table 問「哪裡最慢」;trace 問「為什麼這時候發生這件事」。

Overhead-bound 與 Compute-bound

Profiler table 最後兩行是核心診斷指標:

場景 CPU time total CUDA time total 狀態
64×64 矩陣 2.314 ms 23.104 µs overhead-bound,GPU 98% 空轉
4096×4096 矩陣 4.908 ms 4.495 ms compute-bound,GPU 是瓶頸

CPU time 單位是 ms、CUDA time 單位是 µs → 立刻確診 overhead-bound:CPU 花在準備和發送 kernel 的時間遠多於 GPU 計算時間。解法是加大運算規模、批次化呼叫或融合 kernel。

Trace 中五個關鍵現象

1. 第一個 ProfilerStep 特別寬

ProfileStep#2 比後續步驟寬了約 227 µs。根源是 record_function 進入後到 aten::matmul 真正 dispatch 之間的「dead window」——cuBLAS 在做 workspace 分配與啟發式選核。加 warmup 無法消除這段成本,只能讓它不被計入 active 視窗。

2. CPU 與 GPU lane 之間的 ~2.5 ms 偏移

GPU lane 的第一個 kernel 比 CPU 晚約 2.5 ms 出現,原因是 Activity Buffer Request:profiler 在 GPU VRAM 分配自身的事件 buffer。看到這個就知道偏移來自 profiler 本身,不是程式碼問題。

3. 兩種 CUDA Runtime 呼叫格局

aten::mm  → cudaOccupancyMaxActiveBlocksPerMultiprocessor + cudaLaunchKernel
aten::add → cudaLaunchKernel(僅此一個)

矩陣乘法前多了 occupancy query:cuBLAS 需要知道硬體能同時放多少個 block,才能選出最適合的 tile 大小與 kernel 變體。加法的資源需求固定(32 registers, 0 shared memory),不需要查詢。

這是通用掃描訊號:有 occupancy query → 重型 GEMM 級 kernel;沒有 → elementwise/reduction

4. 同一個 kernel 每次執行時間不同

在 20 次迭代的 4096×4096 trace 中,同一個 ampere_bf16_s16816gemm... kernel 跑出不同耗時。原因包括 GPU clock boost/idle 週期、散熱、電源管理、driver 背景作業。只看平均會得到「matmul 約 1 ms」的錯誤心理模型;trace 才能看到「通常 580 µs,偶爾因 GPU 狀態跑到 1000+ µs」的真實分佈。

5. torch.compile 的融合是 dispatcher 層,不是 kernel 層

compile 後 aten::add + aten::mmaten::addmm,但 GPU 執行的 kernel 名稱完全一樣。實際發生的是:

  1. cudaMemcpyAsync(DtoD):把 bias 複製到輸出 buffer(32 MB, ~33 µs)
  2. cudaLaunchKernel:GEMM 把 bias add 折入 epilogue

不僅沒省掉 memcpy,CPU overhead 還因為每次都要走 Dynamo cache lookup → AOTAutograd → Inductor stack 而增加了約 2 倍torch.compile 對單一 op 是純稅;要跨多個 op 才能攤銷成本。

關鍵要點

  • 快速診斷:CPU time(ms)遠大於 GPU time(µs)→ overhead-bound,優先加大矩陣規模或批次化
  • 第一個 ProfilerStep 偏寬是冷啟動稅,加 warmup 讓它不進入 active 視窗,而非消除它
  • 有 occupancy query = GEMM 級 kernel;沒有 = elementwise/reduction,可用來快速掃描 trace hotspot
  • 不要只看平均:kernel 執行時間有硬體抖動,trace 的分佈才是真相
  • torch.compile 對多 op 模型有效,對 ≤2 個 op 的函數反而增加 CPU 開銷
  • Perfetto 鍵盤導航:W/A/S/D 放大縮小移動,比滑鼠快很多

實務應用

測試矩陣大小要接近生產規模:用小矩陣測試會顯示 overhead-bound,但生產的 LLM 推論幾乎都是 compute-bound,兩個結論天差地別。

warmup 策略:在 profiler 外部跑數次 warmup,讓 cuBLAS 完成 workspace 分配與 kernel 載入;schedulewait=1, warmup=1 再過濾一次 profiler 內的冷啟動。

trace 讀法速查表

訊號 意義
Activity Buffer Request in GPU lane profiler 自身分配 buffer,非程式碼問題
第一個 step 特別寬 冷啟動(workspace、cuBLAS heuristic)
cudaDeviceSynchronize 很長 profiler flush;其長度 ≈ GPU idle 時間,是 overhead-bound 症狀
兩個 kernel 之間有空隙 可能是 buffer request;多跑幾步確認只出現一次
Memcpy DtoD + GEMM 連在一起 addmm 的 bias copy + epilogue,「融合」在 dispatcher 層

延伸觀點

來源:PyTorch 官方 Profiler Recipe(pytorch.org)

記憶體 profiling:超出 compute 維度的診斷

torch.profilerprofile_memory=True 參數可以同時追蹤記憶體分配。輸出的 table 會新增 Self CPU MemCPU Mem 兩欄,分別代表「該 op 自身分配的記憶體」與「包含子 op 的總記憶體」。

這個區別實務上很重要:如果你看到某個高層函數的 CPU Mem 很大但 Self CPU Mem 很小,表示記憶體壓力來自它的子呼叫鏈,而不是它自己。這在診斷 OOM 或記憶體碎片時能快速縮小範圍——不需要逐行看 trace,直接在 table 排序 CPU Mem 欄位就能找到真正的記憶體大戶。

Call stack 整合:把 trace 對應到模型架構

Profiler 支援 with_stack=True 選項,讓 trace 中的每個 op 都帶上完整的 Python call stack。在 Perfetto 中點擊任一 op,就能看到它是從模型的哪一層、哪個方法發起的。

原文的矩陣乘法範例太簡單,看不出這個功能的威力。但在多層 Transformer 或複雜 CNN 中,同一個 aten::mm 可能被 attention、FFN、projection 等不同模組呼叫,耗時各異。沒有 call stack 就只能看到「matmul 佔了 X ms」,有了 call stack 才知道是哪個 attention layer 的哪個 projection 是瓶頸。

這兩個維度(記憶體 + call stack)都是 Hugging Face 這篇入門文未涵蓋的,屬於 torch.profiler 中級用法,適合在確認 compute hotspot 後的下一步分析。

非同步連續批次推論:LLM 推論的 CPU GPU 並行加速 | AWS 基礎模型訓練與推論四層架構 | vLLM V0 升級 V1:強化學習訓練的後端正確性優先原則

反向連結

以下頁面引用了本頁: