Skip to content

[deepseek-v4] Bump sgl-deep-gemm for packed FP4×FP4 mega-MoE (W4A4) - #28210

Open
Romaosir wants to merge 2 commits into
sgl-project:mainfrom
Romaosir:w4a4-packed-fp4-bump
Open

[deepseek-v4] Bump sgl-deep-gemm for packed FP4×FP4 mega-MoE (W4A4)#28210
Romaosir wants to merge 2 commits into
sgl-project:mainfrom
Romaosir:w4a4-packed-fp4-bump

Conversation

@Romaosir

@Romaosir Romaosir commented Jun 14, 2026

Copy link
Copy Markdown

Bumps the pinned sgl-deep-gemm version (pyproject + docker ARG) to the release carrying the new sm100_fp4_mega_moe packed FP4×FP4 (W4A4) mega-MoE kernel, and propagates the selection.

In _apply_mega_moe_dg_env, when SGLANG_OPT_DEEPGEMM_MEGA_MOE_USE_FP4_ACTS=1os.environ.setdefault("DG_MEGA_MOE_FP4", "1") (same setdefault pattern as the existing flags; an external DG_* override still wins). With this, the existing W4A4 config transparently routes to the packed kernel — no other code changes.

Depends on

Measured impact (4× B200, DSV4-Flash, TP4+DP4)

Routing the existing W4A4 config to the packed kernel yields, end-to-end at production concurrency (c ≥ 128): prefill input-throughput +2–7% vs the 0.1.0 W4A4 kernel (binding-controlled), mixed throughput +2%, median TPOT down to −11% at c=512. Accuracy preserved (GSM8K 0.966 / HumanEval 0.898, ≥ W4A8 baseline). Full kernel-level detail in the paired DeepGEMM PR sgl-project/DeepGEMM#44 (sgl-project/DeepGEMM#44); the measured data this routing change unlocks is reproduced below.


Appendix: measured data (the kernel this routing change selects)

All on 4× B200, DSV4-Flash (hidden=4096 / inter=2048 / E=256 / topk=6), TP=4 + DP=4. W4A4-Ours = the packed sm100_fp4_mega_moe selected by this PR's DG_MEGA_MOE_FP4; W4A4-old = the existing 0.1.0-release W4A4 kernel (sgl-project/DeepGEMM#27) — both measured through the same DeepGEMM binding (DG_FORCE_FP8FP4 toggle) so the comparison isolates the kernel from packaging; W4A8 = the FP8-acts reference (0.1.0 release).

A. End-to-end 5-concurrency sweep

Sequential, same GPU group, one server at a time; 1 measured rep/cell after a per-concurrency warmup; num_prompts raised for steady state.

Prefill (in=1024, out=1) — input throughput (tok/s) · median TTFT (ms):

conc Ours tput old tput W4A8 tput Δ tput Ours/old Ours TTFT old TTFT Δ TTFT Ours/old
1 5847 5886 5683 −0.7% 156.8 154.8 +1.4%
4 14802 14478 10095 +2.2% 289.5 292.2 −0.9%
32 54114 54755 52762 −1.2% 417.6 402.2 +3.8%
128 126557 117838 123033 +7.4% 1110.6 1145.8 −3.1%
512 133116 130299 124708 +2.2% 4575.7 4672.6 −2.1%

Mixed (in=1024, out=256) — total throughput (tok/s) · median TPOT (ms):

conc Ours tput old tput W4A8 tput Δ tput Ours/old Ours TPOT old TPOT Δ TPOT Ours/old
1 342 338 340 +1.1% 14.02 14.17 −1.1%
4 1268 1291 1274 −1.8% 14.55 14.82 −1.8%
32 8212 8055 8149 +2.0% 17.82 18.11 −1.6%
128 24982 24544 24204 +1.8% 22.40 22.63 −1.0%
512 51295 50201 50363 +2.2% 33.50 37.81 −11.4%

The win surfaces at production concurrency (c ≥ 128) where mega-MoE's share is largest; low concurrency is within the e2e noise band (decode is ~82% the shared sm100_tf32_hc_prenorm_gemm, ~8% mega-MoE).

B. NCU Speed-of-Light (1-rank isolation, binding-controlled, ntok=512)

metric W4A4-Ours (packed, bk256) W4A4-old (fp8_fp4, bk128) W4A8 (fp8 acts)
Duration (µs) 537.6 712.5 753.2
Compute (SM) % 33.6 28.6 42.5
Memory % 85.2 64.3 61.4
DRAM throughput (TB/s) 6.54 4.93 4.71
L2 % 56.1 42.4 40.3
Achieved occupancy % 22.5 22.4 22.5
Registers/thread 128 128 128

Ours +24.6% vs W4A4-old, +28.6% vs W4A8 (1-rank). Same occupancy/regs → the gain is the data path: block_k=256 sustains 85% memory SOL / 6.54 TB/s DRAM vs W4A4-old's 64% / 4.93. 1-rank over-weights the memory regime; the e2e sweep above is the realistic magnitude.

C. Accuracy (eval_v2, DSV4-Flash)

eval W4A4-Ours W4A4-old W4A8
GSM8K 0.966 0.962 0.962
HumanEval (pass@1) 0.898 0.888 0.887

At or above both baselines on every eval — the routing change preserves accuracy.

@github-actions github-actions Bot added documentation Improvements or additions to documentation dependencies Pull requests that update a dependency file labels Jun 14, 2026

@gemini-code-assist gemini-code-assist Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Code Review

This pull request updates the sgl-deep-gemm dependency from version 0.1.0 to 0.1.3 in the Dockerfile and Python dependencies. It enables routing to DeepGEMM's packed FP4xFP4 mega-MoE kernel (sm100_fp4_mega_moe) by setting the DG_MEGA_MOE_FP4 environment variable when FP4 activations are enabled. Additionally, the documentation and code comments have been updated to reflect this new behavior. There are no review comments to address, and I have no further feedback to provide.

Important

The consumer version of Gemini Code Assist on GitHub is being sunset. Starting June 18, 2026, new organization installations will be blocked, and all code review activity will officially cease on July 17, 2026.
For more details on the timeline and next steps, please review the Help Documentation.

agent and others added 2 commits June 15, 2026 01:05
Routes the existing W4A4 mode (SGLANG_OPT_DEEPGEMM_MEGA_MOE_USE_FP4_ACTS=1
+ USE_MXF4_KIND=1) to the new packed FP4xFP4 sm100_fp4_mega_moe kernel in
sgl-deep-gemm 0.1.3: packed E2M1 operands with tcgen05.mma.kind::mxf4,
halving activation bytes and SMEM versus the FP8-acts path. No sglang code
changes are required - the env plumbing and pre-dispatch call already
support it; this PR only bumps the pinned version and updates docs.

Measured on 4x B200 (DSV4-Flash, TP4/DP4): kernel-level +8-12 percent vs
the FP8-acts (W4A8) path at ntok 512-8192 and parity with the 0.1.0
release W4A4 mode; end-to-end vs that release: prefill input throughput
+5.6 to +8.8 percent across the sweep, mixed-workload median TTFT -11.2
percent at unchanged throughput; GSM8K/HumanEval parity within 1pp.

Depends on: sgl-project/DeepGEMM PR (packed FP4xFP4 mega-MoE kernel) and
a new sgl-deep-gemm release; the 0.1.3 version here is a placeholder to
be aligned with the actual release tag.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
When SGLANG_OPT_DEEPGEMM_MEGA_MOE_USE_FP4_ACTS=1, also setdefault
DG_MEGA_MOE_FP4=1 so DeepGEMM dispatches its packed FP4xFP4 mega-MoE kernel
(reached through the same fp8_fp4_mega_moe entry point, env-gated). No other
sglang change is needed — the existing mega-MoE call path is unchanged.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

dependencies Pull requests that update a dependency file documentation Improvements or additions to documentation

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant