From 3030119ef7cbb2ee3c8666f0aeae9f23135ba104 Mon Sep 17 00:00:00 2001 From: Xinyi Song Date: Fri, 5 Jun 2026 18:26:38 -0700 Subject: [PATCH] [bugfix][AMD] AttributeError and warp mask bugs in DeepSeek V4 FP4 indexer (#27152) Co-authored-by: HaiShaw --- .../jit_kernel/csrc/deepseek_v4/fused_norm_rope_v2.cuh | 6 +++++- .../srt/layers/attention/deepseek_v4_backend_hip_radix.py | 4 +++- 2 files changed, 8 insertions(+), 2 deletions(-) diff --git a/python/sglang/jit_kernel/csrc/deepseek_v4/fused_norm_rope_v2.cuh b/python/sglang/jit_kernel/csrc/deepseek_v4/fused_norm_rope_v2.cuh index a9cac1754..5ffe8fdde 100644 --- a/python/sglang/jit_kernel/csrc/deepseek_v4/fused_norm_rope_v2.cuh +++ b/python/sglang/jit_kernel/csrc/deepseek_v4/fused_norm_rope_v2.cuh @@ -321,7 +321,11 @@ INDEXER_KERNEL void fused_norm_rope_indexer_fp4(const __grid_constant__ FusedNor 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); +#ifndef USE_ROCM + const float other = __shfl_xor_sync(kFullMask, data[i], mask, kWarpThreads); +#else + const float other = __shfl_xor(data[i], mask, kWarpThreads); +#endif data[i] = (lane_id & mask) ? (other - data[i]) : (data[i] + other); } } diff --git a/python/sglang/srt/layers/attention/deepseek_v4_backend_hip_radix.py b/python/sglang/srt/layers/attention/deepseek_v4_backend_hip_radix.py index 3571813af..0ad3e35e8 100644 --- a/python/sglang/srt/layers/attention/deepseek_v4_backend_hip_radix.py +++ b/python/sglang/srt/layers/attention/deepseek_v4_backend_hip_radix.py @@ -373,7 +373,9 @@ class DeepseekV4HipRadixBackend( self.c4_topk = getattr( model_runner.model_config.hf_text_config, "index_topk", C4_TOPK ) - + 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 assert self.topk in [0, 1], "MTP Topk > 1 not supported for DeepSeek V4" self.mtp_enabled = self.topk > 0