[bugfix][AMD] AttributeError and warp mask bugs in DeepSeek V4 FP4 indexer (#27152)

Co-authored-by: HaiShaw <hixiao@gmail.com>

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);
       }
     }

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