From 7443073dc617092d80bc3349236f767dca82d194 Mon Sep 17 00:00:00 2001 From: David Orman Date: Sun, 12 Jul 2026 15:45:55 -0500 Subject: [PATCH 1/4] fix(dsv4): allow fused MHC opt-in on SM120 SM120 disables the standalone TileLang mhc_pre path during argument resolution. The fused post/pre gate also required that flag, so SGLANG_OPT_FUSE_MHC_POST_PRE was silently ineffective even though the fused FMA kernel is a separate supported path. Keep the existing TileLang-post requirement and bypass only the standalone-pre requirement on SM120. Signed-off-by: David Orman --- python/sglang/srt/models/deepseek_v4.py | 12 +++++++++--- 1 file changed, 9 insertions(+), 3 deletions(-) diff --git a/python/sglang/srt/models/deepseek_v4.py b/python/sglang/srt/models/deepseek_v4.py index 2c13f6f51e0a..58b8645086be 100644 --- a/python/sglang/srt/models/deepseek_v4.py +++ b/python/sglang/srt/models/deepseek_v4.py @@ -153,6 +153,7 @@ log_info_on_rank0, make_layers, ) +from sglang.srt.utils.common import is_sm120_supported from sglang.srt.utils.custom_op import register_custom_op from sglang.srt.utils.hf_transformers_utils import get_rope_config @@ -204,12 +205,17 @@ def _get_mhc_ops() -> MhcOps: def _is_fused_mhc_post_pre_enabled() -> bool: - # The fused path directly reuses TileLang mhc_post/mhc_pre kernels and their - # tensor layout assumptions, so keep it disabled when either dependency is off. + # SM120 disables the standalone TileLang mhc_pre path because its split-K + # kernel is unsupported there. The fused post/pre kernel has a separate FMA + # implementation that is supported on SM120, so do not let the standalone + # pre-path override silently disable this explicit opt-in. return ( envs.SGLANG_OPT_FUSE_MHC_POST_PRE.get() - and envs.SGLANG_OPT_USE_TILELANG_MHC_PRE.get() and envs.SGLANG_OPT_USE_TILELANG_MHC_POST.get() + and ( + envs.SGLANG_OPT_USE_TILELANG_MHC_PRE.get() + or is_sm120_supported() + ) ) From eafb65607c8c4fb3c0bdb46208981075d24d09bb Mon Sep 17 00:00:00 2001 From: David Orman Date: Sun, 12 Jul 2026 18:49:38 -0500 Subject: [PATCH 2/4] test(dsv4): cover fused MHC enable policy Signed-off-by: David Orman --- python/sglang/srt/models/deepseek_v4.py | 5 +- .../test_deepseek_v4_fused_mhc_policy.py | 74 +++++++++++++++++++ 2 files changed, 75 insertions(+), 4 deletions(-) create mode 100644 test/registered/unit/models/test_deepseek_v4_fused_mhc_policy.py diff --git a/python/sglang/srt/models/deepseek_v4.py b/python/sglang/srt/models/deepseek_v4.py index 58b8645086be..54281587e7d8 100644 --- a/python/sglang/srt/models/deepseek_v4.py +++ b/python/sglang/srt/models/deepseek_v4.py @@ -212,10 +212,7 @@ def _is_fused_mhc_post_pre_enabled() -> bool: return ( envs.SGLANG_OPT_FUSE_MHC_POST_PRE.get() and envs.SGLANG_OPT_USE_TILELANG_MHC_POST.get() - and ( - envs.SGLANG_OPT_USE_TILELANG_MHC_PRE.get() - or is_sm120_supported() - ) + and (envs.SGLANG_OPT_USE_TILELANG_MHC_PRE.get() or is_sm120_supported()) ) diff --git a/test/registered/unit/models/test_deepseek_v4_fused_mhc_policy.py b/test/registered/unit/models/test_deepseek_v4_fused_mhc_policy.py new file mode 100644 index 000000000000..b53df01c6bf0 --- /dev/null +++ b/test/registered/unit/models/test_deepseek_v4_fused_mhc_policy.py @@ -0,0 +1,74 @@ +"""Unit tests for the DeepSeek-V4 fused-MHC enable policy.""" + +from unittest.mock import patch + +import sglang.srt.models.deepseek_v4 as deepseek_v4 +from sglang.srt.environ import envs +from sglang.test.ci.ci_register import register_cpu_ci +from sglang.test.test_utils import CustomTestCase + +register_cpu_ci(est_time=2, suite="base-a-test-cpu") + + +class TestDeepseekV4FusedMHCPolicy(CustomTestCase): + def _is_enabled( + self, + *, + fuse: bool, + tilelang_pre: bool, + tilelang_post: bool, + sm120: bool, + ) -> bool: + with ( + envs.SGLANG_OPT_FUSE_MHC_POST_PRE.override(fuse), + envs.SGLANG_OPT_USE_TILELANG_MHC_PRE.override(tilelang_pre), + envs.SGLANG_OPT_USE_TILELANG_MHC_POST.override(tilelang_post), + patch.object(deepseek_v4, "is_sm120_supported", return_value=sm120), + ): + return deepseek_v4._is_fused_mhc_post_pre_enabled() + + def test_sm120_allows_fused_opt_in_with_standalone_pre_disabled(self): + self.assertTrue( + self._is_enabled( + fuse=True, + tilelang_pre=False, + tilelang_post=True, + sm120=True, + ) + ) + + def test_other_platform_still_requires_tilelang_pre(self): + self.assertFalse( + self._is_enabled( + fuse=True, + tilelang_pre=False, + tilelang_post=True, + sm120=False, + ) + ) + self.assertTrue( + self._is_enabled( + fuse=True, + tilelang_pre=True, + tilelang_post=True, + sm120=False, + ) + ) + + def test_fusion_opt_in_and_tilelang_post_remain_required(self): + self.assertFalse( + self._is_enabled( + fuse=False, + tilelang_pre=False, + tilelang_post=True, + sm120=True, + ) + ) + self.assertFalse( + self._is_enabled( + fuse=True, + tilelang_pre=False, + tilelang_post=False, + sm120=True, + ) + ) From cb1aa4267359065f222746958868a62f606db546 Mon Sep 17 00:00:00 2001 From: David Orman Date: Sat, 25 Jul 2026 05:12:18 -0500 Subject: [PATCH 3/4] fix(dsv4): correct fused-MHC gate comment and register test entrypoint Replace the gate comment's split-K rationale, which was inaccurate: the split-K TileLang kernel in mhc_pre is only reached for num_tokens <= 2048, and split-K is not itself unsupported on SM120 (mhc_fused_post_pre_fma_tilelang uses a dynamic split_k). State only what the code does instead. Add the unittest __main__ block so CI can execute the registered test. --- python/sglang/srt/models/deepseek_v4.py | 10 ++++++---- .../unit/models/test_deepseek_v4_fused_mhc_policy.py | 5 +++++ 2 files changed, 11 insertions(+), 4 deletions(-) diff --git a/python/sglang/srt/models/deepseek_v4.py b/python/sglang/srt/models/deepseek_v4.py index 54281587e7d8..f56e91d5855c 100644 --- a/python/sglang/srt/models/deepseek_v4.py +++ b/python/sglang/srt/models/deepseek_v4.py @@ -205,10 +205,12 @@ def _get_mhc_ops() -> MhcOps: def _is_fused_mhc_post_pre_enabled() -> bool: - # SM120 disables the standalone TileLang mhc_pre path because its split-K - # kernel is unsupported there. The fused post/pre kernel has a separate FMA - # implementation that is supported on SM120, so do not let the standalone - # pre-path override silently disable this explicit opt-in. + # SM120 post-processing clears SGLANG_OPT_USE_TILELANG_MHC_PRE and + # SGLANG_OPT_DEEPGEMM_HC_PRENORM. mhc_fused_post_pre selects its GEMM on + # SGLANG_OPT_DEEPGEMM_HC_PRENORM alone and never reads the pre flag, so with both + # clear it runs mhc_fused_post_pre_fma_tilelang (built with TL_DISABLE_TMA_LOWER + # and TL_DISABLE_WARP_SPECIALIZED). The pre flag must not veto this opt-in. + # Whether the standalone pre path runs stays a server_args decision. return ( envs.SGLANG_OPT_FUSE_MHC_POST_PRE.get() and envs.SGLANG_OPT_USE_TILELANG_MHC_POST.get() diff --git a/test/registered/unit/models/test_deepseek_v4_fused_mhc_policy.py b/test/registered/unit/models/test_deepseek_v4_fused_mhc_policy.py index b53df01c6bf0..d5e1c3518cae 100644 --- a/test/registered/unit/models/test_deepseek_v4_fused_mhc_policy.py +++ b/test/registered/unit/models/test_deepseek_v4_fused_mhc_policy.py @@ -1,5 +1,6 @@ """Unit tests for the DeepSeek-V4 fused-MHC enable policy.""" +import unittest from unittest.mock import patch import sglang.srt.models.deepseek_v4 as deepseek_v4 @@ -72,3 +73,7 @@ def test_fusion_opt_in_and_tilelang_post_remain_required(self): sm120=True, ) ) + + +if __name__ == "__main__": + unittest.main() From 1c0741044a7b11686e00bf4d754ad18f132490cf Mon Sep 17 00:00:00 2001 From: ormandj Date: Sat, 25 Jul 2026 07:59:02 -0500 Subject: [PATCH 4/4] Update python/sglang/srt/models/deepseek_v4.py Co-authored-by: Mohammad Miadh Angkad <176301910+mmangkad@users.noreply.github.com> --- python/sglang/srt/models/deepseek_v4.py | 9 +++------ 1 file changed, 3 insertions(+), 6 deletions(-) diff --git a/python/sglang/srt/models/deepseek_v4.py b/python/sglang/srt/models/deepseek_v4.py index f56e91d5855c..4014e4e6c89b 100644 --- a/python/sglang/srt/models/deepseek_v4.py +++ b/python/sglang/srt/models/deepseek_v4.py @@ -205,12 +205,9 @@ def _get_mhc_ops() -> MhcOps: def _is_fused_mhc_post_pre_enabled() -> bool: - # SM120 post-processing clears SGLANG_OPT_USE_TILELANG_MHC_PRE and - # SGLANG_OPT_DEEPGEMM_HC_PRENORM. mhc_fused_post_pre selects its GEMM on - # SGLANG_OPT_DEEPGEMM_HC_PRENORM alone and never reads the pre flag, so with both - # clear it runs mhc_fused_post_pre_fma_tilelang (built with TL_DISABLE_TMA_LOWER - # and TL_DISABLE_WARP_SPECIALIZED). The pre flag must not veto this opt-in. - # Whether the standalone pre path runs stays a server_args decision. + # SM120 disables the standalone TileLang pre path. mhc_fused_post_pre does + # not read that flag and dispatches independently for both small and large + # token batches, so the standalone pre flag must not veto the fused opt-in. return ( envs.SGLANG_OPT_FUSE_MHC_POST_PRE.get() and envs.SGLANG_OPT_USE_TILELANG_MHC_POST.get()