先说结论#

decode 一律只算生成阶段的 token 速度，不含 prompt 处理。draft 3 是单流最优；draft 数再往上接受率掉得比批次收益快。同一套文件在 M4 Pro 24GB 的 Mac 上 Metal 后端跑，MTP 反而把 12.9 tok/s 拖到 7.6。

为什么不是 vLLM#

第一反应是用 vLLM。原因很朴素：vLLM 的批处理吞吐通常远高于 llama.cpp，而且 Qwen 官方 recipe 页面写着 --speculative-config '{"method":"mtp","num_speculative_tokens":3}' 一行就能开 MTP。

在那台机器上装完 vLLM 0.28.0 之后发现两个硬事实。

第一，vLLM 0.28.0 已经没有 GGUF loader 了。整个包里 grep -rli gguf 只命中 3 个文件，全是顺带提到，没有 gguf_loader.py，也没有装 gguf 依赖包。手头这两个 GGUF 文件在 vLLM 上没有任何入口，独立的 MTP GGUF 就更不可能当 draft 用。

第二，vLLM 走 MTP 只能用 HF 权重，而 27B 的 int4 HF 权重装不进 16GB。查了一圈 HuggingFace 上的量化版本：

int4 版本普遍 19-20GB，因为 embedding、lm_head、vision encoder 都留 BF16，光 embedding + lm_head 就 5GB（5120 × 248320 词表 × 2 字节 × 2）。3-bit 的 GSQ 版能装下，但它要求 vLLM 0.27.1 加一个源码补丁，README 又明说不支持投机解码。

结论：这两个 GGUF 文件只有 llama.cpp 一条路。vLLM 上要玩 MTP 得换 HF 权重和更大的卡。

llama.cpp 的 MTP 参数#

本机 brew 装的 llama.cpp（build 10621）已经带 MTP 支持，这是 2026 年 5 月 PR #22673 合并进去的。相关参数：

1
llama-server -m Qwen3.8-27B-UD-IQ3_XXS.gguf \
2
  --spec-type draft-mtp \
3
  -md mtp-Qwen3.8-27B-Q4_0.gguf \
4
  --spec-draft-n-max 3 \
5
  -ngl 999 -ngld 999 -fa on

--spec-type 的可选值里有 draft-mtp，-md 指定 MTP 头文件，-ngld 控制 draft 模型的 GPU 层数。

那台 Ubuntu 机器上没有 curl、gcc、cmake，只有 wget。用 llama.app/install.sh 装预编译版，脚本依赖 curl，我写了个 curl 到 wget 的 shim 骗过去，它探测到 sm89 后下载了 CUDA 版（build 10679）。这一步埋了后面最大的坑。

统一的测速脚本#

llama-bench 不支持投机解码，所以写了个客户端：起 llama-server，等 /health，发固定的 4 条 prompt（贪心，256 token），从响应的 timings 字段读数据，再发 4 条并发。

timings.predicted_per_second 是 llama-server 自己算的 decode 速度，只包含生成阶段。draft_n 和 draft_n_accepted 给出投机解码的接受率。并发那一组我一开始用”总 token / 最长那条请求的 eval 时间”来算合计吞吐，后来发现这个算法在请求没有真正重叠时会严重高估（有一组算出 210 tok/s，实际墙钟只有 100），改成按墙钟算。并发吞吐只能看墙钟。

第一轮：Mac 上 MTP 反而变慢#

M4 Pro 24GB 上先跑：

接受率和 CUDA 上完全一样，但速度倒退。MTP 每一步要跑 3 次 draft 头前向（每次都要过 248K 词表的 lm_head），再用主模型验证 4 个 token。Metal 上这套开销没有被批处理收益抵消。IQ4_XS + MTP 在 32K 上下文下还会 Metal 显存不足（kIOGPUCommandBufferCallbackErrorOutOfMemory），4 槽 × 4K 也不够，只有单槽才起得来。

Mac 这条线到此为止：Metal 上别开 MTP。

第二轮：CUDA 上 IQ3_XXS 神秘崩溃#

CUDA 那台上 IQ4_XS 基线一次跑通，单流 39.8 tok/s。IQ3_XXS 却在第一个请求就崩：

1
llamacpp/ggml/src/ggml-cuda/ggml-cuda.cu:107: CUDA error
2
E CUDA error: the resource allocation failed
3
E   current device: 0, in function cublas_handle at ggml-cuda/common.cuh:1503
4
E   cublasCreate_v2(&cublas_handles[device][curr_stream_no])

怪在几点：IQ3_XXS 比 IQ4_XS 小 3GB，显存不可能更紧；手动起 server 用 5 个 token 的 prompt 正常生成，换成十几个 token 的 prompt 就崩；GGML_CUDA_FORCE_MMQ=1 无效。

这个组合指向一个方向：某种张量类型没有 MMQ 内核，批次一大就退到 dequant + cuBLAS，而这个预编译包的 cuBLAS 初始化是坏的。IQ4_XS 全程走 MMQ 从没碰过 cuBLAS，所以没事。

用 gguf-py 把两个文件的张量类型分布 dump 出来：

1
=== Qwen3.8-27B-UD-IQ3_XXS.gguf
2
IQ3_XXS     120
3
IQ3_S       106
4
Q8_0         98
5
IQ4_XS       45
6
IQ2_S        35
7
Q4_K         26
8
IQ2_XXS      24
9
IQ2_XS       12
10
Q2_K         10
11
IQ1_S         8
12
Q5_K          7
13
Q3_K          6
14
Q6_K          6
15
IQ1_M         3   e.g. blk.13.ffn_gate.weight, blk.14.ffn_down.weight, blk.15.ffn_down.weight
16
=== Qwen3.8-27B-UD-IQ4_XS.gguf
17
（没有 IQ1_M）

unsloth 的 UD 量化是按层敏感度混用类型的，IQ3_XXS 这一档里有 3 个张量被压到了 IQ1_M。而 llama.cpp 的 ggml-cuda/mmq.cuh 支持的类型列表里有 IQ1_S 却没有 IQ1_M。三个 18MB 的张量，把整个文件在这台机器上判了死刑。

1
-ot 'blk\.13\.ffn_gate\.weight=CPU,blk\.14\.ffn_down\.weight=CPU,blk\.15\.ffn_down\.weight=CPU'

1
D tensor blk.13.ffn_gate.weight (18 MiB iq1_m) buffer type overridden to CUDA_Host

1
llama-quantize --allow-requantize \
2
  --tensor-type 'blk\.13\.ffn_gate\.weight=Q2_K' ... in.gguf out.gguf COPY

1
for t in reader.tensors:
2
    if t.tensor_type == T.IQ1_M:
3
        f32 = gguf.quants.dequantize(t.data, t.tensor_type)
4
        q = gguf.quants.quantize(f32, T.Q8_0)
5
        converted[t.name] = q
6
        writer.add_tensor_info(t.name, q.shape, q.dtype, q.nbytes, T.Q8_0)
7
    else:
8
        writer.add_tensor_info(t.name, t.data.shape, t.data.dtype, t.data.nbytes, t.tensor_type)

1
OK prompt_n 58 pp tok/s 418.6 decode tok/s 46.6
2
cuda errors: 0

第三轮：参数扫描#

写了个 sweep 脚本，每组配置起一次 server 测一遍，一组 40 秒左右。

GDN 循环状态缓存：显存公式#

扫描里 draft 5 以上和 MTP np8 全部 OOM，一开始以为是运行间隙显存没释放，加了等显存回落的逻辑还是 OOM。verbose 日志里找到了真相：

1
draft-n-max 3, np 4:  CUDA0 RS buffer size =  2394.00 MiB
2
draft-n-max 5, np 4:  CUDA0 RS buffer size =  3591.00 MiB

RS 是 recurrent state 缓存。Qwen3.8-27B 是 64 层混合架构，48 层 Gated DeltaNet 加 16 层全注意力。GDN 层没有 KV cache，只有一份固定大小的循环状态，每个序列约 150MB。投机解码验证 draft 时如果某个 token 被拒绝，状态要回滚，所以 llama.cpp 给每个序列留了 (draft + 1) 份状态快照：

1
RS 显存 = 槽数 × (draft + 1) × 150MB

2394 / 4 / 4 = 150，3591 / 4 / 6 = 150，严丝合缝。--ctx-checkpoints 管的是另一套东西，对它没作用。

整块卡的显存账（IQ3_XXS 重写版）：

所以 draft 数和并发数互相挤占：单槽 draft 3 只要 600MB RS，剩下约 3.5GB 全给 KV，q8_0 能放 80K 上下文（96K 实测 OOM，差 384MB）。

长上下文下的 decode#

部署配置（单槽、draft 3、KV q8_0、80K 上下文）实测：

掉速比普通 dense 模型温和得多，因为只有 16 层全注意力随上下文线性增长，48 层 GDN 是常量。

内存 offload 是死路#

那台机器有 62GB 内存，自然会想能不能拿来换上下文。两种方式都测了。

KV cache 留内存（--no-kv-offload，上下文开到 262K）：

短上下文就从 84 掉到 26。这个开关把 GDN 的循环状态也一起放到了 CPU（显存正好少了 RS + KV 那 1GB），48 层 GDN 每步都要跨 PCIe 读写状态。

权重部分放内存： -ngl 56（8 层放 CPU）开 128K 还差 512MB OOM；-ngl 48（16 层放 CPU）能起，decode 1.9 tok/s。CPU 层加上 MTP 每步 4 个 token 的验证批次，内存带宽直接被打穿。

结论：这块卡上这个模型，KV、GDN 状态、权重必须全在显存里。想超过 80K 只有换更小的量化（UD-IQ2_XXS 约 8GB，估计能到 170K，质量会掉）或者砍 draft（每少一个省 150MB，约多 4-5K 上下文）。

功耗#

用 nvidia-smi 每秒采样 GPU，RAPL 计数器读 CPU 包功耗，decode 512 token 一轮：

一个反直觉的点：MTP 下的 decode 不是纯显存带宽瓶颈。显存频率四档都是 10251 MHz，SM 频率一降速度就跟着掉，说明 draft 头前向和多 token 验证批次让它偏算力受限。普通单 token decode 那套”压功耗不掉速”的经验在这里不成立。220 W 是能效甜点，速度掉 7%，每 token 能耗降 14%。

最终部署#

1
COMMON=(-m Qwen3.8-27B-UD-IQ3_XXS-noiq1m.gguf --spec-type draft-mtp -md mtp-Qwen3.8-27B-Q4_0.gguf
2
        -ngl 999 -ngld 999 -fa on -ctk q8_0 -ctv q8_0
3
        --host 0.0.0.0 --port 8080 --cache-reuse 256 --no-webui --metrics)
4
case "$MODE" in
5
  single) exec llama serve "${COMMON[@]}" -np 1 -c 81920 --spec-draft-n-max 3 ;;  # 80K，84 tok/s
6
  dual)   exec llama serve "${COMMON[@]}" -np 2 -c 65536 --spec-draft-n-max 3 ;;  # 2×32K
7
  batch)  exec llama serve "${COMMON[@]}" -np 6 -c 36864 --spec-draft-n-max 2 ;;  # 6×6K，178 tok/s 合计
8
esac

套上 systemd 服务，显存占 15.7GB，实测请求 decode 87 tok/s、接受率 65%。

经验总结#

• vLLM 0.28 没有 GGUF loader。 手上是 GGUF 就别想 vLLM，MTP 也只认 HF checkpoint 里自带的 MTP 层。
• UD 混合量化里可能藏着没有 MMQ 内核的类型。 遇到”小文件反而崩、短 prompt 正常长 prompt 崩”，先 dump 张量类型分布对照 mmq.cuh 的支持列表。IQ1_M 是目前唯一一个 IQ 系列里没有 MMQ 的。
• llama-quantize 的 COPY 模式不做任何量化，--tensor-type 在 COPY 下无效。 改个别张量用 gguf-py 直接重写更可靠。
• -ot ...=CPU 在 CUDA 后端会落到 CUDA_Host，张量仍由 GPU 算，不能用来绕开某个内核。
• 混合架构模型 + 投机解码的显存公式：RS = 槽数 × (draft+1) × 每序列状态大小。 这决定了 draft 数、并发数、上下文三者的取舍。
• 并发吞吐只看墙钟。 用 per-request 的 eval 时间去推合计吞吐，在请求没有真正重叠时会高估一倍。
• Metal 上别开 MTP。 同样的接受率，速度倒退 40%。
• MTP decode 偏算力受限，压功耗墙会掉速；220 W 附近是能效甜点。