這篇是重賽版。

原版 Day 4 有一段「今天很雷包，但文章還是要繼續」的心情開場。那段不是不要了，只是這次先收去文末註解。因為今天要講的不是我的心情，是一個更欠揍的現實：

你想做高精度 GEMM
你手上是 RTX 4060
你按下 FP64
GPU：你確定嗎.jpg

Day 2、Day 3 我們先用普通 GEMM 學 GPU：一個 thread 算一個 C[i,j]，再用 Shared Memory tiling 把 data movement 改好。

但真實世界裡，同一句：

C = A x B

後面可以接很多條路：

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

其中有一條路很麻煩：高精度 GEMM。

不是所有 GEMM 都要 double。深度學習常常 FP16、BF16、TF32 就可以跑得很開心。但科學計算、數值線性代數、量化金融、或任何一路累加很深、誤差邊界要說清楚的地方，直覺會先想：

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

然後你在 RTX 4060 上開 FP64。

畫面大概像這樣：

我：我要高精度。
4060：可以。
我：那我要快。
4060：剛剛不是已經說可以了嗎？

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

為什麼偏偏用 RTX 4060

先講清楚：選 RTX 4060 不是因為它是最適合高精度計算的卡。

剛好相反。

我想用它，是因為它很適合把矛盾放大給你看：
(才不是 明明就是因為我只有這張卡 哈哈哈哈)

1. 它是消費級 Ada Lovelace，CUDA capability 是 8.9。
2. 它有第四代 Tensor Core，低精度矩陣乘法很肥。
3. 它的 FP64 非常瘦，瘦到不像是同一張卡。
4. 它便宜、常見、很可能就是一般人手上真的有的卡。

如果一開始就拿 A100 來講，故事會變得太舒服。A100 本來就是資料中心卡，FP64 是一等公民。可是很多人手上的不是 A100，是一張遊戲卡或筆電卡。

所以 4060 很適合當這系列的反派。

不是因為它爛。

而是因為

而是因為它會逼你承認：

硬體最會的事情
    不是你數學上最想做的事情

這也是 AdaptiveGEMM 會長成 Ozaki + CRT + INT8 Tensor Core 的原因。

完整程式線在這裡，後面幾天會把它拆開來講：

AdaptiveGEMM

先把規格攤開

NVIDIA 官方 RTX 4060 規格列的是：

來源是 NVIDIA 的 RTX 4060 官方產品頁。它沒有把每一種 Tensor Core mode 都拆成 dense / sparse / INT8 / INT4 給你看，只寫「242 AI TOPS」。所以這裡要把口徑講清楚：

官方頁的 242 AI TOPS
    是 marketing-level AI TOPS
    不是「我們這顆 s8 x s8 GEMM 一定能跑到 242 TOPS」

在 Ada Tensor Core 的常見拆法裡，可以粗略這樣看：

這張表有兩種來源混在一起，所以我直接標明：

FP32
    3072 CUDA cores x 2 ops/clock x 2.46 GHz
    ~= 15.1 TFLOPS

FP64
    Ada sm_89 每 SM 是 128 個 FP32 core，但只有 2 個 FP64 core
    所以 FP64:FP32 ~= 1:64
    15.1 / 64 ~= 0.236 TFLOPS

INT8 / INT4 Tensor Core
    由官方 242 AI TOPS 和常見 dense/sparse 口徑推回來
    實際 GEMM 還要看 kernel、layout、TGP、clock、occupancy

這裡先停一下。

原版 Day 4 的表把 RTX 4060 寫成：

FP32 ~242 GFLOPS
FP64 ~3.8 GFLOPS

那個數字少算了大約 64 倍。比較合理的是：

FP32 ~15.1 TFLOPS
FP64 ~236 GFLOPS

但重點沒有變。

重點是比例：

FP32 : FP64 ~= 64 : 1
dense INT8 TC : FP64 ~= 121 TOPS : 0.236 TFLOPS

如果只看「operation count」的理論峰值，INT8 Tensor Core 對 FP64 CUDA Core 是幾百倍等級的差距。當然 TOPS 跟 TFLOPS 不能直接當成同一種科學計算能力，INT8 也不會神奇變成 double。可是它告訴我們一件事：

這張卡真正寬的路
    是低精度 Tensor Core

這張卡最窄的路之一
    是 FP64 CUDA Core

梗圖位：

FP64：我可以精準。
Tensor Core：我可以很快。
4060：你們兩個可以不要同時要求我嗎？

跟 A100 比，差距更明顯

A100 是另一個世界。NVIDIA A100 官方規格列：

你會看到一個很有趣的反差。

RTX 4060 的 FP32 跟 A100 不是差到完全不能看。可是 FP64 一比下去，A100 直接拉開到數十倍，而且 A100 還有 FP64 Tensor Core。

也就是說，當你寫：

cublasDgemm(...)

在 A100 上，這是在叫資料中心卡做它本來就擅長的事。

在 RTX 4060 上，這比較像是在遊戲機旁邊貼一張紙：

即日起改營業項目：高精度科學計算

GPU 不是不會做。它會。

但它不是為這條路長大的。

CUDA 文件其實已經把答案寫在表裡

CUDA Programming Guide 對 compute capability 8.9 的 SM 描述很直接：

每個 SM：
    128 FP32 cores
    2 FP64 cores
    4 fourth-generation Tensor Cores

這就是 64:1 的硬體根源。

所以不要先懷疑自己是不是少加了一個 #pragma unroll。

Day 3 的主題是 data movement：

Global Memory 太常來回
=> 用 Shared Memory tiling
=> 讓資料重用

Day 4 的主題不是這個。

今天的大頭是 execution pipe：

高精度 GEMM 走 FP64
但 4060 的 FP64 pipe 本來就很窄

這時候你再怎麼 tiling，都改變不了：

FP64 單元就是少

這也是為什麼我覺得 Day 4 必須先講硬體，而不是直接跳 Ozaki 或 CRT。你如果不知道這條路為什麼被封，後面看到「把 double 切成整數薄片，再拆成 7 個小質數餘數，最後丟 INT8 Tensor Core」會覺得作者是不是太晚睡。

先說，作者確實常常太晚睡。

但這次不是因為太晚睡才這樣寫。

我們自己的 microbench 看起來更殘酷

規格表還是有點抽象。所以我們在同一張卡上，讓同一批 thread 分別狂打不同運算路徑。這不是完整 GEMM，也不是 cuBLAS 公平賽，只是問：

同一張卡、同一種 thread 壓力，打不同 execution unit，速度差多少？

當時量到的相對速度大概是：

[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

這個數字不是官方規格，是我們自己的 microbenchmark。它不該拿來取代官方 peak，也不該拿來說「所有 GEMM 都會差 138 倍」。

它只支援一個很前面的判斷：

在 4060 上，低精度 Tensor Core 跟 FP64 CUDA Core
不是同一個量級的路

這跟一些近期 GPU microbenchmark 論文的精神也一致：官方白皮書會給高層規格，但實際 kernel 會卡在 instruction issue、memory hierarchy、Tensor Core pipeline、同步成本、layout 等等細節。DrGPU 是 profiling 方法論；Jarmusch 等人的 Blackwell microbenchmark / performance modeling 論文則提醒我們，很多「看起來是 peak」的東西，真的要靠微基準和模型拆開。

不過 Day 4 這裡要守住範圍：那些 Blackwell 論文不是拿來推 RTX 4060 數字的。它們只是提醒我們：

官方公布值
microbenchmark 實測值
架構推估值

這三種東西要分開標。

不然表格會長得很好看，但讀者會被你帶去撞牆。

所以為什麼不是直接用 FP16 或 INT4

這裡先預告，不展開證明。

既然 4060 的低精度 Tensor Core 很快，那最直覺的救法是：

double 太慢
=> 轉 FP16
=> Tensor Core 跑一跑
=> 轉回 double
=> 收工

如果你只是要「大概差不多」，這有時候可以。

但如果你要的是「比普通 FP32 更接近 FP64 GEMM」，問題就來了：

FP16 rounding 之後
丟掉的 bits 回不來

INT4 更極端。INT4 的理論吞吐更高，官方 242 AI TOPS 也可以用 dense INT4 的角度理解。但 INT4 每個 operand 只有 4-bit payload，拿來直接表示高精度數字更不可能。

所以我們後面不是說：

INT8 比 FP16 精準
INT4 又更神

不是。

INT8 很不精準。INT4 更不精準。

真正的重點是：

低精度整數可以拿來裝「餘數」
餘數可以用 CRT 拼回比較大的整數

這是 Day 5 要講的事。

今天先記住一個方向就好：

FP16 / INT4 / INT8 本身不是答案
Tensor Core 的吞吐才是誘因

Ozaki + CRT 的作用
是把高精度問題改造成 Tensor Core 願意吞的形狀

先把話講保守

這裡很重要。

我們現在這條 Phase 22，不是 full FP64 precision support。

比較準確的說法是：

FP64 input/output
約 34-bit mantissa 的 Ozaki/CRT 實驗路徑
在 2048^3、U(-1,1) 測試上 max error 約 1e-9
比普通 FP32 好
但不是 cublasDgemm 的 ulp
也不是 1e-15

我們要講的是：

消費級 Ada 的 FP64 很窄
低精度 Tensor Core 很寬
所以我們嘗試把部分高精度 GEMM 需求
搬到 INT8 Tensor Core 上

這已經很有趣了，不需要吹成魔法。

這幾天打算怎麼走

這次不要一次把 Ozaki、CRT、七個質數、mma.sync、ldmatrix、bank conflict 全部倒在讀者臉上。

夏天那場 talk 我就犯過這個錯。講者覺得自己很努力，聽眾只覺得：

資訊量：滿
理解度：空

所以重賽版跟 Day 3 一樣：一天只改一個東西。

Day 4   RE: 動機、目錄、RTX 4060 硬體差距     <= 今天
Day 5   數學：為什麼切 + 拼行得通
Day 6   把數學拆成三步：建 CRT、求餘數、拼回來
Day 7   Roadmap (a)：建構 CRT
Day 8   INT8 GEMM：naive，先對答案
Day 9   INT8 GEMM：tile + mma.sync
Day 10  把 INT8 乘法吃滿，對上 cuBLAS emu

Day 10 之後還會有對打、NCU、底層工具，以及一些走錯的路。

但今天先不要貪心。

今天只要把這個問題釘住：

為什麼高精度 GEMM 在 RTX 4060 上，不能只靠「再寫一顆比較勤勞的 FP64 kernel」解決？

答案是：

因為規格本身就把 FP64 路修得很窄。

今天先停在這裡

所以這系列要救的，不是「再寫一顆比較快的 FP32 matmul」。

它真正想問的是：

GEMM 有一條路需要高精度
消費級 Ada 把 FP64 做得很窄
低精度 Tensor Core 卻很寬

那我們能不能把數字切開、拆成餘數
讓 Tensor Core 做它擅長的乘法
最後再把結果拼回來？

明天會問一個更討厭的問題：

既然 FP16 / INT4 / INT8 Tensor Core 都很快，為什麼不直接降精度？為什麼還要 Ozaki 切、CRT 拼？

今天就先看到這裡。

4060 沒有錯。

它只是很誠實地告訴你：

我是一張消費級卡。
你要我做資料中心卡的事，可以。
但你最好準備一點數學。

機哩瓜喇說了一大堆 唉 結果能看得沒兩句?
所以我說人生就是在裝忙。你知道整篇文章好像寫了一堆狗屁不通的東西，也沒有到狗屁不通吧，但我不知道欸，會有一種很裝忙的感覺。

有時候我得去講一下這個硬體性能，有時候得講一下數學，有時候又得強調這東西根本就不是標題講得那麼浮誇。

當然啦，我們最後還是會讓它用得跟標題一樣浮誇，應該啦。對，沒錯，不然的話就對不起標題了嘛，反正啊
我不知道拉~

註：原版 Day 4 的心情開場

原版開頭大概是這樣：

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

這段我還是留著。因為它其實很符合這條線的狀態。

AdaptiveGEMM 不是一開始就很漂亮。中間有很多講錯、量錯、表格寫錯、以為懂了結果沒懂的地方。這篇重賽版把心情放到後面，是因為前面要先把硬體帳講乾淨。

雷包可以放註解。

數字不要雷包。

參考來源

• NVIDIA GeForce RTX 4060 官方規格
• NVIDIA CUDA GPU Compute Capability 表
• CUDA C++ Programming Guide，compute capability 8.9 的 SM 組成與 arithmetic throughput
• NVIDIA A100 官方規格
• NVIDIA Ada Architecture / Tuning Guide
• DrGPU: A Top-Down Profiler for GPU Applications, ICPE 2023
• Jarmusch and Chandrasekaran, Microbenchmarking NVIDIA's Blackwell Architecture, arXiv:2512.02189v3, 2026-03-02
• Jarmusch and Chandrasekaran, Microbenchmark-Driven Analytical Performance Modeling Across Modern GPU Architectures, arXiv:2605.04178v1, 2026-05-05
• Zhihu 彙整文