V100 換引擎
生產推理棧由 IQ4_XS GGUF 換為 v100-skinny 引擎 + 自家 Uncensored NVFP4 檢查點。一晚三扇維護窗完成驗證、7 配置矩陣掃描與切換;同 IP、同 port、同模型名,消費端零改動。
- 離線定罪編譯鏈,裝完整 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%,差距只會被低估。
| 配置 | essay | code | prose | 中文 |
|---|---|---|---|---|
| k=7 · MNS=1 | 101.6 | 127.4 | 100.4 | 98.7 |
| k=7 · MNS=2 | 101.7 | 127.5 | 100.5 | 98.9 |
| k=7 · MNS=4 | 101.1 | 127.2 | 100.5 | 98.1 |
| k=7 · MNS=8 | 101.6 | 127.2 | 101.0 | 98.5 |
| k=3 · MNS=1 | 109.8 | 126.9 | 111.2 | 106.1 |
| k=3 · MNS=4 | 109.2 | 126.3 | 110.7 | 105.6 |
| k=1 · MNS=1 | 85.1 | 89.6 | 82.7 | 82.0 |
| 舊棧 IQ4_XS | 67.7 | — | — | — |
換 k 後首個請求有一次性 JIT 稅(k=3/k=1 首發 60.5/52.0),第二發即穩態;部署已預熱。
| 配置 | c2 | c4 | c8 | c8 最慢單條 |
|---|---|---|---|---|
| k=7 · MNS=1 | 109.2 | 105.3 | 105.5 | 19.4s |
| k=7 · MNS=2 | 92.0 | 135.2 | 149.8 | 13.7s |
| k=7 · MNS=4 | 114.4 | 17.9 | 19.0 | 107s |
| k=7 · MNS=8 | 132.0 | 18.8 | 27.9 | 57s |
| k=3 · MNS=1 | 111.2 | 109.9 | 110.0 | 18.6s |
| k=3 · MNS=4 | 93.3 | 19.8 | 20.6 | 96s |
| k=1 · MNS=1 | 83.6 | 82.4 | 82.5 | 24.8s |
| 舊棧 IQ4_XS | — | — | 46.4 | — |
| 配置 | 2,908 tok TTFT | 11,601 tok TTFT | 長文後解碼 |
|---|---|---|---|
| k=7 | 0.60s | 0.81s | 115–123 |
| k=3 | 0.55s | 0.68s | 112–124 |
| k=1 | 0.55s | 0.68s | 83–90 |
使用者實際走的那條路(8/28 補測)
上面三張表量的是基準路徑:直連後端的 /v1/completions,長生成、原始續寫。
使用者實際走的是 chat 端點 + relay(bot → relay → 後端),生成較短且帶思考段。
同一台機器、同一組配置,兩條路的數字差很多,所以分開列:
| 路徑 | c4 | c8 | c8 最慢 |
|---|---|---|---|
| completions · 直連(基準) | — | 149.0 | — |
| chat · 直連 | 95.9 | 102.2 | 12.9s |
| chat · relay(cap 8) | 92.9 | 95.8 | 15.1s |
| chat · relay(cap 2) | 77.5 | 95.3 | 14.2s |
| chat · relay(cap 1,原設定) | 68.3 | 63.9 | 19.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
7.0前情數據:換引擎之前試過什麼
量測方法與本頁矩陣不同(舊棧用 tpsQ 腳本、DFlash2 用教師強制評測),數字只在各自區塊內可比,不能跨表直讀。
7.1 舊棧基準與 DFlash2 推測解碼(2026-08-25 收官)
| 配置 | essay 單流 | c8 混合 aggregate | 判定 |
|---|---|---|---|
| 舊棧基準 | 67.7 | 46.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%——官方草稿器本來就訓到位,缺口是結構性的。
| 領域 | 接受率 | vs 損益平衡 0.85 |
|---|---|---|
| essay | 0.855 | 達標 |
| 真實混合任務 | 0.54 | 差距大 |
| zh_chat | 0.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/。