Agent Telemetry

V100 推論叢集 · DGX 訓練研究
快照2026-08-17_1728UTC+8 · 每 3 小時更新
現況 · 2026-08-17
線現在的數字狀態
生產 · Qwen3.8-27B IQ4_XS · TP4
100.78.29.29:8601 · fable-27b · 設計上限八個併發使用者		單人 68.1 · 八人合計 260(每人 32.6,同時產出時)
混合負載、錯開到達:合計 155 · 八人各持 21.6K 文件:55 秒全部完成		上線中 · 工具閘門 8/8
一小時 soak 過:0 失敗 · 記憶體零增長
換掉 Codex 部署的 Q4_K_Mc4 合計 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 那次變更之前量的 —— 當時是對的,但不代表現在的生產。被推翻的結論一律保留並標明,不刪 —— 看得到我錯在哪比看不到有用。

最新 · 2026-08-17

否定把權重解碼提出 NCOLS 迴圈:程式碼是對的,速度是零。九形狀正確性全過、誤差與基準逐格相同且非零、兩臂都確認 BMMVQ 4/4 真的掛上。結果 c8 全員串流 260.15 → 261.90(+0.7%)、c4 合計 107.40 → 107.28、c1 持平 —— 全在雜訊裡(同組態不同輪已量到 8% 變異)。「權重在迴圈裡被重讀八次」描述的是源碼的形狀,不是時間花掉的地方 —— L1 本來就吸收掉了。
我自己的前提就錯了我派這個任務時說「效率只有 28%,有 72% 可拿」—— 那個數字用的正是被稀釋的 155。用正確的 260 重算:理想 8×68.1 = 545、實際 260、效率 48% 不是 28%,剩餘空間是 2.1 倍不是 3.5 倍。而 68 → 260 表示每次讀權重現在服務 3.8 個 token,攤提確實在發生。我拿一個被高估的缺口去派工,委員會照著找,找到一個源碼上成立、執行上不成立的解釋。四輪派工的根本問題不是附件也不是引號,是問題本身不準。
派工四輪的失敗都在我① 同時要診斷和 patch,預算被診斷吃光;② 叫它「不確定就自己寫」,它去反推位元佈局;③ 附的參考 header 來自另一棵 llama.cpp,vec_dot 三參數變四參數、__dp4a 變 ggml_cuda_dp4a,編不過;④ 提示裡的雙引號被 PowerShell 吃掉,argparse 拒收。委員會的診斷從第一輪就沒變過。
擋住了驗收鏈按設計運作:編譯錯誤擋一次、缺 pybind 擋一次、每次自動還原,生產全程沒被碰過。而且每臂都印 BMMVQ x/4 —— kernel 沒掛上時伺服器照樣起得來、照樣有數字,那個數字看起來會像「改動沒效果」而不是「改動沒生效」。
更正我把八人的容量少報了 40%。Wayne 問「一個人 35,八個人怎麼只有 155」,那個問題直接指出錯在哪。兩個數字的分母不同:每人的 decode 是「首末 token 之間」(不含 TTFT、不含沒在跑的時間),aggregate 是「全部 token ÷ 整段牆鐘」。而量測工具刻意把到達時間散在兩秒、輸出長度隨機取 100–400,所以牆鐘由最長的請求決定、token 總數卻不是 —— 後半段只剩兩三個人在跑。
量到八個人同時產出時,合計是 260 不是 155。改成同長度、同時放行、只取全員都在串流的視窗:c1 68.08、c4 每人 46.51 合計 186.2、c8 每人 32.61 合計 260.2 —— 和 8 × 32.61 = 260.87 幾乎完全相等,樸素相乘本來就是對的。兩個數字都要留但要標清楚:260 是機器容量,155 是混合負載下的實際吞吐,前者答「撐得住多少」,後者答「跑起來像什麼」。
第五次整天的 SSE 全滅終於查清楚:串流的欄位叫 reasoning,非串流叫 reasoning_content —— 同一個東西兩個名字。delta.content 一直存在但是空字串,所以只試那兩個非串流名稱的檢查會安靜地數到 0 個 token。把第一 KB 用 repr() 印出來花十秒,而我先前猜了兩次格式。
驗收一小時八人 soak 全過,而且記憶體一格沒動。12 個視窗,解碼 p50 32.36–32.60(變異 0.74%)、合計 152.9–167.2、TTFT 0.91–1.20s、0 次失敗、GPU 記憶體 17,162 MiB 從第一分鐘到第六十分鐘完全不變。四種失敗模式(速率漂移 / 記憶體增長 / 尾端請求開始壞 / kernel 脫落)全數通過。這個部署可以交出去了。
量到席次砍到 8 沒有任何好處,而且我的理由是錯的。我以為砍掉 56 個席次和三組 CUDA graph 能還記憶體給 KV pool —— KV pool 兩者都是 2,170,595 token,一個 byte 都沒差。席次不從 KV pool 扣:vLLM 照 util 先切好 KV,席次只管排程上限。c8 合計 158.4 對 154.7、長文件 60.2s 對 61.6s,都在雜訊內(同組態不同輪已量到 8% 變異)。排程類設定至此已在兩種量化上各掃一輪,結論一致:全部無效。
能力Qwen3.8 的視覺是可用的,六項全中。448×448 四色象限加中央黑方塊,它答出左上紅、右上綠、左下藍、右下黃,以及中央黑方塊「跨在四個象限交界上」。兩個前提:只能走 llama.cpp(vLLM 的 GGUF 路徑不吃 mmproj),以及必須加 --image-min-tokens 1024 —— 不加的話 96×96 的圖只換到 81 個 prompt token,加了之後是 1,103。
更早的 81 條 —— 含已撤回的結論,保留備查
量到八個人各拿一份 21.6K 文件:55 秒全部拿到答案,而共用文件沒有任何好處。八份不同文件牆鐘 54.8s、八份逐位元組相同(且先暖過)55.0s —— 差 0.2 秒。3.6 上量到的「不共用 prefill」在 3.8 加了 bt 65536 之後仍然成立。合計吞吐 18.7 tok/s,對短提示八人的 155 —— 長 context 併發是八倍代價,時間幾乎全在 prefill。
變好了但排隊變成並行了。3.6 那次八人問同一份文件是「一人 31 秒、最後一個等 245 秒」;現在八個人一起在 55 秒完成,延遲差距只有 0.3 秒。對使用者體驗來說,「大家一起等 55 秒」比「有人 31 秒有人 245 秒」好得多。(兩次文件大小不同,不能比秒數;可比的是排隊 vs 並行。)
第三次了視覺第一輪回報空答案,而模型其實成功了 —— 日誌寫著 loaded multimodal model、生成 120 個 token。我讀 content 讀到空,內容在 reasoning_content。本 session 第三次讀錯欄位,而我還為這件事寫過記憶。
量到解碼隨 context 掉得很溫和:提示放大 250 倍,速率只掉 37%。258 tok 67.52 → 3,618 65.55 → 9,618 62.24 → 21,618 56.36 → 43,218 48.46 → 64,818 42.73。而 64.8K 的 42.73 和 Codex 在 65,536 量到的 42.8 幾乎完全吻合 —— 兩個人、兩套探針、兩種量法,獨立得到同一個數字。
量到冷 64K TTFT 90.71 秒,對 Codex 的 162.6 秒 —— −44%,是 max-num-batched-tokens 65536 換來的。
自己的坑踩第二次那張表的「冷 TTFT」欄只有最後一列可信。六個長度是巢狀前綴(同一段文字重複到不同長度),所以量到 43K 時前面 21K 早在快取裡 —— 破綻是 43,218 的冷 TTFT(4.63s)比 21,618 的(6.24s)還低。這個陷阱我 8/7 就寫進本頁過(「這一輪的冷數字不能用,我先講」),今天又踩。差別是這次在讀數字時抓到,沒拿去下結論。暖 TTFT 與解碼那兩欄不受影響。
量到併發到 64 仍然沒有轉折點。合計 c32 247 → c48 315 → c64 357,每人 15.3 / 11.4 / 10.0,TTFT 都在 6 秒內。KV pool 2,170,595 token,64 人各分 3.4 萬,容量還沒到極限。這台四卡 V100 同時服務 64 個人是可用的。(注意 c32 兩輪量到 266.7 與 246.7,capture 集合不同差 8% —— 單次差異別當成改善。)
上線Qwen3.8-27B 換成 IQ4_XS 之後併發吞吐 +24%,而且一行 code 都沒寫。Codex 部署的是 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 檔原封不動。
量到到 32 併發還沒有天花板。合計 c1 68 → c8 155 → c16 186.7 → c24 217.7 → c32 266.7,每人仍有 14.63 tok/s、TTFT 1.64s。三十二個人同時用是可用的。正在往 48 / 64 推。
撤回「併發設定對齊能補回 25%」—— 錯的。max-num-seqs(16 vs 8)、capture 集合、util(0.60 vs 0.50)四種組合,c4 合計全部落在 86.5–86.9,c8 全部 133.0–133.6。那 25% 從頭到尾都是量化檔,不是排程設定。
澄清Codex 的 TPS 難看是量法不是部署。他用 65,536 token 的提示量端到端,所以 17.4 / 10.17 / 5.61 裡絕大部分是 prefill(TTFT 162.6 秒)。拿他正在跑的那個服務改用短提示是 69.16,和我調了一整晚的 Qwen3.6(68.88)一模一樣。他的解碼一直都是好的。
有效--max-num-batched-tokens 16384 → 65536:21.6K 冷 prefill 19.8s → 15.4s(−22%),decode 不動。和我先前在 106K 上量到的 −26.3% 同一個量級,已套進生產。
撤回「MTP 接受率低是因為 KV 位元寬」—— 錯的,我自己的六臂實驗把它推翻了。同一顆權重、同一組 kernel,只掃 KV 精度,單人接受率是 2-bit 51.5% / 3-bit 50.0% / 4-bit 52.3% / k8v4 51.5% / f16 51.5% —— 完全不動。連完全不壓的 f16 都一樣,所以 llama.cpp 的 72.4% 不是 KV 換來的。先前那個 47.0→87.1% 的觀察是單一提示的單次量測,沒有重複、沒有跨提示,被我當成因果寫了上來。四併發下 KV 精度確實有效(2-bit 50.3% → 4-bit 62.4% 後打平),但吞吐紋風不動(24.85 → 27.10),因為綁住的是 c 不是 α。
為什麼量不到七次插樁全部失敗,原因是結構性的,不是手滑。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 不等於伺服器起得來,一定要重跑一臂。
結案Muse 這條線到此為止,不再測。它的任務是當對照組,而它做到了 ——遞迴假設已經被殺掉,再測下去是換模型的決策,不是優化決策。權重和自建的 llama.cpp 留在機器上(/mnt/data/backups/windows-v100/models_muse、/mnt/data/llama.cpp-new),要用隨時能起。
量到vLLM 單卡也是 1.238,拓樸只佔超額的 27%。TP1 base 33.52 → MTP 22.69(α 51.5%)→ 步 66.77 ms 對基準 29.83 ms。超額的 1.305 裡 0.948(73%)是引擎、0.357(27%)是 TP。架構、拓樸、草稿機制三個變數全部排除完,病灶在 vLLM 的推測路徑本身。
天花板就算把 c 修到 llama.cpp 的 0.29,配現在的 α=0.52,S = 1.52/1.29 = 1.18 → 81 tok/s,比基準 68.88 好 18%,而且只在單人。四併發合計 34.80 對基準 108.80(0.32 倍)。這條線的上限現在算得出來了。
推翻主假設遞迴狀態不是驗證變貴的原因 —— 對照組把我整晚的前提殺掉了。同一個 llama.cpp、同一張卡:Qwen3.6-27B(65 層裡 48 個帶 GDN)每多驗一格付 0.298 個基準步;Muse Glimmer 30B(52 層全注意力、零遞迴)付 0.287 —— 差 4%。而 vLLM 在同一顆 Qwen3.6 上每格要 1.595,5.4 倍,而且和架構無關。我一直把「48/65 帶遞迴」當成 c=2.01 的根因,給 fusion 委員會的知識庫也是那樣寫的,委員會的三個候選全部建立在這個錯誤前提上。這讓 spec_sequence_masks is None 那個守衛從「候選之一」變成主嫌 —— 它正好是一條純引擎側、與模型架構無關的路徑選擇。(限制:llama.cpp 兩顆都是單卡、vLLM 是 TP4,跨了引擎也跨了拓樸;但 llama.cpp 內部那組同引擎同卡數的比較是乾淨的,足以殺掉遞迴假設。)
量到Muse Glimmer + 官方 DFlash 草稿器:1.12 倍,公式第三次命中。base 32.63 → DFlash 預設 36.55 tok/s;每步 2.069 個 token ÷ (1+0.847) = 1.120,實測 36.55/32.63 = 1.120。加大區塊一樣是壞事(和 MTP 的 γ 同形狀):n-max 8 → 31.55(接受率 20.4%)、n-max 16 → 31.05(12.2%,草稿量從 428 暴增到 1565)。日誌:n_max=16 exceeds the trained block size 16 -- clamping to 15。預設最好。
又一次300 個 token 全在 reasoning_content 而不是 content(打到 finish_reason: length 時還在推理段),所以我的文本檢查印了 uniq_word_ratio 0.00 DEGENERATE —— 吞吐數字是有效的,壞的是我讀錯欄位。實際內容 1585 字元、連貫英文。
作廢Muse 的前兩輪數字全部作廢,兩次都是我的錯,而且第二次騙過了我自己寫的檢查。第一輪:--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、不准印速率。
學到DFlash 不是一顆自迴歸小模型,是「遮罩 token 的區塊平行草稿器」。伺服器 log: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。
注意c1 的 TTFT 這輪是 1.78s,掃描那輪是 0.68s —— 同一組態、不同輪次差 2.6 倍。prefill 的量測變異比我以為的大,那個 0.68s 不能當定值引用。
建好了新的 llama.cpp 建起來了(commit 8e7f22b、gcc 12.4.0、CUDA arch 70、ninja -j2),編譯日誌裡有 common_speculative_impl_draft_dflash —— DFlash 推測支援確實在裡面。/opt/llama.cpp 覆核過完好,對照組數字仍可比。
修好「解碼的 15.5% 是拿 prefill 換的」那個代價有解:兩個旗標一起開。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,不是解碼主力。
新玩具Meta Muse Glimmer 30B 已下載並驗過(23 GB,三顆檔案的 byte 數與遠端 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。
量到八個推測路徑旗標,只有兩個動得了 c,而且落在同一個值。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。排序要靠量,不靠命名。
找到推測一開,GDN 的快速解碼路徑就結構性地拿不到 —— 這寫在守衛裡,不是推論。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 為真,也會關掉這條快速路徑。
委員會給對的TP4 下 torch profiler 為什麼永遠是空的:worker 是獨立 process,start_profile 只武裝 rank 0 的 context,跨不過 process 邊界。這解釋了我三次失敗,替代方案是 nsys --output=…%p 每個 PID 各自出檔。委員會的價值在候選假設與這條路,不在優先順序。(Codex 這次 thread start failed 無實質產出,融合時整份排除,實際是 GLM-5.2 與 MiniMax-M3.0 兩份。)
換路要看清那 29.2 ms 花在哪,我連三次拿不到 profile,停手換方法。第一次沒設 VLLM_TORCH_PROFILER_DIR;第二次設了但讀取端掃的是自己的預設目錄;第三次目錄對上了,TP4 下 /start_profile 之後那個目錄仍然是空的,fork 自己的 phase report 也沒進 log。三次失敗長同一個樣子 = 問題在假設不在實作(RULES 5),所以不再想辦法「看進步驟裡面」,改直接量哪一種實作最便宜 —— 這個 fork 帶了八個切換推測路徑的旗標,一臂一個,c 用同一 session 自己的無推測臂算,不從記憶拿。
委員會把這題交給 fusion 委員會(codex / glm / m3,附完整 KB)。問題陳述現在夠乾淨可以外包了:公式在兩台引擎上獨立驗證過、十二個假設各有具體死因、fork 旗標清單已列。該召開的時機我錯過了兩次 —— KV/γ/p-min 連三輪同一目標失敗那次,以及 c=2.04 掛在報告上很久卻沒人從架構角度問「1/65 的工作量為什麼花 2 個基準步」那次。
定案兩台引擎的帳算平了,而且 S = (1+α)/(1+c) 對兩邊都命中到三位數。llama.cpp 單卡:基準 39.16、α 0.722、c 0.298 → 預測 1.327、實測 1.327。vLLM TP4:基準 68.86、α 0.515、c 2.010 → 預測 0.503、實測 0.503。四個獨立量到的數字各自命中,這個模型不是事後湊的。
價碼把 llama.cpp 的 α 和 c 搬到 vLLM 的 68.86 基準上 = 91.4 tok/s。拆開看:只修 c(2.01→0.30)即使 α 停在 0.515 也有 80.4;只修 α(0.515→0.722)而 c 停在 2.01 只有 39.4 —— 連損益平衡都到不了。所以 c 是壓倒性的槓桿,我原本要先追的草稿頭餵法要往後排。
撤回「llama.cpp 的 72.4% 是 p-min 篩出來的」—— 也錯了,是我下一個被自己數據打掉的假設。把 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。)
線索vLLM 餵給草稿頭的是過完最終 norm 的 trunk 狀態: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 後面。
修好把 GDN 推測狀態合約裡的 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 才開。
量到基準確認 68.86(capture 含 1),和先前的 69.10 對得上 —— KV 掃描那個 59.25 確實是 capture 集合的假象,不是退化。
剩下的病灶草稿頭是 65 層裡的 1 層,它多做的那次 forward 該花基準步的 1/65。實測花掉 2 個基準步 —— 差 130 倍。這不是算力,是管線。今天證明其中 0.21 是同步造成的(c 2.22→2.01),剩下 2.01 還在。先前的 forward 拆解指向 elementwise 佔 42%,那是 48 個遞迴 block 的狀態搬運與回捲,下一刀在那。
失誤llama.cpp 那三臂第一次跑全部 NO NUMBERS —— 我用了 -ctk turbo2,但 /opt/llama.cpp 是上游 b0.15.3 不是 atomic-llama fork,turbo2 是 fork 專有型別。旗標檢查有做、KV 型別沒檢查。改 f16 重跑中。
量到加長草稿只有壞處,γ 不是我們少開的那個開關。vLLM γ=1/2/3 的 position 0 接受率是 50.4 / 44.8 / 43.5% —— 第一個位置不隨 γ 改變,而總接受率被後面的位置一路拉低到 51.5 → 31.1 → 23.7%,吞吐跟著 33.21 → 30.81 → 27.86。逐位置條件接受率 pos1 33.1% / pos2 37.8%,離對照組的 95 / 91% 差一個量級。
下一刀所以 llama.cpp 的 72.4% 更可疑了:同一顆權重,vLLM 在 γ=3 只有 20.6%。只剩兩種解釋,而且可以一刀切開 —— ① p-min 0.5 讓它信心不足時根本不提草稿,那些注定被拒的位置沒進 draft_n,百分比是在篩過的子集上量的;② vLLM 真的把草稿頭餵錯了。把 llama.cpp 的 p-min 降到 0 重量,掉到 ~50% 就是①(那接受率從頭到尾都不是差別,差別是 c),留在 ~72% 就是②(那是個值得修的 bug)。跑了。
再更正我先前寫「MTP 在混合遞迴架構上必輸,建議收掉」—— 那個結論下太早,撤回。一份獨立實測在 同一個架構的 Qwen3.6-27B 上,MTP γ=3 拿到 2.40 倍。回捲稅是真的,但它把 c 從 0.20 推到 0.30 —— 足以把小贏變小輸,離我們的 c = 2.04 差得很遠。那個差距底下還有東西,還沒找到。另外那篇說 Qwen3.6-27B 是「純 dense、沒有 GDN」也是錯的 —— 本機七顆變形(含 OFFICIAL 那顆)metadata 全是 qwen35、65 blocks、48 個帶 SSM、ssm.state_size 128。看檔名不算,要看 metadata。
回捲稅有解混合模型的成本有名字也有 PR:草稿被拒絕時,舊版引擎會完整回捲到最近的檢查點(因為 GDN 遞迴被融進單一 kernel),業界叫它 gated DeltaNet rewind tax。vLLM PR 22400 改成 ring buffer 保留每 token 中間狀態,只回捲被拒絕的後綴,把懲罰從 O(γ) 降到 O(γ−k)。我們 fork 裡有 spec_state_slot_selectors 這類機制,但沒找到 ring buffer —— 是不是半套還沒確認。
病因開 MTP 額外的成本來自這顆模型有 48 層帶遞迴狀態。推測解碼的整個前提是「多驗一個位置幾乎免費」—— 對 attention 成立(多寫一格 KV,拒絕就丟掉),對 linear attention / GDN 不成立:它只保留最後一個 token 的狀態。vLLM 自己的 docstring 寫著,有草稿時每個位置的狀態都得留著,等驗證結果出來再依接受數平移 48 層的狀態,而那組緩衝區只在「推測解碼 ＋ 混合模型」時才配置。實際比例:65 個 block 裡 48 個帶 SSM 狀態、17 個純 attention。trace 裡那些在不開 MTP 時前 25 名根本看不到的 kernel —— memcpy32_post 1879 次、scatter_gather 268、indexSelect 268、CatArrayBatchedCopy 268 —— 全是狀態簿記,不是算術。論文的 1.8x / 3.4x 全部量在純 transformer 上。
壓測過了27B 一小時、四併發、混合提示:12 個窗口速率 46.64–47.04,變動不到 1%,零失敗。記憶體 15688 → 16678 MiB 全發生在前 15 分鐘、之後九個窗口完全持平(配置器安頓,不是洩漏);一小時負載後工具閘門仍 ALL_PASS。
最大一筆批次化 MMVQ:c8 每人 20.58 → 32.31(+57%)、c4 36.74 → 48.47(+32%)、c2 48.52 → 57.27(+18%),c8 合計吞吐 104.7 → 143.3。fork 的 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 一致到印出的精度。
定案MTP 在這台機器上不可能贏,而且原因不是接受率。Leviathan/Kalman/Matias(ICML 2023)的加速式是 (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 最大。
根因c = 2.04 不是草稿頭的成本。DeepSeek-V3 報告說 MTP 是一層 transformer、重用主幹的 hidden state,固有 c ≈ 0.015,正好在論文範圍內。我們的 c 被灌大,是因為論文 Section 3.3 的前提 —— 「同時算 γ+1 個位置不增加牆鐘時間」—— 在我們機器上是假的:M=2 前向要 1.85 倍。而那又是因為 M=1 時矩陣乘只佔 8.4 ms、非矩陣乘佔 9.0 ms(52%),一次前向 2,297 顆 kernel —— decode 根本不是權重頻寬受限。同一個病灶同時壓住單人 69 tok/s 和 MTP。
撤回我先前說「草稿頭被壓成 4-bit 是接受率低的原因」—— 錯的。實測 eh_proj 是 Q8_0 的那顆模型,接受率更低(39.4% 對 47.8%),而 S 兩者都落在 0.45–0.49。也一併撤回「drafter 的 CUDA graph 是瓶頸」(全關掉只差 0.5%)與「decode 是 launch-bound」(GPU 其實忙到 72–76%,我先前用錯了分母)。
結案 2「27B 過不了工具呼叫閘門」是假的,而且是閘門的錯。閘門只用 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。
根因27B 開 MTP 之後在 16→24 併發之間掉懸崖,是 capture 集合對不上。推測解碼每步送的是 席次 ×(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 張。
修了capture 改成 [1,2,4,8,16,32,48,64] 後捕捉數 7/7,c24 每人 4.28 → 8.17(合計 97 → 181)、c32 4.27 → 7.60(合計 128 → 225)。
結案三臂同場對照:c8 三臂完全一樣(9.71 / 9.89 / 9.89),所以先前那個「c8 從 17.97 掉到 9.82」不是 capture 造成的 —— 17.97 在同 session 裡重現不出來,那是跨 session 的漂移。t3_tool_call 也是三臂全失敗含舊集合,同樣與 capture 無關。所以 capture 對齊是乾淨的淨勝:c24 3.76 → 8.16(2.17x)、c32 3.66 → 7.59(2.07x),c1/c8/c16 在雜訊內不動。
留著但那代表同一組態在不同 session 之間,c8 可以差兩倍(17.97 對 9.71,都是舊 capture)。原因未知。凡是要拿來下結論的併發數字,都得同 session 取。
證實ggml_mul_mat_a8 對 IQ4_XS 完全沒有運算。profiler 顯示它只發射 quantize_q8_1,之後沒有任何 matmul,輸出緩衝區沒被寫入。它先前之所以通過數值比對,是 torch 的 caching allocator 把剛算完參考值的那塊記憶體回收再配給它 —— 讀到的是上一次的殘值。因此把 IQ4_XS 加進 MMQ_QUANT_TYPES 來修批次崩塌,會讓模型安靜地吐垃圾而且不報錯。
量到IQ4_XS 批次路徑的實際成本(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,偏保守。
做出來了IQ4_XS 的批次 kernel 寫好、驗過、接上了。端到端(同 session、只切一個環境變數、MTP draft=1):c8 每人 9.72 → 14.71(+51%)、c16 9.58 → 12.54(+31%)、c24 8.15 → 9.85(+21%)、c32 7.60 → 9.00(+18%);合計 c32 225 → 266。兩臂產出的文字逐字相同。
關鍵第一版全錯(相對誤差 ~1.0),因為我沿用了標頭的 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。
沒解釋掉開 kernel 之後單人 c1 少 1.5%(29.97–30.01 對 30.45–30.47,三輪六次讀數一致),而 c1 的 M=2 根本沒進 kernel 的窗口。我猜是 dispatch 裡的 os.getenv 每次矩陣乘都跑,提到模組層級後只移動 0.3% —— 猜錯了。而且用算的就排除得掉:條件短路後開關只差兩次整數比較(約 25 µs),實際差距是每步 0.6 ms。成因在別處,還沒找到。操作結論:只有單人就別開,有併發就開。
改了驗法正確性門檻本來是拍腦袋的絕對值,在 M=1 給過一次假警報 —— 同一顆張量兩次隨機輸入是 9e-4 和 3.8e-2,差別在分母是單列最大值。改成拿 mmvq 當標尺問「比要取代的那條路徑差嗎」、固定亂數種子:九種權重形狀 × 五種 M,誤差跟 mmvq 一致到印出的精度。
找到fork 裡早就有手寫的批次 kernel(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.cpp 的 KV dtype 受控組。llama-bench 的 -ctk/-ctv 是位置配對,先前那批其實只量到 V —— 重跑後:V 量化花 8% 生成速度,K 量化再花約 30%(K 在 QKT 內圈)。
測了llama.cpp 四卡三種 split-mode:layer 38.9(等於單卡)、row 22.0(慢 43%)、tensor NCCL 崩潰。llama.cpp 在這台沒有可用的 TP4,所以它只能當參照組。
排除flash attention 不是變數(f16 開關 FA:703.5 對 710.5)。先前 llama-bench 中止是因為 tbq3 必須要有 FA。
找到一份公開的 3090 紀錄跑同一顆 27B、同樣在 vLLM 上,MTP 逐位置接受率 97/95/91%、接受長度 3.4–3.8、125K 下 85 tok/s。我們是 50.5/15.2%、接受長度 1.50。差的是配置,不是架構。
更正我先前寫「只有一顆 MTP 頭,所以 position 1 崩到 15.2% 是架構決定的」。那個解釋錯了 —— 對照組同樣一顆 nextn 頭,position 3 是 91%。
否定draft 頭的量化不是低接受率的原因。四種量化(Q4_0 / Q8_0 / IQ4_XS)全部落在 40–49%,彼此在雜訊內,而且量化最重的 IQ4_XS 反而最高。把 blk.64 gate 到 Q8_0 重做 requant 這條路不用走了。
否定也不是我 8/6 關掉 GDN 旗標造成的。原始碼裡那個旗標只是 force,另有 auto 會自行武裝;8/8 與 8/9 的伺服器 log 都印了 guard armed (force=False, auto=True)。
修好llama.cpp 的 -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 就跑起來了。
撤回所有 llama.cpp 四卡的速度數字,以及「四卡贏不了單卡」這個結論。-ts 是斜線分隔,我用了逗號 —— 在 llama-bench 那是「掃多組設定」,結果四個臂全部跑在 GPU 0。破綻是 pp512 四種配置差 0.15%。重跑中,這次會驗每張卡有沒有吃到權重。
已撤回MTP 接受率低的原因是 KV 位元寬。2-bit → 3-bit,JSON 從 47.0% 到 87.1%。 8/12 的六臂掃描推翻了它(見本頁最上方)。留著這條是為了記錄我當時憑一次量測就寫了因果。同一段裡「3-bit 在 64K 仍有 64.65 倍併發餘裕」是對的,而且六臂又量了一次:2-bit 池 3.19M tok、3-bit 2.01M、f16 只有 0.55M。
量到KV 確實分散到四卡,而且超線性:TP1 370,085 tok → TP4 5,751,747 tok,15.5 倍不是 4 倍。TurboQuant 2-bit 給的併發是 f16 的 5.3 倍。
解決8 併發 30 以上做到了:Q4_K_S,c8 34.4,同一顆新權重、單人不輸 IQ4_XS。i-quant 在這個 build 裡沒有批次 kernel,ggml_mul_mat_a8 對它是空轉返回(M=2 到 512 耗時全是 0.03 ms,換算 3000 TFLOPS,V100 峰值 125)。
修正整天「27B 過不了工具呼叫閘門」是我自己的假警報 —— 測試腳本少了 --enable-auto-tool-choice 和 --tool-call-parser,生產的啟動腳本一直都有。
驗證fp16 對 Q4 的交叉點在 M = 32–64,和從算術強度推的 M ≈ 35 吻合。解碼端(M=1)量化快 8.8 倍,「fp16 更快」只在 prefill 區間成立。
做了本頁改成分類 + 類內最新在上,補上 TPS / TTFT 總表、27B 完整矩陣、KV 容量表,並把外部數據移到獨立表格。

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

生產與吞吐10 段
2026-08-17

Qwen3.8-27B:併發到 64 沒有天花板,解碼隨 context 只掉 37%

配置

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/sTTFT p50
168.15—1.52s
835.3155.30.28s
1618.8186.70.46s
3215.3246.73.69s
4811.4314.74.37s
6410.0356.92.57s
但生產只會到八人

32 以上那三列是為了找天花板量的,不是產品承諾 —— Wayne 定的上限是八個併發使用者。留著是因為它們回答了「還有沒有餘裕」(有,而且到 64 都沒轉折),但對外要引用的是 c8 那一列。席次已改回 8、capture 改回 [1,2,4,8],省下的記憶體回到 KV pool。

提示長度暖 TTFT解碼 tok/s
2580.25s67.52
3,6181.04s65.55
9,6183.17s62.24
21,6186.21s56.36
43,218—48.46
64,81823.09s42.73
八個人同時,每人一份 21.6K 文件牆鐘每人延遲合計 tok/s
八份不同的文件(沒有共用前綴)54.8s54.7 – 54.8s18.70
八份逐位元組相同的文件(先暖過)55.0s54.9 – 55.0s18.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 秒。

2026-08-12

批次化 MMVQ:權重讀一次,不是每個 token 位置讀一次

fork 的 mul_mat_vec_q 把 token 位置放在 blockIdx.y,所以每個位置各開一組 block、各自把 45 MiB 的權重讀過一遍。M=2 就是讀兩次。

profiler 的破綻

同一顆 kernel,無推測 2489 次 launch / 42.24 ms,有推測 2468 次 / 67.55 ms —— 次數幾乎相同、時間多 60%。那不是「呼叫變多」,是「每次呼叫做了兩倍的搬運」。

Mfork 每位置 ms批次版 每位置 ms總時間倍數
10.0860.0841.02x
20.0790.0551.45x
40.0760.0382.01x
80.0740.0322.34x

fork 那一欄幾乎是平的 —— 完全沒有重用。M=1 持平代表可以無條件取代,不傷單人路徑。九種 IQ4_XS 權重形狀 × 每個 M,誤差與原 kernel 一致到印出的精度(兩者共用同一套 q8_1 activation 量化)。

併發改前改後
c167.8969.09+1.8%
c248.5257.27+18%
c436.7448.47+32%
c820.5832.31+57%
c8 合計104.7143.3+37%

異質混合工作負載、同 session、只切一個環境變數,四臂工具呼叫閘門全 ALL_PASS。

一小時壓測:過了

時間速率 p50TTFT p50失敗GPU 最大 MiB
5m47.040.65s016450
15m46.780.65s016678
30m46.750.64s016678
45m46.890.66s016678
60m46.980.64s016678

四併發、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 —— 跟健康無關。

2026-08-12

27B 的推薦組態 —— 以及同構併發高估了 1.9 倍

完整條件,不是「那顆 27B」

模型 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
c130.1928.230.50—27.30.28s
c815.06103.215.0618.0861.60.51s
c1612.54187.512.4113.79101.21.12s
c328.99266.19.2210.05136.71.80s

異質那組是 12 種不同提示、錯開到達、輸出長度 100–400,32 個請求全數完成零失敗。每人速率兩者差不多,合計卻差 1.9 倍 —— 真實流量的請求長度不一,尾巴會拖長 wall-clock,而同構測法讓所有請求同時開始、同時結束。先前所有併發數字都該讀成同構上界,這是那句話的量化版本。

順帶更正:MTP 接受率不是 80%

先前每個併發都量到「剛好 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 請求是尾巴,不是平均。

2026-08-06

併發天花板是我們自己設的

先排除記憶體

伺服器自報 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 合計合計增益
1123.2123.571.471.5—
868.468.3261.9286.0+9%
1667.855.4228.8403.2+76%
3267.736.9227.0619.0+173%
那條「打平在 68」的平原是排隊,不是極限

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 完全相同,所以它只會在人多時起作用,不會犧牲單人速度。

2026-08-06

每卡 29GB 是怎麼來的 · 砍到 9.5GB 沒有代價

佔用的不是權重,是 KV 預留

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/sutil 0.30 decode/s差
200 / c1123.6123.0−0.5%
200 / c479.879.7−0.1%
200 / c868.168.2+0.1%
28000 / c1106.0106.00.0%
28000 冷 TTFT6.58s6.58s0.0%
已套用(commit 8aed933)

全部落在噪音以內,所以 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 分鐘重啟,而且量到的就是要上線的東西,不是它的手打近似版。

2026-08-06

生產 +24% TPS:一個開著的旗標本身就是瓶頸

找法:對固定成本做消去,而不是再推論一次

三個嫌疑各做一臂。--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/s123.4 / 123.3153.1 / 152.8+24.0%
c8 decode/s68.0 / 68.076.3 / 76.8+12.5%
c8 總吞吐272.8 / 273.3275.7 / 272.7持平
退化指標1.001.00—
工具呼叫ALL_PASSALL_PASS—
已上生產(commit 9972c07)

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 —— 那個看起來不可縮減的固定成本,有一大塊其實就是這個旗標。

2026-08-06

27B 也一樣,而且長 context 併發下最明顯

context / 併發GDN=1(兩趟)GDN=0(兩趟)差
200 / c154.3 / 54.365.3 / 66.5+21.4%
200 / c828.0 / 28.130.8 / 31.3+10.7%
28000 / c145.9 / 46.454.2 / 54.2+17.4%
28000 / c817.0 / 17.021.7 / 21.6+27.4%
28000 / c8 總吞吐47.472.7+53.7%
28000 這兩列是第一次量到

先前它們被報成 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 一致。

2026-08-05

完整矩陣 · 一次跑完

為什麼重測全部

今天有六次「反常結果」最後都是量測方法的問題,不是被測物。零散修補會留下不同批次混用的數字,所以整個矩陣用同一套工具重跑一次:4 個配置 × 3 種 context × 5 種併發。

關鍵改變:prefill 與 decode 分開量。先前的每 token 速率是 tokens / 總耗時,把一次性的冷 prefill 攤進了每 token 成本。現在用串流,TTFT 與間隔中位數各自成欄,冷暖也分開,並附伺服器自報的 prefix 命中率。

配置context併發冷 TTFT暖 TTFTdecode tok/s總吞吐prefix退化
9B · TP4~200c10.23s0.16s123.170.80%0.91
c20.23s0.32s103.192.30%0.91
c40.23s0.33s79.7140.10%0.87
c80.23s0.36s67.9286.60%0.91
c160.23s0.96s68.0199.20%0.83
9B · TP4~8Kc11.13s0.76s117.926.757%0.91
c21.13s1.51s97.158.857%0.80
c41.13s2.85s73.846.957%0.82
c81.13s4.43s60.973.857%0.81
c161.13s6.43s58.970.857%0.79
9B · TP4~28Kc18.93s0.57s106.232.598%0.92
c28.93s1.05s83.341.198%0.92
c48.93s1.86s63.4105.098%0.78
c88.93s3.28s45.9105.198%0.79
c168.93s3.89s44.2101.798%0.74
9B · TP2~200c10.29s0.22s96.258.80%0.93
c20.29s0.44s72.072.40%0.93
c40.29s0.40s57.4126.90%0.91
c80.29s0.66s50.0175.70%0.87
c160.29s1.18s49.7198.30%0.91
9B · TP2~8Kc12.59s1.39s90.015.657%0.91
c22.59s3.27s66.232.357%0.76
c42.59s5.36s51.243.757%0.80
c82.59s8.41s41.230.657%0.83
c162.59s12.22s40.337.157%0.80
9B · TP2~28Kc117.39s0.89s77.930.698%0.95
c217.39s1.70s55.146.898%0.82
c417.39s3.21s42.067.498%0.80
c817.39s5.99s27.996.798%0.72
c1617.39s7.90s27.993.998%0.72
9B · 單卡~200c10.44s0.36s69.740.80%0.93
c20.44s0.72s47.742.50%0.91
c40.44s0.59s39.887.50%0.91
c80.44s0.73s33.0138.20%0.91
c160.44s2.00s32.9123.50%0.91
9B · 單卡~8Kc13.68s2.59s63.98.757%0.91
c23.68s5.13s42.821.157%0.76
c43.68s10.08s33.715.257%0.82
c83.68s19.67s25.626.557%0.81
c163.68s23.60s25.322.457%0.76
9B · 單卡~28Kc133.83s1.57s53.313.298%0.92
c233.83s3.10s34.015.898%0.92
c433.83s5.94s24.232.098%0.82
c833.83s11.43s18.844.498%0.77
c1633.83s14.30s16.443.198%0.80
27B · TP4~200c11.30s0.35s60.452.20%0.76
c21.30s0.70s45.473.20%0.70
c41.30s0.70s32.0109.30%0.74
c81.30s0.80s30.9208.60%0.70
27B · TP4~8K / ~28K—失敗待重跑:8K 回 0 個 content delta(疑似 max_tokens=128 被 think 吃光,即已知的串流空白陷阱);28K 回 HTTP 500(疑似超過我設的 maxlen 32768)。兩者皆未證實。
純 TPS(decode tok/s)

這是不含 prefill 的逐字生成速度,也就是使用者實際感受到的吐字速率。先前報的 111.5 是 tokens / 總耗時,被一次性的 prefill 稀釋過;當時真正的單人天花板是 123.1。這個數字已被取代：關掉 GDN 融合之後是 152.8，見本頁上方。

context單卡TP2TP4單卡→TP4
~200 · 1 人69.796.2123.1+77%
~8K · 1 人63.990.0117.9+85%
~28K · 1 人53.377.9106.2+99%
~200 · 8 人33.050.067.9+106%
~8K · 8 人25.641.260.9+138%
~28K · 8 人18.827.945.9+144%

TPS 的最大殺手是併發,不是 context。TP4 從 1 人的 123.1 掉到 8 人的 67.9(−45%),而同樣是 TP4、從短提示拉到 28K 只掉 14%(123.1 → 106.2)。所以要提升 TPS,該攻的是併發時的排程與批次效率。

另一個方向:卡加得越多,在越吃緊的情境越值得 —— 單卡→TP4 的增益從短提示單人的 +77% 一路升到 28K 八人的 +144%。

附帶發現:TP 在 prefill 的增益比 decode 更大

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 才是穩健的指標(串流間隔中位數,不受回應長度影響);總吞吐只適合看數量級。

2026-08-04 約

併發曲線 · 舊工具 · GDN 融合開啟時期(保留供對照)

配置1 人2 人4 人8 人8 人合計量測工具
9B · TP4 · 生產111.588.864.660.5393.1v2
9B · TP285.162.747.744.6282.2v2
9B · 單卡65.744.635.129.8177.8v2
27B · TP4 · 8K · 捕捉集合僅 [1,2]56.240.95.25.140.7v1 · 待重量
更正:並行增益沒有遞減

先前我說「TP2→TP4 只有 +19%,all-reduce 已是主導項,再加卡不會有回報」—— 那是錯的,是拿兩套不同量測工具的數字相比得出的。四個拓樸現在都在 v2 下量過:單卡→TP2 +29.5%、TP2→TP4 +31.0%,每翻倍的增益一致。單卡→TP4 的正確數字是 +70%,不是先前報的 +61%。

27B TP4 的 c4 斷崖是我的設定錯誤

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,兩者本來就不能直接比。

2026-08-04 約

長 context · 生產 TP4 · GDN 融合開啟時期

提示長度1 人4 人8 人16 人8 人合計
~900 tok106.451.849.634.8347.4
~7K tok40.317.99.69.189.8
~28K tok13.24.65.32.732.7
撤回:「長 context 慢八倍」是量測假象

上表的每 token 速率是 completion_tokens / 總耗時,把一次性的 prefill 攤進了每 token 成本,而每次只生成 128 個 token —— 8.85 秒的冷 prefill 除以 128,就把整個數字壓垮了。那量到的不是「長 context 很慢」,是「冷啟動成本被除進了速率」。Wayne 指出 prefill 只有第一次要付,後續輪次是逐步增長的 context,這是對的。

用串流把兩段拆開重量(同一份 context 連發三輪):

context輪次TTFT每 token 間隔decode tok/sprefix 命中
~8Kr11.16s0.0085s117.90%
~8Kr2–r30.75s0.0085s117.958%
~28Kr18.85s0.0094s106.316%
~28Kr2–r30.49s0.0094s106.398%

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 相當。

併發拐點確實在 8

這一條不受上述假象影響,因為它是同一列內的比較:短提示 c8 合計 347.4、c16 反而 332.4;28K 下 c8 32.7、c16 33.4 幾乎不動。超過 max-num-seqs 8 之後請求只是排隊,總吞吐不再增加。

量測工具 v1 的三個缺陷(已修)

舊工具的比例並不自洽 ——「每人 × 人數」與「合計」相差達 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 量的,重量需要停生產的維護窗口。

27B TP4 為何沒有數字

那一臂在切換前執行,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 當提示,掩蓋了輸出退化,而退化的輸出解碼更快,反而灌高數字。基準測試現已內建真實文本的詞彙重複度斷言。

長 context · 128K5 段
2026-08-07

目標:27B Fable · 128K · 8 併發 —— 全線驗證通過

先確認這不是品質問題

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。

第一次啟動失敗,但修法是改設定不是改 kernel

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 pool4,877,750 tokens = 131,072 之下 37.21x 併發(需要 8)
每卡記憶體19 GB / 32(util 0.60)
105,856 token 提示 · c135.1 / 35.0 tok/s(兩趟)
105,856 token 提示 · c88.1 / 8.1 tok/s
冷 TTFT390.9s
暖 TTFT(prefix 命中 96%)31.1s(c1) / 247.5s(c8)
needle 檢索 · 115,645 token5/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 是真的長,不是標籤寫大而已。

量測工具第五個缺陷:context 標籤從來不是 token 數

拿 tokenizer 實際數過之後:matrix 的標籤高估 1.13 倍(~120000 其實是 105,856),而 needle 探針的英文填充文字差 4.65 倍 —— 它第一次要 120000、實際只拿到 25,804,我差點把 26K 的五連過報成 128K 通過。差別在填充文字:matrix 用中文(約 1.4 字元/token),needle 用英文(約 6.5)。探針現在用真的 tokenizer 長到目標,並把實際拿到多少印出來。

2026-08-07

128K 的成本全在 prefill,而其中一半可以買回來

105,856 token 提示batched-tokens 16384batched-tokens 65536差
冷 TTFT412.42s304.04s−26.3%
暖 TTFT · c131.08s31.14s不變
暖 TTFT · c8247.73s246.34s完全不變
decode c135.035.0不變
decode c87.57.5不變
猜對一半:冷 prefill 是分塊綁的,c8 的懲罰不是

暖起來之後 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 那套配置真的架起來時要帶的設定。

2026-08-07

八個人問同一份 128K 文件:每人還是 31 秒,而且不共用

先前的解釋是錯的

我把 c8 的八倍懲罰歸給「量測工具給每個併發請求不同的尾巴,所以 prefix 只命中 96%,缺的約 4,200 個 token 各自要對整段做注意力」。照這個說法,八個人問同一份文件(逐位元組相同)應該很便宜。不是。

105,856 token · 完全相同的提示TTFTprefix 命中
第 1 次(冷)303.23s—
第 2 次30.95s96%
第 3 次30.95s96%
8 人同時問同一份最快 31.05s · 中位數 245.13s · 最慢 245.14s96%
兩件事,都是操作面的

一、暖起來的 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 變體都是同一種混合架構。先記錄,不再往下猜 —— 今天已經有四個從架構往下推的假設被實測打掉。

2026-08-07

那 31 秒不是固定成本 —— 它長得比文件還快

文件大小暖 TTFT(兩次)對前一列的成長
24,704 tokens1.79 / 1.78 s—
49,4085.90 / 5.90 stoken ×2.00 → 時間 ×3.31(n^1.72)
74,11212.58 / 12.59 s×1.50 → ×2.13(n^1.87)
105,85630.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² 還快,所以答案在文件大小,不在硬體。

2026-08-07

把 KV pool 調小不會讓 TPS 變快 —— 而且有個下限

27B · 128K 配置KV pool併發餘裕106K · c8 decode/s
util 0.60 · bt 163844,877,75037.21x7.5
util 0.60 · bt 655362,523,13619.25x7.5
util 0.25 · bt 65536起不來 —— No available memory for the cache blocks
7.5 是不動的

今天轉過的每一個旋鈕 —— 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 那節)、或換小模型。

MTP · 推測解碼12 段
2026-08-13

對照組:遞迴不是驗證變貴的原因

配置

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草稿接受接受率每步 tokenc
base(不開推測)32.63———1.000—
DFlash 預設36.5542815536.2%2.0690.847
DFlash n-max 831.5590118420.4%——
DFlash n-max 1631.051,56519112.2%——
把成本除以「每步多驗幾格」遞迴 block多驗幾格c每格的 c
Qwen3.6-27B · llama.cpp 單卡 · MTP48 / 6510.2980.298
Muse Glimmer 30B · llama.cpp 單卡 · DFlash0 / 522.950.8470.287
Qwen3.6-27B · vLLM TP4 · MTP(最佳旗標)48 / 6511.5951.595
0.298 對 0.287

同一個引擎、同一張卡,一顆 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 內部那兩列是同引擎、同卡數、兩種極端不同的架構,那組就足以殺掉遞迴假設。

2026-08-12

KV 精度掃描:接受率一動也不動

配置

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

KVKV 池c1 解碼c1 接受率c4 解碼c4 接受率
2-bit · 不開 MTP3,191,98859.25—48.52—
turboquant 2-bit2,513,30533.8051.5%24.8550.3%
turboquant 3-bit2,010,64433.4150.0%26.4961.5%
turboquant 4-bit1,615,14533.4752.3%27.1062.4%
K 8-bit · V 4-bit1,230,79834.7351.5%27.9561.9%
f16(完全不壓)546,61333.7151.5%27.7761.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 不是基準。

2026-08-12

八個旗標消融:只有兩個動得了 c

配置

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 解碼接受率推測步 mscTTFT
不開推測68.88—14.52—0.43s
MTP 原樣34.2951.5%44.182.0431.46s
FULL_FORWARD34.0350.0%44.082.0362.86s
SPEC_CORE_OP39.6249.3%37.681.5954.42s
003_SPEC_CORE_OP38.4944.9%37.651.5932.88s
PACKED_RECURRENT_DECODE34.7251.5%43.632.0051.36s
GDN_DECODE_FLASHQLA34.3751.5%44.082.0360.77s
FUSED_SIGMOID_MIXED_QKV34.5049.3%43.281.9813.54s
疊加c1 解碼接受率cTTFT
FLASH_GRAPH 單獨34.5551.5%2.0201.55s
FLASH_GRAPH + SPEC_CORE_OP39.9652.3%1.6250.68s
SPEC_CORE_OP + FUSED_SIGMOID38.7045.6%1.613.61s
SPEC_CORE_OP + PACKED_DECODE39.8152.3%1.6351.29s
兩個贏家落在同一個 c

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%),所以這個代價不必付。

2026-08-12

兩台引擎的完整帳:公式對兩邊都準

配置

同一顆 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αcS 預測S 實測MTP tok/s
llama.cpp · 單卡39.1625.530.7220.2981.3271.32751.97
vLLM · TP468.8614.520.5152.0100.5030.50334.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 的優點搬過來αcStok/s(基準 68.86)
現況0.5152.0100.50334.64
只修 α0.7222.0100.57239.4 · 仍然虧
只修 c0.5150.2981.16780.4
兩個都修0.7220.2981.32791.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 的差距不可能是算力。

2026-08-12

拿掉合約裡的 nonzero():MTP 單人 +6.8%

配置

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.8648.59—ALL_PASS
MTP draft=1 · NOSYNC=032.4224.4551.5%ALL_PASS
MTP draft=1 · NOSYNC=134.6425.0551.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。

2026-08-12

草稿長度掃描:第一個位置就只有五成

配置

同上一節,只換 num_speculative_tokens;capture 集合跟著改成 (1+γ) 的倍數(γ=1 [2,4,8,16] / γ=2 [3,6,12,24] / γ=3 [4,8,16,32]),否則會被靜默丟掉

γ草稿接受pos0pos1 條件pos2 條件總接受率c1 解碼
173036850.4%——50.4%33.21
21,37441044.8%33.1%—29.8%30.81
32,03441943.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_S21.903.44156.9 ms9.1x
spec3 · Q4_K_S · SPEC_CORE_OP=123.863.49146.2 ms8.5x
spec3 · Q4_K_S · 003_SPEC_CORE_OP=123.553.45146.4 ms8.5x
spec3 · IQ4_XS18.313.61197.4 ms11.5x
spec3 · IQ4_XS · SPEC_DECODE_PIECEWISE=15.003.49698.2 ms40.6x
spec3 · Q4_K_S · TP112.053.49289.5 ms16.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,…>9641,292430x
_scatter_gather_elementwise_kernel08,160全新
fused_sigmoid_gating_delta_rule_update(GDN 本體)28,7048,0640.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,重跑中。

fp16 真的比 Q4 快嗎 · 交叉點實測

配置

權重 17408 × 5120(此模型的 ffn 形狀)· V100-SXM2-32GB · cuda event 計時 · 30 次取平均 · fp16 那欄包含它必須付的解量化成本

列數 M量化 kernel解量化 + fp16贏家
10.074 ms0.651 ms量化 8.8x
80.177 ms0.656 ms量化 3.7x
160.328 ms0.664 ms量化 2.0x
320.647 ms0.686 ms量化 1.1x
641.286 ms0.790 msfp16 1.6x
1282.584 ms1.020 msfp16 2.5x
2565.153 ms1.068 msfp16 4.8x
102420.612 ms2.304 msfp16 8.9x
204841.325 ms4.368 msfp16 9.5x
交叉點在 32–64 之間

質疑是「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 倍,門檻設得保守。

TTFT 總表 · 首 token 延遲

冷 = prefix cache 沒命中,暖 = 同一份提示再問一次。長 context 的代價幾乎全在冷啟動,所以 128K 能不能用,取決於同一份文件會不會被重複問。

情境context冷 TTFT暖 TTFTprefix 命中時期註記
9B · TP4 · c1~200 tok0.23s0.16s0%GDN 開
9B · TP4 · c1~8K tok1.13s0.76s57%GDN 開
9B · TP4 · c1~28K tok8.93s0.57s98%GDN 開
9B · TP4 · c8~28K tok8.93s3.28s98%GDN 開
27B · TP4 · c1~200 tok1.30s0.35s0%GDN 開
27B · TP4 · c1105,856 tok390.90s31.10s96%現行128K 配置
27B · TP4 · c8105,856 tok390.90s247.50s96%現行每人
27B · TP4 · c1115,645 tok492.80s16.70s命中現行needle 探針同一份提示
冷啟動買回來的部分 · 只改設定沒改 kernel

max-num-batched-tokens 16384 → 65536:冷 TTFT 412.42s → 304.04s(−26.3%)。block-size 16 → 64:暖 TTFT −15%。

服務

fable-9b · V100:8600
無回應 ()
Qwopus-27B regen · V100:8900
無回應 ()
工具呼叫
已上線 · 8/8 + 7/7
Hermes gateway
待啟用

背景產線

DSpark drafter 訓練 · DGX GB1013440 / 13440 · 100.0%
1400 樣本子集 × 10 epoch,梯度累積 64。記憶體已打平,約 2.53 s / micro-step
DSpark target cache 建置 · DGX4948 / 5093 · 97.2%
已完成,4948 樣本 / 103GB
Qwopus-27B regen 資料生成 · V100 TP28931 / 8931 · 100.0%
可斷點續傳,供訓練資料擴量

吞吐量實測

9B-Q4_K_M · 64K · TP4×1人 · GDN 融合關閉
現行生產
152.8 tok/s
9B-Q4_K_M · 64K · TP4×1人 · GDN 融合開啟
2026-08-06 前
123.4 tok/s
9B-Q4_K_M · 64K · TP2×1人 · TQ-2bit s96
90.5 tok/s
9B-Q4_K_M · 64K · 單卡×1人 · TQ-2bit s96
切換前
66.6 tok/s
9B-Q4_K_M · 64K · TP4×8人 · GDN 融合關閉
每人
76.8 tok/s
27B-Q4_0 · 8K · TP2×1人 · KV auto
37.2 tok/s
27B-Q4_0 · 8K · TP2×4人 · KV auto
每人
22.5 tok/s
0每使用者 tok/s153
2026-08-12

推測解碼把 capture 集合的需求乘上了 (1+draft),沒人檢查捕捉張數

症狀:16 到 24 併發之間掉 2.9 倍

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 張164.28
MTP draft=2[1,2,4,8,16,32]0 張—連啟動都沒過
MTP draft=1[1,2,4,8,16,32,48,64]7 張328.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]
c130.2430.3830.47
c89.719.899.89
c169.589.589.56
c243.767.878.16
c323.667.597.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 的比較在這台機器上不成立,要下結論的對照一律同場跑完。

2026-08-12

推測解碼的損益平衡是 α > c,而我們的 c 是 2.04

論文給的是閉式解,不用猜

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 127B IQ4_XS,TP4 V100
α 範圍0.53 – 0.820.39 – 0.53
c 範圍0.007 – 0.0732.04
加速1.4x – 3.4x0.49x

論文自己的數據就在證明 α 不是槓桿:Table 2 裡接受率最高的一列 —— T5-large,α = 0.82 —— 加速卻是最差的 1.7x,因為它的 c 最大(800M/11B)。而 α 只有 0.75 的 T5-small 拿到 3.4x。

我們的 c 不是草稿頭的成本

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 必賠,是同一個病灶。

病因:這顆模型有 48 層帶遞迴狀態

推測解碼的整個前提是「多驗一個位置幾乎免費」。對 attention 成立 —— 多寫一格 KV,被拒絕就丟掉,零成本。對 linear attention / GDN 不成立:遞迴層沒有「丟掉一格」這種操作,它只保留最後一個 token 的狀態。

vLLM 自己的原話

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 名根本看不到,開了之後全部冒出來 —— 而且它們是搬資料的,不是算數的:

kernelnospecmtp做什麼
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。

2026-08-10

外部對照組把「單頭」這個解釋否掉了

再更正一次,這次是我上面那段

我在上一節寫: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.81.50
不開推測38 tok/s66.5(四卡)
開 MTP63.8 → 85(加 TurboQuant)30.2(倒扣)
模型格式AutoRound INT4,mtp.fc 解量化成 BF16GGUF,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-requantQ4_041.7%42.85
Q4_0-ehq8-requantQ8_040.5%42.35
IQ4_XSIQ4_XS47.2%47.19
Q4_K_MQ8_049.4%35.19
(a) 死了:draft 頭的量化不是原因

分離對照沒有分岔 —— 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 就能判定,不必重跑。

官方 vs 合併:合併有代價,但只佔一小半

官方 safetensors 直轉的 Qwen3.6-27B-IQ4_XS-mtp,同量化型別、同引擎、同提示、同 session,各兩趟:

模型接受率tg
官方 Qwen3.6-27B60.7% / 60.7%52.8 / 53.9
合併 Fable-Fus-71147.2% / 46.2%47.1 / 47.0
但真正的大頭是我出的題目

60.7% 離對照組的 97/95/91% 還差 34 個百分點,而那段幾乎全在量測條件。同一顆合併模型,只換內容類型:

內容接受率tg實際產出開頭
高重複99.0%77.10The quick brown fox jumps over…
結構化 JSON85.4%69.56[ { "id": 1, "name": "Alice Johnson"…
程式碼70.9%60.68from __future__ import annotations…
散文(原本下結論用的那題)47.2%47.35—
所以先前「MTP 不划算」是在最不利的那一格判的

同一顆模型從 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 比;另有兩格回傳空字串,標為無效,原因未查。

2026-08-08

MTP 的兩層問題,以及模型作者本人給的對照

先更正我自己的說法

我先前寫「推測解碼在這裡結構上不可能贏,加速比 = τ/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 0position 1
229.9%1.6044.7%15.2%
150.5%1.5050.5%—
第二層問題:spec-tokens 2 在架構上就不對

作者的 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 同一個量級。

修正給 Codex 的驗收目標

接手檔原本寫「kernel 修好之後 τ ≈ 1.5–2.5x」。實測 τ 是 1.45–1.60,所以實際期望是 1.5–1.6x。已更正 —— 給錯的驗收目標比不給更糟,會讓一個成功的修正看起來像失敗。

2026-08-03 約

推測解碼:找到根因,修掉大半,但還沒對

症狀與根因

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算術複誦序列工具呼叫degdecode/s
CTL_spec3(修正前)2錯錯錯0/30.2121.7
FIX_spec28錯對對3/31.0024.9
BASE_nospec(不開推測)—錯對對3/31.0054.2
FIX_spec38錯對對3/31.0020.8
修正前是全面崩壞

threshold 2 之下 12345×6789 答成 8333333333…、複誦只吐出一個 X、質數序列變成 2, 3, 5, 5, 5, 5, …1111111…,三個工具呼叫一個都沒發出來。threshold 改成 8 之後複誦、序列、三個工具呼叫全部通過,退化指標從 0.21 回到 1.00。

更正:算術那題是模型不會,不是 bug

我原本從「兩個 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 這一步。這條路沒有東西可以省。

2026-08-02 約

推測解碼 sweep · GPU3

假設

MTP 在自由文本上是淨損,但工具呼叫的輸出高度重複(JSON 骨架、欄位名、<tool_call> 標籤幾乎每次相同)—— 那是 n-gram 命中率最高的地形。因此每個臂同時測工具呼叫與中文散文兩種負載:前者驗證假設,後者當對照組,確認它不會拖累一般對話。三種方法都不需要額外的 draft 模型。

方法工具 c1工具 c4 合計散文 c1uniq備註
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
中文正常輸出 ×30.92–1.001.000
英文正常輸出0.5780.713
單字重複 300 次0.0031.000
短句迴圈0.0221.000
兩句交替0.0341.000
英文詞迴圈0.0080.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 世代。

引擎 · 量化 · 對照5 段
2026-08-12

一顆什麼都沒算的 kernel 通過了數值比對

誤差剛好是零,那不是好消息

驗 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 把剛算完參考值、剛被釋放的那塊記憶體配給了它,所以比對讀到的是自己上一步的答案。

路徑實際發射的 kernelM=64 時間
反量化整顆 + cuBLASdequantize_block_iq4_xs 0.663 + cutlass 0.4411.106 ms
mmvqmul_mat_vec_q<block_iq4_xs>4.169 ms
mmq只有 quantize_q8_1不存在
這留下一個地雷

任何人若順手把 IQ4_XS 加進 MMQ_QUANT_TYPES 來修批次崩塌,模型會安靜地吐垃圾,而且不會報錯。目前生產沒中,只是因為 gguf.py 正好把 i-quant 排除在那條分支外。

做法上的教訓:正確性閘門要毒化輸出緩衝區,並且斷言誤差不為零 —— 對上獨立實作,零誤差是可疑不是優秀。再配一道 roofline(搬的位元組 ÷ 時間 對照卡的峰值),因為沒運算的 kernel 一定快得不合理,那是最便宜的破綻。

2026-08-12

批次 kernel 早就寫好了,只是被鎖在 Q4_0

IQ4_XS 在 M>8 掉進「反量化整顆權重再 cuBLAS」,固定 1.106 ms 且不隨 M 變;同尺寸的 K-quant 走真 MMQ 在 M=16 只要 0.355 ms。mmvq 是線性的,與反量化路徑的交叉點在 M≈15,而 gguf.py 的門檻寫死在 8。

Mmmvq反量化 + GEMM目前會選
10.0861.112mmvq
80.5851.098mmvq
120.8711.101反量化(慢 26%)
161.1571.076反量化
644.4081.214反量化

寫出來了 —— 端到端 c8 +51%

M新 kerneldispatcher 現行會選的倍數
40.213mmvq 0.2831.33x
80.214mmvq 0.5292.48x
160.205反量化 1.0174.95x
320.421反量化 1.0832.57x
640.616反量化 1.2021.95x
961.005反量化 1.4101.40x

端到端,同 session、只切 VLLM_GGUF_IQ4XS_MMQ、MTP draft=1:

併發關開每人合計
c130.4229.91−1.7%(對照組,M=2 兩邊都走 mmvq)—
c89.7214.71+51%72.7 → 108.8
c169.5812.54+31%140.2 → 186.4
c248.159.85+21%182.4 → 217.7
c327.609.00+18%225.3 → 265.9

兩臂產出的文字逐字相同。工具呼叫閘門兩邊一樣失敗 —— 那個問題在 kernel 關閉、也在每一組 capture 集合下都存在,不是這次引入的。

單人那 1.5% 我沒解釋掉

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 —— 見上一節,做完了。

2026-08-10

llama.cpp 對照 · MTP 會贏,但 KV 量化與四卡都輸

為什麼還是留在 vLLM

llama.cpp 是參照組,不是候選。它證明了 MTP 在這顆模型上本來就該賺,所以 vLLM 這邊虧是我們的配置有問題 —— 這個證據值錢。但它自己不能當產線。

項目llama.cppvLLM(我們的 fork)
MTP38.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)pp512tg128
單卡14645 / 10 / 10 / 10855.5439.24
layer 四卡4041 / 3629 / 3835 / 4365863.4238.46
row 四卡4891 / 4681 / 4685 / 505588.9922.12
tensor 四卡4519 / 4313 / 4313 / 43132333.0933.61
重跑之後,兩個結論都反過來

提示處理是會縮放的:tensor 模式 2333 對單卡 855,2.7 倍。我先前說「多的三張卡什麼都沒做」是錯的。

row 模式是病態的,不是「不會變快」:pp512 掉到 88.99,比單卡慢十倍。先前那個假的 37.82 剛好把這件事整個蓋住。

撐過來的只有一項:解碼確實不隨卡數變快,三種模式都輸給單卡。但這次是真的量到的,而且理由清楚 —— batch 1 每層都要四路 all-reduce,65 層的同步次數吃掉頻寬省下的量。成本在屏障的數量,不在它的大小。

唯一沒被波及的:tensor 模式崩潰的根因

-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。所以四卡到底能不能加速,現在是未知而不是「不能」—— 重跑中。

apples-to-apples 的 KV 修正

先前的 llama.cpp 數字用 f16 KV,對上 vLLM 的 2-bit —— 那不是比較。補測之後,llama.cpp 開 TurboQuant KV 反而更慢,MTP 的增益也從 +26% 塌到 +2%。

受控組(llama-bench 的 -ctk/-ctv 是位置配對,所以這三列其實只量到 V):

type_ktype_vpp24576tg128
f16f16711.639.41
f16tbq3_0330.3(−54%)36.30(−8%)
f16tbq4_0377.3(−47%)36.38(−8%)
所以 K 才是貴的那一半

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 都建不起來。

2026-08-05

解碼不是頻寬綁的 —— 砍掉 18% 的位元組只換到 5%

起點:兩個模型都只跑到記憶體屋頂的四分之一

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。)

於是找到一個看起來很肥的目標:沒量化的 LM head

output.weight 是 BF16,而其他所有權重都是 Q4_K/Q6_K。詞表 248,320 很大,所以這一個張量就是 9B 的 29.2%(1940 MB)、27B 的 13.8%(2425 MB),而且每產生一個 token 就要完整讀一次。

臂大小c1 decode/sc4 decode/sthink%deg
Q4_0 · bf16head6,359 MB128.985.930.98
Q4_0 · q6head5,198 MB(−18.3%)136.188.41000.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% 思考的臂放在一起比,本來就該標出來。

2026-08-05

每個 token 有 6.02 毫秒不隨卡數縮減 —— 而且不是通訊

拓樸c1 decode/s該拓樸的記憶體屋頂屋頂效率每 token固定成本佔比
TP169.7138.750.3%14.35 ms42%
TP296.2277.334.7%10.40 ms58%
TP4(生產)123.4554.722.2%8.10 ms74%
擬合:T(n) = 6.02 + 8.32/n 毫秒

加卡只把可縮減的 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.4153.4123.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 與取樣。下一輪量這個。

基礎設施 · 根因2 段
2026-08-06

每次「V100 離線」的根因:它掛在公司 WiFi 上,網路孔沒插線

機器沒事,線路有事

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 中繼)
08-06 00:53 那次對得上分秒

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。

追查二:能不能同時連兩個 WiFi?不能,硬體就做不到

網卡是 TP-Link Archer T2U PLUS(RTL8821AU,單天線 AC600),驅動 rtl8821au。iw phy 明講 interface combinations are not supported —— 單一 radio、不支援介面組合,沒辦法同時掛兩個 STA 關聯。整台機器只有這一個 phy,沒有第二支網卡也沒有 PCI 無線。要第二條路就必須有第二個 radio:再一支 USB 網卡、一支 4G dongle,或者最省事的 —— 那條沒插的網路線。

但找到一個可能更大的槓桿:它待在 DFS 頻道上

管制區是 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 無線開省電本身也是掉線的常見原因。)

2026-08-06

修法與實測

修法修法前修法後說明
num_anchors 512 → 128跑不到第一個 micro-step穩定推進Q_LEN 3584 → 896,scores 面積少 6.2 倍。硬體逼出的偏離,τ 需註明在 128 anchors 下測得。
expandable_segments49→99 步流失 12.3GB流失 1.7GB治裝置端碎片。有效但不足以收斂。
pin_memory=False步 464 剩 12.3GB 後陣亡步 900 仍有 37.4GB治主機端釘選快取。GB10 主機與裝置共用實體 RAM,釘選換不到 DMA 好處 —— 每步耗時反而從 5.0s 降到 2.53s,全程預估從 19 小時縮到 9.4 小時。
DGX · DSpark2 段
2026-08-09

DSpark 記憶體診斷 · DGX

為什麼四層才到底

前三個假設都被實測推翻,而且它們都很有說服力 —— 留在這裡是為了不再重走。真正的轉折是不再靠推論、改看哪一套帳本對不上:當 cuda_reserved 與程序 RSS 同時持平、系統可用記憶體卻持續下降,這個組合只指向一個地方。

推翻
27B 目標模型全載吃掉 54GB
lowmem 路徑早已只從 safetensors 取出 embed_tokens 與 lm_head。它只在失敗時列印訊息,所以「log 裡沒有訊息」正好代表它有生效。
推翻
103GB 的 target cache 被 mmap 讀爆
輾轉當下 buff/cache 只有 1GB。mmap 的頁面計入 page cache 且可回收,不會逼出 swap。為此建的 1400 樣本子集沒有解決問題,但保留 —— 它讓每輪實驗快很多。
推翻
凍結的 embed / lm_head 仍被配了 Adam 狀態
optim.py 已濾掉 requires_grad=False 的參數,而且用的是 bitsandbytes AdamW8bit。
成立
① 孤兒 CUDA 子程序扣住 59GB
kill 父程序不會殺掉 multiprocessing.spawn 的子程序;它孤兒化後仍持有 CUDA context。GB10 的 GPU 記憶體從系統 RAM 切出,驅動那份不計入 /proc/meminfo 任何類別(所有具名項目加總只有 23GB / 121GB),nvidia-smi 回報 [N/A],該程序 RSS 僅 2.4GB —— ps 和 nvidia-smi 都看不見它。基準線因此從 121GB 悄悄掉到 61GB。
成立
② 注意力 scores 矩陣被實體化
drafter 的注意力是 Q=3584 × KV=7680。sm_121a 拿不到融合核心:flex_attention 不能編譯,換 SDPA 又因自訂遮罩退回 math 後端。兩條路都會實體化整個 fp32 scores 矩陣。
成立
③ 釘選主機記憶體快取無限成長
最後一層,也最隱蔽:cuda_allocated、cuda_reserved 與程序 RSS 三者同時持平,系統可用記憶體卻在 365 步內掉了 9.2GB。釘選區塊不在任何一套帳本裡,而 torch 依尺寸快取且從不釋放 —— 這份資料每個樣本長度都不同,每個新形狀就留下一塊永久快取。
2026-08-08

DSpark drafter 評估

checkpointτ 接受長度verify rate位置 0位置 1位置 2位置 3+auc 整體auc@0
step_210 · 2 層 · 128 anchors · 6.8GB1.560.1950.4910.0450.015<0.0040.5280.861
step_70 · 5 層 · 512 anchors · 9.0GB1.240.1540.1800.0340.013<0.0050.4850.397
輕量 drafter 每一項都贏

τ 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。

尾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 筆已收齊 → 資料擴量重訓
自動產生 · 自有硬體推論與訓練研究狀態2026-08-17_1728