Day 4｜GEMM 要高精度的時候，4060 先把路封了

知恥而後勇，是我的一種信念。
但有時候我也覺得自己已經不是勇不勇，是有點傻了，哈哈哈哈。
言而總之，今天是很雷包的一天。
但是文章還是要繼續。

昨天我們說了很多。雖然還沒完整看過 NCU，也還沒帶你看看怎麼點 NCU。
不過我想 blog 這東西，也不是按照順序的。
你想怎麼看，就怎麼看。如果這個剛好有用，管他前幾天寫的是蝦米碗糕。
如果這篇有用，我就看。也沒有什麼因果鏈。

就是這麼的簡單粗暴：有幫助 就會有觀眾 不管是很吸引眼球的幫助 還是很引起共鳴的內容。
所以今天開始會放飛一點，想寫啥就寫啥，還望多多包涵。

我記得今年夏天有報名 COSCUP，主題是 Adaptive GEMM。
對，就是我們 Day 3 的例子：GEMM。但是這次寫了很多顆 kernel，然後用了一些我自己也沒有很懂的技術，哈哈哈。
只是礙於不是很會講，所以我覺得那時候講得不是很好。
畢竟就不好講啊 O3O。
這次就利用這個機會，多多展開來說說。

完整的連結是這個。後面幾天會把裡面的東西拆開來講：
AdaptiveGEMM

裡面大概寫了……很多。但今天不打算把 repo 打開來從頭掃一次。

因為如果一開始就講 Ozaki、CRT、七個質數、mma.sync，大家大概會跟我夏天那時一樣：聽完以後只記得「好像很厲害」，卻不知道為什麼要這麼繞 而且講不清楚 是講者的問題 不是觀眾的問題。

所以 Day 4 只做三件事：動機、硬體差距、後面幾天的目錄。

先從 GEMM 看起

GEMM 有很多種算法可以去算。

Day 2、Day 3 我們走的是最普通的那條：一個 thread 負責一個 C[i,j]，再改成 Shared Memory tiling。那是為了看懂 GPU。真實世界裡，同一句 C = A × B 後面其實可以換成很多條路：

naive  FP32
tiled  FP32
cuBLAS FP32
cuBLAS FP16 / TF32 / Tensor Core
cuBLAS FP64
以及一堆手刻 kernel

其中有一條路，需要很高的精度。

不是所有 GEMM 都要 double。深度學習常常 FP16 就夠了。但科學計算、數值線性代數、量化金融、或任何會一路累加、誤差邊界要講得清楚的場景，直覺還是會走 FP64：

累加深
誤差要可控
對答案要對得回 double
→ 先開 FP64

然後你在消費級 Ada 上按下執行，會遇到一個很現實的狀況：

FP64 的效能被閹割得很慘。

這不是 kernel 寫得醜。Day 3 那種「改 tiling 就會變好」在這裡不夠用。因為瓶頸在硬體規格本身。

Day 1 講過：沒有比較對象的 100 ms，不是快也不是慢。
那現在比較對象其實很清楚。不是「我再寫一顆比較勤勞的 FP64 kernel」，而是：

這張卡的 FP64 單元，本來就沒打算讓你拿來做大事。

那怎麼救？其實有幾套方法。有人直接換資料中心的卡；有人接受降精度；我們這條路比較繞——先把數字切開，丟進這張卡真正有算力的地方，再拼回來。

今天先把「為什麼要救」講清楚。怎麼切、怎麼拼，是後面幾天的事。

這幾天打算怎麼走

夏天那場 talk 不好講，有很大一部分是因為我把整條路一次倒出來 而且我講得很兒戲。
這次改成跟 Day 3 一樣：一天只改一個東西。

Day 4   動機、目錄、硬體差距          ← 今天
Day 5   數學：為什麼切 + 拼行得通
Day 6   把數學拆成三個步驟、三顆 kernel
Day 7   (a) 建構 CRT
Day 8   (b) 拆解、求餘數
Day 9   (c) 重建、把數字拼回來
Day 10  三顆融在一起：fused kernel

Day 7、Day 8 會穿一些我們踩過的坑。Ada 的 PTX 手冊有些地方寫得不清楚，甚至對不太上，實作時都撞過。

Day 10 之後還有對打、NCU、以及別的題。那些先不佔今天。技術主線先走到 fused 為止。

硬體上到底差多少

RTX 4060 是 Ada Lovelace，SM 8.9。如果只看深度學習常用的路徑，它其實不算弱。FP32、INT8 Tensor Core 都還有一口飯吃。

但如果看 FP64，數字會變得很難看：

同一張 4060 上，FP32 跟 FP64 的 peak 差了大約 64:1。
這不是我量測量歪，Nsight Compute 自己也會寫這句話：

The ratio of peak float (fp32) to double (fp64) performance on this device is 64:1.

消費級卡把 FP64 砍得很兇。資料中心的 A100 則不是這樣。所以當 GEMM 這條路堅持要 double，在 4060 上會得到一個很現實的結論：

演算法要 double
GPU 最會的是 Tensor Core
cuBLAS FP64 在消費級卡上
常常被規格直接卡死

可以先想成 Day 1 那間工廠。

工廠裡有很多機器。有的機器一天能做很多螺絲，有的機器理論上能做更精密的螺絲，但工廠只肯放很少台。
你如果堅持所有訂單都走那條精密產線，不是師傅不認真，是產線一開始就那麼窄。

4060 上真正寬的路
    INT8 Tensor Core

4060 上被砍很兇的路
    FP64 CUDA Core

規格表看完還是有點抽象。所以我們在同一張卡上，讓同一批 thread 分別狂打不同的運算單元。不是完整 GEMM，只問相對速度。量出來大概是這樣：

[AdaptiveOzaki] Hardware Profile (vs INT8 TC)
  FP16 TC :    0.98x slower
  TF32 TC :    2.01x slower
  INT32 CC:  114.46x slower
  FP32 CC :  131.12x slower
  FP64 CC :  138.52x slower

讀法很粗暴：

INT8 Tensor Core 如果是 1
FP64 CUDA Core 大概只剩 1/138

這不是跟 cuBLAS 公平對打，也還不是最終 GEMM 的 TFLOP/s。
它只回答一個更前面的問題：

同一張卡、同一批 thread，走不同的運算單元，速度差多少？

差到這個程度，後面開始做一些看起來很神經的事情，就比較說得通了。

不是因為 INT8 比較高貴。
是因為 這張卡會的事情，跟高精度 GEMM 要的精度，不是同一條路。

官方函式庫也救不了這件事。cublasDgemm 在 4060 上可以把 FP64 pipe 打到 90% 以上，看起來很忙；但 Memory 幾乎沒在動，occupancy 也很低。意思是資料不是搬不進來，是精密產線本身太窄。

Day 3 的大頭在 data movement：大家一直來回 Global Memory，所以 tiling 有意義。
今天的大頭是：高精度 GEMM 走了這張卡最弱的那條 execution pipe。

你再怎麼 tiling，也改變不了 64:1。

今天先停在這裡

所以這系列要救的，不是「再寫一顆比較快的 FP32 matmul」。
是：GEMM 有一條路需要高精度，而消費級卡把那條路封了。我們要想辦法，把答案救回這張卡真正有吞吐的地方。

救法有幾套。最直覺的是直接降成 FP16——昨天 microbench 也顯示它幾乎跟 INT8 一樣快。

明天會問一個比較討厭的問題：

既然 FP16 Tensor Core 這麼快，為什麼不直接把 FP64 降成 FP16？為什麼還要切開再拼回來？

今天就看到這裡就好。雷包的一天，能把動機跟目錄講清楚，就算有賺到。