fix review & fix moe attribute error

Signed-off-by: whx-sjtu <xiaowang990929@gmail.com>

Author: whx-sjtu

vllm/model_executor/layers/fused_moe/oracle/mxfp4.py

@@ -65,6 +65,8 @@ class Mxfp4MoeBackend(Enum):
     MARLIN = "MARLIN"
     # ROCm AITER backends
     AITER_MXFP4_BF16 = "AITER_MXFP4_BF16"  # W4A16: CK kernel
+    # Keep the legacy name as an alias while the ROCm split backend rename settles.
+    AITER = "AITER_MXFP4_BF16"
     AITER_MXFP4_FP8 = "AITER_MXFP4_FP8"  # W4A8: triton kernel
     # Triton
     TRITON = "TRITON"
@@ -255,7 +257,7 @@ def _get_priority_backends() -> list[Mxfp4MoeBackend]:
     backend-level ``is_supported_config`` check filters by device capability).
     """
     if current_platform.is_rocm():
-        return [Mxfp4MoeBackend.AITER]
+        return [Mxfp4MoeBackend.AITER_MXFP4_BF16]
     _AVAILABLE_BACKENDS = [
         Mxfp4MoeBackend.FLASHINFER_TRTLLM_MXFP4_MXFP8,
         Mxfp4MoeBackend.DEEPGEMM_MXFP4,
@@ -555,7 +557,7 @@ def _return_or_raise(
     ):
         priority_backends = [
             Mxfp4MoeBackend.TRITON_UNFUSED,
-            Mxfp4MoeBackend.AITER,
+            Mxfp4MoeBackend.AITER_MXFP4_BF16,
         ]
     else:
         priority_backends = _get_priority_backends()
@@ -1269,7 +1271,7 @@ def convert_weight_to_mxfp4_moe_kernel_format(
             w2_bias,
         )
 
-    elif mxfp4_backend == Mxfp4MoeBackend.AITER:
+    elif mxfp4_backend == Mxfp4MoeBackend.AITER_MXFP4_BF16:
         from vllm._aiter_ops import rocm_aiter_ops
 
         if w13_bias is not None:

vllm/utils/deep_gemm.py

@@ -473,11 +473,7 @@ def tf32_hc_prenorm_gemm(
     """
     _lazy_init()
     if _tf32_hc_prenorm_gemm_impl is None:
-        out.zero_()
-        sqrsum.zero_()
-        out[0].copy_(torch.matmul(x.to(torch.float32), fn.t().to(torch.float32)))
-        sqrsum[0].copy_(x.to(torch.float32).square().sum(dim=-1))
-        return out
+        return _missing()
     return _tf32_hc_prenorm_gemm_impl(
         x,
         fn,

vllm/v1/attention/ops/rocm_aiter_mla_sparse.py

@@ -10,12 +10,16 @@
 
 from vllm.forward_context import get_forward_context
 from vllm.platforms import current_platform
-from vllm.platforms.rocm import _ON_GFX942
 from vllm.triton_utils import tl, triton
 from vllm.utils.torch_utils import LayerNameType
 from vllm.v1.attention.backends.mla.indexer import DeepseekV32IndexerMetadata
 from vllm.v1.attention.ops.common import pack_seq_triton, unpack_seq_triton
 
+if current_platform.is_rocm():
+    from vllm.platforms.rocm import _ON_GFX942
+else:
+    _ON_GFX942 = False
+
 
 @triton.jit
 def _indexer_k_quant_and_cache_kernel(