| 線 | 現在的數字 | 狀態 |
|---|---|---|
| 生產 · Qwen3.8-27B IQ4_XS · TP4 100.78.29.29:8601 · fable-27b · 設計上限八個併發使用者 · context 131,072(改了 kernel 才拿到,模型宣告 262,144)· prefill 預算 131072 @ util 0.85 | 單人 68.1 · 八人合計 260(每人 32.6,同時產出時) 混合負載、錯開到達:合計 155 文件問答:八人各 21.6K → 55 秒 · 八人各 64K → 5 分 47 秒 八人同時互動的文件上限抓 20–25K | 上線中 · 工具閘門 8/8 一小時 soak 過:0 失敗 · 記憶體零增長 |
| 換掉 Codex 部署的 Q4_K_M | c4 合計 86.7 → 107.8(+24%) · c8 133 → 154(+16%) | 零行 code · 換檔而已 |
| IQ4_XS 自寫 kernel(batched-MMVQ + tile-MMQ) | 兩顆都在 4/4 worker arm | 上線中 · 一小時 soak 過 |
| 推測解碼(MTP) | 最佳 39.96,仍是基準的 0.58 倍;四併發 0.32 倍 | 已收 · 預設關 |
| Qwen3.6-27B(前一顆生產) | 單人 68.88 · c4 合計 108.8 | 留著,未刪 |
Qwen3.8 和 Qwen3.6 是同一個架構(qwen35、65 blocks、48 個帶 GDN、866 個張量),所以 3.6 上調出來的設定全部直接適用。唯一真正的槓桿是量化檔:Q4_K_M 讓自寫的批次 kernel 掛不上去,換成同一個 repo 的 IQ4_XS 就拿到併發 +24%。而 max-num-seqs、capture 集合、util 三個旋鈕四種組合完全沒有作用 —— c4 合計全都落在 86.5–86.9。
推測解碼這條線的上限已經算得出來:c 修到 llama.cpp 的 0.29、α 維持實測的 0.52 → S = 1.18 → 81 tok/s,比基準好 18%,而且只在單人。目前不建議投入。
本頁分成七個大類,每類裡最新的在最上面,每段標了產出日期。打「約」的日期是回填的,可能差一天。總表裡標 GDN 開 的列是 8/6 那次變更之前量的 —— 當時是對的,但不代表現在的生產。被推翻的結論一律保留並標明,不刪 —— 看得到我錯在哪比看不到有用。
deploy38.sh 已改,備份 .bak_66560。(256, 512, 1024),而分割是從 context 長度推導的,131,072 推出 2048 → Unsupported decode partition_size: 2048。在 LAUNCH_BY_PARTITION 加一個 case 2048 重建 flash_attn_v100_cuda...so 就過了。三組同 session 對照 c1 68.24 / 67.78 / 68.00、全員串流 260.14 / 260.02 / 259.35、c8 合計 154.67 / 153.73 / 153.64 —— 全在 ±1% 噪音內。我原本擔心的 occupancy 代價沒出現:那三個 shared array 是 D*4 + PARTITION_SIZE*12 bytes,2048 仍在 48KB 內。重啟後生產實測 c1 68.19、c8 合計 152.92、工具閘門 8/8、KV 2,747,830 token。deploy38.sh 備份 .bak_81920_pre131k、原版 kernel 留在 .so.bak_pre2048。shm_broadcast.py 的 TimeoutError → EngineDeadError。源頭 envs.py:514 的 VLLM_EXECUTE_MODEL_TIMEOUT_SECONDS 預設 300 秒,而 128K 的 prefill 要 500.2 秒。調到 2400 直接過。deploy38.sh.bak_65536_util060。同時補了兩個工程修正:就緒檢查改成等自己 log 的 GPU KV cache size(原本 poll /health,殘留伺服器會立刻回答),以及 teardown 按 PID 清場 —— 「假設機器是乾淨的」今天害我撞了三次。--max-num-batched-tokens 釘在 65536:64K 剛好一塊裝得下,96K 和 128K 要切兩塊。把預算提到 131072(util 0.60→0.85 才有記憶體)之後:64K 56.6s(1,146)、96K 64.6s(1,487)、128K 132.5s(978)—— 128K 快了 3.8 倍,而本來就沒被切的 64K 完全沒變。拿掉分塊之後 prefill 大致是線性的,三個長度都落在 1,000–1,500 tok/s。一般流量零代價:c1 68.46、c8 合計 154.26。KV 2,719,744,仍高過八人各 131K 所需的 1,048,576。--scheduling-policy priority,而且 /v1/responses 帶得了每請求的 priority(/v1/chat/completions 沒這個欄位)。同一台伺服器兩輪對照:四份 64K 進行中,互動請求 priority −1 花 296.7s、priority 0 花 258.5s —— 而兩輪的「文件完成時間減請求耗時」都是 20.0 秒,正好等於我探針的發送延遲。不管什麼優先權,請求都在文件做完的那一刻才完成,和 fcfs 一樣。短提示流量沒受影響(c8 153.35 對 fcfs 153.45),但同樣四份 64K 文件 316.7s 對 fcfs 的 124.1s(各量一次)。這個 fork 沒有任何排程手段能讓互動請求繞過進行中的 prefill —— 分塊粒度和優先權都試過,都死,分塊還是嚴格劣化。唯一有效的是讓 prefill 本身更快,而那條已經拿到 3.7 倍。kill -9 之前,驅動還沒把 VRAM 還回來就啟動 → Free memory on device cuda:3 (17.32/31.73 GiB) on startup is less than desired。差點被我讀成「priority 不支援」。② 我在腳本正執行 deploy38.sh 時去改它 —— bash 是邊讀邊執行的,byte offset 一移位還原就死在半路,留下四個孤兒 TP worker 各扣 14 GiB,生產停機 27 分鐘。執行中的 shell 腳本絕對不能改,要改就複製成新檔名。No available memory for the cache blocks)。activation buffer 隨預算等比長,util 0.85 才買得回空間。VLLM_FLASH_V100_DECODE_PARTITION_SIZE=1024 是無效的。deploy38.sh 一直釘著它,所以我先問「釘 1024 是不是就不用改 kernel」。釘 1024 的 131,072 跑出來跟 2048 逐項一樣(128K 文件 499.7s 對 500.2s、c1 67.97 對 68.46、c8 154.26 對 154.27),數字分不出是「pin 生效」還是「pin 被忽略」。唯一能分的是換 .so:原版 kernel 在同樣釘 1024 下照樣報 Unsupported decode partition_size: 2048。vLLM 從 context 推導的值壓過環境變數,那行設定完全沒作用。deploy38.sh 的就緒檢查是 poll /health,而殘留的舊伺服器會立刻回答。結果它 1 秒就判定就緒,再去 grep 剛啟動的 log,報 IQ4XS_BMMVQ_READY on 0 workers <-- NOT ARMED —— 假陰性。改成等它自己 log 裡的 GPU KV cache size,第一次上場就正確擋下 DEPLOY FAILED。--tensor-parallel-size,每卡要讀的 bytes 差 4 倍:TP4 68.16(14.671 ms/步,3.66 GB/卡)、TP2 49.66(20.137 ms,7.31 GB)、TP1 32.85(30.441 ms,14.62 GB)。擬合 t = F + bytes/BW,三組配對的斜率 1.410/1.439/1.498 ms/GB、截距 9.19/9.41/9.83 ms。只用 TP4+TP2 擬合去預測 TP1 是 32.2,實測 32.85 —— 差 2%,而 TP1 沒有參與擬合。mul_mat_vec_iq4xs_batched 438 次 20.6 ms(約 700 GB/s)、Q5_K 65 次 2.85 ms(628)、Q6_K 輸出投影 1 次 1.57 ms(669)。三族 GEMV 都是頻寬有效率的,沒東西可撿 —— Q5_K 不是慢路徑,只是張量大(attn_qkv ×48 是 [5120,10240],output.weight 是 12.7 億參數)。FillFunctor<int> 每步 241 次、2.1 ms」是假的。我拿 200-token 那次的總呼叫數除以 208 步,但沒驗證「均勻分佈在每一步」這個假設 —— 我只對 GDN kernel 驗過(48.2 對 48 個 block)。改跑 50 token(58 步)對照:每步比值應該是 0.279。GEMV 0.289、quantize 0.292、triton 0.256、GDN 0.282、Q5_K 0.292、RMSNorm 0.263 —— 六個全中。而 FillFunctor 是 50,095 → 160,比值 0.003。它是一次性爆發被我平均進每一步,不是每步成本。quantize_q8_1_b 432(完全 1:1)、Q5_K 64(每層一次)、GDN 48.0(正好 48 個 block)。時間:IQ4_XS GEMV ~20.6 ms、Q5_K 2.67、Q6_K 1.57、triton_poi_fused_0 2.09、quantize_q8_1_b 0.96、GDN 0.71、triton_ 0.57、RMSNorm 0.63、conv1d 等 0.35。可攻擊的總量從 7.5 ms 修正為約 5.5 ms,而最大的一項是 triton_poi_fused_0(78 次/步,約每層 1.2 次),不是我原本說的整數填充。INPUT_PROJECTION_OP 680 B(不同)68.43 / 153.04;INPUT_CORE_OP 667 B(不同)68.54 / 154.14;輸入+輸出一起開 670 B(不同)68.35 / 156.24。三個全部改變數值,單人速率全在雜訊內,最好的情況是只在八人合計上 +1.8%。和輸出側單獨開的 +1.2% 是同一個故事,也就是它們預設關的原因。QWEN_GDN_OUTPUT_PROJECTION_OP=1(把每層的 reshape/gated RMSNorm/v-reorder/out_proj 融成一個自訂 op):輸出變成 687 bytes / 140 詞,不同,合計只 +1.2%。拿數值變動換 1.2% 不是吞吐選項。GDN_Z_CONTIGUOUS + MIXED_QKV_CONTIGUOUS:輸出崩潰成 2 bytes,而速率一點沒動 —— 連取捨都不算。GDN_EMPTY_CORE_OUT 就是我前一晚手工改 torch.zeros→torch.empty 的那個開關,早就存在而且早就預設關。我重新發明了一個已被否決的東西 —— 改模型檔之前先盤點旗標。warps2 兩次 153.47 / 157.41。所有配置間的差距都躺在這個散佈裡,而且方向還在兩輪之間翻轉。checksum 全程相同,所以計算沒變,變的是量測條件。沒有那三個重複 arm,我會報「warps=2 快 19%」並附一個很像樣的 Volta 解釋。(B=8, T=1),結果每個配置都在 kernel 第 177 行掛掉 —— 包括 default。我當下讀成「BV 是假旋鈕」,實際意思是「我的輸入形狀錯了」。真實解碼路徑會把變長攤平:帶 cu_seqlens 時 q.shape[0] 必須是 1,還要 ssm_state_indices 和 use_qk_l2norm_in_kernel=True。對照組也失敗,就不是被測物失敗。qwen3_5.py:471 的 torch.zeros 換成 torch.empty 之後:生成直接斷掉,而且速度一點沒變。同提示、temperature 0:665 bytes 的完整答案變成 **First 12 —— 10 bytes、兩個詞就停。填充列沒有被遞迴核心寫滿,垃圾會傳出去,那個清零是承重的。而吞吐 c1 68.47 對 67.61、c8 合計 153.75 對 153.60 —— 噪音。empty 真的生效了 —— 既然生效卻沒有加速,那 core_attn_out 就不是那 2.09 ms。那個數字是把 triton_poi_fused_0 的呼叫數加總來的,而 inductor 的 kernel 是每個圖各自從 0 編號 —— 名字不唯一。78 次/步是好幾個不同清零 kernel 的總和,我拿掉的只是其中便宜的一個。2.09 ms 是真的,但來源未歸屬,而我唯一的線索(kernel 名字)已證實不具鑑別力。quantize_q8_1_b:它 blockDim.y = 1、每個 block 一列,amax 是沿列做 width-32 shuffle —— 列與列獨立,填充列的 inf 污染不到別列的尺度。那個檢查是對的,但失敗來自下游而不是量化。查對前提不保證查到的是對的前提。triton_poi_fused_0 的 Triton IR 是純 tt.store 寫 0.0 到 f16 緩衝區、沒有任何 load;而 inductor 自己的呼叫端寫著 rand_strided((8, 12, 128), dtype=torch.float16) = 12,288 元素 =(捕捉尺寸 8, 48 v_head ÷ TP4, head_dim 128),正是 qwen3_5.py:471 的 core_attn_out。注意 inductor 的 kernel 是每個圖各自從 0 編號,所以那個名字不唯一,78 次/步是多個清零 kernel 的總和。不是安全的一字改成 torch.empty:緩衝區按填充後的捕捉尺寸 8 配置,實際 token 可能只有 1,填充列不會被遞迴核心寫到。if mx > 1e-3 —— nan > 1e-3 在 Python 是 False,一路放行,max(0.0, nan) 也保留 0.0 把最壞值報成 0。我整天在講「守門要拒絕出數字」,然後自己寫了一個對 NaN 無感的守門。quantize_q8_1_b 每步 440 次,和 438 次 GEMV 是 1:1 —— 每個矩陣向量積前都重新量化一次激活,而 llama.cpp 那邊早就融進 mmvq。單檔、自足、驗證條件明確,已附線上真正 runtime 編譯的 iq4_mmvq_batched.cu 與正確的參考 header 派給委員會。整數填充沒派,而那個決定後來被證明對得莫名其妙 —— 我當時的理由是「來源未定位」,實際上它根本不是每步成本。不派自己說不清的東西這條規矩救了我一次,但救的方式和我以為的不同。1→2→4→8→8→4→2→1 連續跑、arm 之間不冷卻:每人 68.13 / 56.36 / 47.87 / 34.01 上去,34.03 / 47.91 / 56.40 / 68.64 回來 —— 差異全在 0.1% 內。合計 59.4 / 83.8 / 107.7 / 152.8:八人時每人只剩一半速度,但總產出是單人的 2.24 倍。使用者數量本身不留後遺症,人走了速度立刻滿血。BMMVQ 4/4 真的掛上。結果 c8 全員串流 260.15 → 261.90(+0.7%)、c4 合計 107.40 → 107.28、c1 持平 —— 全在雜訊裡(同組態不同輪已量到 8% 變異)。「權重在迴圈裡被重讀八次」描述的是源碼的形狀,不是時間花掉的地方 —— L1 本來就吸收掉了。vec_dot 三參數變四參數、__dp4a 變 ggml_cuda_dp4a,編不過;④ 提示裡的雙引號被 PowerShell 吃掉,argparse 拒收。委員會的診斷從第一輪就沒變過。BMMVQ x/4 —— kernel 沒掛上時伺服器照樣起得來、照樣有數字,那個數字看起來會像「改動沒效果」而不是「改動沒生效」。reasoning,非串流叫 reasoning_content —— 同一個東西兩個名字。delta.content 一直存在但是空字串,所以只試那兩個非串流名稱的檢查會安靜地數到 0 個 token。把第一 KB 用 repr() 印出來花十秒,而我先前猜了兩次格式。--image-min-tokens 1024 —— 不加的話 96×96 的圖只換到 81 個 prompt token,加了之後是 1,103。bt 65536 之後仍然成立。合計吞吐 18.7 tok/s,對短提示八人的 155 —— 長 context 併發是八倍代價,時間幾乎全在 prefill。loaded multimodal model、生成 120 個 token。我讀 content 讀到空,內容在 reasoning_content。本 session 第三次讀錯欄位,而我還為這件事寫過記憶。max-num-batched-tokens 65536 換來的。Q4_K_M,我先前為 IQ4_XS 寫的兩顆批次 kernel 掛不上去;同一個 HF repo 就有 IQ4_XS 版(15.3 GB)。換檔後兩顆 kernel 在 4/4 worker arm:c4 合計 86.74 → 107.75、c8 133.14 → 153.86,單人 69.42 → 68.15(−1.8%,量化較小的代價)。工具閘門 8/8、真實文本連貫。回滾:/home/mavis/codex_qwen38_launch.saved,Q4_K_M 檔原封不動。max-num-seqs(16 vs 8)、capture 集合、util(0.60 vs 0.50)四種組合,c4 合計全部落在 86.5–86.9,c8 全部 133.0–133.6。那 25% 從頭到尾都是量化檔,不是排程設定。--max-num-batched-tokens 16384 → 65536:21.6K 冷 prefill 19.8s → 15.4s(−22%),decode 不動。和我先前在 106K 上量到的 −26.3% 同一個量級,已套進生產。torch._dynamo.exc.Unsupported: Attempted to call function marked as skipped — module: time, qualname: perf_counter。整條 GDN 路徑在 torch.compile 的 fullgraph 捕捉區內,任何 Python 層計時器都會讓捕捉失敗、引擎起不來。加上 vLLM 自己的 /start_profile 在 TP4 下寫不出東西(worker 是獨立 process,只武裝 rank 0)—— 這台機器上要看進推測步裡面只剩 nsys --output=…%p 一條路,那不是備案,是唯一解。已還原並重跑驗過 39.66 對插樁前的 39.84,NOSYNC 八處仍在 —— 能 parse 不等於伺服器起得來,一定要重跑一臂。/mnt/data/backups/windows-v100/models_muse、/mnt/data/llama.cpp-new),要用隨時能起。spec_sequence_masks is None 那個守衛從「候選之一」變成主嫌 —— 它正好是一條純引擎側、與模型架構無關的路徑選擇。(限制:llama.cpp 兩顆都是單卡、vLLM 是 TP4,跨了引擎也跨了拓樸;但 llama.cpp 內部那組同引擎同卡數的比較是乾淨的,足以殺掉遞迴假設。)n_max=16 exceeds the trained block size 16 -- clamping to 15。預設最好。reasoning_content 而不是 content(打到 finish_reason: length 時還在推理段),所以我的文本檢查印了 uniq_word_ratio 0.00 DEGENERATE —— 吞吐數字是有效的,壞的是我讀錯欄位。實際內容 1585 字元、連貫英文。--spec-type 預設是 none,我只給了 --model-draft —— 兩個草稿臂會在推測關閉下跑成基準。第二輪:對 instruct 模型打了 /completion,沒有東西套 chat template,生成到第 2 個 token 就停(predicted_n=2,內容是 " -"),而那三組 tok/s 是在兩個 token 上量出來的。我的退化檢查放行了 —— 一個字的輸出 uniq_word_ratio 是 1.00,滿分。那個檢查是為了抓重複寫的,而重複是相反的失敗。已改成:先斷言實際生成的 token 數,不到 100 就標 INVALID、不准印速率。block_size=16, mask_token_id=201818, n_extract=5 —— 它一次提一整段、靠 mask token 填空。我拿量 MTP 的條件(n-max 1)去套等於把它掐死在一格,所以 accepted 0 / generated 1。重跑改成給它自己的區塊大小。另外這顆官方 GGUF 有個真警告:special_eot_id is not in special_eog_ids,tokenizer 設定可能有問題。FLASH_GRAPH+SPEC_CORE_OP:c1 39.84(掃描時 39.96,可複現)、接受率 52.3%、工具閘門 8/8 PASS、真實文本 uniq_word_ratio 0.79 沒有退化。但 c4 合計吞吐只有 34.80,無推測基準是 108.80 —— 0.32 倍,TTFT 5.03s。批次已經把成本攤平的地方,推測解碼只會添亂。就算做到最好,這條線也只在單人有意義,而單人現在也才 0.58 倍。生產維持不開推測的 68.88。8e7f22b、gcc 12.4.0、CUDA arch 70、ninja -j2),編譯日誌裡有 common_speculative_impl_draft_dflash —— DFlash 推測支援確實在裡面。/opt/llama.cpp 覆核過完好,對照組數字仍可比。FLASH_V100_0DOT3_COMPILE_GRAPH 單獨對解碼沒用(34.55,c 2.020),但疊在 SPEC_CORE_OP 上把它造成的傷害修掉 —— c1 39.96、接受率回到 52.3%(全場最高,與不開推測同水準)、TTFT 從 4.42s 回到 0.68s。委員會把這個旗標排第一是對的,但理由錯了:它救的是 graph 與 prefill,不是解碼主力。content-length 逐一相等,731 個張量目錄讀到底)。它是我這條線缺的對照組:52 個 block 全是注意力、沒有 SSM(對照 Qwen3.6 的 65 block 裡 48 個帶 GDN 遞迴),KV head 32:2、滑動窗 2048、context 131072,而且官方另出一顆 DFlash 草稿模型。一顆沒有遞迴狀態又自帶草稿模型的 30B,正好能分辨 c=2.01 是架構特有還是引擎通病。官方沒有純 Q4,只有 kquant-17gb(16.8 GB,為 24 GB 消費卡而生)和 kquant-dynamic(19.7 GB)—— 取後者,我們是 32 GB×4,沒有理由為遷就別人的卡吃品質損失。/opt/llama.cpp 是 7/6 建的,muse-glimmer 支援是 8/10 才合併(#26841,b10353 起)—— 認不得這個 arch。要跑得建新的,而且不動 /opt/llama.cpp:今晚 llama.cpp 那組 α 0.722 / c 0.298 的數字全出自它,重建會毀掉可比性。新的建到 /mnt/data/llama.cpp-new。SPEC_CORE_OP c1 39.62(c 2.043→1.595,+15.5%)、003_SPEC_CORE_OP 38.49(c 1.593)。兩條路徑的 c 幾乎同值不是巧合,應該接到同一個改動;SPEC_CORE_OP 吞吐較高只是因為它的接受率沒被拖低(49.3% vs 44.9%)。其餘五個全在 c≈2.0,是雜訊。FULL_FORWARD 排在 SPEC_CORE_OP 前面(理由是名稱暗示同時修候選 ①②)—— 實測反過來。FULL_FORWARD 不但沒用(34.03,c 2.036),TTFT 還從 1.46s 惡化到 2.86s。排序要靠量,不靠命名。qwen_gdn_linear_attn.py 的 4365 與 5706 兩處都有 and attn_metadata.spec_sequence_masks is None,所以用 packed_decode 那條路在推測步永遠不會被選中,一律落到通用的 _forward_core,而那裡把 token 拆成 spec / non-spec 兩組分別跑。conv 那半有守(num_prefills>0 / elif num_decodes>0 / else None),單人時沒有空跑;成本在後面的 gating 與遞迴段。同一個守衛還帶出委員會不知道的交互作用:FUSED_SIGMOID_MIXED_QKV=1 會讓 mixed_qkv_decode_requested 為真,也會關掉這條快速路徑。start_profile 只武裝 rank 0 的 context,跨不過 process 邊界。這解釋了我三次失敗,替代方案是 nsys --output=…%p 每個 PID 各自出檔。委員會的價值在候選假設與這條路,不在優先順序。(Codex 這次 thread start failed 無實質產出,融合時整份排除,實際是 GLM-5.2 與 MiniMax-M3.0 兩份。)VLLM_TORCH_PROFILER_DIR;第二次設了但讀取端掃的是自己的預設目錄;第三次目錄對上了,TP4 下 /start_profile 之後那個目錄仍然是空的,fork 自己的 phase report 也沒進 log。三次失敗長同一個樣子 = 問題在假設不在實作(RULES 5),所以不再想辦法「看進步驟裡面」,改直接量哪一種實作最便宜 —— 這個 fork 帶了八個切換推測路徑的旗標,一臂一個,c 用同一 session 自己的無推測臂算,不從記憶拿。p-min 降到 0、完全不篩,llama.cpp 在 γ=1 仍有 72.2%,草稿/步 = 115/(200−83) = 0.98,等於 γ,篩子根本沒動。vLLM 同樣 γ=1 是 50.4%。同一顆 GGUF、同一顆 MTP 頭、同一張卡 —— vLLM 真的把草稿頭餵得比較差,那是個可定位的 bug 不是架構宿命。(順帶更正我自己的量測腳本:它印的「drafted per generated token」分母錯了,該用步數 —— 接受的草稿會讓一步產兩個 token,把分母灌大,害兩個沒篩的臂被標成 FILTER IS ACTIVE。)target_hidden_states = hidden_states[...],而 qwen3_next.py:953 是 hidden_states, _ = self.norm(...) 之後才 return;MTP 頭接著又對它做一次 pre_fc_norm_hidden。llama.cpp 的對應圖在 src/models/qwen35.cpp(源碼在機器上),下一輪對照。這條排在 c 後面。nonzero() 拿掉,MTP 單人 +6.8%(32.42 → 34.64),四人 +2.5%,接受率不動(51.5%),工具閘門 ALL_PASS。那段用布林遮罩挑推測列,每個 x[mask] 都是 nonzero() —— 它得把命中數讀回 host 才知道輸出多大,是一次完整的裝置同步,而且一次 forward 好幾個。批次同質時遮罩全 True,那個挑選就只是複製一份。判斷用的計數呼叫端已經在 CPU 上算好,問它不花錢。c 從 2.22 降到 2.01。預設關,VLLM_SM70_GDN_NOSYNC=1 才開。-ctk turbo2,但 /opt/llama.cpp 是上游 b0.15.3 不是 atomic-llama fork,turbo2 是 fork 專有型別。旗標檢查有做、KV 型別沒檢查。改 f16 重跑中。p-min 0.5 讓它信心不足時根本不提草稿,那些注定被拒的位置沒進 draft_n,百分比是在篩過的子集上量的;② vLLM 真的把草稿頭餵錯了。把 llama.cpp 的 p-min 降到 0 重量,掉到 ~50% 就是①(那接受率從頭到尾都不是差別,差別是 c),留在 ~72% 就是②(那是個值得修的 bug)。跑了。OFFICIAL 那顆)metadata 全是 qwen35、65 blocks、48 個帶 SSM、ssm.state_size 128。看檔名不算,要看 metadata。spec_state_slot_selectors 這類機制,但沒找到 ring buffer —— 是不是半套還沒確認。memcpy32_post 1879 次、scatter_gather 268、indexSelect 268、CatArrayBatchedCopy 268 —— 全是狀態簿記,不是算術。論文的 1.8x / 3.4x 全部量在純 transformer 上。mul_mat_vec_q 把 token 位置放在 blockIdx.y —— 每個位置各開一組 block、各讀一次 45 MiB 的權重。profiler 的證據無可辯駁:同一顆 kernel 在有無推測下 launch 次數幾乎相同(2489 對 2468),時間卻差 60% —— 同樣次數的呼叫、每次做兩倍的搬運。改成一個 block 帶 NCOLS 個位置、權重靠 L1 重用後,每位置成本從 0.086→0.074(幾乎不降)變成 0.084→0.032。M=1 持平所以不傷單人;九種權重形狀 × 每個 M 的誤差與原 kernel 一致到印出的精度。(1−α^(γ+1)) / ((1−α)(γc+1)),Corollary 3.9 的損益平衡是 α > c。我們實測 α = 0.478、c = 2.04 —— 差 4.3 倍。他們論文裡 c 的範圍是 0.007–0.073,我們是他們畫過最壞情況的 28 倍。而且 Table 2 自己就在證明 α 不是槓桿:接受率最高的 T5-large(0.82)加速最差(1.7x),因為它的 c 最大。tool_choice=required 問。兩顆模型正好相反:生產 9B 在 required 能用、auto 回空;27B 反過來。兩顆都不壞,各自在一種模式下吐出完全正確的呼叫。失敗那側都是同一個樣子:tool_calls 空的,而 reasoning 欄位裡模型正在說「我應該呼叫 lookup_inventory」—— reasoning parser 把開頭那段吃掉,tool parser 就什麼都不剩。閘門已改成兩種都問、任一成立即通過並指名哪一種;生產與 27B 現在都 ALL_PASS。我為這個假訊號往 kernel 和 capture 各查了一輪,那些成因不存在。t5a_no_think_plain 在生產通過,但模型開頭就是 Thinking Process:\n\n1. **Analyze the Request:** —— 那個檢查只找字面的 <think> 標籤,這種寫法直接走過去。沒有動它:把生產監控轉紅是關於既有行為的決定,不是量測 bug。席次 ×(1+draft) 個 token,而 vLLM 只保留能被那個倍數整除的尺寸。同一組 [1,2,4,8,16,32]:不開 MTP 捕捉到 6 張 decode 圖,draft=1 只剩 5 張(1 不能被 2 整除)。真正可用的席次上限因此是 16,c24 和 c32 全走 eager —— 兩者步時間 233.5 / 234.1 ms 幾乎同值,就是這個原因。draft=2 要 3 的倍數,那組裡一個都沒有,所以它捕捉到 0 張。[1,2,4,8,16,32,48,64] 後捕捉數 7/7,c24 每人 4.28 → 8.17(合計 97 → 181)、c32 4.27 → 7.60(合計 128 → 225)。t3_tool_call 也是三臂全失敗含舊集合,同樣與 capture 無關。所以 capture 對齊是乾淨的淨勝:c24 3.76 → 8.16(2.17x)、c32 3.66 → 7.59(2.07x),c1/c8/c16 在雜訊內不動。ggml_mul_mat_a8 對 IQ4_XS 完全沒有運算。profiler 顯示它只發射 quantize_q8_1,之後沒有任何 matmul,輸出緩衝區沒被寫入。它先前之所以通過數值比對,是 torch 的 caching allocator 把剛算完參考值的那塊記憶體回收再配給它 —— 讀到的是上一次的殘值。因此把 IQ4_XS 加進 MMQ_QUANT_TYPES 來修批次崩塌,會讓模型安靜地吐垃圾而且不報錯。blk.0.ffn_up,17408×5120):反量化整顆權重 0.663 ms + cuBLAS 0.441 ms = 1.106 ms,而且不隨 M 變;同尺寸的 K-quant 走真 MMQ 在 M=16 只要 0.355 ms。mmvq 則是線性的,M=1 0.086 ms、M=64 4.169 ms,與反量化路徑的交叉點在 M≈15 —— 而 gguf.py 的門檻寫死在 8,偏保守。QR4_XS/QI4_XS —— 那是 8 和 8,描述的是 MMVQ 怎麼把超區塊分給執行緒,不是 MMQ 的 tile 幾何。tiling 要的是 QK=256 / QR=2 / QI=32,現在綁了 static_assert 讓它不能再漂走。kvalues_iq4nl 查表從 vec-dot 搬到載入時做一次(x tile 會被重用 mmq_x 次,等於同一組查表重做 64 遍)又多 1.5 倍:對比 dispatcher 現行選擇,M=8 2.48x、M=16 4.95x、M=32 2.57x、M=64 1.95x。os.getenv 每次矩陣乘都跑,提到模組層級後只移動 0.3% —— 猜錯了。而且用算的就排除得掉:條件短路後開關只差兩次整數比較(約 25 µs),實際差距是每步 0.6 ms。成因在別處,還沒找到。操作結論:只有單人就別開,有併發就開。mul_mat_q40_dec_y、mul_mat_q40_x64,註記 vs fork 1.26–2.35x),runtime 編譯、env 控制、shape gate 全都在,但都卡在 qweight_type == 2,也就是 Q4_0 專用。缺的只有 load_tiles_iq4_xs 和 vec_dot_iq4_xs_q8_1_mul_mat 兩個函式 —— IQ4_XS 的 32 元素子區塊剛好等於一整個 Q4_0 區塊(都是 4 個 int、低 nibble 都是元素 0–15),所以 tile 幾何可以原封不動沿用。已寫好,待驗。llama-bench 的 -ctk/-ctv 是位置配對,先前那批其實只量到 V —— 重跑後:V 量化花 8% 生成速度,K 量化再花約 30%(K 在 QKT 內圈)。layer 38.9(等於單卡)、row 22.0(慢 43%)、tensor NCCL 崩潰。llama.cpp 在這台沒有可用的 TP4,所以它只能當參照組。llama-bench 中止是因為 tbq3 必須要有 FA。blk.64 gate 到 Q8_0 重做 requant 這條路不用走了。force,另有 auto 會自行武裝;8/8 與 8/9 的伺服器 log 都印了 guard armed (force=False, auto=True)。-sm tensor 崩潰根因是 NCCL 版本:系統的 2.18.3+cuda12.0 對上 CUDA 13 driver,查 entry point 失敗被誤譯成 「CUDA driver is a stub library」。LD_PRELOAD vLLM venv 的 2.27.5+cuda12.9 就跑起來了。-ts 是斜線分隔,我用了逗號 —— 在 llama-bench 那是「掃多組設定」,結果四個臂全部跑在 GPU 0。破綻是 pp512 四種配置差 0.15%。重跑中,這次會驗每張卡有沒有吃到權重。ggml_mul_mat_a8 對它是空轉返回(M=2 到 512 耗時全是 0.03 ms,換算 3000 TFLOPS,V100 峰值 125)。--enable-auto-tool-choice 和 --tool-call-parser,生產的啟動腳本一直都有。