执行摘要

• 一句话：修复 AMD 上 DSV4 FP4 indexer 的属性错误和 warp mask 编译错误
• 推荐动作：值得快速查阅，尤其是关注跨平台 warp shuffle 兼容性处理模式。代码改动量小（+8/-2），逻辑清晰，适合作为 AMD 特殊修复的参考范例。

功能与动机

DeepSeek V4 FP4 indexer 在 AMD GPU 上无法正常运行：

1) DeepseekV4HipRadixBackend 初始化时未从 server_args 读取 enable_deepseek_v4_fp4_indexer，后续访问该属性会抛出 AttributeError。
2) CUDA 核函数 fused_norm_rope_indexer_fp4 中 warp shuffle 传递了 0xFFFFFFFFu（32 位无符号整数），而 ROCm 的 __shfl_xor_sync 要求 64 位 mask，导致静态断言编译失败。

实现拆解

1. 添加缺失的属性赋值：在 deepseek_v4_backend_hip_radix.py 的 __init__ 方法中，从 model_runner.server_args 读取 enable_deepseek_v4_fp4_indexer 并赋值为实例属性。
2. 修复 warp mask 兼容性：在 fused_norm_rope_v2.cuh 的 fused_norm_rope_indexer_fp4 函数中，将 warp shuffle 调用替换为平台感知的写法：对 CUDA 使用 __shfl_xor_sync(kFullMask, ...)（kFullMask 为 64 位），对 ROCm 使用 __shfl_xor（无需 mask 参数，避免宽度问题）。该模式与同一文件中 fused_norm_rope_indexer 函数（L180-184）保持一致。
3. lint 修复与提交整理：最终由维护者通过一次额外 commit 修正代码格式。

关键文件：

• python/sglang/srt/layers/attention/deepseek_v4_backend_hip_radix.py（模块 注意力后端；类别 source；类型 core-logic）: 修复 AttributeError 的核心改动：在 __init__ 中正确初始化 enable_deepseek_v4_fp4_indexer 属性。
• python/sglang/jit_kernel/csrc/deepseek_v4/fused_norm_rope_v2.cuh（模块 JIT 内核；类别 other；类型 core-logic）: 修复 warp mask 编译错误：将硬编码的 0xFFFFFFFFu 替换为平台适应的 warp shuffle 调用，使用 kFullMask 并区分 CUDA/ROCm。

关键符号：未识别

关键源码片段

python/sglang/srt/layers/attention/deepseek_v4_backend_hip_radix.py

修复 AttributeError 的核心改动：在 __init__ 中正确初始化 enable_deepseek_v4_fp4_indexer 属性。

# python/sglang/srt/layers/attention/deepseek_v4_backend_hip_radix.py ( 部分 )
​
class DeepseekV4HipRadixBackend:
    def __init__(
        self,
        model_runner: ModelRunner,
        skip_prefill: bool = False,
        speculative_step_id=0,
        topk=0,
        speculative_num_steps=0,
    ):
        # ... 其他初始化 ...
        self.c4_topk = getattr(
            model_runner.model_config.hf_text_config, "index_topk", C4_TOPK
        )
        # 修复：之前缺少这行赋值导致 AttributeError
        self.enable_deepseek_v4_fp4_indexer: bool = (
            model_runner.server_args.enable_deepseek_v4_fp4_indexer
        )
        self.topk = model_runner.server_args.speculative_eagle_topk or 0
        # ... 后续初始化 ...

python/sglang/jit_kernel/csrc/deepseek_v4/fused_norm_rope_v2.cuh

修复 warp mask 编译错误：将硬编码的 0xFFFFFFFFu 替换为平台适应的 warp shuffle 调用，使用 kFullMask 并区分 CUDA/ROCm。

// python/sglang/jit_kernel/csrc/deepseek_v4/fused_norm_rope_v2.cuh ( 部分 )
​
// 在 fused_norm_rope_indexer_fp4 核函数中的 warp reduce 循环
for (uint32_t mask = 1; mask < kWarpThreads; mask <<= 1) {
#pragma unroll
    for (int i = 0; i < kVecSize; ++i) {
        // 修复前：const float other = __shfl_xor_sync(0xFFFFFFFFu, data[i], mask, kWarpThreads);
        // 修复后：使用 kFullMask（64 位）并区分平台
#ifndef USE_ROCM
        const float other = __shfl_xor_sync(kFullMask, data[i], mask, kWarpThreads);
#else
        // ROCm 上 __shfl_xor 不要求 mask 参数，避免 64 位 mask 问题
        const float other = __shfl_xor(data[i], mask, kWarpThreads);
#endif
        data[i] = (lane_id & mask) ? (other - data[i]) : (data[i] + other);
    }
}

评论区精华

1. gemini-code-assist 关于 c4_sparse_raw_indices 的警告：指出新增字段可能导致 CUDA graph 回放时的 AssertionError，但最终 PR 并未引入该字段（仅为 bot 误判），未产生实际影响。
2. HaiShaw 指出 ROCm 分支与 PR 描述不一致：PR 描述声称用 kFullMask 替换 0xFFFFFFFFu，但实际实现中对 ROCm 使用了不同函数 __shfl_xor，作者回复“fixed”后保持一致。
3. DarkSharpness 询问能否直接复用 kFullMask：作者解释其实现复制自同一文件 L180-184 的已有模式，以保持一致性。

• 暂无高价值评论线程

风险与影响

• 风险：风险极低：修改集中在两条冷路径——AMD 特定 indexer 后端的初始化和一个 CUDA/ROCm 共享的 warp shuffle 调用。不涉及其他平台或模型。主要风险在于若未来修改 kFullMask 定义或 server_args 字段名，需要同步更新此处；但属性赋值模式与周围代码保持一致，不易出错。
• 影响：影响范围：仅影响 DeepSeek V4 模型在 AMD GPU（gfx950）上启用 FP4 indexer 的路径。影响程度：消除了阻断性 bug，使 AMD 用户可以顺利运行推理。对其他平台（NVIDIA）无行为变化。
• 风险标记：平台兼容性修复, 弱测试覆盖

关联脉络

• 暂无明显关联 PR