Agent Telemetry

V100 推論叢集 · DGX 訓練研究
快照2026-10-02_0254UTC+8 · 每 3 小時更新
尾2 段
2026-08-10

本期紀事

修復
TPFIX2 — 多卡輸出從此正確
此 vLLM fork 的張量並行輸出從來沒對過(9B / 27B 都是亂碼)。逐層 dump 定位到 out_proj 的 rank 切片與各卡持有的 head 集合對不上,修好後 TP2 / TP4 首次產出正確文字。
修復
SPECFIX — CUDA graph capture 缺陷根治
capture 時把 CPU 純量烘進圖裡,導致「快但輸出壞」與「正確但很慢」二選一。改成 verify-as-decode 後兩難消失。
翻案
27B 從判死到可用
單卡 14.5 tok/s(判定不可用)→ TP4 單人 38.4,首次越過 30 的互動門檻。
交付
fable-9b 工具呼叫上線
hermes parser + qwen3 reasoning + 支援 tools 的對話模板。非串流 8/8、串流 7/7 通過;含 tool_calls JSON、結果回灌、think 不外洩。
研究
DSpark 訓練方向修正
比對論文後找出三個方向性錯誤:drafter 容量過大、資料規模不足、confidence 排程未部署。已用輕量 drafter 重啟訓練。
事故
V100「掛掉」三小時 — 但它從未離線
21:35 起完全連不上,一度判定是機器斷電、準備請 Wayne 到辦公室查 PSU。00:54 恢復後查證:uptime 2 天 11 小時<b>從未重開</b>、tailscaled PID 未變、核心無網卡 link 事件、四張 GPU 一路載著模型、生產服務從未中斷。斷的只是它到 Tailscale 控制平面的連線。我在這件事上連錯四次,最嚴重的是拿本機區網的 192.168.1.17 做 ARP 測試當成「網卡未通電」的硬證據 —— 那個 IP 根本不是 V100,是我沒問它在哪就假設了,還讓 Wayne 跟著下錯結論。監視器已改用應用層健康檢查:走 DERP 中繼時 ICMP 與裸 TCP 會假陰性。
解鎖
TP2+PP2 不是不相容,是被自家的除錯訊息擋住
委員會的第二建議是改測 TP2+PP2(TP2→TP4 只有 +19%,all-reduce 已是主導項)。查到 07-17 其實試過並崩潰:'IntermediateTensors' object has no attribute 'shape'。但那是 PP 的正確行為 —— 非最後一個 stage 本來就回傳 IntermediateTensors。崩潰點在 fork 自己加的 _sm70_profile_trace,不在 TurboQuant backend,也不在 vLLM 的 PP 支援(get_pp_group / is_last_rank / IntermediateTensors 處理都完整)。它在任何推論發生前就死了,所以 PP 與 TurboQuant 相不相容從未被驗證到。已修並驗過,等窗口。
治理
生產 vLLM fork 納入版本控管
這個 fork 有 21 個檔案被手工修補過卻完全不在 git,只有零散的 .bak/.pre_ 副本,沒有任何記錄說每個修補改了什麼 —— 生產推論堆疊出事時無人能重建。已建立基準(2444 檔 / 30MB),只納管原始碼:.so 與 .pyc 佔 666MB 中的 579MB 且從不手改。
否定
委員會的頭號優化建議不成立 — 反量化早已融合
建議是「把 2-bit KV 反量化融進 attention kernel,可擠 15-20%」。讀原始碼判定:解碼路徑 _tq_decode_stage1 是 Triton kernel,直接讀打包位元組、在 kernel 內解 2/3-bit 索引、重建 scale/zero、查 centroid、累加到 fp32 —— 從未物化 FP16 的 KV 張量,沒有可融合的東西。物化只發生在分塊 prefill(k_full = torch.empty 後串接餵 SDPA),影響 TTFT 與長 context 首塊,不是穩態每 token 吞吐。省下一個維護窗口。
交付
生產切換 TP4 — 每人 66.6 → 107.5 tok/s
量測鏈在 regen 釋放 GPU 後自動完成 27B TP2 / 9B TP2 量測並切換生產,帶健康檢查與工具呼叫閘門、失敗自動回滾。切換後 8/8 驗收全過(純文字對話、tool_calls JSON、結果回灌、think 不外洩),8 人併發每人仍有 60.3、合計 391.4。
診斷
DSpark 訓練記憶體 — 四層才到底
訓練一啟動就把 121GB 吃光並開始輾轉。前三個歸因(27B 全載 / mmap cache / Adam 狀態)都被實測推翻。真正的原因有三層:孤兒 CUDA 子程序扣住 59GB、注意力 scores 矩陣被實體化、釘選主機快取無限成長。修好後記憶體打平,每步耗時反而從 5.0s 降到 2.53s。
研究
推測解碼 sweep 啟動
單卡旋鈕已測到盡頭(全在 ±3% 雜訊內),TP4 又被 regen 擋著,而 GPU3 整晚閒置。改測三種免 draft 模型的推測方法,針對工具呼叫這種高重複輸出。
事故
PSU 陣亡並更換
多模型同時載入的主板軌尖峰疑似壓垮電供,停機 15 小時。已改為序列載入紀律;四卡與載板無損。
2026-08-10

下一步

  • WayneHermes gateway 端對端測試,確認後啟用工具呼叫
  • agentDSpark 訓練完成 → 以 accepted length 對比舊 checkpoint,裁決容量修正是否正確
  • agent維護窗口:改測 TP2+PP2 拓樸。TP2→TP4 只有 +19%,all-reduce 已是主導項,而單 token 解碼下 PP 的點對點延遲通常低於 TP 的 all-reduce
  • agent推測解碼:工具已重建(spec_isolate2.sh),等一張空卡即可跑,並看實際回應內容
  • agent27B TP4 補量(需先讓出 GPU0),確認 38.4 tok/s 是否仍成立
  • agentregen 8931 筆已收齊 → 資料擴量重訓

未分類的舊段落

這些是還沒標上分類的早期段落,內容有效,只是沒歸位。

TPS 總表 · 每人 decode

每一列都是量過的,沒有推算值。併發欄是同時的請求數,tok/s 是每人而非合計。「參照」列是別的引擎或別台機器,只用來對照,不是我們的產出。

模型 · 量化拓樸context併發tok/s時期註記
9B · Q4_K_MTP4~200 tok1152.8現行現行生產
9B · Q4_K_MTP4~200 tok876.8現行每人
27B · IQ4_XSTP4~200 tok168.2現行64K 配置
27B · IQ4_XSTP4~200 tok820.0現行每人 · 合計 160
27B · IQ4_XSTP4~8K tok817.7現行每人
27B · IQ4_XSTP4~28K tok814.1現行每人
27B · IQ4_XSTP4105,856 tok137.7現行128K 配置
27B · IQ4_XSTP4105,856 tok87.0現行每人 · 端到端只有 3.4
27B · Q4_0TP4~200 tok175.1現行requant · 舊版權重
27B · Q4_0TP4~200 tok834.5現行每人 · 合計 276 · 唯一過 30 的
27B · IQ4_XSTP4~200 tok130.2現行開 MTP spec-1 · 淨損
27B · Q4_K_STP4~200 tok2433.6現行每人 · 合計 241 · util 0.75 · 不開 MTP
27B · Q4_K_STP4~200 tok3233.7現行每人 · 合計 256 · c8 到 c32 是平的
27B · IQ4_XS + 批次 MMVQTP4混合 · 100–400 out169.1現行<b>不開 MTP</b> · 異質 p50 · 這是目前最佳單人
27B · IQ4_XS + 批次 MMVQTP4混合 · 100–400 out448.5現行不開 MTP · 異質 p50 · 合計 108 · 改前 36.74
27B · IQ4_XS + 批次 MMVQTP4混合 · 100–400 out832.3現行不開 MTP · 異質 p50 · <b>合計 143</b> · 改前 20.58
27B · IQ4_XSTP4混合 · 100–400 out820.6現行不開 MTP · 未改 kernel · 這是被修掉的那條
27B · IQ4_XS + MTP1 + 新 kernelTP4~200 tok815.1現行每人 · 合計 103 · capture 對齊 ＋ <code>VLLM_GGUF_IQ4XS_MMQ=64</code>
27B · IQ4_XS + MTP1 + 新 kernelTP4~200 tok1612.5現行每人 · 合計 188
27B · IQ4_XS + MTP1 + 新 kernelTP4~200 tok329.0現行每人 · <b>合計 266</b> · 同構上界
27B · IQ4_XS + MTP1 + 新 kernelTP4混合 · 100–400 out815.1現行<b>異質</b> p50 · p95 18.08 · 合計 61.6 · TTFT p50 0.51s
27B · IQ4_XS + MTP1 + 新 kernelTP4混合 · 100–400 out329.2現行<b>異質</b> p50 · p95 10.05 · <b>合計 136.7</b> · TTFT p50 1.80s · 32/32 完成
27B · IQ4_XS + MTP1TP4~200 tok248.2現行capture 對齊 <code>(1+draft)×席次</code> · 合計 181 · 未開新 kernel
27B · IQ4_XS + MTP1TP4~200 tok327.6現行同上 · 合計 225
27B · IQ4_XS + MTP1TP4~200 tok243.8現行capture 未對齊 · 掉出 CUDA graph 走 eager · 這是被修掉的缺陷
27B · Q4_K_MTP4105,856 tok135.1GDN 開8/7 舊量測
27B · Q4_K_MTP4105,856 tok88.1GDN 開8/7 舊量測
9B · Q4_K_MTP4~200 tok1123.1GDN 開8/6 之前
9B · Q4_K_MTP2~200 tok196.2GDN 開
9B · Q4_K_M單卡~200 tok169.7GDN 開
9B · Q4_K_MTP4~28K tok1106.2GDN 開
9B · Q4_K_MTP4~28K tok845.9GDN 開每人
27B · Q4_0TP4~200 tok160.4GDN 開
27B · Q4_0TP4~200 tok830.9GDN 開每人

屋頂線: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_XSllama.cpp · 單卡~200 tok139.24不開推測
27B · IQ4_XSllama.cpp · 四卡 layer~200 tok138.46等於單卡,管線切分買容量不買速度
27B · IQ4_XSllama.cpp · 四卡 row~200 tok122.12pp512 只有 88.99,比單卡慢十倍,病態
27B · IQ4_XSllama.cpp · 四卡 tensor~200 tok133.61pp512 2333 對單卡 855,提示處理 2.7 倍
27B · IQ4_XSllama.cpp · 單卡 + MTP n=3~200 tok155.90MTP 在 llama.cpp 會賺,在 vLLM 不會
27B · IQ4_XSllama.cpp · 單卡24,704 tok133.40f16 KV
27B · IQ4_XSllama.cpp · 單卡24,704 tok125.50tbq3 KV,反而更慢
27B · AutoRound INT4外部 · vLLM · 3090 單卡125K tok185.00別人的機器,不是我們的成績

llama.cpp 四卡的三種模式都驗過每張卡真的吃到權重(tensor 是 4519 / 4313 / 4313 / 4313 MiB),不是只看旗標有沒有被接受 —— 第一次量的時候 -ts 用了逗號,llama-bench 把它讀成「掃多組設定」,四個臂其實都跑在 GPU 0,整批數字作廢重跑。

27B 完整矩陣 · 64K 配置

配置

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/sprefix 命中
~200c10.27s0.27s68.243.80%
~200c40.27s0.82s32.866.90%
~200c80.27s0.69s20.0101.10%
~8Kc14.19s0.96s63.925.385%
~8Kc44.19s4.10s29.657.785%
~8Kc84.19s6.62s17.760.185%
~28Kc125.33s1.79s55.340.297%
~28Kc425.33s6.68s25.147.897%
~28Kc825.33s12.90s14.158.997%

27B 完整矩陣 · 128K 配置

配置

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
~200c10.26s0.26s67.945.2
~200c80.26s0.69s20.098.1
~105Kc1209.08s16.35s37.72.1
~105Kc8209.08s130.03s7.03.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%。

批次縮放 · 為什麼 8 併發掉這麼多

配置

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 步時間
c168.214.7 ms75.113.3 ms
c247.621.0 ms51.619.3 ms
c432.630.7 ms34.628.4 ms
c820.050.0 ms34.528.5 ms
i-quant 每多一個序列固定加 4.83 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 仍然走它。

MTP 在 27B 上是淨損

配置

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.168.4
spec-tokens 129.730.2
spec-tokens 221.125.7

內容類型沒有救到它 —— 工具呼叫那格和散文一樣爛。llama.cpp 上 85–99% 的接受率在 vLLM 這裡不轉換成速度,因為瓶頸是推測步本身的成本。另外 temp 0 之下三個臂的生成長度不同(272 / 271 / 201),推測路徑會改變輸出,那比速度更嚴重。

MTP 接受率的根因 · KV 位元寬 已推翻 8/12

這節的結論是錯的

後來的六臂掃描顯示單人接受率從 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 080.6%97.6%97%
position 144.4%90.4%95%
position 216.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 倍加速。接受率修好了,推測機制本身還是虧的 —— 這兩件現在乾淨分開了。

KV 是否分散到四卡 · 以及 TurboQuant 換到多少併發

配置

27B IQ4_XS · max-model-len 65536 · util 0.75 · block-size 32 · 只讀引擎啟動時報的 pool,不做生成 · 此模型 head_count_kv=4,vLLM 按 head 切,TP4 是分散上限

拓樸KV pool64K 下併發倍數每卡記憶體
TP1370,085 tok5.65x20.8 GB(只有卡 0)
TP22,174,253 tok33.18x23.0 / 22.8
TP45,751,747 tok87.76x23.8 × 4

KV 確實分散,而且是超線性:TP1 → TP4 是 15.5 倍不是 4 倍,因為權重也跟著切,每張卡騰出更多空間給 KV。llama.cpp 四卡 layer 也分散(5439 / 5027 / 5233 / 6049 MiB,單卡是 19177),它按層切,不受 head 數限制。

KV 格式 · TP4KV pool64K 下併發倍數
turboquant_2bit_nc5,751,74787.76x
turboquant_3bit_nc4,237,04464.65x
turboquant_4bit_nc3,480,42953.11x
f16(auto)1,088,04616.60x

TurboQuant 2-bit 給的併發是 f16 的 5.3 倍。這就是為什麼 KV 格式必須寫進每一份配置 —— 它直接決定能放多少人。

新權重 + 8 併發 30 以上 · Q4_K_S

配置

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 · 三個臂都是現行的新合併版,只差量化型別

量化c1c2c4c8批次增益
Q4_K_S68.852.233.934.44.00x
Q4_K_M66.446.933.331.13.75x
IQ4_XS68.447.832.119.92.33x
Q4_0(舊版權重的 requant)75.151.634.634.53.74x
i-quant 在這個 build 裡沒有批次 kernel

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 × 序列數 的來源。

生產修好一個懸崖 · 一行設定換 4.2 倍

配置

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](現行)
c1152.65152.4
c876.3 / 合計 60776.3 / 合計 556
c1614.5 / 合計 23261.0 / 合計 899
c24—40.4 / 合計 897
c32—39.6 / 合計 1188
排程器放 32 個人進來,但只錄到 8 卷帶子

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。調了上限就要量到上限。

自己寫融合 kernel · 五版之後的否定結論

配置

權重 17408 × 5120(此模型 ffn 形狀)· Q4_0 · sm_70 · cuda event 計時 30 次取平均 · 每版都先過正確性再看時間

版本改了什麼M=16
v1一 warp、逐位元組拆 nibble1.19
v2四 warp、N_TILE 64(委員會的分塊)0.63
v3全執行緒參與、K_TILE 128、half20.51
v4位元技巧解量化(免轉換指令)0.44
v5M_TILE 64、多累加器0.72
dp4a(要打敗的)—0.331
屋頂線50.2 MB ÷ 900 GB/s0.056
走不通,而且是 API 層級的

五版全部正確(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=1B=2B=4B=8
IQ4_XS · 單卡38.7963.0086.24110.55
Q4_K_S · 單卡34.4556.5868.9672.97
IQ4_XS · 四卡 layer38.2762.3988.57111.75
IQ4_XS · 四卡 tensor———241.96
Q4_K_S · 四卡 tensor———209.58
更正一:llama.cpp 四卡在高併發下是會縮放的

我寫過「四卡對 llama.cpp 的解碼沒有幫助」。那是 batch 1 量的。八個序列同時在跑時,tensor 是單卡的 2.2 倍(241.96 對 110.55),提示處理更是 1828 對 690。layer 確實不縮放(111.75 對 110.55),那部分成立。

更正二:i-quant 的批次缺陷是我們 fork 特有的

在我們的 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)。兩個都在,而我只拿前者當標竿。

引擎之爭的收束 · 量過之後留在 vLLM

配置

同一顆 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 每人合計
c166.5357.3080.0%69.6—
c430.57110.0780.0%34.1136
c817.24123.0580.0%34.3275
MTP 在 llama.cpp 是真的有效,但不夠

單人從 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」的理由只是生產在上面,那不是技術理由。

自動產生 · 自有硬體推論與訓練研究狀態歷史快照 581 份 →2026-10-02_0254