DeepInfer 源码拆解（二十六）：多 token 的账本——变长验证的扁平行布局与 RequestOutput 对象池

为什么需要 first_token_positions

先看一个具体例子。假设本步调度了三条请求：

• 请求 A：调度器决定验证 3 个草稿，另有 1 个主 token 行，共 seqlen_q = 4；
• 请求 B：验证 1 个草稿，共 2 行；
• 请求 C：没有可用草稿，只验证主 token，共 1 行。

那么 total_q = 4 + 2 + 1 = 7，logits 是一个 [7, vocab_size] 的张量。GPU 内核只认识行号 0 到 6，它不知道什么叫「请求 A」。要把行号翻译成请求语义，DeepInfer 使用三张长度都为 total_q 的表，外加一张长度 req_cnt 的偏移表：

1
def prepare_verification_metadata_ref(ftp, seqlen_q, total_q):
2
    row_to_req = torch.repeat_interleave(
3
        torch.arange(ftp.shape[0], dtype=ftp.dtype, device=ftp.device),
4
        seqlen_q,
5
        output_size=total_q,
6
    )
7
    local_step = torch.arange(total_q, dtype=ftp.dtype, device=ftp.device) - ftp[row_to_req.long()]
8
    is_verify = local_step != seqlen_q[row_to_req.long()] - 1
9
    return row_to_req, local_step, is_verify

这五行代码值得逐行解释，因为后面所有账本都建立在它们之上：

ftp（first token positions） 是每条请求的第一行在扁平张量里的偏移。对上面的例子，ftp = [0, 4, 6]。它满足两个约束：ftp[0] == 0；ftp[i+1] == ftp[i] + seqlen_q[i]。换句话说，ftp 就是 seqlen_q 的前缀和（exclusive prefix sum）。源码在测试里正是用 ftp[1:] = seqlen_q.cumsum(0)[:-1] 构造它。

row_to_req 把每个行号映射回请求号。repeat_interleave 的语义是「把第 i 个元素重复 seqlen_q[i] 次」，所以 row_to_req[0..3] = 0、row_to_req[4..5] = 1、row_to_req[6] = 2。这就是行号到请求号的换算表。

local_step 是行在它所属请求内部的偏移：请求 A 的 4 行是 0、1、2、3，请求 B 的 2 行是 0、1，请求 C 是 0。它由「全局行号减去该请求的 ftp」得到。

is_verify 标记这一行是不是「验证行」。每个请求的最后一行是主 token 行（local_step == seqlen_q - 1），其余行都是草稿验证行。因此 is_verify = local_step != seqlen_q - 1。这个布尔量在下游至少有两个用途：融合采样内核据此决定这一行是做拒绝采样的验证，还是做一次普通采样；最终打包时据此把「被接受的草稿」和「主 token」分开统计。

（图源：DSpark 图 4，Qwen3-4B 目标模型；左三栏分别为 Math、Code、Chat 三个领域，右栏为顺序头的额外延迟）

这张图解释了一个前提：草稿块可以很长，但真正进入验证的前缀长度由调度器逐请求决定。DSpark 的硬件感知前缀调度器会根据置信度与当前吞吐曲线，把每条请求的验证长度压到不同的值；落到引擎里，就是 seqlens_q[i] 在 batch 内取 1、2、…、nextn+1 等不同值。如果没有一张按行组织的索引表，GPU 内核和后面的 C++ 收账都无从知道「第几行是谁的」。

一张图看清扁平布局

（本站自制图，依据 tilelang_ops/sampler.py 中 prepare_verification_metadata_ref 与 fused_sample 的行布局绘制）

这张图对应上面的例子。左边是三条请求各自的验证长度，中间是压平后的 7 行，右边三列是 row_to_req、local_step、is_verify 的取值。注意最右边请求 C 只有一行，local_step = 0，而 seqlen_q - 1 = 0，所以 is_verify = False——它直接走普通采样，不参与拒绝采样。这个细节很重要：「没有草稿」不是特殊分支，而是「seqlen_q = 1 的退化情形」。同一套内核、同一套索引表就能处理，不需要为「纯 decode」另写一条路径。

为什么不用二维张量

一个自然的疑问是：既然每条请求的验证长度不同，为什么不直接用一个 [req_cnt, max_nextn + 1] 的二维张量，短的行用 padding 补齐？这在很多推理框架里确实是常规做法，但 DeepInfer 选择扁平化，原因有三层：

1. kernel 的负载均衡。变长验证的 seqlens_q 分布很不均匀：有的请求验证 5 个草稿，有的只验证 1 个。二维布局会让同一列的行承载完全不同的计算量，而扁平布局把所有行当成独立元素，GPU 的线程块可以均匀地按行分片，不会因为某一列大量 padding 而空转。DSpark 论文在部署部分也明确写了这一点：把不同请求的 token 全部压平、当成独立元素处理，序列内依赖通过稀疏注意力里的 marker 张量表达，而不是靠二维 mask。
2. 与 CUDA Graph 的形状契约一致。DeepInfer 的 decode 走分档 CUDA Graph，图按 graph_size（总行数）选择，而不是按 req_cnt × 固定列数。扁平化让「图的形状」只依赖一个整数，捕获与重放的档位管理简单得多。这一点在CUDA Graph 捕获与重放里已经讲过。
3. 跨语言的接口更窄。C++ 调度器只需要接收一个 [total_q] 的接受结果张量和一个 [req_cnt] 的 first_token_positions，就能还原每条请求的接受前缀；不需要在 C++ 侧再实现一套二维索引逻辑。接口越窄，跨语言边界上能出的错越少。

代价是索引表必须由生产方和消费方共同遵守。first_token_positions 不是可选提示，而是协议：build_request_outputs 在收到 accepted_token_ids 时，如果没有同时收到 first_token_positions，会直接 TORCH_CHECK 失败。这类显式失败比「悄悄按错误偏移读」要好得多。

用一次前缀和表达「前缀存活」

1
def get_accept_result_ref(raw_accept, golden, draft_token_ids,
2
                          candidate_token_ids, is_verify, ftp, row_to_req):
3
    accept = raw_accept
4
    tokens = candidate_token_ids
5
    if golden is not None:
6
        use_golden = golden != -1
7
        accept = torch.where(use_golden, golden == draft_token_ids, accept)
8
        tokens = torch.where(use_golden, golden, tokens)
9
    blocked = ~(is_verify & accept)
10
    prefix = torch.cat([blocked.new_zeros(1, dtype=torch.int32), blocked.cumsum(0, dtype=torch.int32)])
11
    alive = prefix[:-1] == prefix[ftp[row_to_req.long()].long()]
12
    effective_accept = alive & is_verify & accept
13
    is_final = alive & blocked
14
    return tokens, effective_accept, is_final

这段代码的妙处在于，它没有对每条请求写循环，而是用一次全局 cumsum 同时处理所有请求。理解它的关键是 blocked 的定义：「不是验证行」或者「验证了但没接受」都算 blocked。于是每条请求内部，blocked 的第一个位置就是「该请求的接受前缀到此为止」的位置：

• 如果请求 A 的三条草稿全部接受，那么前三行的 blocked = False，第四行是主 token 行（is_verify = False），blocked = True；
• 如果请求 A 的第二条草稿被拒，那么第二行 blocked = True，第三行虽然可能 accept = True，但它已经在前缀之外。

cumsum 之后，prefix[k] 表示「第 0 到 k-1 行里有多少个 blocked」。对请求 i，prefix[ftp[i]] 是它开始之前的全局 blocked 计数。某一行 k 属于请求 i 时，如果 prefix[k] == prefix[ftp[i]]，说明从请求起点到这一行之前没有任何 blocked——也就是这一行仍在「存活前缀」内。alive 就是这个判断。

于是：

• effective_accept = alive & is_verify & accept：真正被接受的草稿行；
• is_final = alive & blocked：每条请求的「第一处断点」，它可能是第一个被拒的草稿行，也可能是永远存在的主 token 行。

is_final 是后面所有计数的锚点。finalize_verification_ref 用它一次性算出每条请求的最终行、接受数量与输出位置：

1
rows = torch.arange(q, dtype=ftp.dtype, device=ftp.device)
2
final_row = rows[is_final]
3
last_valid_offset = final_row - ftp          # 每条请求的最终行相对起点的偏移
4
...
5
num_generated_tokens.scatter_add_(0, step_index, accepted_mask.to(torch.int32))
6
num_verified_tokens.scatter_add_(0, step_index, is_verify.to(torch.int32))

last_valid_offset 的语义是「本步这条请求的最后一个有效 token 位于它自己的第几行」。对上面的例子：请求 A 若接受 2 条草稿，最终行是第 2 行（主 token），last_valid_offset = 2；请求 B 若第一条草稿就被拒，最终行是第 1 行，last_valid_offset = 1。这个值随后会被传给草稿头，作为下一轮生成草稿的起点。

-1 哨兵：被拒绝的行怎么表达

finalize_verification_ref 里还有一行不起眼但很关键的代码：

1
verified_token_ids=torch.where(scalar_mask, effective_token_ids, -1).to(torch.int32),

被接受的草稿行写入真实 token id，没被接受（包括前缀之外的草稿行）写入 -1。为什么不用 0？因为 0 是合法的 token id（DeepSeek 的词表里 0 是 BOS）。用一个不可能出现在正常输出里的负数作哨兵，下游就能用「遇到 -1 就停」来表达前缀语义，而不需要再读一次 effective_accept 布尔张量。

这个哨兵在 C++ 侧被消费了两次：

1
if (accepted_token_ids_base_ptr[flat_idx] != -1)
2
    num_accepted++;
3
else
4
    break;

第一次是在统计 accept_tok 时；第二次是在真正追加 token 时：

1
if (token_id == -1)
2
    break;
3
appended_tokens++;

两处都是「遇到 -1 立刻停」，而不是「跳过 -1 继续读」。这正是前缀语义在数据层面的表达：哨兵不是过滤条件，而是终止符。如果哪天有人把它改成 continue，错误会非常隐蔽——被拒绝的草稿后面如果碰巧有被接受的草稿，就会被错误地追加进序列，而前缀约束已经被破坏。

为什么 verified_token_ids 是 [total_q] 而不是 [req_cnt, nextn]

这与扁平布局是同一个设计选择的一体两面。[req_cnt, nextn] 需要每条请求都补齐到最大草稿数，而实际接受的草稿数可能从 0 到 5 不等。补齐会带来两个问题：一是下游必须知道「哪些位置是 padding」，二是 MTP 的草稿池本身已经按 slot 存了定长 nextn 行，再在输出侧补一次会造成两套长度语义。扁平加哨兵则只有一种长度语义：读到 -1 为止。

池的形状与 slot 0

MTPInfoPool 的文档字符串把布局写得很清楚：

1
Stable GPU storage for MTP proposals indexed by scheduler slots.
2
Each slot owns nextn rows. Slot 0 is reserved for padding; all other
3
slots store one row per MTP proposal.
4
pool_rows = (max_num_seqs + 1) * nextn

也就是说，池子按「请求槽位」组织，每个槽位固定 nextn 行；槽位 0 永远不用，留给 padding。调度器为每条进入运行队列的请求分配一个槽位：

1
void allocate_mtp_info_slot(Sequence &seq) {
2
    if (nextn == 0 || seq.mtp.info_slot != 0) return;
3
    TORCH_CHECK(!free_mtp_info_slots.empty(), "No free MTP info slot for request ", seq.request.request_id);
4
    seq.mtp.info_slot = free_mtp_info_slots.back();
5
    free_mtp_info_slots.pop_back();
6
}

槽位号从 1 开始，随请求的生命周期分配与释放。请求结束时 release_mtp_info_slot 把槽位还回 free list，并清空 Sequence 里的记录。这个设计的直接好处是：草稿的位置与请求的物理槽位绑定，而不是与 batch 内的临时下标绑定。即使下一步请求的排列顺序变了、batch 组成变了，只要请求还活着，它的草稿就还在原来的槽位里。

写一步、读一步

草稿头在本步生成下一轮的草稿后，把结果写进池子：

1
def store(self, write_slots, prediction, logprobs, logprobs_index):
2
    store_mtp_info(
3
        write_slots,
4
        prediction.token_ids,
5
        logprobs,
6
        logprobs_index,
7
        prediction.confidence_rates,
8
        self.token_ids,
9
        self.logprobs,
10
        self.logprobs_index,
11
        self.confidence,
12
    )

write_slots 是 [req_cnt] 的槽位号；store_mtp_info 的内核把每条请求的 nextn 行写到 slot * nextn + j。写到槽位 0 的请求会被跳过，从而保留槽位 0 的 padding 内容。下一步验证时，Sampler._verification 用 mtp_info_read_rows 把池子里的对应行读回来：

1
logprobs, ... = process_sampling(
2
    logits,
3
    driver_inputs.kernel,
4
    self.force_topk,
5
    row_to_req=row_to_req,
6
    local_step=local_step,
7
    is_verify=is_verify,
8
    last_nextn_token_logprobs=mtp_info_pool.logprobs if speculative_inputs is not None else None,
9
    ...
10
    mtp_info_read_rows=speculative_inputs.mtp_info_read_rows if speculative_inputs is not None else None,
11
)

（本站自制图，依据 mtp_info_pool.py 与 scheduler.hpp::allocate_mtp_info_slot 绘制）

这张图把三件事画在一起：左侧是本步的草稿生成，中间是常驻 GPU 的槽位池，右侧是下一步的验证读取。注意图中标注的「槽位 0」：它不是某条请求，而是所有「本步没有草稿」请求的公共替身。写入时跳过它，读取时它的 logprobs 是 -inf（除 padding token），这样验证内核即使读到它也不会产生有效的接受判定。

为什么不能直接返回张量

一个自然的替代方案是：草稿头把本步生成的 token_ids、logprobs 直接作为 SamplerOutput 的一部分返回，下一步再传回来。DeepInfer 没有这么做，原因有三点：

1. ZOS 的延迟消费。在零开销调度下，本步的输出要到下一步才被消费；如果草稿张量是临时分配，生命周期会跨 step，Python 侧必须持有引用并保证不被回收，这会把内存管理复杂度从「一块常驻池」扩散到「每个 step 的临时对象」。
2. CUDA Graph 的地址稳定性。草稿头的图在捕获时绑定了池子的地址；如果每步换一块新张量，图就必须重新捕获，或者退化成 eager 路径。常驻池让地址在捕获前后保持不变。
3. 槽位与请求生命周期对齐。请求被抢占、重算、恢复时，槽位可以跟着 Sequence 一起保存/释放；如果草稿是随 step 返回的临时张量，重算时就得额外决定「上一轮的草稿是否还属于这条请求」。MtpState 里的 reset_for_recompute()、take_draft_info_slot() 正是为这种生命周期准备的。

一、“golden” 这个状态名

读 process_outputs 之前，先看 Sequence 里一个容易让人困惑的成员：MtpState::last_output_token。它保存的是「上一步输出的最后一个 token」，源码里的做法是：

1
seq.mtp.last_output_token = output_token_ids_data[i];

而注释直接称它为「本轮的 golden token」。为什么叫 golden？因为目标模型在本步验证草稿时，第一条主 token 行的输入恰恰就是这个 token：上一轮已经确定的输出，构成了这一轮第一个位置的确定性输入。草稿头的「串行头」也以它为起点继续生成下一轮的草稿——所谓 anchor（锚点）。

理解了这一点，sampled_token_ids_gpu、predict_token_ids_gpu 这些字段为什么是「GPU 常驻副本」就清楚了：ZOS 下，下一步的输入里 target 位置用的是上一步采样结果，draft 位置用的是上一步草稿池。这两份数据都在 GPU 上，下一步的前向直接做 D2D 拷贝，不需要经 CPU 绕一圈。

乐观推进与提交

先看账面。调度时，调度器已经按 seqlens_q 乐观地推进了 KV 账本：

1
seq.cache.advance_optimistically(scheduler_outputs.seqlens_q[i]);

advance_optimistically 会把 previous_num_tokens 记成推进前的值，把 num_tokens 加上 seqlen_q。这样做的原因是：KV 块在调度阶段就要按最坏情况分配，不能等采样结果回来再决定；否则显存可能不够。采样结果回来之后，真正被接受的 token 数可能小于 seqlen_q（草稿被拒），于是要「回填」：

1
// decode：把 cache.num_tokens 从乐观推进值回填为实际接受数（未被接受的 draft 不占 KV）
2
seq.cache.commit_processed_tokens(token_cnt);

commit_processed_tokens 的实现只有一行：

1
void commit_processed_tokens(int token_count) { num_tokens = previous_num_tokens + token_count; }

这行代码是整本账的收口：先按上限记账，再按实际用量结算。previous_num_tokens 是推进前的稳定基线，token_cnt 是本步真正追加的 token 数。中间那个「多算的部分」就是被拒绝的草稿——它们曾经占用过 KV 块的形状预算，但没有被写进序列，所以在提交时被扣掉。

追加循环与停止条件

真正写 token 的循环如下（省略了日志与分支）：

1
token_ids = accepted_token_ids_vec;
2
seq.mtp.last_output_token = output_token_ids_data[i];
3
token_ids.push_back(seq.mtp.last_output_token);
4

5
int token_cnt = 0;
6
for (auto token_id : token_ids) {
7
    seq.append_token_id(token_id);
8
    seq.append_compressed_token_id(map_engram(token_id));
9
    token_cnt++;
10
    _check_stop(seq);
11
    if (seq.is_finished()) break;
12
}
13
if (seq.is_finished()) release_mtp_info_slot(seq);
14
pushed_tokens_before_finish.push_back(token_cnt);

这段代码回答了一个很实际的问题：如果一条请求在接受的前缀中途就触发了停止条件，后面的 token 怎么办？ 答案是不追加。_check_stop 在每次追加后检查（EOS、max_tokens、结构化输出失败等），一旦命中就 break，token_cnt 停在命中位置。于是 token_cnt 可能小于「接受前缀 + 1」，这正是 pushed_tokens_before_finish 名字的由来：它是「结束之前实际被 push 进去的 token 数」。

这个值随后被 build_request_outputs 用来决定本步要发出几个 RequestOutput、哪一个带 finish_reason。如果没有它，输出构造就只能假设「接受多少就发多少」，于是用户会看到 EOS 之后的 token——一个非常典型、也非常难在离线测试里发现的 bug。

被抢占与已中止的请求

process_outputs 还有两个提前返回的分支：

1
// seq is already aborted
2
if (seq.is_finished()) {
3
    release_mtp_info_slot(seq);
4
    pushed_tokens_before_finish.push_back(0);
5
    ...
6
    continue;
7
}
8

9
// ZOS 异步：本步尚在途时就被 _preempt 抢占重算的 seq
10
if (!scheduler_outputs.prompt_run && seq.schedule.status == SequenceStatus::WAITING) {
11
    pushed_tokens_before_finish.push_back(0);
12
    ...
13
    continue;
14
}

两个分支都写入 0，含义是「本步这条请求没有任何 token 可以发给用户」。前一个是请求已经结束（比如被 abort），后一个是请求在本步前向还在飞的时候就被抢占、放回等待队列，本步结果作废。后者是 ZOS 特有的竞态：采样结果回来时，Sequence 的状态可能已经变了。pushed_tokens_before_finish = 0 就是一个显式的「本步无输出」标记，避免把作废的 token 追加进去造成重复输出。

接受草稿先发，主 token 后发

对每条请求，函数先处理被接受的草稿，再处理主 token：

1
if (has_accepted_tokens) {
2
    for (int j = 0; j < num_draft_i; j++) {
3
        int flat_idx = ftp_i + j;
4
        int token_id = accepted_token_ids_base_ptr[flat_idx];
5
        if (token_id == -1) break;
6
        appended_tokens++;
7
        if (appended_tokens > pushed_tokens) break;
8
        ...
9
        bool is_finished = appended_tokens < pushed_tokens ? 0 : seq->is_finished();
10
        int num_output_tokens = seq->get_output_len() - pushed_tokens + appended_tokens;
11
        request_outputs.emplace_back(request_output_pool->allocate(...));
12
        golden_output = ...;
13
    }
14
}
15
if (appended_tokens < pushed_tokens) {
16
    appended_tokens++;
17
    // 主 token 的 RequestOutput
18
}

这里有三个不变量：

1. appended_tokens 从 1 开始计，且最多到 pushed_tokens。草稿循环里每发一个就加一，发到 pushed_tokens 为止；主 token 只在 appended_tokens < pushed_tokens 时才发。这保证了「发出去的 token 数」严格等于「实际追加进序列的 token 数」。
2. finish_reason 只落在本步最后一个输出上。is_finished = appended_tokens < pushed_tokens ? 0 : seq->is_finished()：前面的草稿输出即使发生在「将要结束」的请求上，也不带结束标记；只有最后一个才带。否则客户端会在收到第一个草稿 token 时就把流关掉。
3. num_output_tokens 是「累计口径」而不是「本步增量」。get_output_len() - pushed_tokens + appended_tokens 表示「本步开始前的输出长度 + 本步已经发出的数量」。这个字段用于 usage 统计和客户端判断进度，不是增量。

golden_output 是一个很有意思的变量：它指向本步这条请求最后发出的那个 RequestOutput，然后被写入 draft_tok、accept_tok 与 confidence_score：

1
golden_output->draft_tok = draft_tok_val;
2
golden_output->accept_tok = accept_tok_val;
3
if (seq->mtp.prev_confidence_raw.size() > 0) {
4
    golden_output->confidence_score = seq->mtp.prev_confidence_raw;
5
}

为什么统计信息只挂在最后一个输出上？因为对客户端来说，一条请求的「本步结果」是一个整体：它想知道这一步验证了几个草稿、接受了几个、置信度误差多大，而不是在收到每个草稿 token 时都重复一遍。把这些字段集中到最后一个输出，既减少重复，也避免了「中间输出被提前消费」时的语义歧义。

-9999 钳位：为什么不是 -inf

输出构造里有大量这样的代码：

1
float token_logprob = std::max(token_logprobs_ptr[i], -9999.0f);

logprobs 在数学上可以取到 -inf：被 top-p 剪掉的 token、词表外的图像 token、强制 top-k 模式下未入选的位置，都会得到负无穷。但 -inf 一旦进入 JSON 序列化，就会变成 null 或直接解析失败；进入某些客户端的浮点显示，又会变成 -Infinity 这种非标准字面量。-9999 是一个「足够小、但仍然是有限数」的钳位值，保证协议层永远拿到合法的 JSON 数字。

这个细节和输出账本的主题看似无关，其实是同一类问题：内部表示可以为了计算方便而使用哨兵/无穷，跨出引擎边界时必须换成协议安全的表示。-1 在内部是接受前缀的终止符，但在 RequestOutput 里不会出现；-inf 在内部是合法的 logprob，但到了输出层会被钳成 -9999。

被忽略的请求：ignored_seqs

还有一类请求不走主循环：ignored_seqs。它们通常是因为 prompt 长度超过限制而被拒绝的请求，本步没有进入前向，也没有采样结果。但客户端仍然需要一个明确的结束原因：

1
for (auto &seq : scheduler_outputs.ignored_seqs) {
2
    request_outputs.emplace_back(request_output_pool->allocate(
3
        seq->request.request_id, seq->get_prompt_len(), seq->get_last_token_id(),
4
        get_finished_reason(seq->schedule.status), seq->is_finished(), ...,
5
        seq->get_output_len() + 1, PromptLogprobs()));
6
}

注意它的 token_id 是 get_last_token_id()（prompt 的最后一个 token），num_output_tokens 是 get_output_len() + 1（保证至少为 1），logprobs 用无效值 1.0 / -1.0 填充。这是一条「没有生成、但必须告诉客户端为什么结束」的输出。get_finished_reason 会把 FINISHED_IGNORED 映射成 "prompt_length_exceeded"，客户端据此显示错误，而不是一直等待一个永远不会到来的 token。

固定 10000 槽的池子

RequestOutput 不是 Python 类，而是一个 C++ 结构体，里面有 std::string、std::vector、c10::optional。如果每个 step 都为每条输出 new 一个对象，再交给 Python 持有，那么在高并发下会产生大量小对象分配与析构，且 Python GC 的时机不可控。DeepInfer 的做法是预先分配一整块内存，用 placement new 在上面构造：

1
class RequestOutputPool {
2
    RequestOutput *pool_{nullptr};
3
    std::vector<RequestOutput *> free_;
4
public:
5
    static constexpr size_t max_size = 10000;
6

7
    RequestOutputPool() {
8
        pool_ = reinterpret_cast<RequestOutput *>(
9
            aligned_alloc(alignof(RequestOutput), sizeof(RequestOutput) * max_size));
10
        ...
11
    }
12

13
    template <typename... Args>
14
    uint64_t allocate(Args &&...args) {
15
        TORCH_CHECK(!free_.empty(), "RequestOutputPool exhausted (max_size=", max_size,
16
                    "); outputs are not being released");
17
        auto p = free_.back();
18
        free_.pop_back();
19
        new (p) RequestOutput(std::forward<Args>(args)...);
20
        return reinterpret_cast<uint64_t>(p);
21
    }
22

23
    void release(py::array_t<uint64_t> ptrs_py) {
24
        const uint64_t *ptrs = ptrs_py.data();
25
        int size = ptrs_py.size();
26
        for (int i = 0; i < size; ++i) {
27
            auto p = reinterpret_cast<RequestOutput *>(ptrs[i]);
28
            std::destroy_at(p);
29
            free_.push_back(p);
30
        }
31
    }
32
};

几个设计点值得拆开：

• max_size = 10000 是硬上限。池子耗尽时直接 TORCH_CHECK 失败，而不是退化成 new。这是一种「宁可显式失败，不要静默降级」的选择：如果输出没有被及时释放，说明调用方的生命周期管理有问题，继续分配只会把问题拖到更晚、更难定位的地方。
• 返回的是裸指针的整数值。allocate 返回 uint64_t，build_request_outputs 把一批指针塞进 py::array_t<uint64_t>。Python 侧看到的是一串整数，只有需要字段时才通过 _C 的绑定解码成对象。这避免了每个输出一次的 Python 对象构造，也避免了跨语言拷贝。
• release 接受指针数组。释放不是逐个调用，而是一次性把一批指针还回 free list。这让「释放本步所有输出」变成一次批量操作，和 DeferredRelease 的批量回收天然对齐。

clone 与 release：谁拥有这些对象

池子的所有权模型可以用一句话概括：池子拥有内存，调用方拥有「本步的租约」。step() 返回的指针在本步有效；下一步 step() 会回收上一批输出。如果调用方要把输出保留到下一步之后（例如异步流式发送），必须显式 clone：

1
for raw_output in step_outputs:
2
    output = RequestOutput.clone(raw_output)
3
    ...

clone 的实现是一行 std::make_unique<RequestOutput>(*raw)——它并不占用池子的槽位，而是在堆上新建一个独立的 RequestOutput，拷贝构造全部字段（包括 std::string 与 std::vector）。这样原始输出可以在下一步被回收，而副本由 Python 持有引用、随引用计数回收。也正因为它不占池子，大量克隆的输出只受 Python 堆的管理；而池子侧要防的，只是「本步产出的输出没有被 DeferredRelease 回收」。

为什么必须 gil_scoped_release

build_request_outputs 在构造循环开始前有一行：

1
py::gil_scoped_release release;

然后在返回前重新获取：

1
py::gil_scoped_acquire acquire;
2
return py::array_t<uint64_t>(request_outputs.size(), request_outputs.data());

这看起来只是性能优化，实际上和 ZOS 的正确性有关：输出构造发生在引擎侧的后处理阶段，这段时间 GPU 正在执行下一拍的前向。如果持有 GIL，Python 侧的其他线程（例如结构化输出的编译线程、流式发送线程）就会被阻塞，CPU 后处理的时间会被拉长，最终把 GPU 拖空。释放 GIL 让构造循环真正并行于其他 Python 工作，也是「把 CPU 工作藏进 GPU 阴影里」这一整套设计的组成部分。