Prhub

#51367 [Kernel] Use Murmur3 RNG for Gumbel sampling

原始 PR 作者 freyfwt 合并时间 2026-09-16 14:52 文件变更 5 提交数 7 评论 30 代码增减 +216 / -26

执行摘要

Gumbel 采样改用 Murmur3,修复尾部概率偏差

PR body 明确写了动机:Replace the Gumbel sampler's Philox-based uniform generation ... The goal is to preserve the target sampling distribution while improving RNG resolution and avoiding backend-specific math paths,并关联 RFC #54143,同时区分于保留 Philox、只多取输出位的 #47524。作者在评论中进一步给出量化证据:24 位截断路径在 vocab=129,280、logit gap 18/20 时尾部质量 z 分数为 -5.89 / -5.16,改用完整 32 位源后收敛到 -1.18 / -1.09,确认截断就是偏差主因;此外 Ascend A3 上 Philox 采样 kernel 比 Murmur3 慢 26-89 倍,属于典型的后端专属数学路径成本。

值得精读,尤其是三处「跨后端数值可移植性」的设计:murmur3_uniform32 的 split-16 精确构造(用两个可精确表示的半部绕开后端 uint32→fp32 转换)、murmur3_uniform64 用 domain=0x9E3779B9 拼 53 位并显式打开右端点、_log1p_neg_stable 的分段多项式替代 libdevice log1p。阅读时建议重点验证四点:h = (seed ^ seed) 与 h ^ 16 是否为有意为之;u → 1 退化分支的噪声下限是否可接受;删除入口测试的「交错完成」断言是否应由别的测试补位;性能数字是否会在最终 head 上复现。对维护者的动作建议:在 release note 中显式提示 seeded 非 greedy 输出序列变更,并补一条 u → 1 端点的数值回归测试。

讨论亮点

讨论区 30 条里绝大多数是 CI/bot 编排(/ci run、mergify 冲突提示、pre-commit 失败提示),真正的技术交锋集中在作者自证的几条长评论上:

作者 freyfwt:Changed fp32 Murmur uniform from 24-bit truncation to full 32-bit source resolution: u = (random32 + 0.5) / 2^32, reconstructed through two exact 16-bit int→fp32 conversions for backend portability. Verified split-16 reconstruction against correctly rounded fp32 on 2M random uint32 values: 0 mismatches.

作者 freyfwt:For vocab=129,280 and 100k draws, tail-mass z-scores were -1.18 at logit gap 18 and -1.09 at gap 20. The old 24-bit path gave -5.89 and -5.16 respectively, confirming the truncation was the material bias source.

作者 freyfwt(性能结论克制):The host was shared with an unrelated 8-GPU workload ... so these results are used as a no-material-regression check, not as a clean speedup claim.

作者 freyfwt(自陈覆盖缺口):the available runtime image predates the is_drafting call-site update. Both variants therefore used the same one-line is_drafting=False harness adapter ... drafting/speculative behavior is outside its scope.

评审侧只有两条机器/单人信号:claude[bot] 在 fork 场景下先说明自动评审被禁用,随后给出「Code review found no issues / No high-confidence issues detected in this change」,Isotr0py 给出 APPROVED。review_comments_count 为 0,意味着没有逐行 diff 讨论、没有张量统计之外的第三方复核,PR 的正当性几乎完全建立在作者自测数据上。

实现拆解

  1. 入口换 RNG 并重排噪声流隔离:vllm/v1/worker/gpu/sample/gumbel.py 的 gumbel_noised_argmax 删除 gumbel_seed = tl.randint(seed, pos) 与 tl_rand32/tl_rand64 调用,改为按 (seed, pos, keys) 直接调用 murmur3_uniform32/64;draft 与 target 的隔离从「偏移 offset」改为「对 pos 加 _DRAFT_NOISE_SALT」,因为 seed 与 pos 现在都直接进入 hash,salt 位置即可保证两条噪声流不一致。同时移除 tldevice 导入(不再需要后端 log1p),_TL_RAND_MIN 的注释改写为仅供 rejection sampling 使用。
  2. 新增 Murmur3 哈希族:_murmur3_rotl32(无内置 rotate,用移位或拼回)、_murmur3_mix(body:乘 0xCC9E2D51、rotl 15、乘 0x1B873593、异或、rotl 13、乘 5 加 0xE6546B64)、_murmur3_fmix32(finalizer)以及对外入口 murmur3_hash32(seed, pos, offset, domain=0)。murmur3_hash32 先把 64 位 seed 拆成高低 32 位混合、再混 pos、最后混词表内 offset,并用 domain 常量在同一坐标上再开一条独立流。
  3. 两种均匀数构造:murmur3_uniform32 用「拆 16 位再转 fp32」的方式实现 (random32 + 0.5) * 2**-32,两个半部都能被 fp32 精确表示,因此不依赖后端 uint32→fp32 转换也能用满 32 位分辨率;murmur3_uniform64 用 domain=0x9E3779B9 取第二个 hash,拼出 53 位有效位后交给 _uniform64_from_random53,并用 tl.where(uniform == 1.0, uniform - 1.1102230246251565e-16, uniform) 显式打开右端点,避免最大中点舍入到 1.0 后 log(1-u) 变成 -inf。
  4. fp32 尾部对数可移植化:新增 _log1p_neg_stable,在 u < 0.25(决定 argmax 胜出的大 Gumbel 噪声尾部)走 8 次 Horner 多项式(绝对误差 < 6e-7),其余区间退化为 tl.log(tl.maximum(1 - u, 5.96e-08)),用分段近似替代 libdevice log1p,从而避免后端专属数学库差异。
  5. 测试与配套改动:tests/v1/worker/test_gpu_gumbel_sample.py 新增最小 Triton kernel _murmur3_known_answer_kernel 与 test_murmur3_known_answers(锁定 6 组已知答案,含 int64 极值 seed 与 1<<30 附近的 pos)、_uniform64_endpoint_kernel 与 test_uniform64_excludes_endpoints(断言结果为 2**-54 与 1 - 2**-53),并把 test_drafting_uses_a_separate_noise_stream 的 docstring 从「salts the Philox offset」更新为「salts the RNG position」。三处输出敏感的入口测试被去随机化:test_include_reasoning.py 四个用例加 temperature=0 并把 max_tokens 从 200 提到 512;test_struct_output_generate.py 中结构化输出用例 temperature=1.0 改 0;test_completion.py 删除 test_parallel_no_streaming 与 test_parallel_streaming 里「子请求必须不同长度」的 output_token_lengths / chunk_lengths 断言。
文件 类型 关键动作
vllm/v1/worker/gpu/sample/gumbel.py 源码 新增 Murmur3 RNG 族与分段 log1p,替换采样内核的 RNG 调用点
tests/v1/worker/test_gpu_gumbel_sample.py 测试 已知答案 + fp64 端点回归,更新噪声流语义的文档字符串
tests/entrypoints/openai/chat_completion/test_include_reasoning.py 测试 改为 temperature=0、max_tokens=512,去掉对采样流的依赖
tests/entrypoints/openai/completion/test_completion.py 测试 删除并行补全交错完成的长度断言
tests/entrypoints/llm/test_struct_output_generate.py 测试 结构化输出用例改贪心,避免 RNG 变更导致 flaky
文件 模块 状态 重要度
vllm/v1/worker/gpu/sample/gumbel.py 采样内核 modified 8.56
tests/v1/worker/test_gpu_gumbel_sample.py 采样测试 modified 7.38
tests/entrypoints/openai/completion/test_completion.py 补全接口测试 modified 4.53
tests/entrypoints/openai/chat_completion/test_include_reasoning.py 推理内容测试 modified 4.33
tests/entrypoints/llm/test_struct_output_generate.py 结构化输出 modified 3.62

关键符号

gumbel_noised_argmax murmur3_hash32 murmur3_uniform32 murmur3_uniform64 _uniform64_from_random53 _log1p_neg_stable _murmur3_mix _murmur3_fmix32 test_murmur3_known_answers test_uniform64_excludes_endpoints

关键源码片段

tests/v1/worker/test_gpu_gumbel_sample.py test-coverage

为新的 Murmur3 路径补已知答案与 fp64 端点回归,并更新 draft 噪声流语义的断言说明,是本次数值正确性的主要保护网

# tests/v1/worker/test_gpu_gumbel_sample.py:新增的已知答案与端点测试
​
@triton.jit
def _murmur3_known_answer_kernel(
    seeds_ptr, positions_ptr, offsets_ptr,
    hashes_ptr, domain_hashes_ptr,
    numel: tl.constexpr, BLOCK_SIZE: tl.constexpr,
):
    # 被测函数无法直接在 Python 侧调用,这里用最小 kernel 把设备端真实实现
    # 搬出来,保证测的是 Triton 里的 murmur3_hash32 而不是 Python 复刻
    idx = tl.arange(0, BLOCK_SIZE)
    mask = idx < numel
    seed = tl.load(seeds_ptr + idx, mask=mask)
    pos = tl.load(positions_ptr + idx, mask=mask)
    offset = tl.load(offsets_ptr + idx, mask=mask)
    hash32 = murmur3_hash32(seed, pos, offset).to(tl.int64)
    domain_hash32 = murmur3_hash32(seed, pos, offset, domain=0x9E3779B9).to(tl.int64)
    tl.store(hashes_ptr + idx, hash32, mask=mask)
    tl.store(domain_hashes_ptr + idx, domain_hash32, mask=mask)
​
​
def test_murmur3_known_answers():
    # 覆盖 0 / 1 / -1 / 跨 32 位 seed / int64 两端极值,pin 住跨后端 hash 一致性
    seeds = torch.tensor(
        [0, 1, -1, 0x123456789ABCDEF, -(1 << 63), (1 << 63) - 1],
        dtype=torch.int64, device=DEVICE,
    )
    positions = torch.tensor(
        [0, 2, 17, 123456789, 1 << 30, (1 << 30) + 7],
        dtype=torch.int64, device=DEVICE,
    )
    offsets = torch.tensor(
        [0, 3, 1023, 199999, 4096, 0xFFFFFFFF],
        dtype=torch.int64, device=DEVICE,
    )
    hashes = torch.empty_like(seeds)
    domain_hashes = torch.empty_like(seeds)
​
    _murmur3_known_answer_kernel[(1,)](
        seeds, positions, offsets, hashes, domain_hashes,
        numel=seeds.numel(), BLOCK_SIZE=8,
    )
​
    assert hashes.tolist() == [
        0x8134CDF8, 0x23B53F91, 0xC8C13C14,
        0x5DDE974F, 0xF1A9355B, 0xD952F9AA,
    ]
    # domain=0x9E3779B9 的独立流同样锁答案,防止 fp64 路径悄悄退化
    assert domain_hashes.tolist() == [
        0xD28AD962, 0xE3AAF818, 0xD223C51C,
        0x0547D2FE, 0x3AED160F, 0x49E8EA7F,
    ]
​
​
def test_uniform64_excludes_endpoints():
    # random53 = 0 与 2**53 - 1 对应最靠近两端的样本:期望 2**-54 与
    # 1 - 2**-53,即区间必须保持开区间,否则 log(1 - u) 会出现 -inf
    random53 = torch.tensor([0, (1 << 53) - 1], dtype=torch.int64, device=DEVICE)
    uniform = torch.empty(2, dtype=torch.float64, device=DEVICE)
​
    _uniform64_endpoint_kernel[(1,)](
        random53, uniform, numel=random53.numel(), BLOCK_SIZE=2,
    )
​
    assert torch.isfinite(uniform).all()
    assert ((uniform > 0.0) & (uniform < 1.0)).all()
    assert uniform.tolist() == [2.0**-54, 1.0 - 2.0**-53]

评论区精华

fp32 均匀数从 24 位截断改为完整 32 位源分辨率 正确性

作者自查发现初始实现只用 24 位导致尾部概率偏差,改用以两个精确 16 位半部重建 `(random32 + 0.5) / 2**32`,并用 2,000,000 个随机 uint32 验证与正确舍入 fp32 完全一致(0 mismatch);Monte Carlo 下 vocab=129,280 的尾部质量 z 分数从 -5.89 / -5.16 收敛到 -1.18 / -1.09。

结论:确认 24 位截断是偏差主因并已修复;旧路径在 far-tail 上系统性低估,新路径回到统计噪声范围内。 · 已解决

fp64 均匀数端点必须保持开区间 正确性

最大中点 `(2**53 - 1 + 0.5) * 2**-53` 在 fp64 上舍入到 1.0,会让 `tl.log(1 - u)` 变成 -inf;修复用 `tl.where(uniform == 1.0, uniform - ulp, uniform)` 显式回退一个 ulp,并新增 `test_uniform64_excludes_endpoints` 断言输出为 `2**-54` 与 `1 - 2**-53`。

结论:端点已闭合为开区间并有回归测试锁定;同时避免依赖 Triton 可能按 fp32 物化的近似 1 字面量。 · 已解决

draft / target 噪声流隔离方式从 offset salt 改为 pos salt 设计

RNG 改为按 (seed, pos, keys) 直接 hash 后,原先对 gumbel_seed 偏移的 salt 机制失效;改为 `pos = pos + _DRAFT_NOISE_SALT`,并同步把测试 docstring 从 salts the Philox offset 改为 salts the RNG position。作者在 merge upstream 的 commit 里专门说明 Preserve the draft noise-stream salt while using Murmur3 directly with the salted position。

结论:隔离语义保持等价,依赖「pos 为 int64 且远不接近 2**30」的既有假设;spec-decode 验证的概率比语义未变。 · 已解决

输出敏感的入口测试去随机化与「交错完成」断言删除 测试

RNG 流变化让推理内容、结构化输出、并行补全三组 HTTP 层测试变得 flaky;作者改 `temperature=0`、放宽 `max_tokens`,并删除 `assert len(output_token_lengths) > 1` / `assert len(chunk_lengths) > 1`,理由是 Staggered child completion is already covered deterministically by the `ParentRequest` unit tests。

结论:去 flaky 的目标达成,但 HTTP/ 流式层的交错完成覆盖净减少;引擎层单测与入口层时序验证不等价,这一取舍没有第三方复核。 · resolved-with-tradeoff

自动化评审与人工 approval 的实际深度 question

claude[bot] 先因 fork 禁用自动评审,随后给出 Code review found no issues;`Isotr0py` 给出 APPROVED 但未留评论,review_comments_count 为 0,30 条评论几乎全是 CI 触发与 bot 提示。

结论:缺少逐行复核与独立数值复算,合入依据主要是作者自测数据与 CI 绿灯(Buildkite #86309 与后续重跑)。 · 待处理

风险与影响

  1. 确定性输出变更(兼容性):body 已声明 seeded 非 greedy 请求会产生不同的确定性 token 序列;作者实测 base/head 只有 1/16 首轮补全完全一致。任何依赖 seed 复现的输出基线(用户侧回归脚本、内部 golden 输出、缓存/重放对比)都会失效,但 PR 声明 No documentation change is needed,对外可见的行为变更缺少 release note。
  2. RNG 统计强度的论证方式:Murmur3 是非密码学哈希,作者用 Pearson z 分数(mean -0.009,std 0.998,max |z| 3.90)、byte/bit/serial 检验与 Monte Carlo 尾部质量来论证,但这些是经验检验,不等价于 Philox 这类 counter-based PRNG 的已知 equidistribution 性质。另外 murmur3_hash32 里 h = (seed ^ seed).to(tl.uint32) 恒为 0(疑似「清零并转 uint32」的遗留写法),_murmur3_fmix32(h ^ 16) 把 16 当长度字节,两处都偏离标准 MurmurHash3 写法,需要确认是刻意设计还是可读性陷阱。
  3. fp32 对数近似的边界:_log1p_neg_stable 只在 value < 0.25 用多项式,其余走 tl.log(tl.maximum(1 - value, 5.96e-08));当 u 接近 1 时 fp32 相减本身损失精度并被 clamp,退化分支给出的噪声下限约为 log(5.96e-08) ≈ -16.7,属于对下尾的显式取舍,但目前没有测试直接覆盖 u → 1 端。
  4. 测试覆盖的净回退:test_completion.py 删除的「子请求在不同长度完成」断言是唯一在 HTTP/流式层观察交错完成的位置,作者以 ParentRequest 单元测试已确定性覆盖为由删除;单请求引擎层测试与端到端流式时序不是同一层,覆盖确实变薄。
  5. 性能结论的证据强度:Ascend A3 的 26-89 倍加速微基准早于最终 head(作者自述 predate the final fp32-uniform, logarithm, and fp64-endpoint changes),模型级 A/B 又跑在与 8 卡无关负载共享的主机上(输出吞吐 CV 11.5%/7.4%),且 drafting 路径未端到端覆盖。当前 head 的 kernel 延迟尚未复测,性能收益不能直接作为既成事实引用。
  6. 安全面:seed 由调用方提供、Murmur3 可逆性远高于 Philox,理论上降低了「seed 泄露即可预测采样」的门槛,但采样 RNG 本身不承担安全职责,风险等级低。

影响面集中在 Model Runner V2 的非 greedy 采样路径:所有使用 seed 的 temperature > 0 请求在升级后得到不同(但分布等价)的输出,可复现性基线与下游 golden 对比需要重建;未设置 seed 的用户无感知。收益侧,非 NVIDIA 后端(Ascend A3 采样 kernel 从 160 ms 级降到 1.8 ms 级)与结构化输出/推理解析等多请求场景最明显,RTX 5090 上模型级 A/B 显示输出吞吐 +4.15%、mean TPOT -5.15%,可读作无实质回归。工程侧影响包括:draft/target 噪声流隔离方式从 offset salt 改为 pos salt,spec-decode 的验证语义保持不变但依赖「pos 永不接近 2**30」的既有假设;入口测试(推理内容、结构化输出、并行补全)因改贪心或删断言而更稳定,但代价是少了一层对采样流的端到端约束。greedy 路径、logits cache 行为、公开 API 与 kernel launch 结构均不变,说明爆炸半径被有意限制在单个采样文件。

确定性输出变更 非密码学哈希充当 RNG 测试覆盖部分回退 性能数据待最终 head 复测 drafting 路径未端到端覆盖

关联 Issue

未识别关联 Issue

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

完整报告

参与讨论