feat(deepep_v2): validate Qwen3.5 MoE for --moe-a2a-backend deepep_v2 - #39408
Open
whn09 wants to merge 1 commit into
Open
feat(deepep_v2): validate Qwen3.5 MoE for --moe-a2a-backend deepep_v2#39408whn09 wants to merge 1 commit into
whn09 wants to merge 1 commit into
Conversation
Qwen3_5MoeForConditionalGeneration is rejected by validate_deepep_v2_model_architecture, so deepep_v2 cannot be used on any Qwen3.5 MoE checkpoint. No model-side change is needed: Qwen3.5's MoE layer is Qwen2MoeSparseMoeBlock, which already routes deepep_v2 to _forward_deepep (qwen2_moe.py:747), drives the EP geometry off the same predicate (:397), skips the post-experts all-reduce via should_skip_post_experts_all_reduce (:811) and treats v2 as a deepep-class backend for the shared-expert 1/ep_size pre-scale (:443, :459). Both conditions the gate's comment names therefore hold, so this is a whitelist entry only. Verified on Qwen/Qwen3.5-35B-A3B-FP8, 1 node x 8 B300, TP=EP=8, against the DeepEP v1 path at matched geometry, chunk size and MoE runner: 0/12 top-1 mismatches, 0 top-5 set or order differences, 10/12 prompts bit-identical, and the two that differ do so by less than the same arm's own run-to-run variance (tl.atomic_add slot claiming in the scatter kernels, which is also why deepep_v2 refuses --enable-deterministic-inference). The text-only Qwen3_5MoeForCausalLM is deliberately not claimed; it shares the same MoE block but has not been run. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Motivation
--moe-a2a-backend deepep_v2is rejected outright on every Qwen3.5 MoE checkpoint, becausevalidate_deepep_v2_model_architecturedoes not list the architecture:Unlike #39080, no model-side change is needed here. Qwen3.5's MoE layer is
Qwen2MoeSparseMoeBlock(models/qwen3_5.py:952and:1204), and that block is alreadyv2-aware on
main@ 66c7bc8. The gate's own comment states the two conditions anarchitecture has to satisfy — "take the A2A MoE path and skip post-expert all-reduce" —
and both hold without touching the model file:
main@ 66c7bc8)qwen2_moe.py:747-751—forward()returns_forward_deepep()whenis_deepep_v2()qwen2_moe.py:397-401— the same predicate drives theep_size/num_experts/top_kblockqwen2_moe.py:811-818— the final all-reduce is guarded byshould_skip_post_experts_all_reduce(is_tp_path=True)1/ep_sizepre-scale not double-appliedqwen2_moe.py:443and:459—is_deepep_class_backend()already covers v2So this PR is a whitelist entry and nothing else.
Modifications
arg_groups/moe_hook.py:"Qwen3_5MoeForConditionalGeneration"added tovalidate_deepep_v2_model_architecture'svalidated_architectures, after"Qwen3MoeForCausalLM".Deliberately not claiming the text-only
Qwen3_5MoeForCausalLM(qwen3_5.py:1962, plusthe wrapper at
qwen3_5_text.py:223). It uses the sameQwen2MoeSparseMoeBlockand wouldvery likely work, but I have only run the ConditionalGeneration checkpoint, and this
whitelist means "someone has actually run it". Same for
InternS2PreviewForConditionalGeneration, which subclasses the Qwen3.5 model but registersits own architecture string.
Accuracy Tests
Numerics against the DeepEP v1 path (
--moe-a2a-backend deepep) on the same weights andprompts. v1 is the right control: it is the established EP dispatcher, so a difference is
attributable to the dispatcher rather than to EP-vs-not. Both arms at matched
TP=EP=8, matched--chunked-prefill-size 4096, matched--moe-runner-backend deep_gemm,so the dispatcher is the only variable.
Qwen/Qwen3.5-35B-A3B-FP8(256 experts, top-8,hidden_size=2048, 128×128 blockwise FP8),1 node × 8 B300, greedy, first-generated-token top-5 logprobs over a fixed 12-prompt set:
The two differing prompts are inside this path's own run-to-run noise, and it is worth
measuring that noise rather than assuming a threshold. Capturing the same v1 server
twice reproduces the same pattern with larger deltas on the same two prompts:
That control also produced a top-5 order change that the v1-vs-v2 comparison did not. The
mechanism is in-tree: both scatter kernels claim each expert's destination slots with
tl.atomic_add(kernels/ops/moe/ep_moe_kernels.py:1120), so token order within an expertvaries per run and DeepGEMM's accumulation order follows. sglang says as much itself —
deepep_v2refuses--enable-deterministic-inferenceon the grounds that it "does notforward deterministic=True to ElasticBuffer, so deterministic sorting remains disabled"
(
moe_hook.py:328-334). Short prompts have little per-expert contention and staybit-reproducible; long ones do not. A fixed absolute threshold applied to this comparison
measures that, not the dispatcher.
Per-prompt criterion used instead:
|Δ|(v1, v2) <= max(1e-3, |Δ|(v1, v1')). 12/12 pass,with 0 top-1 mismatches.
Sanity check that the flag is actually doing something after the patch — this line appears on
all 8 ranks, and does not appear at all on unpatched
main(the server never gets that far):Reproducing
deepep_v2requires a 128×128 blockwise FP8 checkpoint with dynamic activation scaling(
_validate_deepep_v2_quant_method,fused_moe_triton/layer.py:242), and there is noload-time conversion path, so a BF16 Qwen3.5 checkpoint is rejected with
weight_block_size=Noneand--quantization fp8gives per-tensor scales. Use the FP8checkpoint:
The env var matters: its default of 128 is below the prefill budget, and the launch fails with
required=16384, capacity=128before reaching anything model-specific.--deepep-v2-mode hybridis required on this fabric (EFA);directkeepsginStrongSignalsRequired=true,which EFA's GIN provider does not implement.
Speed Tests and Profiling
Not included. I have a single 8-GPU node, so both arms would run entirely over NVLink and the
resulting throughput numbers would say nothing about the a2a path at the scale
deepep_v2targets — I would rather not present a number that reads as a benchmark of it. For a
whitelist widening, the accuracy comparison above is the load-bearing evidence. Multi-node
deepep_v2throughput for a different architecture is in #39080.Checklist
deepep_v2server on a blockwise-FP8 MoE checkpoint, which CI does not have.Review and Merge Process
I cannot add the
run-cilabel myself, so this needs someone with write access to start CI.Independent of this change, but a
deepep_v2server needs it: #37211 —ep_scatter_from_psumomits two positional kernel arguments, so any
deepep_v2prefill dies withTypeError: dynamic_func() missing 2 required positional argumentsbefore the effect of thiswhitelist entry is observable. All numbers above were measured with that fix applied. Note
there are now three open PRs for the same one-line fix — #37211, #38518, #39399 — worth
deduplicating.
Touches the same tuple as #39080 (
"MiMoV2ForCausalLM"); the two conflict textually and aretrivially rebasable in either order.
🤖 Generated with Claude Code
CI States
Latest PR Test (Base): ❌ Run #34833429638
Latest PR Test (Extra): ❌ Run #34833429452
Latest PR Test (AMD ROCm 10): ❌ Run #34833429450