[SM120] Allow fused MHC opt-in with standalone TileLang pre disabled

[ormandj]

Motivation

SM120 model post-processing sets SGLANG_OPT_USE_TILELANG_MHC_PRE=False. _is_fused_mhc_post_pre_enabled() also required that flag, so SGLANG_OPT_FUSE_MHC_POST_PRE=1 had no effect on SM120.

A 60-step C1 trace with the fusion flag set contained, per decode cycle, 86 mhc_post, 86 hc_split_sinkhorn, and 86 hc_combine launches, and zero mhc_fused_post_pre_fma_tilelang launches.

mhc_fused_post_pre never reads SGLANG_OPT_USE_TILELANG_MHC_PRE. It dispatches on its own token-count threshold, taking mhc_fused_post_pre_fma_tilelang for small batches and a separate branch for large ones, selecting the GEMM there on SGLANG_OPT_DEEPGEMM_HC_PRENORM. Both regimes are handled independently of the standalone pre path, so that flag should not gate the fusion opt-in.

Modifications

• Allow the fused-MHC opt-in to bypass the standalone TileLang-pre flag on SM120.
• Keep standalone TileLang pre disabled on SM120.
• Keep TileLang post required.
• Keep the gate unchanged on other architectures.
• Add a registered CPU unit test covering SM120, non-SM120, fusion opt-in, and TileLang-post policy.

This follows the fused path added in #25976. It does not change the fused kernel or its numerical behavior.

Accuracy Tests

• Registered CPU unit test (test/registered/unit/models/test_deepseek_v4_fused_mhc_policy.py): 3 cases pass.
• Isolated fused-op reference checks: M=1, 2, 8, 32, 384, and 8192 passed.
• GSM8K smoke: 20/20.

Effect

Hardware: 2x RTX PRO 6000 Blackwell Max-Q at 300 W, TP2, DeepSeek-V4-Flash, FP8 KV, FlashInfer MXFP4 MoE, full decode CUDA graphs, EAGLE MTP2.

With SGLANG_OPT_USE_TILELANG_MHC_PRE at its SM120 default (off), the non-fused path falls back to hc_pre_torch_impl, a torch-level F.linear/rsqrt implementation. The fusion opt-in replaces it. Median target-verify decode graph:

	fused off	fused on
Kernel nodes	3,306	2,796
Summed kernel time	16.230 ms	14.587 ms

The 510-node difference is the hc_pre_torch_impl fallback. Per-category (fused minus non-fused): gemm_cutlass -767 us (492 -> 407 kernels), other -434 us, gemm_other -320 us (148 -> 63), norm -165 us (130 -> 45), fill -97 us, quant +68 us, sparse_attention +74 us.

Kernel-node counts are structural and do not vary run to run. Throughput is not reported here: the available A/B ran with EAGLE MTP2, whose acceptance rate drifts run to run, and the fused cell spanned 6.9% across five 30 s samples. A non-speculative re-measurement under the settled harness is pending.

Checklist

• pre-commit run --files passed on both changed files.
• Added and registered a CPU unit test under test/registered/unit/models/.
• Existing environment variable and interface are unchanged.
• Followed the SGLang code-style guidance.

---

CI States

Latest PR Test (Base): ✅ Run #30159293986
Latest PR Test (Extra): ❌ Run #30159293903

[gemini-code-assist]

Warning

You have reached your daily quota limit. Please wait up to 24 hours and I will start processing your requests again!

[ormandj]

@JoeLee314 This fixes SM120 activation of the fused path added in #25976. Could you review and trigger /tag-and-rerun-ci extra if the change looks correct?

[JoeLee314]

@JoeLee314 This fixes SM120 activation of the fused path added in #25976. Could you review and trigger /tag-and-rerun-ci extra if the change looks correct?

#25976 was submitted by me, but I’m only a contributor and don’t have permission to trigger /tag-and-rerun-ci. Could you ask a maintainer to review it and trigger /tag-and-rerun-ci extra?

[ormandj]

Measured this on SM120 hardware (2x RTX PRO 6000 Blackwell, TP2, DeepSeek-V4-Flash, MTP2 EAGLE steps=2/topk=1/draft-tokens=3, FP8 KV, FlashInfer MXFP4 MoE, full decode CUDA graphs, temperature 0). Five consecutive 30 s C1 runs per cell on a settled server, medians below:

SGLANG_OPT_USE_TILELANG_MHC_PRE	SGLANG_OPT_FUSE_MHC_POST_PRE	median C1 tok/s
unset (SM120 default)	off	172.78
unset (SM120 default)	on (this PR)	188.16
1	off	187.13
1	on	187.60

Reproducibility note: rows 3 and 4 required a local patch wrapping the SM120 defaults in if not ...is_set(). On stock main SGLANG_OPT_USE_TILELANG_MHC_PRE=1 is overwritten by the unconditional .set(False), so those two rows collapse onto rows 1 and 2. Rows 1 and 2 are stock main, row 2 with this PR's gate change.

The second row is what this PR unlocks: +8.9% over the SM120 default, without needing to force SGLANG_OPT_USE_TILELANG_MHC_PRE=1 (which is how we had been working around the gate).

Profiler traces explain the mechanism. Median target-verify decode graph, first row vs second:

	fused off	fused on
Kernel nodes	3,306	2,796
Summed kernel time	16.230 ms	14.587 ms
GPU span	12.287 ms	11.549 ms

Per-category deltas (fused minus non-fused): CUTLASS GEMM -767 us (492 -> 407 kernels), other GEMM -320 us (148 -> 63), norm -165 us (130 -> 45), fill -97 us, misc -434 us; quant +68 us and sparse attention +74 us. With MHC_PRE unset the non-fused path falls back to hc_pre_torch_impl, and those 510 extra nodes are that fallback; the fused kernel replaces it.

Two observations for reviewers:

1. The fused path and the standalone TileLang pre path are alternatives, not complements. Rows 2, 3 and 4 are within run-to-run noise of each other, so enabling both gains nothing over either alone.
2. On SM120 the current default (row 1) is ~8.5% slower than any of the other three configurations, so this gate change is what makes the fast path reachable without an unsupported-by-default workaround.

[ormandj]

@ch-wan @Fridge003 review request when you have a moment.

@ch-wan, you merged #25976 which added the fused MHC post/pre path. This is a three-line fix to its enable gate: on SM120 SGLANG_OPT_USE_TILELANG_MHC_PRE is cleared by model post-processing, and because the gate requires that flag, setting SGLANG_OPT_FUSE_MHC_POST_PRE=1 on SM120 is currently a silent no-op. The change lets the explicit opt-in take effect without also forcing the standalone pre path on.

Measured impact on SM120 hardware is in the comment above: +8.9% C1 decode over the SM120 default (172.78 -> 188.16 median of five runs), with the profiler showing 510 fewer kernel nodes and 1.64 ms less summed kernel time per decode graph. Default behavior on other architectures is unchanged, and a unit test covers the policy matrix.

[mmangkad]

The gate comment only covers decode (prefill uses different kernels), and the test is missing the __main__ block CI needs to run it

[ormandj]

Thanks @mmangkad. Both addressed.

The gate comment now states only what the code does, and the PR body is updated to match. The test has import unittest and an if __name__ == "__main__": unittest.main() block.

On your prefill note: is the concern comment accuracy only, or should the gate itself distinguish prefill from decode? use_fused_mhc_post_pre is resolved once in __init__ and applies to both, so if prefill needs a different policy on SM120 I would rather handle it here than leave it implicit.

Rebased onto current main.

[mmangkad]

@ormandj how did you test the SGLANG_OPT_USE_TILELANG_MHC_PRE=1 rows?

[ormandj]

Via SGLANG_OPT_USE_TILELANG_MHC_PRE=1 in the pod env, on an image that patches the SM120 block to if not ...is_set(). Those two rows are therefore not reproducible on stock main, where the .set(False) overwrites the override.

The two rows the PR rests on do not need that patch: MHC_PRE unset with FUSE off is stock main at 172.78, and MHC_PRE unset with FUSE on is stock main plus this PR's gate change at 188.16. I also edited the table comment above to note which rows need the patch.

The same unconditional .set() also breaks --enable-deepseek-v4-fp4-indexer on SM120: it passes validation at server_args.py:7314, then raises at indexer.py:703 because the SM120 block already forced SGLANG_OPT_USE_TILELANG_INDEXER on, and setting it to 0 is overwritten. Separate PR if wanted.

[mmangkad]

I also patched the override but hit a compilation error. Anyway, it is not reachable on main or needed here, so this looks good.

[mmangkad]

Verified on SM120 with an A/B/A test and saw a repeatable 6%+ gain with the expected fused kernel launches