feat(deepep_v2): expanded (do_expand=True) prefill dispatch for DeepEP-V2 - #37261
feat(deepep_v2): expanded (do_expand=True) prefill dispatch for DeepEP-V2#37261cyhdmjzzy wants to merge 6 commits into
Conversation
b5dee26 to
230be08
Compare
57e1c12 to
888b45e
Compare
ch-wan
left a comment
There was a problem hiding this comment.
Summary
This PR unlocks DeepEP-V2's expanded (do_expand=True) prefill dispatch, replaces the ep_scatter bridge with fill_m_indices_from_psum, and folds top-k weights into down_proj's fp32 input scale. The expand/non-expand split, MiMoV2 A2A wiring, and ElasticBuffer prebuild look internally consistent with the existing decode masked path.
The main correctness risk is the new prefill token-budget math, which divides by attn_tp_size even when tokens are not scattered.
Two of the PR's stated DeepEP premises do not hold against upstream: ElasticBuffer.dispatch has mapped num_sms=0 to get_theoretical_num_sms since 2.0.0, so the "zero-SM hang" derivation reproduces DeepEP's own default; and DeepEP 2.1.0 (#674) accepts 1-D expanded weights in combine, which would remove the pre-combine weighting (fold, ue8m0 fallback, scale_expanded_rows_) entirely. The version dependency should be recorded now and a gated follow-up planned.
The two extra positionals passed to ep_scatter_from_psum are a standalone fix: main's only deepep_v2 prefill path has raised a missing-argument error since #35758. That deserves disclosure in the PR body and a unit test, since CI never reaches the wrapper.
Remaining items are style and conventions: domain logic inlined into the frozen model_runner.py, env-var naming/sentinel conventions, a duplicate startup check, and a few comments that misdescribe the contracts they sit next to.
Issue counts by severity
- bugs: 1
- suggestions: 10
- nits: 1
| view.max_prefill_tokens or 0 | ||
| ) | ||
| prefill_tokens = -(-prefill_tokens // attn_tp_size) | ||
| if prefill_tokens > capacity: |
There was a problem hiding this comment.
[bug] The new prefill budget divides max_prefill_buffer_tokens by attn_tp_size unconditionally (attn_dp_size = dp_size if enable_dp_attention else 1, so with DP-attention off attn_tp_size == tp_size). Tokens are only sharded across attn-TP ranks when DP-attention is on; without it each EP rank still dispatches the full chunk. Example: --tp 8 without --enable-dp-attention, chunked_prefill_size=2048, cap=512 now validates (ceil(2048/8)=256) and then dies in _validate_common on a 2048-token dispatch. The formula also omits attn_cp_size (derive_attention_widths uses tp / dp / cp), so it disagrees with the live get_parallel().attn_tp_size used in prebuild_deepep_v2_buffers. The new unit test only covers enable_dp_attention=True.
Suggestion: Divide only when tokens are actually scattered (DP-attention / CP). Reuse derive_attention_widths (or get_parallel().attn_tp_size) instead of recomputing tp_size // attn_dp_size. Add a case with enable_dp_attention=False, tp_size>1 that still requires the undivided chunk.
There was a problem hiding this comment.
Fixed. The division now only applies when tokens are actually scattered:
if view.enable_dp_attention or view.attn_cp_size > 1:
attn_dp_size = view.dp_size if view.enable_dp_attention else 1
attn_tp_size = max(1, view.tp_size // attn_dp_size // view.attn_cp_size)
prefill_tokens = -(-prefill_tokens // attn_tp_size)
So with DP-attention off each EP rank keeps the full chunk (no division), and the width now mirrors derive_attention_widths (tp / attn_dp / cp) including attn_cp_size, matching get_parallel().attn_tp_size used in prebuild. Added a regression case (enable_dp_attention=False, tp_size=16) that keeps the undivided 2048-token chunk and still trips the cap.
| and chunked_prefill_size > 0 | ||
| and attn_tp_size > 0 | ||
| ): | ||
| min_tokens_per_rank = chunked_prefill_size // attn_tp_size |
There was a problem hiding this comment.
[suggestion] Prefill-node prebuild uses chunked_prefill_size // attn_tp_size (floor) and ignores dynamic-chunking's 1.25x probe, while validate_deepep_v2_dispatch_token_budget uses ceil(max_prefill_buffer_tokens / attn_tp_size). Remainder tokens (2049 / 8 -> 257 vs 256) can pass this check and still exceed the ElasticBuffer cap at dispatch. The check is also redundant with the server-args validator when the two formulas agree.
Suggestion: Drop the duplicate check, or share one helper with the validator (ceil of max_prefill_buffer_tokens, same attn-TP width).
There was a problem hiding this comment.
Dropped the duplicate check. The server-args validator (validate_deepep_v2_dispatch_token_budget) is now the single source of truth for the prefill budget — ceil of max_prefill_buffer_tokens over the same attn-TP width, so the floor/1.25x-probe discrepancy is gone. prebuild no longer recomputes min_tokens_per_rank.
| return get_exec().moe.deepep_v2_mode == "hybrid" | ||
| # Multi-node needs scale-out (hybrid); direct is NVLink-only and hangs across | ||
| # nodes. Honor an explicit "hybrid" too so a single node can opt in. | ||
| return get_exec().moe.deepep_v2_mode == "hybrid" or get_parallel().nnodes > 1 |
There was a problem hiding this comment.
[suggestion] _get_allow_hybrid_mode is mode == "hybrid" or nnodes > 1, so --deepep-v2-mode direct is silently ignored on multi-node. Forcing hybrid is the right hang-avoidance default, but the CLI then does not mean what it says (the PR text also claims it remains overridable).
Suggestion: Treat nnodes > 1 as the default when mode is unset/auto, and either honor an explicit direct or reject it with a clear error that NVLink-only mode cannot run across nodes.
There was a problem hiding this comment.
Done. _get_allow_hybrid_mode now treats unset/auto as "hybrid when nnodes > 1", honors an explicit "hybrid", and rejects an explicit "direct" across nodes with a clear error instead of silently overriding it:
if mode == "direct" and nnodes > 1:
raise ValueError(
"--deepep-v2-mode direct is NVLink-only and cannot run across "
f"nodes (nnodes={nnodes}); use hybrid or leave it unset (auto)."
)
I also removed the stale "remains overridable" wording from the PR text.
| torch.cuda.synchronize() | ||
| logger.warning("DeepEP v2 expanded contig activation returned") | ||
|
|
||
| # down_proj is linear and fp8 dequant is q*scale, so folding the per-row topk |
There was a problem hiding this comment.
[suggestion] The three-line block restates why folding w into down_input_scale is exact and what ue8m0 does instead. That design rationale already lives in the PR; next to the if it is history, not a line-local constraint.
Suggestion: Keep at most one trap (ue8m0 is a power of two, so do not fold into the scale) attached to the not DEEPGEMM_SCALE_UE8M0 guard.
There was a problem hiding this comment.
Trimmed to a single line-local trap on the not-ue8m0 guard (ue8m0 scales are powers of two, so folding into the scale would break them). The design rationale stays in the PR body.
| deepep_v2_expert_alignment, | ||
| ) | ||
| ep_expand_init_m_indices_from_psum(psum_num_recv_tokens_per_expert, m_indices) | ||
| # Leave hidden_states_scale_tma_aligned at its default False: DeepEP's |
There was a problem hiding this comment.
[suggestion] The comment narrates an omitted DeepGemmRunnerInput argument (hidden_states_scale_tma_aligned left at default False) and the GEMM's follow-on tma_align_input_scale. The default is already False; the load-bearing fact is that expanded DeepEP recv scales are not in the contiguous-GEMM TMA layout.
Suggestion: One line on the recv-scale layout constraint, or drop it if _run_contiguous_gemm already always aligns when the flag is False.
There was a problem hiding this comment.
Reduced to one line stating the load-bearing constraint: expanded DeepEP recv scales are not in the contiguous-GEMM TMA layout, so the flag stays False and _run_contiguous_gemm aligns them. Dropped the narration of the default value.
| # Weight the expanded rows before combine (skipped when already folded | ||
| # into down_input's scale); combine ignores topk_weights in expand mode. | ||
| from sglang.kernels.ops.moe.ep_moe_kernels import scale_expanded_rows_ | ||
|
|
||
| scale_expanded_rows_(hidden_states, topk_weights) |
There was a problem hiding this comment.
[suggestion] "combine ignores topk_weights in expand mode" is true for DeepEP 2.0.0 (its docstring: "Not used in expand mode"), but DeepEP 2.1.0 (#674, 2026-07-03) changed the contract: combine(topk_weights=...) accepts a 1-D [num_expanded_tokens] tensor in expand mode when the buffer was built with allow_multiple_reduction=True, which is the constructor default DeepEPv2Buffer.get_buffer already inherits (csrc/elastic/buffer.hpp takes the get_shape<1> path; tests/elastic/test_ep.py:227-229). On 2.1.0 the whole pre-combine weighting (the down-scale fold, the ue8m0 fallback and scale_expanded_rows_) collapses to buffer.combine(hidden_states, handle=..., topk_weights=recv_topk_weights) with no extra kernel.
Suggestion: State the version dependency in this comment and in the PR body now, and plan a follow-up that gates on the DeepEP version and hands recv_topk_weights straight to combine. That needs a bench, since allow_multiple_reduction changes the reduction epilogue (DeepEP documents a precision / transfer-size trade-off for it).
There was a problem hiding this comment.
Recorded the version dependency in the comment and will note it in the PR body:
# DeepEP 2.0.0 combine ignores topk_weights in expand mode; 2.1.0+ (#674)
# accepts a 1-D weight there, so a follow-up can gate on version and pass
# recv_topk_weights straight to combine.
Agreed this collapses the down-scale fold + ue8m0 fallback + scale_expanded_rows_ on 2.1.0. I'll keep it as a gated follow-up rather than in this PR, since allow_multiple_reduction changes the reduction epilogue and DeepEP documents a precision/transfer-size trade-off, so it needs its own bench.
| decode_runner_captured = ( | ||
| self.decode_cuda_graph_runner is not None | ||
| and not isinstance(self.decode_cuda_graph_runner, EagerRunner) | ||
| ) | ||
| if not decode_runner_captured: | ||
| prebuild_deepep_v2_buffers( | ||
| model=self.model, | ||
| disaggregation_mode=get_disagg().disaggregation_mode, | ||
| chunked_prefill_size=get_schedule().chunked_prefill_size, | ||
| attn_tp_size=get_parallel().attn_tp_size, | ||
| ) |
There was a problem hiding this comment.
[suggestion] model_runner.py is a frozen orchestration-only file (.claude/skills/large-class-style/SKILL.md §1.3-1.5). decode_runner_captured is domain knowledge, not coordination: it encodes that a captured decode graph already built the ElasticBuffer and that EagerRunner means no capture happened. That reasoning belongs with the DeepEP-v2 setup code, and inlining it here also removes an override point for downstream forks.
Suggestion: Move the gate into the collaborator, e.g. maybe_prebuild_deepep_v2_buffers(model=..., decode_cuda_graph_runner=self.decode_cuda_graph_runner, disaggregation_mode=..., chunked_prefill_size=..., attn_tp_size=...) doing the EagerRunner check itself, and leave init_cuda_graphs with the single delegate call (or a maybe_init_deepep_v2_buffers helper method).
There was a problem hiding this comment.
Moved. The EagerRunner gate now lives in the collaborator as maybe_prebuild_deepep_v2_buffers(model=..., decode_cuda_graph_runner=...), which does the capture check itself. model_runner.py keeps only the single delegate call, so the frozen file stays orchestration-only and downstream forks get their override point back.
| # Prefill: True uses DeepEP's expanded layout (skips ep_scatter); False uses | ||
| # the non-expand ep_scatter/ep_gather path. | ||
| SGLANG_DEEPEP_V2_PREFILL_DO_EXPAND = EnvBool(True) | ||
| # GPU-side ElasticBuffer barrier timeout (seconds); raise it so idle ranks | ||
| # tolerate the first-request JIT compile. 0 keeps DeepEP's default (100s). | ||
| SGLANG_DEEPEP_V2_GPU_TIMEOUT_SECS = EnvInt(0) |
There was a problem hiding this comment.
[suggestion] Three points against .claude/skills/env-var-conventions/SKILL.md:
SGLANG_DEEPEP_V2_PREFILL_DO_EXPANDis a boolean without a verb category (Rule 4:ENABLE_/DISABLE_/USE_...).SGLANG_DEEPEP_V2_ENABLE_PREFILL_EXPAND(defaultTrue) reads at the call site as "if enabled". Renaming later costs an alias entry, so it is cheapest now.SGLANG_DEEPEP_V2_GPU_TIMEOUT_SECS = EnvInt(0)uses0as an "unset" sentinel; Rule 2 asks forEnvInt(None)when unset must be distinguishable from a real value.- The knob only raises
num_gpu_timeout_secs. Prefill dispatch runs withdo_cpu_sync=True, so the wait that actually times out while a peer rank is JIT-compiling is DeepEP's CPU-sidenum_cpu_timeout_secs(300 s default). If the goal is surviving first-request stalls, that timeout matters at least as much, or the comment should say why only the GPU barrier needs raising.
Also, this file's convention is bare declarations; the rationale for both flags is already in the PR body, so the two comment blocks can go.
There was a problem hiding this comment.
All three addressed:
(1) Renamed to SGLANG_DEEPEP_V2_ENABLE_PREFILL_EXPAND (default True), reads as "if enabled" at the call site.
(2)+(3) Rather than fix the sentinel, I removed SGLANG_DEEPEP_V2_GPU_TIMEOUT_SECS entirely and no longer pass num_gpu_timeout_secs to ElasticBuffer (it uses DeepEP's default). You're right that the GPU barrier wasn't the right lever. The stall it papered over comes from ranks finishing their lazy DeepGEMM warmup at different times, so the first-layer dispatch barrier waits past the default. The correct fix is to remove that skew via DeepGEMM precompilation (pre-run sglang.compile_deep_gemm, or set SGLANG_JIT_DEEPGEMM_PRECOMPILE=0 at deploy time) so every rank hits a warm cache and the default timeout is ample — no knob needed.
Also dropped both comment blocks per the bare-declaration convention.
| # psum is the alignment-padded exclusive prefix sum; padding rows get a real | ||
| # expert id (not a sentinel) so every m_indices entry stays valid. | ||
| e = tl.program_id(0) | ||
| prev_end = tl.load(psum_ptr + e - 1, mask=e > 0, other=0) | ||
| start = ((prev_end + ALIGN - 1) // ALIGN) * ALIGN | ||
| end = tl.load(psum_ptr + e) | ||
| seg_end = ((end + ALIGN - 1) // ALIGN) * ALIGN | ||
| if e == num_local_experts - 1: | ||
| seg_end = total_rows |
There was a problem hiding this comment.
[nit] Two inaccuracies that will mislead the next editor. (1) The psum is not an exclusive prefix sum: DeepEP's handle docstring defines expand-mode psum[i] = align(psum[i-1]) + unaligned_count_i (inclusive, only the earlier experts aligned), which is exactly why both start and seg_end are rounded up here. (2) DeepEP sizes the expanded recv_x as the sum of aligned per-expert counts under do_cpu_sync=True (num_recv_tokens_per_expert_list holds aligned counts and buffer.hpp accumulates them into num_expanded_tokens), so total_rows == align(psum[-1]) and the last-expert seg_end = total_rows branch is a no-op. If it is meant to guard a capacity-sized buffer (do_cpu_sync=False), labelling that tail with a real expert id would make DeepGEMM compute the whole buffer instead of skipping -1 rows.
Suggestion: Fix the comment, and replace the branch with the total_rows % ALIGN == 0 assert the old kernel had.
There was a problem hiding this comment.
Both fixed. The comment now says the psum is DeepEP's inclusive per-expert count (psum[i] = align(psum[i-1]) + count_i, only earlier experts aligned), which is why both start and seg_end round up. And I replaced the no-op last-expert branch with the total_rows % expert_alignment == 0 assert the old kernel had, since do_cpu_sync=True sizes recv_x as the sum of aligned counts (total_rows == align(psum[-1])).
…P-V2 MoE Add a higher-performance prefill path for the DeepEP-V2 (ElasticBuffer) MoE A2A backend that uses DeepEP's expanded dispatch layout, gated by SGLANG_DEEPEP_V2_PREFILL_DO_EXPAND (default True). Community's deepep_v2 already had a latent "expanded non-masked" prefill branch but it was locked off at dispatch and used slow fallbacks; this change opens and optimizes it, while leaving the decode (expand + masked slab) path untouched. Core idea --------- Use DeepEP's dispatch_copy_epilogue (which performs the per-(token,expert) expansion inline with the all-to-all TMA copy) instead of a separate ep_scatter full-hidden copy, and push the un-expand + reduce down into combine. The sglang side is left with only two light ops -- build m_indices, and weight the rows -- so the expanded path saves one full HBM round trip (the ep_scatter_2 read of num_recv*H plus write of sum(cnt_e)*H) versus every non-expand path. Measured on DeepSeek-class shapes this is ~51us/layer net (saving ~58us of ep_scatter_2 for ~6.5us more in row weighting) plus two fewer host readbacks. Changes on top of the community deepep_v2 (all within the expand branch) ------------------------------------------------------------------------ - token_dispatcher/deepep_v2.py: unlock expand at dispatch. Community bound use_masked = use_expand_layout = not is_extend, pinning prefill to non-expand. Split them: use_masked = is_decode; use_expand_layout = is_decode or _prefill_expand_enabled, so prefill runs expanded but never masked. Also quantize the expanded recv scale column-major only under ue8m0, so non-ue8m0 stays row-major for the GEMM's own tma_align. - moe_runner/deep_gemm.py (pre-permute, expanded non-masked): replace ep_expand_init_m_indices_from_psum + torch.full(-1) sentinel with fill_m_indices_from_psum, labeling every row (incl. alignment padding) with a real expert id straight from the device psum -- no dependence on DeepGEMM's negative-value skip. Leave hidden_states_scale_tma_aligned False so the contiguous GEMM still runs tma_align_input_scale on the expanded recv scale. - moe_runner/deep_gemm.py (post-permute, expanded non-masked): replace hidden_states * topk_weights.unsqueeze(-1) (stride-0 broadcast -> scalar elementwise) with scale_expanded_rows_, a coalesced triton kernel (~2.8x). - ep_moe_kernels.py: add fill_m_indices_from_psum and scale_expanded_rows_. - model_runner.py + moe_ep_setup.py: prebuild the ElasticBuffer at deployment time (in init_cuda_graphs, after capture and before serving), skipping it when decode graph capture already built it, so the first request does not pay the ~2GB symmetric alloc + cross-rank NCCL barrier lazily. Correctness / robustness ------------------------- - SGLANG_DEEPEP_V2_PREFILL_DO_EXPAND=False strictly falls back to the community path: prefill runs do_expand=False (ep_scatter_from_psum + ep_gather), byte-for- byte equivalent to the pre-change community behavior. All expand changes live under use_expand_layout / is_expanded branches that are not entered when the flag is off, so the non-expand and decode paths are untouched. - Derive num_sms from get_theoretical_num_sms when SGLANG_DEEPEP_V2_NUM_SMS is 0 (dispatch has no zero-SM fallback, unlike combine, and would hang the NVLink barrier otherwise). - Use ElasticBuffer hybrid mode automatically when nnodes > 1 (direct is NVLink-only and hangs across nodes); still overridable via --deepep-v2-mode. - models/mimo_v2.py + arg_groups/moe_hook.py: include is_deepep_v2() in the MiMoV2 EP / _enable_a2a_moe branches and add MiMoV2ForCausalLM / MiMoV2FlashForCausalLM to the validated architectures; the prefill token-budget check divides by attn_tp_size (a per-DP chunk scatters across the DP group's attn-TP ranks). - environ.py: add SGLANG_DEEPEP_V2_PREFILL_DO_EXPAND and SGLANG_DEEPEP_V2_GPU_TIMEOUT_SECS. The per-rank dispatch-token cap keeps the existing SGLANG_DEEPEP_V2_NUM_MAX_DISPATCH_TOKENS_PER_RANK unchanged.
Fold the expand-mode topk weighting into down_proj's fp32 input scale instead of scaling the bf16 down_output before combine. down_proj is linear and fp8 dequant is fp8_q * scale, so scaling the per-row input scale is numerically exact and shrinks the ~112us pass over the [expanded, hidden] bf16 tensor to ~4us over the [expanded, hidden/128] scale. ue8m0 scale is a power of two, so it keeps weighting down_output via scale_expanded_rows_.
…he topk-weight fold Fix a latent bug in the community's do_expand=False prefill path: the ep_scatter_from_psum call to _fwd_kernel_ep_scatter_2 was missing the expert_start / num_experts positional args the kernel gained, so any SGLANG_DEEPEP_V2_PREFILL_DO_EXPAND=False run crashed with a TypeError. Pass expert_start=0 and num_experts, matching the sibling ep_scatter call site. Always fold the expand-mode topk weight into down_proj's fp32 input scale (the best-performing path) and drop the SGLANG_DEEPEP_V2_FUSE_WEIGHT_INTO_SCALE knob; ue8m0 still weights down_output via scale_expanded_rows_. scale_expanded_rows_ is now stride-aware so it handles the column-major input scale. Trim comments to the load-bearing ones (why-not-obvious constraints and traps).
- fix attn-tp prefill token rounding in moe_hook buffer sizing - rename SGLANG_DEEPEP_V2_PREFILL_DO_EXPAND to ENABLE_PREFILL_EXPAND - drop the GPU barrier timeout knob; rely on DeepGEMM precompilation - pass num_sms straight through to ElasticBuffer - reject explicit direct mode across nodes - tighten scale_expanded_rows_ contract to 1-D row weights - simplify DeepEP v2 buffer prebuild delegation and comments
Without DP-attention or CP every EP rank dispatches the whole prefill chunk, so the per-rank budget must not be divided by tp_size. The earlier code divided unconditionally, letting an over-capacity chunk (2048 > cap) pass as 2048/16. Assert the pure-TP case still raises.
16d6095 to
d9efaaf
Compare
|
Thanks for the very thorough review. I've addressed all 12 points; a per-item summary:
Details inline. The GPU-side timeout knob is removed entirely rather than fixed — see the env-var thread for the reasoning. |
|
https://aie-shared-objects.cnbj1.mi-fds.com/aie-shared-objects/prefill-DeepEP-V2-weighted-before-down_proj.nsys-rep |
| if mode == "direct": | ||
| if nnodes > 1: | ||
| raise ValueError( | ||
| "--deepep-v2-mode direct is NVLink-only and cannot run across " |
There was a problem hiding this comment.
[suggestion] Prior finding 3 is only half-fixed. Explicit --deepep-v2-mode direct now raises on nnodes > 1 (good), but the comment, the error text ("leave it unset (auto)"), and the final return nnodes > 1 still describe an auto/unset hybrid default that does not exist. deepep_v2_mode is Literal["direct", "hybrid"] and defaults to "direct", so the fallback branch is dead, and following the error's "unset (auto)" advice still raises.
Suggestion: Drop the auto/unset wording and the dead return nnodes > 1. The error should tell the user to pass --deepep-v2-mode hybrid. If multi-node should auto-select hybrid, that belongs in resolution (change the default / add "auto"), not in this helper.
There was a problem hiding this comment.
You're right — deepep_v2_mode is Literal["direct","hybrid"] with default "direct", so the fallback and the "unset (auto)" wording described a state that can't occur. Dropped the dead return nnodes > 1; explicit direct on nnodes>1 now raises asking for --deepep-v2-mode hybrid, otherwise it returns mode == "hybrid". Auto-selection belongs in resolution (default/"auto"), which I can do as a follow-up if wanted.
| deepep_v2_expert_alignment, | ||
| ) | ||
| ep_expand_init_m_indices_from_psum(psum_num_recv_tokens_per_expert, m_indices) | ||
| # Expanded DeepEP recv scales are not in the contiguous-GEMM TMA layout; |
There was a problem hiding this comment.
[suggestion] Prior finding 5 is still present after the field revert. The comment narrates that hidden_states_scale_tma_aligned is omitted and stays False. _run_contiguous_gemm already TMA-aligns when DEEPGEMM_NEED_TMA_ALIGNED_SCALES is set (and that flag is False under UE8M0), so the omitted kwarg is the dataclass default, not a local invariant that needs restating.
Suggestion: Delete the comment. If the default-False choice is the trap, a one-liner on the DeepGemmRunnerInput field is the right home, not a narration of a dropped kwarg.
There was a problem hiding this comment.
Deleted. You're right — with the field reverted, hidden_states_scale_tma_aligned=False is just the dataclass default, and _run_contiguous_gemm already TMA-aligns when DEEPGEMM_NEED_TMA_ALIGNED_SCALES is set, so the comment only restated the default. Removed it rather than moving it, since the field name is self-describing and the alignment is generic, not a per-field trap.
| "deepep_v2_weight_prefused", False | ||
| ): | ||
| # Weight the expanded rows before combine (skipped when already folded | ||
| # into down_input's scale). DeepEP 2.0.0 combine ignores topk_weights in |
There was a problem hiding this comment.
[suggestion] Prior finding 4 is trimmed around the scale-fold site, but post-permute still embeds design history and an unowned follow-up: DeepEP 2.0.0 vs 2.1.0+ (#674), and "a follow-up can gate on version". That belongs in the PR body, not next to the call.
Suggestion: Keep at most the live constraint ("combine ignores topk_weights in expand mode, so rows must be weighted here unless already folded into the down-proj scale"). Drop the version changelog and the unowned follow-up, or make it TODO(<owner>): if it is planned work.
There was a problem hiding this comment.
Trimmed to just the live constraint (combine ignores topk_weights in expand mode, so rows are weighted here unless already folded into the down-proj scale). Dropped the 2.0.0-vs-2.1.0 changelog and the unowned follow-up from the call site — the version dependency stays in the PR body, and I'll track the 2.1.0 gating there rather than as an inline note.
- _get_allow_hybrid_mode: deepep_v2_mode is Literal[direct, hybrid] with no auto/unset state, so drop the dead nnodes>1 fallback and the misleading "unset (auto)" wording; explicit direct across nodes now asks for hybrid. - deep_gemm: drop the TMA-layout comment that just restated the dataclass default; the contiguous GEMM already aligns recv scales when needed. - deep_gemm: trim the pre-combine weighting comment to the live constraint, dropping the DeepEP version changelog and the unowned follow-up note. - environ: drop the leftover rationale block on ENABLE_PREFILL_EXPAND to match the bare-declaration convention.
Motivation
Add a faster prefill path for the DeepEP-V2 (ElasticBuffer) MoE A2A backend
(
--moe-a2a-backend deepep_v2) that uses DeepEP's expanded dispatch layout, gated bySGLANG_DEEPEP_V2_PREFILL_DO_EXPAND(defaultTrue). DeepEP-V2 usesElasticBuffer:one buffer serves both prefill and decode, switching behavior at runtime via
do_cpu_sync/do_expand.The existing deepep_v2 backend already had an "expanded non-masked" prefill branch, but
it was locked off at dispatch and used a slower fallback. This change opens and
optimizes that branch so prefill runs on the expanded layout, while leaving the decode
(expand + masked slab) path unchanged. The goal is a faster prefill MoE all-to-all
with identical model outputs. In production PD-disaggregated (prefill/decode split)
serving, decode runs on the V1
low_latencybackend in a separate process (rationale inModifications → Why decode stays on V1); this PR's optimization deliberately targets
prefill only.
Modifications
Core
All changes live inside the expand branch; the non-expand and decode paths are
untouched.
token_dispatcher/deepep_v2.py— unlock expand at dispatch. The backend bounduse_masked = use_expand_layout = not is_extend, pinning prefill to non-expand; splitthem into
use_masked = is_decodeanduse_expand_layout = is_decode or _prefill_expand_enabledso prefill can run expandedbut never masked. Quantize the expanded recv scale column-major only under ue8m0
(non-ue8m0 stays row-major for the GEMM's own
tma_align). Derivenum_smsfromget_theoretical_num_smswhenSGLANG_DEEPEP_V2_NUM_SMSis 0 (dispatch has nozero-SM fallback and would hang the NVLink barrier otherwise); use ElasticBuffer
hybrid mode automatically when
nnodes > 1(direct is NVLink-only), still overridablevia
--deepep-v2-mode.kernels/ops/moe/ep_moe_kernels.py— addfill_m_indices_from_psum(label rowsfrom the device
psum, no cumsum / H2D / sentinel) andscale_expanded_rows_(in-place row weighting, stride-aware so it also handles the column-major fp8 input
scale).
moe_runner/deep_gemm.py— the expanded non-masked pre/post permute use the twonew kernels; leave
hidden_states_scale_tma_alignedFalse so the contiguous GEMM stillruns
tma_align_input_scaleon the expanded recv scale. The expand path applies thetop-k weight before combine (which ignores
topk_weightsin expand mode) by foldingthe weight into down_proj's fp32 input scale — down_proj is linear and fp8 dequant is
fp8_q * scale, so a per-row scale factor is numerically exact, shrinking a ~112 µspass over the
[expanded, H]bf16 tensor to ~4 µs over the[expanded, H/128]scale.(Under ue8m0 the scale is a power of two and multiplying by the weight would need
re-rounding, so that case instead weights
down_outputwithscale_expanded_rows_before combine.)
model_runner.py+model_runner_components/moe_ep_setup.py— prebuild theElasticBuffer at deployment time (in
init_cuda_graphs, after capture and beforeserving), skipping it when decode graph capture already built it, so the first request
does not pay the ~2 GB symmetric alloc + cross-rank NCCL barrier.
Wiring
models/mimo_v2.py— includeis_deepep_v2()in the EP /_enable_a2a_moebranches. This is required: without it
_enable_a2a_moestaysFalse, MoE runsforward_normalwhose extratensor_model_parallel_all_reduce(whichforward_deepepskips, since
dispatcher.combinealready reduces across ranks) hangs the firstcross-node request.
arg_groups/moe_hook.py— addMiMoV2ForCausalLM/MiMoV2FlashForCausalLMtothe validated deepep_v2 architectures; divide the prefill token-budget check by
attn_tp_size(a per-DP chunk scatters across the DP group's attn-TP ranks beforedispatch, so a per-DP chunk of N needs
N / attn_tp_sizeper EP rank, not N).environ.py— addSGLANG_DEEPEP_V2_PREFILL_DO_EXPANDandSGLANG_DEEPEP_V2_GPU_TIMEOUT_SECS. The per-rank dispatch-token cap keeps the existingSGLANG_DEEPEP_V2_NUM_MAX_DISPATCH_TOKENS_PER_RANKunchanged.SGLANG_DEEPEP_V2_PREFILL_DO_EXPAND(default: on)Prefill dispatch has two layouts; the only differing dispatch argument is
do_expand(do_cpu_syncstaysTruein both, so shapes are exact and neither usesCUDA graph):
do_expand=False— dispatch returnsrecv_x=[num_recv, H], one row perdeduplicated token, not grouped by expert. A bridge kernel
ep_scatter(inherited from V1) physically expands it into per-(token, expert) rows grouped by
expert for the contiguous grouped GEMM.
ep_scatteris two triton kernels: a lightlayout kernel (
ep_scatter_1, ~1.7 µs) and a heavy full-hidden copy (ep_scatter_2,~58 µs), plus one H2D of the per-expert counts. This path's operator chain is
identical to the community's original non-expand design (
ep_scatter→ contiguousGEMM →
ep_gather), for easy A/B comparison againstdo_expand=True(here weightingis folded into
ep_gather, with no down_proj scale folding).do_expand=True(default) — DeepEP'sdispatch_copy_epilogueproduces theper-(token, expert) rows grouped by expert within dispatch's own TMA copy, so
ep_scatteris not needed. The heavyep_scatter_2(~58 µs) and that H2D bothdisappear; the light layout kernel is replaced by
fill_m_indices_from_psum(~1.5 µs,reusing DeepEP's device-side
psuminstead of recomputing a cumsum). It removes onefull HBM round-trip per MoE layer at no accuracy cost.
dispatch_copy_epilogue_implproduces different outputs based on thedo_expandtemplate parameter (nsys traces):
V1 dispatch:

V2 dispatch (do_expand=False):

V2 dispatch (do_expand=True):

V2 combine (do_expand=True, topk weighted before combine):

V2 combine (do_expand=True, topk weight folded into down_proj's input scale) —

_fwd_kernel_scale_expanded_rowsmoved from before combine to before down_proj, actingonly on the small fp8 input scale:
In expand mode ElasticBuffer's combine refuses
topk_weights(it does the un-expand andreduction itself inside the kernel), so the top-k weighting must be done by SGLang before
combine. Instead of scaling down_proj's bf16 output, this design folds each row's
routing weight
winto down_proj's fp32 input scaledown_input_scale=[M, I/128](
[~18432, 16]) before the GEMM: down_proj is linear and fp8 dequant isvalue = q × scale, andwis one positive scalar per row, so it merges into the inputscale, the output carries the weight automatically, and the result is numerically exact.
The weighted tensor shrinks from
[M, H]bf16 ([~18432, 6144], ~453 MB) to[M, I/128]fp32 (~1.2 MB), and the multiply folds into the pre-GEMM handling thatalready runs — measured ~112 µs → ~4 µs (ue8m0-packed scales fall back to weighting
the output).
Default is
Truebecause it is strictly less work thanFalsewith identical results;set
0to fall back to the validated non-expanded path for A/B comparison. "Less work"here means
do_expand=Trueversusdo_expand=False:do_expand=Falseneedsep_scatterto physically expand the deduplicatedrecv_xinto the per-expert layout; its heavy
ep_scatter_2reads and writes all of hiddenonce (a ~58 µs full HBM round-trip), plus one H2D copying the per-expert counts
back to the device.
do_expand=Truelets DeepEP'sdispatch_copy_epiloguedo the expansion withindispatch's own TMA copy, so
ep_scatter_2(~58 µs) and that H2D both disappear,leaving only a ~1.5 µs
fill_m_indices_from_psumto label rows. That removes one fullhidden HBM round-trip per MoE layer (fewer kernels, one fewer large-tensor pass), with
the GEMM and expert compute unchanged and results identical.
Why decode stays on V1
low_latency(not DeepEP-V2)This PR's optimization targets prefill only. The deepep_v2 decode (expand + masked slab)
path is left untouched; in production PD-disaggregated serving, decode runs on V1's
low_latencymode in a separate process.Decode requires a masked GEMM (fixed-height slab
[num_local_experts, max_m, H]) to beCUDA-graph-capturable:
low_latencydirectpacked_recv_x = [num_local_experts, num_ranks×cap, H]— fixed-height gride × max_mexpand_to_masked_slab(before GEMM) +masked_slab_to_expand(after)V1's
low_latencyoutput is natively the masked-GEMM input, so decode needs no bridgekernels. V2's
directoutput must be reshaped into a slab and back — two extra HBMround-trips every layer, every step, with no offsetting benefit in SGLang (unlike vLLM,
whose V2 decode gain comes from a decode-only activation kernel SGLang does not use).
Accuracy Tests
Verified with
--disaggregation-mode null(single process; deepep_v2 drives bothprefill and decode MoE dispatch — a stricter check than the production split), GSM8K
5-shot, 1319 questions, greedy decoding, MiMoV2, tp16/dp2:
deepep_v1(baseline)deepep_v2The 0.006 gap (~8 of 1319 questions) is within greedy-decoding noise, so DeepEP-V2 has
no accuracy regression. (Production uses V2 for prefill + V1
low_latencyfordecode; both paths are independently validated.)
Speed Tests
End-to-end prefill throughput, MiMoV2, TP8/EP16/DP2, chunk 32k per DP,
bench_serving'sInput token throughput (tok/s). "single DP throughput" = all-throughput / 2. A
three-way comparison: V1 baseline, the community's existing deepep_v2 (
do_expand=False),and this PR's final design (
do_expand=Truewith the topk weight folded into down_proj'sinput scale, weighted-before-down_proj).
Δis the gain over V1.do_expand=False(tok/s)Two takeaways:
do_expand=False) is already ~2–8% faster thanV1; this PR pushes the gain to ~4–15% (~13–15% at short/medium sequence lengths,
narrowing to ~4–6% at very long ones as attention's O(L²) cost dominates and the MoE
all-to-all share shrinks).
(e.g. 8k: 27028 → 29107, +7.7%; 512k: 12767 → 13028, +2.0%), from the
ep_scatter_2full-hidden copy that
do_expand=Trueremoves plus folding the topk weight intodown_proj's input scale (shrinking a ~112 µs large-tensor weighting to ~4 µs). The
rationale for each is in the
SGLANG_DEEPEP_V2_PREFILL_DO_EXPANDsection above and inweighted-before-down_proj.md.Checklist
Notes for reviewers
Runtime environment
Developed and validated on (2-node × 8-GPU, Hopper-class, ~140 GB HBM per GPU):
cu12)ncclGinRequest_t/ginTrafficClass) needs 2.30.x; the bundled 2.29.x lacks these symbols, so the image overrides everylibnccl.so(including torch's) with 2.30.7deep_ep.ElasticBuffer), built against torch 2.13 / NCCL 2.30.7,cp312Communication backend: DeepEP-V2 (ElasticBuffer) communicates over NCCL GIN
(GPU-Initiated Networking, reusing PyTorch's NCCL communicator), not NVSHMEM/IBGDA —
the V2 path has no NVSHMEM dependency (NVSHMEM is the DeepEP-V1 transport).
NCCL_CUMEM_ENABLE=1is required by ElasticBuffer and is set automatically. Single nodeuses
--deepep-v2-mode direct(NVLink); multi-node needshybrid(auto-selected whennnodes > 1, or set explicitly).CI States
Latest PR Test (Base): ❌ Run #34617090627
Latest PR Test (Extra): ❌ Run #34617089426
Latest PR Test (AMD ROCm 10): ❌ Run #34617090277