一次只搬一个字节：把 CS336 的推理课搬进 DeepSeek DeepInfer 的源码里

说明 本文有两条底稿。一条是本博客早先的《Stanford CS336 第十讲精读：推理》（Percy Liang 主讲，课堂上的直觉、账本和五种 KV 瘦身法），另一条是 DeepSeek 在 2026-10-16 开源的 deepseek-ai/DeepInfer（Apache-2.0）——我按 2026-10-20 的 main 通读了它的 Python 引擎、C++ 调度器和 TileLang 内核目录。文中所有代码引用都给出了仓库内路径，可以直接对照。 需要预先声明三件事：① 模型结构的具体数字（40 层、384 专家取 6、1M 上下文、Engram 196.6B 等）来自本系列早期文章所引的官方 recipe，不是我这次独立核实的；② 这次开源只提供离线推理，只支持 DeepSeek-V4.1-Flash，且明确列出 KV cache 系统、PD 分离、跨 DP 副本负载均衡“未包含在本次发行中”——所以我们读到的是一台拆开的引擎，不是一套完整服务；③ 凡是我从代码推断、而非代码直接写明的结论，我都会用“推断”标出来。

开场：请把自己想象成一个搬运工

伽莫夫写过一本《物理世界奇遇记》，主角汤普金斯先生是个银行小职员，一做梦就掉进各种被放大的物理世界里。这一篇我也想借他一个梦。

请把自己想象成 GPU 里的一个线程。你的工作不是算数——算数很快，快得几乎免费——你的工作是把数据从显存里搬过来。显存是个大仓库，你手上只有一个小小的操作台（寄存器、共享内存），一次只能搬一点点。你搬过来的每个字节，都花掉了时间。

于是整个推理系统的命运，就落在同一个问题上：每搬一个字节，我能换到多少次浮点运算？

这个比值有个名字叫算术强度。它是一把尺子，CS336 第十讲几乎所有的结论都是用它量出来的，DeepInfer 里几乎每一个开关也都是为它服务的。这篇文章只有一件事：拿着这把尺子，把 DeepInfer 的源码从头量一遍，看看哪些地方它只是把课上讲的东西做成了工程，哪些地方它走出了课堂。

TL;DR

• 尺子：CS336 给的算式是“矩阵乘的算术强度约等于 batch size B”，而 H100 的加速器强度约 295 flops/byte。B > 295 是计算受限，B = 1 是访存受限。推理的 generation 阶段永远在 B=1 这一侧。
• DeepInfer 的账：官方 benchmark 用 4 张 SM100 量 prefill、用 64 张量 decode；一颗稳定的 decode 步是“64 个请求 × 每请求 1+5 个 token = 384 个 token”。一步只有三百多个 token 在过矩阵乘法——剩下的全是在搬权重、搬 KV。
• 课堂上的五种 KV 瘦身法，DeepInfer 是四个一起上：单 KV 头 + MLA 低秩投影、128 的滑窗、压缩 KV latent + 稀疏 indexer、跨层共享 KV 源。它还在课堂之外加了一个 YOCO：后半数的层根本不自己算全局 KV，直接复用前半截的隐藏态。
• 课堂说“想想你的操作系统课”，DeepInfer 抄得比操作系统还狠：分页、引用计数、抢占重算、连续批处理都在，还多了一个课堂没讲的 零开销调度（ZOS）——CPU 在 GPU 跑当前步的时候，已经把下一步的输入张量准备好了。
• 课堂没讲的三件事，我认为是这次开源的真正分量：① 把“可复现”当成一等公民（批次无关推理 + 确定性采样）；② 把“记忆”从神经网络里拆出来（Engram 哈希表）；③ 把“验证”从算法变成调度问题（DSpark 的置信度调度 + 在线时延模型）。
• 一个反直觉的收束：CS336 说“最大的机会在架构层，而不是系统层”。DeepInfer 恰好是这句话的注脚——它把系统层的技巧全都做了，但真正让它成为 V4.1 官方推理栈的，是那些架构级的省（YOCO、压缩 KV、Engram、超连接、DSpark）。

一、第一幕：这把尺子怎么量一台真实的引擎

先复习那把尺子。注意力这一层最要命的地方，课堂上讲得比谁都直白：

MLP 的权重是所有人共享的，所以大批量很划算；而 attention 的 KV cache 是“每条序列自己的雪花”，把请求拼成 batch 并不会把它摊薄。

所以，批量化救不了 attention。MLP 的强度是 B·T，prefill 时 T 很大、强度很够；generation 时 T=1，只能靠并发请求数 B 撑。而 attention 的强度是 S·T/(S+T)，永远不超过 1，而且完全不依赖 B。

现在把它搬到 DeepInfer 的参数上。README 写着性能基准“measured on 4 SM100 GPUs for prefill and 64 for decode”——注意这两个数字的落差：prefill 用 4 张卡就够了，decode 要 64 张。这不是随便选的。prefill 是计算受限的，4 张卡就能把算力吃满；decode 是访存受限的，单张卡的带宽早就见底了，唯一的办法是拿更多的卡去摊。

再看 examples/README.md 里那句口径：

The decode output length is derived from an internal step-count constant … the 64 requests of one decode step account for 384 tokens in total. … decode throughput (tokens/s/GPU): avg_SPS × decode_num_requests × (1 + nextn) / --cp, under full-acceptance accounting.

64 × (1 + 5) = 384。一颗 decode 步，整台 64 卡的机器，总共只有 384 个 token 在过矩阵乘法。 而这一步要把几百 GB 的权重从显存里搬一遍。这就是那把尺子量出来的结果，也是全文所有设计的出发点：

• 要么少搬一点（KV 瘦身、量化、稀疏）；
• 要么让同一批搬运服务更多 token（投机解码、连续批处理）；
• 要么把搬运和别的事重叠起来（异步调度、CUDA graph）。

DeepInfer 的每一个开关，都可以归类到这三条里。

二、第二幕：KV cache 的瘦身法，从“五选一”到“四个一起上”

课堂第六节列了五种瘦身法：GQA（少几个 KV 头）、MLA（投影到低维隐空间）、CLA（跨层共享）、局部注意力（只留最近 K 个 token）、以及把它们组合起来。DeepInfer 的处理方式几乎是“能叠的全叠上”，而且叠得很干净。

2.1 先从配置形态看：一层一个“态度”

打开 deepinfer/hugging_face/config.py，会发现这份模型配置字典里最显眼的是一个叫 compress_ratios 的列表——每一层一个数，这个数决定这层对 KV 的“态度”：

def _build_attn_module_list(text: dict[str, Any], shared: AttrDict) -> list[Any]:
    """n_layer backbone layers + 1 lm head placeholder (None) + nextn DFlash layers."""
    ratios = text["compress_ratios"]
    ...
    specs: list[Any] = [
        _win_spec(text) if ratio == 0 else _compress_spec(text, ratio, shared)
        for ratio in ratios[:n_layer]
    ]

• ratio == 0 → 这一层是纯滑窗层，sparse_stage=0，注意力只看最近一小段；
• ratio > 0 → 这一层是带压缩的层，除了滑窗还能回溯到更远的位置，靠的是压缩后的 KV latent。

而且配置层会拒绝一切它不支持的形态，粗暴得令人安心：

def _assert_supported(hf: dict[str, Any]) -> None:
    text = hf["text_config"]
    # The runtime requires SwiGLU, bias-free attention with zero dropout,
    # one shared KV head and BF16 model parameters.
    for key, expected in {
        "hidden_act": "silu",
        "attention_bias": False,
        "attention_dropout": 0.0,
        "num_key_value_heads": 1,
    }.items():
        if text[key] != expected:
            raise ValueError(...)

num_key_value_heads == 1——整台模型只有一个 KV 头。这不是 GQA 的“少几个头”，而是把 KV 的表示彻底换成了 MLA 那一套低维隐向量。课堂里说 MLA 是“把 KV 投影到低维隐空间”，这里是它的极端形态：KV 不再是 [头数, 维度] 的堆叠，而是一份共享的隐向量，配合 RoPE 分开存：

# HF's shared KV vector includes RoPE; the internal rank excludes it.
kv_lora_rank=text["head_dim"] - text["qk_rope_head_dim"],

2.2 MLA 的低秩投影长什么样

deepinfer/model_executor/models/v4_modules/mla_attn.py 里，MLA 的投影被写成一个很朴素的类：

class ColumnParallelLora(nn.Module):
    """`x -> A -> (norm) -> B` MLA low-rank projection;
    when `out_features == 0` there is only A (e.g. Wkv)."""

注意那句注释里的 out_features == 0 就只剩 A——KV 那一侧根本没有“B”，它只做一个降维投影，产出的就是那根要长期缓存、反复搬运的低维 latent。少一层反投影，就少一次搬运。这是 MLA 省显存的全部秘密，被压缩成一句参数注释。

2.3 128 的滑窗：把“看哪”写成一个索引数组

滑窗那一级的物理形态很朴素：每个 token 只精确地看最近一段。但它在两种模式下要用不同的方式表达——训练/推理时的稠密窗口，和 KV 缓存被“压缩 + 分块”之后的稀疏窗口。于是有了 deepinfer/model_executor/layers/tilelang_ops/swa_index.py：

def expand_swa_index(pivots: torch.Tensor, width: int, out: torch.Tensor | None = None) -> torch.Tensor:
    """Expand two KV segments per token into dense prefill SWA indices.

    Args:
        pivots: Contiguous int32 [num_tokens, 4], containing first offset,
            first length, second offset, second length. ...

    Returns:
        Int32 [num_tokens, width] indices into the gathered KV buffer, with
        unoccupied columns set to -1. CP reordering is applied by the caller.
    """

每个 token 的可见 KV 被描述成两段：一段是最近的滑窗，一段是远处的压缩背景。整个稀疏注意力，在实现层面就是“把这两段翻译成一个合法的列索引数组，空位填 -1，然后交给一个支持稀疏索引的 attention kernel”。

课堂里那两句直觉（“近处老老实实看”、“远处看摘要”），在代码里就是这四元组 [first offset, first length, second offset, second length]。

2.4 稀疏 indexer：先粗筛出一小块，再精看

“看摘要”也要挑。全部摘要都看，等于没省。于是有一个小型 indexer 负责打分。deepinfer/model_executor/models/v4_modules/index_attn.py 的开头就写明了它的职责：

"""Sparse indexer runtime: score compressed KV per query and select the top-k blocks.

`MQA` computes the `relu(q @ k) * weights` logits that drive sparse attention, and keeps the
prefill and decode metadata, candidate-block restriction and cross-layer index-K sharing in
sync with `cross_layer.py`.
"""

relu(q @ k) * weights——一个不带 softmax 的内积打分器，这就是所谓 lightning indexer 的味道：它不要求概率语义，只需要“谁更相关”的排序。

真正有意思的是“挑块”这一步。candidate_blocks.py 里先把每个块的分数压成一个标量：

def candidate_block_max(logits, valid_lens, block_size):
    """Compute block scores and valid block counts for candidate selection.

    Each score is the maximum valid token score in the block, capped at the
    largest finite bf16 value. The last valid block gets +inf so
    that candidate selection always includes it. ...
    """

三个细节，每个都是经验写成的：

1. 块内取 max 而不是取和：只要块里有一个强相关的 token，这个块就值得被看——注意力本来就是“被最相关的那一个拉过去”，不是被整块的平均值拉过去。
2. 分数被钳到 bf16 的最大有限值：因为后面要拿它做 top-k 比较，不能让一个 inf 污染排序。
3. 最后一个有效块被强制设成 +inf：保证“最近的一小段”永远入选。这是把架构假设（近处必须看）硬编码进打分器，而不是靠学出来的权重去保证。这是我很喜欢的一处工程味：能靠结构保证的事，就不要交给训练去赌。

2.5 跨层共享与 YOCO：把“谁负责写 KV”变成一张拓扑图

课堂讲的 CLA 是“连层与层之间也共享”。DeepInfer 把它做成了一件更结构化的事：配置里给出若干个源层，其余层按“最近的、不晚于自己的那个源层”去认领：

def _source_map(source_ids: list[int], n_layer: int) -> dict[int, int]:
    """Map each layer to the nearest source layer that is not later than that layer."""
    ids = sorted(source_ids)
    return {layer: max(s for s in ids if s <= layer) for layer in range(ids[0], n_layer)}

然后这套映射被塞进 SharedLayerTopology（v4_modules/cross_layer.py），成为运行时判断“我这一层到底要不要自己算 KV、要不要写缓存”的依据。mla_attn.py 的文档字符串把这件事说得最清楚：

a layer with layer_map[l] == l is a source layer, and the remaining layers reuse the modules of the source layer.

到这里都还在课堂范围内。但接下来这一步走出了课堂。

2.6 YOCO：CS336 没有讲的那一级省法

config.py 里有一个不起眼的函数：

def _n_prefill_layer(text: dict[str, Any]) -> int:
    """The first layer with ratio==1: it and everything before it form the
    prefill side, the layers after it are yoco consumer layers."""
    return next(i for i, ratio in enumerate(ratios) if ratio == 1)

n_prefill_layer。它一旦小于总层数，模型就进入了 YOCO（You Only Cache Once） 模式。v4_modules/yoco.py 的开头解释了它的两条谓词：

def yoco_enabled(cfg) -> bool:
    """Whether the model enables YOCO
    (layers after n_prefill_layer are consumer layers)."""

def is_yoco_consumer_layer(cfg, layer_idx) -> bool:
    """YOCO consumer layer: [n_prefill_layer, n_layer) in the main model."""

而真正说明它为什么值钱的是第三个函数的长注释：

def yoco_shared_hidden_required(cfg, shared_topology) -> bool:
    """Whether the YOCO shared hidden (the attn-norm hidden of the first consumer layer)
    must be captured/passed on, i.e. whether there is a consumer layer that consumes it.
    There are two consumption points:
      1. Layers with SWA under yoco_swa_share_hidden: kv_swa is projected from the shared hidden;
      2. Packed kv-source layers after the first consumer: the shared hidden computes the new
         compressed KV. ...
    """

读懂了：后一半的层不再自己计算全局 KV，而是消费前一半留下的那个隐藏态（shared hidden）。按早期文章里的说法，V4.1 的 40 层被切成前 20 层“编码器”和后 20 层“解码器”，解码器的全局 KV 直接从编码器最后一层的隐藏状态投影出来。

为什么要这么干？还是那把尺子：prefill 一次算出的隐藏态，如果能让 20 层共享，就等于把“写全局 KV”这件事的搬运量砍到几分之一。代价是表达能力受限（后一半层的全局视野全都来自同一个投影），收益是长上下文真的能跑起来。

这是课堂上完全没有的一类想法。 课堂里的五招都在改“KV 长什么样、存多少”；YOCO 改的是**“KV 由谁负责算”**——它把 KV cache 从“每层各自维护”变成了“整台模型维护一份、按需分发”。我个人的判断（推断）：这是这次开源里最容易被忽视、但对长上下文最实惠的一招。

2.7 那一份 KV 记录到底有多小

抽象说了这么多，落到字节上才踏实。deepinfer/model_executor/layers/tilelang_ops/cache_common.py 里有一张格式表：

V41_FORMAT = KVCacheFormat(
    quantized_dtype="e4m3",
    quantized_elements=512,
    elements_per_sf=32,
    bf16_elements=0,
    do_cast=True,
    sf_format="ue8m0",
    note="V41, main/swa kv.",
    token_interleaved=True,
)

一个 token 的主/滑窗 KV 记录 = 512 个 FP8(e4m3) 数值 + 512/32 = 16 个 ue8m0 缩放字节 = 528 字节。 而这 528 字节里，没有一个是 BF16——bf16_elements=0，注释甚至写着一句“mixed quantized/BF16 KV caches are no longer supported”，混合精度 KV 已经被从这个系统里删掉了。

作为对照：如果按 BF16 存一份常规的多头 KV，同样的位置大概是几千字节量级。早期文章引官方 recipe 给出的“每 token 全局 KV 约 890 字节”这个数字，量级上是同一回事（那份账还包括 indexer 的 FP4 索引 KV，FP4_INDEX_KV_FORMAT）。

到这里，物理量已经很小了。于是第二幕的结论其实是个拐点：KV 一旦不再是主要矛盾，主要矛盾就重新回到了“搬运”本身——你每生成一个字，还是要把 511 GB 的权重从显存里搬一遍。这就是第三幕要打的仗。

三、第三幕：投机解码——课堂讲算法，DeepInfer 讲调度

课堂第九节的投机解码讲得很漂亮，而且必须承认它漂亮：用便宜的方式猜，再用贵的方式并行验证；因为验证是一次 prefill（计算受限、很快），所以“检查比生成快”；而它是带修改的拒绝采样，因此保证得到目标模型的精确样本。课上给的经验数字是约 2×，Medusa 让草稿并行猜、Eagle 让草稿吃目标模型的隐层特征。

DeepInfer 把这套东西实现成了一个叫 DFlash 的 MTP（multi-token prediction）模块，并且在上面加了三层课堂没讲的东西。

3.1 草稿头：噪声 token、Markov 头、置信度头

deepinfer/model_executor/models/mtp/dflash.py 里，草稿头的构造透露出很多设计意图：

self.noise_token_id = cfg.dflash_cfg.dspark_noise_token_id
self.dflash_block_size = self.max_MTP_value          # 每个 seq 生成的 propose 个数
...
self.markov_rank = cfg.dflash_cfg.get("dspark_markov_rank", 0)
self.enable_markov = self.markov_rank > 0
...
self.target_layer_ids = cfg.dflash_cfg.dspark_target_layer_ids
self.num_target_feature_layers = len(self.target_layer_ids)
self.projection = ReplicatedLinear(
    cfg.hidden_size * self.num_target_feature_layers,
    cfg.hidden_size,
    weight=modules[0].projection.weight,
)

三件事值得停下来看：

1. noise_token_id——草稿位置要填一个“噪声”token。 这一步几乎是必须的：既然草稿是并行猜出来的（不是一个一个自回归跑出来的），那第 k 个位置的输入就必须是“上一轮的猜测结果”或者一个占位符。用一个专门的噪声 token 明确告诉模型“这里是猜的，别当真实历史”。（推断：这是 DeepSeek 侧对“并行草稿”的处理方式。）
2. Markov 头：markov_w1（embedding）+ markov_w2（投回词表，use_fp32_output=True）——它把上一个草稿 token 的信息加进当前草稿位置的 logits。换句话说，dspark_markov_rank 那点小参数，就是在并行草稿里找回一部分自回归的链式依赖。这和 Medusa 的思路是一脉相承的：并行出草稿，但用便宜的小修正量把 token 之间的关联补回来。
3. dspark_target_layer_ids：投影的输入是目标模型的若干中间层特征拼起来的。这正是 Eagle 的思想——草稿模型不应该是一个独立训练的小模型，而应该“寄生”在目标模型的隐层上。

然后是最有意思的那个头：

class ConfidenceHead(nn.Module):
    """Predict the acceptance probability of each draft token."""

它的前向就是：

def confidence_head_fused(hidden, weight, temperature, markov_emb=None, out=None):
    """Fused ``sigmoid(cat([hidden, markov_emb], -1) @ weight / temperature)``."""

一个会说出“我有多确定”的头。 输入是隐藏态（必要的话再拼上 Markov 嵌入），输出是一个 0 到 1 的接受概率估计。默认温度是 1.0（DEEPINFER_DFLASH_CONFIDENCE_HEAD_TEMPERATURE，也可以在环境变量里按 nextn 个位置各给一个温度）。

课堂上的投机解码，K 是个全局常数；DeepInfer 里有了一张可以预测每个草稿 token 会不会被接受的成绩单。这张成绩单就是下一节的主角。

3.2 置信度调度：把“验证预算”当成一种会被花掉的资源

这是我认为整个 DSpark 里最值得写的一段。deepinfer/model_executor/layers/tilelang_ops/schedule_mtp_varlen.py：

def compute_seqlens_q_by_confidence_rates(
    total_q: int,
    remap: torch.Tensor,
    prev_confidence_rates: torch.Tensor,
    seqlens_q: torch.Tensor,
) -> None:
    """Distribute total_q tokens across sequences based on MTP confidence rates.

    Fuses remap scatter: confidence_rates[i] = prev_confidence_rates[remap[i]] if remapped else 0.0.

    Args:
        total_q: total query token budget.
        remap: [num_seqs], int32, prev batch index or -1 for fresh seqs.
        prev_confidence_rates: [n_prev_seq, nextn], float32, per-position confidence from prev MTP head.
        seqlens_q: [num_seqs], int32 (out), per-sequence query lengths summing to total_q.
    """

请慢慢读这段签名，它信息量非常大：

• 有一个总预算 total_q——这一步全机器一次能验证多少行 query，是有限的；
• 有个上一轮的成绩单 prev_confidence_rates，形状 [序列数, nextn]，每个草稿位置各有一个置信度；
• 输出 seqlens_q 是每个序列这一步能验证多少个位置，且它们之和恰好等于总预算。

也就是说：置信度高的序列，多给它几个位置去验证；置信度低的，少给。 课堂上的 K 是一个常数，DeepInfer 的 K 是一个按请求动态分配、总和守恒的预算。

还有一个细节：remap 把“上一轮的批”和“这一轮的批”对齐，-1 表示新来的序列（置信度按 0 记）。批次是会变的——上一轮在这个位置上的请求，这一轮可能已经走了，可能被挤到别的位置。所以这个函数不只是分配，它内部顺手把 remap 的散射也融进同一个 kernel 里做掉了。这就是“调度”这个词在真实系统里的温度：它必须处理“世界在变”。

3.3 验证的落地：乐观写入，拒绝是廉价的

投机解码的验证，在 deepinfer/model_executor/layers/tilelang_ops/sampler.py 里被拆成两步：

def prepare_verification_metadata(ftp, seqlen_q, total_q, num_steps):
    """Build metadata for request rows packed along the ``total_q`` dimension.
    Returns:
        row_to_req: [total_q], int32. Request index for each packed row.
        local_step: [total_q], int32. Row offset within its request.
        is_verify: [total_q], bool. True for verification rows and false for
            each request's final fresh-sampling row.
    """

is_verify 这一列说明了一切：一个请求的 seqlens_q[i] 行里，前面几行是“验证草稿”，最后一行是“新鲜采样”。因为无论前面接受了几个、在哪一个位置被拒绝，最后总要有一次从修正分布里重采样——这正是课堂那段伪代码里“从 max(0, q-p) 里采一个然后收工”的位置。

而 finalize_verification 的签名里有一列叫 effective_accept（生效的接受结果）——“effective”这个词很讲究：拒绝之后，从那个位置往后的草稿即使本身是对的，也已经作废了。

最能说明工程成熟度的是 update_main_cache 的那段注释：

"""Write every DFlash layer's main-token keys from the target-layer features; ...
Decode 时随主模型 CUDA graph 一起 replay，此时接受结果尚未产生，全部 token 都按
slot_mapping_store_swa 写入：将被拒绝的 draft 位置正是 forward_drafts 随后写 draft key
的位置，在任何读取之前就会被覆盖，因此无需按接受结果屏蔽。
"""

翻译一下：验证还没出结果的时候，就已经把所有 token（包括可能被拒绝的草稿）的 KV 都写进缓存了。 因为被拒绝的那几个槽位，稍后写草稿 KV 时会按同样的地址覆盖掉，在任何读取发生之前就已经是干净的了。所以系统根本不需要“等结果、再决定写什么”——拒绝是廉价的，脏数据自己会被冲掉。

这是把“乐观并发”的思路用在了 KV 缓存上。课堂里投机解码的伪代码是串行叙述的（先验证、再决定）；真实系统里，写入发生在决定之前。

3.4 varlen：当每个请求的草稿长度都不一样

课堂的伪代码里，所有请求的 K 一样。DeepInfer 打开 DEEPINFER_ENABLE_MTP_VARLEN=1 之后，每个请求这一步的验证长度各不相同（就是上面那个 seqlens_q）。这带来一串连锁反应：

• DEEPINFER_USE_NEXTN_DECODE 默认 5，是“最多能猜几个”；显式设 0 就关掉投机解码。
• DEEPINFER_MAX_BACKTRACKING_TOKENS 的默认计算是 max(值, nextn + 1)——必须至少能回退一个整轮的草稿长度，否则“全部被拒”这种情况就没法收拾。
• YOCO 的截断约束（environ.py）写得更直白：

"""YoCo consumer prefill 截断开启时保留的末尾 query token 数。...
必须满足 KEEP_TOKENS >= window_size + DEEPINFER_MAX_BACKTRACKING_TOKENS；
默认 0，启用截断时必须显式设置。"""

一回事：只要你允许“猜了又反悔”，你就必须保证反悔的空间还在缓存里。 这是投机解码与所有“有界缓存”设计（滑窗、YOCO、PagedAttention）之间最本质的耦合——课堂里它是一条注释，这里是三个必须互相满足的不等式。

3.5 这一步的结论

课堂把投机解码讲成一个算法（拒绝采样 + 精确性证明）。DeepInfer 把它讲成一个调度问题：有预算、有成绩单、有守恒律（sum(seqlens_q) == total_q）、有“猜错也要能收拾”的空间约束。算法是死的，预算是活的。

四、第四幕：操作系统的课，DeepInfer 又上了一次——还加了一个在线辨识器

课堂第十节说得很直接：“想想你的操作系统课，那些东西可以直接搬到推理上来。” 于是就有了连续批处理（火车不等你）、选择性批处理（attention 单算、MLP 拼起来）、以及 PagedAttention（分页 + 写时复制 + 引用计数）。

DeepInfer 的对应物是 csrc/engine/ 下那套 C++ 调度器。

4.1 分页、抢占、重算

deepinfer/worker/cache_engine.py 里有一句很朴素的常量：

KVSTORE_GRAIN = 128
"""kvstore 落盘粒度：只在 token 数为 128 的倍数处落盘"""

和 csrc/engine/block_manager.hpp 一起，这就是 PagedAttention 的那套账。而 csrc/engine/scheduler.hpp 里的 decode 调度写得很克制：

/// decode（completion）调度：先把 prompted 队列中的 seq 移入 running，
/// 再按 LCFS 挑选本轮 decode 的 seq，block 不足时抢占 output 最短的 seq。

LCFS（后进先出）而不是 FIFO，这个选择值得停一下：在一个每一步都要重算耗时的系统里，先服务“最短的那些序列”能让它们尽快结束、腾出块来，从而让整个批次的压力下降。这和操作系统的短作业优先是同一个直觉。

“block 不足时抢占”则是那个“重算”（recompute）机制：被抢的序列 KV 被丢掉，之后重新 prefill 回来。于是有了 seq->reset_for_recompute() 这种必须把所有滚动状态（cache.num_tokens、在途计数、MTP 状态）一起复位的地方——一个状态漏了，就会变成那种“跑一万次才出一次”的 bug。

4.2 步型投票：全组必须“同型”

有一个工程细节，课堂上不会有，但真实的多卡系统里必须有。NORMAL 引擎每一步要么做 prefill、要么做 decode，而DP 组内的所有 rank 必须做同一种，否则集合通信会对不上。于是调度器搞了一次投票：

// 两类尝试点：
//   ALL 点（每 prefill_attempt_interval 个 decode step）：全员都有 prefill 活才做，
//     防止一个 rank 干活、其余 rank 陪跑空 prompt step。
//   ANY 点（每 prefill_any_attempt_interval 个 decode step）：一个 rank 有活就一起做，
//     给孤立 rank 兜底——它的 prefill 最多等这么多个 decode step。

用一次 allreduce(MIN) 决定全组这一步做不做 prefill。默认 70 和 100 这两个间隔（DEEPINFER_PREFILL_ATTEMPT_INTERVAL / _ANY_）背后是一个很现实的取舍：“陪跑”也是要花时间的——一个 rank 做 prefill，其余 rank 空跑一步，等于把整组的吞吐拉低。于是有了这套“多等等、一起干”的节奏。

这类东西永远不会出现在论文里，但它决定了集群的实际吞吐。读源码读到的，一半是这种“不得不处理的世界”。

4.3 零开销调度：CPU 在 GPU 跑的时候，已经在准备下一步

这是我这次读代码最喜欢的一个设计，README 里只用一句“overlap CPU batch scheduling with device execution”带过，但它其实是整个引擎的骨架。

先说问题。decode 每一步要做的事有两件：CPU 侧准备这一批的输入（token id、序列长度、块表、采样参数……），GPU 侧执行前向和采样。如果串行，CPU 准备的那段时间 GPU 就闲着——而 decode 一步只有几毫秒，CPU 的准备时间完全不可忽略。

ZOS 的做法是：让 CPU 始终领先 GPU 一步。开启之后（NORMAL / COMPLETION 引擎默认开，engine_type.py 里的 default_zero_overhead_scheduling 就是干这个的），调度器在 GPU 还在跑第 t 步时，就把第 t+1 步的输入准备好。

难点在哪？在于第 t+1 步的输入依赖于第 t 步的结果（因为第 t 步会接受几个草稿 token，改变序列长度）。这个鸡生蛋的问题，靠一个叫 remap 的东西解开：csrc/engine/zos.cpp 里的 prepare_zos_inputs 拿“这一步”和“上一步”两份序列元数据做对齐，预先算出“上一步准备好的张量里，哪一行该被搬到哪一行”。剩下的交给 GPU 上的 scatter。

代价值得注意，而且它被诚实地记在账上——deepinfer/config.py：

def _get_runtime_max_model_len(max_model_len: int) -> int:
    """Return the exclusive upper bound of every runtime position."""
    nextn = envs.DEEPINFER_USE_NEXTN_DECODE
    scheduler_overhead = nextn + (nextn + 1 if envs.DEEPINFER_ZERO_OVERHEAD_SCHEDULING else 0)
    draft_overhead = nextn
    return max_model_len + scheduler_overhead + draft_overhead

“提前一步”是有预算的：每个按 position 索引的缓冲区（RoPE cache 之类）都得多留 nextn + (nextn+1) 行位置。省下来的时间，是用多留一点空间换的。这种把影子成本写进配置的地方，是我判断一个系统是否成熟的标准之一。

另一个细节来自 llm_engine.py：ZOS 开启时，“还有没有未完成的请求”这个判断要跑在 CPU 组上（PG.cpu_data_parallel_group()），因为它必须和 CPU 侧的流水线同步；不开 ZOS 时则跑在设备组上。流水线改变了“同步发生在哪里”——这几乎是所有异步系统的通病，也同样是所有异步系统的精髓。

4.4 调度器在线辨识自己的机器

如果说前面还只是“抄操作系统”，这一步就抄到了控制论。csrc/engine/mtp_varlen.hpp 里有两个估计器：

// 维护滑动窗口取中位数，用于根据运行时反馈动态调整 decode step 的 sampler 部分耗时估计。
class SamplerRuntimeEstimator { ... };

// forward step 的启动曲线没有考虑上下文长度的影响，动态拟合一个 slope_ms，给选中的 graph size 的
// 耗时估计增加 slope_ms * graph_size。输入为 (forward_ms - baseline_ms) / graph_size。
// 维护最近的 _window_size 个 forward step，先按 graph size 聚合取中位数，再对这些中位数取中位数。
// 只考虑有至少 _min_samples 个样本的 graph size，各 graph 等权。
class ForwardRuntimeEstimator { ... };

注意这里的统计手法，它比“取平均”讲究得多：

• 两层中位数：先在同一个 CUDA graph size 内部取中位数，再对这些中位数取中位数。目的是不被离群点带跑——在一个和别的租户共享的集群上，偶发的抖动是常态，“均值”会被抖动吃干。
• 只采用样本足够的 graph size：样本太少的估计不是估计，是噪声。
• slope_ms 是拟合出来的：前向耗时对上下文长度有斜率，而这个斜率在线拟合。

读出来的味道是：调度器不相信一个离线标定的常数，它相信刚刚发生的事。 这已经不是在“抄操作系统”，而是给调度器装了一个在线辨识器——它一边跑一边估计自己的机器有多快，然后据此决定这一轮能塞多少活。

这里我要说一句诚实的判断：课堂上所有关于“延迟 vs 吞吐”的图都是静态的（给定 batch size，算出的延迟曲线）。真实系统的调度器必须处理参数会漂移这件事。DeepInfer 用两个中位数窗口来解决它，朴素但有效。

五、第五幕：课堂没讲的四件事

前面四幕，大体是“课堂给了地图，DeepInfer 修了路”。这一幕换个方向：这些东西课堂上一点都没提，而我认为它们是这次开源真正的分量。

5.1 批次无关：把“可复现”当成一等公民

想象一个很普通的场景：你用 temperature=0（贪心）跑同一个 prompt 两次，一次单独跑，一次和另外 63 个请求拼在一个 batch 里跑。结果会不会一样？

在很多引擎里，答案是不一定。原因不在模型，而在实现：

• 归约的顺序不同（sum 一棵树还是两棵树，浮点结果不一样）；
• kernel 的形状不同（batch 大小变了，tile 划分变了，累加的中间量就变了）；
• CUDA graph 的尺寸不同（DEEPINFER_CUDAGRAPH_BATCH_SIZE_INTERVAL 默认 8，意味着批大小会被对齐到 8 的倍数，padding 进来的行会影响某些 kernel 的规约行为）；
• attention 的 softmax 归约在不同 head 数/不同 tile 下有不同顺序。

这些差异通常小到小数点后好几位，但一旦进入“采样”这一步，它就可能翻转一个 argmax，或者让一个几乎等概率的采样落到另一个 token 上。于是“同输入不同输出”。

CS336 整讲没有提这件事。而 DeepInfer 把它当成一等公民，有专门的开关：

# README
- **Batch-invariant inference:** keep model results identical across different batch compositions.
  Enable with `DEEPINFER_BATCH_INVARIANT=1` (NVIDIA only).
- **Deterministic sampling:** produce identical outputs for identical input tokens and sampling
  parameters. Set `DEEPINFER_BATCH_INVARIANT=1 DEEPINFER_ENABLE_MTP_VARLEN=0` and pass a fixed
  `seq_seed` to `SamplingParams` (NVIDIA only).

而实现路径在代码里是看得见的两处：

其一，attention 换成一个批次无关的实现。 deepinfer/model_executor/layers/attention_backend.py：

output, lse = flash_mla_with_kvcache(
    q=query,
    k_cache=key_cache,
    ...
    enable_batch_invariant=batch_invariant,
)

batch_invariant 这个开关直接透传进 FlashMLA 的 kernel——这意味着批次无关不是靠“算完再对齐”糊上去的，而是 kernel 层的一种可选语义。

其二，采样不再用“一个全局随机流”。 这是最漂亮的一步。deepinfer/model_executor/layers/tilelang_ops/random_sampling.py：

def random_sampling(probs, seq_seeds, sample_positions, is_nextn_predict=False):
    """Sample one column per row, weighted by probs with each row summing to at most 1.

    Deterministic in (seq_seeds, sample_positions, is_nextn_predict): independent
    of batch size and row order; zero-probability columns are never sampled.
    """

“Deterministic in (seq_seeds, sample_positions, is_nextn_predict)”——这份确定性只依赖三个东西：这一条序列的种子、这是它第几个位置、以及“这一行是不是草稿预测”。一个批次里有几条序列、它们在行里怎么排、谁和谁挨着——全都不影响。

对照传统的做法（一个 RNG 流，按顺序取数），这两者的差别是本质的：流式随机会因为“前面多抽了一个数”而整体错位；而 (seed, position) 的哈希不会。SamplingParams 里的 seq_seed（uint64）就是那条序列的种子：

seq_seed: int(uint64), the random seed of this sequence for sampling.

再补一块拼图：sum_smaller_probs.py 里的 get_sum_smaller_probs，算的是“按概率升序排列后的包含式 CDF”，并且等概率时按列索引打破平局：

"""Inclusive CDF after sorting tokens by ascending probability. Per row,
sum probabilities through the sampled token; equal-probability tokens are
ordered by column index.
"""

为什么要算这个？因为这是拒绝采样里那把“尺子”需要的一个量：不仅要知道“抽到了谁”，还要知道“这个选择代表的概率质量是多少”。而“是多少”必须与顺序无关，否则整个确定性链条就断在最后一步。 连“等概率怎么排序”都被规定了——这就是把可复现做到底的样子。

我为什么认为这是重大创新？因为它改变的不只是一个数字，而是系统的承诺：

• 评测和线上可以用同一套数字（否则你永远在怀疑“这个涨点是评测噪声还是线上噪声”）；
• prefix cache 命中与否、批大小怎么变，都不改变输出（否则“缓存”这件事会引入一类极难排查的随机性）；
• “确定性采样”这个能力本身（README 里单列的一条特性）就建立在这上面。

代价是真实的：你要 padding 到固定形状、要避开原子归约、要放弃一部分 kernel 优化空间。Deepinfer 愿意付这个代价——这是工程价值观，不是算法技巧。 这也解释了为什么这两个开关都标着 “NVIDIA only”：把它做到位需要 kernel 层的配合。

5.2 Engram：把“记忆”从神经网络里拆出来

这一条我打算从一个更朴素的问题讲起：注意力是不是唯一的“取信息”方式？

注意力是“按内容寻址”：我拿一个 query 去和所有 key 比相似度，然后按相似度加权取回 value。它很强大，但有个成本——为了取回一点信息，你必须把所有的 key 都搬过来比一遍（或者说，必须让它们参与计算）。

还有另一种取信息的方式：按索引查表。我不比相似度，我直接算一个下标，去表里取出那一行。

Engram 就是这个东西。deepinfer/model_executor/models/v4_modules/engram_metadata.py 里，构建过程写得很清楚：

class CompressedVocabulary:
    """Build the HF tokenizer's original-token -> compressed-token mapping on CPU.

    Decode one token at a time, retaining special tokens. Invalid UTF-8 pieces
    use their raw vocabulary spelling so distinct byte pieces cannot collapse.
    """
    ...
    normalizer = normalizers.Sequence([
        normalizers.NFKC(),
        normalizers.NFD(),
        normalizers.StripAccents(),
        normalizers.Lowercase(),
        normalizers.Replace(Regex(r"[ \t\r\n]+"), " "),
        ...
    ])

先做一次词表压缩：把原始 token 映射到一个归一化后的等价类（NFKC/NFD/去重音/小写/空白折叠）。直觉是——对于“查词组”这件事，“Apple”和“apple”应该是同一个东西。

然后按 n-gram 建表。config.py 里：

def _build_engram(hf, hf_dir):
    layer_ids = text.get("engram_layer_ids")
    if not layer_ids:
        return None
    n_ngram_tables = text["engram_max_ngram_size"] - 1  # 2..max_ngram_size gives max-1 tables in total

每个 n 一张表（2-gram、3-gram、4-gram……），每张表用哈希索引去取一行嵌入。运行时的职责写在 engram.py 的开头：

"""Engram runtime: hash n-grams, gather the sharded tables and gate the fused projection.

`CompressedTokenizer` maps raw token ids onto the compressed vocabulary. `Engram` and
`EngramRuntime` gather the sharded n-gram tables by hash across ranks, project them through
the fused KV projection and gate the result into the ...
"""

“gather the sharded tables by hash across ranks”——表本身是被分片存放在不同 rank 上的，查表要跨 rank 取。而它最终的作用是“gate 进残差流”：查到的向量不是直接替换什么，而是被一个门控加进主干。

为什么这是“课堂没讲”的？因为这不是神经网络里的一种层，而是外挂的一块记忆。按早期文章的账，Engram 占 196.6B 参数、约 183 GiB，是权重表里仅次于专家的一大块——但它只做查表，不做矩阵乘法，所以可以名正言顺地放在 CPU 内存里，用互连带宽换显存（DEEPINFER_ENGRAM_FORCE_GPU_STORAGE 在 Ascend 上把它强制拉回设备内，说明有些平台上这条路径不通）。

把它放回那把尺子里看，Engram 的意义就清楚了：它把一部分“参数搬运”换成了“索引查找”。 查表不需要把所有 key 都搬出来比一遍，只需要拿到那一行。对 decode 这种每一步都要搬一遍权重的阶段，这是一次量级上的省。

我在课堂上没有听过这类讨论。课堂里“记忆”就等于“KV cache”和“参数”，而 Engram 提出第三种：可以被寻址、可以被放在慢一点的地方、可以被整体共享的显式记忆。

5.3 EPLB：MoE 的负载均衡，训练一次，推理一直做

课堂第四讲讲了 MoE 的均衡：f·P 辅助损失、以及 DeepSeek V3 那套“无辅助损失”的路由偏置。那是训练时的均衡——让 router 学会均摊。

但真实服务里，路由器学会了均摊，不代表机器这盘棋是平局。序列是千变万化的，热门 token 会在某几步里扎堆挤向同一个专家，于是“理论均衡、实际倾斜”。DeepInfer 的做法是在推理时一直搬家：EPLB（Expert Parallel Load Balancer）。

deepinfer/model_executor/eplb/policy.py 的文档字符串给出了整个策略的骨架：

def rebalance_experts_global(tokens_per_expert, num_physical_experts, num_local_physical_experts):
    """
    全局负载均衡：给热点 logical expert 分配冗余副本，再把所有副本均衡到各 GPU。

    Parameters:
        tokens_per_expert: [num_steps, num_moe_layers, num_logical_experts]

    return:
        physical_to_logical_map: [num_moe_layers, num_physical_experts]
        logical_to_physical_map: [num_moe_layers, num_logical_experts, num_redundant_experts + 1]
        logical_count: [num_moe_layers, num_logical_experts]，各 logical expert 的副本数
    """

两件事：给热门的专家复制冗余副本（logical 是模型定义的那个专家，physical 是显存里真实躺着的那些副本），再把所有副本摆平到各 GPU 上。

分配副本的那几行代码，是全仓库我认为最漂亮的一段：

# 贪心分配冗余副本：每轮选一个 expert 复制。
# score[..., j] 是 expert j 当前的每副本负载，score_added 是给 j 加一个副本后的每副本负载；
# 把 argmax 位置替换成其 score_added 后，values[..., j] 即"副本给 j"时的最大每副本负载
# （j 非 argmax 时最大值不变，j 为 argmax 时是替换后重算的 max），按 step 求和取 argmin 得最优复制对象。
tokens_per_expert_all_diff = tokens_per_expert + arange_logical * 1e-4
for i in range(num_redundant_experts):
    score = tokens_per_expert_all_diff / logical_count
    score_added = tokens_per_expert / (logical_count + 1)
    values, indices = score.max(-1, keepdim=True)

我读了三遍才读出味道：

• score 是当前每个副本的平均负载，score_added 是“再复制一份之后”的平均负载；
• 目标是最小化最热的那一个副本的负载（min-max）——因为全局耗时由最慢的那个 rank 决定，而不是由平均值决定；
• 技巧在于“把 argmax 的位置替换成它加了副本之后的值再重算 max”——这样 values[j] 就恰好是“把这一份副本给 j 时，全局最大负载会变成多少”，然后取 argmin 就是最优的复制对象。一次减法就绕过了“给每个候选各试一遍”的循环。

还有一个我必须点出来的细节：tokens_per_expert + arange_logical * 1e-4。加这么一点点噪声，是为了让“分数互不相同”，否则 argmax 在并列时会在 GPU 上产生不确定的选择——为确定性而加扰动，是个很成熟的习惯。

至于搬家的执行，eplb/shuffle.py 用的是 torch.distributed._symmetric_memory（对称内存）——也就是说，权重的整体搬家走的是 NVLink 级别的通道，不是普通通信。这是一个很实在的选择：搬家本身不能比不搬家还贵。

那为什么说这是“课堂没讲”？因为课堂讲的均衡是一次性的、训练期的，而 EPLB 是持续的、推理期的，并且它承认了一个残酷的事实：路由是学出来的，但流量是长出来的。 两者永远对不齐，所以你只能一直搬。

（另外，搬家也要有节制。environ.py 里的 DEEPINFER_MOE_LB_STAT_INTERVAL（默认 128）、DEEPINFER_EXPERT_COUNT_WINDOW_SIZE（默认 100）、DEEPINFER_EXPERT_REARRANGEMENT_STEP_INTERVAL 这些开关，本质上都是“多久统计一次、多久搬一次”的节流阀——搬家的通信开销也是要还的。）

5.4 结构化输出：一个把“语法”变成 mask 的小世界

最后一件课堂没讲的事：约束生成。让模型的输出必须是一个正则表达式、一个 JSON schema、或者一个 EBNF 语法。README 把它列为一条核心特性，例子在 examples/_structured_output.py：

cases = [
    ("...", {"schema_type": "regex", "schema": r"[A-Z]{2}-[0-9]{3}"}),
    ("...", {"schema_type": "json",  "schema": json.dumps(json_schema)}),
    ("...", {"schema_type": "ebnf",  "schema": 'root ::= "answer=" ("yes" | "no") ";level=" [1-3]'}),
]

实现放在 deepinfer/structure_output/。三个文件，三种角色：

• compiler.py 用 xgrammar 把语法编译成 token 级的 bitmask，而且是带缓存的编译器，缓存大小默认 10 GiB，还专门警告了缓存无上限时的 OOM 风险（引用了一个 CVE 编号）——说明这块踩过坑；
• matcher.py 负责跟踪当前走到语法的哪一步，并产出“下一步哪些 token 合法”；
• mask.py 把 mask 应用回 logits。

mask.py 里这句注释揭示了它与投机解码的正面相撞：

def apply_vocab_mask_by_step_inplace(logits, vocab_mask, indices, ftp=None, step=0) -> None:

by_step——按步应用。因为 ZOS + MTP 一步会在同一批里同时处理“验证草稿”和“新鲜采样”，语法状态也必须能按步推进，不能整批一起刷。

更有意思的是 matcher.py 里这个常量：

MAX_ROLLBACK_TOKENS = envs.DEEPINFER_USE_NEXTN_DECODE + 2

语法匹配器必须能回退 nextn + 2 步。

为什么？因为猜出来的草稿可能被拒绝。如果草稿已经推进了语法状态（“我读了 5 个 token，语法走到第 5 步”），而目标模型拒绝了其中 2 个，那语法状态就必须退回去。所以一个可靠的 matcher 必须自带回滚能力，而且回滚深度必须覆盖“最坏情况下会被撤销几步”。

于是 matcher.py 里就出现了一个看起来完全不属于“推理引擎”的算法——KMP 的失败函数：

def _kmp_failure(seq: list[int]) -> list[int]:
    failure = [0] * len(seq)
    k = 0
    for i in range(1, len(seq)):
        while k and seq[i] != seq[k]:
            k = failure[k - 1]
        ...

（它用来在历史里找一个“触发序列”的下次可能匹配位置，以便回滚时不至于从头重扫。）

我特意把这一节放最后，因为它最能说明一个问题：特性之间会互相咬。 单独看，投机解码是“猜+验”；单独看，结构化输出是“mask”；但放在一起，“一步猜 6 个 token”这件事就逼着语法引擎长出回滚能力，而回滚能力又逼出了一个字符串匹配算法。真实系统的难度不在单个算法，在算法之间的接口上。 这也是为什么“读懂一个开源推理框架”比“读懂一篇论文”要难得多——论文里每个模块都是自洽的，仓库里每个模块都是互相欠债的。

5.5 顺带一提：一份语义，两种硬件方言

还有一件课堂不会讲、但工程上极其昂贵的事。deepinfer/model_executor/layers/tilelang_ops/ 下面，cuda/ 和 asc/ 两个目录几乎是一一对应的——同一套算子，要在 NVIDIA 的 SM100 和华为 Ascend 950 上各实现一遍，用 TileLang 写。

而且硬件的差别会反过来影响算法。schedule_mtp_varlen.py 里有一行断言：

if PLATFORM.is_npu:
    assert num_seqs * nextn * 4 <= 32 * 1024, f"rates_ub exceeds 32KB UB: num_seqs*nextn={num_seqs * nextn}"

Ascend 的 UB（unified buffer）只有 32 KB，所以“一次能有几条序列 × 每条猜几个”这件事，在 NPU 上被一个 reuse buffer 的容量直接卡死。同样一个置信度调度算法，在两种硬件上会有不同的可行域。

这一条我不打算展开（它本身够写一篇），但要留一个记号：当一个系统的算法开始被硬件的 buffer 大小约束时，它就真的到了“系统”这一层。

六、尾声：伽莫夫的梦，和费曼的那块黑板

让我们回到开头那个搬运工。

这一趟走下来，我发现 CS336 第十讲其实只给了三样东西：

1. 一把尺子：算术强度。用它量一量，就知道该省算力还是省搬运；
2. 一张图：prefill 计算受限、decode 访存受限。所有技巧都是在攻击后半张图；
3. 一句箴言：“想想你的操作系统课。”

DeepInfer 把这三样东西全都实现了——而且实现得比课堂上讲的更严苛。但它真正有意思的地方，是在三个方向上走出了课堂，而这三个方向恰好对应三种不同的学科：

• 把不确定性当敌人——这是实验物理的价值观。 批次无关推理、确定性采样、(seq_seed, sample_position, is_nextn) 的哈希采样、等概率按列索引打破平局……这些不是“技巧”，而是对“可复现”这件事做出的承诺。费曼说 “What I cannot create, I do not understand”；对一个推理系统，我愿意把它改成：“无法逐位复现的东西，我没有真正理解它。” DEEPINFER_BATCH_INVARIANT 这个开关，就是对这句话的一次实践。
• 把记忆从计算里拆出来——这是架构级的省。 Engram 不是更快的矩阵乘法，而是“根本不做矩阵乘法”。YOCO 也不是更快的 KV cache，而是“让一半的层不写 KV”。CS336 结尾那句“最大的机会在模型与架构层，而不是系统层”，在这个仓库里有了具体的样子。
• 把验证变成调度、把调度变成在线辨识——这是控制论。 DSpark 用置信度决定“这一步验证谁、验证几个”；ZOS 让 CPU 领先 GPU 一步，并用 remap 解开依赖；两个中位数窗口的估计器在线拟合自己机器的耗时。课堂上的延迟/吞吐曲线是静态的；真实系统的调度器必须相信刚刚发生的事。

最后，说一点这篇文章想避免的习惯。上面这些名字——MLA、YOCO、DSpark、EPLB——听起来各不相同，但如果你只记住那把尺子，你会发现它们全都在回答同一个问题：每搬一个字节，我能换到多少次浮点运算？

MLA 是“少搬几个字节”，滑窗是“少搬一段”，YOCO 是“少搬一半的层”，DSpark 是“一次搬运服务六个 token”，ZOS 是“搬的时候别闲着”，Engram 是“把搬运换成索引”，批次无关是“让每一次搬运的答案都一样”。它们不是七种奇技淫巧，而是同一句话的七种说法。

伽莫夫那本书里，汤普金斯先生最后醒过来，发现世界还是那个世界，只是他现在看得见了。读完 DeepInfer 之后大概也是这样：推理还是那个推理，只是你再看它的时候，看见的已经是一台不停在搬东西的机器——而所有的巧思，都长在那把尺子上。

关于本文的边界

这次开源提供的是离线推理，只支持 DeepSeek-V4.1-Flash 一个模型，并且官方明确列出了三项“未包含”：KV cache 系统、VE/Prefill/Decode 分离、以及 prefill 与 decode 的跨 DP 副本负载均衡。所以本文读到的调度器、批处理、投机解码是一台完整的引擎，但不是一套完整的服务。此外，很多能力是通过环境变量开关的（DEEPINFER_*，见 deepinfer/environ.py），这本身就是一份很好的“设计意图清单”——一个系统的取舍，往往就写在它的开关列表里。