# PR #27152 完整报告

- 仓库：`sgl-project/sglang`
- 标题：[bugfix][AMD] AttributeError and warp mask bugs in DeepSeek V4 FP4 indexer
- 合并时间：2026-06-06 09:26
- 原文链接：http://prhub.com.cn/sgl-project/sglang/pull/27152

---

# 执行摘要

- 一句话：修复 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
# 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。

```cpp
// 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