2026 先导杯实战：别只盯着大算子，把 Qwen3.5 vLLM 单 Token 输入准备搬上 DCU

2026 先导杯实战：别只盯着大算子，把 Qwen3.5 vLLM 单 Token 输入准备搬上 DCU

在海光 BW1000 DCU（gfx936 架构）上，我们把 Qwen3.5 vLLM 稳态 decode
的单 Token 输入准备组件从约
0.259 ms 降到 0.101 ms，组件延迟降低约 60.9%，加速约 2.56 倍。
这篇文章不讲玄学调参，而是完整拆解一个很容易被忽略的问题：模型算完一个
token 后，系统为了计算下一个 token，到底还在忙些什么？

关键词：Qwen3.5、vLLM、海光 BW1000 DCU、gfx936、Triton、单请求推理、
单 Token decode、DCU 常驻、输入准备

项目背景：来自 2026 先导杯真实赛题

本文优化来自 2026 先导杯，即全国大学生计算机系统能力大赛智能计算创新设计赛
“基于国产加速卡的千问大模型推理服务优化”赛题。比赛要求在统一软硬件环境中，
使用单张海光 BW1000 DCU（gfx936 架构）运行 Qwen3.5-27B BF16 模型，并通过
vLLM 对外提供推理服务。

赛题的工作负载与常见高并发服务不同：评测固定为单请求并发、batch size 为 1，
分别统计 4K-8K、8K-16K 和 16K-32K 三档输入的输出吞吐，同时设置 SLA 和回答
精度约束。这意味着优化不能只追求某个微基准跑得快，还必须同时满足：

• 不改变模型权重、模型结构、tokenizer 和采样语义；
• 不通过量化、剪枝、投机解码、提前退出或答案缓存绕过计算；
• 不修改max-model-len、max-num-seqs、max-num-batched-tokens以及 batch
  scheduler；
• 输出文本和生成 Token 数需要与基线一起检查，避免把数值漂移误当成性能提升；
• 不满足专用条件时必须回退 vLLM 原始路径。

固定 batch=1 一方面减少了通用动态批处理带来的优化空间，另一方面也让执行形态
非常稳定：进入稳态 decode 后，每轮通常只有一个请求和一个新 Token。本文正是利用
这个确定性，把通用输入准备中的重复工作变成 BW1000 上的单 Token 专用快路径，
没有触碰比赛禁止修改的调度参数。

1. 为什么要优化一个看起来只有零点几毫秒的步骤？

做大模型推理优化时，我们通常先盯着 GEMM、Attention、RMSNorm 这些“大算子”。
这很合理，因为它们承担了绝大部分浮点计算。但当模型算子已经经过多轮优化，
尤其是在batch size 为 1、每轮只生成一个 token的在线推理场景中，另一类
开销会逐渐显眼：

• CPU 生成位置、序列长度等元数据；
• CPU 根据 KV cache 页表计算 slot mapping；
• 多个很小的 Host-to-Device（H2D）拷贝；
• DCU 上已有的数据被再次搬运；
• 每个 decode step 都重复执行的 Python、NumPy 和调度准备工作。

这些操作单次只花零点几毫秒，看起来远小于一次完整的模型前向。但 decode 不是
只执行一次。生成 512 个 token 时，相同准备流程大约要重复 511 次；生成越长，
固定开销累计得越明显。

可以把大模型 decode 想成一条流水线：

上一个 token 采样完成 | v 准备 input_ids、position、序列元数据、KV cache 写入位置 | v 执行模型前向 | v 计算 logits 并采样下一个 token | +--------------------> 进入下一轮

只优化中间的“模型前向”，却放任每一轮开头做一串细碎的 CPU 工作，就像把高速
公路修宽了，却没有处理收费站。本文的优化目标正是这个“收费站”。

2. 背景：vLLM 异步调度已经把 Token 留在 DCU，为什么还不够？

本文讨论的 vLLM 推理引擎在源码中位于vllm/v1/目录。V1只是 vLLM
内部的新引擎架构名，不是 Qwen3.5 的版本号。为了让非 vLLM 开发者也容易理解，
后文统一称为“vLLM 推理引擎”。

vLLM 异步调度有一个重要特征：上一轮采样得到的 Token 可以保留在 DCU，下一轮
不必先把 Token 取回 CPU，再传回 DCU。代码中的prev_sampled_token_ids就保存了
这份 DCU 结果。

这已经避免了一次明显的 DCU -> CPU -> DCU 往返，但通用输入准备路径仍要兼容
许多复杂场景：

• 同时调度多个请求；
• 每个请求调度不同数量的 token；
• chunked prefill；
• speculative decode；
• 请求插入、移除和重排；
• LoRA、多模态输入和各种并行方式；
• 一个模型包含多组不同的 KV cache。

为了覆盖这些情况，通用路径会执行 request 展开、前缀和、位置生成、token 索引
查找、slot mapping 计算和多份元数据拷贝。这套逻辑很稳健，但对下面这个固定场景
来说有些“用力过猛”：

请求数 = 1 本轮 scheduled token = 1 请求顺序 = 没有变化 上一轮 sampled Token = 已经在 DCU 序列位置 = 上一位置 + 1 无 speculative decode、LoRA 和 encoder 输入

在这个场景里，大部分结果都能直接写出来：

• input_ids[0]就是上一轮 DCU 上的 sampled Token；
• position[0]就是当前已计算 token 数；
• query_start_loc固定为[0, 1, 1, ...]；
• seq_lens固定为[position + 1, 0, 0, ...]；
• 当前请求应被采样，因此discard_request_mask全为 false；
• logits_indices永远是[0]；
• KV cache slot 只需一次整数除法、取余和查表。

问题的本质因此变成：数据已经在 DCU，计算也只是几个标量和数组写入，却仍然
绕到 CPU 组织一遍，再通过多次小拷贝送回 DCU。

3. 原始路径到底做了什么？

以下流程是对 vLLM 通用输入准备逻辑的简化。它不是错误实现，而是为通用性
支付的合理成本。

CPU DCU --- --- 展开 req_indices 计算 cumsum / arange 计算 position 查找本轮 input token sampled token 已在显存 根据 block table 计算 slot mapping 填写 query_start_loc 填写 seq_lens 填写 discard_request_mask 复制 block table -------------------------> block table 复制 slot mapping ------------------------> slot mapping 复制/散射 sampled token ------------------> input_ids 复制 M-RoPE / position -------------------> positions 复制 query_start_loc ---------------------> query_start_loc 复制 seq_lens ----------------------------> seq_lens 复制 discard mask ------------------------> discard mask

这里有两个值得注意的细节。

3.1 小拷贝不等于免费

一次只复制一个 token、几个整数或一小段元数据，数据量确实不大。但小拷贝的成本
不只来自带宽，还包括运行时调用、stream 排队、同步关系和 Python 调度。对这类
数据，传输“启动成本”往往比真正搬运字节的成本更重要。

3.2 slot mapping 是一个简单但高频的计算

Paged KV cache 需要把逻辑 token 位置转换为物理 cache slot。单组 cache 的公式是：

logical_block = position // block_size block_number = block_table[logical_block] block_offset = position % block_size slot = block_number * block_size + block_offset

公式非常简单，但旧路径要先在 CPU/NumPy 中计算，再把结果复制到 DCU。单请求、
单 Token 时，真正的工作量小到不值得走这一圈。

4. 核心思路：让 decode 的数据闭环留在 DCU

优化后的原则可以概括成一句话：

已经在 DCU 上的数据，不要为了准备下一轮输入而回到 CPU；能由一次 kernel launch
完成的元数据写入，不要拆成多次小拷贝。

新路径做了四件事：

1. 直接读取 DCU 上一轮的sampled_token_ids；
2. 用一次 Triton kernel launch 写完输入和 attention 元数据；
3. 在 DCU 上读取 block table 并计算每个 KV cache group 的 slot；
4. 给 block table 增加 dirty tracking，内容没变时不再重复 H2D。

优化后的数据流如下：

常驻 DCU 显存 +--------------------------------+ 上一轮采样 ---> sampled_token_ids | | block_tables | | output buffers | +----------------+---------------+ | v 一次 Triton kernel launch | +-----------------+------------------+ | | | v v v input_ids position/M-RoPE attention metadata | v 各 KV group slot mapping

注意，这里说的是“一次 kernel launch”，不是整个任务只启动一个 Triton program。
实际 grid 大小等于 KV cache group 数：group_id == 0的 program 负责公共元数据，
每个 group 对应的 program 再分别计算自己的 slot mapping。这样既减少 launch 数，
又能兼容 Qwen3.5 中存在多组 KV cache 的情况。

5. 实现一：一个 Triton kernel 写完本轮 DCU 输入

专用实现位于
vllm/v1/worker/gpu_decode_input_prep.py。

下面是删去类型和边界细节后的核心逻辑：

@triton.jitdefprepare_single_token_decode(...):group_id=tl.program_id(0)# 公共元数据只写一次ifgroup_id==0:token_id=tl.load(sampled_token_ids_ptr)tl.store(input_ids_ptr,token_id)tl.store(position_ptr,position)tl.store(query_start_loc_ptr,[0,1,1,...])tl.store(seq_lens_ptr,[position+1,0,0,...])tl.store(discard_request_mask_ptr,0)# 每组 KV cache 都有自己的 block size、页表和 slot 输出block_size=tl.load(block_sizes+group_id)block_table=load_pointer(block_table_ptrs+group_id)block_number=tl.load(block_table+position//block_size)slot=block_number*block_size+position%block_size slot_mapping=load_pointer(slot_mapping_ptrs+group_id)tl.store(slot_mapping,slot)

真实实现还处理了两个重要分支：

• 模型使用 M-RoPE 时，一次写入三路位置；
• 启用 prompt embeddings 时，同时写is_token_ids。

这些分支都是 Triton 的编译期常量，不会在热路径里引入复杂的动态判断。

5.1 为什么要保存“指针数组”？

Qwen3.5 可能有多个 KV cache group，每组都有独立的 block table 和 slot mapping。
Python 初始化阶段会把这些张量的显存地址整理为 DCU 上的uint64数组：

self.block_table_ptrs=torch.tensor([table.block_table.gpu.data_ptr()fortableinblock_tables],dtype=torch.uint64,device=device,)self.slot_mapping_ptrs=torch.tensor([table.slot_mapping.gpu.data_ptr()fortableinblock_tables],dtype=torch.uint64,device=device,)

kernel 根据group_id读取对应地址，再把它转换为正确类型的指针。这样每一轮
decode 只传入固定的持久化状态，不必由 Python 为每个 group 分别发起操作。

5.2 为什么复用logits_indices=[0]？

单请求、单 token 下，本轮唯一需要计算 logits 的行就是第 0 行。因此
logits_indices不必每次由query_start_loc[1:] - 1动态生成，可以在初始化时
创建一个 DCU 常量张量并持续复用。

这是一个很小的优化，但它体现了同一个原则：稳态 decode 中不变的对象，不要每轮
重新构造。

6. 实现二：block table 没变，就不要反复复制

block table 记录逻辑 KV block 到物理 block 的映射。通用实现为了保证 DCU 状态
最新，会在输入准备时提交 block table。但在单请求稳态 decode 中，许多连续步骤
都在同一个已分配 block 内，页表内容没有变化。

为此，我们在
vllm/v1/worker/block_table.py
中加入 dirty tracking：

defappend_row(...):self._block_table_dirty=True...defmove_row(...):self._block_table_dirty=True...defswap_row(...):self._block_table_dirty=True...defcommit_block_table(num_reqs):ifnotself._block_table_dirty:returnself.block_table.copy_to_gpu(num_reqs)self._block_table_dirty=False

这不是“永远不复制 block table”，而是把提交语义从“每轮都复制”改为“有变化才
复制”：

• 新增 cache block：标脏并复制；
• 请求行新增、移动或交换：标脏并复制；
• 稳态 decode 且映射不变：直接跳过；
• 清空表：CPU/DCU 同时清零并恢复干净状态。

dirty flag 的价值不在于复杂，而在于把控制面已经知道的事实告诉数据面：内容没变，
就不必再做一次看似安全、实际冗余的 H2D。

7. 实现三：快路径必须“窄”，不能拿正确性换速度

专用优化最危险的做法，是看到 batch=1 就把所有通用逻辑直接删除。真实服务里可能
发生请求替换、抢占、prefill/decode 混合或状态重排。一旦错误地进入快路径，写错的
不是一个无关紧要的统计量，而是 KV cache 的物理位置，后续输出会整体失真。

因此，快路径在
vllm/v1/worker/gpu_model_runner.py
中设置了两层门控。

7.1 初始化能力门控

只有同时满足以下条件，才创建专用 preparer：

• 海光 BW1000 DCU（gfx936 架构）；
• Qwen3_5ForConditionalGeneration；
• 开启 vLLM 异步调度；
• tensor parallel、pipeline parallel 和 decode context parallel 均为 1；
• 没有 speculative decoding；
• 没有 LoRA；
• 不使用 XD-RoPE。

这一步回答“当前服务配置有没有资格使用快路径”。

7.2 每轮状态门控

即使服务配置满足要求，每个 engine step 仍要检查：

• 当前只有一个请求；
• 本轮恰好只调度一个 token；
• 没有 encoder 输入和 speculative token；
• 上一轮 sampled token 存在于同一设备，形状和类型正确；
• 当前请求在上一轮仍位于 index 0；
• 当前已经进入 decode，而不是仍在处理 prompt；
• 请求持有的 token 数与当前位置连续；
• CPU 侧上一轮seq_len与当前 position 对得上。

最后一项尤其重要。它能识别抢占、请求替换或非单步推进后留下的陈旧状态。任意一项
不满足，系统都会完整回退原始 vLLM 输入准备路径。

用伪代码表示就是：

ifservice_supports_fast_pathandcurrent_step_is_strict_single_token_decode:prepare_on_gpu()else:prepare_with_original_vllm_path()

这个优化没有修改 scheduler，没有改变模型数学计算，没有调整采样参数，也没有要求
用户增加新的启动环境变量。它只是给一个能被严格证明等价的高频场景增加专用实现。

8. 为什么 DCU 写完了，还要维护 CPU 镜像？

这是实现中一个容易被忽略的点。

positions、query_start_loc和seq_lens等缓冲区通常同时存在 CPU 镜像和 DCU
张量。DCU kernel 可以直接写模型执行需要的张量，但后续 attention metadata builder
仍可能读取 CPU 镜像来决定 kernel 参数或构造 Python 对象。

如果只更新 DCU，不更新 CPU，当前前向也许能跑，下一处读取 CPU 元数据时却可能
得到旧值。这样的 bug 很隐蔽，常表现为“短输入正常、长输入偶发错”或“请求切换后
才出错”。

因此，专用路径仍在 CPU 侧维护必要的镜像：

query_start_loc_cpu=[0,1,1,...]seq_lens_cpu=[position+1,0,0,...]discard_request_mask_cpu=[False,False,...]

区别在于，这些 CPU 值不再通过多次copy_to_gpu()驱动本轮模型输入；DCU 目标
张量由融合 kernel 一次写完。换句话说，CPU 镜像用于保持软件状态一致，DCU kernel
负责热路径数据准备。

9. 如何证明新路径真的等价？

性能优化不能只看“能跑”。尤其是 slot mapping，一旦错一个 block，模型可能仍能
输出文字，但答案已经不可信。

我们为此编写了独立组件基准：

• test_scripts/bench_v1_decode_input_prep.py
• test_scripts/run_bench_v1_decode_input_prep.sh

基准默认模拟 16K-32K 长上下文中的一个位置：

9.1 逐项做整数精确校验

基准先运行旧路径，再运行新路径，并同步设备后检查：

• input_ids中的 token id；
• 三路 M-RoPE position；
• 完整query_start_loc；
• seq_lens及填充区；
• discard_request_mask；
• 两个 KV cache group 的 slot mapping。

这些值都是离散元数据，因此没有“允许一点浮点误差”的说法，必须逐项完全一致。

9.2 使用 A/B/A，而不是只跑一次 A/B

这类组件只有零点几毫秒，很容易受到频率、温度和后台任务干扰。测试顺序采用：

A：旧路径第一次 B：DCU 融合路径 A：旧路径第二次

如果前后两个 A 差异很大，就说明环境漂移明显，不能轻易把中间 B 的变化归因于代码。
本次测试中两个 A 的漂移约为 0.25%，结果较稳定。

10. 实测结果：0.259 ms 降到 0.101 ms

海光 BW1000 DCU（gfx936 架构）上的组件 A/B/A 结果如下：

取更快的旧路径 A2 作为基线：

单步节省 = 0.258436 - 0.101048 = 0.157388 ms 加速比 = 0.258436 / 0.101048 = 2.5576x 组件延迟降低比例 = 1 - 0.101048 / 0.258436 = 60.9%

10.1 这不代表整机吞吐提升 2.56 倍

必须强调：2.56 倍是输入准备组件的加速，不是整个 Qwen3.5 推理的加速。
完整 TPOT 还包括 Attention、GEMV、GDN、MLP、RMSNorm、lm_head 和采样等大量工作。

这项优化更接近“每个稳态 decode step 固定省下约 0.157 ms”。理论累计值可以这样
估算：

之所以是“生成 token 数减一”，是因为第一个输出 token 通常由 prefill 末端产生，
最后生成后也不一定再进入下一轮。这张表只是根据组件差值计算的上界估计，真实端到端
收益还受输出长度、DCU 频率和其它流水线开销影响。

举例来说，如果某系统原始 TPOT 为 45 ms，那么单步减少 0.157 ms 对 TPOT 的直接
贡献约为 0.35%，而不是 60.9%。这个数字看起来不夸张，但它是每个 token 都能获得、
不改变模型精度的固定收益。多项这类固定成本优化叠加后，才会形成可观察的吞吐提升。

10.2 这项优化对先导杯成绩意味着什么？

先导杯统计的是三档测试集上的实际输出吞吐，不是单个 kernel 的理论算力。单 Token
输入准备处在每轮 decode 的必经路径，因此，只要请求连续生成，它的固定节省会在
4K-8K、8K-16K 和 16K-32K 三档中重复累积。它更直接影响 TPOT 和端到端生成时间，
而不是由长 prompt 计算主导的 TTFT。

这也解释了为什么组件加速 2.56 倍，最终比赛吞吐通常只会小幅提高：输入准备只是
完整 decode 的一部分。对于 16K-32K 样例，prefill 和长上下文 attention 仍占据大量
时间；对于 4K-8K 和 8K-16K，decode 固定开销的相对占比可能更高。这项优化的价值
在于收益稳定、三档都能使用，并且不改变输出长度来“做大分子”。

正式判断是否提高比赛得分时，仍应使用同一批样例进行端到端 A/B，并同时比较：

• 三档输出吞吐与 TTFT、TPOT；
• Total generated tokens和每条请求的输出长度；
• 生成文本是否出现差异；
• 官方 accuracy、SLA 与精度扣分。

只有吞吐提高且生成行为、SLA 和精度检查都通过，才能把收益归因于真实执行效率，
而不是输出 Token 数变化或数值误差。

11. 如何复现实验？

在已配置好 DTK/ROCm 环境的测试机上进入项目根目录，执行：

bash./test_scripts/run_bench_v1_decode_input_prep.sh\--position24576\--max-model-len32768\--max-num-reqs128\--max-num-tokens4096\--block-sizes16,16\--warmup100\--iters2000\--output./test_scripts/v1_decode_input_prep.csv

包装脚本会执行三件重要的事：

1. 从当前 worktree 重新构建并安装 vLLM wheel；
2. 比较关键源码的 SHA-256，确认测试的安装包与当前代码一致；
3. 从/tmp启动 Python，避免项目源码目录意外遮蔽已安装 wheel。

参考输出：

=== Running V1 single-token input-prep A/B/A benchmark === device=BW block_sizes=[16, 16] baseline_before/candidate/baseline_after= 0.259076/0.101048/0.258436 ms speedup=2.5576x saved=0.157388 ms

在 ROCm 版 PyTorch 中，设备 API 仍沿用torch.cuda命名空间，因此测试脚本里的
torch.cuda.synchronize()并不表示测试运行在 NVIDIA 显卡上；本文实测设备是
海光 BW1000 DCU。

11.1 如何确认完整服务命中了快路径？

服务日志中应先后出现：

[dcu-v1-input-gpu] reserved gfx936 Qwen3.5 decode input state [dcu-v1-input-gpu] dispatched Qwen3.5 single-token decode input

第一行表示当前模型和服务配置具备能力，第二行表示至少一个实际 engine step 通过了
运行时状态门控。只有第一行而没有第二行，说明工作区已预留，但真实请求没有满足
单请求、单 token、状态连续等条件。

12. 哪些场景适合，哪些场景不适合？

适合

• batch size 为 1 的在线生成；
• 请求在连续 decode step 之间保持不变；
• 上一轮 sampled Token 已保留在 DCU；
• 长输出或固定开销占比较高的模型；
• 模型算子已经较快，希望继续压缩 TPOT 尾部开销。

不适合直接套用

• 高并发、多请求动态 batch；
• speculative decoding；
• 请求频繁抢占、替换和重排；
• LoRA、多级并行或 context parallel；
• encoder-decoder 混合输入；
• 尚未证明 metadata 形状与单 token 假设一致的其它模型。

不适合并不意味着永远不能优化，而是需要为对应场景重新设计数据布局和门控，不能
直接删除通用路径。

13. 这个优化真正值得复用的三条经验

13.1 Profile 不要只看“大核”

算子 profile 很容易把注意力吸引到耗时最长的 kernel，但端到端延迟还包含 kernel
之间的空隙、CPU 准备、小拷贝和 launch。模型越快、batch 越小，这些固定成本越值得
关注。

13.2 先问“数据现在在哪里”

本次优化的突破口不是发明更快的数学公式，而是发现 sampled Token 已经在 DCU。
只要下一轮所需状态也能由 DCU 上的 Token、position 和 block table 推导，就应该
优先建立 DCU 内部闭环，而不是继续优化 CPU 到 DCU 的搬运姿势。

13.3 专用快路径的质量取决于回退设计

“牺牲通用性”不应等于破坏其它场景。更稳妥的方式是：

• 保留原始通用实现；
• 用严格条件证明专用路径成立；
• 在状态不连续时立即回退；
• 独立验证所有离散元数据；
• 最后再用完整推理检查生成 token、文本和精度。

快路径可以很激进，门控必须很保守。

14. 总结

这项 2026 先导杯优化没有修改 Attention 或 GEMM，也没有通过量化或改变输出 Token
数换速度。它做的是一件更朴素的事：把已经处于 DCU 闭环中的 decode 状态继续留在
DCU，把多次小拷贝和 CPU slot mapping 合并为一次 Triton launch，并通过 block
table dirty tracking 跳过没有意义的重复提交。

最终，海光 BW1000 DCU（gfx936 架构）上该组件从约 0.259 ms 降到 0.101 ms，
组件加速约 2.56 倍，
每个稳态 decode token 固定减少约 0.157 ms。它不会单独带来数量级的整机吞吐提升，
但给出了一个很有价值的方向：

当大算子已经足够快时，下一步不一定是继续调 block size。沿着一枚 token 的完整
生命周期，找到那些反复跨越 CPU 和 DCU 的小数据，往往更容易获得稳定、可叠加且
不损失精度的收益。