[Feature] NVFP4 Marlin fallback for non-Blackwell GPUs (SM75+)

python/sglang/jit_kernel/csrc/gemm/marlin_moe/marlin_template.h

@@ -1243,17 +1243,17 @@ __global__ void Marlin(
       }
     }
 
-    // Commented out FP4/FP8 scale dequantization since we don't generate
-    // kFE2M1f kernels to reduce compilation time
-    // if constexpr (w_type == host::kFE2M1f) {
-    //   int s_quant_0 = reinterpret_cast<int*>(frag_s[k2])[0];
-    //   int s_quant_1 = reinterpret_cast<int*>(frag_s[k2])[1];
-    //
-    //   dequant_fp8_scales<scalar_t2, s_type_id>(
-    //       s_quant_0, reinterpret_cast<scalar_t2*>(&frag_s[k2]));
-    //   dequant_fp8_scales<scalar_t2, s_type_id>(
-    //       s_quant_1, reinterpret_cast<scalar_t2*>(&frag_s[k2]) + 2);
-    // }
+    // Convert FP8 per-group scales to BF16/FP16 before applying them.
+    // Required for kFE2M1f (NVFP4): frag_s holds raw float8_e4m3fn bytes;
+    // without this conversion scale<scalar_t> would misinterpret them as
+    // BF16/FP16, producing NaN/Inf multipliers.
+    if constexpr (w_type == host::kFE2M1f) {
+      int s_quant_0 = reinterpret_cast<int*>(frag_s[k2])[0];
+      int s_quant_1 = reinterpret_cast<int*>(frag_s[k2])[1];
+
+      dequant_fp8_scales<scalar_t2, s_type_id>(s_quant_0, reinterpret_cast<scalar_t2*>(&frag_s[k2]));
+      dequant_fp8_scales<scalar_t2, s_type_id>(s_quant_1, reinterpret_cast<scalar_t2*>(&frag_s[k2]) + 2);
+    }
 
 // We have the m dimension as the inner loop in order to encourage overlapping
 // dequantization and matmul operations.

python/sglang/srt/layers/moe/fused_moe_triton/fused_marlin_moe.py

@@ -23,6 +23,12 @@ def get_scalar_type(num_bits: int, has_zp: bool):
         return scalar_types.uint4b8 if num_bits == 4 else scalar_types.uint8b128
 
 
+def _get_fp4_scalar_type():
+    from sgl_kernel.scalar_type import scalar_types
+
+    return scalar_types.float4_e2m1f
+
+
 @register_custom_op(out_shape="hidden_states")
 def fused_marlin_moe(
     hidden_states: torch.Tensor,
@@ -46,6 +52,8 @@ def fused_marlin_moe(
     is_k_full: bool = True,
     inplace: bool = False,
     routed_scaling_factor: Optional[float] = None,
+    w1_global_scale: Optional[torch.Tensor] = None,
+    w2_global_scale: Optional[torch.Tensor] = None,
 ) -> torch.Tensor:
     """
     This function computes a Mixture of Experts (MoE) layer using two sets of
@@ -76,6 +84,9 @@ def fused_marlin_moe(
     """
     from sglang.srt.layers.moe.fused_moe_triton import moe_align_block_size
 
+    # Detect FP4 Marlin mode (when global scales are provided)
+    _is_fp4_marlin = w1_global_scale is not None
+
     assert hidden_states.shape[0] == gating_output.shape[0], "Number of tokens mismatch"
     assert hidden_states.shape[1] == w1.shape[1] * 16, "Hidden size mismatch w1"
     assert hidden_states.shape[1] == w2.shape[2] // (
@@ -85,12 +96,14 @@ def fused_marlin_moe(
     assert w1.is_contiguous(), "Expert weights1 must be contiguous"
     assert w2.is_contiguous(), "Expert weights2 must be contiguous"
     assert hidden_states.dtype in [torch.float16, torch.bfloat16]
-    assert (
-        hidden_states.dtype == w1_scale.dtype
-    ), f"moe_wna16_marlin_gemm assumes hidden_states.dtype ({hidden_states.dtype}) == w1_scale.dtype ({w1_scale.dtype})"
-    assert (
-        hidden_states.dtype == w2_scale.dtype
-    ), f"moe_wna16_marlin_gemm assumes hidden_states.dtype ({hidden_states.dtype}) == w2_scale.dtype ({w2_scale.dtype})"
+    # For FP4 Marlin, scales are in special float8_e4m3fn format (not input dtype)
+    if not _is_fp4_marlin:
+        assert (
+            hidden_states.dtype == w1_scale.dtype
+        ), f"moe_wna16_marlin_gemm assumes hidden_states.dtype ({hidden_states.dtype}) == w1_scale.dtype ({w1_scale.dtype})"
+        assert (
+            hidden_states.dtype == w2_scale.dtype
+        ), f"moe_wna16_marlin_gemm assumes hidden_states.dtype ({hidden_states.dtype}) == w2_scale.dtype ({w2_scale.dtype})"
     assert num_bits in [4, 8]
 
     M, K = hidden_states.shape
@@ -121,8 +134,13 @@ def fused_marlin_moe(
             max_workspace_size, dtype=torch.int, device=device, requires_grad=False
         )
 
-    scalar_type1 = get_scalar_type(num_bits, w1_zeros is not None)
-    scalar_type2 = get_scalar_type(num_bits, w2_zeros is not None)
+    # FP4 Marlin uses float4_e2m1f scalar type (not uint4b8/uint8b128)
+    if _is_fp4_marlin:
+        scalar_type1 = _get_fp4_scalar_type()
+        scalar_type2 = _get_fp4_scalar_type()
+    else:
+        scalar_type1 = get_scalar_type(num_bits, w1_zeros is not None)
+        scalar_type2 = get_scalar_type(num_bits, w2_zeros is not None)
 
     intermediate_cache2 = torch.empty(
         (M * topk_ids.shape[1], N),
@@ -150,7 +168,7 @@ def fused_marlin_moe(
         w1,
         None,  # b_bias_or_none
         w1_scale,
-        None,  # global_scale_or_none
+        w1_global_scale,  # None for INT4/INT8, tensor for FP4 Marlin
         w1_zeros,
         g_idx1,
         sort_indices1,
@@ -184,7 +202,7 @@ def fused_marlin_moe(
         w2,
         None,  # b_bias_or_none
         w2_scale,
-        None,  # global_scale_or_none
+        w2_global_scale,  # None for INT4/INT8, tensor for FP4 Marlin
         w2_zeros,
         g_idx2,
         sort_indices2,

python/sglang/srt/layers/moe/moe_runner/marlin.py

@@ -69,8 +69,13 @@ class MarlinMoeQuantInfo(MoeQuantInfo):
     w13_qzeros: Optional[torch.Tensor] = None
     w2_qzeros: Optional[torch.Tensor] = None
 
-    # Optional
+    # FP4 Marlin specific (Optional)
+    w13_global_scale: Optional[torch.Tensor] = None
+    w2_global_scale: Optional[torch.Tensor] = None
+
+    # EP support (Optional)
     expert_map: Optional[torch.Tensor] = None
+    global_num_experts: int = -1
 
 
 @register_fused_func("none", "marlin")
@@ -106,6 +111,7 @@ def fused_experts_none_to_marlin(
         gating_output=topk_output.router_logits,
         topk_weights=topk_output.topk_weights,
         topk_ids=topk_output.topk_ids,
+        global_num_experts=quant_info.global_num_experts,
         expert_map=quant_info.expert_map,
         g_idx1=quant_info.w13_g_idx,
         g_idx2=quant_info.w2_g_idx,
@@ -118,6 +124,8 @@ def fused_experts_none_to_marlin(
         is_k_full=quant_info.is_k_full,
         inplace=runner_config.inplace,
         routed_scaling_factor=runner_config.routed_scaling_factor,
+        w1_global_scale=quant_info.w13_global_scale,
+        w2_global_scale=quant_info.w2_global_scale,
     ).to(hidden_states.dtype)
 
     return StandardCombineInput(

python/sglang/srt/layers/quantization/compressed_tensors/schemes/compressed_tensors_w4a4_nvfp4.py

@@ -34,7 +34,7 @@ def __init__(self):
 
     @classmethod
     def get_min_capability(cls) -> int:
-        return 100
+        return 75  # SM75+ (Turing) supports Marlin FP4 fallback; SM100 for native FP4
 
     def create_weights(
         self,
@@ -47,6 +47,7 @@ def create_weights(
     ):
         output_size_per_partition = sum(output_partition_sizes)
         layer.logical_widths = output_partition_sizes
+        layer.params_dtype = params_dtype
         layer.input_size_per_partition = input_size_per_partition
         layer.output_size_per_partition = output_size_per_partition
 
@@ -91,6 +92,25 @@ def create_weights(
         layer.register_parameter("input_global_scale", input_global_scale)
 
     def process_weights_after_loading(self, layer) -> None:
+        from sglang.srt.layers.quantization.fp8_utils import is_blackwell_supported
+        from sglang.srt.layers.quantization.marlin_utils_fp4 import (
+            is_fp4_marlin_supported,
+            prepare_fp4_layer_for_marlin,
+        )
+
+        if not is_blackwell_supported() and is_fp4_marlin_supported():
+            # Marlin FP4 fallback: consolidate global scale then repack weights
+            global_scale = layer.weight_global_scale.max().to(torch.float32)
+            layer.weight_global_scale = Parameter(global_scale, requires_grad=False)
+            prepare_fp4_layer_for_marlin(
+                layer,
+                weight_attr="weight_packed",
+                weight_scale_attr="weight_scale",
+                weight_global_scale_attr="weight_global_scale",
+            )
+            layer.use_marlin_fallback = True
+            return
+
         global_input_scale = layer.input_global_scale.max().to(torch.float32)
         layer.input_global_scale = Parameter(global_input_scale, requires_grad=False)
 
@@ -136,6 +156,22 @@ def apply_weights(
         x: torch.Tensor,
         bias: Optional[torch.Tensor] = None,
     ) -> torch.Tensor:
+        if getattr(layer, "use_marlin_fallback", False):
+            from sglang.srt.layers.quantization.marlin_utils_fp4 import (
+                apply_fp4_marlin_linear,
+            )
+
+            return apply_fp4_marlin_linear(
+                input=x,
+                weight=layer.weight_packed,
+                weight_scale=layer.weight_scale,
+                weight_global_scale=layer.weight_global_scale,
+                workspace=layer.marlin_workspace,
+                size_n=layer.output_size_per_partition,
+                size_k=layer.input_size_per_partition,
+                bias=bias,
+            )
+
         output_dtype = x.dtype
         w_n, _ = layer.weight_packed.shape
         output_shape = [x.shape[0], w_n]

python/sglang/srt/layers/quantization/compressed_tensors/schemes/compressed_tensors_w4a4_nvfp4_moe.py

@@ -38,19 +38,30 @@
 class CompressedTensorsW4A4Nvfp4MoE(CompressedTensorsMoEScheme):
 
     def __init__(self):
+        self.group_size = 16
         if not is_blackwell_supported():
-            raise ValueError(
-                "Current platform does not support NVFP4"
-                " quantization. Please use Blackwell and"
-                " above."
+            from sglang.srt.layers.quantization.marlin_utils_fp4 import (
+                is_fp4_marlin_supported,
             )
-        self.group_size = 16
-        self.use_flashinfer_trtllm = get_moe_runner_backend().is_flashinfer_trtllm()
+
+            if not is_fp4_marlin_supported():
+                raise ValueError(
+                    "Current platform does not support NVFP4"
+                    " quantization. Please use SM75+ (Turing or newer)."
+                )
+            logger.warning_once(
+                "GPU is not Blackwell (SM100+). Using Marlin FP4 fallback kernel "
+                "for MoE layers. Weights remain compressed in FP4 format."
+            )
+            self.use_marlin_fallback = True
+            self.use_flashinfer_trtllm = False
+        else:
+            self.use_marlin_fallback = False
+            self.use_flashinfer_trtllm = get_moe_runner_backend().is_flashinfer_trtllm()
 
     @classmethod
     def get_min_capability(cls) -> int:
-        # Requires sm100(blackwell) architecture
-        return 100
+        return 75  # SM75+ (Turing) supports Marlin FP4 fallback; SM100 for native FP4
 
     def create_weights(
         self,
@@ -64,6 +75,7 @@ def create_weights(
         from sglang.srt.layers.moe.fused_moe_triton import FusedMoeWeightScaleSupported
 
         layer.params_dtype = params_dtype
+        layer.intermediate_size_per_partition = intermediate_size_per_partition
 
         w13_weight = torch.nn.Parameter(
             torch.empty(
@@ -175,6 +187,25 @@ def process_weights_after_loading(self, layer: torch.nn.Module) -> None:
         )
         delattr(layer, "w2_weight_packed")
 
+        if self.use_marlin_fallback:
+            from sglang.srt.layers.quantization.marlin_utils_fp4 import (
+                prepare_moe_fp4_layer_for_marlin,
+            )
+
+            # CompressedTensors checkpoint: global_scale is stored as the inverse.
+            # Actual dequant scale = 1 / stored_value. We create w*_weight_scale_2
+            # with the actual scale before calling prepare_moe_fp4_layer_for_marlin().
+            layer.w13_weight_scale_2 = torch.nn.Parameter(
+                (1.0 / layer.w13_weight_global_scale).to(layer.params_dtype),
+                requires_grad=False,
+            )  # [E, 2]
+            layer.w2_weight_scale_2 = torch.nn.Parameter(
+                (1.0 / layer.w2_weight_global_scale).to(layer.params_dtype),
+                requires_grad=False,
+            )  # [E]
+            prepare_moe_fp4_layer_for_marlin(layer)
+            return
+
         if self.use_flashinfer_trtllm:
             w, s = reorder_w1w3_to_w3w1(
                 layer.w13_weight.data, layer.w13_weight_scale.data, dim=-2
@@ -303,7 +334,10 @@ def create_moe_runner(
         self, layer: torch.nn.Module, moe_runner_config: MoeRunnerConfig
     ):
         self.moe_runner_config = moe_runner_config
-        self.runner = MoeRunner(MoeRunnerBackend.TRITON, moe_runner_config)
+        if self.use_marlin_fallback:
+            self.runner = MoeRunner(MoeRunnerBackend.MARLIN, moe_runner_config)
+        else:
+            self.runner = MoeRunner(MoeRunnerBackend.TRITON, moe_runner_config)
 
     def apply_weights(
         self,
@@ -313,6 +347,33 @@ def apply_weights(
 
         from sglang.srt.layers.moe.token_dispatcher import StandardCombineInput
 
+        if self.use_marlin_fallback:
+            from sglang.srt.layers.moe.moe_runner.marlin import MarlinMoeQuantInfo
+
+            expert_map = None
+            global_num_experts = -1
+            if hasattr(layer, "dispatcher") and hasattr(
+                layer.dispatcher, "local_expert_mapping"
+            ):
+                expert_map = layer.dispatcher.local_expert_mapping
+                if expert_map is not None:
+                    global_num_experts = self.moe_runner_config.num_experts
+
+            quant_info = MarlinMoeQuantInfo(
+                w13_qweight=layer.w13_weight,
+                w2_qweight=layer.w2_weight,
+                w13_scales=layer.w13_weight_scale,
+                w2_scales=layer.w2_weight_scale,
+                w13_g_idx_sort_indices=None,
+                w2_g_idx_sort_indices=None,
+                weight_bits=4,
+                w13_global_scale=layer.w13_weight_scale_2,
+                w2_global_scale=layer.w2_weight_scale_2,
+                expert_map=expert_map,
+                global_num_experts=global_num_experts,
+            )
+            return self.runner.run(dispatch_output, quant_info)
+
         x = dispatch_output.hidden_states
         topk_output = dispatch_output.topk_output