一台 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， 同时把多模态开了回来。下面把这台机器上真正发生的事、每条数字背后的算术、 以及踩过的每一个坑写清楚。所有数字都是本机实测，公开数据会单独标注。

TL;DR

一句话结论：第 2 步是最大单项收益（2.58x）；第 4 步不提升速度，它是解锁项 —— 不做它，KV 只剩 5.7 GiB，MTP 之后连多模态的 3 GiB 都塞不下。

一、硬件：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，看衰减曲线。

顺带一个测量陷阱：别用 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 是带宽限制），而这正好就是缺失的原生 FP4 MoE 内核所影响的区间 —— 1035 TFLOPS 的 FP4 张量核在 MoE 上闲置，退回 Marlin （权重解压后再算）。同理，并发上去之后每个 token 的算力摊薄，compute-bound 的成分上升， 缺内核的代价会被放大。这条差距是真实且可归因的。

小结

单流      → 已到带宽屋顶（100%），配置级优化没有空间了
并发      → 被我们自己的 max-num-seqs=2 锁死，是"人为"差距，可立刻解除
预填充/高并发 → 真实的算力缺口，只能靠补内核
跨平台比较  → Thor(sm_110, 有 tcgen05) vs Spark(sm_121, 无 tcgen05)，不是同一档

十、接下来做什么

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

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

1. 解除并发上限：--max-num-seqs 2 → 8（或多档扫 4/8/16）， --max-num-batched-tokens 2048 → 8192。KV 池有 446k token 容量（32k 上下文 13.6 路并发）， 完全支撑得起。这是当前最大的、零成本的收益项：参照 rittim 的 4 路 51 与 MiaAI 的 8 路 162.9，我们的聚合吞吐应该从 43 量级跳到 100+ 量级。
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. 原生 sm_110a FP4 grouped-MoE 内核。torch / FlashInfer 里已经有 sm_110a 的 FP4 CUTLASS GEMM（59 个 kubins），断掉的是 block-scaled grouped MoE 那一条。 收益主要在 prefill 与高并发（compute-bound 区间），单流基本无感 —— 别指望它提升 33。
5. 编译 fused GDN MTP post-conv 的 SM110 cubin（Leibniz LIMITATIONS 第 2 条）。 现在草稿路径走 Triton 回退，每步草稿开销偏高，限制了"每步产出 token"的性价比。
6. 修 QSA cooperative Top-K 在 SM110 上的 cluster 几何（第 3 条）。 顺带能去掉 persistent_topk 这条非确定性风险源。
7. 排查 Thor 上 PLE GPU gather（ATS）的非法访问。修好之后抓行可以留在 GPU 侧， 省掉每步 host/device 往返（第 5 条指向同一处）。

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

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

一句话

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

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

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

1. 没有原生 sm_110 FP4 MoE 内核 —— 这是唯一还有 2x 空间的地方。 CUTLASS 的 run_fp4_blockwise_scaled_group_mm_sm100 在 sm_110 上初始化失败 （status=7），于是目标 MoE 退回 Marlin 路径（权重解压），1035 TFLOPS 的 FP4 张量核在 MoE 上基本闲置。torch/FlashInfer 里明明有 sm_110a 的 FP4 CUTLASS GEMM，但 grouped-MoE 那条路是断的。
2. Thor 的 ATS GPU gather 非法访问，PLE 抓行只能走 CPU（约 93% 的加载路径会踩到）。
3. 融合版 MTP GDN 内核在 SM110 上不可用：fused_gdn_decode_post_conv_mtp 没有 SM110 cubin，只能退回 VLLM_GDN_DECODE_KERNEL=triton 通用路径， 草稿路径的每步开销因此偏高（Leibniz 的 LIMITATIONS 第 2 条是同一个缺口）。 注意别把这条过度解读：MTP 深度本身是可用的 —— 我们实测 K=3 确实产生了 3 个草稿 token，逐位接受率也拿得到；K>1 不划算的原因是接受率衰减，不是内核坏了。
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 tok/s 量级；4 并发 43 tok/s 量级

两处已知待调：--max-num-seqs（本篇用的是早期省内存时的 2，KV 池其实支撑 8-16 路） 和 --max-num-batched-tokens（2048 → 8192）。详见第十节 P0-1。

结语

这台机器的故事其实是软件适配的价值远大于硬件规格的标准案例： 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 名字 —— 全都白纸黑字写在它自己的注释里，只是没人会先去读。