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