[Backport][NVFP4] ds4-sm120-* hardcodes Mxfp4MoEMethod — DeepSeek-V4-Flash-NVFP4 fails to load on SM120 (works via ModelOpt→Marlin)

[danielwoz]

Summary

On ds4-sm120-preview (and also ds4-sm120, ds4-sm120-full, dsv4, ds4-sm120-next, codex/ds4-sm120-min-enable), DeepseekV4FP8Config.get_quant_method returns Mxfp4MoEMethod for any expert_dtype == "fp4". That only handles DeepSeek's native MXFP4 experts (group-32, e8m0 scales). NVIDIA's nvidia/DeepSeek-V4-Flash-NVFP4 repackages the experts as ModelOpt NVFP4 (group-16, e4m3 block scales + per-tensor weight_scale_2 / input_scale; config.json: producer.name=modelopt, moe_quant_algo=NVFP4, group_size=16). Loading that checkpoint through the MXFP4 path fails:

1. KeyError: 'layers.0.ffn.experts.w13_input_scale', then
2. after guarding that, a 128 vs 256 shape mismatch in FusedMoE._load_w13 (the group-16 vs group-32 ratio).

This is a backport, not a new design (prior art / not-a-duplicate)

I want to be upfront that this duplicates work that already exists elsewhere:

• Your closed PR deepseek_v4: route NVFP4-modelopt experts to ModelOptNvFp4FusedMoE #18 (@vedcsolution) proposed the same concept (route NVFP4-modelopt experts to ModelOptNvFp4FusedMoE), but the landed commit took the FlashInfer CUTLASS / B12x path and was validated on GB10 (sm_121a). It did not leave a deepseek_v4.py-level route in any current branch, so the gap is still open on these branches.
• Upstream vllm-project/vllm has since merged equivalent support: Add NVFP4 MOE support for Deepseek V4. vllm-project/vllm#42209 ("Add NVFP4 MOE support for Deepseek V4") and [NVFP4 MoE/Deepseek V4] Marlin: wire SwiGLU clamp + allow it for clamped models on non-Blackwell vllm-project/vllm#45836 ("[NVFP4 MoE/Deepseek V4] Marlin: wire SwiGLU clamp + allow it for clamped models on non-Blackwell"), which fixes the exact No NvFp4 MoE backend supports the deployment configuration error on non-SM100 cards. These fork branches predate that work and haven't inherited it (your SM12x enablement is still tracked in the open upstream PR [New Model][Nvidia] Add SM12x support for DeepSeek V4 Flash with essential fixes vllm-project/vllm#41834).

So this is really a backport request. Filing as an issue (rather than a PR onto a frozen branch) so you can decide whether to take a minimal backport on the preview/active SM120 branches or just rebase them onto current upstream main (which subsumes it).

Minimal fix that works on pure SM120 (2× RTX PRO 6000 Blackwell)

Two changes in vllm/model_executor/models/deepseek_v4.py:

1. In get_quant_method, when the checkpoint is ModelOpt-NVFP4 (producer==modelopt and moe_quant_algo==NVFP4), build a ModelOptNvFp4Config and return ModelOptNvFp4FusedMoE instead of Mxfp4MoEMethod. That method registers the correct two-level NVFP4 tensors and auto-selects the Marlin NVFP4 MoE backend, which runs on SM120 with no SM100 CUTLASS / FlashInfer dependency.
2. A guard in the expert-loading loop that continues when name_mapped not in params_dict, so the checkpoint's per-tensor *_input_scale tensors (unused by this kernel path) are skipped instead of raising KeyError.

Applies cleanly to ds4-sm120-preview (git apply --check OK), 2 hunks, +51 lines, 0 deletions:

diff

diff --git a/vllm/model_executor/models/deepseek_v4.py b/vllm/model_executor/models/deepseek_v4.py
--- a/vllm/model_executor/models/deepseek_v4.py
+++ b/vllm/model_executor/models/deepseek_v4.py
@@ -203,11 +203,55 @@ class DeepseekV4FP8Config(Fp8Config):
             ):
                 return UnquantizedFusedMoEMethod(layer.moe_config)
             if self.expert_dtype == "fp4":
+                # NVIDIA's DeepSeek-V4-Flash-NVFP4 repackages the experts as
+                # ModelOpt NVFP4 (group_size=16, e4m3 block scales + per-tensor
+                # weight_scale_2/input_scale) instead of the native MXFP4
+                # (group 32, e8m0) layout Mxfp4MoEMethod expects. Route those
+                # experts to vLLM's ModelOpt NVFP4 FusedMoE method, which
+                # selects a Marlin (SM120-capable) backend.
+                nvfp4_method = self._maybe_modelopt_nvfp4_moe_method(layer)
+                if nvfp4_method is not None:
+                    return nvfp4_method
                 return Mxfp4MoEMethod(layer.moe_config)
             # expert_dtype == "fp8": fall through to Fp8Config which
             # returns Fp8MoEMethod with block-wise float32 scales.
         return super().get_quant_method(layer, prefix)
 
+    def _modelopt_nvfp4_quant_config(self):
+        """Return the raw checkpoint quantization_config dict, or {}."""
+        try:
+            hf_cfg = get_current_vllm_config().model_config.hf_config
+        except Exception:
+            return {}
+        qc = getattr(hf_cfg, "quantization_config", None)
+        if qc is None:
+            return {}
+        if hasattr(qc, "to_dict"):
+            try:
+                qc = qc.to_dict()
+            except Exception:
+                pass
+        return qc if isinstance(qc, dict) else {}
+
+    def _maybe_modelopt_nvfp4_moe_method(self, layer):
+        qc = self._modelopt_nvfp4_quant_config()
+        producer = (qc.get("producer") or {}).get("name", "")
+        moe_algo = qc.get("moe_quant_algo")
+        if not (producer == "modelopt" and moe_algo == "NVFP4"):
+            return None
+        from vllm.model_executor.layers.quantization.modelopt import (
+            ModelOptNvFp4Config,
+            ModelOptNvFp4FusedMoE,
+        )
+        if getattr(self, "_nvfp4_moe_cfg", None) is None:
+            self._nvfp4_moe_cfg = ModelOptNvFp4Config(
+                is_checkpoint_nvfp4_serialized=True,
+                kv_cache_quant_algo=None,
+                exclude_modules=list(getattr(self, "ignored_layers", []) or []),
+                group_size=int(qc.get("group_size", 16) or 16),
+            )
+        return ModelOptNvFp4FusedMoE(self._nvfp4_moe_cfg, layer.moe_config)
+
     def is_mxfp4_quant(self, prefix, layer):
         return isinstance(layer, FusedMoE) and self.expert_dtype == "fp4"
 
@@ -1452,6 +1496,13 @@ class DeepseekV4Model(nn.Module):
                         name_mapped = name.replace(weight_name, param_name)
                         if is_pp_missing_parameter(name_mapped, self):
                             continue
+                        if name_mapped not in params_dict:
+                            # The selected MoE backend (e.g. MXFP4/Marlin block
+                            # microscaling on SM120) does not register this scale.
+                            # NVIDIA's NVFP4 checkpoint ships per-tensor
+                            # *_input_scale tensors that this kernel path does not
+                            # use; skip them instead of raising KeyError.
+                            continue
                         param = params_dict[name_mapped]
                         # We should ask the weight loader to return success or not
                         # here since otherwise we may skip experts with other

Why Marlin on SM120

ModelOptNvFp4FusedMoE's backend selector picks the Marlin NVFP4 MoE kernel on SM120 (compute 12.0), avoiding the SM100-only TMEM/tcgen05 CUTLASS and FlashInfer-TRTLLM paths that aren't available on RTX PRO 6000 / 5090-class parts. This matches the resolution merged upstream in vllm-project#45836.

Test evidence

• HW: 2× NVIDIA RTX PRO 6000 Blackwell (SM120 / compute 12.0), tensor-parallel 2.
• Model: nvidia/DeepSeek-V4-Flash-NVFP4, --kv-cache-dtype fp8.
• Result: model loads and serves; ~59 tok/s single-stream decode; correct outputs on math and text prompts; tool-calling works via the Anthropic /v1/messages endpoint.
• Backend selected at runtime: Marlin NVFP4 MoE (no SM100 CUTLASS / FlashInfer required).

AI-assistance disclosure

This change and the investigation behind it were produced with AI assistance (Claude Code). Validation was end-to-end serving on the 2× RTX PRO 6000 host (model load + generation + tool-call), as above; the full pytest kernel/unit suite was not run on this SM120 host. Happy to open a PR against whichever branch you prefer, or to adjust the approach.

[jasl]

Will check later

[jasl]

Thanks for the thorough write-up — and you called it. The 06-30 upstream merge (a5ebb5f66b) subsumed this on the two active branches (codex/ds4-sm120-min-enable and ds4-sm120-preview-dev).

On current HEAD it now:

• routes ModelOpt-NVFP4 experts to ModelOptNvFp4FusedMoE in vllm/models/deepseek_v4/quant_config.py — keyed on moe_quant_algo == "NVFP4" (note: this checkpoint reads producer == None, so a producer == "modelopt" guard wouldn't fire — keying on moe_quant_algo is what triggers);
• loads the w1/w3 input_scale tensors properly in RoutedExperts.weight_loader (via _to_scalar) rather than skipping them, so there's no w13_input_scale KeyError and no 128-vs-256 shape mismatch;
• picks up upstream [NVFP4 MoE/Deepseek V4] Marlin: wire SwiGLU clamp + allow it for clamped models on non-Blackwell vllm-project/vllm#45836, so MARLIN is now a valid SM120 NVFP4 MoE candidate as well.

Verified end-to-end on 2× RTX PRO 6000 (SM120, TP=2, --kv-cache-dtype fp8): nvidia/DeepSeek-V4-Flash-NVFP4 loads with 0 errors, resolves expert_dtype='fp4', and selects:

Using 'FLASHINFER_CUTLASS' NvFp4 MoE backend out of ['FLASHINFER_TRTLLM', 'FLASHINFER_CUTLASS', 'MARLIN']
Model loading took 77.95 GiB  |  GPU KV cache 85,120 tokens

Generation is correct (e.g. "60 km / 1.5 h = 40 km/h") at ~53–90 tok/s single-stream — consistent with your numbers. (Auto-select lands on FLASHINFER_CUTLASS here rather than Marlin, but both are SM120-capable now.)

The other branches you listed (ds4-sm120-preview, ds4-sm120, ds4-sm120-full, dsv4, ds4-sm120-next) are stale April–May snapshots and won't be backported — the active branches (PR + dev) are where development continues.

Leaving this open for now in case you want to confirm on your end; happy to close once you've had a look. Thanks again for the careful report and the prior-art links.