一台 128 GB 统一内存的 Jetson AGX Thor T5000，一个 98.57 GiB 的 NVFP4 大模型， 一条没人走通过的路（sm_110 + 内联 NVFP4 的 PLE 表）。

结果是单流 9.53 → 33.23 tok/s（3.5 倍），KV 池从 36k token 扩到 446k token， 同时把多模态开了回来。下面把这台机器上真正发生的事、每条数字背后的算术、 以及踩过的每一个坑写清楚。所有数字都是本机实测，公开数据会单独标注。

2026-09-15 补记：原文定稿后又做完了三件事 —— ① 把 §14.3 那个"未量化的 bf16 输出头" 真的换成了本地造的 NVFP4 头（c=1 +15% / c=2 +25%，质量判分持平）；② 上线了一个带 多 key、按 key 计量、按 key 排队限流的转发网关（8000 网关 → 8001 vLLM），并把 uncensored 权重挂上同一台服务；③ 把 agent 侧每轮的真实 token 开销量了出来 （ZeroClaw 16,009 token vs pi 1,073 token，差一个数量级）。 这三部分在文末新增的第十五 ~ 十七节，原文一字未改。

TL;DR

一句话结论：第 2 步是最大单项收益（2.58x）；第 4 步不提升速度，它是解锁项 —— 不做它，KV 只剩 5.7 GiB，MTP 之后连多模态的 3 GiB 都塞不下。第 7 步是本篇唯一 "先验证机制、再被实测证实"的一条：把草稿层的 MoE 从 Marlin 权重解压换到原生内核后， K=3 从全面落后变成 c=1/c=8 的最好成绩。

一、硬件：Thor T5000 的真实底牌

先把宣传数字拆开。

2070 这个数字是怎么来的

T5000 datasheet（DS11945001）里 2070 TFLOPS = FP4(E2M1) + 2:4 稀疏，@MAXN 1.575 GHz。 注意是 TFLOPS 不是 TOPS，而且带稀疏。同门的 dense 数字：

几何核算对得上：20 SM × 2560 CUDA 核 × 2 (FMA) × 1.575 GHz = 8.064 TFLOPS FP32。 实测运行时钟 1386 MHz（落在 120 W 功耗档）—— 所以日常跑不到 1.575 GHz。

真正决定 decode 速度的是带宽，不是算力

实测（torch 微基准，128 GB 统一 LPDDR5X）：

官方 273 GB/s 是理论峰值，实测纯读只有 193.6 —— 差 29%。 这一条决定了后面所有的 估算：判断一个模型在 Thor 上能跑多快，用「每步活跃权重字节 ÷ 实测带宽」估上界， 不要用 TOPS。

tcgen05 确实能用（这是 sm_110 的关键疑点）

sm_110 是不是真 Blackwell 张量核？业界一直有争议（因为它是 sm_110 而非 sm_120）。 我们直接编内核验证：

1. 编译期：PTX .version ≥ 9.0 下，sm_110a / sm_110f 接受 tcgen05.* 指令； 基础 sm_110、sm_120a、sm_121a 全部拒收。
2. 运行期：写了一个 tcgen05.alloc / relinquish / dealloc 的裸内核，在真机上 launch + sync 全部返回 cudaSuccess。
3. 指令映射（反汇编自己的 cubin 得到）：

4. 生态里的实际情况（扫已编译的二进制）：

最后一行值得记住：vLLM 官方 aarch64 二进制里没有任何 sm_110a 特化的内核， 所以它跑在这台机器上时走的是通用路径。这也解释了后面 MoE 的窘境。

但 tcgen05 对 dense 单流没用 —— 因为 dense decode 是带宽瓶颈（搬权重），不是算力瓶颈。 算力闲置是真，但把它用起来不会让单流变快。

二、模型：98.57 GiB 装不进 128 GB 的算术

Qwen3.8-Flash-Next-NVFP4：98.57 GiB，301,730 个张量，34 个分片。

结构：48 层（12 层 full attention + 36 层 linear attention/GDN）、512 专家 top-10、 hidden 2560、2 个 KV 头 × head_dim 256、MTP 头 1.49 GiB、vision 0.38 GiB。

每步到底要读多少字节

每步活跃权重 ≈ 9.23 GiB（开 MTP）/ 7.74 GiB（不开）。

按纯读 193.6 GB/s 算：9.23 GiB ÷ 193.6 GB/s ≈ 51 ms/步。太慢？ 不 —— MTP 每步能吐多个 token，所以真正的上界要看「每步产出 token 数 ÷ 步长」。 单流带宽上界大约 21 tok/s（MTP0 时）。实测 eager 9.53 → 利用率 45%， 而 Spark 上的公开实现约 70%。

内存墙（这一节是全篇的核心）

总内存            122.86 GiB (MemTotal)
启动时实际可分配   113.85 GiB
权重              96.8 - 97.4 GiB     ← 98.57 GiB 的 checkpoint，加载后
KV                ~71 KB/token
─────────────────────────
权重 + 最小可用 KV(9 GiB) = 106 GiB > 可用 113.85 减去宿主/驱动所需

实测结论：gmu（--gpu-memory-utilization）不存在可行区间。 0.92（KV 12.69 GiB）→ 引擎加载期死；0.89（KV 8.99 GiB）→ 死； 只有 0.86（KV 5.61 GiB）能起来。死亡方式是：

NVRM: GPU0 nvCheckOkFailedNoLog: Check failed: Out of memory
[NV_ERR_NO_MEMORY] (0x00000051) returned from _memdescAllocInternal(pMemDesc)

而且 EngineCore 是无 traceback 静默死亡的，日志里只有 API server 那句 RuntimeError: Engine core initialization failed。要看真因必须从特权容器里读 dmesg：

docker run --rm --privileged --entrypoint sh <IMAGE> -c 'dmesg | tail -20'

所以 PLE 表 offload 不是优化项，是前置条件 —— 26.82 GiB 的 N-gram 表常驻显存， 就是那堵墙本身。

三、四级提升，一级一级拆

第 1→2 级：关 eager、上 PIECEWISE CUDA 图（最大单项，2.58x）

起始配置为了稳妥开了 --enforce-eager（每层 ops 都不进图）。改成：

--compilation-config '{"cudagraph_mode":"PIECEWISE"}'

外加一条容易被忽略的前置动作：

sync; echo 1 > /proc/sys/vm/drop_caches

为什么需要 drop_caches：vLLM 启动时的内存检查把可回收的 page cache 算作已用， 于是报 Free memory on device (115.17/122.86 GiB) on startup is less than desired， 看起来像 gmu 设太高。清一下缓存（Cached 35 GB → 246 MB）就过了。 我一开始误判成"gmu 上限 0.937"，白折腾了几轮。

结果：9.53 → 24.61 tok/s（三次 24.36 / 24.61 / 24.72，很稳），prefill 1.06k → 1.52k tok/s。 —— 顺手追平了公开 Thor 最好水平（Leibniz-HBI 的 23-24，而且他们那 24 是开了 MTP1 才拿到的）。

第 2→3 级：MTP1（+32%）

--speculative-config '{"method":"mtp","num_speculative_tokens":1}'

24.61 → 32.46 tok/s。但 K 不能随便加，见第五节。

第 3→4 级：PLE 表 offload —— 速度不变，容量翻 12 倍

这是本篇技术含量最高、也踩坑最多的一段。

背景：NVIDIA 官方的 /opt/qwen38-thor.patch 里塞了 Qwen4ExpPLENVFP4EmbeddingMethod —— 运行时把 PLE 表当成一个 合并的 320,001,536 行参数，在显存里两次 F.embedding 再反量化。

我们的 checkpoint 与公开配方不同：公开实现（blazux / MiaAI-Lab）假设 PLE 表是 独立的 FP8 分片文件、每行 160 字节；而我们这份是内联在分片 1~19 里的 NVFP4：

128 个头 × 2,500,012 行/头 × 90 字节/行
90 B = 80 B packed U8 权重 + 10 B F8_E4M3 尺度
总计 320,001,536 行 = 26.82 GiB

行空间进一步验证：checkpoint 里的 ngram_heads_offsets = [0, 20,000,003, 40,000,026, ...] （16 个头，每头约 2000 万行），最大合法 id 320,001,446 < 总行数 320,001,536。

做法：自己做 mmap 按行抓取适配器（这也是比公开配方更优的一点：我们不下载新 checkpoint， 且 90 B/行 比 160 B/行 省带宽），三处 hook：

1. Qwen4ExpPLENVFP4EmbeddingMethod.create_weights → 只分配 1 行占位（省 25.6 GiB + 3.2 GiB 尺度）
2. Qwen4ExpNGramEmbedding.load_weights → 跳过 256 个 ngram_embedding.shard_* 张量，记录布局
3. Qwen4ExpPLENVFP4EmbeddingMethod.embedding → 自定义 op vllm::ple_mmap_lookup_ids

PLE offload 的四个坑（每条都是 vLLM 自己 docstring 里写明的约束）

坑 1：光把 op 名写进 splitting_ops 不够。

第一次尝试（offload 生效了，Model loading took 71.73 GiB，确实省下 26 GiB）但崩在 CUDA 图捕获：

RuntimeError: Cannot copy between CPU and CUDA tensors during CUDA graph capture
unless the CPU tensor is pinned.

因为 op 里有一次 H2D 拷贝，被捕获进图了。正确做法是给 op 包上 vLLM 的装饰器：

from vllm.compilation.breakable_cudagraph import eager_break_during_capture

@eager_break_during_capture
def _impl(row_ids, out): ...

它做三件事：结束当前图段 → 在这个流上 eager 执行这个 op → 再开一个新段。 （另外注意：is_breakable_cudagraph_enabled() 依赖环境变量 VLLM_USE_BREAKABLE_CUDAGRAPH， 好在 NVIDIA 这个镜像默认是开的。）

坑 2：输出张量必须地址稳定 + 原地写 —— 否则引擎会静默死亡。

我一开始在每次调用里 torch.empty() 一个新输出张量。捕获用的地址和重放时的地址不一致， 图重放会写进已释放的内存，结果是：

Application startup complete.          ← 起来 3 秒后
Process manager: send sigterm to process EngineCore
EngineDeadError: EngineCore encountered an issue.

没有任何 Python traceback（原生层崩溃），docker inspect 也看不出所以然。 修法是按尺寸缓存持久缓冲：

_BUFFERS: dict[tuple[int, int], torch.Tensor] = {}

def _out_buffer(device, n_rows):
    key = (device.index or -1, n_rows)
    if key not in _BUFFERS:
        _BUFFERS[key] = torch.empty((n_rows, 90), dtype=torch.uint8, device=device)
    return _BUFFERS[key]

坑 3：捕获期的假输入会产生越界行号，而官方 CUDA 路径根本不校验索引。

我们第一次成功捕获后报 KeyError: 3689274810368 —— 我算出的 shard 索引是 3.69e12。 反推发现 3689274810368 × 2,500,012 = 2^63，也就是 id 是个垃圾值（约 INT64_MAX）。 原因是图捕获用的 dummy batch 里，N-gram 行号是垃圾；而官方路径用 F.embedding（CUDA） 不检查越界，只是读到垃圾行、对 dummy 输出毫无影响，所以它从不报错 —— 只有我这边 Python 侧严格查表才发现。

修法干净：ids %= total_rows。合法 id 取模后不变，垃圾 id 被折进合法范围。

坑 4：op 名必须用这个镜像的真实注册名。

公开配方里写的是 vllm::qwen3_8_flash_next_ple_short_conv 之类， 但 NVIDIA 这个 Jetson 镜像里注册的是 qwen4_exp 命名：

vllm::qwen4_exp_ple_short_conv
vllm::qwen4_exp_qsa_with_output

照抄公开配方的话，splitting_ops 里这些名字一个都匹配不上，你会以为"写了没用"。

offload 的效果

权重 + 非 torch    97.0 GiB  →  71.73 GiB
KV 池              5.72 GiB  →  22.91 GiB（144,640 token，4k 上下文并发 35.3 路）
抓行开销           冷缓存 1.5~4.2 µs/行；解码每步 32 行 ≈ 可忽略
单流               32.46 → 33.23 tok/s（没有付出速度代价）

四、投机解码：接受率决定一切

MTP 的 K 值扫描（同配置，只改 K）：

K=1 最优，K 越大越慢。 为什么？从 /metrics 拿到逐位接受率（K=3 时）：

草稿次数 730 / 草稿 token 2190（= 730×3）/ 接受 1028
整体接受率 47.0%
逐位：位置 0 = 494/730 = 67.7%
      位置 1 = 314/730 = 43.0%
      位置 2 = 220/730 = 30.1%

算一下就明白了：

K=1 每步产出 = 1 + 0.677 = 1.68 个 token
K=3 每步产出 = 1 + 1.409 = 2.41 个 token      ← 只多 1.43 倍
但 K=3 的成本 = 3 次草稿前向（读 1.49 GiB 草稿层权重 + 各自 KV）≈ 2 倍
→ 1.43x 收益 < 2x 成本，必亏

根因是逐位接受率衰减太快（68% → 43% → 30%）。 所以调投机解码别拍脑袋改 K，先读 vllm:spec_decode_num_accepted_tokens_per_pos_total，看衰减曲线。

后续在同一高可预测负载下的复核（K=1 vs K=3）

上面的 K 扫描是在混合负载上做的，于是又补了一轮代码改写负载（对投机最有利的场景） 和并发维度的测量，结论更硬：

即使代码改写负载的第一位接受率高达 85%，K=3 依然全面落后：

MTP1:  30.48 tok/s, 1.90 tok/步 → 16.0 步/s  → 62.4 ms/步
MTP3:  25.10 tok/s, 3.19 tok/步 →  7.9 步/s  → 127 ms/步
→ 每步 token 只涨 1.68 倍，步长却翻了一倍 → 净亏 18%

还顺便量到两件有意思的事：

1. 接受率随并发下降（批内序列越多，越难预测）：

   代码改写 K=3 c=1 [0.852, 0.704, 0.630] → c=8 [0.764, 0.599, 0.511] 散文 K=3 c=1 [0.760, 0.442, 0.250] → c=8 [0.675, 0.399, 0.215]

2. 为什么 Spark 上 MTP3 能赚、在我们这里不能：每步的 token 产出双方差不多量级， 但我们每步的草稿成本更高 —— 日志显示草稿层的 MoE 走的是 MARLIN 权重解压 （见第十节 P1-5），而 Spark 的 12.x 家族有可用的快速内核。也就是说这不是负载问题， 是草稿前向的相对开销问题。

换成原生内核之后，K 的最优点跟着并发布移动（最终版扫描）

把草稿层的 MoE 从 Marlin 换到 humming 之后重扫一遍，并且这次用步长而不是 tok/s 来比较：

拿三点做一个线性模型，误差只有 3%：

步长 = 42.3 ms（基础：读 7.74 GiB 主干） + K × 13.7 ms（每次草稿前向）

基础段 42.3 ms vs 物理下界 7.74 GiB ÷ 194 GB/s = 39.9 ms  → 利用率 94%
草稿段 13.7 ms vs 物理下界 1.49 GiB ÷ 194 GB/s =  7.9 ms  → 每次白花 5.8 ms
验算：MTP2 预测 42.3 + 2×13.7 = 69.7 ms，实测 67.5 ms

这就把"K 的最优点为什么会移动"讲清楚了：单流时草稿前向永远是 M=1 的小 GEMM（读 1.49 GiB 却只算一个 token，带宽吃不满，13.7 ms 里 5.8 ms 是浪费），所以 K 越大越亏； 多流时批内序列多、草稿前向的 M 变大、效率上来，K=2/3 才开始赚。

选型规则：

多路 agent / 并发服务  → MTP2 或 MTP3 都行（c=8 两者噪声内相当：118-130 对 124-129），
                          但 MTP2 的 c=1 更稳（33.9-34.3）且 KV 池最大（493k）
单流交互（一个人写代码）→ MTP3 可以再搏 8%（36.89），代价是抖动到 [28.7-37.1]
能确立的只有：MTP1 在 c=8 明显低 8-10%

顺带一个测量陷阱：别用 SSE chunk 数当 token 数。vLLM 会把一步内接受的所有 token 打包进同一个 delta，chunk 数其实是步数，这样测会得出"接受率 97% 但速度只有 1 tok/s"的荒谬结论。 永远用 usage.completion_tokens。

五、内存墙的正确解法：显式留出宿主余量

offload 成功后我犯了一个新错误：腾出来的 26 GiB 被 vLLM 的 KV 自动分配一口吃光了。

gmu 0.88 的预算            108.12 GiB
权重 + 非 torch             74.22
激活峰值                      1.17
CUDA 图                       0.59
KV（自动分配的！）          +32.73
──────────────────────────
合计                       108.71 GiB（顶到预算）
启动时可用                  113.85 GiB  → 宿主只剩 ~5 GiB
→ 又一次 NV_ERR_NO_MEMORY，引擎静默死亡

这和 MiaAI-Lab 仓库里警告的失败模式完全一样（"KV 给太多、宿主只剩 6.9-8.8 GiB，驱动开始拒绝分配"）。

正确做法：显式留宿主余量。 gmu 降到 0.80：

预算            98.29 GiB
权重+非torch    74.18 / 加多模态后 77.03
激活峰值         1.20 / 2.48
CUDA 图          0.61 / 0.38
KV              22.91 / 18.78
→ 宿主保留 ~15-17 GiB

注意一点：mmap 的 PLE 表需要 page cache 才便宜，所以饿死宿主不只是稳定性问题， 也是性能问题。

一个反直觉的副产品

把 --max-model-len 从 4096 提到 32768，KV 的 GiB 没涨（22.91 → 18.78，因为改用大 block）， 但 token 容量翻了 3 倍（144,640 → 446,223）。原因是 vLLM 为对齐 mamba page 会调整 block_size（4k 时 1568，32k 时 1600），padding 开销变小。

同时"恢复多模态"的真实代价不是 4 GiB 而是 ~3 GiB（编码器缓存），却换来了更大的 token 容量和可用的视觉输入。这笔交易很划算。

六、其他坑（速查）

七、评测的坑（如果有做模型对比，这些会各花掉你一小时）

用 lm-eval 对两个同规模模型（27B dense / 35B-A3B MoE）做能力对比时踩的坑：

1. 思考模型返回空 content。 带 reasoning parser 时答案落在 reasoning 字段， 而 lm-eval 只读 content → 每条都是 null（API returned null content）、全 0 分假象。 评测时必须 --default-chat-template-kwargs '{"enable_thinking": false}'， 并且先看一条响应的 content 非空再信分数。
2. lm-eval 接本地端点用的是 auth_token= 而不是 api_key=。 实际 header 从 OPENAI_API_KEY 环境变量拼 —— 401 且 header 显示 'Bearer ' 就是环境变量没设。 还有个坑：脚本里 export OPENAI_API_KEY=$KEY 写在 KEY=$(cat ...) 之前，会静默拿到空 key。
3. HumanEval 对聊天模型基本不可用（as shipped）。 任务自带 until=['\ndef','\nclass',...]，模型一写出 def 就被切断，得到 40-70 字符残片和 ~0 分。 要么用 humaneval_instruct 并且覆盖停止串 （--gen_kwargs "until=[] max_gen_toks=1400"），要么自己做执行判分。
4. --limit N 是按子任务算的。 mmlu_pro 有 14 个子科目，--limit 200 = 2800 次请求（几小时）。 --limit 15 → 210 题，够看趋势。
5. GPQA 是 gated 数据集，没 HF 授权会直接 DatasetNotFoundError。
6. 重跑前一定清掉/改名旧的 --output_path 日志 —— 上一次失败留下的 samples_*.jsonl 和一次成功的结果长得一模一样。
7. 还有一条不算 lm-eval 的锅：GSM8K 的 samples 每个问题存两遍 （strict-match / flexible-extract 两个 filter 各写一次），统计错题时会虚高一倍 （我看到 20/200 实际是 10/100）。

八、别人的工作：链接、数字，以及我们用了他们什么

这条路不是没人走过 —— 恰恰相反，前后有十几个项目。先把原文和数字列清楚， 再说我们站在谁的肩膀上。

和我们的实测完全对得上的部分

这种"独立复现出同一组数字"的吻合，比任何单方面的宣称都可信：

• Leibniz 的 eager MTP0 = 9.5 tok/s，我们独立实测 9.53 tok/s（同配置口径）。
• 他们的 LIMITATIONS.md 第一条就是"原生 SM110 NVFP4 MoE 是最大的性能缺口"， 与我们扫二进制得出的结论一致（CUTLASS grouped-MoE 在 sm_110 上 status=7）。
• 第二条"融合版 MTP GDN 内核在 SM110 上无法执行，只能 VLLM_GDN_DECODE_KERNEL=triton" —— 我们遇到的是同一个内核缺口。
• 第三条"QSA cooperative Top-K 在 SM110 上会以非法 cluster 几何启动，配方把它改成 persistent Top-K" —— 这正是我们镜像里 persistent_topk 的由来（也是非确定性风险的来源）。
• 第五条"PLE mmap 仍有 host/device 同步开销" —— 我们在每步都能量到那次同步。

一个必须先纠正的常见误解：Spark 和 Thor 不是同一档芯片

网上（包括我最初的理解）常把 DGX Spark 的 GB10 和 Jetson AGX Thor 当成"同一种 sm_110"。 这是错的，而这个错误会让所有结论都跑偏：

来源：arnon.dk 的 arch 对照表、 conda-forge 的 GPU 归类讨论 （"Jetson Thor is sbsa+sm110/11.0 … DGX Spark (GB10) is sbsa+sm121/12.1"）、 Level1Techs 论坛的对比。

为什么这件事很关键：整个 CUDA 生态（flashinfer 的 b12x 系列、大量 CUTLASS GEMM） 是为 12.x 消费级家族编译的，因为那才是有量的设备；而 Thor 的 sm_110 是一条 只有它自己的架构线（CUDA 13 才把 sm_101 改名 sm_110）。 结果就是本文反复出现的那句话：Thor 的硬件更强（有 tcgen05、矩阵吞吐 2 倍）， 但生态里没有它的内核。Spark 的 48.7 tok/s 是在没有 tcgen05 的芯片上、 靠一套能跑的内核 + 更深的投机解码拿到的。

九、我们走到了哪一步，以及为什么没到人家的速度

走到了哪一步

一句话：在 Thor 这个平台上，逐档我们都在最前面；没达到的只有 Spark 那两个数字。

为什么没到 —— 三条归因，按证据强度排序

① 我们把自己限在了 2 路并发（这是我引入的瓶颈，不是硬件限制）。

启动脚本里写的是 --max-num-seqs 2 —— 这是早期权重 97 GiB、KV 只剩 5.7 GiB 时为省内存设的上限。 后来 PLE offload 腾出 26 GiB、KV 涨到 446k token 之后，我忘了把这个上限提上去。 所以本文所有"4 并发"测量实际上是被压成 2 路跑的：

4 请求 / 43.09 tok/s 聚合  →  实际是 2 路并行，每路 ~21.5 tok/s

对照 MiaAI 的 8 路 162.9（每路 20.4）—— 我们的"每路"效率其实和它同一水平， 差的只是并行路数被自己锁死了。

② 单流已经贴在内存带宽屋顶上，所以内核优化在单流上根本不会有效果。

这是本文最该被记住的一段算术：

MTP1 每步产出 = 1 + 接受率(位置0) = 1 + 0.677 = 1.68 token
33.23 tok/s ÷ 1.68 tok/步 = 19.78 步/秒  →  50.6 ms/步
每步活跃权重 9.23 GiB ÷ 50.6 ms = 196 GB/s
实测纯读带宽                     = 193.6 GB/s
→ 利用率 ≈ 100%

也就是说：每步要读的字节已经被读满了。 单流唯一的杠杆不是"算得更快"， 而是每步多产出 token（更深的投机解码）或每步少读字节（更激进的量化）。 这解释了为什么把 K 从 1 加到 3 没用（接受率衰减使每步产出只从 1.68 涨到 2.41，成本却翻倍）， 也解释了为什么"补上 FP4 MoE 内核"不会让单流变快 —— 单流根本不在算力限制区。

对照：MiaAI 单流 48.7 tok/s
若按接受长度 2.8-3.0（MTP3）算 → 每步产出 ~3.8-4.0 token → 约 12.2 步/秒 → 82 ms/步
他们每步比我们慢（82 vs 50.6 ms），但每步产出是我们的 2.3 倍
→ 快的是"每步产出 token 数"，不是"每步算得快"

但这里有个必须说明的口径问题：他们那个 2.8-3.0 的接受长度很可能来自高可预测负载 （代码改写/重复文本），而我们报的 47% 整体接受率来自混合散文基准。 同一份模型上我们早前在 EDIT-heavy 负载测到过 97.7% 的接受率 —— 投机解码的收益完全由负载决定，两边的数字不构成同口径对比。 想公平比较， 必须用同一组提示词重测。

③ prefill / 高并发是真正的差距所在，而那里确实是缺内核。

我们      prefill ~1,516 tok/s（2.5k 输入）；长上下文衰减到 ~800 tok/s（20k 输入）
MiaAI     prefill ~2,200 tok/s（8k 输入）

prefill 是算力限制的（不像 decode 是带宽限制），而这里正好是内核短板所在。 我们的启动日志把内核分配说得非常清楚（原文摘录）：

[nvfp4.py:304]  Using 'VLLM_CUTLASS' NvFp4 MoE backend ...     ← 目标 MoE：原生 FP4 ✓
[__init__.py]   Using MarlinNvFp4LinearKernel for NVFP4 GEMM   ← 线性层：Marlin 回退
[__init__.py]   Using MarlinMxfp8LinearKernel for MXFP8 GEMM   ← 线性层：Marlin 回退
[marlin.py:34]  Your GPU does not have native support for FP4 computation but FP4
                quantization is being used. Weight-only FP4 compression will be used
                leveraging the Marlin kernel. This may degrade performance for
                compute-heavy workloads.
[nvfp4.py:304]  Using 'MARLIN' NvFp4 MoE backend ...           ← 草稿层 MoE：Marlin 回退

也就是说：目标 MoE 的专家层其实拿到了原生 FP4，真正被降级的是 线性层（NVFP4 / MXFP8 GEMM）与草稿层（MTP）的 MoE —— 它们跑的是 Marlin 的 "权重解压后再算"（日志原话 weight-only FP4 compression，并明确警告会拖慢 compute-heavy 负载）。prefill 和高并发正是 compute-heavy 区间，所以这条差距 是真实且可归因的。

那"别人说的 41 tok/s"到底怎么来的

同类硬件（GB10/Spark，同为 273 GB/s 规格、无 tcgen05）上的独立报告给的就是这个区间 —— NVIDIA 开发者论坛的 Flash-Next 长帖里：

"34 to 42 tok/s decode with the MTP head"，262K 原生上下文 + vision + 前缀缓存（单台 GB10）
"35 t/s with no degradation over context length"（FP8 + 750k KV，但 agentic 负载只有 15-17）
双 Spark FP8 MTP=3：tg1024 @ d4096 = 35.02 t/s

所以在 Thor 上我们实测的步长拆解，可以把这个 41 讲得很清楚。 关键实测：步长与负载无关。

两种负载的步长几乎一样（85.1 vs 86.2 ms），所以单流速度 = token/步 ÷ 85 ms：

41 tok/s 需要 41 × 0.0851 = **3.49 token/步**
我们最好的负载拿到 3.08 → 36.15 tok/s，是 41 的 88%
MTP3 时每步只吃掉 121 GB/s，而实测纯读是 193.6 GB/s → **步长还有约 40% 余量**

差距分两块，而且都有明确的下一步：

1. 命中率（负载相关）：3.08 → 3.49 之间的部分靠更强的 drafter，或更适合投机的负载。
2. 步长本身：把每步的字节拆开算，就能看清 85 ms 是怎么来的。

步长拆解（这里纠正了一个流传很广的说法）

先纠正一处：网上（包括 Leibniz 的 LIMITATIONS）说"融合版 GDN MTP 内核在 SM110 上不可用、 只能退回 Triton"。在我们这套 build 上不成立 —— 启动日志明确写着：

INFO [qwen_gdn_linear_attn.py:528] GDN decode kernel: cuda

这段门禁（_fused_gdn_decode_unsupported_reason）只要求 BF16 + K=V=128 + compute capability ≥ 8.0， sm_110 轻松通过。所以"单流慢是因为 GDN 回退"这个解释是错的。

真正的账是草稿层被读了 K 遍：

MTP1 每步读 = 主干 7.74 GiB + 草稿层 1.49 GiB × 1 =  9.23 GiB
              理论 9.91 GB ÷ 193.6 GB/s = 51.2 ms，实测 50.6 ms → **100% 带宽饱和**

MTP3 每步读 = 7.74 GiB + 1.49 GiB × 3 = 12.21 GiB
              理论 13.11 GB ÷ 193.6 GB/s = 67.7 ms，实测 85.1 ms → **多出 17.4 ms**

多出来的 17 ms 是草稿前向的小批量低效：草稿是自回归的（第 n 个 token 依赖第 n-1 个）， 单流时永远是 M=1 的 GEMM —— 要读 1.49 GiB 权重却只算一个 token，带宽吃不满。

这解释了 MTP3 为什么在多流下才显出优势：

c=1（草稿 M=1，低效）：MTP3 32.18 vs MTP1 33.23 → 基本打平
c=8（草稿 M=8，高效）：MTP3 129.08 vs MTP1 117.01 → +10%

所以选型的规则很简单：单流自用 → MTP1/MTP3 差别不大；要多路并发 → 必须 MTP3 + 原生内核。

结论：单流不是"机器不行"，而是每步的字节数（草稿层读 K 遍）与那份文本的命中率决定的。 想再往上走，两个方向：让草稿前向更省字节（更激进的量化／共享），或者提高命中率。 另外别把 RTX PRO 6000 的 180-226 tok/s 拿来比 —— 那是 1.8 TB/s 的显存，完全不同档。

想够到 41 的最后两条路（都试了，都堵死）

① 换更强的 drafter —— 没有可用的。 吞吐瓶颈既然在"每步 token 数"，最直接的办法就是换接受率更高的草稿器。 查下来：RadixArk/Qwen3.8-27B-DSpark（1.86B draft，配 SGLang 的 DSPARK）确实存在，但那是 27B 的； Flash-Next 没有公开的 DSpark/EAGLE speculator（社区也在问同样的问题，答案都是"先等官方 EAGLE 头"）。 而且 vLLM 只对 config.json 里带 dspark_* 块的 checkpoint 支持 dspark —— 现有 Flash-Next 权重没有这个块。 → 除非自己训一个 speculator，这条路走不通。

② ngram（prompt-lookup）投机 —— 被 QSA 的 ring capacity 断言挡死。 理论上它最合适：不需要任何草稿权重、不需要草稿前向，每步字节从 12.21 GiB 回到 7.74 GiB （步长 ~50 ms 而不是 85 ms），而"改代码 / 重写文件"这类负载本来就是大量复制上下文， 正是 prompt-lookup 的甜点区。实测直接崩在启动：

AssertionError: QSA ring capacity 12 must divide the attention block size 1616

根因（读 models/qwen4_exp/common/qsa_cache.py 得到）：

span     = compress_ratio + num_speculative_tokens        # 6 + 5 = 11
capacity = compress_ratio * ceil(span / compress_ratio)   # 6 * 2 = 12
assert block_size % capacity == 0                          # 1616 % 12 != 0 → 挂

ngram 会强制切到 V1 model runner（日志原话 "Model Runner V2 does not yet support ngram/ngram_gpu speculative decoding"），而 V1 算出的 block_size=1616 不能被 12 整除；MTP 路径下是 1584，能整除，所以没事。 → 这是 V1 runner 与 QSA 的 block-size 协商不兼容，不是配置写错。

③ 但还有第三条路，而且是做完 kernel 级归因才发现的（见第十四节）： 输出投影 lm_head 是 bf16、1.27 GB，且每步被读 4 次（1 次目标验证 + 3 次草稿， 草稿复用同一份权重）。仅这一项就是 20.3 ms/步 = 单步的 24%。 实测把草稿那 3 次换成 FP8 副本后，步长稳定降 7.5-10 ms（83 → 74.8 ms）—— 按此推算，若量化粒度做得够细（按 K 分块，MXFP8 风格）使接受率不受影响， 同负载单流就是 约 40 tok/s。这是目前唯一一条既不改内核、也不训模型的路。

小结

单流      → 带宽与步长层面已到顶（MTP1 达实测纯读带宽的 100%）；提速只能靠"每步多出 token"。
            MTP3 每步多出的 ~17 ms 是草稿前向的小批量低效（M=1 的 GEMM 吃不满带宽）
并发      → 曾被我们自己的 max-num-seqs=2 锁死（人为差距，已解除：43 → 129 tok/s）
预填充/高并发 → 真实的算力缺口，只能靠补内核（humming 原生内核是第一个进展，但只拿到持平）
投机深度   → 一旦草稿层离开 Marlin，K=3 从"全面落后"变成最好（+28% @c=1、+19% @c=8）
跨平台比较  → Thor(sm_110, 有 tcgen05) vs Spark(sm_121, 无 tcgen05)，同带宽级但不是同一档芯片

十、接下来做什么

按"性价比"排序，不是按"听起来高级"排序。

P0 — 配置层（今天就能做，预期收益最大）

1. 解除并发上限 —— 已实测完成：--max-num-seqs 2 → 8、 --max-num-batched-tokens 2048 → 8192。结果（同机同模型）：

单流零代价（符合"已贴带宽屋顶"的结论），聚合吞吐 2.3-2.7 倍。 KV 池有 415k token 容量（32k 上下文 12.7 路并发），完全支撑得起。 这一项同时校正了我此前所有"并发"数字 —— 它们其实都是 2 路的结果。 2. 用同口径负载重测接受率：准备两组提示词（高可预测的代码改写 / 混合散文）， 分别测 K=1/2/3。现在的 47% 只代表混合负载，不能用来判定 K 的最优值 —— 在 EDIT-heavy 负载上 K=16 曾在别的模型上从 12.3 干到 66 tok/s。 3. 显式控制 KV 目标（--kv-cache-memory 或用 KV_TARGET_GIB 的思路）， 把"自动分配吃满"变成"留出指定余量"；省下的显存换更长上下文或更深的图。

P1 — 内核层（按实测日志重排，工作量最大）

4. 把线性层从 Marlin 换到原生 FP4 CUTLASS。证据：MarlinNvFp4LinearKernel for NVFP4 GEMM 与 MarlinMxfp8LinearKernel for MXFP8 GEMM，而 Marlin 自己警告 "weight-only FP4 会拖慢 compute-heavy 负载"；同时 torch 自带 59 个 sm_110a 的 CUTLASS block-scaled FP4 GEMM。 影响每一层每一步，收益主要在 prefill 与高并发 —— 单流基本无感，别指望它提升 33。
5. 让草稿层的 MoE 脱离 Marlin。日志里 NvFp4 MoE backend 被选了两次：目标 MoE = VLLM_CUTLASS，草稿层 = MARLIN。草稿层的 Marlin 开销正是让 K>1 亏本的那笔成本 （见第四节），把它修好可能让更深的投机重新变得划算。

先记下一个已排除的捷径：直接强制 --moe-backend cutlass 行不通，实测启动即失败：

   ValueError: NvFp4 MoE backend 'VLLM_CUTLASS' does not support the deployment
   configuration since kernel does not support quantization scheme
   QuantKey(u8, scale(f8e4m3fn, static, GroupShape(row=1, col=16)))

注意这个报错的含义：降级不是"忘了打开一个旗标"，而是那个量化方案在这套 build 里 根本没有对应内核 —— auto 的选择已经是候选列表里能拿到的最优组合 （候选池：FLASHINFER_TRTLLM / FLASHINFER_CUTEDSL / FLASHINFER_CUTEDSL_BATCHED / FLASHINFER_CUTLASS / VLLM_CUTLASS / MARLIN / HUMMING / EMULATION）。 要改善只有两条路：补内核，或者换量化方案（例如把 MXFP8 那部分requantize成 有内核的方案）。这解释了为什么这个坑在社区里挂了很久 —— 它不是配置问题。 6. 试 Marlin 自己提示的开关：VLLM_MARLIN_USE_ATOMIC_ADD=1 —— 已实测，不采纳。 日志原话是 "can achieve better performance for small size_n"，听起来正对草稿层和小批量 解码的形状，但在这台机器上（代码改写负载，同一提示词）：

两个档位都慢约 12%，而接受率还略高一点（0.85/0.79 vs 0.83/0.75），说明是内核本身 变慢了，不是统计抖动。散文负载上结果混合（c=4 慢、c=8 快），不足以翻盘。 教训：框架自己打印的性能建议也要实测，它是在另一套形状/另一代架构上得出的。 7. 解开 FlashInfer cuTe-DSL MoE 在 sm_110 上的门禁：flashinfer_cutedsl_moe.py:81 只放行 is_device_capability_family(100)，而候选 backend 列表里 FLASHINFER_CUTEDSL 是存在的 —— 也就是说可能只差一个能力判断就能多拿一档。 8. （已修正）不是"编译 GDN 内核"：实测日志显示 GDN decode kernel: cuda，融合内核在 sm_110 上本来就启用（Leibniz 的 LIMITATIONS 第 2 条只适用于他们那份 pinned build）。 真正的杠杆是减少草稿层的每步字节：MTP3 时草稿层权重 1.49 GiB 被读 3 遍， 且单流下是 M=1 的小 GEMM（实测每步多出 ~17 ms）。方向：草稿层更激进量化、 或让草稿前向在多流下天然批量更大（这已经是 c=8 时 MTP3 领先的原因）。 9. 修 QSA cooperative Top-K 在 SM110 上的 cluster 几何（第 3 条）。 顺带能去掉 persistent_topk 这条非确定性风险源。 10. 排查 Thor 上 PLE GPU gather（ATS）的非法访问。修好后抓行可留在 GPU 侧， 省掉每步 host/device 往返（第 5 条指向同一处）。

P2 — 容量层（如果目标是超长上下文）

11. fp8 KV：KV 显存减半 → 同显存下 token 容量约 2 倍（MiaAI 用它做到 512k 上下文 + 8 路）。
12. 更激进的 PLE 常驻/量化：把表的一部分按热度常驻显存，兼顾延迟与容量。

一句话

单流已经到底了（100% 带宽利用率），接下来的收益全在"每步多产出 token"和 "并发/预填充的算力侧"—— 前者靠投机解码的负载匹配，后者靠补上 sm_110 的内核。 而且第一件该做的事，是把我自己设的 max-num-seqs 2 拿掉。

十一、动手记录：把 humming 的原生 NVFP4/MXFP8 内核救活

第十节 P1 的第 4 条（把线性层从 Marlin 换到原生内核）在写这份笔记的当天就往前推了一大步， 过程比结论有意思，值得单独记一节。

发现：镜像里其实装着一条没人用的原生路径

镜像里装了一个叫 humming 的库（v0.1.12），vLLM 里有对应的线性层内核 （HummingNvFp4LinearKernel / HummingMxfp8LinearKernel），而且它的 is_supported() 只要求 has_device_capability(75) —— 没有排除 sm_110。 更关键的是它带 nvrtc_compile：这是运行时按设备 JIT 编译的库， 不是只支持特定几代的预编译包。

所以就有了一个很朴素的尝试：

--linear-backend humming        # 合法取值之一，文档说"没实现的层自动回退"

结果内核选择确实变了：

Using HummingMxfp8LinearKernel for MXFP8 GEMM     ← 原来是 MarlinMxfp8LinearKernel
Using HummingNvFp4LinearKernel for NVFP4 GEMM     ← 原来是 MarlinNvFp4LinearKernel

但选中之后启动即崩，然后是一连串"只差一行"的缺口 —— 每一处都指向同一件事： 库作者认真处理过 DGX Spark（GB10/sm_121）的统一内存情形，只是漏了 Thor（sm_110）。

① humming/tune/__init__.py
   heuristics_map = {75,80,86,87,89,90,100,103,120,121}    ← 没有 110
   → KeyError: 110
   （注意 103 已经直接复用 Sm100Heuristics，说明 sm_10x 这条数据中心线共享同一档启发式）

② humming/utils/device.py   calculate_gpu_bandwidth
   if (major, minor) == (12, 1):
       # NVIDIA GB10 (DGX Spark, sm_121): unified LPDDR5X memory with no
       # discrete memory-clock domain, so NVML_CLOCK_MEM raises
       # NVML_ERROR_NOT_SUPPORTED ... Use the platform's known memory bandwidth instead.
       return 273.0
   ← 这段特例写得很清楚。只是 (11,0) 也是统一内存、也在 NVML 上查不到显存时钟

③ humming/utils/device.py   estimate_tensorcore_max_tops
   nvmlDeviceGetMaxClockInfo(handle, NVML_CLOCK_SM) 在 Jetson 上同样 NOT_SUPPORTED，
   而且 ops_map 里也没有 110

修法：运行时补丁，不改源码

用 .pth 在解释器启动时打补丁（API server 与 EngineCore 进程都会生效）：

heuristics_map[110] = Sm100Heuristics          # 与 sm_103 同款处理
带宽查询失败 → 回落 273 GB/s                    # 与库内 GB10 特例同值（Thor 规格同为 273）
张量核峰值失败 → 20 SM × 8192 ops/clk × 1.575 GHz = 232 TFLOPS FP16 dense
                 （与 datasheet 的 FP8 dense 517 TFLOPS 自洽：FP16 是 FP8 的一半）

补完这三处，服务起来了：模型的完整前向（含多模态编码器）在 humming 内核上跑通， KV 池 452,005 token，--linear-backend humming 稳态可用。

正确性先于性能

换内核必须验输出（新路径可能算错，而且错得很安静）：

日本的首都是哪里？           → 东京 ✓
写 fib(n) 函数               → 语法正确的 Python ✓
三句话解释带宽瓶颈            → 恰好三句、连贯、内容正确 ✓
无退化、无重复、无乱码 ✓

性能：说实话，没测出确定的收益（用重复取中位数复测后仍然如此）

第一轮用单次测量，结果自相矛盾；改成每个档位重复 3 次取中位数之后（Marlin 基线用的是 同一套脚本、同一组提示词）：

结论：基本持平，没有可以写在标题上的收益。 但它的意义不在速度：

1. 这是一条真正在 sm_110 上跑起来的原生 FP4/MXFP8 路径，不再是"权重解压后再算"；
2. 它运行在 JIT 编译路径上，意味着调优/自动调参是有空间的（目前的配置来自启发式估算， 而不是针对这块芯片实测搜索出来的）；
3. 它把草稿层的 MoE 也一起从 Marlin 换了出来（--moe-backend humming）， 而草稿层的开销正是决定"更深投机是否划算"的那笔账（见第四节）。

关于噪声来源，还是那句话：换了内核 → 浮点累加顺序变了 → 同一提示词、温度 0 下生成的文本 本身就不一样（实测接受率 0.90 → 0.76），于是 tok/s 里混进了"这次写的是什么"的影响。 所以换后端类实验必须看与生成内容无关的步长，并重复取中位数。这轮之后我把基准脚本固定成 了三重复制 + 中位数 + 区间。

还没试完的部分 → 已经试完：把草稿层拉出 Marlin 让 K=3 回本

日志显示 NvFp4 MoE backend 有两次选择（目标 MoE = VLLM_CUTLASS，草稿层 = MARLIN）， 而 vLLM 里有 humming 的 MoE 实现（fused_humming_moe、NvFp4MoeBackend.HUMMING）。 于是把 --moe-backend humming 也打开，让草稿层的 MoE 离开 Marlin 权重解压， 再测 K=1 / K=3（同一个代码改写负载，重复取中位数）：

关键对比（同一 K，只换草稿层后端）：

K=3, c=1:  25.10 → 32.18   （+28%）
K=3, c=8: 108.30 → 129.08  （+19%）

也就是说：第四节里"K>1 不划算"的结论，根因并不在草稿质量，而在草稿层的 Marlin 开销。 草稿层一旦跑上原生内核，K=3 从"全面落后"变成 c=1 和 c=8 的最好成绩。 这条把两个看似独立的发现（投机解码的接受率曲线、Marlin 回退的内核缺口）接上了， 也是这份笔记里唯一一个"先有机制预测、后被实测证实"的结果。

（诚实标注：c=1 的区间是 [24.9–33.0]、c=8 是 [118.6–141.1]，抖动仍然不小， c=4 还略低于 MTP1+Marlin。所以"最好"是中位数意义上的最好，需要更多重复才能下最终结论。）

顺带一笔必须记住的运维账：启动慢的真凶既不是 humming，也不是编译缓存，而是视觉（多模态）的 Triton warmup。 时间轴来自一次真实启动：

05:34:58  KV 池建好
… 14.7 分钟完全静默（1-2 个 CPU 核钉在 100%、GPU 功耗仅 6W、全盘零新文件、无 IO）…
05:49:39  Warmed position embedding and vision rotary kernels on grids=[...]
05:49:39  Warmed M-RoPE Triton kernels / Warmed Mamba batch_memcpy_kernel
05:49:50  init engine took 938.96 s

KV 池之后到 READY 之间，日志里唯一的一行就是视觉 warmup 的完成行 —— 那 14.7 分钟全花在里面， 而它只在结束时打印，所以整段过程看起来像卡死（其实 GPU 在空载等 CPU 编译内核）。

对照（引擎初始化时间）：

Marlin + 关多模态        394-412s
humming + 开多模态        881-939s

下面几个猜测都被实测证伪了，别再重复：

挂 humming 编译缓存        无缓存 26 min → 挂上 19.2 min，但后续两次 21.9 / 20.0 min，落在抖动内
挂 Triton 编译缓存         无收益（98 MB / 1879 个文件，白挂）
关 enable_jit_warmup      日志确认生效，但启动仍 20 分钟
关 enable_cutedsl_warmup  同上
关 enable_flashinfer_autotune
                          日志明确 "Skipping FlashInfer autotune because it is disabled"，启动仍 20 分钟
py-spy 抓 CPU 栈           容器缺 CAP_SYS_PTRACE → Permission denied（要 --cap-add SYS_PTRACE）

实用取舍：

纯文本（--limit-mm-per-prompt '{"image":0,"video":0}'）→ 启动回到个位数分钟，KV 池还更大
要视觉                                              → 接受约 15 分钟的一次性 warmup

十二、还没解决的根源问题（诚实清单）

这些不是调旗标能解决的，是真正的空白：

1. 线性层与草稿层没有原生 FP4 内核，只能跑 Marlin 权重解压 —— 这是最大的内核缺口。 注意别把这句说成"MoE 没有原生 FP4 内核"：实测启动日志显示目标 MoE 的专家层走的是 VLLM_CUTLASS 原生 NVFP4，被降级的是 NVFP4/MXFP8 的线性层 （MarlinNvFp4LinearKernel / MarlinMxfp8LinearKernel）和草稿层（MTP）的 MoE （第二次 backend 选择结果 = MARLIN）。而 torch 自带的 libtorch_cuda.so 里就有 59 个 sm_110a 的 CUTLASS block-scaled FP4 GEMM —— 说明"这个内核在这块芯片上编得出来"， 缺的是把它接进 vLLM 的线性层与草稿层。
2. Thor 的 ATS GPU gather 非法访问，PLE 抓行只能走 CPU（约 93% 的加载路径会踩到）。
3. 草稿层的每步字节数：MTP3 时草稿层权重（1.49 GiB）被读 3 遍，占总字节的 1/3； 而草稿前向在单流下是 M=1 的小 GEMM，实测每步因此多出 ~17 ms（理论 67.7 ms vs 实测 85.1 ms）。 注意：这不是"缺内核" —— 启动日志显示 GDN decode kernel: cuda，融合内核在 sm_110 上是 启用的（网上流传的"SM110 只能用 Triton 回退"在这套 build 上不成立）。 真正的方向是减少草稿层的读取字节（更激进的量化）或提高其批量效率。
4. persistent_topk 非确定性：GB10 上有人实测 temp=0 时 13/50 次发散 （atomicAdd 槽位，对应 vllm#51782）。我们这条 Jetson 镜像路径下 3 提示 × 6 次 = 18/18 字节完全一致，没复现 —— 但样本量小，长期当 agent 用还得再验。

十三、复现清单

# 0. 前置：清 page cache（否则启动内存检查假失败）
docker run --rm --privileged --entrypoint sh <IMAGE> -c 'sync; echo 1 > /proc/sys/vm/drop_caches'

# 1. 起服务
docker run -d --name vllm-flashnext --runtime nvidia --network host \
  --ulimit memlock=-1 --ulimit stack=67108864 --shm-size=32g \
  -e VLLM_API_KEY="$KEY" \
  -e VLLM_PLE_MMAP=1 -e VLLM_PLE_MMAP_DIR=/model \
  -v /path/flash-next:/model:ro \
  -v /tmp/fixed_config.json:/model/config.json:ro \        # model_type=qwen4_exp
  -v ~/.cache/vllm:/root/.cache/vllm \
  -v ~/.cache/flashinfer:/root/.cache/flashinfer \
  <IMAGE>:plemmap /model \
  --served-model-name qwen3.8-flash-next --host 0.0.0.0 --port 8000 \
  --gpu-memory-utilization 0.80 \                          # 显式留 ~15 GiB 给宿主
  --max-model-len 32768 --max-num-seqs 2 \
  --max-num-batched-tokens 2048 \
  --safetensors-load-strategy lazy --trust-remote-code \
  --speculative-config '{"method":"mtp","num_speculative_tokens":1}' \
  --compilation-config '{"cudagraph_mode":"PIECEWISE","splitting_ops":[
      "vllm::unified_attention_with_output",
      "vllm::mamba_mixer2","vllm::short_conv","vllm::linear_attention",
      "vllm::qwen_gdn_attention_core","vllm::sparse_attn_indexer",
      "vllm::qwen4_exp_ple_short_conv","vllm::qwen4_exp_qsa_with_output",
      "vllm::ple_mmap_lookup_ids"]}'

关键检查点（照着看日志，每一步都对得上再往下走）：

权重 + 非 torch   应约 71.7-77.0 GiB（不做 offload 是 ~97-101 GiB）
CUDA 图           4-6 个 PIECEWISE 图捕获成功
KV 池             18.8-22.9 GiB / 144k-446k token
宿主可用内存      ≥ 15 GiB（低于这个数就会撞 NV_ERR_NO_MEMORY）
单流              33-37 tok/s（MTP1/2/3，视负载）；8 并发 124-154 tok/s
启动耗时          首次 ~6 分钟；之后引擎初始化 30-48 s（缓存在 ~/.cache/humming）

三处已知待调/可调： ① --max-num-seqs（本篇早期用的是省内存时的 2，实际支撑 8）与 --max-num-batched-tokens （2048 → 8192）—— 见第十节 P0-1。 ② 关多模态 --limit-mm-per-prompt '{"image":0,"video":0}' —— 能把启动从 20 分钟压到 4.5 分钟 （多模态视觉 warmup 占约 750 s），代价是失去视觉能力。 ③ 可选实验开关：-e VLLM_MTP_LOGITS_FP8=1 —— 让草稿的输出投影走 FP8 副本 （见第十四节；目前步长省 7.5-10 ms，但量化粒度待改进才能真正转化为吞吐）。

十四、拆开一个 decode 步：kernel 级归因，以及一个未量化的 lm_head

前面十节全在算"每步搬多少字节"，但草稿那 13.7 ms 到底花在哪个 kernel 上， 只有 profiler 能回答。这一节是实测记录，包含两次自我更正和一个意外发现。

14.1 怎么在这套栈上抓 kernel 级 trace

vLLM 自带 torch profiler，不用 nsys/ncu（镜像里也没有）：

--profiler-config '{"profiler":"torch","torch_profiler_dir":"/root/prof","torch_profiler_with_stack":false}'
# 然后 POST /start_profile → 跑几次生成 → POST /stop_profile
# （路由在 vllm/entrypoints/serve/profile/api_router.py；VLLM_TORCH_PROFILER_DIR 已废弃，只会 warning）

三个坑，每个都值一次完整的重启周期：

① trace 是 pretty-print 的 JSON，字段分行、冒号后有空格。我第一版流式解析器写的是 紧凑正则 "cat":"kernel"，于是一个 kernel 事件都没匹配到 —— 而那一份 1.48 GB 的 trace 里其实有 273,941 个 kernel 事件。我据此得出了"CUDA 图把 kernel 藏起来了"的错误结论， 又花一次启动去用 --enforce-eager 重抓。

② 图模式（PIECEWISE）里 kernel 是可见的，不需要 eager。而且 eager 会歪曲你要测的东西：

图模式（生产配置）: 单步 86.0 ms，其中 kernel 时间 81.6 ms（95% 忙）
--enforce-eager  : 单步 166.7 ms，其中 kernel 时间 83.5 ms（49% 忙）

eager 把一半时间变成了 CPU 发射空隙（2153 个 kernel/步要一个个发），恰好放大了你正在 研究的"小算子低效"。要看真实占比，必须测图模式。

③ 解析要流式：380 token 的 decode 解压后 1.48 GB，而 gpu_memory_utilization 0.8 下宿主只剩约 3 GB 可用，整体 json.load 会 OOM。

14.2 单步解剖（MTP3，80 步平均）

单步墙钟      86.0 ms
  ├─ kernel   81.6 ms（95%）
  └─ 空隙      4.4 ms
kernel 数     2137 个/步

前 60%（≈52 ms）  目标模型 48 层堆栈：humming 内核，每个 0.02-0.44 ms，合计约 40 ms
后 40%（≈34 ms）  草稿阶段：结构完全不同，被**一个** cuBLASLt kernel 家族主导

每步 kernel 时间 Top：
  21.28 ms × 15.9  nvjet_sm110_tst_128x8_64x12_2x1_v_bz_TNT      ← 全步的 25%
  17.02 ms × 50.2  humming <Shape<0,1280,2560>, <16,256,64>>
   6.93 ms × 50.2  humming <Shape<0,2560,640>,  <16,256,64>>
   6.21 ms × 35.4  humming <Shape<0,16384,2560>,<16,256,64>>

14.3 真凶：未量化的 bf16 lm_head，每步被读 4 次

把那个 kernel 的时长和 grid 配对之后，它自己分成了两群：

grid[1940,1,1]   520 次   中位 4958 µs   ← 慢的全部在这里
grid[20,1,1]    1421 次   中位    53 µs   ← 便宜（每 SM 一个 block）
grid[100,1,1]    127 次   中位   249 µs

1940 = 248320 ÷ 128 —— 而 checkpoint 里

lm_head.weight : BF16 [248320, 2560] = 1,271,398,400 字节（1.27 GB）

即：没有量化的输出投影，每步被读 4 次（1 次目标验证 + 3 次草稿），每次约 5 ms：

4 × 1.27 GB = 5.09 GB/步  ≈ 20.3 ms/步 = 单步的 24%

而且草稿没有自己的头：models/qwen4_exp/nvidia/mtp.py 把 mtp.shared_head.head. 重映射到 lm_head.（第 107-116 行），Qwen4ExpMTP.compute_logits 直接 self.logits_processor(self.lm_head, hidden_states) —— 草稿和目标共用同一份 1.27 GB。

顺带解释了两个此前没想明白的现象： - 为什么 c=8 能到 130-154 tok/s：8 路序列共享一次 lm_head 读取，这块固定成本被摊薄了； - 为什么单流那么难上 40：单流时这 20 ms/步里每次只产出 1 个 token 的 logits。

14.4 试修：给草稿的那个 head 换 FP8 副本

思路：目标验证仍走 bf16（保输出分布），只把 3 次草稿的读取换成 FP8 副本（1.27 → 0.64 GB）。 先做微基准验证收益存在：

bf16  F.linear : 4.82 ms
fp8   _scaled_mm: 2.49 ms   → 1.94×，误差 mean 0.0155 / max 0.09

补丁做法（mtp_logits_fp8.py，运行时 .pth 加载，VLLM_MTP_LOGITS_FP8=1 才启用）： 包装 Qwen4ExpMTP.compute_logits，在调用期间把 self.lm_head.quant_method 换成 FP8 替身、调用完再换回来 —— 这样即使草稿和目标的 head 是同一个对象也不会串味 （目标验证在时序上不会与草稿并发）。权重按 8192 行分块量化，避免 float32 中间量爆显存。

实测（同负载，代码改写）：

步长       bf16 82.3-84.8 ms  →  FP8 74.8 ms      省 7.5-10 ms ✓ 与预测的 3×2.33 ms 吻合
接受率     bf16 2.60 / 2.80 / 2.80 / 2.91（四次采样）  vs  FP8 2.58  → 在噪声带内，无证据变差
吞吐       bf16 33.0-35.4 tok/s  vs  FP8 34.5 tok/s（单次）→ 分不出胜负

两处必须自己更正的地方（写在这里因为它们都是"看起来对"的测量陷阱）：

1. 我打印的"FP8 每次 0.11 ms"是假的：torch._scaled_mm 是异步的，我只量到了 CPU 入队时间。跨配置比较必须看端到端的步长，不能信单次调用的计时。
2. "只改草稿、输出逐字节不变"是错的。vLLM 的 rejection_sample_method 默认 "standard"（概率式拒绝采样，见 vllm/config/speculative.py），被拒绝位置的替代 token 从 max(0, p_target − p_draft) 残差里重采样 —— 输出确实依赖草稿分布，实测同一组 temp=0 提示词 3 条里 2 条分叉。任何"改草稿侧不影响输出"的说法都必须实测验证。

结论：带宽收益是真的（步长稳定降 7.5-10 ms），但我的量化太粗（per-row absmax → 一个 2560 长的点积只有 1 个 scale，logit 误差 ~0.09，而 top 候选之间常常只差 0.1）， 接受率刚好把收益抵消到测不出。要吃到这 7.5 ms，得改成按 K 分块量化（MXFP8 风格， 每 32-128 元素一个 scale），把误差降一个量级。

14.5 另一个测量陷阱：冷启动

同一配置、同一负载，启动后第一次测量总是显著偏慢：

第一次（刚启动完）  25.53 tok/s，步长 101.9 ms
之后同配置         33.01 / 34.60 / 35.36 tok/s，步长 80.9-84.8 ms

原因是 PLE 表的 mmap 冷页：第一次抓行要从磁盘读。这解释了本文档里几处"同配置数字对不上" 的现象 —— 永远不要拿冷的第一跑和热的第二跑对比。

结语

这台机器的故事其实是软件适配的价值远大于硬件规格的标准案例： 2070 TFLOPS 的宣传数字和 273 GB/s 的理论带宽，在日常 decode 里一个都没用上 —— 真正起作用的是 193.6 GB/s 的实测读带宽、CUDA 图捕获、投机解码的接受率曲线， 以及"把 26.82 GiB 的表从显存里搬出去、但记得给宿主留 15 GiB"这种琐碎的内存算术。

而这个故事还有一个更讽刺的维度：Thor 的硬件比 DGX Spark 更强 （sm_110 有 tcgen05、矩阵吞吐约 2 倍），但整个 CUDA 生态是为 12.x 消费级家族 （含 sm_121 的 Spark）编译的 —— 于是"硬件更强的那台"反而拿不到能跑满它的内核， 只能退回 Marlin 和 Triton 回退路径。硬件在，软件不在，这就是全部差距。

一路上最能省时间的经验只有一条：vLLM（以及所有这类框架）的 docstring 里写的约束 是真的约束。我踩的四个坑 —— 装饰器、地址稳定、原地写、op 名字 —— 全都白纸黑字写在它自己的注释里，只是没人会先去读。

如果只留一个数字作结，我会留这个：每步 86 ms 里，有 20 ms 花在把一个 1.27 GB 的 bf16 输出投影读 4 遍（第十四节）。它不是算力瓶颈、不是内核缺失，只是"这个张量没人量化"。 在统一内存的机器上，这类"没人做过的小事"往往比 2070 TFLOPS 的宣传数字更能决定体验 —— 这和本文开头的结论是同一件事：软件在不在，比硬件强不强重要。

2026-09-15 补记

十五、上线：把模型变成一个可计量、可限流、可多租的服务

为什么不能用 vLLM 自带的 key

--api-key 只接受一把静态 key，给多把也只是"等价的多把"（鉴权处做的是全等比较）， 于是不能按 key 限速、限并发、分开计量；更麻烦的是它是启动参数 —— 在这台机器上加一个 key 要重启，而重启一个 98 GiB 的权重是 5~12 分钟。

结论：多租控制必须自己加一层。模型服务负责算，网关负责管。

拓扑（端口对调后）

客户端 / harness / 各个 agent 实例
        │  Authorization: Bearer <各自的 key>
        ▼
  0.0.0.0:8000   ← 转发网关（零依赖 Node，挂在 tmux 里）
        │  换成上游那把 key 再转发
        ▼
  0.0.0.0:8001   ← vLLM（qwen3.8-flash-next / -uncensored）

对调只做了一次：8000 从模型改成网关卡，模型挪到 8001。客户端侧零改动 —— 把原来打 8000 的那把 key 注册成网关的一个身份，客户端仍然打 8000、仍然用那把 key， 只是从"直连模型"变成"过网关"，配额与账本立刻生效。

网关做了什么

同一身份（max_concurrent = 1）连发 3 条请求，排队是真在起作用，而且可计量：

（端到端一次真实对话 1.09 s，那一次携带了 16K 的固定前缀。）

计量口径的一个坑

非流式响应里有 usage，好办。流式响应默认没有 —— vLLM 只在客户端带 stream_options: {include_usage: true} 时才在最后一个 chunk 里回 usage。 所以网关的对策是：有 usage 就按 token 记；没有就只计次，token 写 null 并标 usage_absent，绝不猜数字。要么让客户端带上这个参数，要么承认这条请求只有次数账。

十六、收尾两件悬案：本地 NVFP4 输出头，以及最终配置下的并发

§14.3 那个"每步被读 4 遍的 bf16 输出头"，换掉了

第十四节把单步 86 ms 里的 20 ms 归因给一个未量化的 bf16 lm_head（1.27 GB × 4 次）。 本地把输出头做成 NVFP4 之后（不动镜像，运行时替换），同一负载下：

质量判分持平（不是"看起来还行"，是同一套判分脚本）。这一步的价值不在平均值， 而在它证明了第十四节的归因是对的：那 20 ms 确实是"没人量化"，不是算力不够。

uncensored 权重挂上同一台服务

同一个 server 暴露两个模型名（qwen3.8-flash-next / qwen3.8-flash-next-uncensored）， 同一套启动参数：SEQS=4、--gpu-memory-utilization 0.73、MTP2、text-only、 humming 原生内核、PIECEWISE CUDA 图、async-scheduling，外加 --reasoning-parser qwen3 与 --tool-call-parser qwen3_xml --enable-auto-tool-choice。

uncensored 权重单流 51.99 tok/s；同一台机器上的并发表（KV 池 1,013,816 token 那一次）：

一个必须诚实记录的现象：KV 池会"抽签"

同一套参数连续三次启动，KV 池拿到的是：1,013,816 / 557,599 / 557,599 token， 而 vLLM 自报的 Actual usage 是 59.5 GiB 和 72.4 GiB —— 差 13 GiB， 不属于启动参数能解释的部分（当时我一度以为是漏了环境变量，实测推翻：带全量 env 重启后 池子依旧是 557,599）。

所以两句话：别指望靠重启"抽到"一个大池子；要稳就调 --gpu-memory-utilization。 以及：报容量时必须报"你要的那个配置下这一次启动"的数字，不要报历史最好值。

顺带一条容量账（以 557,599 token 的池子为例）：

这一版 vLLM 的前缀缓存是开着的（日志里 Mamba cache mode is set to 'align' ... when prefix caching is enabled）。这一条对多租服务很重要：十几个 agent 共用的那段 固定前缀（工具定义 + 系统提示，见下节）在 KV 里只存一份，不会按用户数翻倍。

十七、两条花掉我最多时间的教训

教训一：复刻一个跑着的容器，不要用白名单过滤环境变量

为了做 A/B，我用"看起来相关"的几个环境变量重建了启动脚本。漏掉了 4 个， 其中一个叫 VLLM_API_KEY —— 后果是新容器完全没有鉴权： /v1/models 不带 key 返回 200。这台机器对局域网可达，等于开了个后门， 而我在自己的测试里一直只验证"带 key 能通"，所以第一轮检测没发现。

同批漏掉的还有三个调优开关（VLLM_MARLIN_USE_ATOMIC_ADD、VLLM_MTP_LOGITS_FP8、 VLLM_PLE_MMAP_STATS），它们会让内存/性能悄悄偏离现场，而不报任何错。

正确的做法只有一种：

docker inspect <name> --format '{{json .}}' > restore/<name>.inspect.json
# 全量导出 env（只排除由镜像自带、改错会坏启动的 PATH / LD_LIBRARY_PATH）
python3 restore/extract_env.py          # 写出 restore/vllm.env（600）
# 启动脚本里引用它，不要再手写 -e
docker run ... --env-file restore/vllm.env ...

以及一条检查清单：起来之后第一件事是验鉴权 —— 不带 key 必须 401。 "带 key 能通"证明不了任何事。

教训二：agent 的提示词固定开销，是多租服务的隐形天花板

把 agent 接上模型之后，第一件让我意外的事是每轮的真实输入 token：

差一个数量级。原因是 ZeroClaw 把 SOP 工作流、定时任务、子代理、跨机消息、技能包、 甚至"硬件访问（Arduino）"这些都作为工具和段落塞进了提示词 —— 对一个 QQ 聊天机器人， 大部分是纯税。而且这份税每一轮都要重新付一次 prefill（前缀缓存能把它在 KV 里 只存一份，但算力和带宽仍要为它做一次前向）。

试着裁剪过：schema 里确实有 excluded_tools（36 个通道级配置类型都带这个键， 描述是"不暴露给模型"），但我把它写进 [channels.webhook] 之后， daemon 直接报 config section ... is malformed and was reset to defaults， 重启后 input_tokens 一点没变（16,009 → 16,009）。所以这条开关的实际效果未知， 我没有验成 —— 记在这里，而不是写成一个结论。

给做多租服务的同行的结论只有一句：在选 agent 内核之前，先把它的固定提示词开销量出来。 模型侧省下来的那 15% 单流性能，很容易被 agent 侧每轮多发的 15K token 吃掉。

以上为 2026-09-15 补记；原文（第一 ~ 十四节 + 结语）保持定稿状态未改动。