[DeepSeek V4] Default FP4 checkpoints to FlashInfer MXFP4 MoE (#35919)
This commit is contained in:
@@ -1164,16 +1164,27 @@ def _deepseek_v4_overrides(server_args: Any, hf_config: Any) -> dict:
|
||||
overrides["swa_full_tokens_ratio"] = 0.1
|
||||
logger.info(f"Setting swa_full_tokens_ratio to 0.1 for {model_arch}.")
|
||||
|
||||
# nvidia/DeepSeek-V4-Pro-NVFP4 uses flashinfer_trtllm_routed MoE runner backend.
|
||||
if (
|
||||
server_args.moe_runner_backend == "auto"
|
||||
and server_args.get_model_config().nvfp4_moe_meta is not None
|
||||
):
|
||||
overrides["moe_runner_backend"] = "flashinfer_trtllm_routed"
|
||||
logger.info(
|
||||
"Use flashinfer_trtllm_routed as MoE runner backend for "
|
||||
f"{model_arch} hybrid FP8+NVFP4 checkpoint."
|
||||
)
|
||||
if server_args.moe_runner_backend == "auto":
|
||||
model_config = server_args.get_model_config()
|
||||
# nvidia/DeepSeek-V4-Pro-NVFP4 uses the routed TRT-LLM runner.
|
||||
if model_config.nvfp4_moe_meta is not None:
|
||||
overrides["moe_runner_backend"] = "flashinfer_trtllm_routed"
|
||||
logger.info(
|
||||
"Use flashinfer_trtllm_routed as MoE runner backend for "
|
||||
f"{model_arch} hybrid FP8+NVFP4 checkpoint."
|
||||
)
|
||||
elif (
|
||||
server_args.device == "cuda"
|
||||
and not is_hip()
|
||||
and server_args.moe_a2a_backend == "none"
|
||||
and not envs.SGLANG_DSV4_FP4_DEQUANT.get()
|
||||
and model_config.is_fp4_experts
|
||||
and (is_sm90_supported() or is_sm100_supported() or is_sm120_supported())
|
||||
):
|
||||
overrides["moe_runner_backend"] = "flashinfer_mxfp4"
|
||||
logger.info(
|
||||
"Use flashinfer_mxfp4 as MoE runner backend for " f"{model_arch}."
|
||||
)
|
||||
return overrides
|
||||
|
||||
|
||||
@@ -1921,20 +1932,6 @@ def _deepseek_v4_kv_cache_dtype(view: Any) -> dict:
|
||||
return {}
|
||||
|
||||
|
||||
@register_post_process
|
||||
def _deepseek_v4_sm120_moe(view: Any) -> dict:
|
||||
"""Default DeepSeek V4 MXFP4 experts to FlashInfer CUTLASS on SM120."""
|
||||
hf_config = view.get_model_config().hf_config
|
||||
if hf_config.architectures[0] != "DeepseekV4ForCausalLM":
|
||||
return {}
|
||||
if is_sm120_supported() and view.moe_runner_backend == "auto":
|
||||
logger.info(
|
||||
"Use flashinfer_mxfp4 as MoE runner backend on SM120 for DeepseekV4"
|
||||
)
|
||||
return {"moe_runner_backend": "flashinfer_mxfp4"}
|
||||
return {}
|
||||
|
||||
|
||||
@_register_for("MuseGlimmerForConditionalGeneration", "MuseGlimmerForCausalLM")
|
||||
def _muse_glimmer_fp4_gemm_runner_overrides(server_args: Any, hf_config: Any) -> dict:
|
||||
if is_sm120_supported() and server_args.fp4_gemm_runner_backend == "auto":
|
||||
|
||||
@@ -5540,15 +5540,6 @@ class ServerArgs:
|
||||
validate_deepseek_v4_cp(self)
|
||||
validate_deepseek_v4_mega_moe_token_budget(self)
|
||||
|
||||
# The SM120 marlin fallback moved to the resolution pipeline
|
||||
# (arg_groups/overrides.py: _deepseek_v4_sm120_moe), invoked here
|
||||
# at its legacy slot.
|
||||
from sglang.srt.arg_groups.overrides import (
|
||||
_deepseek_v4_sm120_moe,
|
||||
run_post_process_pass,
|
||||
)
|
||||
|
||||
run_post_process_pass(self, _deepseek_v4_sm120_moe)
|
||||
if is_sm120_supported():
|
||||
# SM120 lacks tcgen05/TMEM: disable features that depend on
|
||||
# DeepGEMM or require >99KB SMEM (topk_v2).
|
||||
|
||||
@@ -1,6 +1,7 @@
|
||||
"""B200 per-commit CI: DeepSeek-V4-Flash FP4 (LowLatency recipe).
|
||||
|
||||
Launches TP=4 with flashinfer_mxfp4 MoE runner + EAGLE speculative decoding.
|
||||
Launches TP=4 with the auto-selected flashinfer_mxfp4 MoE runner and EAGLE
|
||||
speculative decoding.
|
||||
Runs 12 ServerSanity probes (correctness, streaming, concurrency, determinism)
|
||||
plus a GSM8K accuracy gate.
|
||||
|
||||
@@ -56,8 +57,6 @@ class TestDSV4FlashFP4B200(
|
||||
"--trust-remote-code",
|
||||
"--tp",
|
||||
"4",
|
||||
"--moe-runner-backend",
|
||||
"flashinfer_mxfp4",
|
||||
"--speculative-algorithm",
|
||||
"EAGLE",
|
||||
"--speculative-num-steps",
|
||||
|
||||
@@ -971,20 +971,28 @@ class TestGoldenModelOverrides(_IsolatedPublish):
|
||||
defaults = dict(
|
||||
device="cuda",
|
||||
swa_full_tokens_ratio=ServerArgs.swa_full_tokens_ratio,
|
||||
moe_a2a_backend="none",
|
||||
moe_runner_backend="auto",
|
||||
get_model_config=lambda: SimpleNamespace(nvfp4_moe_meta=None),
|
||||
get_model_config=lambda: SimpleNamespace(
|
||||
is_fp4_experts=True, nvfp4_moe_meta=None
|
||||
),
|
||||
)
|
||||
defaults.update(kw)
|
||||
return SimpleNamespace(**defaults)
|
||||
|
||||
self.assertEqual(
|
||||
_deepseek_v4_overrides(_args(), hf),
|
||||
{
|
||||
"attention_backend": "dsv4",
|
||||
"page_size": 256,
|
||||
"swa_full_tokens_ratio": 0.1,
|
||||
},
|
||||
)
|
||||
with (
|
||||
envs.SGLANG_DSV4_FP4_DEQUANT.override(False),
|
||||
patch.object(overrides_module, "is_sm100_supported", return_value=True),
|
||||
):
|
||||
self.assertEqual(
|
||||
_deepseek_v4_overrides(_args(), hf),
|
||||
{
|
||||
"attention_backend": "dsv4",
|
||||
"moe_runner_backend": "flashinfer_mxfp4",
|
||||
"page_size": 256,
|
||||
"swa_full_tokens_ratio": 0.1,
|
||||
},
|
||||
)
|
||||
# NPU pool geometry
|
||||
self.assertEqual(
|
||||
_deepseek_v4_overrides(_args(device="npu"), hf)["page_size"], 128
|
||||
@@ -994,45 +1002,82 @@ class TestGoldenModelOverrides(_IsolatedPublish):
|
||||
"swa_full_tokens_ratio",
|
||||
_deepseek_v4_overrides(_args(swa_full_tokens_ratio=0.5), hf),
|
||||
)
|
||||
# An explicit user choice takes precedence over the model default.
|
||||
self.assertNotIn(
|
||||
"moe_runner_backend",
|
||||
_deepseek_v4_overrides(_args(moe_runner_backend="triton"), hf),
|
||||
)
|
||||
# FlashInfer MXFP4 only supports the standard (non-A2A) dispatcher.
|
||||
with (
|
||||
envs.SGLANG_DSV4_FP4_DEQUANT.override(False),
|
||||
patch.object(overrides_module, "is_sm100_supported", return_value=True),
|
||||
):
|
||||
self.assertNotIn(
|
||||
"moe_runner_backend",
|
||||
_deepseek_v4_overrides(_args(moe_a2a_backend="deepep"), hf),
|
||||
)
|
||||
# Runtime FP4-to-FP8 dequantization must retain the generic FP8 runner.
|
||||
with (
|
||||
envs.SGLANG_DSV4_FP4_DEQUANT.override(True),
|
||||
patch.object(overrides_module, "is_sm100_supported", return_value=True),
|
||||
):
|
||||
self.assertNotIn(
|
||||
"moe_runner_backend",
|
||||
_deepseek_v4_overrides(_args(), hf),
|
||||
)
|
||||
# FP8 checkpoints and non-CUDA platforms keep their platform-specific
|
||||
# auto-resolution paths.
|
||||
fp8_model_config = lambda: SimpleNamespace(
|
||||
is_fp4_experts=False, nvfp4_moe_meta=None
|
||||
)
|
||||
self.assertNotIn(
|
||||
"moe_runner_backend",
|
||||
_deepseek_v4_overrides(_args(get_model_config=fp8_model_config), hf),
|
||||
)
|
||||
self.assertNotIn(
|
||||
"moe_runner_backend",
|
||||
_deepseek_v4_overrides(_args(device="npu"), hf),
|
||||
)
|
||||
with patch.object(overrides_module, "is_hip", return_value=True):
|
||||
self.assertNotIn(
|
||||
"moe_runner_backend",
|
||||
_deepseek_v4_overrides(_args(), hf),
|
||||
)
|
||||
# Unsupported NVIDIA architectures keep the generic auto-resolution
|
||||
# path instead of selecting a FlashInfer kernel that cannot launch.
|
||||
with (
|
||||
patch.object(overrides_module, "is_sm90_supported", return_value=False),
|
||||
patch.object(overrides_module, "is_sm100_supported", return_value=False),
|
||||
patch.object(overrides_module, "is_sm120_supported", return_value=False),
|
||||
):
|
||||
self.assertNotIn(
|
||||
"moe_runner_backend",
|
||||
_deepseek_v4_overrides(_args(), hf),
|
||||
)
|
||||
# SM120 uses the same model hook; no later pass is needed.
|
||||
with (
|
||||
envs.SGLANG_DSV4_FP4_DEQUANT.override(False),
|
||||
patch.object(overrides_module, "is_sm90_supported", return_value=False),
|
||||
patch.object(overrides_module, "is_sm100_supported", return_value=False),
|
||||
patch.object(overrides_module, "is_sm120_supported", return_value=True),
|
||||
):
|
||||
self.assertEqual(
|
||||
_deepseek_v4_overrides(_args(), hf)["moe_runner_backend"],
|
||||
"flashinfer_mxfp4",
|
||||
)
|
||||
# nvfp4 hybrid checkpoint routes the MoE runner
|
||||
self.assertEqual(
|
||||
_deepseek_v4_overrides(
|
||||
_args(
|
||||
get_model_config=lambda: SimpleNamespace(nvfp4_moe_meta=object())
|
||||
get_model_config=lambda: SimpleNamespace(
|
||||
is_fp4_experts=False, nvfp4_moe_meta=object()
|
||||
)
|
||||
),
|
||||
hf,
|
||||
)["moe_runner_backend"],
|
||||
"flashinfer_trtllm_routed",
|
||||
)
|
||||
|
||||
def test_deepseek_v4_sm120_moe_pass(self):
|
||||
from sglang.srt.arg_groups.overrides import (
|
||||
ResolvedView,
|
||||
_deepseek_v4_sm120_moe,
|
||||
)
|
||||
|
||||
def _view(arch="DeepseekV4ForCausalLM", **kw):
|
||||
hf = SimpleNamespace(architectures=[arch])
|
||||
defaults = dict(moe_runner_backend="auto")
|
||||
defaults.update(kw)
|
||||
return ResolvedView(
|
||||
SimpleNamespace(
|
||||
get_model_config=lambda: SimpleNamespace(hf_config=hf), **defaults
|
||||
)
|
||||
)
|
||||
|
||||
with patch.object(overrides_module, "is_sm120_supported", return_value=True):
|
||||
self.assertEqual(
|
||||
_deepseek_v4_sm120_moe(_view()),
|
||||
{"moe_runner_backend": "flashinfer_mxfp4"},
|
||||
)
|
||||
self.assertEqual(
|
||||
_deepseek_v4_sm120_moe(_view(moe_runner_backend="triton")), {}
|
||||
)
|
||||
self.assertEqual(_deepseek_v4_sm120_moe(_view(arch="LlamaForCausalLM")), {})
|
||||
with patch.object(overrides_module, "is_sm120_supported", return_value=False):
|
||||
self.assertEqual(_deepseek_v4_sm120_moe(_view()), {})
|
||||
|
||||
def test_nemotron_h_overrides_at_callable_level(self):
|
||||
from sglang.srt.arg_groups.overrides import _nemotron_h_overrides
|
||||
|
||||
|
||||
Reference in New Issue
Block a user