DeepInfer 源码完全拆解（二）：异步调度与零开销调度 ZOS——把 CPU 后处理移出关键路径

为什么「延后一拍」能换来重叠

答案藏在第 2 步的同步点上。_post_process_step 会 synchronize() 上一步的 sample_end_evt，这个等待在 ZOS 下发生的时候，GPU 已经因为第 1 步的提交而在忙碌。CPU 等待上一步采样完成的同时，GPU 正在算这一步的前向——等待的这段时间不再是纯空泡。

换个角度说：非 ZOS 的时间线是「GPU forward → GPU sample → CPU 等采样完成 → CPU 后处理 → GPU 下一轮 forward（必须等 CPU 把新 metadata 备好）」，串成一条链。ZOS 的时间线是「GPU forward(t) → [CPU 后处理(t-1) 与 GPU forward(t) 并行] → GPU sample(t) → [CPU 后处理(t) 与 GPU forward(t+1) 并行]」。后处理从关键路径上的一个串行段，变成了挂在 GPU 工作上的一个并行影子。

代价也很清楚：

• 延迟增加一步。 第 t 步的输出在第 t+1 步才被回传出去。对离线批处理接口 LLM.generate 来说这无所谓，因为它是循环 step() 直到跑完；但对在线流式服务来说，每个 token 晚一拍返回，这是一个实打实的延迟成本。README 里写得很直白：DeepInfer 是离线推理的开放版本，LLM 类是离线批处理接口。这也解释了为什么 ZOS 的代价在这个框架里可以被接受。
• 在途状态变多。 prev_step_state 持有上一步的序列元数据对象和输出指针，它们必须活到下一步的后处理结束才能释放。这就是 DeferredRelease 存在的理由。
• 正确性变窄。 下一步的输入元数据是拿旧账本构造的，必须有修正路径兜住。

下面这张图把非 ZOS 和 ZOS 两条时间线并排放在一起。

（图源：本站依据 deepinfer/worker/worker.py 的 execute_model_step 编排绘制。方框宽度不代表真实耗时比例，只表达先后与重叠关系。）

空步也是步：占位序列与数据并行的 lockstep

step() 的末尾有一句容易读漏的判断——把 is_placeholder=scheduler_outputs.is_empty() 传给 worker。也就是说，当调度器这一步没调度出任何序列时，引擎照样走一遍执行流程：worker 构造一条占位序列（model_runner._placeholder_seq_metadata(is_prompt)），照常前向、照常采样，只是产出被丢弃。

_placeholder_seq_metadata 的构造很朴素：request_id=PLACEHOLDER_REQUEST_ID、seq_id=0、一个 token 的 SequenceData、seqlen_q=1、seqlen_k=1、block_table=[0]。它的作用不是产出，而是占位。

为什么不干脆跳过这一步？因为 DeepInfer 是数据并行（Attention DP）部署：多个 DP rank 各自持有完整的注意力权重与自己的调度器，但 MoE 层需要跨 rank 做 All-to-All（DeepEP 的 dispatch/combine），EPLB 的统计还依赖周期性的 EP all_gather。这些集合通信要求所有 rank 同步进入——如果某个 rank 因为队列空就跳过一步，其余 rank 会在集合通信上等到超时。

占位步保证所有 rank 的步数一致，把「队列空」变成一个代价可控的伪步。同一个问题在本站已有的文章里出现过：BalanceRoute 完全拆解：大规模 LLM 数据并行服务如何在线负载均衡 是在外层做请求路由均衡，DeepInfer 是在引擎内部用空步兜底。两者的约束来源相同：DP 副本之间必须对齐。

值得一提的是，占位步也会走 ZOS 的分支。_correct_seqlens_qk 的开头就有一句：

1
prev = self.prev_step_state
2
if prev is None or prev.is_placeholder:
3
    return None

如果上一步是占位步，就没有真实产出可供修正，直接跳过。同样的，占位序列自己也不会被修正——prepare_zos_inputs 里对 seq->is_placeholder 的行直接 continue，填 0 了事。

滞后的是什么：K 长度账

先明确一个朴素的事实：decode 时每条序列的注意力输入是两个长度——seqlen_q（本步这条序列要算几个 query）和 seqlen_k（这条序列在 KV 缓存里一共可见多少个 key）。

• seqlen_q 在普通解码里恒为 1（每步生成一个 token）；在投机解码（MTP）下等于 1 + 本步验证的草稿数，而草稿数是可以逐请求不同的（DEEPINFER_USE_NEXTN_DECODE 默认 5，即最多验证 5 个草稿）。
• seqlen_k 是「已缓存长度 + 本步新增的 q 数」，而「已缓存长度」在 ZOS 下的 CPU 账本里是上一步开始前的旧值。

为什么是旧值？因为更新这个账本的代码在 scheduler.process_outputs 里，而它跑在 _post_process_step 里——被推迟到了下一步。第 t 步构造输入时，调度器手上只有第 t-1 步开始前的缓存长度。

于是第 t 步的 K 长度必须由两部分拼出来：

其中 $\text{cache\_len}^{(t-1)}$ 是第 t-1 步开始前已缓存的 token 数（这个 CPU 是知道的），$q^{(t)}$ 是本步调度的 query 数（CPU 也知道），而 $\text{accepted}^{(t-1)}$ 是第 t-1 步投机解码实际被接受、因而真正写进 KV 的 token 数——这个值只有 GPU 知道，因为它取决于第 t-1 步采样验证的结果。

这就是整个修正链路要搬回来的那一个数。用代码里的写法（compute_seqlens_k_zos 的文档字符串）：

1
seqlens_k[i] += seqlens_q[i]
2
seqlens_k[i] -= nextn - prev_num_accepted_tokens[j]   (if remap[i] == j != -1)

展开一下就是：

代入 $\text{nextn} = q^{(t)}-1$（本步 q 数减 1 就是草稿数上限），化简后恰好得到上面那个公式：

直觉是这样的：旧的 base 是「上一步开始前已缓存的 token 数」。上一步实际往 KV 里写了「1 个 golden token + accepted 个草稿 token」，其中那 1 个就是本步的输入。所以本步的 K 长度 = base + 本步 q 数 + 上一步接受数。

remap：把本步批次对回上一步批次

修正的前提是把「本步的第 i 个序列」和「上一步的第 j 个序列」对上。这件事不平凡，因为两次调度的批次组成可以不一样：上一步被抢占的序列不在本步的批次里，本步可能有新入队的序列。

DeepInfer 在 C++ 侧用一个 remap 数组表达这个对应关系。prepare_zos_inputs 的逻辑分两支：

• 快路径：如果两次批次的长度相同并且逐位置 seq_id 一致，直接 remap_indices[i] = i，has_overlap = true。
• 慢路径：否则构建一个 unordered_map<seq_id → prev_index>，再逐位置查表；查不到就填 -1。

-1 表示这条序列是新入队的（代码里叫 fresh），它在上一步没有产物可用。remap 数组在 zos_status.py 里的注释写得很精确：

1
remap[i] = j if curr_batch[i] matches prev_batch[j], else -1

三个数组一起从 C++ 侧产出，都是 pinned 内存：

1
auto remap    = torch::empty({n_seq_with_padding}, pinned(torch::kInt32));
2
auto seqlens_q = torch::empty({n_seq_with_padding}, pinned(torch::kInt32));
3
auto seqlens_k = torch::empty({n_seq_with_padding}, pinned(torch::kInt32));

其中 seqlens_q 和 seqlens_k 是待修正的初值，prepare_zos_inputs 对每个位置填的是：

1
seqlens_q_ptr[i] = seq->seqlen_q;
2
seqlens_k_ptr[i] = seq->seqlen_k - (nextn > 0 ? seq->seqlen_q : 0);

也就是把 CPU 侧那个滞后的 seqlen_k 先减去本步 q 数，归一化成「已缓存长度」。这样 GPU 侧只要再加回本步 q 数和上一步接受数即可，语义清晰。

padding 行（为了对齐 CUDA Graph 的桶尺寸而多出来的空转序列）被填成 remap = -1、seqlens_q = 1、seqlens_k = 1，不会参与任何有效计算。关于 CUDA Graph 的桶尺寸与 padding，总览篇已经拆过，可回看 DeepInfer 源码完全拆解（一）：整体架构与模块全景 里「CUDA Graph 的尺寸桶与预热-缓存闭环」一节。

三个 TileLang 内核

修正不是在 Python 里用 tensor 算子拼的，而是三个专门写的 TileLang kernel。这里逐个看。

第一个是 compute_seqlens_q_by_confidence_rates，只在投机解码变长（DEEPINFER_ENABLE_MTP_VARLEN）时启用。它解决的是「本步每个请求应该验证几个草稿」的问题。总览篇提过 DeepInfer 用置信度头给每个草稿位打分，这里就是它的落地：每个草稿位有一个累乘后的接受概率，全局按概率排名，排名靠前的位才被分配进本步的 token 预算。

CUDA 版的实现分两个 phase：

1
# Phase 1: compute cumprod, reading confidence rates via remap (0.0 for fresh seqs)
2
with T.Kernel(num_seqs, threads=128) as pid:
3
    acc = 1.0
4
    prev_idx = remap[pid]
5
    if prev_idx == -1:
6
        acc = 0.0
7
    for j in T.serial(nextn):
8
        if prev_idx != -1:
9
            acc = acc * prev_confidence_rates[prev_idx, j]
10
        cum_rates[pid, j] = acc
11

12
# Phase 2: each block reads precomputed rates, does parallel rank counting
13
with T.Kernel(num_seqs, threads=128) as pid:
14
    my_rates = ...  # 读取自己的 nextn 个累乘概率
15
    rank_reducer = T.alloc_reducer((nextn,), T.int32, replication="all")
16
    for other_idx in T.Parallel(num_seqs * nextn):
17
        other_seq = other_idx // nextn
18
        other_pos = other_idx % nextn
19
        other_rate = cum_rates[other_seq, other_pos]
20
        for j in T.serial(nextn):
21
            my_flat_idx = pid * nextn + j
22
            if other_rate > my_rates[j]:            beats = 1
23
            if other_rate == my_rates[j] and other_idx < my_flat_idx: beats = 1
24
            rank_reducer[j] += beats
25
    T.finalize_reducer(rank_reducer)
26
    # 排名在前 (total_q - num_seqs) 名之内的草稿位被接受
27
    extra_budget = total_q - num_seqs
28
    num_tokens = sum(1 for j in range(nextn) if rank_reducer[j] < extra_budget)
29
    seqlens_q[pid] = 1 + num_tokens

Phase 1 的语义是：对每个序列，把它上一步 nextn 个草稿位的置信度做累乘。累乘的含义在 C++ 侧 process_outputs 的注释里有说明——「modeling 返回的 confidence rate 是每个 token 独立的，使用前要做一个累乘」。第 j 位的累乘概率表示「前 j 个草稿全被接受」的概率，这是一个逐位递减的序列。fresh 序列（remap == -1）的 acc 被强制设为 0.0，等于这些位自动排在最后，不会被选中。

Phase 2 是一个并行排名：每个 query 位置把自己的累乘概率和全局所有位置比一遍，数出「有多少个位置比我高」，这个计数就是它的排名；排名小于本步的额外预算 total_q - num_seqs（注意 total_q 是本步已经定好的 token 预算，num_seqs 是每个请求保底的 1 个 golden token）的位置被接受。并列时用扁平下标做 tie-break（other_idx < my_flat_idx），保证结果确定。

这个内核的 ASC 版本（tilelang_ops/asc/schedule_mtp_varlen.py）有 5131 字节，比 CUDA 版的 3216 字节更长——昇腾版本需要显式处理更多同步与布局细节。总览篇说过 tilelang_ops/cuda/ 与 tilelang_ops/asc/ 是同一组逻辑算子的两份实现，这里是一个典型样本。

第二个是 compute_seqlens_k_zos，就是前面推导过的那个修正。CUDA 版只有 37 行：

1
@tilelang.jit(pass_configs={...})
2
def compute_seqlens_k_zos_tl(nextn: int) -> Any:
3
    n_seq = T.dynamic("n_seq")
4
    n_prev_seq = T.dynamic("n_prev_seq")
5
    BLOCK_SIZE = 128
6

7
    @T.prim_func
8
    def _compute_seqlens_k_zos(
9
        seqlens_k: T.Tensor((n_seq,), "int32"),          # in/out
10
        seqlens_q: T.Tensor((n_seq,), "int32"),
11
        remap: T.Tensor((n_seq,), "int32"),              # -1 or prev_batch_index
12
        prev_num_accepted_tokens: T.Tensor((n_prev_seq,), "int32"),
13
    ):
14
        with T.Kernel(T.ceildiv(n_seq, BLOCK_SIZE), threads=BLOCK_SIZE) as bx, \
15
             T.Parallel(BLOCK_SIZE) as tx:
16
            sk = T.alloc_var("int32")
17
            seq_idx = bx * BLOCK_SIZE + tx
18
            T.pdl_sync()
19
            if seq_idx < n_seq:
20
                sk = seqlens_k[seq_idx] + seqlens_q[seq_idx]
21
                prev_idx = remap[seq_idx]
22
                if prev_idx != -1:
23
                    sk = sk - nextn + prev_num_accepted_tokens[prev_idx]
24
                seqlens_k[seq_idx] = sk
25
    return _compute_seqlens_k_zos

两个细节值得注意。

其一是 T.pdl_sync()。PDL 是 CUDA 的 Programmatic Dependent Launch——它允许后续 kernel 在前一个 kernel 还没完全结束时就开始执行前导部分（preamble），依赖数据时才等待。这里三个修正内核在同一条流上连续发射，PDL 让它们之间不需要完整的 kernel 边界等待。对一个「每步都要跑、每次只算几百个 int32」的微内核来说，省下的就是纯 launch 开销。

其二是它同时被当作 seqlens_k 的原地更新器（in-place，out_idx=None 且写回输入张量）。源码里配了一个 PyTorch 参考实现 compute_seqlens_k_zos_ref 做逐元素比对测试，测试参数覆盖 has_remap ∈ {True, False} 与 nextn ∈ {3, 7}，n_seq 从 64 到 1024 分档跑 benchmark。这是 DeepInfer 对微内核的标准做法：先写参考实现，再写 kernel，测试保证两者逐元素相等。

第三个是 gather_input_ids_zos，负责把本步前向要喂的 token id 拼出来。它的逻辑是：对每个批次位置，如果 remap >= 0（是上一步的延续），本步的输入 token 就是上一步在 GPU 上产出的东西——第 0 位取 prev_sampled_token_ids[pidx]，第 j>0 位取 prev_predict_token_ids[pidx, j-1]；如果 remap == -1（fresh），则从 CPU 侧构造好的 fresh_tokens 里取。

1
for j in T.Parallel(max_seqlen_q):
2
    if j < sq:
3
        if pidx >= 0:
4
            if j == 0:
5
                out_input_ids[q_base + j] = prev_sampled_token_ids[pidx]
6
            else:
7
                out_input_ids[q_base + j] = prev_predict_token_ids[pidx, j - 1]
8
        else:
9
            out_input_ids[q_base + j] = fresh_tokens[seq_idx, j]

这就是 ZOS 下「input_ids 不经过 CPU」的机制。在非 ZOS 里，上一步采样出的 token 会被后处理写回 Sequence，下一步构造输入时再从 Python 侧的 token 列表里读出来、H2D 拷上去。ZOS 把这条 CPU 中转掐掉了：token 从头到尾留在 GPU 上，由一个 gather kernel 直接在设备侧把上一步的产物拼成本步的输入。

这带来一个连锁收益——prepare_tensor_input_zos.py 里的注释明确写了两个分支的差别：

• remap is None（无重叠）：CPU 侧把完整打包好的 input_ids 张量一次性 H2D 拷贝。
• 有重叠：CPU 只准备 fresh_input_ids（形状 [n_seq, max_seqlen_q]，只填 fresh 行，其余补 0），匹配行由 GPU 的 gather kernel 填。

也就是说，有重叠的场景下 H2D 的数据量变小了：只传新序列那几行，而不是整个 batch。这在长上下文、批次大部分延续的场景下是明显的节省。

端到端的修正链路

把上面几块拼起来，第 t 步的修正链路是这样的：

（图源：本站依据 zos.cpp 的 prepare_zos_inputs、zos.cpp 的 update_metadata_for_sampling_zos、worker.py 的 _correct_seqlens_qk 与 tilelang_ops/cuda/* 三个内核绘制。）

对应到 _correct_seqlens_qk 的源码，顺序是：

1
remap_cpu, seqlens_q_cpu, seqlens_k_cpu, has_overlap = prepare_zos_inputs(
2
    seq_metadata_list, prev.seq_metadata_list, n_seq_with_padding, self.nextn
3
)
4
if not has_overlap:
5
    return None
6
remap    = remap_cpu.cuda(non_blocking=True)
7
seqlens_q = seqlens_q_cpu.cuda(non_blocking=True)
8
seqlens_k = seqlens_k_cpu.cuda(non_blocking=True)
9

10
if envs.DEEPINFER_ENABLE_MTP_VARLEN:
11
    compute_seqlens_q_by_confidence_rates(
12
        total_q, remap[:n_seq], prev.sampler_output.confidence_rates_gpu, seqlens_q[:n_seq])
13
if self.nextn > 0:
14
    compute_seqlens_k_zos(
15
        seqlens_k[:n_seq], seqlens_q[:n_seq], remap[:n_seq],
16
        prev.sampler_output.num_accepted_tokens_gpu, nextn=self.nextn)
17

18
d2h_evt = None
19
if envs.DEEPINFER_ENABLE_MTP_VARLEN:
20
    # H2D, correction, and D2H are ordered on the current stream;
21
    # only varlen rewrites this pinned source.
22
    seqlens_q_cpu[:n_seq].copy_(seqlens_q[:n_seq], non_blocking=True)
23
    d2h_evt = torch.cuda.Event()
24
    d2h_evt.record()

最后那三行是链路的关键收尾，源码注释也点明了：H2D、修正、D2H 三步在同一条流上被排序，所以 d2h_evt 记录时，seqlens_q_cpu 里已经是修正后的值。而且只有 varlen 才需要回读——因为只有变长投机解码下，本步每个请求的 q 长度和 CPU 预想的（统一 pad 到 nextn+1）不一致；非 varlen 时 CPU 侧的初值就是对的，不需要 D2H。

注意 remap_cpu 与 seqlens_q_cpu 是同一块 pinned 内存：remap_cpu 是 H2D 的源（要被拷到设备上去），seqlens_q_cpu 是 D2H 的目的（修正后的值要写回来）。MetadataZOS 的字段注释把这个双重身份写出来了：

1
remap_cpu: C++-generated pinned remap: H2D source and CPU-side
2
           fresh/previous dispatch input for this step.

改一处要考虑另一处，这是这类「pinned 内存复用」设计的固有代价。总览篇把它列进了「工程上的坑」一节，这里补上了具体机制。

回读之后：修什么、为什么必须修

D2H 回来的 seqlens_q_cpu 交给 update_metadata_for_sampling_zos（C++ 侧，也在 zos.cpp）：

1
const auto *seqlens_q_ptr = seqlens_q.const_data_ptr<int32_t>();
2
const auto *remap_ptr = remap.const_data_ptr<int32_t>();
3
scheduler_outputs.seqlens_q.resize(n_seq);
4
for (size_t i = 0; i < n_seq; ++i) {
5
    auto *metadata = seq_metadata_list[i];
6
    auto &seq = *scheduler_outputs.scheduled_seqs[i];
7
    metadata->seqlen_q = seqlens_q_ptr[i];
8
    // EOS 后的额外 forward 可能未将 golden token 追加到 Sequence，
9
    // 采样侧以已缓存长度为准。
10
    metadata->seqlen_k = metadata->seqlen_q + seq.cache.num_tokens;
11
    if (remap_ptr[i] >= 0) {
12
        metadata->refresh_for_sampling(seq, nextn);
13
    }
14
    scheduler_outputs.seqlens_q[i] = metadata->seqlen_q;
15
}

三件事：把修正后的 seqlen_q 写进 metadata；按 seqlen_q + seq.cache.num_tokens 重算 seqlen_k；对 remap >= 0（有上一步产物）的序列调 refresh_for_sampling——把上一步的真实产出（最后采样出的 token、接受的草稿 token 列表）补进采样侧的元数据。

为什么必须修？ 因为采样器要按 seqlen_q 给每个请求分配 logits 的行。如果 CPU 侧以为本步每条序列的 q 长度是 1 + nextn（统一 pad），而 GPU 实际按置信度分配了不同长度（有的请求只验证 2 个，有的验证 5 个），那么采样器读 logits 时的行切片就会错位，取到的根本不是这条序列该采样的那个位置。这不是精度问题，是语义错误。

refresh_for_sampling 同理：采样时需要知道「上一步这条序列到底产出了什么 token」，才能把它作为已生成序列的一部分进行处理（比如判断是否命中 stop string）。ZOS 把这部分产出的处理推迟了，所以在采样侧必须显式把它补回来。

正确性的边界：哪些环节会打破这个链路

拆到这里，ZOS 的修正链路是完整的。但它的正确性依赖几条前提，源码里也有明确的警示：

第一，任何绕过修正路径、直接相信调度器算出的 seqlens_k 的下游代码都会错。 在 ZOS 下，SequenceMetadata 上的 seqlen_k 在不同阶段含义不同：调度器刚产出时是「已缓存长度」的归一化值，修正后是「本步真实 K 长度」。Worker 的注释也点明了 MetadataZOS.remap_cpu 的双重身份。下游代码必须清楚自己拿到的 metadata 处在哪个阶段。

第二，DEEPINFER_ENABLE_MTP_VARLEN 与确定性采样互斥。 变长验证的草稿长度由置信度排名决定，而置信度来自上一步的运行时估计，会引入非确定性。要确定性就必须 DEEPINFER_ENABLE_MTP_VARLEN=0，这一点总览篇已经列在坑清单里，这里从机制上确认了原因。

第三，无重叠时整条链路被短路。 prepare_zos_inputs 在 has_overlap == false 时返回 {nullopt, nullopt, nullopt, false}，_correct_seqlens_qk 随即 return None，后续的所有 kernel 都不发射，input_ids 走 CPU 侧的完整 H2D 路径。这是一个重要的性能事实：新请求大规模涌入、批次完全换血的场景下，ZOS 的修正开销自动归零，只剩下「延后一拍」的编排本身。

第四，被抢占的序列需要特殊对待。 process_outputs（C++ 侧）里有一句：

1
// ZOS 异步：本步在途期间被 _preempt 抢占重算的 seq，MTP 计数已整组复位，
2
// 本步记账一并作废（判据与 process_outputs 的跳过一致；否则 process_cnt
3
// 抢跑 schedule_cnt，MTPVarlen::schedule 的配对断言会崩）。
4
if (!scheduler_outputs.prompt_run && seq.schedule.status == SequenceStatus::WAITING) {
5
    continue;
6
}

被抢占的序列在下一步会整组重算，它的 MTP 计数（schedule_cnt / process_cnt）被复位，所以这一步的记账必须整条跳过。否则两个计数器会失配，触发 MTPVarlen::schedule 里的那个断言：

1
TORCH_CHECK(mtp.process_cnt <= mtp.schedule_cnt && mtp.schedule_cnt <= mtp.process_cnt + 1,
2
    "schedule_cnt: ", mtp.schedule_cnt, ", process_cnt: ", mtp.process_cnt, "，不应该。");

这个「不超过 1」的强约束，正是 ZOS 在 MTP 记账上留下的指纹：正常路径下调度和处理的配对最多差一轮，而这个差正好对应「在途一步」。

两套估计器

第一套是 SamplerRuntimeEstimator，估计采样阶段的耗时。它是一个 9 槽的滑动窗口，至少攒够 5 个样本才给出估计，给出的值是窗口内的中位数：

1
class SamplerRuntimeEstimator {
2
public:
3
    void update(double value_ms) {
4
        _samples[_next] = value_ms;
5
        _next = (_next + 1) % _window_size;
6
        _size = std::min(_size + 1, _window_size);
7
        auto sorted = _samples;
8
        std::sort(sorted.begin(), sorted.begin() + _size);
9
        const size_t middle = _size / 2;
10
        _median_ms = _size % 2 == 0
11
            ? (sorted[middle - 1] + sorted[middle]) * 0.5
12
            : sorted[middle];
13
    }
14
    std::optional<double> get() const {
15
        return _size >= _min_samples ? std::optional<double>(_median_ms) : std::nullopt;
16
    }
17
private:
18
    static constexpr size_t _window_size = 9;
19
    static constexpr size_t _min_samples = 5;
20
    ...
21
};

采样阶段被单独估计的原因，源码注释给了：「Sampler 没有启动曲线」——也就是采样耗时不吃 GPU 前向那条随 batch size 上升的启动曲线，它主要取决于批内序列数。所以 DeepInfer 按 num_seqs 分桶维护多个 SamplerRuntimeEstimator，并有全局回退：

1
// 对于 sampler，如果相同 num_seqs 已经收集足够样本，就用相同 num_seqs 过往几步的时间估计
2
// 否则回退到用全局过往耗时估计
3
auto sampler = _sampler_time_by_num_seqs.find(num_seqs);
4
auto estimate = sampler != _sampler_time_by_num_seqs.end()
5
    ? sampler->second.get() : std::nullopt;
6
if (!estimate) { estimate = _sampler_time_fallback.get(); }

第二套是 ForwardRuntimeEstimator，估计前向耗时随 batch 变化的斜率。它的输入是归一化残差：

其中 baseline_ms 是从「解码性能曲线」里查出来的：_decode_perf_curve 是一串 (batch_size, forward_ms) 点，用 std::lower_bound 找到不小于当前 graph size 的那个点取它的耗时作为基线。这个曲线本身也是实测产物，来自 model_runner.profile_decode_perf()——对每个 CUDA Graph 桶跑 1 次 warmup、5 次带事件计时的测量，去掉最大最小后取截尾均值，再跨 EP group 做一次 all_reduce 取平均。它还可以被缓存到磁盘跨进程复用（DecodePerfCache，哈希键包含 time_unit/model/batch_sizes/parallel_config）。

ForwardRuntimeEstimator 的结构比 Sampler 那套复杂，因为它要做两层中位数：

1
维护最近的 5000 个 forward step，先按 graph size 聚合取中位数，
2
再对这些中位数取中位数。只考虑有至少 10 个样本的 graph size，各 graph 等权。

第一层用 MedianTracker（两个 multiset 维护动态中位数）按 graph size 分组；第二层 _graph_medians 也是同一个 MedianTracker，把所有「组的组内中位数」再取一次中位数。更新是增量的：每步只更新新增/淘汰样本所属的 graph，以及它们在第二层中位数中的贡献，复杂度 $O(\log W)$。

为什么是「两层中位数」而不是「一个总中位数」

因为不同 graph size 的样本数量差异很大。小 batch 的桶（1、2、4、8）几乎每步都在用，样本很快攒到几千个；大 batch 的桶（比如 512）可能跑一整天才出现几十次。如果把这些样本混在一起取一个全局中位数，小桶的样本会压倒性地主导结果，大桶的斜率被淹没。

两层结构的处理方式是：先在每个 graph size 内部取中位数（把样本数压平），再对这些中位数取中位数（各 graph 等权）。源码注释把这句写得很直白——「各 graph 等权」。这是一个很典型的统计设计权衡：外层等权保证了不同 batch 规模对斜率估计的贡献相同，内层中位数保证了单个 graph 内部的异常样本（比如某步撞上了 EPLB 统计 flush）不会污染该 graph 的代表值。

而「为什么是中位数不是均值」，答案在于异常值的处理方式。这里的样本是运行时实测的 kernel 耗时，而运行时会周期性插入一些抖动源：EPLB 的统计需要跨 rank 做 EP all_gather，注释里说它会让个别步抖动十几毫秒（总览篇的 benchmark 口径一节里提到，会与 EPLB flush 撞上的区间是要整段丢弃的）。均值会被这种偶发的大值显著拉高，中位数不会——它是稳健统计量（robust statistic），对少量极端值不敏感。

这在调度决策上的意义是：估值偏高会让调度器过分保守（少验证草稿，牺牲吞吐），估值偏低会让它过分激进（多验证，拉长步时间）。中位数给出的是一个「典型步耗时」，正好是吞吐决策想要的量。

估计值怎么变成调度决策

MTPVarlen::schedule 是消费者。它先按接受概率对所有候选草稿位排序，然后做一个贪心搜索：从「只有 num_seqs 个 golden token」的起点开始，每次把一个草稿位加进预算，算出新的预期产出与新的估计耗时，取吞吐（预期产出/耗时）最高的那个点：

1
auto update_best = [&]() {
2
    while (cur_bsz_idx < num_bsz_selections && cur_bsz > _decode_perf_curve[cur_bsz_idx].first) {
3
        ++cur_bsz_idx;
4
    }
5
    if (cur_bsz_idx == num_bsz_selections) return;
6
    const auto &[graph_size, baseline_ms] = _decode_perf_curve[cur_bsz_idx];
7
    const double time_ms = baseline_ms + slope_ms * graph_size + sampler_ms;
8
    const double cur_throughput = cur_expected / time_ms;
9
    if (cur_throughput > best_throughput) {
10
        best_throughput = cur_throughput;
11
        best_bsz = cur_bsz;
12
    }
13
};

那个 time_ms 的表达式就是估计器的全部输出被消费的地方：

三项分别对应：这项硬件/模型在该形状下的基础耗时（实测曲线查表）、超出基线的形状增长（斜率估计）、以及采样侧的固定开销（按 num_seqs 分桶的中位数）。slope_ms 还被 std::max(0.0, ...) 夹住了下界：

1
// 没有 graph 达到样本门槛时使用启动曲线；所有候选共用斜率，按各自 graph size 修正。
2
// 以启动曲线为下界，避免负斜率外推到更大的 graph 后产生非正耗时。
3
const double slope_ms = std::max(0.0, _forward_time_slope.get());

负斜率的物理意义是「batch 变大反而更快」，这在某些区间可能真实存在（cache 命中、SM 占用率），但外推到更大的 graph 会产生非正耗时，所以直接夹到 0。

选完 best_bsz 后，第二个问题（预算怎么分）用一个独立的排序队列回答：

1
std::sort(acc_queue_for_seqs_q.begin(), acc_queue_for_seqs_q.end(), by_rate_desc);
2
seqlens_q.assign(num_seqs, 1);
3
for (int i = 0; i < best_bsz - num_seqs; ++i) {
4
    seqlens_q[acc_queue_for_seqs_q[i].second] += 1;
5
}

即：按接受概率从高到低给草稿位排序，把多出来的预算逐个发给概率最高的位。每个请求至少拿 1（golden token），最多不超过它自己的草稿数。

一个被反复强调的时序坑：用第 N-2 轮的接受率定预算

MTPVarlen::schedule 的注释里有一段很长的说明，是这类运行时反馈系统的典型难点：

1
需要保证 confidence_{i+1, ..., nextn} 不影响 draft_token_i 是否被验证。
2
假设此时在第 N 轮，我们期望使用 N-2 轮的 next_n_token_accept_rate 去计算 best_bsz (total_q)；
3
并使用 N-1 轮的 next_n_token_accept_rate 去分配每个请求的 seqlen_q
4
(non-ZOS 在这里做，ZOS 在 GPU 上修正)。

为什么总预算要用两轮前的数据？因为预算决定的是「本步要发射多大的 batch」，而这个问题必须在调度时刻回答——那一刻第 N-1 轮的结果还在途。而分批预算可以用第 N-1 轮的数据，因为——注意注释的后半句——ZOS 下这件事被挪到 GPU 上做了，由 compute_seqlens_q_by_confidence_rates 消费 prev.sampler_output.confidence_rates_gpu，那正是第 N-1 轮的产物，此刻已经在显存里了。

为了在 CPU 侧区分「第 N-1 轮的值」和「第 N-2 轮的值」，代码用了两个计数器：

1
// 1. schedule_cnt == process_cnt，说明请求每次调度上 engine 后都产生了 accept_rate：
2
//    accept_rates 是 N-1 轮的值，prev_accept_rates 是 N-2 轮的值。
3
// 2. schedule_cnt == process_cnt + 1，说明请求最新一次 forward 还没产生 accept_rate（如 ZOS）：
4
//    accept_rates 是 N-2 轮的值，prev_accept_rates 是 N-3 轮的值。

这正是前面那个「配对差不超过 1」的断言所保护的不变量。两个计数器看起来只是簿记，实际上承担的是在异步系统里精确定位某个值何时产生这件难事——而这恰恰是 ZOS 让一切都复杂起来之后的必然代价。

输出的生命周期规则也变了

LLMEngine.step() 的 docstring 里有一段专门讲返回值：

1
The returned RequestOutput objects are managed by RequestOutputPool.
2
RequestOutput will be automatically released in the next step() call.
3
If you need to retain them beyond the current step, you must call
4
pool.clone(list[RequestOutput]) to obtain copies, and release them via
5
pool.release(list[RequestOutput]).

这是 ZOS 对调用方接口的直接后果。因为输出对象被延迟释放，调用方在这一步拿到的 RequestOutput 指针，在下一步就可能指向被复用的内存。想留用必须显式 clone。这不是设计缺陷，而是「指针 + 池 + 延迟释放」这套组合的必然代价——它换来的正是 step 热路径上零对象构造与析构。总览篇把这种模式称为「DeepInfer 性能文化的缩影」，这里补上了它为什么必须配上一条显式的生命周期契约。