Ampere 的 cp.async 怎麼藏 HBM 延遲

在 NVIDIA A100 上，HBM2e 一次載入大約要 450 到 600 cycles。這個數字很殘酷。你如果什麼都不做，一個 warp 可能就卡在那邊發呆。

Ampere 的 cp.async 很有意思。它把資料直接搬進 shared memory。它不先佔住 register，也不會把 long scoreboard 拉滿。講白了，就是讓你先做別的事，再回頭收資料。

這篇文章要談的，不只是指令本身。重點是思維切換。你不再想「先 load，再 compute」。你要想的是「資料先飛，算力先跑」。

A100 的記憶體階層，才是主角

A100 的效能，不是只看 CUDA core 數量。真正決定速度的，常常是記憶體階層。Registers 很快，但每個 thread 只有 255 個上限。Shared memory 和 L1 共享 192 KB。L2 cache 有 40 MB。HBM2e 理論頻寬可到 2 TB/s，但實戰通常沒那麼漂亮。

這些數字不是規格表裝飾品。它們直接決定 kernel 會不會翻車。Register spill 會掉到 local memory。那就是 global memory 等級的痛。Shared memory 如果打到 bank conflict，warp 也會被迫排隊。L2 miss 太多，延遲就會飆上去。

所以你看到 source code 很順，不代表跑起來就順。GPU 最愛在這種地方打臉人。尤其是資料路徑一長，問題就會被放大。

• Register file：每個 SM 約 256 KB
• Shared memory bank：32 個 bank，每個 4 bytes 寬
• L2 cache：A100 上是 40 MB
• HBM2e 理論頻寬：2 TB/s
• HBM2e 延遲：大約 450 到 600 cycles

所以 cp.async 的目的很明確。它不是消滅延遲。它是把延遲藏起來。這兩件事差很多。

cp.async 到底改了什麼

傳統 global load 會先進 register。這代表 warp 要等資料回來，才能繼續用那些目的暫存器。硬體會把這段等待算進 long scoreboard。你就只能乾等。

cp.async 不一樣。它把資料直接從 global memory 搬到 shared memory。中間不經過目的 register。這樣一來，warp 發出指令後，可以立刻去做其他運算。

這個差異看起來很小，實際上很兇。因為它把 load 和 compute 拆開了。你可以在算上一批 tile 的時候，讓下一批資料自己飛進來。這就是 overlap。

“Latency hiding is the name of the game.” — Mark Harris

這句話很老派，但一直有效。Mark Harris 一直在講同一件事。GPU 程式設計的核心，不是讓記憶體變快。是讓算力不要閒著。

我覺得 cp.async 厲害的地方，就在這裡。它不是魔法。它是把原本硬碰硬的等待，改成排程問題。

commit、wait、double buffer 才是實戰

cp.async.commit_group 和 cp.async.wait_group 這組搭配，才是實戰重點。前者只是做分組記帳。後者則是等到還剩幾組在飛。你如果設成 wait_group 1，就代表允許一組還在路上。

這樣就能做 double buffer。A buffer 在算，B buffer 在載。下一輪再交換。Kernel 不需要把 memory 變快。它只要讓 machine 一直忙。

這種做法很像工廠產線。不是每個工人都等同一個零件。是把流程拆開，讓每個站都不空轉。GPU 很吃這套。

• 傳統路徑：load 到 register，再等資料回來
• cp.async 路徑：直接進 shared memory
• commit_group：把一批 async copy 分組
• wait_group 1：保留一組在飛，其他先算

但這裡有代價。Shared memory 佔用會增加。stage 數一多，occupancy 可能掉。這不是免費午餐。你如果 kernel 算術密度不夠，可能反而賠。

所以像 CUTLASS 這種 library，才會把 pipeline depth 當成可調參數。不是每個 kernel 都適合同一種 stage 數。這種事真的只能量。

Profiler 看到的差別很直接

你如果想知道 kernel 有沒有吃到 cp.async 的好處，別先看感覺。直接看 profiler。傳統 load-heavy kernel，常常是 long scoreboard stall 佔大頭。你會看到 warp 很多時間都在等資料。

改成好的 pipelined 版本後，情況會變。long scoreboard 會明顯下降。FMA pipe 會更忙。這才是你要的畫面。不是「理論頻寬很高」，而是「實際有在算」。

這裡有個常見誤區。很多人只盯著 bandwidth。其實 kernel 快不快，不只看搬多少 GB/s。更重要的是，搬資料的時候，有沒有順便把 compute 填滿。

• 改前：smsp__warp_issue_stalled_long_scoreboard 常見 40% 到 70%
• 改後：long scoreboard 可能降到 5% 以下
• 調好後：smsp__pipe_fma_cycles_active 可到 70% 到 90%
• A100 L2 帶寬：約 4 TB/s aggregate

如果你想自己看，NVIDIA CUDA Samples 是很好的起點。先看原版 kernel，再做一版 tiled + async copy。差異通常很明顯。

工具面也別省。NVIDIA Nsight Compute 的 stall reason 和 issue activity，真的值得看。沒有這些數據，很多優化都只是猜。

跟 Hopper 比，Ampere 還差哪裡

如果你把視野拉大，Ampere 只是中繼站。Hopper 又往前推了一步。它有 Tensor Memory Accelerator，也就是 TMA。這東西把資料搬運再往硬體化推進。

這代表什麼？代表資料移動越來越不像 blocking load。它更像一個可排程的搬運任務。程式設計師還是要想資料布局，但不用把每次搬運都當成同步事件。

我自己的看法很直接。你如果還在寫那種「load、wait、compute、repeat」的 kernel，通常還有空間可以挖。尤其在 A100 這種卡上，cp.async 很值得試。

• Ampere：靠 cp.async 做 overlap
• Hopper：再往前，加入 TMA
• 競品面：AMD ROCm 也在推資料搬運優化，但 API 路線不同
• 實務面：GEMM、convolution、stencil 類 kernel 最常吃到好處

但別亂上。不是每個 kernel 都適合 async copy。資料量太小、算術密度太低，或 occupancy 已經很差的時候，硬上只會更亂。先量，再改。

這件事其實是 CUDA 老問題的新解法

CUDA 很多年來都在講 overlap。只是早期工具沒那麼順。你要自己拆 load、自己控同步、自己顧 pipeline。現在 cp.async 只是把這套做得更自然。

這也解釋了為什麼很多高效能 library 都很愛它。像 GEMM、attention、卷積這些工作，資料搬運本來就很重。只要能把搬運藏到計算後面，整體效率就會好看很多。

台灣做 AI 軟體的人，很多都只盯模型。其實底層 kernel 才是血肉。模型跑得快，不只是 Transformer 參數多。還要看資料怎麼走。這點很現實，也很煩，但就是事實。

下一步怎麼做

如果你手上有 A100 或其他 Ampere GPU，我會建議你先挑一個熱點 kernel。看它是不是被 long scoreboard 卡住。再試一版 double-buffer 的 cp.async 寫法。不要一次改太多。

如果 stall 比例下降，FMA 利用率上升，那就代表方向對了。若沒有，問題可能在資料布局、shared memory bank conflict，或 occupancy 本身。這時候別硬拗，回頭看 profiler。

說到底，cp.async 的價值很務實。它不是讓 HBM 變不慢。它是讓你少等一點。對做 CUDA 的人來說，少等 100 個 cycles，常常就夠有感了。

你如果現在就在調 kernel，我的建議很簡單：先量 long scoreboard，再試 async pipeline。別先信直覺。GPU 很少照直覺走。