revert(moe): revert #3738 SM90 CUTLASS MoE backend (+ dependents #4025, #4080) on release-v0.6.17

csrc/fused_moe/cutlass_backend/cutlass_fused_moe_instantiation.cu

@@ -44,22 +44,17 @@ template class CutlassMoeFCRunner<__nv_fp8_e4m3, cutlass::uint4b_t, __nv_bfloat1
 #endif
 #endif
 #ifdef ENABLE_FP4
-template class CutlassMoeFCRunner<Fp4Type, Fp4Type, half>;
-template class CutlassMoeFCRunner<Fp4Type, Fp4Type, half, half>;
-template class CutlassMoeFCRunner<__nv_fp8_e4m3, Fp4Type, half>;
-template class CutlassMoeFCRunner<__nv_fp8_e4m3, Fp4Type, half, half, half, false,
-                                  Sm90Wfp4Afp8ScaleMode::kHummingPreMmaE8M0>;
-// PHASE3_POST_MMA_PLACEHOLDER: future post-MMA MXFP4 x FP8/MXFP8 paths should
-// use FP16/BF16 external activation input plus online FP8/MXFP8 quantization,
-// not the transitional direct-FP8 input runner.
-template class CutlassMoeFCRunner<half, Fp4Type>;
+template class CutlassMoeFCRunner<__nv_fp4_e2m1, __nv_fp4_e2m1, half>;
+template class CutlassMoeFCRunner<__nv_fp4_e2m1, __nv_fp4_e2m1, half, half>;
+template class CutlassMoeFCRunner<__nv_fp8_e4m3, __nv_fp4_e2m1, half>;
+template class CutlassMoeFCRunner<__nv_fp8_e4m3, __nv_fp4_e2m1, half, half>;
+template class CutlassMoeFCRunner<half, __nv_fp4_e2m1>;
 #ifdef ENABLE_BF16
-template class CutlassMoeFCRunner<Fp4Type, Fp4Type, __nv_bfloat16>;
-template class CutlassMoeFCRunner<Fp4Type, Fp4Type, __nv_bfloat16, __nv_bfloat16>;
-template class CutlassMoeFCRunner<__nv_fp8_e4m3, Fp4Type, __nv_bfloat16>;
-template class CutlassMoeFCRunner<__nv_fp8_e4m3, Fp4Type, __nv_bfloat16, __nv_bfloat16,
-                                  __nv_bfloat16, false, Sm90Wfp4Afp8ScaleMode::kHummingPreMmaE8M0>;
-template class CutlassMoeFCRunner<__nv_bfloat16, Fp4Type>;
+template class CutlassMoeFCRunner<__nv_fp4_e2m1, __nv_fp4_e2m1, __nv_bfloat16>;
+template class CutlassMoeFCRunner<__nv_fp4_e2m1, __nv_fp4_e2m1, __nv_bfloat16, __nv_bfloat16>;
+template class CutlassMoeFCRunner<__nv_fp8_e4m3, __nv_fp4_e2m1, __nv_bfloat16>;
+template class CutlassMoeFCRunner<__nv_fp8_e4m3, __nv_fp4_e2m1, __nv_bfloat16, __nv_bfloat16>;
+template class CutlassMoeFCRunner<__nv_bfloat16, __nv_fp4_e2m1>;
 #endif
 #endif