| 線 | 現在的數字 | 狀態 |
|---|---|---|
| 生產 · Qwen3.8-27B IQ4_XS · TP4 100.78.29.29:8601 · fable-27b · 設計上限八個併發使用者 · context 131,072(改了 kernel 才拿到,模型宣告 262,144)· prefill 預算 131072 @ util 0.85 | 單人 68.1 · 八人合計 260(每人 32.6,同時產出時) 混合負載、錯開到達:合計 155 文件問答:八人各 21.6K → 55 秒 · 八人各 64K → 5 分 47 秒 八人同時互動的文件上限抓 20–25K | 上線中 · 工具閘門 8/8 一小時 soak 過:0 失敗 · 記憶體零增長 |
| 換掉 Codex 部署的 Q4_K_M | c4 合計 86.7 → 107.8(+24%) · c8 133 → 154(+16%) | 零行 code · 換檔而已 |
| IQ4_XS 自寫 kernel(batched-MMVQ + tile-MMQ) | 兩顆都在 4/4 worker arm | 上線中 · 一小時 soak 過 |
| 推測解碼(MTP) | 最佳 39.96,仍是基準的 0.58 倍;四併發 0.32 倍 | 已收 · 預設關 |
| Qwen3.6-27B(前一顆生產) | 單人 68.88 · c4 合計 108.8 | 留著,未刪 |
Qwen3.8 和 Qwen3.6 是同一個架構(qwen35、65 blocks、48 個帶 GDN、866 個張量),所以 3.6 上調出來的設定全部直接適用。唯一真正的槓桿是量化檔:Q4_K_M 讓自寫的批次 kernel 掛不上去,換成同一個 repo 的 IQ4_XS 就拿到併發 +24%。而 max-num-seqs、capture 集合、util 三個旋鈕四種組合完全沒有作用 —— c4 合計全都落在 86.5–86.9。
推測解碼這條線的上限已經算得出來:c 修到 llama.cpp 的 0.29、α 維持實測的 0.52 → S = 1.18 → 81 tok/s,比基準好 18%,而且只在單人。目前不建議投入。
本頁分成七個大類,每類裡最新的在最上面,每段標了產出日期。打「約」的日期是回填的,可能差一天。總表裡標 GDN 開 的列是 8/6 那次變更之前量的 —— 當時是對的,但不代表現在的生產。被推翻的結論一律保留並標明,不刪 —— 看得到我錯在哪比看不到有用。
deploy38.sh 已改,備份 .bak_66560。(256, 512, 1024),而分割是從 context 長度推導的,131,072 推出 2048 → Unsupported decode partition_size: 2048。在 LAUNCH_BY_PARTITION 加一個 case 2048 重建 flash_attn_v100_cuda...so 就過了。三組同 session 對照 c1 68.24 / 67.78 / 68.00、全員串流 260.14 / 260.02 / 259.35、c8 合計 154.67 / 153.73 / 153.64 —— 全在 ±1% 噪音內。我原本擔心的 occupancy 代價沒出現:那三個 shared array 是 D*4 + PARTITION_SIZE*12 bytes,2048 仍在 48KB 內。重啟後生產實測 c1 68.19、c8 合計 152.92、工具閘門 8/8、KV 2,747,830 token。deploy38.sh 備份 .bak_81920_pre131k、原版 kernel 留在 .so.bak_pre2048。shm_broadcast.py 的 TimeoutError → EngineDeadError。源頭 envs.py:514 的 VLLM_EXECUTE_MODEL_TIMEOUT_SECONDS 預設 300 秒,而 128K 的 prefill 要 500.2 秒。調到 2400 直接過。deploy38.sh.bak_65536_util060。同時補了兩個工程修正:就緒檢查改成等自己 log 的 GPU KV cache size(原本 poll /health,殘留伺服器會立刻回答),以及 teardown 按 PID 清場 —— 「假設機器是乾淨的」今天害我撞了三次。--max-num-batched-tokens 釘在 65536:64K 剛好一塊裝得下,96K 和 128K 要切兩塊。把預算提到 131072(util 0.60→0.85 才有記憶體)之後:64K 56.6s(1,146)、96K 64.6s(1,487)、128K 132.5s(978)—— 128K 快了 3.8 倍,而本來就沒被切的 64K 完全沒變。拿掉分塊之後 prefill 大致是線性的,三個長度都落在 1,000–1,500 tok/s。一般流量零代價:c1 68.46、c8 合計 154.26。KV 2,719,744,仍高過八人各 131K 所需的 1,048,576。--scheduling-policy priority,而且 /v1/responses 帶得了每請求的 priority(/v1/chat/completions 沒這個欄位)。同一台伺服器兩輪對照:四份 64K 進行中,互動請求 priority −1 花 296.7s、priority 0 花 258.5s —— 而兩輪的「文件完成時間減請求耗時」都是 20.0 秒,正好等於我探針的發送延遲。不管什麼優先權,請求都在文件做完的那一刻才完成,和 fcfs 一樣。短提示流量沒受影響(c8 153.35 對 fcfs 153.45),但同樣四份 64K 文件 316.7s 對 fcfs 的 124.1s(各量一次)。這個 fork 沒有任何排程手段能讓互動請求繞過進行中的 prefill —— 分塊粒度和優先權都試過,都死,分塊還是嚴格劣化。唯一有效的是讓 prefill 本身更快,而那條已經拿到 3.7 倍。kill -9 之前,驅動還沒把 VRAM 還回來就啟動 → Free memory on device cuda:3 (17.32/31.73 GiB) on startup is less than desired。差點被我讀成「priority 不支援」。② 我在腳本正執行 deploy38.sh 時去改它 —— bash 是邊讀邊執行的,byte offset 一移位還原就死在半路,留下四個孤兒 TP worker 各扣 14 GiB,生產停機 27 分鐘。執行中的 shell 腳本絕對不能改,要改就複製成新檔名。No available memory for the cache blocks)。activation buffer 隨預算等比長,util 0.85 才買得回空間。VLLM_FLASH_V100_DECODE_PARTITION_SIZE=1024 是無效的。deploy38.sh 一直釘著它,所以我先問「釘 1024 是不是就不用改 kernel」。釘 1024 的 131,072 跑出來跟 2048 逐項一樣(128K 文件 499.7s 對 500.2s、c1 67.97 對 68.46、c8 154.26 對 154.27),數字分不出是「pin 生效」還是「pin 被忽略」。唯一能分的是換 .so:原版 kernel 在同樣釘 1024 下照樣報 Unsupported decode partition_size: 2048。vLLM 從 context 推導的值壓過環境變數,那行設定完全沒作用。deploy38.sh 的就緒檢查是 poll /health,而殘留的舊伺服器會立刻回答。結果它 1 秒就判定就緒,再去 grep 剛啟動的 log,報 IQ4XS_BMMVQ_READY on 0 workers <-- NOT ARMED —— 假陰性。改成等它自己 log 裡的 GPU KV cache size,第一次上場就正確擋下 DEPLOY FAILED。--tensor-parallel-size,每卡要讀的 bytes 差 4 倍:TP4 68.16(14.671 ms/步,3.66 GB/卡)、TP2 49.66(20.137 ms,7.31 GB)、TP1 32.85(30.441 ms,14.62 GB)。擬合 t = F + bytes/BW,三組配對的斜率 1.410/1.439/1.498 ms/GB、截距 9.19/9.41/9.83 ms。只用 TP4+TP2 擬合去預測 TP1 是 32.2,實測 32.85 —— 差 2%,而 TP1 沒有參與擬合。mul_mat_vec_iq4xs_batched 438 次 20.6 ms(約 700 GB/s)、Q5_K 65 次 2.85 ms(628)、Q6_K 輸出投影 1 次 1.57 ms(669)。三族 GEMV 都是頻寬有效率的,沒東西可撿 —— Q5_K 不是慢路徑,只是張量大(attn_qkv ×48 是 [5120,10240],output.weight 是 12.7 億參數)。FillFunctor<int> 每步 241 次、2.1 ms」是假的。我拿 200-token 那次的總呼叫數除以 208 步,但沒驗證「均勻分佈在每一步」這個假設 —— 我只對 GDN kernel 驗過(48.2 對 48 個 block)。改跑 50 token(58 步)對照:每步比值應該是 0.279。GEMV 0.289、quantize 0.292、triton 0.256、GDN 0.282、Q5_K 0.292、RMSNorm 0.263 —— 六個全中。而 FillFunctor 是 50,095 → 160,比值 0.003。它是一次性爆發被我平均進每一步,不是每步成本。quantize_q8_1_b 432(完全 1:1)、Q5_K 64(每層一次)、GDN 48.0(正好 48 個 block)。時間:IQ4_XS GEMV ~20.6 ms、Q5_K 2.67、Q6_K 1.57、triton_poi_fused_0 2.09、quantize_q8_1_b 0.96、GDN 0.71、triton_ 0.57、RMSNorm 0.63、conv1d 等 0.35。可攻擊的總量從 7.5 ms 修正為約 5.5 ms,而最大的一項是 triton_poi_fused_0(78 次/步,約每層 1.2 次),不是我原本說的整數填充。INPUT_PROJECTION_OP 680 B(不同)68.43 / 153.04;INPUT_CORE_OP 667 B(不同)68.54 / 154.14;輸入+輸出一起開 670 B(不同)68.35 / 156.24。三個全部改變數值,單人速率全在雜訊內,最好的情況是只在八人合計上 +1.8%。和輸出側單獨開的 +1.2% 是同一個故事,也就是它們預設關的原因。QWEN_GDN_OUTPUT_PROJECTION_OP=1(把每層的 reshape/gated RMSNorm/v-reorder/out_proj 融成一個自訂 op):輸出變成 687 bytes / 140 詞,不同,合計只 +1.2%。拿數值變動換 1.2% 不是吞吐選項。GDN_Z_CONTIGUOUS + MIXED_QKV_CONTIGUOUS:輸出崩潰成 2 bytes,而速率一點沒動 —— 連取捨都不算。GDN_EMPTY_CORE_OUT 就是我前一晚手工改 torch.zeros→torch.empty 的那個開關,早就存在而且早就預設關。我重新發明了一個已被否決的東西 —— 改模型檔之前先盤點旗標。warps2 兩次 153.47 / 157.41。所有配置間的差距都躺在這個散佈裡,而且方向還在兩輪之間翻轉。checksum 全程相同,所以計算沒變,變的是量測條件。沒有那三個重複 arm,我會報「warps=2 快 19%」並附一個很像樣的 Volta 解釋。(B=8, T=1),結果每個配置都在 kernel 第 177 行掛掉 —— 包括 default。我當下讀成「BV 是假旋鈕」,實際意思是「我的輸入形狀錯了」。真實解碼路徑會把變長攤平:帶 cu_seqlens 時 q.shape[0] 必須是 1,還要 ssm_state_indices 和 use_qk_l2norm_in_kernel=True。對照組也失敗,就不是被測物失敗。qwen3_5.py:471 的 torch.zeros 換成 torch.empty 之後:生成直接斷掉,而且速度一點沒變。同提示、temperature 0:665 bytes 的完整答案變成 **First 12 —— 10 bytes、兩個詞就停。填充列沒有被遞迴核心寫滿,垃圾會傳出去,那個清零是承重的。而吞吐 c1 68.47 對 67.61、c8 合計 153.75 對 153.60 —— 噪音。empty 真的生效了 —— 既然生效卻沒有加速,那 core_attn_out 就不是那 2.09 ms。那個數字是把 triton_poi_fused_0 的呼叫數加總來的,而 inductor 的 kernel 是每個圖各自從 0 編號 —— 名字不唯一。78 次/步是好幾個不同清零 kernel 的總和,我拿掉的只是其中便宜的一個。2.09 ms 是真的,但來源未歸屬,而我唯一的線索(kernel 名字)已證實不具鑑別力。quantize_q8_1_b:它 blockDim.y = 1、每個 block 一列,amax 是沿列做 width-32 shuffle —— 列與列獨立,填充列的 inf 污染不到別列的尺度。那個檢查是對的,但失敗來自下游而不是量化。查對前提不保證查到的是對的前提。triton_poi_fused_0 的 Triton IR 是純 tt.store 寫 0.0 到 f16 緩衝區、沒有任何 load;而 inductor 自己的呼叫端寫著 rand_strided((8, 12, 128), dtype=torch.float16) = 12,288 元素 =(捕捉尺寸 8, 48 v_head ÷ TP4, head_dim 128),正是 qwen3_5.py:471 的 core_attn_out。注意 inductor 的 kernel 是每個圖各自從 0 編號,所以那個名字不唯一,78 次/步是多個清零 kernel 的總和。不是安全的一字改成 torch.empty:緩衝區按填充後的捕捉尺寸 8 配置,實際 token 可能只有 1,填充列不會被遞迴核心寫到。if mx > 1e-3 —— nan > 1e-3 在 Python 是 False,一路放行,max(0.0, nan) 也保留 0.0 把最壞值報成 0。我整天在講「守門要拒絕出數字」,然後自己寫了一個對 NaN 無感的守門。quantize_q8_1_b 每步 440 次,和 438 次 GEMV 是 1:1 —— 每個矩陣向量積前都重新量化一次激活,而 llama.cpp 那邊早就融進 mmvq。單檔、自足、驗證條件明確,已附線上真正 runtime 編譯的 iq4_mmvq_batched.cu 與正確的參考 header 派給委員會。整數填充沒派,而那個決定後來被證明對得莫名其妙 —— 我當時的理由是「來源未定位」,實際上它根本不是每步成本。不派自己說不清的東西這條規矩救了我一次,但救的方式和我以為的不同。1→2→4→8→8→4→2→1 連續跑、arm 之間不冷卻:每人 68.13 / 56.36 / 47.87 / 34.01 上去,34.03 / 47.91 / 56.40 / 68.64 回來 —— 差異全在 0.1% 內。合計 59.4 / 83.8 / 107.7 / 152.8:八人時每人只剩一半速度,但總產出是單人的 2.24 倍。使用者數量本身不留後遺症,人走了速度立刻滿血。BMMVQ 4/4 真的掛上。結果 c8 全員串流 260.15 → 261.90(+0.7%)、c4 合計 107.40 → 107.28、c1 持平 —— 全在雜訊裡(同組態不同輪已量到 8% 變異)。「權重在迴圈裡被重讀八次」描述的是源碼的形狀,不是時間花掉的地方 —— L1 本來就吸收掉了。vec_dot 三參數變四參數、__dp4a 變 ggml_cuda_dp4a,編不過;④ 提示裡的雙引號被 PowerShell 吃掉,argparse 拒收。委員會的診斷從第一輪就沒變過。BMMVQ x/4 —— kernel 沒掛上時伺服器照樣起得來、照樣有數字,那個數字看起來會像「改動沒效果」而不是「改動沒生效」。reasoning,非串流叫 reasoning_content —— 同一個東西兩個名字。delta.content 一直存在但是空字串,所以只試那兩個非串流名稱的檢查會安靜地數到 0 個 token。把第一 KB 用 repr() 印出來花十秒,而我先前猜了兩次格式。--image-min-tokens 1024 —— 不加的話 96×96 的圖只換到 81 個 prompt token,加了之後是 1,103。bt 65536 之後仍然成立。合計吞吐 18.7 tok/s,對短提示八人的 155 —— 長 context 併發是八倍代價,時間幾乎全在 prefill。loaded multimodal model、生成 120 個 token。我讀 content 讀到空,內容在 reasoning_content。本 session 第三次讀錯欄位,而我還為這件事寫過記憶。max-num-batched-tokens 65536 換來的。Q4_K_M,我先前為 IQ4_XS 寫的兩顆批次 kernel 掛不上去;同一個 HF repo 就有 IQ4_XS 版(15.3 GB)。換檔後兩顆 kernel 在 4/4 worker arm:c4 合計 86.74 → 107.75、c8 133.14 → 153.86,單人 69.42 → 68.15(−1.8%,量化較小的代價)。工具閘門 8/8、真實文本連貫。回滾:/home/mavis/codex_qwen38_launch.saved,Q4_K_M 檔原封不動。max-num-seqs(16 vs 8)、capture 集合、util(0.60 vs 0.50)四種組合,c4 合計全部落在 86.5–86.9,c8 全部 133.0–133.6。那 25% 從頭到尾都是量化檔,不是排程設定。--max-num-batched-tokens 16384 → 65536:21.6K 冷 prefill 19.8s → 15.4s(−22%),decode 不動。和我先前在 106K 上量到的 −26.3% 同一個量級,已套進生產。torch._dynamo.exc.Unsupported: Attempted to call function marked as skipped — module: time, qualname: perf_counter。整條 GDN 路徑在 torch.compile 的 fullgraph 捕捉區內,任何 Python 層計時器都會讓捕捉失敗、引擎起不來。加上 vLLM 自己的 /start_profile 在 TP4 下寫不出東西(worker 是獨立 process,只武裝 rank 0)—— 這台機器上要看進推測步裡面只剩 nsys --output=…%p 一條路,那不是備案,是唯一解。已還原並重跑驗過 39.66 對插樁前的 39.84,NOSYNC 八處仍在 —— 能 parse 不等於伺服器起得來,一定要重跑一臂。/mnt/data/backups/windows-v100/models_muse、/mnt/data/llama.cpp-new),要用隨時能起。spec_sequence_masks is None 那個守衛從「候選之一」變成主嫌 —— 它正好是一條純引擎側、與模型架構無關的路徑選擇。(限制:llama.cpp 兩顆都是單卡、vLLM 是 TP4,跨了引擎也跨了拓樸;但 llama.cpp 內部那組同引擎同卡數的比較是乾淨的,足以殺掉遞迴假設。)n_max=16 exceeds the trained block size 16 -- clamping to 15。預設最好。reasoning_content 而不是 content(打到 finish_reason: length 時還在推理段),所以我的文本檢查印了 uniq_word_ratio 0.00 DEGENERATE —— 吞吐數字是有效的,壞的是我讀錯欄位。實際內容 1585 字元、連貫英文。--spec-type 預設是 none,我只給了 --model-draft —— 兩個草稿臂會在推測關閉下跑成基準。第二輪:對 instruct 模型打了 /completion,沒有東西套 chat template,生成到第 2 個 token 就停(predicted_n=2,內容是 " -"),而那三組 tok/s 是在兩個 token 上量出來的。我的退化檢查放行了 —— 一個字的輸出 uniq_word_ratio 是 1.00,滿分。那個檢查是為了抓重複寫的,而重複是相反的失敗。已改成:先斷言實際生成的 token 數,不到 100 就標 INVALID、不准印速率。block_size=16, mask_token_id=201818, n_extract=5 —— 它一次提一整段、靠 mask token 填空。我拿量 MTP 的條件(n-max 1)去套等於把它掐死在一格,所以 accepted 0 / generated 1。重跑改成給它自己的區塊大小。另外這顆官方 GGUF 有個真警告:special_eot_id is not in special_eog_ids,tokenizer 設定可能有問題。FLASH_GRAPH+SPEC_CORE_OP:c1 39.84(掃描時 39.96,可複現)、接受率 52.3%、工具閘門 8/8 PASS、真實文本 uniq_word_ratio 0.79 沒有退化。但 c4 合計吞吐只有 34.80,無推測基準是 108.80 —— 0.32 倍,TTFT 5.03s。批次已經把成本攤平的地方,推測解碼只會添亂。就算做到最好,這條線也只在單人有意義,而單人現在也才 0.58 倍。生產維持不開推測的 68.88。8e7f22b、gcc 12.4.0、CUDA arch 70、ninja -j2),編譯日誌裡有 common_speculative_impl_draft_dflash —— DFlash 推測支援確實在裡面。/opt/llama.cpp 覆核過完好,對照組數字仍可比。FLASH_V100_0DOT3_COMPILE_GRAPH 單獨對解碼沒用(34.55,c 2.020),但疊在 SPEC_CORE_OP 上把它造成的傷害修掉 —— c1 39.96、接受率回到 52.3%(全場最高,與不開推測同水準)、TTFT 從 4.42s 回到 0.68s。委員會把這個旗標排第一是對的,但理由錯了:它救的是 graph 與 prefill,不是解碼主力。content-length 逐一相等,731 個張量目錄讀到底)。它是我這條線缺的對照組:52 個 block 全是注意力、沒有 SSM(對照 Qwen3.6 的 65 block 裡 48 個帶 GDN 遞迴),KV head 32:2、滑動窗 2048、context 131072,而且官方另出一顆 DFlash 草稿模型。一顆沒有遞迴狀態又自帶草稿模型的 30B,正好能分辨 c=2.01 是架構特有還是引擎通病。官方沒有純 Q4,只有 kquant-17gb(16.8 GB,為 24 GB 消費卡而生)和 kquant-dynamic(19.7 GB)—— 取後者,我們是 32 GB×4,沒有理由為遷就別人的卡吃品質損失。/opt/llama.cpp 是 7/6 建的,muse-glimmer 支援是 8/10 才合併(#26841,b10353 起)—— 認不得這個 arch。要跑得建新的,而且不動 /opt/llama.cpp:今晚 llama.cpp 那組 α 0.722 / c 0.298 的數字全出自它,重建會毀掉可比性。新的建到 /mnt/data/llama.cpp-new。SPEC_CORE_OP c1 39.62(c 2.043→1.595,+15.5%)、003_SPEC_CORE_OP 38.49(c 1.593)。兩條路徑的 c 幾乎同值不是巧合,應該接到同一個改動;SPEC_CORE_OP 吞吐較高只是因為它的接受率沒被拖低(49.3% vs 44.9%)。其餘五個全在 c≈2.0,是雜訊。FULL_FORWARD 排在 SPEC_CORE_OP 前面(理由是名稱暗示同時修候選 ①②)—— 實測反過來。FULL_FORWARD 不但沒用(34.03,c 2.036),TTFT 還從 1.46s 惡化到 2.86s。排序要靠量,不靠命名。qwen_gdn_linear_attn.py 的 4365 與 5706 兩處都有 and attn_metadata.spec_sequence_masks is None,所以用 packed_decode 那條路在推測步永遠不會被選中,一律落到通用的 _forward_core,而那裡把 token 拆成 spec / non-spec 兩組分別跑。conv 那半有守(num_prefills>0 / elif num_decodes>0 / else None),單人時沒有空跑;成本在後面的 gating 與遞迴段。同一個守衛還帶出委員會不知道的交互作用:FUSED_SIGMOID_MIXED_QKV=1 會讓 mixed_qkv_decode_requested 為真,也會關掉這條快速路徑。start_profile 只武裝 rank 0 的 context,跨不過 process 邊界。這解釋了我三次失敗,替代方案是 nsys --output=…%p 每個 PID 各自出檔。委員會的價值在候選假設與這條路,不在優先順序。(Codex 這次 thread start failed 無實質產出,融合時整份排除,實際是 GLM-5.2 與 MiniMax-M3.0 兩份。)VLLM_TORCH_PROFILER_DIR;第二次設了但讀取端掃的是自己的預設目錄;第三次目錄對上了,TP4 下 /start_profile 之後那個目錄仍然是空的,fork 自己的 phase report 也沒進 log。三次失敗長同一個樣子 = 問題在假設不在實作(RULES 5),所以不再想辦法「看進步驟裡面」,改直接量哪一種實作最便宜 —— 這個 fork 帶了八個切換推測路徑的旗標,一臂一個,c 用同一 session 自己的無推測臂算,不從記憶拿。p-min 降到 0、完全不篩,llama.cpp 在 γ=1 仍有 72.2%,草稿/步 = 115/(200−83) = 0.98,等於 γ,篩子根本沒動。vLLM 同樣 γ=1 是 50.4%。同一顆 GGUF、同一顆 MTP 頭、同一張卡 —— vLLM 真的把草稿頭餵得比較差,那是個可定位的 bug 不是架構宿命。(順帶更正我自己的量測腳本:它印的「drafted per generated token」分母錯了,該用步數 —— 接受的草稿會讓一步產兩個 token,把分母灌大,害兩個沒篩的臂被標成 FILTER IS ACTIVE。)target_hidden_states = hidden_states[...],而 qwen3_next.py:953 是 hidden_states, _ = self.norm(...) 之後才 return;MTP 頭接著又對它做一次 pre_fc_norm_hidden。llama.cpp 的對應圖在 src/models/qwen35.cpp(源碼在機器上),下一輪對照。這條排在 c 後面。nonzero() 拿掉,MTP 單人 +6.8%(32.42 → 34.64),四人 +2.5%,接受率不動(51.5%),工具閘門 ALL_PASS。那段用布林遮罩挑推測列,每個 x[mask] 都是 nonzero() —— 它得把命中數讀回 host 才知道輸出多大,是一次完整的裝置同步,而且一次 forward 好幾個。批次同質時遮罩全 True,那個挑選就只是複製一份。判斷用的計數呼叫端已經在 CPU 上算好,問它不花錢。c 從 2.22 降到 2.01。預設關,VLLM_SM70_GDN_NOSYNC=1 才開。-ctk turbo2,但 /opt/llama.cpp 是上游 b0.15.3 不是 atomic-llama fork,turbo2 是 fork 專有型別。旗標檢查有做、KV 型別沒檢查。改 f16 重跑中。p-min 0.5 讓它信心不足時根本不提草稿,那些注定被拒的位置沒進 draft_n,百分比是在篩過的子集上量的;② vLLM 真的把草稿頭餵錯了。把 llama.cpp 的 p-min 降到 0 重量,掉到 ~50% 就是①(那接受率從頭到尾都不是差別,差別是 c),留在 ~72% 就是②(那是個值得修的 bug)。跑了。OFFICIAL 那顆)metadata 全是 qwen35、65 blocks、48 個帶 SSM、ssm.state_size 128。看檔名不算,要看 metadata。spec_state_slot_selectors 這類機制,但沒找到 ring buffer —— 是不是半套還沒確認。memcpy32_post 1879 次、scatter_gather 268、indexSelect 268、CatArrayBatchedCopy 268 —— 全是狀態簿記,不是算術。論文的 1.8x / 3.4x 全部量在純 transformer 上。mul_mat_vec_q 把 token 位置放在 blockIdx.y —— 每個位置各開一組 block、各讀一次 45 MiB 的權重。profiler 的證據無可辯駁:同一顆 kernel 在有無推測下 launch 次數幾乎相同(2489 對 2468),時間卻差 60% —— 同樣次數的呼叫、每次做兩倍的搬運。改成一個 block 帶 NCOLS 個位置、權重靠 L1 重用後,每位置成本從 0.086→0.074(幾乎不降)變成 0.084→0.032。M=1 持平所以不傷單人;九種權重形狀 × 每個 M 的誤差與原 kernel 一致到印出的精度。(1−α^(γ+1)) / ((1−α)(γc+1)),Corollary 3.9 的損益平衡是 α > c。我們實測 α = 0.478、c = 2.04 —— 差 4.3 倍。他們論文裡 c 的範圍是 0.007–0.073,我們是他們畫過最壞情況的 28 倍。而且 Table 2 自己就在證明 α 不是槓桿:接受率最高的 T5-large(0.82)加速最差(1.7x),因為它的 c 最大。tool_choice=required 問。兩顆模型正好相反:生產 9B 在 required 能用、auto 回空;27B 反過來。兩顆都不壞,各自在一種模式下吐出完全正確的呼叫。失敗那側都是同一個樣子:tool_calls 空的,而 reasoning 欄位裡模型正在說「我應該呼叫 lookup_inventory」—— reasoning parser 把開頭那段吃掉,tool parser 就什麼都不剩。閘門已改成兩種都問、任一成立即通過並指名哪一種;生產與 27B 現在都 ALL_PASS。我為這個假訊號往 kernel 和 capture 各查了一輪,那些成因不存在。t5a_no_think_plain 在生產通過,但模型開頭就是 Thinking Process:\n\n1. **Analyze the Request:** —— 那個檢查只找字面的 <think> 標籤,這種寫法直接走過去。沒有動它:把生產監控轉紅是關於既有行為的決定,不是量測 bug。席次 ×(1+draft) 個 token,而 vLLM 只保留能被那個倍數整除的尺寸。同一組 [1,2,4,8,16,32]:不開 MTP 捕捉到 6 張 decode 圖,draft=1 只剩 5 張(1 不能被 2 整除)。真正可用的席次上限因此是 16,c24 和 c32 全走 eager —— 兩者步時間 233.5 / 234.1 ms 幾乎同值,就是這個原因。draft=2 要 3 的倍數,那組裡一個都沒有,所以它捕捉到 0 張。[1,2,4,8,16,32,48,64] 後捕捉數 7/7,c24 每人 4.28 → 8.17(合計 97 → 181)、c32 4.27 → 7.60(合計 128 → 225)。t3_tool_call 也是三臂全失敗含舊集合,同樣與 capture 無關。所以 capture 對齊是乾淨的淨勝:c24 3.76 → 8.16(2.17x)、c32 3.66 → 7.59(2.07x),c1/c8/c16 在雜訊內不動。ggml_mul_mat_a8 對 IQ4_XS 完全沒有運算。profiler 顯示它只發射 quantize_q8_1,之後沒有任何 matmul,輸出緩衝區沒被寫入。它先前之所以通過數值比對,是 torch 的 caching allocator 把剛算完參考值的那塊記憶體回收再配給它 —— 讀到的是上一次的殘值。因此把 IQ4_XS 加進 MMQ_QUANT_TYPES 來修批次崩塌,會讓模型安靜地吐垃圾而且不報錯。blk.0.ffn_up,17408×5120):反量化整顆權重 0.663 ms + cuBLAS 0.441 ms = 1.106 ms,而且不隨 M 變;同尺寸的 K-quant 走真 MMQ 在 M=16 只要 0.355 ms。mmvq 則是線性的,M=1 0.086 ms、M=64 4.169 ms,與反量化路徑的交叉點在 M≈15 —— 而 gguf.py 的門檻寫死在 8,偏保守。QR4_XS/QI4_XS —— 那是 8 和 8,描述的是 MMVQ 怎麼把超區塊分給執行緒,不是 MMQ 的 tile 幾何。tiling 要的是 QK=256 / QR=2 / QI=32,現在綁了 static_assert 讓它不能再漂走。kvalues_iq4nl 查表從 vec-dot 搬到載入時做一次(x tile 會被重用 mmq_x 次,等於同一組查表重做 64 遍)又多 1.5 倍:對比 dispatcher 現行選擇,M=8 2.48x、M=16 4.95x、M=32 2.57x、M=64 1.95x。os.getenv 每次矩陣乘都跑,提到模組層級後只移動 0.3% —— 猜錯了。而且用算的就排除得掉:條件短路後開關只差兩次整數比較(約 25 µs),實際差距是每步 0.6 ms。成因在別處,還沒找到。操作結論:只有單人就別開,有併發就開。mul_mat_q40_dec_y、mul_mat_q40_x64,註記 vs fork 1.26–2.35x),runtime 編譯、env 控制、shape gate 全都在,但都卡在 qweight_type == 2,也就是 Q4_0 專用。缺的只有 load_tiles_iq4_xs 和 vec_dot_iq4_xs_q8_1_mul_mat 兩個函式 —— IQ4_XS 的 32 元素子區塊剛好等於一整個 Q4_0 區塊(都是 4 個 int、低 nibble 都是元素 0–15),所以 tile 幾何可以原封不動沿用。已寫好,待驗。llama-bench 的 -ctk/-ctv 是位置配對,先前那批其實只量到 V —— 重跑後:V 量化花 8% 生成速度,K 量化再花約 30%(K 在 QKT 內圈)。layer 38.9(等於單卡)、row 22.0(慢 43%)、tensor NCCL 崩潰。llama.cpp 在這台沒有可用的 TP4,所以它只能當參照組。llama-bench 中止是因為 tbq3 必須要有 FA。blk.64 gate 到 Q8_0 重做 requant 這條路不用走了。force,另有 auto 會自行武裝;8/8 與 8/9 的伺服器 log 都印了 guard armed (force=False, auto=True)。-sm tensor 崩潰根因是 NCCL 版本:系統的 2.18.3+cuda12.0 對上 CUDA 13 driver,查 entry point 失敗被誤譯成 「CUDA driver is a stub library」。LD_PRELOAD vLLM venv 的 2.27.5+cuda12.9 就跑起來了。-ts 是斜線分隔,我用了逗號 —— 在 llama-bench 那是「掃多組設定」,結果四個臂全部跑在 GPU 0。破綻是 pp512 四種配置差 0.15%。重跑中,這次會驗每張卡有沒有吃到權重。ggml_mul_mat_a8 對它是空轉返回(M=2 到 512 耗時全是 0.03 ms,換算 3000 TFLOPS,V100 峰值 125)。--enable-auto-tool-choice 和 --tool-call-parser,生產的啟動腳本一直都有。每一列都是量過的,沒有推算值。併發欄是同時的請求數,tok/s 是每人而非合計。「參照」列是別的引擎或別台機器,只用來對照,不是我們的產出。
| 模型 · 量化 | 拓樸 | context | 併發 | tok/s | 時期 | 註記 |
|---|---|---|---|---|---|---|
| 9B · Q4_K_M | TP4 | ~200 tok | 1 | 152.8 | 現行 | 現行生產 |
| 9B · Q4_K_M | TP4 | ~200 tok | 8 | 76.8 | 現行 | 每人 |
| 27B · IQ4_XS | TP4 | ~200 tok | 1 | 68.2 | 現行 | 64K 配置 |
| 27B · IQ4_XS | TP4 | ~200 tok | 8 | 20.0 | 現行 | 每人 · 合計 160 |
| 27B · IQ4_XS | TP4 | ~8K tok | 8 | 17.7 | 現行 | 每人 |
| 27B · IQ4_XS | TP4 | ~28K tok | 8 | 14.1 | 現行 | 每人 |
| 27B · IQ4_XS | TP4 | 105,856 tok | 1 | 37.7 | 現行 | 128K 配置 |
| 27B · IQ4_XS | TP4 | 105,856 tok | 8 | 7.0 | 現行 | 每人 · 端到端只有 3.4 |
| 27B · Q4_0 | TP4 | ~200 tok | 1 | 75.1 | 現行 | requant · 舊版權重 |
| 27B · Q4_0 | TP4 | ~200 tok | 8 | 34.5 | 現行 | 每人 · 合計 276 · 唯一過 30 的 |
| 27B · IQ4_XS | TP4 | ~200 tok | 1 | 30.2 | 現行 | 開 MTP spec-1 · 淨損 |
| 27B · Q4_K_S | TP4 | ~200 tok | 24 | 33.6 | 現行 | 每人 · 合計 241 · util 0.75 · 不開 MTP |
| 27B · Q4_K_S | TP4 | ~200 tok | 32 | 33.7 | 現行 | 每人 · 合計 256 · c8 到 c32 是平的 |
| 27B · IQ4_XS + 批次 MMVQ | TP4 | 混合 · 100–400 out | 1 | 69.1 | 現行 | <b>不開 MTP</b> · 異質 p50 · 這是目前最佳單人 |
| 27B · IQ4_XS + 批次 MMVQ | TP4 | 混合 · 100–400 out | 4 | 48.5 | 現行 | 不開 MTP · 異質 p50 · 合計 108 · 改前 36.74 |
| 27B · IQ4_XS + 批次 MMVQ | TP4 | 混合 · 100–400 out | 8 | 32.3 | 現行 | 不開 MTP · 異質 p50 · <b>合計 143</b> · 改前 20.58 |
| 27B · IQ4_XS | TP4 | 混合 · 100–400 out | 8 | 20.6 | 現行 | 不開 MTP · 未改 kernel · 這是被修掉的那條 |
| 27B · IQ4_XS + MTP1 + 新 kernel | TP4 | ~200 tok | 8 | 15.1 | 現行 | 每人 · 合計 103 · capture 對齊 ＋ <code>VLLM_GGUF_IQ4XS_MMQ=64</code> |
| 27B · IQ4_XS + MTP1 + 新 kernel | TP4 | ~200 tok | 16 | 12.5 | 現行 | 每人 · 合計 188 |
| 27B · IQ4_XS + MTP1 + 新 kernel | TP4 | ~200 tok | 32 | 9.0 | 現行 | 每人 · <b>合計 266</b> · 同構上界 |
| 27B · IQ4_XS + MTP1 + 新 kernel | TP4 | 混合 · 100–400 out | 8 | 15.1 | 現行 | <b>異質</b> p50 · p95 18.08 · 合計 61.6 · TTFT p50 0.51s |
| 27B · IQ4_XS + MTP1 + 新 kernel | TP4 | 混合 · 100–400 out | 32 | 9.2 | 現行 | <b>異質</b> p50 · p95 10.05 · <b>合計 136.7</b> · TTFT p50 1.80s · 32/32 完成 |
| 27B · IQ4_XS + MTP1 | TP4 | ~200 tok | 24 | 8.2 | 現行 | capture 對齊 <code>(1+draft)×席次</code> · 合計 181 · 未開新 kernel |
| 27B · IQ4_XS + MTP1 | TP4 | ~200 tok | 32 | 7.6 | 現行 | 同上 · 合計 225 |
| 27B · IQ4_XS + MTP1 | TP4 | ~200 tok | 24 | 3.8 | 現行 | capture 未對齊 · 掉出 CUDA graph 走 eager · 這是被修掉的缺陷 |
| 27B · Q4_K_M | TP4 | 105,856 tok | 1 | 35.1 | GDN 開 | 8/7 舊量測 |
| 27B · Q4_K_M | TP4 | 105,856 tok | 8 | 8.1 | GDN 開 | 8/7 舊量測 |
| 9B · Q4_K_M | TP4 | ~200 tok | 1 | 123.1 | GDN 開 | 8/6 之前 |
| 9B · Q4_K_M | TP2 | ~200 tok | 1 | 96.2 | GDN 開 | |
| 9B · Q4_K_M | 單卡 | ~200 tok | 1 | 69.7 | GDN 開 | |
| 9B · Q4_K_M | TP4 | ~28K tok | 1 | 106.2 | GDN 開 | |
| 9B · Q4_K_M | TP4 | ~28K tok | 8 | 45.9 | GDN 開 | 每人 |
| 27B · Q4_0 | TP4 | ~200 tok | 1 | 60.4 | GDN 開 | |
| 27B · Q4_0 | TP4 | ~200 tok | 8 | 30.9 | GDN 開 | 每人 |
屋頂線:27B IQ4_XS 是 15.13 GB,單卡 900 GB/s 給出 59.5 tok/s 的單 token 上限,四卡 238。llama.cpp 單卡不開推測拿到 65%,開 MTP 拿到 94% —— 單卡幾乎榨乾了。我們的 vLLM 四卡在 29%,差距不在頻寬。
這張表刻意和上面分開。先前它們混在同一張表裡、只用一個小標籤標「參照」,結果 3090 那列的 85 被讀成我們的 27B 成績。我們自己的 27B 在 vLLM 上的最高紀錄是 75.1(Q4_0 · TP4 · KV turboquant_2bit_nc · ~200 tok · c1),本機所有 log 裡沒有更高的。
| 模型 · 量化 | 引擎 · 拓樸 | context | 併發 | tok/s | 註記 |
|---|---|---|---|---|---|
| 27B · IQ4_XS | llama.cpp · 單卡 | ~200 tok | 1 | 39.24 | 不開推測 |
| 27B · IQ4_XS | llama.cpp · 四卡 layer | ~200 tok | 1 | 38.46 | 等於單卡,管線切分買容量不買速度 |
| 27B · IQ4_XS | llama.cpp · 四卡 row | ~200 tok | 1 | 22.12 | pp512 只有 88.99,比單卡慢十倍,病態 |
| 27B · IQ4_XS | llama.cpp · 四卡 tensor | ~200 tok | 1 | 33.61 | pp512 2333 對單卡 855,提示處理 2.7 倍 |
| 27B · IQ4_XS | llama.cpp · 單卡 + MTP n=3 | ~200 tok | 1 | 55.90 | MTP 在 llama.cpp 會賺,在 vLLM 不會 |
| 27B · IQ4_XS | llama.cpp · 單卡 | 24,704 tok | 1 | 33.40 | f16 KV |
| 27B · IQ4_XS | llama.cpp · 單卡 | 24,704 tok | 1 | 25.50 | tbq3 KV,反而更慢 |
| 27B · AutoRound INT4 | 外部 · vLLM · 3090 單卡 | 125K tok | 1 | 85.00 | 別人的機器,不是我們的成績 |
llama.cpp 四卡的三種模式都驗過每張卡真的吃到權重(tensor 是 4519 / 4313 / 4313 / 4313 MiB),不是只看旗標有沒有被接受 —— 第一次量的時候 -ts 用了逗號,llama-bench 把它讀成「掃多組設定」,四個臂其實都跑在 GPU 0,整批數字作廢重跑。
27B Fable · TP4 · 4×V100-SXM2-32GB NV2 · KV turboquant_2bit_nc(K2V2)· max-model-len 65536 · util 0.75 · kv-splits 96 · block-size 32 · batched-tokens 16384 · max-num-seqs 8 · prefix caching 開 · MTP 關 · GDN full-forward 關 · 模型 IQ4_XS 14.09 GiB · 每卡 23.8 GB / 32
| 提示 | 併發 | 冷 TTFT | 暖 TTFT | 每人 tok/s | 端到端 tok/s | prefix 命中 |
|---|---|---|---|---|---|---|
| ~200 | c1 | 0.27s | 0.27s | 68.2 | 43.8 | 0% |
| ~200 | c4 | 0.27s | 0.82s | 32.8 | 66.9 | 0% |
| ~200 | c8 | 0.27s | 0.69s | 20.0 | 101.1 | 0% |
| ~8K | c1 | 4.19s | 0.96s | 63.9 | 25.3 | 85% |
| ~8K | c4 | 4.19s | 4.10s | 29.6 | 57.7 | 85% |
| ~8K | c8 | 4.19s | 6.62s | 17.7 | 60.1 | 85% |
| ~28K | c1 | 25.33s | 1.79s | 55.3 | 40.2 | 97% |
| ~28K | c4 | 25.33s | 6.68s | 25.1 | 47.8 | 97% |
| ~28K | c8 | 25.33s | 12.90s | 14.1 | 58.9 | 97% |
27B Fable · TP4 · KV turboquant_2bit_nc · max-model-len 131072 · util 0.60 · kv-splits 160 · block-size 64 · batched-tokens 65536 · max-num-seqs 8 · prefix caching 開 · MTP 關 · 模型 IQ4_XS · 每卡 13.5 GB / 32
| 提示 | 併發 | 冷 TTFT | 暖 TTFT | 每人 tok/s | 端到端 tok/s |
|---|---|---|---|---|---|
| ~200 | c1 | 0.26s | 0.26s | 67.9 | 45.2 |
| ~200 | c8 | 0.26s | 0.69s | 20.0 | 98.1 |
| ~105K | c1 | 209.08s | 16.35s | 37.7 | 2.1 |
| ~105K | c8 | 209.08s | 130.03s | 7.0 | 3.4 |
兩欄不一樣:「每人 tok/s」是 token 開始吐之後的速率,「端到端」把 TTFT 也算進去。105K 那兩列端到端只有 2.1 / 3.4,因為 256 個 token 前面卡著 209 秒的冷 prefill。那格的體感由 TTFT 決定。對照 8/7 的 Q4_K_M:冷 390.9s、暖 c1 31.1s —— 新模型的 prefill 快了 47%。
27B Fable · TP4 · 4×V100-SXM2-32GB NV2 · KV turboquant_2bit_nc(K2V2)· max-model-len 65536 · util 0.75 · kv-splits 96 · block-size 32 · batched-tokens 16384 · max-num-seqs 8 · prefix caching 開 · MTP 關 · GDN full-forward 關 · 提示 ~200 tok · 每個臂在同一個伺服器內重測 c1 當錨
| 併發 | IQ4_XS 每人 | IQ4_XS 步時間 | Q4_0 每人 | Q4_0 步時間 |
|---|---|---|---|---|
| c1 | 68.2 | 14.7 ms | 75.1 | 13.3 ms |
| c2 | 47.6 | 21.0 ms | 51.6 | 19.3 ms |
| c4 | 32.6 | 30.7 ms | 34.6 | 28.4 ms |
| c8 | 20.0 | 50.0 ms | 34.5 | 28.5 ms |
IQ4_XS 的步時間精確符合 11.3 ms + 4.83 ms × 序列數 —— c4 代入 30.7(實測 30.7)、c8 代入 50.0(實測 50.0)。而 Q4_0 從 c4 到 c8 多了四個序列,步時間只多 0.1 ms,已經飽和。
那條直線是 i-quant 專屬的病,不是這台機器的性質:IQ4_XS 用非均勻碼表,每多一列都要重新查表拆包;Q4_0 是均勻 4-bit 分塊,批次幾乎免費。
不是頻寬:4.83 ms 折算成流量是每序列每步 17 GB,比整顆 14.09 GiB 的模型還大。
GDN_DECODE_FLASHQLA、MIXED_QKV_CONTIGUOUS+Z_CONTIGUOUS、block-size 64、capture [1,2,4,8,16]+max-num-seqs 16 —— 四個臂的 c1/c2/c4/c8 一位小數都不差。不是效果小,是零效果。
我對這件事給過三個機制,三個都被自己的量測推翻:① GDN 遞迴狀態頻寬 —— 算出來 0.056 ms,差三個數量級;② GDN 對序列迴圈 —— profiler 顯示 GDN kernel 在 c8 的啟動次數是 c1 的 1.02x,沒有迴圈;③ ticket18 中等批次 kernel —— 開與關量出 34.48 對 35.13,關掉還略好。
真正有效的變動只有一個:換量化型別。profiler 指出主力 kernel 是 mul_mat_vec_q,那是矩陣×向量、為 batch 1 設計的,而 c8 仍然走它。
27B Fable · TP4 · 4×V100-SXM2-32GB NV2 · KV turboquant_2bit_nc(K2V2)· max-model-len 65536 · util 0.75 · kv-splits 96 · block-size 32 · batched-tokens 16384 · max-num-seqs 8 · prefix caching 開 · MTP 關 · GDN full-forward 關 · 模型 IQ4_XS · --spec-method mtp · c1
| 推測 | 散文 | 工具呼叫形狀 |
|---|---|---|
| 關 | 68.1 | 68.4 |
| spec-tokens 1 | 29.7 | 30.2 |
| spec-tokens 2 | 21.1 | 25.7 |
內容類型沒有救到它 —— 工具呼叫那格和散文一樣爛。llama.cpp 上 85–99% 的接受率在 vLLM 這裡不轉換成速度,因為瓶頸是推測步本身的成本。另外 temp 0 之下三個臂的生成長度不同(272 / 271 / 201),推測路徑會改變輸出,那比速度更嚴重。
後來的六臂掃描顯示單人接受率從 2-bit 到 f16 完全不動(51.5 / 50.0 / 52.3 / 51.5 / 51.5%)。下面的數字是真的量到的,但那是單一提示的單次量測,我把它寫成了因果。保留原文備查。
27B Fable · TP4 · 4×V100-SXM2-32GB NV2 · KV turboquant_2bit_nc(K2V2)· max-model-len 65536 · util 0.75 · kv-splits 96 · block-size 32 · batched-tokens 16384 · max-num-seqs 8 · prefix caching 開 · MTP 關 · GDN full-forward 關 · 模型 IQ4_XS · --spec-method mtp --spec-tokens 3(對照組用的深度)· max-model-len 32768 · c1 · /no_think
| KV 格式 | 散文 | 程式碼 | JSON |
|---|---|---|---|
turboquant_2bit_nc(我們一直在用) | 14.7% | 34.9% | 47.0% |
turboquant_3bit_nc(對照組用的) | 31.1% | 58.4% | 87.1% |
| 逐位置(JSON) | 我們 2-bit | 我們 3-bit | 外部對照組 |
|---|---|---|---|
| position 0 | 80.6% | 97.6% | 97% |
| position 1 | 44.4% | 90.4% | 95% |
| position 2 | 16.1% | 73.5% | 91% |
MTP 的 draft 頭從 hidden state 預測下一個 token,而 hidden state 來自對 KV cache 的注意力。2-bit 把餵給它的東西壓爛了 —— 這不會讓輸出變差(target 每個 token 都會驗),它讓草稿變爛。愈後面的位置愈依賴前面草稿推出的狀態,所以 position 2 塌得最兇。
而我們負擔得起 3-bit:見下面的容量表,3-bit 在 64K 仍有 64.65 倍併發餘裕,而我們只需要 8。2-bit 換到的是我們根本用不到的容量,代價是 MTP 的接受率。
還沒解決的:即使 87.1% 接受率,解碼仍然只有 22.65 對不開推測的 68。接受長度算起來約 3.5,理論上該有 1.67 倍加速。接受率修好了,推測機制本身還是虧的 —— 這兩件現在乾淨分開了。
27B IQ4_XS · max-model-len 65536 · util 0.75 · block-size 32 · 只讀引擎啟動時報的 pool,不做生成 · 此模型 head_count_kv=4,vLLM 按 head 切,TP4 是分散上限
| 拓樸 | KV pool | 64K 下併發倍數 | 每卡記憶體 |
|---|---|---|---|
| TP1 | 370,085 tok | 5.65x | 20.8 GB(只有卡 0) |
| TP2 | 2,174,253 tok | 33.18x | 23.0 / 22.8 |
| TP4 | 5,751,747 tok | 87.76x | 23.8 × 4 |
KV 確實分散,而且是超線性:TP1 → TP4 是 15.5 倍不是 4 倍,因為權重也跟著切,每張卡騰出更多空間給 KV。llama.cpp 四卡 layer 也分散(5439 / 5027 / 5233 / 6049 MiB,單卡是 19177),它按層切,不受 head 數限制。
| KV 格式 · TP4 | KV pool | 64K 下併發倍數 |
|---|---|---|
turboquant_2bit_nc | 5,751,747 | 87.76x |
turboquant_3bit_nc | 4,237,044 | 64.65x |
turboquant_4bit_nc | 3,480,429 | 53.11x |
f16(auto) | 1,088,046 | 16.60x |
TurboQuant 2-bit 給的併發是 f16 的 5.3 倍。這就是為什麼 KV 格式必須寫進每一份配置 —— 它直接決定能放多少人。
27B Fable · TP4 · 4×V100-SXM2-32GB NV2 · KV turboquant_2bit_nc(K2V2)· max-model-len 65536 · util 0.75 · kv-splits 96 · block-size 32 · batched-tokens 16384 · max-num-seqs 8 · prefix caching 開 · MTP 關 · GDN full-forward 關 · 提示 ~200 tok · 三個臂都是現行的新合併版,只差量化型別
| 量化 | c1 | c2 | c4 | c8 | 批次增益 |
|---|---|---|---|---|---|
| Q4_K_S | 68.8 | 52.2 | 33.9 | 34.4 | 4.00x |
| Q4_K_M | 66.4 | 46.9 | 33.3 | 31.1 | 3.75x |
| IQ4_XS | 68.4 | 47.8 | 32.1 | 19.9 | 2.33x |
| Q4_0(舊版權重的 requant) | 75.1 | 51.6 | 34.6 | 34.5 | 3.74x |
Q4_K_S 是答案:同一顆新模型、單人不輸(68.8 對 68.4)、c8 從 19.9 拉到 34.4,超過 30 的目標,不必退回舊權重。
機制有兩個獨立佐證。派發器裡三條中等批次路徑全部要求 qweight_type == 2(Q4_0),i-quant 一條都進不去;而微基準顯示 ggml_mul_mat_a8 對 IQ4_XS 根本沒在算 —— M=2 到 M=512 的耗時全是 0.03 ms 不動,換算 3000 TFLOPS,V100 峰值只有 125,物理上不可能,那條路徑是空轉返回。
所以 i-quant 在批次時只能留在逐列的向量路徑,這就是那條 11.3 ms + 4.83 ms × 序列數 的來源。
9B Q4_K_M · TP4 · KV turboquant_2bit_nc · max-model-len 65536 · util 0.30 · block 32 · batched 16384 · max-num-seqs 32 · MTP 關 · 對線上生產直接量,未停機
| 併發 | capture [1,2,4,8](原本) | capture [1,2,4,8,16,32](現行) |
|---|---|---|
| c1 | 152.65 | 152.4 |
| c8 | 76.3 / 合計 607 | 76.3 / 合計 556 |
| c16 | 14.5 / 合計 232 | 61.0 / 合計 899 |
| c24 | — | 40.4 / 合計 897 |
| c32 | — | 39.6 / 合計 1188 |
CUDA graph 是把一步解碼的幾千次 kernel 發射錄成一份、之後一次重播,但錄影綁死批次形狀。捕捉集合只到 8,第 9 個併發就沒有東西可以墊,整步退回 eager。
代價量過了:+34 MiB/卡、+31 秒啟動,c1–c8 完全沒退化,工具呼叫每個臂 ALL_PASS。已上線,回滾是 cp fable9b_start.sh.bak_capture 再重啟。
這個不一致是 8/6 我自己造成的 —— 那天把 max-num-seqs 從 8 調到 32,只驗了 c1 和 c8,沒量 c16。調了上限就要量到上限。
權重 17408 × 5120(此模型 ffn 形狀)· Q4_0 · sm_70 · cuda event 計時 30 次取平均 · 每版都先過正確性再看時間
| 版本 | 改了什麼 | M=16 |
|---|---|---|
| v1 | 一 warp、逐位元組拆 nibble | 1.19 |
| v2 | 四 warp、N_TILE 64(委員會的分塊) | 0.63 |
| v3 | 全執行緒參與、K_TILE 128、half2 | 0.51 |
| v4 | 位元技巧解量化(免轉換指令) | 0.44 |
| v5 | M_TILE 64、多累加器 | 0.72 |
| dp4a(要打敗的) | — | 0.331 |
| 屋頂線 | 50.2 MB ÷ 900 GB/s | 0.056 |
五版全部正確(max abs err 0.0000),但沒有一格贏過。而且時間對 M 幾乎不變(v4 在 M=1 是 0.439、M=16 是 0.437)—— tensor core 是免費的,成本全在把權重物化成 fp16。
原因:WMMA 的 load_matrix_sync 只能從記憶體載入 fragment,所以解量化結果一定要先寫進 shared memory,我因此多搬了 178 MB。dp4a 在暫存器裡邊解邊乘,完全不物化。Marlin 那類能贏的 kernel 用的是原始 mma.sync PTX 手工排暫存器,繞過這個限制。
所以「用 WMMA C++ API 做融合量化 GEMM」這條在這個問題上是死的 —— 不是調得不夠好。這個否定結論擋掉一整個方向,值得留著。
微基準顯示 M=64 時 dp4a(1.290)已經輸給 dequant+cuBLAS(0.787),而門檻 DEQUANT_CUBLAS_MIN_ROWS 是 256,看起來 48–255 那段跑在慢路上。實測三個門檻:256 與 48 每格相同(解碼 M ≤ 32、prefill 分塊 ≫ 256,那段碰不到),而 16 讓 c16 掉 39%(解碼批次被推去走解量化路徑)。維持 256,這條劃掉。
27B · llama.cpp llama-batched-bench · -c 32768 -b 2048 -ub 512 · -npp 256 -ntg 128 · KV f16 · 欄位是 aggregate 生成吞吐(B 是同時在跑的序列數)
| 拓樸 | B=1 | B=2 | B=4 | B=8 |
|---|---|---|---|---|
| IQ4_XS · 單卡 | 38.79 | 63.00 | 86.24 | 110.55 |
| Q4_K_S · 單卡 | 34.45 | 56.58 | 68.96 | 72.97 |
IQ4_XS · 四卡 layer | 38.27 | 62.39 | 88.57 | 111.75 |
IQ4_XS · 四卡 tensor | — | — | — | 241.96 |
Q4_K_S · 四卡 tensor | — | — | — | 209.58 |
我寫過「四卡對 llama.cpp 的解碼沒有幫助」。那是 batch 1 量的。八個序列同時在跑時,tensor 是單卡的 2.2 倍(241.96 對 110.55),提示處理更是 1828 對 690。layer 確實不縮放(111.75 對 110.55),那部分成立。
在我們的 vLLM fork,IQ4_XS 在 8 併發塌到 19.9 而 Q4_K_S 是 34.4。在 llama.cpp 剛好相反 —— 同一顆檔案,IQ4_XS 在 B8 是 110.55、Q4_K_S 只有 72.97,i-quant 快 51%。
所以 i-quant 不是天生不適合批次,是我們的 GGUF 載入路徑沒有實作:微基準顯示 ggml_mul_mat_a8 對 IQ4_XS 空轉返回(M=2 到 512 耗時全是 0.03 ms,換算 3000 TFLOPS,V100 峰值 125)。參考實作就在同一台機器上的 llama.cpp 裡,這把「自己寫 kernel」變成「從上游搬」。
8/10 早上就量到 MTP 在 llama.cpp 賺 44%(38.8 → 55.9)、在 vLLM 虧。我接著花了一整天修 vLLM 那個虧的,而不是去看那條會賺的路能不能撐起產線。把 vLLM 當預設答案的理由只是「生產在上面」,那不是技術理由。
來源也該分清楚:逐位置 97/95/91% 那組是那篇 3090 的 vLLM 紀錄;模型作者 DavidAU 自己的 MTP 數字是 llama.cpp 的(75 → 90+ t/s)。兩個都在,而我只拿前者當標竿。
同一顆 27B IQ4_XS · llama.cpp llama-server · -c 65536 -np 8 -sm layer · KV f16 · --spec-type draft-mtp --spec-draft-n-max 3 · 對照組是 vLLM Q4_K_S · TP4 · KV turboquant_2bit_nc · maxlen 65536 · 不開 MTP
| 併發 | llama.cpp + MTP 每人 | 合計 | 接受率 | vLLM 每人 | 合計 |
|---|---|---|---|---|---|
| c1 | 66.53 | 57.30 | 80.0% | 69.6 | — |
| c4 | 30.57 | 110.07 | 80.0% | 34.1 | 136 |
| c8 | 17.24 | 123.05 | 80.0% | 34.3 | 275 |
單人從 38.27 拉到 66.53(+74%),接受率 80%。技術本身沒問題,問題是即使開了有效的 MTP,也只追平 vLLM 不開推測的水準,而八併發差兩倍以上。
llama.cpp 唯一會縮放的拓樸是 tensor(batched-bench 量到 B8 241.96 合計),但 llama-server 不支援它:llama_params_fit is not implemented for SPLIT_MODE_TENSOR, abort —— 伺服器起得來、健康檢查會過、但產不出任何內容。所以那個數字不可服務,而且它還是不開 MTP 的。
產線留在 vLLM,現在是結論不是預設。先前「留在 vLLM」的理由只是生產在上面,那不是技術理由。
Qwen3.8-27B-Uncensored IQ4_XS(15.3 GB,JonathanColetti repo,MTP 頭保留)· vLLM 1.2.1 SM70 fork · TP4 V100-SXM2-32GB · KV turboquant_2bit_nc · max-model-len 66560 · util 0.60 · block 32 · max-num-batched-tokens 65536 · seqs 64 · capture [1,2,4,8,16,32,64] · prefix caching 開 · flash_attn_version 2 · decode partition 1024 · 自寫 batched-MMVQ 與 tile-MMQ 在 4/4 worker arm · KV pool 2,170,595 token · 不開推測
| 併發 | 每人 tok/s | 合計 tok/s | TTFT p50 |
|---|---|---|---|
| 1 | 68.15 | — | 1.52s |
| 8 | 35.3 | 155.3 | 0.28s |
| 16 | 18.8 | 186.7 | 0.46s |
| 32 | 15.3 | 246.7 | 3.69s |
| 48 | 11.4 | 314.7 | 4.37s |
| 64 | 10.0 | 356.9 | 2.57s |
32 以上那三列是為了找天花板量的,不是產品承諾 —— Wayne 定的上限是八個併發使用者。留著是因為它們回答了「還有沒有餘裕」(有,而且到 64 都沒轉折),但對外要引用的是 c8 那一列。席次已改回 8、capture 改回 [1,2,4,8],省下的記憶體回到 KV pool。
| 提示長度 | 暖 TTFT | 解碼 tok/s |
|---|---|---|
| 258 | 0.25s | 67.52 |
| 3,618 | 1.04s | 65.55 |
| 9,618 | 3.17s | 62.24 |
| 21,618 | 6.21s | 56.36 |
| 43,218 | — | 48.46 |
| 64,818 | 23.09s | 42.73 |
| 八個人同時,每人一份 21.6K 文件 | 牆鐘 | 每人延遲 | 合計 tok/s |
|---|---|---|---|
| 八份不同的文件(沒有共用前綴) | 54.8s | 54.7 – 54.8s | 18.70 |
| 八份逐位元組相同的文件(先暖過) | 55.0s | 54.9 – 55.0s | 18.63 |
兩列差 0.2 秒。八個人問同一份文件,不會因為前綴相同而變快 —— 3.6 上量到的這件事,在 3.8 加上 bt 65536 之後仍然成立。對外不能承諾「大家問同一份手冊會比較快」。
合計 18.7 對短提示八人的 155:長 context 併發是八倍代價,時間幾乎全在 prefill。
但有一件事變好了:3.6 那次是嚴格排隊(一人 31 秒、最後一個 245 秒),現在八個人並行前進、延遲差距 0.3 秒。兩次文件大小不同不能比秒數,可比的是排隊變成並行 —— 「大家一起等 55 秒」比「有人 31 有人 245」好得多。
448×448 四色象限加中央黑方塊,六項全中:左上紅、右上綠、左下藍、右下黃、中央黑方塊,連「跨在四個象限交界上」都答對。只能走 llama.cpp(vLLM 的 GGUF 路徑不吃 mmproj),而且必須 --image-min-tokens 1024 —— 不加時 96×96 只換到 81 個 prompt token,加了之後 1,103。
64,818 token 下解碼 42.73,而 Codex 用他自己的探針在 65,536 下量到 42.8。兩個人、兩套工具、兩種量法。這比任何單邊的數字都可信。
提示放大 250 倍,解碼只掉 37%。意思是「單人 68」這個數字對真實工作負載不會嚴重高估 —— 32K 文件下仍有 56,64K 下仍有 43。
冷 TTFT 那一欄只有最後一列可信。六個長度是巢狀前綴,量到 43K 時前面 21K 早在快取裡。破綻是 43,218 的冷 TTFT(4.63s)低於 21,618 的(6.24s)。這個陷阱本頁 8/7 就寫過,我今天又踩了一次 —— 差別是這次在下結論之前抓到。暖 TTFT 與解碼不受影響。真正的冷 64K 是 90.71 秒,對 Codex 原設定的 162.6 秒。
fork 的 mul_mat_vec_q 把 token 位置放在 blockIdx.y,所以每個位置各開一組 block、各自把 45 MiB 的權重讀過一遍。M=2 就是讀兩次。
同一顆 kernel,無推測 2489 次 launch / 42.24 ms,有推測 2468 次 / 67.55 ms —— 次數幾乎相同、時間多 60%。那不是「呼叫變多」,是「每次呼叫做了兩倍的搬運」。
| M | fork 每位置 ms | 批次版 每位置 ms | 總時間倍數 |
|---|---|---|---|
| 1 | 0.086 | 0.084 | 1.02x |
| 2 | 0.079 | 0.055 | 1.45x |
| 4 | 0.076 | 0.038 | 2.01x |
| 8 | 0.074 | 0.032 | 2.34x |
fork 那一欄幾乎是平的 —— 完全沒有重用。M=1 持平代表可以無條件取代,不傷單人路徑。九種 IQ4_XS 權重形狀 × 每個 M,誤差與原 kernel 一致到印出的精度(兩者共用同一套 q8_1 activation 量化)。
| 併發 | 改前 | 改後 | |
|---|---|---|---|
| c1 | 67.89 | 69.09 | +1.8% |
| c2 | 48.52 | 57.27 | +18% |
| c4 | 36.74 | 48.47 | +32% |
| c8 | 20.58 | 32.31 | +57% |
| c8 合計 | 104.7 | 143.3 | +37% |
異質混合工作負載、同 session、只切一個環境變數,四臂工具呼叫閘門全 ALL_PASS。
| 時間 | 速率 p50 | TTFT p50 | 失敗 | GPU 最大 MiB |
|---|---|---|---|---|
| 5m | 47.04 | 0.65s | 0 | 16450 |
| 15m | 46.78 | 0.65s | 0 | 16678 |
| 30m | 46.75 | 0.64s | 0 | 16678 |
| 45m | 46.89 | 0.66s | 0 | 16678 |
| 60m | 46.98 | 0.64s | 0 | 16678 |
四併發、12 個窗口、速率變動不到 1%、零失敗、一小時負載後工具閘門仍 ALL_PASS。記憶體 15688 → 16678 MiB 全部發生在前 15 分鐘,之後九個窗口完全持平 —— 那是 caching allocator 遇到新的 batch 形狀後安頓,不是洩漏。
第一次:啟動器拒絕交付,因為看不到 IQ4XS_BMMVQ_READY。kernel 是好的 —— 二十分鐘前同一個 build 才量到 c8 35.71 —— 但讓它出聲的補丁寫好、傳上去、忘了執行。閘門的行為是對的:它拒絕的不是「慢」,是「無法確認快的那條有在跑」。
第二次:記憶體規則把配置器的正常擴張判成洩漏(單一窗口 500 MiB 就判死),而暖身只有四個字、壓測最長的提示有九十段填充文字。改成基準取第二個窗口之後、要連續三個窗口都漲且總量超過 1 GiB。同時工具閘門硬寫 fable-9b,對上掛成 fable-27b 的伺服器吃 404 —— 跟健康無關。
模型 Qwen3.6-27B-Fable-Fus-711-UnHeretic-NM-DAU-NEO-MAX-NEO-LOW-MTP · 量化 IQ4_XS · 拓樸 TP4(四張 V100-SXM2-32GB,NVLink NV2 全網) · KV turboquant_2bit_nc,block 32,prefix caching 開 · max-model-len 65536,max-num-batched-tokens 16384 · max-num-seqs 32,util 0.50 · capture [1,2,4,8,16,32,48,64] · MTP draft=1 · VLLM_GGUF_IQ4XS_MMQ=64
結果:KV pool 2,513,305 token,64K 請求下 38.35 倍併發餘裕,每卡 16.4 GB,decode graph 7/7,工具呼叫閘門 ALL_PASS(works via auto)。
| 併發 | 同構 每人 | 同構 合計 | 異質 p50 | 異質 p95 | 異質 合計 | TTFT p50 |
|---|---|---|---|---|---|---|
| c1 | 30.19 | 28.2 | 30.50 | — | 27.3 | 0.28s |
| c8 | 15.06 | 103.2 | 15.06 | 18.08 | 61.6 | 0.51s |
| c16 | 12.54 | 187.5 | 12.41 | 13.79 | 101.2 | 1.12s |
| c32 | 8.99 | 266.1 | 9.22 | 10.05 | 136.7 | 1.80s |
異質那組是 12 種不同提示、錯開到達、輸出長度 100–400,32 個請求全數完成零失敗。每人速率兩者差不多,合計卻差 1.9 倍 —— 真實流量的請求長度不一,尾巴會拖長 wall-clock,而同構測法讓所有請求同時開始、同時結束。先前所有併發數字都該讀成同構上界,這是那句話的量化版本。
先前每個併發都量到「剛好 80.0%」,那是同一個提示複製 N 份、temperature 0、同一個 seed 的產物 —— 一條序列的數字印了五遍。真實混合工作負載下是 51.3 / 55.1 / 56.2%,而且隨併發上升。離對照組的逐位置 97/95/91% 還有距離,那個缺口是 KV 位元寬(2-bit 對 3-bit),不是這裡的問題。
c32 最慢的那筆是 qa_long_ctx 的 2.66 tok/s —— 長 context 請求是尾巴,不是平均。
伺服器自報 GPU KV cache size: 17,794,448 tokens、Maximum concurrency for 65,536 tokens per request: 271.52x —— KV 空間可以撐 271 個滿長度請求,而我們設 max-num-seqs 8。那不是記憶體限制,是人為上限。(這也順帶關掉一條路:DeepSeek V4 那套 CSA/HCA 序列軸壓縮是為了把 KV 壓到能落 disk,我們的 KV 已經過剩 271 倍,壓它換不到任何東西。)
| 併發 | seqs=8 每人 | seqs=32 每人 | seqs=8 合計 | seqs=32 合計 | 合計增益 |
|---|---|---|---|---|---|
| 1 | 123.2 | 123.5 | 71.4 | 71.5 | — |
| 8 | 68.4 | 68.3 | 261.9 | 286.0 | +9% |
| 16 | 67.8 | 55.4 | 228.8 | 403.2 | +76% |
| 32 | 67.7 | 36.9 | 227.0 | 619.0 | +173% |
seqs=8 時 c16、c32 的每人速度紋風不動(67.8 / 67.7),因為超過 8 的請求只是在等。放寬到 32 之後,總吞吐在 c32 從 227 升到 619,2.7 倍。
但這是取捨不是純賺:每人從 67.7 掉到 36.9,同時服務的人數變四倍。佐證在 chk 欄(decode×併發 ÷ 總吞吐):從 9.55 降到 1.91 —— 牆鐘時間真的花在解碼而不是排隊了。
所以要單人最快就壓低併發(123.5),要總產能最大就放寬上限(619)。兩者不可兼得,取決於這台機器要服務幾個人。已套用 max-num-seqs 32(2026-08-06,commit 8aed933):Wayne 給的使用者上限是 16,而 32 在 c1–c8 量到與 8 完全相同,所以它只會在人多時起作用,不會犧牲單人速度。
9B 的模型檔是 6.5GB,TP4 之下每卡約 1.6GB 權重,但 nvidia-smi 每卡讀到 29GB。差額幾乎全是 --gpu-memory-utilization 0.90 預留的 KV pool:17,794,448 tokens,而 maxlen 65536 × max-num-seqs 8 最多只用得到 524,288 —— 過剩 33.9 倍。
| ctx / 併發 | util 0.90 decode/s | util 0.30 decode/s | 差 |
|---|---|---|---|
| 200 / c1 | 123.6 | 123.0 | −0.5% |
| 200 / c4 | 79.8 | 79.7 | −0.1% |
| 200 / c8 | 68.1 | 68.2 | +0.1% |
| 28000 / c1 | 106.0 | 106.0 | 0.0% |
| 28000 冷 TTFT | 6.58s | 6.58s | 0.0% |
全部落在噪音以內,所以 0.90 → 0.30 已進生產:每卡 29.0GB → 9.5GB,四張卡共釋出約 78GB。KV pool 降到 4,319,677 tokens,仍是滿長度請求需求的 65.91 倍。驗收:health 200、工具呼叫 ALL_PASS。回滾:cp fable9b_start.sh.bak_util fable9b_start.sh 後重啟。
我原本的理由是釋出記憶體之後實驗不必再停生產。那個理由不成立 —— 擋住並存的不是記憶體,是 PSU:先前疑似燒掉電源的正是「四張卡上同時載入兩個模型」的功耗尖峰,記憶體有餘裕並不解除這條紀律。
真正換到的是故障半徑。先前一次沒收乾淨的實驗留下四個 VLLM::Worker 孤兒各佔 29GB,生產重啟每次都撞 Free memory 2.83 GiB < 28.56,在 Restart=always 之下變成 34 分鐘的失敗迴圈。現在生產只需要 9.5GB 就能起來,同樣的孤兒擋不住它。
第一版 probe 一個臂都沒跑完就結束:local u=$1 label="util${u/./}" —— bash 在 local 賦值之前就展開了整行所有的詞,所以 ${u/./} 展開時 u 還沒有值,在 set -u 之下函式直接失敗,EXIT trap 23 秒後把生產放回去。第二版順便改了方法:生產本身就是 0.90 那個臂,直接量它、再改一個值重啟,停機從兩個 20 分鐘窗口降到兩次 4 分鐘重啟,而且量到的就是要上線的東西,不是它的手打近似版。
三個嫌疑各做一臂。--enforce-eager 得到 10.3 tok/s(基準 123.6)—— CUDA graph 每 token 省掉約 89 毫秒的 launch 成本,而且 log 顯示解碼是 Capturing CUDA graphs (decode, FULL),單一完整 graph,所以「分段啟動成本」這條也死了。意外的是第三臂:關掉 VLLM_SM70_QWEN_GDN_FULL_FORWARD 反而快。
| 指標 | GDN=1(原生產) | GDN=0 | 差 |
|---|---|---|---|
| c1 decode/s | 123.4 / 123.3 | 153.1 / 152.8 | +24.0% |
| c8 decode/s | 68.0 / 68.0 | 76.3 / 76.8 | +12.5% |
| c8 總吞吐 | 272.8 / 273.3 | 275.7 / 272.7 | 持平 |
| 退化指標 | 1.00 | 1.00 | — |
| 工具呼叫 | ALL_PASS | ALL_PASS | — |
A/B 用的是生產的完整配置(maxlen 65536、max-num-seqs 32、TP4、turboquant 2-bit KV),每臂跑兩趟。上線後再驗兩趟:c1 152.6 / 152.9、c8 76.5 / 77.1、deg 1.00、ALL_PASS。
性質要說清楚:這是單人延遲的改進,不是總產能的 —— c8 的總吞吐兩邊都是 273 左右。回滾:cp fable9b_start.sh.bak_gdn fable9b_start.sh 後重啟。
第一次套用時量到 c8 只有 14.6 tok/s,我判定「探針配置沒轉移到生產」並退回。那是一個沒有重複驗證的樣本 —— 它的冷 TTFT 是 0.43,而其他每一次 run 都是 0.16,典型的剛重啟未暖。在生產完整配置上各跑兩趟之後,+24% 完全重現。凡是要決定生產變更的量測,一律每臂兩趟,這條從今天起是預設。
順帶修正前一節的模型:T(n) = 6.02 + 8.32/n 是在融合路徑開著時擬合的。關掉之後一個 token 是 6.53 ms —— 那個看起來不可縮減的固定成本,有一大塊其實就是這個旗標。
| context / 併發 | GDN=1(兩趟) | GDN=0(兩趟) | 差 |
|---|---|---|---|
| 200 / c1 | 54.3 / 54.3 | 65.3 / 66.5 | +21.4% |
| 200 / c8 | 28.0 / 28.1 | 30.8 / 31.3 | +10.7% |
| 28000 / c1 | 45.9 / 46.4 | 54.2 / 54.2 | +17.4% |
| 28000 / c8 | 17.0 / 17.0 | 21.7 / 21.6 | +27.4% |
| 28000 / c8 總吞吐 | 47.4 | 72.7 | +53.7% |
先前它們被報成 only 0 content deltas 並被我當成模型失敗。那是量測工具數錯欄位(見前一節),把 reasoning 補上、max_tokens 提到 512 之後就有數字了。27B 不在生產上,所以沒有東西要套用 —— 但之後任何一次跑 27B,這個旗標都該是關的。
28000 / c1 的退化指標 GDN=1 是 1.00、GDN=0 是 0.61,看起來變差。但同一格的 think% 是 100 —— GDN=0 那臂把整個 512 token 預算都花在思考上,所以指標評的是思考文字而不是答案,而思考文字本來就比較重複。這是量測混淆,不是已證實的品質退化,但我也不能宣稱兩邊等價。要判定得把思考與答案分開評分,探針還沒有這個能力。
可判別的正確性檢查(複誦、質數序列、三個工具呼叫、toolcall_tests)兩臂全過,與 9B 一致。
今天有六次「反常結果」最後都是量測方法的問題,不是被測物。零散修補會留下不同批次混用的數字,所以整個矩陣用同一套工具重跑一次:4 個配置 × 3 種 context × 5 種併發。
關鍵改變:prefill 與 decode 分開量。先前的每 token 速率是 tokens / 總耗時,把一次性的冷 prefill 攤進了每 token 成本。現在用串流,TTFT 與間隔中位數各自成欄,冷暖也分開,並附伺服器自報的 prefix 命中率。
| 配置 | context | 併發 | 冷 TTFT | 暖 TTFT | decode tok/s | 總吞吐 | prefix | 退化 |
|---|---|---|---|---|---|---|---|---|
| 9B · TP4 | ~200 | c1 | 0.23s | 0.16s | 123.1 | 70.8 | 0% | 0.91 |
| c2 | 0.23s | 0.32s | 103.1 | 92.3 | 0% | 0.91 | ||
| c4 | 0.23s | 0.33s | 79.7 | 140.1 | 0% | 0.87 | ||
| c8 | 0.23s | 0.36s | 67.9 | 286.6 | 0% | 0.91 | ||
| c16 | 0.23s | 0.96s | 68.0 | 199.2 | 0% | 0.83 | ||
| 9B · TP4 | ~8K | c1 | 1.13s | 0.76s | 117.9 | 26.7 | 57% | 0.91 |
| c2 | 1.13s | 1.51s | 97.1 | 58.8 | 57% | 0.80 | ||
| c4 | 1.13s | 2.85s | 73.8 | 46.9 | 57% | 0.82 | ||
| c8 | 1.13s | 4.43s | 60.9 | 73.8 | 57% | 0.81 | ||
| c16 | 1.13s | 6.43s | 58.9 | 70.8 | 57% | 0.79 | ||
| 9B · TP4 | ~28K | c1 | 8.93s | 0.57s | 106.2 | 32.5 | 98% | 0.92 |
| c2 | 8.93s | 1.05s | 83.3 | 41.1 | 98% | 0.92 | ||
| c4 | 8.93s | 1.86s | 63.4 | 105.0 | 98% | 0.78 | ||
| c8 | 8.93s | 3.28s | 45.9 | 105.1 | 98% | 0.79 | ||
| c16 | 8.93s | 3.89s | 44.2 | 101.7 | 98% | 0.74 | ||
| 9B · TP2 | ~200 | c1 | 0.29s | 0.22s | 96.2 | 58.8 | 0% | 0.93 |
| c2 | 0.29s | 0.44s | 72.0 | 72.4 | 0% | 0.93 | ||
| c4 | 0.29s | 0.40s | 57.4 | 126.9 | 0% | 0.91 | ||
| c8 | 0.29s | 0.66s | 50.0 | 175.7 | 0% | 0.87 | ||
| c16 | 0.29s | 1.18s | 49.7 | 198.3 | 0% | 0.91 | ||
| 9B · TP2 | ~8K | c1 | 2.59s | 1.39s | 90.0 | 15.6 | 57% | 0.91 |
| c2 | 2.59s | 3.27s | 66.2 | 32.3 | 57% | 0.76 | ||
| c4 | 2.59s | 5.36s | 51.2 | 43.7 | 57% | 0.80 | ||
| c8 | 2.59s | 8.41s | 41.2 | 30.6 | 57% | 0.83 | ||
| c16 | 2.59s | 12.22s | 40.3 | 37.1 | 57% | 0.80 | ||
| 9B · TP2 | ~28K | c1 | 17.39s | 0.89s | 77.9 | 30.6 | 98% | 0.95 |
| c2 | 17.39s | 1.70s | 55.1 | 46.8 | 98% | 0.82 | ||
| c4 | 17.39s | 3.21s | 42.0 | 67.4 | 98% | 0.80 | ||
| c8 | 17.39s | 5.99s | 27.9 | 96.7 | 98% | 0.72 | ||
| c16 | 17.39s | 7.90s | 27.9 | 93.9 | 98% | 0.72 | ||
| 9B · 單卡 | ~200 | c1 | 0.44s | 0.36s | 69.7 | 40.8 | 0% | 0.93 |
| c2 | 0.44s | 0.72s | 47.7 | 42.5 | 0% | 0.91 | ||
| c4 | 0.44s | 0.59s | 39.8 | 87.5 | 0% | 0.91 | ||
| c8 | 0.44s | 0.73s | 33.0 | 138.2 | 0% | 0.91 | ||
| c16 | 0.44s | 2.00s | 32.9 | 123.5 | 0% | 0.91 | ||
| 9B · 單卡 | ~8K | c1 | 3.68s | 2.59s | 63.9 | 8.7 | 57% | 0.91 |
| c2 | 3.68s | 5.13s | 42.8 | 21.1 | 57% | 0.76 | ||
| c4 | 3.68s | 10.08s | 33.7 | 15.2 | 57% | 0.82 | ||
| c8 | 3.68s | 19.67s | 25.6 | 26.5 | 57% | 0.81 | ||
| c16 | 3.68s | 23.60s | 25.3 | 22.4 | 57% | 0.76 | ||
| 9B · 單卡 | ~28K | c1 | 33.83s | 1.57s | 53.3 | 13.2 | 98% | 0.92 |
| c2 | 33.83s | 3.10s | 34.0 | 15.8 | 98% | 0.92 | ||
| c4 | 33.83s | 5.94s | 24.2 | 32.0 | 98% | 0.82 | ||
| c8 | 33.83s | 11.43s | 18.8 | 44.4 | 98% | 0.77 | ||
| c16 | 33.83s | 14.30s | 16.4 | 43.1 | 98% | 0.80 | ||
| 27B · TP4 | ~200 | c1 | 1.30s | 0.35s | 60.4 | 52.2 | 0% | 0.76 |
| c2 | 1.30s | 0.70s | 45.4 | 73.2 | 0% | 0.70 | ||
| c4 | 1.30s | 0.70s | 32.0 | 109.3 | 0% | 0.74 | ||
| c8 | 1.30s | 0.80s | 30.9 | 208.6 | 0% | 0.70 | ||
| 27B · TP4 | ~8K / ~28K | — | 失敗待重跑:8K 回 0 個 content delta(疑似 max_tokens=128 被 think 吃光,即已知的串流空白陷阱);28K 回 HTTP 500(疑似超過我設的 maxlen 32768)。兩者皆未證實。 | |||||
這是不含 prefill 的逐字生成速度,也就是使用者實際感受到的吐字速率。先前報的 111.5 是 tokens / 總耗時,被一次性的 prefill 稀釋過;當時真正的單人天花板是 123.1。這個數字已被取代：關掉 GDN 融合之後是 152.8，見本頁上方。
| context | 單卡 | TP2 | TP4 | 單卡→TP4 |
|---|---|---|---|---|
| ~200 · 1 人 | 69.7 | 96.2 | 123.1 | +77% |
| ~8K · 1 人 | 63.9 | 90.0 | 117.9 | +85% |
| ~28K · 1 人 | 53.3 | 77.9 | 106.2 | +99% |
| ~200 · 8 人 | 33.0 | 50.0 | 67.9 | +106% |
| ~8K · 8 人 | 25.6 | 41.2 | 60.9 | +138% |
| ~28K · 8 人 | 18.8 | 27.9 | 45.9 | +144% |
TPS 的最大殺手是併發,不是 context。TP4 從 1 人的 123.1 掉到 8 人的 67.9(−45%),而同樣是 TP4、從短提示拉到 28K 只掉 14%(123.1 → 106.2)。所以要提升 TPS,該攻的是併發時的排程與批次效率。
另一個方向:卡加得越多,在越吃緊的情境越值得 —— 單卡→TP4 的增益從短提示單人的 +77% 一路升到 28K 八人的 +144%。
28K 的冷 TTFT 是 單卡 33.83s → TP2 17.39s → TP4 8.93s(3.8 倍),大於同條件下 decode 的 2.0 倍。這不影響 TPS 結論,但先前完全沒量過。暖起來後 28K TTFT 只剩 0.57s(prefix 命中 98%),所以多輪對話裡長 context 的代價幾乎只在第一次。
它是 content delta 數 / 牆鐘,而每則回應的長度本來就不一樣,所以同一列會出現非單調(例如 8K 的 c4 46.9 低於 c2 58.8)。decode tok/s 才是穩健的指標(串流間隔中位數,不受回應長度影響);總吞吐只適合看數量級。
| 配置 | 1 人 | 2 人 | 4 人 | 8 人 | 8 人合計 | 量測工具 |
|---|---|---|---|---|---|---|
| 9B · TP4 · 生產 | 111.5 | 88.8 | 64.6 | 60.5 | 393.1 | v2 |
| 9B · TP2 | 85.1 | 62.7 | 47.7 | 44.6 | 282.2 | v2 |
| 9B · 單卡 | 65.7 | 44.6 | 35.1 | 29.8 | 177.8 | v2 |
| 27B · TP4 · 8K · 捕捉集合僅 [1,2] | 56.2 | 40.9 | 5.2 | 5.1 | 40.7 | v1 · 待重量 |
先前我說「TP2→TP4 只有 +19%,all-reduce 已是主導項,再加卡不會有回報」—— 那是錯的,是拿兩套不同量測工具的數字相比得出的。四個拓樸現在都在 v2 下量過:單卡→TP2 +29.5%、TP2→TP4 +31.0%,每翻倍的增益一致。單卡→TP4 的正確數字是 +70%,不是先前報的 +61%。
c1 56.2、c2 40.9,c4 直接掉到 5.2(牆鐘 15.7s → 63.9s)——階梯狀而非漸進,根因在 log 裡:cudagraph_capture_sizes: [1, 2]。我建量測鏈時漏傳 --compilation-config 給 27B 那一臂,用了 fork 的預設,所以 c4、c8 掉出捕捉集合退回 eager。生產的 9B 設 [1,2,4,8],曲線就平滑。這不是 27B 的效能特性,是量測設定錯誤,該臂需重跑。
另外先前記錄的 27B TP4 = 38.4 是 64K context 量的,這次是 8K,兩者本來就不能直接比。
| 提示長度 | 1 人 | 4 人 | 8 人 | 16 人 | 8 人合計 |
|---|---|---|---|---|---|
| ~900 tok | 106.4 | 51.8 | 49.6 | 34.8 | 347.4 |
| ~7K tok | 40.3 | 17.9 | 9.6 | 9.1 | 89.8 |
| ~28K tok | 13.2 | 4.6 | 5.3 | 2.7 | 32.7 |
上表的每 token 速率是 completion_tokens / 總耗時,把一次性的 prefill 攤進了每 token 成本,而每次只生成 128 個 token —— 8.85 秒的冷 prefill 除以 128,就把整個數字壓垮了。那量到的不是「長 context 很慢」,是「冷啟動成本被除進了速率」。Wayne 指出 prefill 只有第一次要付,後續輪次是逐步增長的 context,這是對的。
用串流把兩段拆開重量(同一份 context 連發三輪):
| context | 輪次 | TTFT | 每 token 間隔 | decode tok/s | prefix 命中 |
|---|---|---|---|---|---|
| ~8K | r1 | 1.16s | 0.0085s | 117.9 | 0% |
| ~8K | r2–r3 | 0.75s | 0.0085s | 117.9 | 58% |
| ~28K | r1 | 8.85s | 0.0094s | 106.3 | 16% |
| ~28K | r2–r3 | 0.49s | 0.0094s | 106.3 | 98% |
decode 幾乎不受 context 長度影響:8K 117.9 → 28K 106.3,只掉 10%。在 TurboQuant 2-bit 下,28K 的 KV 相對於每步都要讀的 5.5GB 權重不算什麼,權重才是主導 —— 所以「每個 token 要重讀整份 KV 導致變慢」的說法也不成立。
prefill 則完全符合直覺:28K 第一次 8.85s,prefix cache 命中 98% 之後降到 0.49s,18 倍。實際的多輪對話裡,長 context 暖起來之後每 token 速度與短 context 相當。
這一條不受上述假象影響,因為它是同一列內的比較:短提示 c8 合計 347.4、c16 反而 332.4;28K 下 c8 32.7、c16 33.4 幾乎不動。超過 max-num-seqs 8 之後請求只是排隊,總吞吐不再增加。
舊工具的比例並不自洽 ——「每人 × 人數」與「合計」相差達 19%。查出三件事:
① 每個併發等級用的題目不同:PROMPTS[i % 8] 讓 c1 只跑第 0 題、c8 跑滿八題,中英文與長度都不同,曲線因此混進了「題目換了」這個變因。
② prefix caching 偏袒 c1:暖機只跑第 0 題,而那正是 c1 唯一量到的請求。
③ 兩欄不可相乘:每人取的是中位數,合計是總 token ÷ 牆鐘。
v2 讓每個等級跑相同的八題、全部預熱、三輪取中位。修正後「平均 × 人數 ÷ 合計」在 c1–c4 都是 1.00 上下,三輪離散 ≤ 0.8。
但頭條數字沒有被灌高 —— 修正後 c1 反而從 107.5 升到 109–111,因為八題全暖後平均更高。真正被破壞的是曲線形狀(c4 由 67.4 修正為 65.0)與兩欄的一致性。TP2 與 27B 兩列仍是 v1 量的,重量需要停生產的維護窗口。
那一臂在切換前執行,GPU0 還被生產佔著:Free memory on device cuda:0 (0.3/31.73 GiB) —— 排程順序問題,不是 27B TP4 的退化。先前量到的 38.4 tok/s 仍然成立。
max_tokens 太小時,推理模型會把整個額度花在思考段,串流下呼叫端會收到完全空白且無錯誤的回應(實測 200 tokens 時 4/6 空白)。請給 1024 以上。
52.5 tok/s / 每人 — 當時以隨機 token 當提示,掩蓋了輸出退化,而退化的輸出解碼更快,反而灌高數字。基準測試現已內建真實文本的詞彙重複度斷言。
GGUF metadata:qwen35.context_length = 262144 —— 模型原生訓練到 256K,128K 不需要 rope 外推,所以它純粹是供應問題。同一份 metadata 還帶出兩件事:full_attention_interval = 4(只有四分之一的層保留 KV,這解釋了為什麼 pool 用 token 數算起來這麼大,也解釋了為什麼那個 GDN 旗標影響這麼大 —— 四分之三的層都走那條路),以及 GQA 24/4、key/value length 各 256。
RuntimeError: Unsupported decode partition_size: 2048。partition_size 是 ceil(maxlen / num_kv_splits) 往上取到合法磚塊,而 131072/96 = 1366 → 2048。flash_attn_interface.py 的 VALID_DECODE_PARTITION_SIZES 裡有 2048,但編譯好的 CUDA 擴充從來沒實例化它 —— 兩邊不同步,所以錯誤要拖到 kernel 呼叫才爆。把 tq_max_kv_splits_for_cuda_graph 從 96 調到 160,per-split 變 820、磚塊回到 1024,問題就沒了。不用重編 kernel。(96 是當初為 64K 挑的:65536/96 = 683,剛好也落在 1024。)
| 項目 | 實測 |
|---|---|
| KV pool | 4,877,750 tokens = 131,072 之下 37.21x 併發(需要 8) |
| 每卡記憶體 | 19 GB / 32(util 0.60) |
| 105,856 token 提示 · c1 | 35.1 / 35.0 tok/s(兩趟) |
| 105,856 token 提示 · c8 | 8.1 / 8.1 tok/s |
| 冷 TTFT | 390.9s |
| 暖 TTFT(prefix 命中 96%) | 31.1s(c1) / 247.5s(c8) |
| needle 檢索 · 115,645 token | 5/5 全過(深度 5/25/50/75/95%) |
| 工具呼叫 | ALL_PASS |
116K 提示第一次要 492.8 秒(8.2 分鐘),prefix cache 命中之後同一份提示只要 16.7 秒。所以 128K 能不能用,取決於同一份長文件會不會被重複問 —— 會的話代價攤掉了,每次換新文件就是每次八分鐘。
prefill 是平方成長的:4.3 倍的 token 換到 16.5 倍的時間(4.3² = 18.5)。這也順帶反證那個 context 是真的長,不是標籤寫大而已。
拿 tokenizer 實際數過之後:matrix 的標籤高估 1.13 倍(~120000 其實是 105,856),而 needle 探針的英文填充文字差 4.65 倍 —— 它第一次要 120000、實際只拿到 25,804,我差點把 26K 的五連過報成 128K 通過。差別在填充文字:matrix 用中文(約 1.4 字元/token),needle 用英文(約 6.5)。探針現在用真的 tokenizer 長到目標,並把實際拿到多少印出來。
| 105,856 token 提示 | batched-tokens 16384 | batched-tokens 65536 | 差 |
|---|---|---|---|
| 冷 TTFT | 412.42s | 304.04s | −26.3% |
| 暖 TTFT · c1 | 31.08s | 31.14s | 不變 |
| 暖 TTFT · c8 | 247.73s | 246.34s | 完全不變 |
| decode c1 | 35.0 | 35.0 | 不變 |
| decode c8 | 7.5 | 7.5 | 不變 |
暖起來之後 c8 是 247.5s、c1 是 31.1s —— 比值 7.97,幾乎剛好是 8,看起來就是八個 prefill 嚴格排隊。而 --max-num-batched-tokens 整個 session 都固定在 16384,是唯一沒動過、又直接管 prefill 進度的旋鈕。
結果:冷 prefill 真的是分塊綁的,加寬到 65536 省下 26.3%(412 → 304 秒)—— 而 128K 的成本正好集中在這個數字,所以這個要拿。但 c8 的懲罰一動也不動(247.73 → 246.34)。那代表八倍是每個請求各自要付的工作,不是排程器扣住產能不放,加寬 prefill 批次藏不掉它。這是本 session 第四個被量測打掉的結構性推論。
16384 那一臂與前一輪獨立跑出來的結果吻合到小數點:暖 c1 31.08 vs 31.14、暖 c8 247.73 vs 247.56。所以這不是單一樣本,是可重複的行為。
目前沒有套用到任何地方 —— 生產跑的是 9B / 64K,冷 TTFT 本來就只有 6.6 秒。這是 27B / 128K 那套配置真的架起來時要帶的設定。
我把 c8 的八倍懲罰歸給「量測工具給每個併發請求不同的尾巴,所以 prefix 只命中 96%,缺的約 4,200 個 token 各自要對整段做注意力」。照這個說法,八個人問同一份文件(逐位元組相同)應該很便宜。不是。
| 105,856 token · 完全相同的提示 | TTFT | prefix 命中 |
|---|---|---|
| 第 1 次(冷) | 303.23s | — |
| 第 2 次 | 30.95s | 96% |
| 第 3 次 | 30.95s | 96% |
| 8 人同時問同一份 | 最快 31.05s · 中位數 245.13s · 最慢 245.14s | 96% |
一、暖起來的 106K 文件,每一次問都要 31 秒,而且永遠不會更便宜。第 2 次和第 3 次吻合到小數點兩位(30.95 / 30.95)—— 那是決定性的重算,不是還沒暖完。
二、八個人同時問同一份文件不會共用那份工作。一個人拿到 31 秒,其餘排隊,牆鐘 245 秒 —— 每多一個併發使用者就多約 31 秒。
最順的候選是 GDN 狀態重掃:full_attention_interval = 4,65 層裡有 49 層是遞迴的,它們的狀態沒辦法放進 KV prefix cache,每次都得重掃。但 SSM 掃描是線性的,而實測的暖 TTFT 是平方的 —— 24,704 → 105,856 token 是 4.28 倍,1.82s → 31.0s 是 17.0 倍(平方預測 18.4,線性預測 4.28)。所以那個解釋不成立。
另一個相容的說法是「固定約 4% 的區塊沒命中,而這些 token 各自要對整段做注意力」,那是 0.04n²,形狀對得上。但為什麼逐位元組相同的提示會有 4% 沒命中,我沒有答案。要定案得拿一個沒有 GDN 層的模型來對照,而這裡每個 Fable 變體都是同一種混合架構。先記錄,不再往下猜 —— 今天已經有四個從架構往下推的假設被實測打掉。
| 文件大小 | 暖 TTFT(兩次) | 對前一列的成長 |
|---|---|---|
| 24,704 tokens | 1.79 / 1.78 s | — |
| 49,408 | 5.90 / 5.90 s | token ×2.00 → 時間 ×3.31(n^1.72) |
| 74,112 | 12.58 / 12.59 s | ×1.50 → ×2.13(n^1.87) |
| 105,856 | 30.98 / 30.97 s | ×1.43 → ×2.46(n^2.53) |
1.72 → 1.87 → 2.53。不只是平方,而且越長越糟。每一組兩次量測吻合到 0.01 秒,所以這是決定性的計算量,不是雜訊。照最後一段的斜率外推,131,072 token 的暖 TTFT 約 53 秒。
實務上的意思:文件放大一倍,等待時間不只放大一倍。128K 不是「64K 的兩倍慢」,是更糟。如果要壓這個成本,把文件切小比買更多卡有效得多。
四個尺寸是巢狀前綴(同一段文字重複到不同長度),所以量到 120K 時前面 84K 早就在快取裡了 —— 這一輪的「冷」欄(216.67s)被污染,真正的冷數字是專門那一輪量的 412.42s(bt 16384)與 304.04s(bt 65536)。暖數字不受影響:它是同一份完整文件被問過之後再問的結果。
依序死掉的是 —— 量測工具尾巴不同(逐位元組相同的提示成本一樣)、prefill 批次寬度(16384 vs 65536 對暖 TTFT 影響為 0)、以及 GDN 狀態沒被快取(排程器裡有 _mamba_block_aligned_split,就是專門在快取它的)。我不再往下猜。今天從架構往下推的假設已經死了五個,而這個數字的形狀對決策就夠用了:成長比 n² 還快,所以答案在文件大小,不在硬體。
| 27B · 128K 配置 | KV pool | 併發餘裕 | 106K · c8 decode/s |
|---|---|---|---|
| util 0.60 · bt 16384 | 4,877,750 | 37.21x | 7.5 |
| util 0.60 · bt 65536 | 2,523,136 | 19.25x | 7.5 |
| util 0.25 · bt 65536 | 起不來 —— No available memory for the cache blocks | ||
今天轉過的每一個旋鈕 —— gpu-memory-utilization、max-num-batched-tokens、GDN 融合開關 —— 都沒能移動 c8 的 7.5。決定它的是每產生一個 token 要掃自己那 106K 的 KV,八個人就是八份,而那份 KV 是實際在用的,不會因為池子變小而變少。池子只是預留:預留少一點會省記憶體,不會省時間(9B 上也量過同一件事,pool 砍到四分之一,123.6 → 123.0)。
下限也找到了:util 0.25 之下,權重 4.3 GB/卡 加上 bt 65536 的 activation workspace 就吃光 8.1 GB 的預算,KV 一格不剩。
順帶:先前那個看起來過剩的 37.21x 已經自己降到 19.25x —— 是 bt 65536 多吃 workspace 順手擠掉的,不用另外調。
要讓 8 人的 TPS 上去,只剩三條路,而且都不在配置裡:人數減半、文件切小(見 n^2.5 那節)、或換小模型。
Muse Glimmer 30B kquant-dynamic(19.7 GB)· llama.cpp 8e7f22b(自建,arch 70,gcc 12.4)· 單卡 V100-SXM2-32GB · -ngl 99 -fa on · KV f16 · -c 8192 · /v1/chat/completions · temp 0 · max_tokens 300 · 三趟取中位數。對照的 Qwen3.6-27B 數字出自 /opt/llama.cpp(未動過的參照樹)
| 臂 | tok/s | 草稿 | 接受 | 接受率 | 每步 token | c |
|---|---|---|---|---|---|---|
| base(不開推測) | 32.63 | — | — | — | 1.000 | — |
| DFlash 預設 | 36.55 | 428 | 155 | 36.2% | 2.069 | 0.847 |
| DFlash n-max 8 | 31.55 | 901 | 184 | 20.4% | — | — |
| DFlash n-max 16 | 31.05 | 1,565 | 191 | 12.2% | — | — |
| 把成本除以「每步多驗幾格」 | 遞迴 block | 多驗幾格 | c | 每格的 c |
|---|---|---|---|---|
| Qwen3.6-27B · llama.cpp 單卡 · MTP | 48 / 65 | 1 | 0.298 | 0.298 |
| Muse Glimmer 30B · llama.cpp 單卡 · DFlash | 0 / 52 | 2.95 | 0.847 | 0.287 |
| Qwen3.6-27B · vLLM TP4 · MTP(最佳旗標) | 48 / 65 | 1 | 1.595 | 1.595 |
同一個引擎、同一張卡,一顆 65 層裡 48 個帶 GDN 遞迴狀態、一顆 52 層全注意力零遞迴,每多驗一格的邊際成本差 4%。遞迴狀態不是驗證變貴的原因。
我今晚一直把「48/65 帶遞迴」當成 c=2.01 的根因,寫進報告、也寫進給 fusion 委員會的知識庫 —— 委員會回來的三個候選全部建立在這個錯誤前提上,包括「qlen=2 對遞迴層是時間軸掃描」那個看起來最合理的。對照組花了兩小時,殺掉的是一整晚的推論方向。
差別在引擎:vLLM 每多驗一格要 1.595 個基準步,llama.cpp 兩顆模型都是 0.29。這讓 spec_sequence_masks is None 那個守衛從候選之一變成主嫌 —— 它是一條純引擎側、與模型架構無關的路徑選擇。
限制要講清楚:llama.cpp 兩顆都是單卡,vLLM 是 TP4,那一列同時跨了引擎和拓樸,不能當成純引擎差異的定量證據。但 llama.cpp 內部那兩列是同引擎、同卡數、兩種極端不同的架構,那組就足以殺掉遞迴假設。
27B Fable-Fus-711…MTP-IQ4_XS · vLLM 1.2.1 SM70 · TP4(V100-SXM2-32GB ×4,NVLink NV2 全網)· max-model-len 65536 · util 0.50 · block 32 · batched 16384 · seqs 8 · prefix caching 開 · MTP draft=1 · capture [2,4,8,16] · IQ4XS batched-MMVQ 與 tile-MMQ 兩顆都開 · 六臂同一 session、只換 --kv-cache-dtype
| KV | KV 池 | c1 解碼 | c1 接受率 | c4 解碼 | c4 接受率 |
|---|---|---|---|---|---|
| 2-bit · 不開 MTP | 3,191,988 | 59.25 | — | 48.52 | — |
| turboquant 2-bit | 2,513,305 | 33.80 | 51.5% | 24.85 | 50.3% |
| turboquant 3-bit | 2,010,644 | 33.41 | 50.0% | 26.49 | 61.5% |
| turboquant 4-bit | 1,615,145 | 33.47 | 52.3% | 27.10 | 62.4% |
| K 8-bit · V 4-bit | 1,230,798 | 34.73 | 51.5% | 27.95 | 61.9% |
| f16(完全不壓) | 546,613 | 33.71 | 51.5% | 27.77 | 61.8% |
單人接受率那一欄從 2-bit 到 f16 是 51.5 / 50.0 / 52.3 / 51.5 / 51.5%。壓到 2 位元和完全不壓,同一個數字。所以 llama.cpp 那邊的 72.4% 不是 KV 換來的,我先前把它寫成因果是錯的。
四併發那一欄確實會動(50.3% → 62.4%,3-bit 之後打平),但吞吐只從 24.85 走到 27.95 —— α 動了、S 沒動,正是 c 綁住時公式預測的樣子。這一輪反而是 c 是主因的正面證據。
f16 的代價很實在:KV 池只剩 546,613 token(2-bit 的 17%),四併發 TTFT 從 1.25 s 掉到 5.11 s。工具閘門六臂全 ALL_PASS。
這一輪的基準臂不能用。我把 capture 設成 [2,4,8,16] 套給所有臂,不開 MTP 時 batch=1 沒被捕捉、掉出 CUDA graph,所以 59.25 比同組態該有的 69.10 低。各 MTP 臂之間仍可比,但那個 59.25 不是基準。
27B IQ4_XS · vLLM 1.2.1 SM70 · TP4 · KV turboquant_2bit_nc · max-model-len 65536 · util 0.50 · block 32 · seqs 8 · capture [2,4,8,16](無推測臂用 [1,2,4,8])· 兩顆 IQ4_XS kernel 都開 · NOSYNC=1 · MTP draft=1 · 八臂同一 session,c 用本 session 自己的無推測臂算,不從記憶拿
| 臂 | c1 解碼 | 接受率 | 推測步 ms | c | TTFT |
|---|---|---|---|---|---|
| 不開推測 | 68.88 | — | 14.52 | — | 0.43s |
| MTP 原樣 | 34.29 | 51.5% | 44.18 | 2.043 | 1.46s |
FULL_FORWARD | 34.03 | 50.0% | 44.08 | 2.036 | 2.86s |
SPEC_CORE_OP | 39.62 | 49.3% | 37.68 | 1.595 | 4.42s |
003_SPEC_CORE_OP | 38.49 | 44.9% | 37.65 | 1.593 | 2.88s |
PACKED_RECURRENT_DECODE | 34.72 | 51.5% | 43.63 | 2.005 | 1.36s |
GDN_DECODE_FLASHQLA | 34.37 | 51.5% | 44.08 | 2.036 | 0.77s |
FUSED_SIGMOID_MIXED_QKV | 34.50 | 49.3% | 43.28 | 1.981 | 3.54s |
| 疊加 | c1 解碼 | 接受率 | c | TTFT |
|---|---|---|---|---|
FLASH_GRAPH 單獨 | 34.55 | 51.5% | 2.020 | 1.55s |
FLASH_GRAPH + SPEC_CORE_OP | 39.96 | 52.3% | 1.625 | 0.68s |
SPEC_CORE_OP + FUSED_SIGMOID | 38.70 | 45.6% | 1.61 | 3.61s |
SPEC_CORE_OP + PACKED_DECODE | 39.81 | 52.3% | 1.635 | 1.29s |
SPEC_CORE_OP 與 003_SPEC_CORE_OP 的 c 是 1.595 和 1.593 —— 幾乎同值,所以兩條路徑應該接到同一個底層改動。前者吞吐較高純粹是接受率沒被拖低(49.3% vs 44.9%)。其餘五個旗標全落在 c≈2.0,是雜訊。
最好的 MTP 現在是 39.62,對基準 68.88 仍是 0.575 倍。損益平衡要 c < α = 0.493,從 1.595 還差 3.2 倍。這一刀把 29.2 ms 的額外成本壓到 23.2 ms,離 llama.cpp 的 7.6 ms 還有 15.6 ms。
注意 SPEC_CORE_OP 單獨用時 TTFT 是 4.42s,基準的十倍 —— 解碼買到的 15.5% 是拿 prefill 換的。疊上 FLASH_GRAPH 就沒事了(見上表:0.68s,接受率還回到 52.3%),所以這個代價不必付。
同一顆 GGUF Qwen3.6-27B-Fable-Fus-711…MTP-IQ4_XS · 同一台 V100-SXM2-32GB。llama.cpp = 上游 b0.15.3、單卡、KV f16、-fa on · -c 8192 · --spec-type draft-mtp · p-min 0(完全不篩)· temp 0 seed 1 · 200 tok。vLLM = 1.2.1 SM70、TP4、KV turboquant_2bit_nc、max-model-len 65536 · util 0.50 · seqs 8 · 兩顆 IQ4_XS kernel 都開 · NOSYNC=1。兩邊都是 γ=1
| 引擎 | 基準 tok/s | 步 ms | α | c | S 預測 | S 實測 | MTP tok/s |
|---|---|---|---|---|---|---|---|
| llama.cpp · 單卡 | 39.16 | 25.53 | 0.722 | 0.298 | 1.327 | 1.327 | 51.97 |
| vLLM · TP4 | 68.86 | 14.52 | 0.515 | 2.010 | 0.503 | 0.503 | 34.64 |
S = (1+α)/(1+c) 對兩台引擎、四個獨立量到的數字各自命中到三位數。損益平衡條件是 α > c:llama.cpp 是 0.722 > 0.298 所以賺,vLLM 是 0.515 < 2.010 所以虧一半。而基準反過來 —— vLLM 的 68.86 是 llama.cpp 39.16 的 1.76 倍,因為它的 all-reduce 在 CUDA graph 裡,而 llama.cpp 的多卡是逐 seam 同步。
| 把 llama.cpp 的優點搬過來 | α | c | S | tok/s(基準 68.86) |
|---|---|---|---|---|
| 現況 | 0.515 | 2.010 | 0.503 | 34.64 |
| 只修 α | 0.722 | 2.010 | 0.572 | 39.4 · 仍然虧 |
| 只修 c | 0.515 | 0.298 | 1.167 | 80.4 |
| 兩個都修 | 0.722 | 0.298 | 1.327 | 91.4 |
c 是壓倒性的槓桿。只修接受率、把 α 拉到 llama.cpp 的水準,在 c=2.01 之下只有 39.4 —— 連基準的一半都不到,還是虧的。反過來只修 c、接受率原封不動,就有 80.4。我原本準備先追「草稿頭被餵錯」那條線,這個算式把順序倒過來了。
c=2.01 表示每個推測步多花 29.2 ms,而草稿頭是 65 層裡的 1 層;llama.cpp 做同一件事只花 7.6 ms。那 21.6 ms 的差距不可能是算力。
27B IQ4_XS · vLLM 1.2.1 SM70 · TP4 · KV turboquant_2bit_nc · max-model-len 65536 · util 0.50 · block 32 · seqs 8 · 兩顆 kernel 都開 · 三臂同一 session,MTP 兩臂只差一個環境變數
| 臂 | c1 解碼 | c4 解碼 | 接受率 | 工具閘門 |
|---|---|---|---|---|
| 不開推測(capture 含 1) | 68.86 | 48.59 | — | ALL_PASS |
MTP draft=1 · NOSYNC=0 | 32.42 | 24.45 | 51.5% | ALL_PASS |
MTP draft=1 · NOSYNC=1 | 34.64 | 25.05 | 51.5% | ALL_PASS |
build_gdn_spec_decode_state_contract 用布林遮罩挑出推測那幾列。每個 x[mask] 編出來都是 nonzero() 加 index(),而 nonzero() 必須把命中數讀回 host 才能決定輸出張量多大 —— 一次完整的裝置同步,一次 forward 好幾個。批次同質時(單人必然,多人只要沒有 prefill 插隊也是)遮罩全 True,x[mask] 就只是複製一份,x[~mask] 是空的,兩者都不需要 nonzero。而判斷同質與否的計數呼叫端已經在 CPU 上算好了,問它零成本。
基準步 14.52 ms,MTP 步從 46.73 掉到 43.74 ms → c 從 2.22 降到 2.01。損益平衡要 c < α = 0.52,還差四倍。接受率完全沒動(51.5%),證明這是純步成本的改動,沒有動到數值。
但這一輪最重要的不是那 6.8%。草稿頭是 65 層裡的 1 層,它多做的那次 forward 照理該花基準步的 1/65。實測花掉 2 個基準步 —— 差 130 倍。這個量級的落差不可能是算力,是管線。同步只佔其中 0.21。
同上一節,只換 num_speculative_tokens;capture 集合跟著改成 (1+γ) 的倍數(γ=1 [2,4,8,16] / γ=2 [3,6,12,24] / γ=3 [4,8,16,32]),否則會被靜默丟掉
| γ | 草稿 | 接受 | pos0 | pos1 條件 | pos2 條件 | 總接受率 | c1 解碼 |
|---|---|---|---|---|---|---|---|
| 1 | 730 | 368 | 50.4% | — | — | 50.4% | 33.21 |
| 2 | 1,374 | 410 | 44.8% | 33.1% | — | 29.8% | 30.81 |
| 3 | 2,034 | 419 | 43.5% | 30.5% | 37.8% | 20.6% | 27.86 |
| 對照組(外部,同架構) | 97% | 95% | 91% | — | — | ||
position 0 的接受率 不隨 γ 改變(50.4 / 44.8 / 43.5%),所以「草稿開太短」不成立。加長只會把後面幾個爛位置算進分母,總接受率一路掉到 20.6%,吞吐跟著掉 16%。γ=1 已經是這台機器上最好的設定。
對照組的 97 / 95 / 91% 和我們的 50 / 33 / 38% 差一個量級,這不是調參數能補的距離。而 llama.cpp 在同一顆權重上報 72.4%,它跑的是 --spec-draft-p-min 0.5 —— 信心不足就不提草稿,那些位置不進分母。把 p-min 降到 0 就能分辨那 72.4% 是更好的草稿頭還是篩過的分母。
27B · TP4 · KV turboquant_3bit_nc · max-model-len 32768 · util 0.75 · block 32 · batched 16384 · seqs 8 · --spec-method mtp · c1 · 每個臂都在同一 session 內量,基準也在同一 session 重測
| 臂 | 解碼 | 接受長度 | 步時間 | 對基準 |
|---|---|---|---|---|
| 不開推測 | 58.13 | — | 17.2 ms | — |
| spec3 · Q4_K_S | 21.90 | 3.44 | 156.9 ms | 9.1x |
spec3 · Q4_K_S · SPEC_CORE_OP=1 | 23.86 | 3.49 | 146.2 ms | 8.5x |
spec3 · Q4_K_S · 003_SPEC_CORE_OP=1 | 23.55 | 3.45 | 146.4 ms | 8.5x |
| spec3 · IQ4_XS | 18.31 | 3.61 | 197.4 ms | 11.5x |
spec3 · IQ4_XS · SPEC_DECODE_PIECEWISE=1 | 5.00 | 3.49 | 698.2 ms | 40.6x |
| spec3 · Q4_K_S · TP1 | 12.05 | 3.49 | 289.5 ms | 16.8x |
接受長度 3.44–3.61,和對照組的 3.4–3.8 同一區間。一個推測步該換到 3.5 個 token,結果步成本是 8.5–11.5 倍,所以淨值是虧的。接受率和步成本是兩件事,現在乾淨分開了。
設定層面全部試過,沒有一個能修:PIECEWISE 讓它慢 3.5 倍、TP1 慢一倍(所以屏障不是問題,平行反而在幫忙)、兩個 spec-core 旗標各買 7%、換掉 i-quant 買 21%。
| kernel | 不開推測 | 開推測 | 比 |
|---|---|---|---|
elementwise_kernel<128,4,…> | 96 | 41,292 | 430x |
_scatter_gather_elementwise_kernel | 0 | 8,160 | 全新 |
fused_sigmoid_gating_delta_rule_update(GDN 本體) | 28,704 | 8,064 | 0.28x |
不開推測跑 300 個 token(約 300 步),開推測只跑 84 個(約 23 步)。換算成每步:0.32 次對 1,795 次,多五千倍。而 GDN 本體的 kernel 反而更少(0.28x)—— 所以問題從來不在遞迴層的計算,在把 token 拆來拆去的搬運。
程式碼裡有一條免拆分的快路,但條件是:
ddtree_tree_gdn_pure_spec = ddtree_requires_branch and num_prefills == 0 and num_decodes == 0
它綁在樹狀推測上,而 MTP 是線性鏈,所以永遠進不去。單一請求開推測時 non-spec 那半根本是空的,卻還是照拆、照跑一遍空張量。
這條線上我給過的機制:遞迴狀態頻寬(算出來差三個數量級)、GDN 對序列迴圈(profiler 顯示啟動次數 1.02x)、ticket18 中等批次 kernel(開關對照,關掉還略好)、PRE_CONV_DIAG 逐層同步(被 _diag_preconv_done 擋住,每層只印一次)。四個全錯,四次都是實測推翻的。
還有一次是量測本身無效:index_select 歸因跑了兩個臂、拿到 454 對 448 看似「沒差」,但 request done 沒印出來 —— 伺服器還在載入就被我的迴圈誤判成死掉,那些計數是模型載入與 graph 捕捉的,不是解碼的。誤判來自拿 setsid 之後立刻退出的 PID 當存活判據。已改成看 pgrep,重跑中。
權重 17408 × 5120(此模型的 ffn 形狀)· V100-SXM2-32GB · cuda event 計時 · 30 次取平均 · fp16 那欄包含它必須付的解量化成本
| 列數 M | 量化 kernel | 解量化 + fp16 | 贏家 |
|---|---|---|---|
| 1 | 0.074 ms | 0.651 ms | 量化 8.8x |
| 8 | 0.177 ms | 0.656 ms | 量化 3.7x |
| 16 | 0.328 ms | 0.664 ms | 量化 2.0x |
| 32 | 0.647 ms | 0.686 ms | 量化 1.1x |
| 64 | 1.286 ms | 0.790 ms | fp16 1.6x |
| 128 | 2.584 ms | 1.020 ms | fp16 2.5x |
| 256 | 5.153 ms | 1.068 ms | fp16 4.8x |
| 1024 | 20.612 ms | 2.304 ms | fp16 8.9x |
| 2048 | 41.325 ms | 4.368 ms | fp16 9.5x |
質疑是「fp16 比 Q4 快不合理」。在解碼那一端完全正確 —— M=1 時量化快 8.8 倍。「fp16 更快」只在 M 超過約 48 之後才成立,而那是 prefill 的區間。
從算術強度推的交叉點是 M ≈ 35(V100 的機器平衡點約 139 FLOP/byte,Q4 在 M 列的強度是 4M),實測落在同一格。贏的不是位元組,是算力等級:量化 kernel 走 dp4a 用不到 tensor core,fp16 GEMM 走 HMMA。
可用的發現:DEQUANT_CUBLAS_MIN_ROWS 目前是 256,但 64–255 這一段 fp16 已經快 1.6–4.8 倍,門檻設得保守。
冷 = prefix cache 沒命中,暖 = 同一份提示再問一次。長 context 的代價幾乎全在冷啟動,所以 128K 能不能用,取決於同一份文件會不會被重複問。
| 情境 | context | 冷 TTFT | 暖 TTFT | prefix 命中 | 時期 | 註記 |
|---|---|---|---|---|---|---|
| 9B · TP4 · c1 | ~200 tok | 0.23s | 0.16s | 0% | GDN 開 | |
| 9B · TP4 · c1 | ~8K tok | 1.13s | 0.76s | 57% | GDN 開 | |
| 9B · TP4 · c1 | ~28K tok | 8.93s | 0.57s | 98% | GDN 開 | |
| 9B · TP4 · c8 | ~28K tok | 8.93s | 3.28s | 98% | GDN 開 | |
| 27B · TP4 · c1 | ~200 tok | 1.30s | 0.35s | 0% | GDN 開 | |
| 27B · TP4 · c1 | 105,856 tok | 390.90s | 31.10s | 96% | 現行 | 128K 配置 |
| 27B · TP4 · c8 | 105,856 tok | 390.90s | 247.50s | 96% | 現行 | 每人 |
| 27B · TP4 · c1 | 115,645 tok | 492.80s | 16.70s | 命中 | 現行 | needle 探針同一份提示 |
max-num-batched-tokens 16384 → 65536:冷 TTFT 412.42s → 304.04s(−26.3%)。block-size 16 → 64:暖 TTFT −15%。
27B IQ4_XS 開 MTP draft=1,每人 decode 從 c16 的 12.22 tok/s 掉到 c24 的 4.28,步時間 81.8 → 233.5 ms,而 batch 只加倍。而且 c24 與 c32 是 233.5 / 234.1 —— 幾乎同值。兩個不同 batch 花一樣的時間,通常代表它們走的是同一條路,而那條路是 eager。
一個推測步驟送出的是 席次 ×(1+draft) 個 token,vLLM 只保留能被那個倍數整除的 capture 尺寸。同一組 [1,2,4,8,16,32] 送進去,兩臂捕捉到的張數不同 —— 這是唯一的破綻,而它就印在 server log 裡:
| 組態 | 要求的 capture | 實際捕捉 | 可用席次上限 | c24 每人 |
|---|---|---|---|---|
| 不開 MTP | [1,2,4,8,16,32] | 6 張 | 32 | — |
| MTP draft=1 | [1,2,4,8,16,32] | 5 張 | 16 | 4.28 |
| MTP draft=2 | [1,2,4,8,16,32] | 0 張 | — | 連啟動都沒過 |
| MTP draft=1 | [1,2,4,8,16,32,48,64] | 7 張 | 32 | 8.17 |
draft=1 丟掉的是尺寸 1(不能被 2 整除),剩下 2,4,8,16,32 個 token,對應 1,2,4,8,16 個席次。draft=2 要 3 的倍數,那組裡一個都沒有。
第一次量到改善時,同一輪 c8 也從 17.97 掉到 9.82、工具呼叫閘門從 ALL_PASS 變成 FAILED。但那兩組數字跨了 session,而本機的規則是只有同 session 的比值算數。所以三臂重跑,只差 capture 集合:
| 併發 | old [1,2,4,8,16,32] | pow2 […,32,64] | full […,32,48,64] |
|---|---|---|---|
| c1 | 30.24 | 30.38 | 30.47 |
| c8 | 9.71 | 9.89 | 9.89 |
| c16 | 9.58 | 9.58 | 9.56 |
| c24 | 3.76 | 7.87 | 8.16 |
| c32 | 3.66 | 7.59 | 7.59 |
| toolcall ×2 | 兩次都 FAILED | 兩次都 FAILED | 兩次都 FAILED |
c8 三臂一樣,所以那個「退步」不存在 —— 17.97 在同 session 裡重現不出來。工具呼叫閘門連舊集合都失敗,所以也不是這次改動造成的。capture 對齊因此是乾淨的淨勝:c24 2.17 倍、c32 2.07 倍,小併發完全不動。含 48 比只含 2 的冪次在 c24 多 4%,沒有代價。
同一組態、同樣的舊 capture,在兩個 session 之間 c8 是 17.97 對 9.71 —— 差兩倍,原因未知。這不是這次改動的問題,但它是所有併發數字的背景雜訊:跨 session 的比較在這台機器上不成立,要下結論的對照一律同場跑完。
Leviathan、Kalman、Matias(Google Research,ICML 2023)的牆鐘加速:
S = (1 − αγ+1) / ((1 − α)(γc + 1))
α = 接受率,γ = 草稿長度,c = 草稿成本 ÷ 目標模型成本。Corollary 3.9 直接給出損益平衡條件:α > c。
| 論文 Table 2 | 我們 | |
|---|---|---|
| 目標模型 | T5-XXL 11B,單張 TPU-v4,batch 1 | 27B IQ4_XS,TP4 V100 |
| α 範圍 | 0.53 – 0.82 | 0.39 – 0.53 |
| c 範圍 | 0.007 – 0.073 | 2.04 |
| 加速 | 1.4x – 3.4x | 0.49x |
論文自己的數據就在證明 α 不是槓桿:Table 2 裡接受率最高的一列 —— T5-large,α = 0.82 —— 加速卻是最差的 1.7x,因為它的 c 最大(800M/11B)。而 α 只有 0.75 的 T5-small 拿到 3.4x。
DeepSeek-V3 報告說 MTP 模組是一層 transformer,重用主幹最後一層的 hidden state、共用 embedding 與 output head,所以它的固有 c ≈ 1/65 ≈ 0.015 —— 正好落在論文的範圍內。
我們的 2.04 來自論文 Section 3.3 的前提被違反。原文是:「we'll assume that we can run γ+1 concurrent evaluations of Mp in parallel without increasing the walltime」。在我們機器上 M=2 的前向要 1.85 倍,那個假設是假的,而差額就是我們的 c。
再往下一層:M=1 時矩陣乘只佔 8.4 ms,非矩陣乘佔 9.0 ms(52%),一次前向發 2,297 顆 kernel。一個 27B 的 decode 該有九成是讀權重,我們只有四成八 —— 所以第二個 token 位置幾乎要付全價。這同時解釋了單人只有 69 tok/s(頻寬的 27%)和 MTP 必賠,是同一個病灶。
推測解碼的整個前提是「多驗一個位置幾乎免費」。對 attention 成立 —— 多寫一格 KV,被拒絕就丟掉,零成本。對 linear attention / GDN 不成立:遞迴層沒有「丟掉一格」這種操作,它只保留最後一個 token 的狀態。
gpu_model_runner._update_states_after_model_execute:
"This is used for MTP/EAGLE for hybrid models, as in linear attention, only the last token's state is kept. In MTP/EAGLE, for draft tokens the state are kept until we decide how many tokens are accepted ... and a shifting is done during the next iteration based on the number of accepted tokens."
而那組緩衝區只在 speculative_config is not None and model_config.is_hybrid 時才配置 —— 純 transformer 開 MTP 完全不會走到這條路。
| 這顆 27B 的 block | 數量 |
|---|---|
帶 SSM 遞迴狀態(ssm_a/alpha/beta/conv1d/dt/norm/out) | 48 |
純 attention(attn_q/k/v/output) | 17 |
| 合計 | 65(64 層 ＋ 1 個 MTP 層) |
第一次我用 startswith("attn_q") 判斷注意力層,結果把 GDN block 的 attn_qkv(那是 GDN 的輸入投影)也算進去,得到「48 對 65、42%」這種加起來超過總 block 數的結果。正確是 48 對 17。而且本機七顆 27B 變形 —— 含 Qwen3.6-27B-OFFICIAL-MTP-IQ4_XS —— metadata 完全一致:qwen35、65 blocks、48 個 SSM、ssm.state_size 128。Qwen3.6-27B 本來就是混合架構。
trace 的指紋完全對上。以下 kernel 在不開 MTP 時前 25 名根本看不到,開了之後全部冒出來 —— 而且它們是搬資料的,不是算數的:
| kernel | nospec | mtp | 做什麼 |
|---|---|---|---|
memcpy32_post | — | 1879 次 / 3.58 ms | 狀態搬移 |
_scatter_gather_elementwise | — | 268 | 依接受數散射 |
indexSelectSmallIndex | — | 268 | 選出要保留的 slot |
CatArrayBatchedCopy | — | 268 | 狀態串接 |
我一度寫成「混合架構所以 MTP 必輸」。一份獨立實測在同一個架構的 Qwen3.6-27B 上,MTP γ=3 拿到 2.40 倍。而那篇量到的回捲稅是把 c 從 0.20 推到 0.30 —— 足以把 1.17x 變成 0.96x,但離我們的 2.04 差了近七倍。
所以回捲稅解釋得了「為什麼混合架構比純 dense 吃虧」,解釋不了「為什麼我們比別人的混合架構還差七倍」。那一段還沒找到,而它才是我們真正該修的東西。
他們算 c 的方法也比我好:用每步讀多少 GB —— 27B Q8 的 verifier 讀約 27 GB、MTP 頭讀約 0.43 GB → c ≈ 0.016,而不是像我從 TPS 反解。
草稿頭精度:eh_proj 是 Q8_0 的那顆模型接受率更低(39.4% 對 4-bit 的 47.8%),S 兩者都是 0.45–0.49。本機 skill 裡「MTP 配方必須 gate eh_proj → q8_0」在這顆模型上不成立。
drafter 的 CUDA graph:把它完全關掉只差 0.5%(30.48 → 30.32)。往差的方向走不花錢,往好的方向自然也賺不到。
launch-bound:GPU 其實忙到 72–76%,eager launch 只有 835 次 / 9.19 ms。我先前算出「62–70% 閒置」是拿 5 個 step 的視窗去除以 16 個 token,分母錯了。
而且就算三者全修好也沒用:c = 2.04 之下,完美草稿(α→1、γ=1)也只有 2 / 3.04 = 0.66x。
我在上一節寫:nextn_predict_layers = 1,只有一顆 MTP 頭,所以 position 1 只有 15.2% 是架構決定的。這個解釋是錯的。
一份公開的單張 RTX 3090 紀錄跑同一顆 Qwen3.6-27B、同樣一顆 nextn 頭、同樣在 vLLM 上、同樣用 TurboQuant KV,拿到的是逐位置 97 / 95 / 91%、平均接受長度 3.4–3.8。單頭如果真的會讓 position 1 崩到 15%,那份紀錄的 position 3 不可能是 91%。架構不是原因,我們這台的配置才是。
| 該紀錄(3090,單卡) | 我們(V100 ×4) | |
|---|---|---|
| 逐位置接受率 | 97 / 95 / 91% | 50.5 / 15.2% |
| 平均接受長度 | 3.4–3.8 | 1.50 |
| 不開推測 | 38 tok/s | 66.5(四卡) |
| 開 MTP | 63.8 → 85(加 TurboQuant) | 30.2(倒扣) |
| 模型格式 | AutoRound INT4,mtp.fc 解量化成 BF16 | GGUF,blk.64 跟著主體量化 |
(a) draft 頭被量化壞了。對照組刻意把 MTP 那層留在 BF16。我們指定使用的 IQ4_XS 版本,blk.64 整層(含 nextn.eh_proj)都是 IQ4_XS。draft 頭失真不會產生錯誤輸出 —— target 每個 token 都會驗 —— 它只會提出一堆被打回票的 token,症狀就是低接受率。
(b) 七路合併位移了 hidden state。這個檔是 DavidAU 七個微調的合併,而 nextn 頭是對原始 base 訓的。若是這個,requant 救不回來,只能換權重檔。
切開的方法:Q4_0-requant 對 Q4_0-ehq8-requant —— 同一次 requant,只差 eh_proj 一個張量(Q4_0 對 Q8_0)。分岔 → 是 (a),修法是把 blk.64 gate 到 Q8_0 重做一次 requant;不分岔 → 是 (b)。
| 模型 | eh_proj | 接受率 | tg |
|---|---|---|---|
| Q4_0-requant | Q4_0 | 41.7% | 42.85 |
| Q4_0-ehq8-requant | Q8_0 | 40.5% | 42.35 |
| IQ4_XS | IQ4_XS | 47.2% | 47.19 |
| Q4_K_M | Q8_0 | 49.4% | 35.19 |
分離對照沒有分岔 —— 41.7% 對 40.5% 在雜訊內,而且量化最重的 IQ4_XS 接受率反而最高。四種量化全部落在 40–49%,和 eh_proj 的精度沒有相關性。所以把 blk.64 gate 到 Q8_0 重做 requant 不會有用,這條路不用走了。
另外兩個獨立佐證都指向同一邊:換引擎不改變答案(vLLM 量到 50.5%,llama.cpp 量到 40–49%),而模型作者自己的 card 就寫著「接受率低於 50% 就不該用 MTP 版」—— 我們正好卡在那條線上。
我一度懷疑是 8/6 為了 +24% 關掉的 VLLM_SM70_QWEN_GDN_FULL_FORWARD 污染了接受率 —— fork 的註解把 full forward 稱為「active-MTP 品質防護」。
不是。原始碼裡那個旗標只是 force,另有 auto 會在偵測到 speculative config 時自行武裝;而 8/8 與 8/9 兩次量測的伺服器 log 都印了 full-forward guard armed (force=False, auto=True, spec_core=False)。防護當時是開著的,用當時的 log 就能判定,不必重跑。
官方 safetensors 直轉的 Qwen3.6-27B-IQ4_XS-mtp,同量化型別、同引擎、同提示、同 session,各兩趟:
| 模型 | 接受率 | tg |
|---|---|---|
| 官方 Qwen3.6-27B | 60.7% / 60.7% | 52.8 / 53.9 |
| 合併 Fable-Fus-711 | 47.2% / 46.2% | 47.1 / 47.0 |
60.7% 離對照組的 97/95/91% 還差 34 個百分點,而那段幾乎全在量測條件。同一顆合併模型,只換內容類型:
| 內容 | 接受率 | tg | 實際產出開頭 |
|---|---|---|---|
| 高重複 | 99.0% | 77.10 | The quick brown fox jumps over… |
| 結構化 JSON | 85.4% | 69.56 | [ { "id": 1, "name": "Alice Johnson"… |
| 程式碼 | 70.9% | 60.68 | from __future__ import annotations… |
| 散文(原本下結論用的那題) | 47.2% | 47.35 | — |
同一顆模型從 47.2% 到 99.0%。我拿來下結論的那題是短提示、新散文、temp 0、冷啟動,而且 200 個 token 全被思考前言吃掉。對照組報的是 warmed 的 code/narrative,而且它自己就把兩者分開報(79.7 對 63.8 TPS)—— 兩邊根本不在同一格。
這對生產有直接意義:我們的工作負載是工具呼叫,那是結構化輸出,落在 85–99% 那一端,不是 47%。
第一次的內容對照整組作廢:官方模型對兩個完全不同的提示回傳一模一樣的 drafted 171 / accepted 141。我的驗證腳本比對輸出 hash、判定「巧合,數字有效」—— 判定是錯的。把生成內容印出來才看見兩份都是 <think> Here's a thinking process: 開頭的同一套樣板:這是推理模型,n_predict 200 根本走不到答案,每個臂量的都是同一段前言。
現在每個臂都印出實際生成的前 80 字。題目沒生效當場就看得出來,不必等某個數字剛好可疑。
還沒解決:官方那份 GGUF 不吃 /no_think(四個臂全都還是思考前言),所以官方欄仍在量前言,不能直接跟合併版的 70.9 / 85.4 / 99.0 比;另有兩格回傳空字串,標為無效,原因未查。
我先前寫「推測解碼在這裡結構上不可能贏,加速比 = τ/uq < 1」。那是錯的。這個技術的前提就是「驗證 k 個 draft 只花約一次 forward」——k+1 個 query 位置共用同一次 KV 掃描,所以加速比是 τ。產生 τ/uq 的是這個 fork 的後端,不是技術本身。
模型作者(DavidAU)的 model card 給了獨立佐證:同一顆模型在 llama.cpp 上,一般 GGUF 約 75 t/s、MTP 版可超過 90 t/s —— 那邊 MTP 贏 20%,我們這邊虧 65%。變數是引擎,不是模型也不是量化。
| spec-tokens | 接受率(greedy) | τ | position 0 | position 1 |
|---|---|---|---|---|
| 2 | 29.9% | 1.60 | 44.7% | 15.2% |
| 1 | 50.5% | 1.50 | 50.5% | — |
作者的 card 說接受率低於 50% 就不該用 MTP 版。我們在 spec-tokens 2 是 29.9%,降成 1 之後是 50.5% —— 剛好跨過門檻。原因在 metadata:qwen35.nextn_predict_layers = 1,這顆模型只有一顆 MTP 頭。要它吐第二個 token,只能拿第一個 draft 的隱藏狀態再推一次,誤差滾雪球 —— position 1 只有 15.2% 就是這麼來的。
溫度不是主因:作者建議的 temp 1.0 之下是 34.1%(spec 2)與 45.4%(spec 1),和 greedy 同一個量級。
接手檔原本寫「kernel 修好之後 τ ≈ 1.5–2.5x」。實測 τ 是 1.45–1.60,所以實際期望是 1.5–1.6x。已更正 —— 給錯的驗收目標比不給更糟,會讓一個成功的修正看起來像失敗。
27B 開 MTP 之後輸出崩壞,而且只在 spec-tokens ≥ 2 時發生。根因在 turboquant_attn.py:reorder_batch_threshold 是 2,所以 qlen 為 2(spec-tokens 1)走 _spec_decode_attention 正常,qlen ≥ 3 掉進 prefill 路徑,每一題都錯。那個 uniform-decode 迴圈寫成 for j in range(uq),本來就通用,只有這個門檻把 3 以上擋在外面。
| 臂 | threshold | 算術 | 複誦 | 序列 | 工具呼叫 | deg | decode/s |
|---|---|---|---|---|---|---|---|
| CTL_spec3(修正前) | 2 | 錯 | 錯 | 錯 | 0/3 | 0.21 | 21.7 |
| FIX_spec2 | 8 | 錯 | 對 | 對 | 3/3 | 1.00 | 24.9 |
| BASE_nospec(不開推測) | — | 錯 | 對 | 對 | 3/3 | 1.00 | 54.2 |
| FIX_spec3 | 8 | 錯 | 對 | 對 | 3/3 | 1.00 | 20.8 |
threshold 2 之下 12345×6789 答成 8333333333…、複誦只吐出一個 X、質數序列變成 2, 3, 5, 5, 5, 5, …1111111…,三個工具呼叫一個都沒發出來。threshold 改成 8 之後複誦、序列、三個工具呼叫全部通過,退化指標從 0.21 回到 1.00。
我原本從「兩個 threshold 8 的臂算術都錯,而且錯成不同的數」推論還有殘留 bug。無推測基準線把這一半推翻了:BASE_nospec 完全不開推測,算術一樣錯(8380905)。這題是 12345×6789,27B 心算五位乘四位本來就會錯 —— 那道題不是有效的正確性判別,複誦、序列與三個工具呼叫才是,而它們在 threshold 8 之下全過。
剩下的疑點沒有消失但弱得多:三個臂給出三個不同的錯答案(8380905 / 83810106 / 83809155),greedy 之下本該逐 token 相同。但推測解碼會把 qlen 從 1 變成 3,GEMM 形狀跟著變、浮點歸約順序跟著變,在接近平手的 argmax 上翻面是合理的 —— 這無法區分殘留 bug 與浮點非決定性。要判定得換測法:同一個提示,把開與不開推測的完整 token 序列做 diff,而不是問它答案對不對。這個探針要改。
同條件的無推測基準線量到 54.2 tok/s,對上 spec-tokens 2 的 24.9 與 spec-tokens 3 的 20.8 —— 退化 −54%。(先前我引的 54.3 來自另一個 session,當時不該那樣比;基準線重現到 54.2,所以那個比較事後看是站得住的,但理由是運氣不是方法。)
慢的機制是清楚的:_spec_decode_attention 把 uq 個驗證位置拆成 uq 次獨立的全 KV 掃描(out[j::uq] = self._decode_attention(...)),每一次都重掃整個 cache,所以成本是 1/uq。而接受長度 τ 實測 1.5–2.5,永遠不可能超過 uq —— 結構上就贏不了。
真正的解法是把 triton_turboquant_decode.py 改成多 query 的 kernel(目前簽章 Q_rot_ptr # [B, Hq, D] 只吃一個 query,累加器是 tl.zeros([BLOCK_D]))。天花板是 τ 倍,即 27B 約 80–135 tok/s。
fusion 委員會建議「把反量化融合進 attention kernel,可再擠 15–20%」。讀 triton_turboquant_decode.py(820 行)之後確認:反量化本來就在 kernel 內 —— 2/3-bit 索引解包、scale/zero 重建、centroid 查表全都在裡面,沒有先物化成 FP16 這一步。這條路沒有東西可以省。
MTP 在自由文本上是淨損,但工具呼叫的輸出高度重複(JSON 骨架、欄位名、<tool_call> 標籤幾乎每次相同)—— 那是 n-gram 命中率最高的地形。因此每個臂同時測工具呼叫與中文散文兩種負載:前者驗證假設,後者當對照組,確認它不會拖累一般對話。三種方法都不需要額外的 draft 模型。
| 方法 | 工具 c1 | 工具 c4 合計 | 散文 c1 | uniq | 備註 |
|---|---|---|---|---|---|
| baseline(無推測) | — | — | — | — | 排隊中 |
| ngram · 3 tokens | — | — | — | — | 排隊中 |
| ngram · 5 tokens | — | — | — | — | 排隊中 |
| ngram_gpu · 5 tokens | — | — | — | — | 排隊中 |
| suffix · 5 tokens | — | — | — | — | 排隊中 |
散文臂帶退化斷言。ngram_gpu 正是它攔下的:散文 25.6 tok/s 是所有推測臂最高,但輸出已是垃圾 —— 沒有閘門會被記成「最快」。但我後來去驗這道閘門,發現它對中文是瞎的。舊寫法用空白切詞,而中文不用空白,純中文輸出常常只切出一個 token,1/1 = 1.00,看起來滿分健康。在生產上實測校準,四種退化有三種拿到滿分:
| 樣本 | 字元 3-gram | 舊的 uniq-word |
|---|---|---|
| 中文正常輸出 ×3 | 0.92–1.00 | 1.000 |
| 英文正常輸出 | 0.578 | 0.713 |
| 單字重複 300 次 | 0.003 | 1.000 |
| 短句迴圈 | 0.022 | 1.000 |
| 兩句交替 | 0.034 | 1.000 |
| 英文詞迴圈 | 0.008 | 0.008 |
改用字元 3-gram(先去掉所有空白,中英文一視同仁):健康最低 0.578、退化最高 0.034,分離度 17 倍,門檻訂 0.40。這也代表先前那批推測臂裡若有純中文的退化,我根本看不見 —— `ngram_gpu` 是因為英文混雜才被抓到。
vLLM 的 ngram proposer 比對的是精確 token ID 序列,不是「這看起來像 JSON」—— 結構相似不等於 token 序列相同,所以工具呼叫的重複性並不會轉成命中率。而且失敗是正確性問題而非速度問題:三個推測臂都產不出合法的 tool_calls。同樣的量化 KV + ngram 失敗在 RTX 3090 上也有回報(vLLM issue #40831),因此不能歸因於 SM70 世代。
驗 ggml_mul_mat_a8 對 IQ4_XS 的正確性時,它跟獨立參考值的 max abs err 是 0.00000,三個 batch 大小都一樣。量化整數 kernel 跟 fp16 反量化再矩陣乘,光是累加順序不同就該有 1e-3 等級的差 —— bit-exact 相同只可能是兩者跑了同一段 code,或讀了同一塊記憶體。
時間說得更大聲:M=64 跑 20 次共 0.56 ms,換算 407 TFLOPS,而 V100 的 fp16 tensor core 峰值是 125。
profiler 給出答案:那一次呼叫只發射了 quantize_q8_1(量化 activation),之後沒有任何 matmul。輸出緩衝區從頭到尾沒被寫入 —— torch 的 caching allocator 把剛算完參考值、剛被釋放的那塊記憶體配給了它,所以比對讀到的是自己上一步的答案。
| 路徑 | 實際發射的 kernel | M=64 時間 |
|---|---|---|
| 反量化整顆 + cuBLAS | dequantize_block_iq4_xs 0.663 + cutlass 0.441 | 1.106 ms |
| mmvq | mul_mat_vec_q<block_iq4_xs> | 4.169 ms |
| mmq | 只有 quantize_q8_1 | 不存在 |
任何人若順手把 IQ4_XS 加進 MMQ_QUANT_TYPES 來修批次崩塌,模型會安靜地吐垃圾,而且不會報錯。目前生產沒中,只是因為 gguf.py 正好把 i-quant 排除在那條分支外。
做法上的教訓:正確性閘門要毒化輸出緩衝區,並且斷言誤差不為零 —— 對上獨立實作,零誤差是可疑不是優秀。再配一道 roofline(搬的位元組 ÷ 時間 對照卡的峰值),因為沒運算的 kernel 一定快得不合理,那是最便宜的破綻。
IQ4_XS 在 M>8 掉進「反量化整顆權重再 cuBLAS」,固定 1.106 ms 且不隨 M 變;同尺寸的 K-quant 走真 MMQ 在 M=16 只要 0.355 ms。mmvq 是線性的,與反量化路徑的交叉點在 M≈15,而 gguf.py 的門檻寫死在 8。
| M | mmvq | 反量化 + GEMM | 目前會選 |
|---|---|---|---|
| 1 | 0.086 | 1.112 | mmvq |
| 8 | 0.585 | 1.098 | mmvq |
| 12 | 0.871 | 1.101 | 反量化(慢 26%) |
| 16 | 1.157 | 1.076 | 反量化 |
| 64 | 4.408 | 1.214 | 反量化 |
| M | 新 kernel | dispatcher 現行會選的 | 倍數 |
|---|---|---|---|
| 4 | 0.213 | mmvq 0.283 | 1.33x |
| 8 | 0.214 | mmvq 0.529 | 2.48x |
| 16 | 0.205 | 反量化 1.017 | 4.95x |
| 32 | 0.421 | 反量化 1.083 | 2.57x |
| 64 | 0.616 | 反量化 1.202 | 1.95x |
| 96 | 1.005 | 反量化 1.410 | 1.40x |
端到端,同 session、只切 VLLM_GGUF_IQ4XS_MMQ、MTP draft=1:
| 併發 | 關 | 開 | 每人 | 合計 |
|---|---|---|---|---|
| c1 | 30.42 | 29.91 | −1.7%(對照組,M=2 兩邊都走 mmvq) | — |
| c8 | 9.72 | 14.71 | +51% | 72.7 → 108.8 |
| c16 | 9.58 | 12.54 | +31% | 140.2 → 186.4 |
| c24 | 8.15 | 9.85 | +21% | 182.4 → 217.7 |
| c32 | 7.60 | 9.00 | +18% | 225.3 → 265.9 |
兩臂產出的文字逐字相同。工具呼叫閘門兩邊一樣失敗 —— 那個問題在 kernel 關閉、也在每一組 capture 集合下都存在,不是這次引入的。
c1 開 kernel 是 29.97–30.01、關是 30.45–30.47,三輪六次讀數完全一致,不是雜訊。而 c1 的 M=2 根本沒進 kernel 的窗口。
我第一個假設是 dispatch 裡那行 os.getenv 每次矩陣乘都跑 —— 把它提到模組層級後只從 29.90 移到 29.99,假設是錯的。而且用算的就能排除:條件在 4 <= x.shape[0] 短路,開跟關差兩次整數比較,60 層算下來約 25 µs,實際差距卻是每步 0.6 ms,差 24 倍。
所以成因在別處 —— 最可能是 graph 捕捉時把 kernel 錄進去、連帶改變記憶體池佈局。我沒有證據指向更細的原因,就先照實記著。操作上的結論很明確:只有單人用就別開,有併發就開,1.5% 換 51%。
第一版全錯,相對誤差 ~1.0,因為我沿用了標頭的 QR4_XS/QI4_XS。那是 8 和 8,描述的是 MMVQ 怎麼把超區塊分給執行緒,不是 MMQ 的 tile 幾何 —— 用了它們,模板每次推進四個超區塊而 load_tiles 只給一個。現在用具名常數並綁了 static_assert。
正確性門檻本來是拍腦袋的絕對值,在 M=1 給過一次假警報:同一顆張量兩次隨機輸入是 9e-4 和 3.8e-2,差別只在分母是單列最大值。改成拿 mmvq 當標尺——「比要取代的那條路徑差嗎」——並固定亂數種子。九種權重形狀 × 五種 M,誤差跟 mmvq 一致到印出的精度。
真正的意外是 fork 裡早就有手寫的批次 kernel:mul_mat_q40_dec_y(M≤32,註記 vs fork 1.26–1.84x)與 mul_mat_q40_x64(32<M≤96,1.6–2.35x),runtime 編譯、env 開關、shape gate 一應俱全,但都卡在 qweight_type == 2 —— Q4_0 專用。
而 IQ4_XS 的 32 元素子區塊剛好等於一整個 Q4_0 區塊:都是 4 個 int,低 nibble 都是元素 0–15、高 nibble 都是 16–31,而泛型模板的 k += vdr=4 讓 k/4 兩邊都正好是子區塊索引。所以 tile 幾何原封不動,只有兩處要換:scale 存 d×(ls−32)(每子區塊一個,粒度與 Q4_0 相同),量化值經 kvalues_iq4nl 查表而不是 q−8。load_tiles_iq4_xs 與 vec_dot_iq4_xs_q8_1_mul_mat —— 見上一節,做完了。
llama.cpp 是參照組,不是候選。它證明了 MTP 在這顆模型上本來就該賺,所以 vLLM 這邊虧是我們的配置有問題 —— 這個證據值錢。但它自己不能當產線。
| 項目 | llama.cpp | vLLM(我們的 fork) |
|---|---|---|
| MTP | 38.8 → 55.9 t/s(+44%,單卡短 context) | 虧,qlen 2 的 target forward 是 qlen 1 的 1.88x |
| TurboQuant KV 解量化 | 沒融合。K 量化在 QKT 內圈,pp 掉 54% | 融合進 attention kernel |
| 四卡 | 四種模式沒有一個贏過單卡,見下 | TP4 可用 |
我先後報過 row 四卡 22.0(慢 43%)、再更正成 37.82、layer 38.9、tensor 34.47,並據此下結論說「四張卡沒有一個贏過單卡、多的三張卡什麼都沒做」。那些數字一個都不成立,結論也不成立。
破綻是 pp512:提示處理是算力綁的,四張卡該接近四倍,我卻拿到單卡 852.93、row2 853.34、row4 853.91、關 peer copy 854.24 —— 差 0.15%。真的分到四張卡不可能和單卡撞得這麼準。
原因在 llama-bench 自己印的表裡,我沒去看:ts 欄是 1.00,不是 1.00/1.00/1.00/1.00,而且 row2 印了兩對數據、row4 印了四對。-ts 的語法是 <ts0/ts1/..> —— 斜線分隔;逗號在 llama-bench 是「掃多組設定」。所以 -ts 1,1,1,1 被讀成「跑四次,每次 tensor-split = 1.00」,也就是全部塞在 GPU 0。四個臂量到的是同一件事。
這是同一天的第三次同類錯誤(還有 -ctk/-ctv 的位置配對、llama-cli 的 22.0):旗標吃進去、沒報錯、但做的不是我以為的事。重跑的版本不再下 -ts(預設就按 VRAM 平均切),並且每個臂在跑的時候取樣四張卡的記憶體,1–3 號卡沒吃到權重就把該臂標成 INVALID —— 不看旗標,看卡上有沒有東西。
| llama.cpp 拓樸 | 各卡峰值 (MiB) | pp512 | tg128 |
|---|---|---|---|
| 單卡 | 14645 / 10 / 10 / 10 | 855.54 | 39.24 |
layer 四卡 | 4041 / 3629 / 3835 / 4365 | 863.42 | 38.46 |
row 四卡 | 4891 / 4681 / 4685 / 5055 | 88.99 | 22.12 |
tensor 四卡 | 4519 / 4313 / 4313 / 4313 | 2333.09 | 33.61 |
提示處理是會縮放的:tensor 模式 2333 對單卡 855,2.7 倍。我先前說「多的三張卡什麼都沒做」是錯的。
row 模式是病態的,不是「不會變快」:pp512 掉到 88.99,比單卡慢十倍。先前那個假的 37.82 剛好把這件事整個蓋住。
撐過來的只有一項:解碼確實不隨卡數變快,三種模式都輸給單卡。但這次是真的量到的,而且理由清楚 —— batch 1 每層都要四路 all-reduce,65 層的同步次數吃掉頻寬省下的量。成本在屏障的數量,不在它的大小。
-sm tensor 一直 abort 在 ggml_backend_cuda_comm_allreduce_nccl,NCCL 印的是 Cuda failure 'CUDA driver is a stub library'。但 libcuda.so.1 指到的是真的 driver(580.173.02,96 MB),不是 stub —— 那句話是舊 NCCL 對「driver entry point 查不到」的誤譯。
版本對不上:系統 NCCL 是 2.18.3+cuda12.0(2023 年),driver 卻回報 cudaDriverVersion 13000。把 vLLM venv 自帶的 2.27.5+cuda12.9 用 LD_PRELOAD 蓋上去,tensor 模式立刻跑起來(exit 0)。
這一條不受上面那個 -ts 錯誤影響:崩潰與不崩潰是二元的,和權重切在哪張卡無關。但它當時測到的速度(34.47)跟其他四卡數字一起作廢。
硬體側也確認過,而且這部分是獨立量的:四卡全 NV2 網狀、nvidia-smi topo -p2p r 全 OK。所以四卡到底能不能加速,現在是未知而不是「不能」—— 重跑中。
先前的 llama.cpp 數字用 f16 KV,對上 vLLM 的 2-bit —— 那不是比較。補測之後,llama.cpp 開 TurboQuant KV 反而更慢,MTP 的增益也從 +26% 塌到 +2%。
受控組(llama-bench 的 -ctk/-ctv 是位置配對,所以這三列其實只量到 V):
| type_k | type_v | pp24576 | tg128 |
|---|---|---|---|
| f16 | f16 | 711.6 | 39.41 |
| f16 | tbq3_0 | 330.3(−54%) | 36.30(−8%) |
| f16 | tbq4_0 | 377.3(−47%) | 36.38(−8%) |
V 量化只花 8% 的生成速度。但 K 與 V 都量化時整體掉到 25.5 t/s,多出來的約 30% 全記在 K 頭上 —— K 在 QKT 的內圈,每個 query 對每個 key 都要重讀重解一次,V 只在加權求和時碰一次。-ctk f16 -ctv tbq3 是還沒測的中間操作點。
順帶排除:flash attention 不是變數(f16 開關 FA 是 703.5 對 710.5)。早先 llama-bench 中止是因為 tbq3 必須要有 FA,-fa 0 連 context 都建不起來。
V100 32GB 的 HBM2 是 900 GB/s。TP4 之下每卡每 token 要串的權重:9B 是 1.63 GB、27B 是 4.31 GB,對應的屋頂是 552 與 209 tok/s。實測 123.6 與 54.2 —— 22.4% 與 25.9%。而且 TP 每翻倍只 +30%(理想 +100%),擬合 T(n)=固定+w/n 得到約 45% 的解碼時間不隨卡數縮減。
我一開始把它講成互連頻寬問題。那是錯的 —— TP 之下權重切開放在各卡,跨卡的只有 activation,而且只在 row-parallel 層(o_proj、down_proj)之後把部分和加起來。payload 只有 hidden×2 bytes ≈ 8 KB,每 token 全部加起來約 1 MB,在 NVLink(NV2,每卡 6 條 ×25.78 GB/s)上是 20 微秒,可忽略。真正的成本是每 token 128 道同步屏障,不是頻寬。(拓樸已確認:四卡兩兩 NV2,不是 PCIe。)
output.weight 是 BF16,而其他所有權重都是 Q4_K/Q6_K。詞表 248,320 很大,所以這一個張量就是 9B 的 29.2%(1940 MB)、27B 的 13.8%(2425 MB),而且每產生一個 token 就要完整讀一次。
| 臂 | 大小 | c1 decode/s | c4 decode/s | think% | deg |
|---|---|---|---|---|---|
| Q4_0 · bf16head | 6,359 MB | 128.9 | 85.9 | 3 | 0.98 |
| Q4_0 · q6head | 5,198 MB(−18.3%) | 136.1 | 88.4 | 100 | 0.76 |
純頻寬模型預測 +22.3%,實測 +5.6%(c1)、+2.9%(c4)。位元組不是綁住解碼的東西 —— 這與「只跑到屋頂 25%」互相印證:真正的瓶頸在反量化的計算量、同步屏障與 kernel 效率,不在記憶體頻寬。所以「把 head 量化掉」這條路關掉:5% 的收益還要付 deg 0.98 → 0.76 的品質代價。
連帶也修正了先前那個 +20% TPS 的推估 —— 那是從位元組比例外推的,不是量出來的,實際只有四分之一。
這一輪有三個結果被丟掉(27B 的 28000/c4、c8,以及整個 q6head 臂),報成 only 0 content deltas,我一度當成模型壞掉。真相:工具只數 delta.content,而思考 token 走別的欄位。第一次修我改成也數 reasoning_content —— 還是 0,因為這個 fork 送的欄位叫 reasoning,不是 OpenAI 慣用的 reasoning_content。抓原始 SSE 才看到。工具現在兩個都吃,並多一個 think% 欄 —— 上表 q6head 的 c1 是 100% 思考,這種臂跟 3% 思考的臂放在一起比,本來就該標出來。
| 拓樸 | c1 decode/s | 該拓樸的記憶體屋頂 | 屋頂效率 | 每 token | 固定成本佔比 |
|---|---|---|---|---|---|
| TP1 | 69.7 | 138.7 | 50.3% | 14.35 ms | 42% |
| TP2 | 96.2 | 277.3 | 34.7% | 10.40 ms | 58% |
| TP4(生產) | 123.4 | 554.7 | 22.2% | 8.10 ms | 74% |
加卡只把可縮減的 8.32 ms 越切越薄,那 6.02 ms 動都不動 —— 所以屋頂效率一路從 50.3% 掉到 22.2%,到了 TP4 每個 token 有 四分之三 是這個固定成本。
我先前說固定成本是 row-parallel 那 128 道 all-reduce 屏障。錯了:TP1 一次 all-reduce 都沒有,固定成本照樣是 6.02 ms。今天被打掉的三個結構性推論依序是 —— 互連頻寬、LM head 位元組、同步屏障。共同模式是我從架構往下推,而不是先量。
順帶更正數字:先前的 TP1→TP2 +29.5%、TP2→TP4 +31.0% 是把兩組不同量測鏈接算的。同 session 直接量是 +38.0% 與 +28.3%,TP1→TP4 是 +77.0% 不是 +70%。
既然邊際報酬這麼差,同樣四張卡拆成四個獨立副本反而多得多:
| 配置 | 4 使用者總吞吐 | 16 使用者總吞吐 | 單人速度 |
|---|---|---|---|
| 1 × TP4(現況) | 123.4 | 153.4 | 123.4 |
| 4 × TP1 副本 | 278.8(2.3x) | 333.6(2.2x) | 69.7 |
單人速度掉 44%(123.4 → 69.7),前面要加一個 router。9B 只有 6.5 GB,一張卡一個副本很寬鬆(util 0.30 之後每卡才 9.5GB)。但啟動必須序列化 —— 四個模型同時載入正是先前疑似燒掉 PSU 的功耗型態,記憶體有餘裕不解除這條紀律。
那 6.02 ms 還沒定位。已知不是頻寬也不是通訊;剩下的嫌疑是 piecewise CUDA graph 在 21 個 splitting_ops(整條 qwen_gdn_* 都在裡面)上被切斷造成的逐層 kernel launch、GDN 本身、以及 248,320 詞表的 logits 與取樣。下一輪量這個。
V100 唯一的對外介面是 USB 無線網卡 wlx6c4cbc1c49f1,SSID Luxshare-Mobile(5GHz,WPA-Enterprise 802.1X,RADIUS 是 Cisco ISE)。主機板的有線埠 enp8s0 存在,但狀態是 NO-CARRIER —— 網路線沒插。開機至今三天沒重開過(boot 8/3 13:41)。
| 指標 | 實測 |
|---|---|
| 3 天內重新關聯次數 | 91(約每 48 分鐘一次) |
| 其中硬斷線 | 11 |
| 訊號強度 | −66 dBm |
| 漫遊經過的 AP 數 | 至少 7 個 BSSID |
| NAT 型態 | MappingVariesByDestIP: true(硬 NAT,被迫走 DERP 中繼) |
00:53:43 tailscaled 同時失去 IPv4 與 IPv6 上游:控制面 no route to host、所有 v6 目標 network is unreachable、DERP hkg 連不上、PollNetMap 失敗。00:54:22 wpa_supplicant 記下 CTRL-EVENT-DISCONNECTED reason=7 —— AP 端早就把它踢掉了,主機晚了 40 秒才發現。00:54:29 接上另一台 AP。一台不會移動的桌機之所以一直在漫遊,是因為 −66 dBm 剛好卡在好幾個 cell 的邊緣。
先前我對這次離線給過四個錯誤診斷,其中最糟的一個是拿家用網段的 192.168.1.17 去做 ARP 測試然後宣告「網卡沒供電」—— 那台機器在公司,網段是 172.30.166.0/23,根本不在那個子網。現在有實證了:純粹是網路,而且是 WiFi。
軟體端沒有解 —— 沒有任何設定能讓 802.1X 無線關聯穩定下來。唯一的修法是實體的:插一條網路線到 enp8s0,人要在辦公室才做得到。
順帶一個量測紀律:對這台主機單發一次 ICMP 或 TCP 探測毫無意義,每天約 30 次的重新關聯裡任何一次都會讓它失敗。監視器已改成連續兩次失敗才報,並且打應用層端點而不是 ping。
網卡是 TP-Link Archer T2U PLUS(RTL8821AU,單天線 AC600),驅動 rtl8821au。iw phy 明講 interface combinations are not supported —— 單一 radio、不支援介面組合,沒辦法同時掛兩個 STA 關聯。整台機器只有這一個 phy,沒有第二支網卡也沒有 PCI 無線。要第二條路就必須有第二個 radio:再一支 USB 網卡、一支 4G dongle,或者最省事的 —— 那條沒插的網路線。
管制區是 TW: DFS-FCC。它現在關聯的 AP 在 5720 MHz(ch 144),而 iw phy 對這個頻率標著 (radar detection) —— 那是 DFS 頻道。DFS 的規定是 AP 一旦偵測到雷達就必須在 10 秒內清空頻道,把所有 client 踢掉。那正好produce「突然、全斷、短暫」的斷線型態。
而且它待的那台訊號只有 59,同一個 SSID 有一台 78:F1:C6:3F:9F:0E 在 5180 MHz(ch 36,非 DFS)、訊號 77。11 次硬斷線裡 8 次是 reason=15(4-way handshake timeout),也就是重新協商金鑰失敗,而不是訊號不足 —— 與 DFS 強制清場後的重連失敗一致。
1. 插網路線到 enp8s0 —— 免費、一勞永逸。
2. 沒線的話,把連線釘在 78:F1:C6:3F:9F:0E(ch 36、非 DFS、訊號 77)。
3. 真的要雙路徑,加第二支 USB 網卡;第二個網路現場有 Luxshare-Office(802.1X,77)可選。
這三件都要人在現場。我連 power_save off 都改不了 —— 那台機器上 sudo 要密碼,我沒有 root。(網卡目前 Power save: on,USB 無線開省電本身也是掉線的常見原因。)
| 修法 | 修法前 | 修法後 | 說明 |
|---|---|---|---|
| num_anchors 512 → 128 | 跑不到第一個 micro-step | 穩定推進 | Q_LEN 3584 → 896,scores 面積少 6.2 倍。硬體逼出的偏離,τ 需註明在 128 anchors 下測得。 |
| expandable_segments | 49→99 步流失 12.3GB | 流失 1.7GB | 治裝置端碎片。有效但不足以收斂。 |
| pin_memory=False | 步 464 剩 12.3GB 後陣亡 | 步 900 仍有 37.4GB | 治主機端釘選快取。GB10 主機與裝置共用實體 RAM,釘選換不到 DMA 好處 —— 每步耗時反而從 5.0s 降到 2.53s,全程預估從 19 小時縮到 9.4 小時。 |
前三個假設都被實測推翻,而且它們都很有說服力 —— 留在這裡是為了不再重走。真正的轉折是不再靠推論、改看哪一套帳本對不上:當 cuda_reserved 與程序 RSS 同時持平、系統可用記憶體卻持續下降,這個組合只指向一個地方。
| checkpoint | τ 接受長度 | verify rate | 位置 0 | 位置 1 | 位置 2 | 位置 3+ | auc 整體 | auc@0 |
|---|---|---|---|---|---|---|---|---|
| step_210 · 2 層 · 128 anchors · 6.8GB | 1.56 | 0.195 | 0.491 | 0.045 | 0.015 | <0.004 | 0.528 | 0.861 |
| step_70 · 5 層 · 512 anchors · 9.0GB | 1.24 | 0.154 | 0.180 | 0.034 | 0.013 | <0.005 | 0.485 | 0.397 |
τ 1.24 → 1.56(+26%)、位置 0 接受率 0.180 → 0.491(2.7 倍)、信心頭 auc@0 從 0.397(比亂猜還差)到 0.861,而且模型小 25%(9.0GB → 6.8GB)。論文方向的修正成立。
更有力的是:新配置是帶著資料劣勢贏的 —— 它只用了約 7% 的訓練訊號(anchors 512→128 是每樣本訊號的 ¼,樣本又只取 4948 的 28%)。若 anchors 與資料量才是主導因素,新的應該要輸。
提議 7+1 個 token 只接受 1.56 個,而且位置 1 就掉到 0.045、位置 2 之後幾乎歸零 —— 鏈條仍然沒有延續性,τ 撐不起 verify 的成本。信心頭也只在第一步有效(auc@1 = 0.355 比亂猜還差,@3–@5 剛好 0.500 等於零訊號)。
下一步是把資料劣勢拿掉:訓練跑完後記憶體全數釋放(閒置 82GB),而當初 512 anchors 吃約 59GB 撞牆是因為那時只剩 61GB 可用。以 512 anchors × 完整 4948 樣本重訓,才是只差「5 層 vs 2 層」的乾淨 A/B。