Day 3 — 先猜，再問硬體

(還在趕草稿 可惡><)

Day 2 我們一路從 Thread、Warp 看到 Block，最後把 naïve GEMM 改成了一個 tiled version：原本每個 Thread 反覆從 global memory 抓資料，現在改成讓整個 Block 合作，把一部分 A、B 搬進 Shared Memory，再讓 Threads 重複使用。

直覺上，這應該會快不少。

到底快多少？

Profiler 先不要開，我們自己算。

只從 source code 數 logical accesses，一個 phase 裡 v0-control 需要 4096 次 FP32 element access，v1-tiled 剩下 320 次：

v0-control: 4096
v1-tiled:    320

4096 / 320 = 12.8×

12.8×。

這筆帳很好看。

但我先不信它跟最後的 runtime 有直接關係。

因為現在數到的只是 source code 裡少了多少 logical access，還不是 GPU memory system 真正少做了多少工作。

---

[圖 1：v0-control，一個 Warp 如何讀 A / B]

---

這也是今天真正想用 Nsight Compute 回答的問題。

不是先打開 profiler，看哪個 bar 最大、哪個數字變紅，而是先從程式碼猜：

如果我現在對 Warp 和 GPU memory access 的理解沒有錯，這個 12.8× 一路走到真正的 hardware 之後，到底還會剩多少？

Profiler 等一下再開。

先算。

先留一個 Control

這裡我不直接拿 Day 2 最原始的 naïve kernel 跟 tiled kernel 比。

因為 tiling 一做下去，通常不只 memory access 改了。Loop shape 可能不同，compiler 的 unrolling opportunity 可能不同，address calculation 也可能跟著改變，再加上 Shared Memory instruction 和 synchronization。

如果最後 tiled version 快了，我當然很高興，但這種 comparison 很難告訴我到底是哪個 change 在付錢。

所以我多留一個 v0-control：

for (int k0 = 0; k0 < K; k0 += 8) {
    #pragma unroll
    for (int kk = 0; kk < 8; ++kk) {
        acc += A[row * K + k0 + kk]
             * B[(k0 + kk) * N + col];
    }
}

它仍然直接從 global memory 讀 A、B，只是 loop structure 盡量跟 v1-tiled 對齊。

tiled 版本則是：

Bs[threadIdx.y][threadIdx.x]
    = B[(k0 + threadIdx.y) * N + col];

if (threadIdx.x < 8) {
    As[threadIdx.y][threadIdx.x]
        = A[row * K + k0 + threadIdx.x];
}

__syncthreads();

#pragma unroll
for (int kk = 0; kk < 8; ++kk) {
    acc += As[threadIdx.y][kk]
         * Bs[kk][threadIdx.x];
}

__syncthreads();

所以今天真正比較的是：

v0-control
vs.
v1-tiled

一次少動一點。

不然最後數字再漂亮，也很難歸因。

---

12.8× 走到 Warp，還剩多少？

前面的 4096 和 320，是從 lane / source code 的角度數出來的。

但 Day 2 已經知道，GPU 真正開始執行 memory instruction 時，我們還得把視角拉到 Warp。

v0-control 一個 phase 有 8 個 kk。每個 kk 裡，一個 Warp 會執行一條 A load、一條 B load：

2 loads × 8 kk
= 16 Warp-level requests

一個 (32, 8) Block 有 8 個 Warps：

16 × 8
= 128 requests / Block / phase

tiled 版本則是一個 Warp 每個 phase 做一次 A cooperative load，以及一次 B cooperative load：

2 × 8 Warps
= 16 requests / Block / phase

所以：

128 → 16

= 8×

好，12.8× 走到 Warp，先變成 8× 了。

這個數字我就比較願意拿去問 hardware。

---

[圖 2：v1-tiled，Block 如何合作把資料搬進 Shared Memory]

---

為什麼是 8×，不是 12.8×？

原因就在昨天看過的 A access。

同一個 Warp 在某個 k 上，32 個 lanes 都需要相同的：

A[row, k]

Source code 看起來像 32 個 lanes 各做了一次 element access，但它們全部指向同一個 address。

所以：

logical element accesses
≠
Warp-level requests

這兩個數字都沒有錯，只是站的位置不一樣。

我們甚至還沒打開 profiler，12.8× 就已經先變成 8× 了。

這也是我今天很想留下的一個感覺：

你在 source code 上看到的 optimization，不會原封不動一路傳到 hardware。

---

再往下一層：Sector

Warp request 還不是 memory system 做的所有工作。

同一個 request 裡，32 個 lanes 的 addresses 最後還會碰到一個或多個 L1/TEX sectors。

按照目前這個 access pattern，我們可以先估：

v0-control
≈ 320 sectors / phase

v1-tiled
≈ 40 sectors / phase

也就是：

320 → 40

≈ 8×

所以打開 NCU 以前，我手上的 prediction 是：

Quantity	v0-control	v1-tiled	Prediction
Logical accesses	4096	320	12.8× ↓
Warp requests	128	16	8× ↓
L1 sectors	~320	~40	~8× ↓
L2 traffic	?	?	不知道
DRAM traffic	?	?	不知道
Runtime	?	?	不知道

到 L2 我就停了。

因為 cache hierarchy 已經進來。光靠 source code，我現在不知道多少 request 會在 L1 hit、多少真的會繼續往 L2，最後又有多少真的跑到 device memory。

這些問號不是忘了算。

是現在沒有資格算。

---

[圖 3：Source Code → Warp Request → L1/TEX → L2 → DRAM → Runtime]

---

好，現在才開 Nsight Compute

第一次開 NCU 很容易被嚇到。

又藍、又綠、又一堆 percentage，旁邊還有好多看起來像「你應該懂」的名詞。

先不要怕。

今天也沒有要把 Nsight Compute 一次學完。我只想拿它回答剛才那兩個問題：

Requests 真的 ↓ 8× 嗎？

Sectors 真的 ↓ 8× 嗎？

其他 metric 先當路人。

NCU 的 Speed Of Light 本來就是 compute / memory throughput 的 high-level overview，而 Memory Workload Analysis 才會再把 L1/TEX、L2、Shared Memory 等 memory resources 拆得更細。(NVIDIA Docs)

先從最上面的 overview 看一眼。

---

4
[圖 4：through_put_hit_rate_compare_table.png]
(註:這個圖是由speed of light (Duaration的部分) & memory analysis (其他部分) 合在一起地 你在NCU裏面是看不到的)

v0-control：

Duration            196.18 ms
Memory Throughput   181.10 GB/s
L1/TEX Hit Rate      87.20%
L2 Hit Rate           23.31%

v1-tiled：

Duration            182.30 ms
Memory Throughput   195.00 GB/s
L1/TEX Hit Rate      ~0.2%
L2 Hit Rate           22.71%

第一個最直接的結果：

196.18 ms → 182.30 ms

runtime 大概只少了 7%。

等等。

說好的 8× 呢？

我們前面不是算得很開心嗎？

如果我偷懶把：

Global-memory work ↓ 8×

直接翻譯成：

Kernel 應該快 8×

那現在已經被第一張圖打臉了。

更奇怪的是 L1/TEX Hit Rate。

v0: 87%
v1: 接近 0%

我明明「優化」了第二版，結果 cache hit rate 直接炸掉。

看到這裡第一反應大概是：

黃仁勳你騙我？
我特地學 Shared Memory，結果 L1 hit rate 給我掉到 0？

先不要退學。

往下看。

---

先看最簡單的地方：V0 到底在幹嘛？

這張是 v0-control。

---

5 A
[圖 5A：V0_Shared_memory_and_cache.png]

---

先不用把整張表看懂，只看我框起來的地方。

上面 Shared Memory：

Shared Load  = 0
Shared Store = 0

完全沒用到。

這跟 source code 一致，因為 v0-control 根本沒有 Shared Memory。

再往下面 L1/TEX Cache 看 Global Load：

Instructions = 4,294,967,296
Requests     = 4,294,967,296
Sectors      = 10,737,418,240
Sectors/Req  = 2.50

這裡先認識幾個 NCU 的詞就好。

Instructions 是實際執行的 SASS memory instructions，以 Warp 為計數尺度；對 Volta 以及更新架構的一般 global/local LSU traffic，一條這類 instruction 會產生一個 L1 request。Sectors 則是這些 requests 實際碰到多少個 L1 sectors，而一個 sector 是 32 bytes。(NVIDIA Docs)

所以這裡其實在告訴我們：

v0 正在瘋狂地向 L1/TEX 發 global-load requests。

4.29 billion 次。

Giga 等級。

---

再看 V1，畫面突然完全不一樣

---

[圖 5B：V1_Shared_memory_and_cache.png]
5 B

v1-tiled 的 Global Load：

Instructions = 536,870,912
Requests     = 536,870,912
Sectors      = 1,342,177,280
Sectors/Req  = 2.50

但 Shared Memory 多出：

Shared Load Instructions
= 2,684,354,560

Shared Store Instructions
= 536,870,912

原本：

Global Load
Global Load
Global Load
Global Load
...

現在變成：

Global Memory
      ↓
先合作搬一次
      ↓
Shared Memory
      ↓
大量重複使用

所以 tiling 並沒有讓 memory operation 蒸發。

它做的是換地方 reuse。

原本大量重複的需求走 global-load path，而且很多被 L1 cache 接住；現在 programmer 自己把這件事寫得更明確：先把 tile staging 到 Shared Memory，後面的 computation 直接在 Block 裡 reuse。

這時候前面那個奇怪的 L1 hit rate 就開始有答案了。

---

L1 Hit Rate 掉到 0，原來不一定是壞事

v0-control 的 L1/TEX hit rate 大約：

87%

tiled 則接近：

0%

如果只記得：

Cache Hit Rate 越高越棒

那這裡一定會看得很痛苦。

但現在把 request 數量一起看：

v0:
4.29G Global Requests

v1:
536.87M Global Requests

v1 根本已經把大量重複的 global requests 拿掉了。

原本：

Global Load
    ↓
 L1 Hit
    ↓
拿到之前用過的資料

現在很多 reuse 變成：

Global Load
    ↓
Shared Memory
    ↓
接下來直接在 Block 裡 reuse

所以 v1 剩下的 global loads，本來就比較接近「把新 tile 搬進來」的需求。

L1 不再是主要幫我們做 reuse 的地方，hit rate 當然可以大幅下降。

這裡真正的 lesson 反而是：

Profiler 裡的數字，不是考試分數。

Hit Rate 不是看到 87 就給 A，看到 0.2 就當掉。

要先問它在描述哪一批 workload。

---

好，那我們猜的 8× 到底有沒有？

有。

而且漂亮到有點誇張。

Global Load Requests：

v0 = 4,294,967,296
v1 =   536,870,912

4,294,967,296 / 536,870,912
= 8

剛好 8×。

我們完全猜對了。

再看 Sectors：

v0 = 10,737,418,240
v1 =  1,342,177,280

= 8×

又是剛好 8×。

所以：

Metric	v0-control	v1-tiled	Change
Global Load Requests	4,294,967,296	536,870,912	8× ↓
L1 Global Sectors	10,737,418,240	1,342,177,280	8× ↓
Sectors / Request	2.50	2.50	不變

這大概是今天最爽的一段。

因為我們真的先從 Warp mapping 猜：

128 → 16
= 8×

再去問硬體。

硬體回答：

4.29G → 536.87M
= 8×

Prediction 活下來了。

---

Sectors/Req = 2.50 其實也很好玩

這欄第一次看到可能不知道在幹嘛。

NCU 的 Sectors/Req 就是平均一個 L1 request 需要多少 sectors。對 32 個 active lanes 各讀一個 aligned FP32 的一般連續存取，理想值會是 4；如果很多 lanes 重複讀相同資料，ratio 也可能低於 4。(NVIDIA Docs)

我們這個案例剛好可以自己解釋 2.50。

一個 Warp 讀 A 時：

大家讀同一個 A[row,k]
≈ 1 sector

讀 B 時：

32 個 FP32 連續讀
≈ 4 sectors

平均：

(1 + 4) / 2
= 2.5

NCU：

2.50

又對上了。

而 tiled 版本 A cooperative load 大約也是 1 sector，B 仍然大約是 4 sectors，所以平均還是 2.50。

這也告訴我們一件事情：

這次 tiling 的主要 improvement 不是讓單次 global request 更 coalesced，而是直接減少 global request 的總數。

Reuse 才是這次的主角。

---

那為什麼不是快 8×？

這才是今天真正有意思的地方。

來看你整理好的 Memory Chart。

---

6 A
[圖 6A：V0_memory_chart(1).png]

---

先找右下角大大的 V0，然後看左上角紅框。

4.30 G Inst
4.29 G Req

很多。

再看左下角 Shared Memory：

0

很符合我們的 control kernel。

接著看中間 L1/TEX：

Hit Rate = 87.20%

v0 雖然瘋狂發 global request，但很多根本就在 L1 hit。

再看右邊 L2 和 Device Memory，device-memory transfer 大概還是在 34.75 GB 這個量級。

接著換 V1。

---

6 B
[圖 6B：V1_memory_chart(1).png]

---

左上 Global：

537.40 M Inst
536.87 M Req

真的只剩原本大約八分之一。

但是左下 Shared Memory 現在開始有東西了：

3.22 G Shared Instructions

2.68 G Shared Load Requests
536.87 M Shared Store Requests

這張圖其實把 tiled GEMM 畫得超清楚。

我們把原本很多 global work 搬到了 Shared Memory。

但再往右看。

Device Memory 還是：

約 34.72 GB

跟 v0 的：

約 34.75 GB

幾乎一樣。

這就是 8× 消失的地方。

我們的：

Global Request ↓ 8×

是真的。

L1 Sector ↓ 8×

也是真的。

可是這些被我們消掉的 requests，有很大一部分原本就會在 L1 hit。

它們本來就沒有每次都一路跑到 L2，更沒有每次都跑到 DRAM。

所以 source code 上看起來消掉一大票工作，到了更遠的 memory hierarchy，差距就迅速縮小。

這也是為什麼最後：

196.18 ms
→
182.30 ms

只有大約 7%。

這裡比單純說：

「因為這部分只佔總時間一小部分。」

更精確一點。

我們現在真的看見了：

我們消掉很多 global-load requests，但它們原本很多就已經被 L1 cache 吃掉，因此真正更昂貴的 L2 / device-memory work 並沒有跟著少 8×。

這就是 hardware counter 幫我們補上的那半個故事。

---

圖上還有 Sector、Peak、Bank Conflict，順便認一下

到這裡 NCU 上還有幾個字一直在旁邊晃。

今天不會全部展開，但至少知道它們在說什麼。

Sector

Sector 是 L1/TEX table 裡的概念。

一個 L1 sector 是 32 bytes；一個 request 可以碰一個或多個 sectors，所以我們才會看到 Sectors 和 Sectors/Req。(NVIDIA Docs)

注意它不是「資料經過 Shared Memory 時也叫 sector」。

Shared Memory 那張表看的比較是：

Instructions
Requests
Wavefronts
Bank Conflicts

兩邊不要混在一起。

% Peak

Shared Memory table 裡面的 % Peak，可以先理解成：

這組 workload 把處理這些 operations 的硬體能力用了多少。

數字高代表 unit 比較忙、可能接近瓶頸，但不代表「效率一定比較好」。NCU 官方對這個欄位的定義也是 utilization 相對於 peak；高 utilization 可以提示 bottleneck，但不是 efficiency score。(NVIDIA Docs)

我們的 v1 Shared Memory Total 大概來到：

47.64% Peak

至少可以確認：Shared Memory 現在真的成為 kernel 裡一個有份量的工作來源了。

Bank Conflict

最後這個名字很兇：

Bank Conflicts

Shared Memory 內部被切成多個 banks。如果同一個 Warp 的不同 Threads 同時存取同一個 bank 裡不同的位置，硬體可能需要把 request 拆開、序列化處理，這就是 bank conflict。NVIDIA 目前的 shared-memory model 是 32 banks，連續 32-bit words 會映射到連續 banks。(NVIDIA Docs)

我們這次的 report 確實可以看到 Bank Conflicts 不是 0。

但是這裡我要忍住一個很誘人的故事：

「抓到了！一定就是 Bank Conflict，所以才只快 7%！」

現在還不能這樣講。

因為 Shared Load 那一列是：

2,684,354,560 instructions

Bank Conflicts
141,516

相對整體 Shared Load workload 其實非常小。

Total 裡雖然還有 Other 帶來另外一批 conflicts，但光靠今天這張圖，還沒有證據說 Bank Conflict 就是吃掉 speedup 的主因。

這反而正好符合 Day 1 到現在一直在做的事情：

看到可疑 metric，不要急著抓戰犯。

Bank Conflict 今天先認識名字。

哪一天真的遇到一顆 kernel，Shared Memory requests 和 wavefronts 被 conflicts 明顯放大，我們再正式把它抓出來處理。

---

所以今天到底學會怎麼看 NCU 了嗎？

很粗略地，算有。

至少下次再看到這張又藍又綠的 Memory Chart，不需要直接關掉。

可以先從左往右看：

Kernel
  ↓
Global / Shared / Local...
  ↓
L1/TEX
  ↓
L2
  ↓
Device Memory

然後問的不是：

哪個數字最大？

而是：

我剛剛改的那一段 code，應該在哪裡留下痕跡？

今天我們改的是：

Global reuse
→
Shared Memory reuse

所以很自然地去找：

Global Requests 有沒有下降？
Shared Memory activity 有沒有出現？
L1 behavior 怎麼變？
差異有沒有一路活到 L2 / Device Memory？
最後 Runtime 又拿回多少？

這就夠了。

Nsight Compute 還有一大堆東西。

Occupancy、Scheduler、Warp Stall、Source Counters、Roofline……

今天全部不管。

後面遇到問題再一個一個學。

---

回頭再看今天整條路：

Source logical accesses
        ↓
      12.8×

Warp-level request prediction
        ↓
        8×

NCU Global Requests
        ↓
        8×  ✓

NCU L1 Sectors
        ↓
        8×  ✓

L1 / Cache behavior
        ↓
差距開始被吃掉

Device Memory
        ↓
幾乎沒變

Runtime
        ↓
約 7%

這個結果我其實比「最後真的快 8×」還喜歡。

因為如果只是快了，我們頂多知道 optimization 成功。

現在我們反而真的看到：

同一個 source transformation，到了不同 abstraction level，會變成完全不同的 performance story。

4096 到 320 沒有錯。

128 到 16 也沒有錯。

NCU 的 4.29G 到 536.87M 更是剛好驗證了我們的 8×。

錯的是如果我們把其中任何一個數字，直接當成最後 runtime 的答案。

所以今天整篇最後其實還是回到很簡單的 workflow：

Source
   ↓
先猜硬體應該做什麼
   ↓
Prediction
   ↓
Nsight Compute
   ↓
Measurement
   ↓
看看 Prediction 活到哪裡
   ↓
找到第一個分岔點
   ↓
重新理解自己的 Kernel

能算的先算。

推不到的就停。

再問硬體。

我反而比較怕一開始什麼 prediction 都沒有，因為那樣 profiler 最後出現什麼數字，都有辦法事後替它講出一個故事。

猜錯沒關係。

至少知道 model 從哪裡開始壞掉。

Day 4，正式開工。