[趣味競賽 進階 #12] 為什麼你的 4-bit 模型在 2080 Ti 上快不起來?——拆開 .so 才發現這代卡只有兩種齒輪

前言

把貨從大箱換成小箱，省下的是倉庫空間;搬運工卻堅持每箱到站台都要拆開，裝回大箱才准搬上車——省的是倉租，不是搬運的工時。你可以換更輕的搬運工，也可以吵著要修規矩，省下來的錢從頭到尾就沒進過工時那一欄。

這是改裝 2080 Ti 22G系列第 12 篇，直接接在 #11 後面。那篇用 INT8 把 MiniMax-H3 跑起來，也順手判死了 W4A4;但有個問題當時沒解:INT8 跟 W4A4 之間，為什麼沒有一條「權重 4-bit、品質不崩」的中間路線?這篇把那個問題一路拆解到硬體指令層。

答案不在軟體，在這張卡的硬體裡。拆開 CUDA 後端的 .so 檔案才看懂:不是誰忘了寫某條加速路徑，是那條路徑從一開始就不存在。

量化在這張卡上省的是空間，不是時間

你的 GGUF Q4 權重比 fp16 小一半，直覺上該更快。實測常常不是這樣:tok/s 沒變，有時候還更慢。

最乾脆的反例是 IQ3_M——比 Q4_K 更激進的量化格式，想說壓得更狠總該跑快一點。三種工作負載的生成速度(tok/s)全部反著來:struct 從 45.9 掉到 37.1 tok/s，慢了 19%;reasoning 從 45.3 掉到 39.9 tok/s，慢了 12%;create 從 29.3 掉到 28.6 tok/s，慢了 3%。檔案更小，反而更慢。

這篇要講的就是為什麼——答案跟「壓得越狠越快」正好相反。

這張卡的整數 tensor core 只有兩種齒輪:s4×s4 跟 s8×s8

模型推論的算力幾乎都花在 GEMM(General Matrix Multiply，矩陣乘法)上;GEMM 跑在 tensor core 這種專門單元上。量化能不能真的變快，取決於 tensor core 支不支援你選的那個位元寬度組合。

拆開 comfy-kitchen 的 CUDA 後端 _C.abi3.so，按架構跟運算元型別分類，結果是這張表——看 s4×s8 跟 s4×f16 那兩欄，全空:

架構	運算元組合	支援?	每種輸出精度的符號數
Sm75(Turing，這張卡)	s4 × s4	✓	32 個
Sm75(Turing，這張卡)	s8 × s8	✓	64 個
Sm75(Turing，這張卡)	s4 × s8 / s4 × f16	✗ 不存在	0 個
Sm80(Ampere)	s8 × s8	✓	96 個
Sm89(Ada)	s4 × s4	✓	56 個

mma.s4.s4 跟 mma.s8.s8 這兩條硬體指令存在;mma.s4.s8 不存在。兩個運算元(相乘的兩邊)必須同寬——你不能拿 4-bit 的權重去乘一個 16-bit 的 activation(活化值，模型算到一半的中間結果)。

「保留 4-bit 權重、把活化拉回 fp16」在硬體層不存在——不是軟體還沒做，是指令不存在。

所以你的 4-bit 權重每一層都在被拆回 fp16

s4×s8 不存在，但你手上的 GGUF Q4 明明能跑。它是怎麼跑起來的?

答案是反量化(dequantize，把壓縮過的低位元權重還原成計算用的浮點數)。GGUF、AWQ、GPTQ、Marlin 全部走同一條路:先把 4-bit 權重展開回 fp16，再丟進純 fp16 的 GEMM 去算。TurboMind 論文(arXiv 2508.15601v2)講得很白:INT4 載入後先做 I2F 轉換、dequant 回 FP16，再進 FP16 Tensor Core 算——跟 GGUF Q4 是同一個「dequant-to-fp16」原理。

IQ3_M 為什麼比 Q4_K 慢，答案也在這裡。Q4_K 是線性分塊量化，反量化只是一個乘加;IQ3_M 是 codebook 查表(非線性)，每個權重的反量化計算比線性塊重很多——瓶頸從 VRAM 頻寬搬到了反量化計算本身。

「檔案小=更快」要成立，得同時滿足兩個前提:(a) 真的卡在頻寬，(b) 反量化便宜。Turing 兩個都不成立——每一層多付的那次反量化，把計算量墊高了。

純 4-bit(W4A4)有原生路，但誤差是 INT8 的 18 倍

s4×s4 是存在的，那全用 4-bit(W4A4，權重跟 activation 都是 4-bit)豈不是真的能跑原生加速?

kernel 級測出來是真的快。blocks.0.attn.out_proj 這一層，eager 模式要 19.256 ms，走原生 cuda kernel 只要 1.373 ms，快 14.02 倍;最大的那個矩陣 eager 模式直接 OOM，cuda kernel 卻能在 3.063 ms 跑完。不過這個倍數是在一張快滿的 GPU 上量的，環境可能放大了差距;端到端的 fp32 生成流程裡實測只快 6.3%。

代價在畫質。microbenchmark 測出的相對誤差:INT8 是 0.0110，W4A4 是 0.2005——差 18 倍。完整的判死過程(縮放掃描、畫面壞法)在 #11 的「W4A4 為什麼不行」一節。

所以這張卡上的 4-bit 只有兩階，沒有中間帶:W4A4 快但畫質會壞，INT8 有原生加速、品質也守得住。

2080 Ti 量化實用選擇

這表看哪格:找到最像你的那一列，直接抄那一列的選擇。

你的情況	選擇	為什麼
模型塞不下 VRAM	GGUF Q4	能跑，但別期待反量化之後還變快
塞得下、要速度	W8A8-INT8	唯一有原生加速的量化，對 FP16 有 2.07–2.83×
遇到 NVFP4 / FP8 選項	別碰	這張卡是軟體模擬:GGUF Q4_K_M 53.083 秒 vs NVFP4 142.434 秒(影片生成實測)，GGUF 反而快 2.68 倍
想榨到極限(W4A4)	只在畫質不重要時	誤差 18 倍，DiT 這種生成模型會畫質崩(詳 #11)

INT8 也不是無腦套。同一批測試裡，Wan 2.2 的 LowNoise 那顆 DiT 換成 W8A8 之後 SSIM 掉到 0.345642，畫面變成彩色碎石雜訊——量化前先跑一眼品質對照，別把「這張卡唯一有加速的路」當成「這張卡永遠該用的路」。

回到貨運那個比喻:這張卡上的量化，省的是倉租。工時那一欄，一毛都沒少付。

---

進階:W4A8 偵探日記

不讀這節不影響使用。上面那張選擇表是結論，這節是怎麼查到「W4A8 為什麼沒有比作者宣稱的快」這個具體問題的過程，包含一次追錯方向、兩顆沿途踩到的地雷，還有我自己七月寫的判斷被推翻兩次的紀錄。

少讀 36% 的權重，為什麼一秒都沒省到?

W4A8(權重 4-bit、activation 8-bit)這個量化格式，comfy-kitchen 的作者宣稱速度是 int8 的 ~1.09 倍。同一顆 DiT(就是 #11 跑的那顆 H3)，W4A8 版本 12.54 GB，比 INT8 的 19.53 GB 少讀 36% 的權重。

我當時的預想很直接:少讀 36%，即使算力沒有差異，光是頻寬省下來的時間就該讓它比 INT8 快，至少不會慢。

實際跑一組固定條件對照(864×480、124 幀、14 步、Spectrum 開、seed 42):

組態	VRAM	秒數	載入狀態
INT8	19.53 GB	238.55	loaded partially，778.39 MB 卸載到系統記憶體
W4A8	12.54 GB	454.19	loaded completely，零卸載

W4A8 慢了 90%。而且它是在零卸載、VRAM 綽綽有餘的最佳條件下慢的——這排除了我原本可能會怪的那個變因:**不是卸載拖慢的，連需要往系統記憶體搬資料的機會都沒有，單純算得比較久。**自己的實驗數字，直接打掉了「少讀 36% 權重至少該省下一點時間」這個預想。慢在計算，不在頻寬，調 --reserve-vram 對這個問題沒用。

這個模式不只在 W4A8 上出現。H3 DiT(影片模型)那邊量過一組類似的對照:Q3 GGUF 全部常駐、完全不卸載要 189.23 秒，INT8 帶卸載反而只要 110 秒——慢了 72%(2026-08-05，ai-lab)。要注意，這組數字不能直接拿來證明「反量化本身就是成本」——變因沒隔離乾淨:兩組測試是 Q3_K_M 跟 INT8-ConvRot 不同量化等級，駐留狀態也不一樣，嚴格成立的只有「Q3 GGUF 全駐仍輸給 INT8 帶卸載」這一句。真正乾淨的隔離證明還是上面 W4A8 那一對:零卸載對零卸載，454.19 秒 vs 238.55 秒。

banner 都印了 asym_w4a8_int8，為什麼還是慢?

第一個懷疑方向是:也許系統根本沒把這套量化當成原生格式處理，只是包裝成 W4A8 的名字，實際上走了某種更慢的通用路徑。

啟動 banner 看起來排除了這個猜測——Native ops: 那行清清楚楚印著 asym_w4a8_int8，像是系統確實把它當原生算子在處理。但翻開程式碼一看，猜測其實是對的:

# comfy_kitchen/backends/cuda/__init__.py:2258 附近
used = _C.cutlass_int8_dequant(...)      # Turing 回 False(_cuda_device_supports_cutlass_int8_dequant 要求 major >= 8)
if not used:
    return eager_w4a8_int8_linear(...)   # ← 掉進慢路

_cuda_device_supports_cutlass_int8_dequant 這道檢查要求 major >= 8，而 Turing 的 compute capability major 是 7。guard 直接回傳 False，W4A8 真的掉進 eager_w4a8_int8_linear——跟一開始懷疑的那條更慢的通用路徑是同一件事。

這一發的教訓不是「猜對了」，是分清楚 banner 在講什麼:Native ops: 那行只代表系統認得這個算子的名字，不代表它有快速路徑可用。而且不只這一次:

看到的字串	為什麼不是證據
Native ops:	那是「沒被列進 disabled」，不是「有硬體加速」
--force-fp16	只管權重儲存格式
Using sage attention	只表示旗標被接受
pip install sageattention==2.2.0 成功	裝得起來 ≠ 這張卡真的用得到那顆 kernel

asym_w4a8_int8 印在 Native ops: 那行，結果慢 90%——同一套邏輯又中一次。

guard 背後是軟體沒做，還是硬體沒有?

找到 arch guard 之後，問題變成:這道 major >= 8 是保守設定(cutlass 模板單純沒為 sm_75 編過)，還是底層真的沒有對應指令?我原本的假設是前者——猜測只是 cutlass 的融合反量化模板沒替 Turing 編譯，補上去應該就能跑。

要驗證這個猜測，得繞過應用層，直接看 .so 檔案裡實際存在哪些 kernel。用 nm -D 列出 comfy_kitchen/backends/cuda/_C.abi3.so 的動態符號表，再用 c++filt 還原成可讀的 C++ 函式簽名，按 cutlass::arch::SmXX 跟運算元型別分類，得到這張完整表:

架構	運算元(A × B → 輸出)	每種輸出精度的符號數
Sm75(Turing)	s4 × s4 → bf16 / f16 / f32	32 個
Sm75(Turing)	s8 × s8 → bf16 / f16 / f32	64 個
Sm80(Ampere)	s8 × s8 → bf16 / f16 / f32	96 個
Sm89(Ada)	s4 × s4 → bf16 / f16 / f32	56 個

三個架構、每個架構檢查過的所有運算元組合裡，沒有任何一格是 s4 × s8 或 s4 × f16。這不是某個架構漏編，是所有架構都沒有——.so 裡根本不存在對應的符號可以連結。

我原本的假設是「cutlass 模板沒編」，這是軟體層的缺口，理論上補上去就能修;實際情況是指令集本身就沒有這個組合，不是編譯設定能解的問題。這兩件事差很多:前者等更新就好，後者只能換方法。

沿途撿到的兩顆地雷

追這條線的路上踩到兩個跟主線無關、但各自燒了時間的坑。

METADATA 沒列的相依套件

comfy-kitchen 0.2.28 的 wheel METADATA 沒有把 torch 列進依賴清單。看起來像「這個套件不吃特定 torch 版本」，實際上它的 na.py 用了 PEP-585 的型別寫法 list[int]，而 torch 2.6.0 的 infer_schema 只吃 typing.List[int] 這種舊式寫法。兩邊都沒報錯，直到實際呼叫那段 schema 推導才炸——線上服務因此中斷了 4 分鐘。METADATA 沒列出依賴，不代表沒有依賴，只代表沒有人把它寫進去。

自家 rescale 腳本擋住 W4A8

追另一次失敗，發現連環境本身都沒問題的組態也跑不動 W4A8。原因出在自己寫的 h3_rescale.py——當 SCALE=1.0 時，程式邏輯上這一步該是空轉，但程式碼還是照樣呼叫了 sc.mul_(INV)。而 W4A8 的 weight_s_rel 存的是 fp8 格式，Turing 沒有 fp8 運算單元，直接丟出 "mul_cuda" not implemented for 'Float8_e4m3fn'。

為舊組態留下的程式碼，會在新組態上變成擋路石，而且它不會說自己是空轉的。

化石與翻案:同一個判斷被自己推翻了兩次

七月的時候我自己寫過一份 2080 Ti 的量化選擇表，結論是:AWQ / W4A16 在 Turing 上用的是原生 mma.sync.s4、Marlin 只支援 sm_80 以上、AWQ 比 GGUF 快 1.5 到 2 倍。

六天後讀到 TurboMind 那篇論文(arXiv 2508.15601v2)，推翻了第一條:AWQ / W4A16 在 Turing 上並不是走原生 mma.sync.s4——論文講的路徑是 INT4 載入後先做 I2F 轉換、dequant 回 FP16，再進 FP16 Tensor Core，跟 GGUF Q4 是同一個 dequant-to-fp16 原理。

一個月後查 vLLM v0.27.1 的文件，又推翻了第二條:Marlin 自約 v0.14 起就把 sm_75 加進支援範圍，不是我原本以為的 sm_80 以上限定。

兩條都被推翻，但底層判斷沒有變。Marlin 支援 sm_75 這件事本身沒有繞過 Turing 缺 s4×f16 這個事實——Marlin 依然是反量化 kernel，只是把反量化這段也做得比 eager 模式快一些。batch=1 的判斷從頭到尾沒動過:這張卡上沒有一種 4-bit 量化格式，是靠硬體原生的混合位寬矩陣乘法在跑。被推翻的是「走哪條軟體路徑」的細節，沒被推翻的是「有沒有原生混合位寬指令」這個更底層的事實。

尾聲:值得等的一個訊號

s8×s8 的 kernel 已經在官方 wheel 裡，W4A8 現在只卡在那一道 major >= 8 的 arch guard 上。如果哪天官方對 sm_75 開放融合反量化，內插推估這組 W4A8(12.54 GB)大約能落在 104 秒，比 INT8 快 12%——這個比例跟作者原本宣稱的 ~1.09× int8 speed 吻合。但這只是內插推估，不是實測。

連驗證這個推估都被版本牆擋住了:自建那版 kernel 的計畫在第一關就被擋下——官方 build 需要 CUDA 12.8 以上，而這台機器的 nvcc 是 12.4，連 .so 都編不出來。想拿到「等 guard 開放後真的會這麼快」的答案，得先等這台機器的 CUDA 升級，或者等官方直接把 Turing 納入支援範圍。這件事目前停在「值得等」，不是「已經驗證」。

---

這個系列的其他文章:

• #9 0xc0000409:當我的 AI 服務靜默暴斃，而 log 把死因吃掉了
• #10 同一張 2080 Ti，一台生圖慢 3.4 倍——兇手是啟動 log 裡一行 dtype fallback
• #11 什麼?2080 Ti 居然跑得動 MiniMax-H3，還生得出 1080p 有聲影片