Day 10 寫了AI。

Day 11 講了更多AI。

Day 12 還是AI，這人帆布凡?
煩 。
但暫時還沒處理好跟AI的關係同時也放棄跟cublass比較了
太難了，他們那個利用率簡直離譜媽媽給離譜開門
離譜到家了...
算了... 希望我之後有心整理吧。

下面正文:

Day 10 補了精度。

Day 11 講了加速方向。

今天要把它們放在同一張表裡看：

正確性
速度
GPU profile

只看正確性不夠，因為正確但慢到不能用，工程價值有限。

只看速度也不夠，因為低精度 shortcut 跑很快，但那不是同一個問題。

所以 Day 12 的核心是：

公平比較時，必須先講清楚 precision budget，再看時間和 profile。

先定義公平比較

公平?你先跟我翻譯翻譯
什麼
叫做
他嗎的
公平?

高精度 GEMM 很容易比錯。

例如下面這幾種東西看起來都輸入 double、輸出 double：

native cublasDgemm
cuBLAS emulation fixed mantissa
cuBLAS emulation dynamic
簡化 INT8/CRT fast path
完整 Scheme II target
custom MMA fused path

但它們不是同一個 precision contract。

公平比較時至少要列出：

target mantissa bits
modulus 數量
slice 數量
pair schedule
是否使用 residual
是否 fallback 到 native FP64
max absolute error
max relative error
RMS error
執行時間

如果一個方法只做 34-bit-ish，另一個方法要做 53-bit mantissa target，那它們不能只拿 ms 直接比較。

比較速度以前，要先問：

你們算的是同一個精度問題嗎？

參考比較對象

這裡可以把比較對象分成幾類。

native DGEMM

這是最直覺的 baseline：

cublasDgemm

它走 FP64 硬體路線。

在資料中心 GPU 上，FP64 throughput 可能很強。

但在很多消費級 GPU 上，FP64 throughput 很弱，所以 native DGEMM 會慢很多。

這也是我們會想走 INT8 Tensor Core + Ozaki/CRT 的原因：

用消費級 GPU 比較強的 INT8 Tensor Core，去補 FP64 throughput 不足。

cuBLAS emulation

cuBLAS 13 裡的 emulation 路線可以當成非常強的參考。

它不是 native DGEMM。

它的目標也是用低精度或混合精度硬體去模擬更高精度的 GEMM，而且實作已經非常成熟。

所以拿它比較時，不能只說：

我們比 native DGEMM 快

還要看：

跟 cuBLAS fixed53 emulation 差多少？
跟 cuBLAS dynamic emulation 差多少？

這才是比較接近的對手。

自己的 fused path

自己的 fused path 要再拆開看。

有些版本只是把 INT8 GEMM 換成自己寫的 MMA kernel。

有些版本把幾個 modulus 合成一組。

有些版本把 partial CRT 提前做到 accumulator 附近。

有些版本嘗試把整個 CRT epilogue 全塞進同一顆 kernel。

這些速度和瓶頸都不一樣。

所以 benchmark table 裡不能只寫一個「fused」。

要寫清楚它到底 fused 了什麼。

一組實測數字怎麼看

下面這組數字來自 RTX 4060 Laptop、CUDA Toolkit 13.0 的一次 M9 測試。

數字會受 GPU 溫度、clock、driver、矩陣大小影響，所以不要把它當成永遠固定的 benchmark。

但它很適合用來理解瓶頸。

後續優化後，custom partial CRT 在同一台機器上又看到：

2048^3: 16.40 ms
4096^3: 123.54 ms，熱機時約 131.7 ms

這些數字可以得到幾個結論。

第一，走 INT8/Ozaki/CRT 的方向是有意義的。

因為它比 native DGEMM 快很多：

4096^3 native DGEMM:       669.75 ms
4096^3 custom partial CRT: 123.54 - 138.42 ms

第二，它還沒有追上 cuBLAS emulation：

4096^3 cuBLAS fixed53 emu: 56.86 ms
4096^3 cuBLAS dynamic emu: 74.95 ms
4096^3 custom partial CRT: 123.54 - 138.42 ms

第三，自己寫的 custom MMA 比 cuBLAS backend 快，代表我們的融合方向有回收一些 staged pipeline 的成本：

4096^3 cuBLAS backend SchemeII: 183.86 ms
4096^3 custom partial CRT:      123.54 - 138.42 ms

但還差 cuBLAS emulation 一大截。

這就是 profile 要回答的問題：

差在哪？

Profile 不能只看 Tensor Core

很多人看到 Tensor Core utilization 不高，第一反應是：

那就是 MMA kernel 寫得不好。

可能是，但不一定。

在這種高精度 GEMM 裡，Tensor Core 只是其中一段。

完整 pipeline 還包含：

scale / exponent scan
residue quantization
INT8 GEMM
mod reduction
CRT reconstruction
undo scaling
write double output

所以 NCU 要拆開看。

如果某顆 kernel 的 Tensor pipe 很低，但 FP64 pipe 或 ALU pipe 很高，代表瓶頸可能不是 MMA，而是 epilogue。

如果 DRAM throughput 很高，代表可能是 global memory traffic。

如果 registers per thread 很高，eligible warps 很低，代表可能是 register pressure 讓 SM 沒辦法塞更多 warp。

一次 NCU 結果怎麼解讀

在 512³ 的 profiling 裡，grouped MMA kernel 曾經看到類似：

registers/thread: 218
shared memory/block: 50176 bytes
occupancy limit: 1 block per SM
Tensor pipe: 約 7.5%
DRAM throughput: 約 6.5%
L2 hit rate: 約 94%
eligible warps: 約 0.27

這代表什麼？

首先，它不像是單純 memory bandwidth 爆掉。

因為 DRAM throughput 不高，而且 L2 hit rate 很高。

真正比較可疑的是：

registers/thread 很高
shared memory/block 不小
eligible warps 很低

也就是 SM 裡能同時跑的 warp 太少。

GPU scheduler 每個 cycle 可以發射的 warp 不夠，Tensor Core 自然就吃不滿。

後來把 partial CRT 的壓力移出一部分，讓 register 壓力下降，occupancy 從 1 往 2 靠，4096³ 的時間也跟著改善。

這是一個很典型的 GPU kernel 調校結果：

不是某個數學公式變少了，而是硬體終於能同時排更多 warp 去藏 latency。

Combine kernel 的瓶頸又不一樣

另一顆 kernel 是 partial CRT combine。

它做的是把每組 modulus 的 partial result 再合成最後結果。

這顆 kernel 的 profile 會長得不一樣。

曾經看到：

FP64 pipe utilization: 約 85%
registers/thread: 約 40
512^3 duration: 約 112 us

這表示它不是 Tensor Core kernel。

它比較像 epilogue / reconstruction kernel。

主要成本在：

wide integer reconstruction
CRT combine
scaling 回 double

如果這顆變慢，不能用「多塞一點 mma.sync」解決。

它需要的是：

減少每個 output 要做的 CRT 工作
固定常見 group count 的 specialized kernel
避免 register 變多
避免 vectorize 後 occupancy 掉下去

這也是為什麼有些看起來合理的優化會被拒絕。

例如一次處理更多 output 看起來可以減少 overhead，但如果 registers/thread 從 40 跳到 68，occupancy 掉下去，最後反而更慢。

為什麼真正 full fused 反而可能慢

最直覺的想法是：

把所有 modulus
所有 CRT
所有 scaling
全部放進同一顆 MMA kernel

聽起來很好。

但 profile 可能會告訴你另一件事：

DRAM throughput 很低
L2 hit rate 很高
Tensor pipe 很低
FP64 / ALU pipe 很高

這代表什麼？

代表你不是被 memory 卡住。

你是把太多 scalar reconstruction 塞到 Tensor Core kernel 裡，讓 warp 做完 MMA 後卡在 CRT / FP64 epilogue。

這種情況下，更多融合不一定更快。

比較好的方向可能是 hierarchical：

每組 modulus 先做 partial CRT
把 group partial 寫出去
最後用 specialized combine kernel 合併

也就是不要讓單一 kernel 同時背所有工作。

這聽起來沒有「full fused」帥，但比較符合硬體。

為什麼 tile-local prologue 沒有直接贏

還有一個很 tempting 的方向：

不要 materialize A_mod/B_mod
每個 CTA 直接讀 double A/B
在 tile 裡 quantize 成 residue
馬上餵 Tensor Core

這很像 FlashAttention 的精神。

但實測上它可能非常慢。

原因是 A/B residue quantization 不是便宜操作：

double load
exponent scaling
round
mod
balanced int8 conversion

如果這些工作在每個 K tile 都重做一次，scalar 成本會非常重。

所以 tile-local prologue 要看情境：

這也是 Day 11 說的：

FlashAttention 精神不是盲目不落地，而是只保留對總時間真的有利的 on-chip 狀態。

用 profile 決定下一步

看到上面的數字，我們下一步不該亂猜。

如果目標是追 cuBLAS emulation，接下來要優先處理：

1. grouped MMA 的 register pressure
2. partial CRT epilogue 的 scalar 指令數
3. A/B residue materialization 的重用策略
4. common group count 的 specialized combine
5. tile shape / stage / swizzle autotune

比較不該做的是：

看到 fused 就開心
看到 tile 變大就以為會更快
看到 Tensor Core utilization 低就只調 MMA
把 residual 或 native FP64 fallback 混進比較裡

profile 的價值就是讓我們不要被直覺騙。

今天的重點

Day 12 的結論可以整理成幾句。

1. 先定義 precision contract，再比較速度。
   不同 mantissa target、不同 modulus budget、不同 residual/fallback 設定，不能混在同一張速度表裡當同一件事。

2. 比 native DGEMM 快，不代表已經追上真正強的對手。
   在消費級 GPU 上，native FP64 很慢；更接近的對手是 cuBLAS emulation。

3. custom MMA 有用，但瓶頸不只在 MMA。
   grouped MMA、CRT epilogue、combine、quantize、global memory round-trip 都會吃時間。

4. NCU 會告訴你瓶頸在哪個硬體單元。
   Tensor pipe、FP64 pipe、ALU、DRAM、L2、occupancy、register pressure 要一起看。

5. 下一步不是追求名義上的 full fusion，而是追求總時間下降。
   有些東西該留在 accumulator 附近，有些東西該拆出去給 specialized kernel，有些東西該預先 materialize。

到這裡，Day 10 到 Day 12 的主線就收起來了：

Day 10：精度怎麼補齊
Day 11：補齊後怎麼加速
Day 12：用 benchmark 和 profile 看差距

接下來如果要繼續往下寫，就可以開始拆更細的 GPU kernel 主題：

shared memory layout
ldmatrix lane mapping
register pressure
occupancy
NCU 指標解讀
CRT epilogue 設計

之後應該會放一些比較簡單的主題? 不知道
砸們且看且走 這次會努力發完吧 大概
我還想著 找時間 把我的PR在這邊解釋一下 之類的
唉 時乃 多事之秋阿?
入秋了 大家記得出門帶把傘