fix compilation error

Author: nv-yunzheq

csrc/nv_internal/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_gemm_template_dispatch.h

@@ -96,9 +96,7 @@ struct genericMoeGemmKernelLauncher {
 
     static_assert(cutlass::platform::is_same<T, WeightType>::value ||
                   cutlass::platform::is_same<WeightType, uint8_t>::value ||
-#if defined(ENABLE_FP4)
                   cutlass::platform::is_same<WeightType, __nv_fp4_e2m1>::value ||
-#endif
                   cutlass::platform::is_same<WeightType, cutlass::uint4b_t>::value);
 
     static_assert(arch::kMinComputeCapability < 90,
@@ -739,42 +737,34 @@ void MoeGemmRunner<T, WeightType, OutputType, ScaleBiasType>::dispatchToArch(
                        "Hopper configuration provided for non-Hopper architecture");
 
   if (sm_ >= 75 && sm_ < 80) {
-#if defined(ENABLE_FP4)
     if constexpr (!std::is_same_v<WeightType, __nv_fp4_e2m1>) {
-#endif
       cutlass_kernels_oss::dispatchMoeGemmToCutlass<T, WeightType, ScaleBiasType,
                                                     cutlass::arch::Sm75, EpilogueTag>(
           inputs, multi_processor_count_);
-#if defined(ENABLE_FP4)
     } else {
       TLLM_THROW("FP4 data type is not supported on SM < 90");
     }
-#endif
   } else if (sm_ >= 80 && sm_ < 90) {
-    if constexpr (use_fp8 || use_w4afp8) {
+    if constexpr (!std::is_same_v<WeightType, __nv_fp4_e2m1>) {
+      if constexpr (use_fp8 || use_w4afp8) {
 #if defined(ENABLE_FP8)
-      static_assert(
-          !std::is_same_v<OutputType, __nv_fp8_e4m3> && !std::is_same_v<OutputType, __nv_fp8_e5m2>,
-          "FP8 GEMM Output not supported");
+        static_assert(!std::is_same_v<OutputType, __nv_fp8_e4m3> &&
+                          !std::is_same_v<OutputType, __nv_fp8_e5m2>,
+                      "FP8 GEMM Output not supported");
 #endif
 
-      TLLM_CHECK_WITH_INFO(sm_ == 89, "For sm >= 80 and < 90, fp8 is only supported with sm == 89");
-      cutlass_kernels_oss::dispatchMoeGemmToCutlass<T, WeightType, ScaleBiasType,
-                                                    cutlass::arch::Sm89, EpilogueTag>(
-          inputs, multi_processor_count_);
-    } else {
-#ifdef ENABLE_FP4
-      if constexpr (std::is_same_v<WeightType, __nv_fp4_e2m1>) {
-        TLLM_THROW("FP4 data type is not supported on SM < 90");
+        TLLM_CHECK_WITH_INFO(sm_ == 89,
+                             "For sm >= 80 and < 90, fp8 is only supported with sm == 89");
+        cutlass_kernels_oss::dispatchMoeGemmToCutlass<T, WeightType, ScaleBiasType,
+                                                      cutlass::arch::Sm89, EpilogueTag>(
+            inputs, multi_processor_count_);
       } else {
         cutlass_kernels_oss::dispatchMoeGemmToCutlass<T, WeightType, ScaleBiasType,
                                                       cutlass::arch::Sm80, EpilogueTag>(
             inputs, multi_processor_count_);
       }
-#else
-      dispatchMoeGemmToCutlass<T, WeightType, ScaleBiasType, cutlass::arch::Sm80, EpilogueTag>(
-          inputs, multi_processor_count_);
-#endif
+    } else {
+      TLLM_THROW("FP4 data type is not supported on SM < 90");
     }
   } else if (sm_ >= 90) {
     // For SM120+ pure FP8 MoE (not FP8 x FP4), redirect to SM89 (Ada) FP8 kernel implementations.
@@ -995,9 +985,6 @@ void MoeGemmRunner<T, WeightType, OutputType, ScaleBiasType>::moeGemmBiasAct(
     case ActivationType::Geglu:
       runGemm<cutlass_extensions::EpilogueOpDefaultFtGelu>(inputs, hopper_inputs);
       break;
-    case ActivationType::Relu2:
-      TLLM_THROW("Relu2 is not supported.");
-      break;
     case ActivationType::InvalidType:
       TLLM_THROW("Activation type for fpA_intB must be valid.");
       break;