Prhub

#47383 [Bugfix][Model Runner V2][Spec Decode] Fix int32 offset overflow in block verification kernels

原始 PR 作者 WoosukKwon 合并时间 2026-07-02 22:32 文件变更 2 提交数 1 评论 0 代码增减 +148 / -4

执行摘要

修复 V2 投机解码 block 验证 kernel int32 索引溢出

当词汇表规模约为 155k 时,logit_idx * vocab 等乘法在 int32 范围内溢出,导致 block 验证内核读取越界 logits,进而破坏接受决策或引起非法内存访问。此问题在审计 V2 spec-decode 采样路径时发现(关联 #47239),虽非该 issue 直接原因,但属于同一类 int32 索引溢出 bug。

值得精读,展示 Triton 内核中因混合 int32/int64 导致的索引溢出典型修复模式,测试设计(参数化溢出场景)值得学习。

讨论亮点

PR 无实质性 review 讨论,仅由 claude[bot] 触发审查提示和 TheEpicDolphin 批准。

实现拆解

  1. _compute_cumulative_log_p_kernel 中,将 req_state_idxstart_idx 的加载结果通过 .to(tl.int64) 提升为 int64。
  2. _compute_local_residual_mass_kernel 中,将 logit_idx(来自 tl.program_id(0))和 req_state_idx 也通过 .to(tl.int64) 提升。
  3. 新增测试文件 tests/v1/worker/test_gpu_rejection_sampler_i64.py,构造两种溢出场景(目标端索引 15000、草稿端索引 8000),验证修复后的输出与无溢出参考一致。
  4. 每个测试用例分配约 5 GiB GPU 内存,并在 CI 中仅对 CUDA 运行。
文件 模块 状态 重要度
vllm/v1/worker/gpu/spec_decode/rejection_sampler_utils.py 投机解码 modified 6.44
tests/v1/worker/test_gpu_rejection_sampler_i64.py 测试 added 6.83

关键符号

_compute_cumulative_log_p_kernel _compute_local_residual_mass_kernel

分析完成后,这里会展示 LLM 生成的相对完整源码片段和详细注释。

评论区精华

无实质性 Review 讨论 other

PR 无 review 评论,仅 claude[bot] 触发审查提示和 TheEpicDolphin 批准合并

结论:直接合并 · 已解决

风险与影响

风险较低,修复直接且被测试覆盖。但需注意其他未检查的 Triton 内核是否也存在类似问题。性能影响可忽略,int64 加载在 GPU 上通常无开销。

影响使用 rejection_sample_method='block' 的投机解码用户,尤其大词汇表模型(如 GLM)。修复后输出正确,避免崩溃或静默错误。对其他路径无影响。

int32 溢出 非默认路径 GPU 内核修复

关联 Issue

未识别关联 Issue

当前没有检测到明确关联的 Issue 链接,后续同步到相关引用后会出现在这里。

完整报告

参与讨论