执行摘要

• 一句话：添加 _skip_rope_for_aiter_fused_mla 避免 gfx950 上的双重复旋转
• 推荐动作：值得精读。该 PR 展示了如何从临时环境变量方案演进为结构性修复，是设计决策的良好案例。特别关注 _skip_rope_for_aiter_fused_mla 方法的定义和它在 forward_absorb_prepare 中的插入点，理解条件判断的边界。

功能与动机

在 MI350X (gfx950) 上使用 Aiter 后端运行 DeepSeek-V3.2 等 MLA 模型时，GSM8K 准确率从 0.953 暴跌至 0.169。根本原因是旋转位置编码（RoPE）被计算了两次：一次在 forward_absorb_prepare 中，另一次在 fused kernel 中。

实现拆解

1. 添加跳过条件判断：在 forward_mla.py 的 forward_absorb_prepare 方法中，新增调用 _skip_rope_for_aiter_fused_mla()，并在原有 if 条件中增加对应的 and not 子句，使得当条件满足时跳过外层 rotary_emb 调用。
2. 新增 _skip_rope_for_aiter_fused_mla 方法：该方法返回 _use_aiter_gfx95 and self.current_attention_backend not in FORWARD_ABSORB_CORE_ATTENTION_BACKENDS，明确指示应在 gfx95 且 backend 不在融合列表中时跳过 RoPE，转由 forward_absorb_core 中的融合内核处理。
3. 清除环境变量方案：原提交尝试使用环境变量 SGLANG_AITER_FUSED_MLA_QK_ROPE 作为开关，经 review 讨论后废弃，改为上述结构性修复，避免了默认行为携带 bug 的问题。

关键文件：

• python/sglang/srt/models/deepseek_common/attention_forward_methods/forward_mla.py（模块 MLA层；类别 source；类型 core-logic；符号 _skip_rope_for_aiter_fused_mla）: 核心修改文件，添加了 _skip_rope_for_aiter_fused_mla 方法和调用条件，修复了 gfx950 上的双重复旋转问题。

关键符号：_skip_rope_for_aiter_fused_mla, forward_absorb_prepare

关键源码片段

python/sglang/srt/models/deepseek_common/attention_forward_methods/forward_mla.py

核心修改文件，添加了 _skip_rope_for_aiter_fused_mla 方法和调用条件，修复了 gfx950 上的双重复旋转问题。

# 在 forward_absorb_prepare 方法中，原有跳过条件后增加新判断
skip_rope_for_nsa_tilelang_fused = self._skip_rope_for_nsa_tilelang_fused()
skip_rope_for_aiter_fused_mla = self._skip_rope_for_aiter_fused_mla() # 新增
if (
    self.rotary_emb is not None
    and (not self._fuse_rope_for_trtllm_mla(forward_batch))
    and (not skip_rope_for_nsa_tilelang_fused)
    and (not skip_rope_for_aiter_fused_mla) # 新增条件
    and (not _use_aiter or not _is_gfx95_supported or self.use_nsa)
):
    q_pe, k_pe = self.rotary_emb(positions, q_pe, k_pe)
​
# 新方法：判断是否跳过 ROPE，让 fused kernel 处理
def _skip_rope_for_aiter_fused_mla(self: DeepseekV2AttentionMLA) -> bool:
    '''
    Skip rope in prepare and let the fused kernel in forward_absorb_core handle it,
    when running aiter-backend MLA on gfx95 (i.e., the `else` branch in forward_absorb_core
    that calls fused_qk_rope_cat_and_cache_mla).
    '''
    return (
        _use_aiter_gfx95
        and self.current_attention_backend
        not in FORWARD_ABSORB_CORE_ATTENTION_BACKENDS
    )

评论区精华

核心设计讨论：

• Jacob0226 坚决反对环境变量方案，指出默认 True 将 bug 作为默认行为，而默认 False 则永久禁用快速路径。建议直接进行结构性修复。
• 作者 amd-mvarjoka 接受了建议，改为添加 _skip_rope_for_aiter_fused_mla 方法，在 prepare 中跳过 RoPE，由 fused kernel 负责。

• 其他评论涉及环境变量命名（由长名改为简写）和模型配置澄清（1am9trash 提醒 DSv3.2 应使用 --attention-backend nsa）。

• 使用环境变量还是结构性修复 (design): 作者接受建议，改由添加新方法 _skip_rope_for_aiter_fused_mla 进行结构化修复，废弃环境变量方案。

风险与影响

• 风险：该变更仅在 _use_aiter_gfx95 为 True 且 current_attention_backend 不在融合后端列表中时生效，对 MI300X 和非 Aiter 后端无影响。主要风险包括：
  • 条件误判：如果未来新增后端忘记更新 FORWARD_ABSORB_CORE_ATTENTION_BACKENDS，可能导致跳过逻辑失效或错误跳过。
  • 测试覆盖不足：本次未添加直接单元测试，依赖 CI 中的集成测试（如 GSM8K 基准）保证质量，可能遗漏边界情况。
  • 路径耦合：forward_absorb_prepare 和 forward_absorb_core 之间的 RoPE 处理逻辑需要保持一致，否则可能导致精度或性能问题。
  • 影响：直接影响 AMD MI350X (gfx950) 用户，修复了使用 Aiter 后端时 DeepSeek 系列模型的精度问题（GSM8K 从 0.169 恢复至 0.953）。对 MI300X 用户无影响。变更范围局限在单个文件（14 行新增），侵入性低。未引入新配置项，用户无需更改启动命令。
  • 风险标记：核心路径变更, 缺少测试覆盖, AMD 专用代码

关联脉络

• 暂无明显关联 PR