执行摘要

• 一句话：修复 ROCm DeepSeek V4 indexer 数值与测试覆盖

• 推荐动作：建议

• 值得精读，尤其是 fused_indexer_q.py 中的显式 FMA 处理和测试条件跳过逻辑，展示跨平台数值一致性的处理方式。

• 关注后续是否会将相似修复应用到其他内核，以及 MXFP4 在非 NVIDIA 平台的支持进展。

功能与动机

PR #40860 遗留了 Triton RoPE 表达式与 HIP 参考实现不同的收缩行为，PR #41015 添加了基于 NVIDIA 专用转换路径的 MXFP4 缓存用例，导致 ROCm 上出现舍入不匹配和 PTX 编译错误（见 Buildkite 输出）。本 PR 修复受支持的 FP8 路径，并准确限制不支持的缓存特性，同时不减少 ROCm 的 FP8 测试覆盖。

实现拆解

实现拆解

1. 修改融合内核：在 vllm/models/deepseek_v4/common/ops/fused_indexer_q.py 的 _fused_indexer_q_rope_quant_kernel 中新增 USE_EXPLICIT_FMA 常量，并在 RoPE 计算中使用显式 FMA 操作以匹配 HIP 参考行为。该常量在调用 fused_indexer_q_rope_quant 时根据 current_platform.is_rocm() 动态设置。

2. 调整测试参数化：在 tests/kernels/test_fused_indexer_q_rope_quant.py 和 tests/kernels/test_compressor_kv_cache.py 中，将 use_fp4 参数化改为条件跳过，仅当平台为 CUDA 且设备能力为 SM100 时运行 MXFP4 用例，避免在 ROCm 上触发不支持的路径。同时引入 current_platform 导入。

3. 保持位精确断言：未放宽数值容差，确保 FP8 断言保持位精确。

4. 测试与 CI：修改了两个测试文件，并确认无性能回归（延迟无显著变化）。

关键文件：

• vllm/models/deepseek_v4/common/ops/fused_indexer_q.py（模块 融合内核；类别 source；类型 core-logic；符号 _fused_indexer_q_rope_quant_kernel, fused_indexer_q_rope_quant）: 核心内核修改，新增显式 FMA 以匹配 HIP RoPE 行为，影响所有平台的数值输出。
• tests/kernels/test_fused_indexer_q_rope_quant.py（模块 测试；类别 test；类型 test-coverage）: 调整 use_fp4 参数化，条件跳过 MXFP4 用例，确保 ROCm 上 FP8 覆盖不受影响。
• tests/kernels/test_compressor_kv_cache.py（模块 测试；类别 test；类型 test-coverage）: 同样调整 use_fp4 参数化，限制 MXFP4 到支持的平台。

关键符号：_fused_indexer_q_rope_quant_kernel, fused_indexer_q_rope_quant

关键源码片段

vllm/models/deepseek_v4/common/ops/fused_indexer_q.py

核心内核修改，新增显式 FMA 以匹配 HIP RoPE 行为，影响所有平台的数值输出。

# vllm/models/deepseek_v4/common/ops/fused_indexer_q.py
# kernel 参数新增 USE_EXPLICIT_FMA，用于控制 RoPE 计算的收缩方式
# 在 ROCm 上启用，匹配 HIP 参考实现，避免数值偏差
def _fused_indexer_q_rope_quant_kernel(
    ...
    USE_FNUZ: tl.constexpr = False,
    USE_EXPLICIT_FMA: tl.constexpr = False, # 新增：显式 FMA 开关
):
    ...
    # RoPE 计算部分
    x_even = tl.load(rot_base + half_offset * 2).to(tl.float32)
    x_odd = tl.load(rot_base + half_offset * 2 + 1).to(tl.float32)
    if USE_EXPLICIT_FMA:
        # 匹配 HIP rotary_embedding 收缩行为，在 bf16 materialization 之前
        r_even = tl.fma(x_even, cos, -(x_odd * sin))
        r_odd = tl.fma(x_odd, cos, x_even * sin)
    else:
        # 默认路径，保持原实现
        r_even = x_even * cos - x_odd * sin
        r_odd = x_odd * cos + x_even * sin
​
# 外层函数根据平台设置 USE_EXPLICIT_FMA
def fused_indexer_q_rope_quant(...):
    ...
    USE_EXPLICIT_FMA=current_platform.is_rocm(), # 仅 ROCm 启用

评论区精华

Review 中仅包含 bot 的自动评论和 tjtanaa 的批准，无实质性技术讨论。

• 自动代码审查 (other): 未产生技术讨论。

风险与影响

• 风险：风险分析

• 内核数值风险：修改 fused_indexer_q.py 中的 RoPE 计算可能会影响数值行为，但已通过位精确断言确保与参考实现一致。

• 平台兼容性：USE_EXPLICIT_FMA 仅在 ROCm 上启用，CUDA 路径保持不变，降低回归风险。
• 测试覆盖：MXFP4 用例在非 SM100 平台被跳过，可能导致这些平台上的覆盖缺失，但该特性本身不受支持，属于合理限制。
• 性能风险：显式 FMA 可能影响 kernel 性能，但 PR 描述称无显著延迟回归。

• 影响：影响分析

• 用户影响：修复了 ROCm 用户使用 DeepSeek V4 indexer 时的潜在数值错误和编译错误，提升功能性。

• 系统影响：内核数值修正可能略微改变输出精度，但目标与参考一致。
• 团队影响：增强了 ROCm 测试覆盖的准确性，为未来平台特定调整提供参考模式。
• 风险标记：平台特定代码, 测试覆盖受限, 内核数值变更

关联脉络

• PR #40860 引入 Triton RoPE 表达式，导致收缩行为与 HIP 不一致: 本 PR 修复了该 PR 遗留的数值偏差问题。
• PR #41015 添加 MXFP4 缓存用例，依赖 NVIDIA 专用转换路径: 本 PR 限制这些用例到 SM100，修复 ROCm 上的编译错误。