PRODUCTION CHANGE REPORT SWITCHED 2026-08-26 · 4×V100-SXM2-32GB

V100 換引擎

生產推理棧由 IQ4_XS GGUF 換為 v100-skinny 引擎 + 自家 Uncensored NVFP4 檢查點。一晚三扇維護窗完成驗證、7 配置矩陣掃描與切換;同 IP、同 port、同模型名,消費端零改動。

單流解碼
101 tok/s
67.7 → 101(+50%)
混合負載 c8
134 tok/s
46.4 → 134(2.9×)
每卡 VRAM
13.9 GB
29.5 → 13.9(−53%)
品質
無損
草稿逐 token 由本尊驗證
  • 離線定罪編譯鏈,裝完整 CUDA 12.8,離線證明全綠
  • 窗 1base 模型驗引擎,boot gate 11 項全過
  • 窗 2自家模型組裝上殼,一次過、三域零損耗
  • 窗 37 配置矩陣,每配置重開機實測
  • 切換上線 8601,記憶體兩輪收斂至 0.40

2.0卡了七次的根因:不是 sm_70 太舊,是 header 太舊

前七次啟動全倒在 tilelang JIT 的 35 個編譯錯誤(__hmax_nan、TMA enum 未定義)。逐項排除逾時、孤兒行程、gcc-13、環境變數、ninja 之後,10 行探針釘死真兇:隨包的「CUDA 12.8」只是 540KB 的 nvcc 殼,nvcc -M 依賴清單顯示 header 全解析到 /usr/include 的 CUDA 12.0 舊檔。新版 header 對 __hmax_nan 的門檻本來就是 ≥530——Volta 從來不是問題。

免 root 裝完整 CUDA 12.8 toolkit 至 /mnt/data/nvfp4/cuda-12.8-full 後,探針、當初報錯的 reduce.h 本尊、torch extension 預熱三段離線全綠,第 8 次啟動一次通過全部 boot gate。

3.0兩路驗證:base 驗引擎,自家模型上殼

Path A(RadixArk base):11 項 boot gate 全過——QPN kernel 路由普查恰 512、XQA tensor-core 解碼、零 scalar fallback、MTP depth=7。

Path B(自家 Uncensored 微調):引擎只認 VL wrapper 結構,離線組裝三項缺口——權重鍵名嵌套 model.language_model.*、自 base 移植 vision(333)與 MTP 草稿頭(15,先前轉換跳過的)、補 401 個 input_scale。硬斷言 2194 tensor 普查與逐 dtype 形狀全同後上機,一次全過。

essay/prose/中文與 base 完全同速(移植草稿頭零接受率損失),code 148 → 128(微調風格分佈所致,仍勝舊棧 89%)。品質閘:uniq-word 0.74、命名多樣性 10/10、中文樣本人工比對通過。微調行為保留是結構保證:trunk 逐 byte 是我們的權重,草稿頭只出候選、一律被 trunk 驗證,vision 無輸入即死碼。

4.0完整矩陣:7 配置、每配置重開機實測

固定條件:NVFP4(MLP+lm_head)+FP8(attention)混合量化、FP16 KV、TP4、max_model_len 32768、greedy;數值為伺服器 usage 統計(含 prefill)。

量測方法

單流
非串流 /v1/completions,temp=0,每領域固定 prompt、max_tokens=384;tps = usage.completion_tokens ÷ 牆鐘(送出到收完,含 prefill,短 prompt 佔比約 7%——保守值)。每領域 2 輪取穩態。
併發
N 執行緒同發、每條 256 tok;aggregate = 全部 completion_tokens ÷ 整批牆鐘。per-request 含排隊等待——MNS=1 時後段請求慢是排隊,不是解碼慢。
長上下文
prompt 實長以 usage.prompt_tokens 為準;TTFT = 串流第一個 chunk 到達時間;長文後解碼 = completion_tokens ÷(串流總時長 − TTFT)。
token 計數
一律用伺服器 usage.completion_tokens。踩過的雷:數 SSE 事件會把 MTP 的一事件多 token 數成 1,曾把 104 量成 38(差 2.7×)。
舊棧數字
67.7/46.4 是換裝前以 tpsQ 腳本量的串流解碼速率(不含 prefill),方法略不同——新數字含 prefill 仍快 50%,差距只會被低估。
單流解碼 tok/s(穩態)
配置essaycodeprose中文
k=7 · MNS=1101.6127.4100.498.7
k=7 · MNS=2101.7127.5100.598.9
k=7 · MNS=4101.1127.2100.598.1
k=7 · MNS=8101.6127.2101.098.5
k=3 · MNS=1109.8126.9111.2106.1
k=3 · MNS=4109.2126.3110.7105.6
k=1 · MNS=185.189.682.782.0
舊棧 IQ4_XS67.7———

換 k 後首個請求有一次性 JIT 稅(k=3/k=1 首發 60.5/52.0),第二發即穩態;部署已預熱。

併發混合負載 aggregate tok/s
配置c2c4c8c8 最慢單條
k=7 · MNS=1109.2105.3105.519.4s
k=7 · MNS=292.0135.2149.813.7s
k=7 · MNS=4114.417.919.0107s
k=7 · MNS=8132.018.827.957s
k=3 · MNS=1111.2109.9110.018.6s
k=3 · MNS=493.319.820.696s
k=1 · MNS=183.682.482.524.8s
舊棧 IQ4_XS——46.4—
長上下文 prefill(MNS=1)
配置2,908 tok TTFT11,601 tok TTFT長文後解碼
k=70.60s0.81s115–123
k=30.55s0.68s112–124
k=10.55s0.68s83–90

使用者實際走的那條路(8/28 補測)

上面三張表量的是基準路徑:直連後端的 /v1/completions,長生成、原始續寫。 使用者實際走的是 chat 端點 + relay(bot → relay → 後端),生成較短且帶思考段。 同一台機器、同一組配置,兩條路的數字差很多,所以分開列:

同一組 8 併發,四條路徑對照
路徑c4c8c8 最慢
completions · 直連(基準)—149.0—
chat · 直連95.9102.212.9s
chat · relay(cap 8)92.995.815.1s
chat · relay(cap 2)77.595.314.2s
chat · relay(cap 1,原設定)68.363.919.9s

歸因(8/28 實測釘死):差距來自推測解碼的接受長度,速度幾乎正比於它—— 同一段 prompt、同樣 384 token,completions 98.9 tok/s ↔ 接受長度 2.77;chat 91.9 tok/s ↔ 2.60: 長度差 6%、速度差 7%,一比一。與思考無關(thinking 開或關都是 87.3)。 接受長度本身會漂(觀測 2.25–2.77),所以單次量測會跳,看 acceptance length: 指標才準。 relay 本身單流開銷為零、併發下約 6%;原本卡在 1 是最大浪費,放到 8 後 c4 從 68.3 → 92.9(+36%), 真正的併發上限由後端 MNS=2 決定,relay 不需要再卡一層。

調過但沒用的旋鈕:max_num_batched_tokens。vLLM 開機警告建議調高, 實測 4096 → 32768 對單流、TTFT、c8 全部無差別(雜訊內),卻把 KV 池從 332K 壓到 183K tokens (prefix cache 容量 −45%);131072 則因活化緩衝吃光預算而完全開不起來 (No available memory for the cache blocks,舊棧能用是因為配 GMU 0.85)。已還原 4096。

五個發現

  • k=3 單流反勝 k=7(essay +8%、prose +11%、中文 +7%):接受長度在 3 附近飽和,k=7 多出的草稿位純付驗證稅。
  • 併發斷崖在「同時執行 >2 條」,與 k 無關:MNS≤2 全健康,MNS≥4 一律崩到每條 4.5–5 tok/s;與 kernel 進入點 gemm_qpn8_mt2(mt2=雙序列)吻合。
  • 11.6K prompt 的 TTFT 不到 1 秒:長文 prefill 在此長度帶完全不咬。
  • MTP 深度是真槓桿:k=1 全域掉到 82–90,證明 100+ 來自推測解碼而非 kernel 本身。
  • 量測教訓:SSE 事件數 ≠ token 數(MTP 一事件多 token,曾把 104 量成 38);一律以 usage.completion_tokens 為準。

5.0記憶體的帳

依矩陣選 k=7 · MNS=2(c8 最佳、單流無損),記憶體按實際需求收斂:

KV 單價
64 KiB/token —— 64 層中只有 16 層 full-attention 存 KV,其餘 48 層 GDN 是固定小狀態
峰值需求
MNS=2 → 同時執行僅 2 條(排隊不佔 KV),2 × 32K 全長 = 4 GiB 全機;8–16 人、同時最多 8 人的情境壓力與 2 人相同,差別只在排隊秒數
設定值
GMU 0.88 → 0.72 → 0.40;實測 KV 池 126,670 tokens(真需求 1.9 倍,餘裕留給 prefix cache)
結果
每卡 29.5 → 13.9 GB,全機釋放約 62 GB;單流 101.4、c8 133.7,速度分毫未損
實際換算
KV 池 351,436 tokens 是四張卡合計的容量(每個 token 的 KV 切成四份分存):每卡 6.67 GiB、合計 26.7 GiB,每 token 約 80 KB(單卡分攤 20 KB)。含權重與活化,每卡佔 17.2 GB/32 GB,四卡共用 69 GB/128 GB。
夠不夠
vLLM 自算 Maximum concurrency for 135,168 tokens per request: 2.60x——可同時容納 2.6 個吃滿 132K 的請求,而併發上限本來就是 MNS=2,剛好夠且有 30% 餘裕。空著的 59 GB 買不到併發(天花板由 kernel 的雙序列快徑決定),只能買更大的 prefix cache。
KV 沒有量化
舊棧用 turboquant_2bit_nc(2-bit),池子 2,719,744 tokens;新棧是 FP16 未量化,池子 351,436。這是刻意的取捨不是漏設:checkpoint 本來要求 FP8 KV,但 V100 沒有 FP8 硬體,一旦壓縮 KV,手寫的 XQA tensor-core 解碼 kernel 就走不了、掉回慢路徑——開機 gate 的 FP16 KV resolved 就在守這條。用 7.7 倍的上下文容量(用不到)換 1.5 倍解碼速度(每天都在用)。真的撞到池子上限時,先升 GMU(還空著 59 GB)而不是開 KV 量化。

6.0營運手冊

服務位址
100.78.29.29:8601(IP、port、/health 與舊棧相同)
模型名
主名 qwen38-27b,另保留 qwen3.8-27b、fable-27b 兩個舊別名,既有消費端零改動
生產配置
k=7 · MNS=2 · GMU 0.50 · max_model_len 135168(132K);上表矩陣在 32K 下量得,生產跑滿 132K(模型原生上限 256K)
前端 relay
DGX 上的 fable-gateway.service(:18777,程式在 ~/bin/fable_gateway.py):工具呼叫正規化、503 退避、per-key 認證。併發上限 8——真正的限流由後端 MNS=2 負責,relay 不再卡第二層。
部署
bash /home/mavis/deploy_skinny.sh(drain → 開機 → 11 項 gate,不過即自動回滾)
回滾
bash /home/mavis/deploy38.sh —— IQ4_XS 舊棧完整保留,一鍵回去
檢查點
v100:/mnt/data/nvfp4/qwen38-uncensored-nvfp4-vl/(21 GB)
原始數據
G:\Projects\dflash2-furnace\kernel2\matrix_results_20260826.tsv
已知缺口與開放項:伺服器由 nohup 啟動,整機重開機不會自動回來(與舊棧相同),需要時補 systemd unit。矩陣唯一未測的有望格子是 k=3 · MNS=2。多卡滿載實驗前依 PSU 紀律先告知。

7.0前情數據:換引擎之前試過什麼

量測方法與本頁矩陣不同(舊棧用 tpsQ 腳本、DFlash2 用教師強制評測),數字只在各自區塊內可比,不能跨表直讀。

7.1 舊棧基準與 DFlash2 推測解碼(2026-08-25 收官)

舊棧 IQ4_XS + turboquant 2bit KV · TP4
配置essay 單流c8 混合 aggregate判定
舊棧基準67.746.4生產線
+ DFlash2(官方草稿器)60.2(−11%)−44%不上生產
+ DFlash2(Option A 蒸餾後)淨退步 −6%—停爐

DFlash2 在舊棧上完整跑通且無損(essay 從 0.000 修到 60.2),但打不平基準。損益平衡需要每步接受長度 ≈3.2(接受率 ≈0.85);essay 探針達 0.855,真實混合任務只有 0.54。拿生產流量蒸餾草稿器(1,719 筆離線特徵、SpecForge 改造、packed-anchor 訓練)反而淨退步 6%——官方草稿器本來就訓到位,缺口是結構性的。

DFlash2 接受率領域地圖(教師強制評測)
領域接受率vs 損益平衡 0.85
essay0.855達標
真實混合任務0.54差距大
zh_chat0.200天花板(中文 token 密度)

這張地圖是換路線的依據:問題不在草稿器權重,在 IQ4_XS 的 verify 稅——LUT 反量化把 kernel 天花板壓死。NVFP4 在暫存器內近乎免費展開成 FP16,量化格式的選擇位於 kernel 天花板的上游。這就是本頁引擎的立論,實測 +50%/2.9× 兌現。

7.2 七月—八月調校戰役(舊棧上打完的單點)

context 131K
解鎖到 131K(被 kernel 白名單擋住的;PARTITION_SIZE 環境變數是假旋鈕,完全無作用)
prefill 3.7×
128K 文件 prefill 快 3.7 倍——牆在分塊預算(batch 131072)不在注意力
解碼牆
每步 9.4ms 中 64% 是固定開銷不是頻寬;讀權重時其實有 77% 峰值——單點優化空間已盡
kernel2 v2
IQ4_XS 反量化提升 kernel:逐位正確但 1.00×——編譯器早已做完同樣的最佳化(誠實陰性)
GDN 旗標
掃三個預設關閉的旗標,三個都有關閉的理由,全部否決

完整原始數據與逐輪紀錄:G:\Projects\dflash2-furnace\(HANDOFF.md + kernel2/ + scratch/),DFlash2 證據檔 accept_base.json / accept_trained.json 在 dgx:~/furnace/。