2026 先导杯实战别只盯着大算子把 Qwen3.5 vLLM 单 Token 输入准备搬上 DCU在海光 BW1000 DCUgfx936 架构上我们把 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 DCUgfx936 架构运行 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以及 batchscheduler输出文本和生成 Token 数需要与基线一起检查避免把数值漂移误当成性能提升不满足专用条件时必须回退 vLLM 原始路径。固定 batch1 一方面减少了通用动态批处理带来的优化空间另一方面也让执行形态非常稳定进入稳态 decode 后每轮通常只有一个请求和一个新 Token。本文正是利用这个确定性把通用输入准备中的重复工作变成 BW1000 上的单 Token 专用快路径没有触碰比赛禁止修改的调度参数。1. 为什么要优化一个看起来只有零点几毫秒的步骤做大模型推理优化时我们通常先盯着 GEMM、Attention、RMSNorm 这些“大算子”。这很合理因为它们承担了绝大部分浮点计算。但当模型算子已经经过多轮优化尤其是在batch size 为 1、每轮只生成一个 token的在线推理场景中另一类开销会逐渐显眼CPU 生成位置、序列长度等元数据CPU 根据 KV cache 页表计算 slot mapping多个很小的 Host-to-DeviceH2D拷贝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 往返但通用输入准备路径仍要兼容许多复杂场景同时调度多个请求每个请求调度不同数量的 tokenchunked prefillspeculative 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 Tokenposition[0]就是当前已计算 token 数query_start_loc固定为[0, 1, 1, ...]seq_lens固定为[position 1, 0, 0, ...]当前请求应被采样因此discard_request_mask全为 falselogits_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完成的元数据写入不要拆成多次小拷贝。新路径做了四件事直接读取 DCU 上一轮的sampled_token_ids用一次 Triton kernel launch 写完输入和 attention 元数据在 DCU 上读取 block table 并计算每个 KV cache group 的 slot给 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_idtl.program_id(0)# 公共元数据只写一次ifgroup_id0:token_idtl.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,[position1,0,0,...])tl.store(discard_request_mask_ptr,0)# 每组 KV cache 都有自己的 block size、页表和 slot 输出block_sizetl.load(block_sizesgroup_id)block_tableload_pointer(block_table_ptrsgroup_id)block_numbertl.load(block_tableposition//block_size)slotblock_number*block_sizeposition%block_size slot_mappingload_pointer(slot_mapping_ptrsgroup_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_ptrstorch.tensor([table.block_table.gpu.data_ptr()fortableinblock_tables],dtypetorch.uint64,devicedevice,)self.slot_mapping_ptrstorch.tensor([table.slot_mapping.gpu.data_ptr()fortableinblock_tables],dtypetorch.uint64,devicedevice,)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 trackingdefappend_row(...):self._block_table_dirtyTrue...defmove_row(...):self._block_table_dirtyTrue...defswap_row(...):self._block_table_dirtyTrue...defcommit_block_table(num_reqs):ifnotself._block_table_dirty:returnself.block_table.copy_to_gpu(num_reqs)self._block_table_dirtyFalse这不是“永远不复制 block table”而是把提交语义从“每轮都复制”改为“有变化才复制”新增 cache block标脏并复制请求行新增、移动或交换标脏并复制稳态 decode 且映射不变直接跳过清空表CPU/DCU 同时清零并恢复干净状态。dirty flag 的价值不在于复杂而在于把控制面已经知道的事实告诉数据面内容没变就不必再做一次看似安全、实际冗余的 H2D。7. 实现三快路径必须“窄”不能拿正确性换速度专用优化最危险的做法是看到 batch1 就把所有通用逻辑直接删除。真实服务里可能发生请求替换、抢占、prefill/decode 混合或状态重排。一旦错误地进入快路径写错的不是一个无关紧要的统计量而是 KV cache 的物理位置后续输出会整体失真。因此快路径在vllm/v1/worker/gpu_model_runner.py中设置了两层门控。7.1 初始化能力门控只有同时满足以下条件才创建专用 preparer海光 BW1000 DCUgfx936 架构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[position1,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.pytest_scripts/run_bench_v1_decode_input_prep.sh基准默认模拟 16K-32K 长上下文中的一个位置参数默认值目的position24576覆盖长上下文 decodemax model length32768对齐 32K 上下文max requests128保留生产缓冲区填充规模max batched tokens4096对齐运行缓冲区KV block sizes16, 16覆盖两个 KV cache groupwarmup100排除首次编译与预热影响iterations2000降低短组件计时抖动9.1 逐项做整数精确校验基准先运行旧路径再运行新路径并同步设备后检查input_ids中的 token id三路 M-RoPE position完整query_start_locseq_lens及填充区discard_request_mask两个 KV cache group 的 slot mapping。这些值都是离散元数据因此没有“允许一点浮点误差”的说法必须逐项完全一致。9.2 使用 A/B/A而不是只跑一次 A/B这类组件只有零点几毫秒很容易受到频率、温度和后台任务干扰。测试顺序采用A旧路径第一次 BDCU 融合路径 A旧路径第二次如果前后两个 A 差异很大就说明环境漂移明显不能轻易把中间 B 的变化归因于代码。本次测试中两个 A 的漂移约为 0.25%结果较稳定。10. 实测结果0.259 ms 降到 0.101 ms海光 BW1000 DCUgfx936 架构上的组件 A/B/A 结果如下路径平均耗时相对旧路径旧路径 A10.259076 ms1.00xDCU 融合路径 B0.101048 ms2.56x旧路径 A20.258436 ms1.00x取更快的旧路径 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 数约触发的后续 decode 准备次数理论累计节省10099约 15.6 ms256255约 40.1 ms512511约 80.4 ms1000999约 157.2 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-16Kdecode 固定开销的相对占比可能更高。这项优化的价值在于收益稳定、三档都能使用并且不改变输出长度来“做大分子”。正式判断是否提高比赛得分时仍应使用同一批样例进行端到端 A/B并同时比较三档输出吞吐与 TTFT、TPOTTotal 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包装脚本会执行三件重要的事从当前 worktree 重新构建并安装 vLLM wheel比较关键源码的 SHA-256确认测试的安装包与当前代码一致从/tmp启动 Python避免项目源码目录意外遮蔽已安装 wheel。参考输出 Running V1 single-token input-prep A/B/A benchmark deviceBW block_sizes[16, 16] baseline_before/candidate/baseline_after 0.259076/0.101048/0.258436 ms speedup2.5576x saved0.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 尾部开销。不适合直接套用高并发、多请求动态 batchspeculative decoding请求频繁抢占、替换和重排LoRA、多级并行或 context parallelencoder-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并通过 blocktable dirty tracking 跳过没有意义的重复提交。最终海光 BW1000 DCUgfx936 架构上该组件从约 0.259 ms 降到 0.101 ms组件加速约 2.56 倍每个稳态 decode token 固定减少约 0.157 ms。它不会单独带来数量级的整机吞吐提升但给出了一个很有价值的方向当大算子已经足够快时下一步不一定是继续调 block size。沿着一枚 token 的完整生命周期找到那些反复跨越 CPU 和 DCU 的小数据往往更容易获得稳定、可叠加且不损失精度的收益。