Add Qwen3.5-35B-A3B (support vision and text input)#180
Open
qingzwang wants to merge 5 commits into
Open
Conversation
Qwen/Qwen3.5-35B-A3B is the MoE flagship of the qwen3_5 family — 35 B total params, ~3 B activated per token via top-8 routing over 256 experts plus a sigmoid-gated shared expert. Same hybrid attention stack as the dense siblings: [3 gated DeltaNet + 1 full GQA] × 10 = 40 layers, head_dim=256, partial_rotary_factor=0.25, mrope_section=[11,11,10]. Uniquely: model_type is qwen3_5_moe_text (not qwen3_5_text), no dense intermediate MLP, and every layer's MLP is a sparse MoE with 256 experts (moe_intermediate_size=512) plus a shared expert with a per-token sigmoid gate. This is the FIRST DeltaNet + MoE integration on Neuron; upstream NxDI ships a Qwen3-MoE (dense-attention MoE) reference and PR aws-neuron#173 provides DeltaNet+dense, but nothing combines the two. The MoE plumbing wraps NxDI's initialize_moe_module (moe_v2) inside a new Qwen35MoEBlock and adds a sigmoid-gated shared expert on top (NxDI's built-in SharedExperts only sums). Modeling deltas vs dense-2B/4B/9B/27B: - Qwen35InferenceConfig.from_pretrained preserves model_type (was hardcoded qwen3_5_text) so qwen3_5_moe_text is detected. - Qwen35InferenceConfig.__init__ auto-populates num_local_experts, n_shared_experts=1, and maps moe_intermediate_size → intermediate_size. - New Qwen35MoEBlock: routed experts via initialize_moe_module + per-token sigmoid-gated shared expert (SwiGLU MLP + sigmoid gate). - NeuronQwen35DecoderLayer routes MLP to Qwen35MoEBlock when _is_moe=True. - convert_qwen35_hf_to_neuron_state_dict transposes stacked expert weights (HF: (E, 2I, H) / (E, H, I) → NxDI: (E, H, 2I) / (E, I, H)) and renames mlp.gate.weight, mlp.experts.*, mlp.shared_expert.*_proj.weight to the NxDI convention (moe.router.linear_router.weight, moe.expert_mlps.mlp_op.gate_up_proj.weight etc, shared_gate_proj etc). Text-only validated on trn2.48xlarge, TP=8, bf16, seq_len=512: prompt tokens TTFT (ms) TPOT (ms) tok/s 16 553.8 7.67 129.9 64 554.0 7.79 128.2 256 553.6 7.74 129.2 Outputs on 5 prompts all qualitatively correct and coherent: "The capital of France is Paris." "Jupiter is a gas giant..." "Water boils at 100°C." autumn haiku ("Crimson drifts down slow, Golden carpet ...") photosynthesis definition. Notable runtime knobs (set automatically by the runner scripts): - MoENeuronConfig with moe_tp_degree=8, moe_ep_degree=1 - blockwise_matmul_config.use_torch_block_wise=True — the DLAMI ships no LNC=2 shard-hidden NKI kernel; torch fallback is functionally correct but slower. - router_config.dtype=fp32, router_config.act_fn=softmax, normalize_top_k_affinities=True. - shared_expert_gate uses plain nn.Linear (scalar output; can't shard). VL not attempted in this contrib — see README follow-ups. HF greedy match also deferred (67 GB CPU bf16 is prohibitively slow). Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Default backend changed from `use_torch_block_wise=True` (pure PyTorch
fallback) to `use_shard_on_intermediate_dynamic_while=True` (LNC=2 NKI
blockwise-matmul kernel). This is functionally correct on the shipped
SDK 2.29 DLAMI and gives ~13 % TTFT improvement at prompt=16:
backend prompt=16 prompt=64 prompt=256
use_torch_block_wise 553.8 554.0 553.6 ms
use_shard_on_intermediate_dynamic 480.0 575.5 667.1 ms ← new default
Trade-off: shard-on-intermediate TTFT scales linearly with prompt length,
while torch fallback is flat (because CPU work dominates and grows with
prompt too, just faster than the linear NKI path at long prompts). At
prompt=16 the NKI path wins; at prompt≥64 the fallback is faster. For a
short-prompt-heavy workload (chat) shard-on-intermediate is the right
default. The truly-preferred `_call_shard_hidden_kernel` LNC=2 kernel
that NxDI's ExpertMLPsV2 would pick automatically is NOT shipped in the
DLAMI (raises "kernel not imported from nkilib") — a future SDK drop
should give a further TTFT reduction.
Also explored moe_ep_degree=2 (expert parallelism, shard 256 experts
across 2 groups). EP works at prefill but NxDI raises
`NotImplementedError: Selective Loading with Expert parallelism is not
supported in token generation.` because our (top_k=8, num_experts=256)
config has per-token expert fraction of 8/256 = 3 % which is below the
DEFAULT_SELECTIVE_LOADING_THRESHOLD (=1.0) that gates the all-experts EP
path. Kept moe_ep_degree=1 as-is.
TPOT unchanged at ~7.7 ms across all three backends (decode is
selective-loading only, doesn't touch the blockwise kernel).
README updated with the two-backend table and the TTFT/TPOT ratio
explanation (35B-A3B has TTFT/TPOT ~62× vs dense 5-6× — the gap is a
consequence of the shipped MoE kernel state, not the architecture).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
…orward
NxDI's default LNC=2 forward MoE kernel (`_call_shard_hidden_kernel`) is a
NotImplementedError stub in the shipped SDK 2.29 because the
`neuronxcc.nki._private.blockwise_mm` module it tries to import doesn't
exist. But the same kernel forward implementation IS present at
`nkilib.experimental.moe.forward.bwmm_shard_on_H.blockwise_mm_baseline_shard_hidden`
— it just isn't wired up. `_patch_nxd_shard_hidden_kernel()` in
`modeling_qwen35.py` monkey-patches the stub at module import time to
delegate to the nkilib kernel with LNC=2.
Net effect: the default backend path is now runnable (flat 567 ms TTFT)
instead of raising. It doesn't beat `shard_on_intermediate` for short
prompts (480 ms at prompt=16) so that remains the default choice in
run_text_smoke.py / run_benchmark.py, but the patched path is chosen
automatically when a caller sets `blockwise_matmul_config={}` (NxD default)
and is now the fastest path for prompts ≥ 256:
backend prompt=16 prompt=64 prompt=256
use_torch_block_wise 553.8 554.0 553.6
use_shard_on_intermediate_dynamic 480.0 575.5 667.1 ← default
patched _call_shard_hidden_kernel 567.0 566.5 566.7 ← new fallback
Also corrected the README description of `use_torch_block_wise` — it runs
on Neuron via XLA (unrolled dynamic-slice loop), NOT on CPU as originally
worded.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
VL benchmark passes across three image sizes on trn2.48xlarge (TP=8, bf16): 512×512 → TTFT 1037 ms, 1024×1024 → 2004 ms, 2048×2048 (2×2 tile) → 5490 ms; TPOT stays flat at ~8.1 ms across sizes. Neuron-compiled vision encoder is shared with the dense siblings (bucket 1024 + 4096); text model recompiled with use_text_only_cte_inputs=False and CTE bucketing. This is the first DeltaNet + MoE + vision-scatter combination on Neuron. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Replace "SDK 2.29 / NKI 0.3.0" placeholders (from earlier drafting) with the real shipped versions: neuronx-cc 2.26.6360, nki 0.5.0, neuronx-distributed-inference 0.10.18399. README compatibility table now lists every relevant package version instead of an aggregate SDK number. Co-Authored-By: Claude Opus 4.7 <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.
Description
Add Qwen3.5-35B-A3B contrib model — NeuronX Distributed Inference port of
Qwen/Qwen3.5-35B-A3B, the MoE flagship of the Qwen3.5 family. 35B total parameters, ~3B activated per token through top-8 routing over 256 experts + 1 sigmoid-gated shared expert, with the same hybrid attention stack as the dense siblings: [3 gated DeltaNet + 1 full GQA] × 10 = 40 layers.This is the first DeltaNet + MoE integration on Neuron. It reuses the DeltaNet + attention path from PR #173 (Qwen3.5-27B dense) and integrates NxDI's
initialize_moe_module(frommoe_v2) via a newQwen35MoEBlock. Both text-only and vision-language inference are validated end-to-end on trn2.48xlarge.Key implementation highlights:
Qwen35MoEBlockwraps NxDI'sinitialize_moe_modulefor 256 routed experts + a per-token sigmoid-gated shared expert (NxDI's built-inSharedExpertslacks per-token gating)ExpertMLPsV2transposed layoutshard_on_intermediate_dynamic_while(default, 13% faster at short prompts),use_torch_block_wise, and monkey-patched_call_shard_hidden_kernelModel Information
Model Name: Qwen3.5-35B-A3B
Model Architecture: Hybrid DeltaNet + GQA decoder with Mixture-of-Experts FFN (256 routed + 1 shared, top-8 sigmoid routing, ~3B active/token), 40 layers,
model_type=qwen3_5_moe_textPurpose: Text generation and vision-language inference
Checklist
Required Components
Accuracy Test (
test/integration/test_model.py)README.md with the following sections:
Qwen/Qwen3.5-35B-A3B(~67 GB bf16)Source Code (
src/)modeling_qwen35.py— DeltaNet + GQA + NEWQwen35MoEBlock(~336 KB)modeling_qwen35_vision.py— ViT encoder wrappermodeling_qwen35_vl.py— VL orchestratorhybrid_apc.py— Automatic Prefix Cachenki_kernels/— DeltaNet NKI kernels (unchanged from dense siblings)Optional Components
test/unit/directory)Folder Structure
Testing
How did you test this change?
Tested on trn2.48xlarge (TP=8, bf16) with:
Test Results:
Text-only performance (TP=8, bf16, seq_len=512,
shard_on_intermediatedefault):Vision-language (TP=8, Neuron VE, CTE buckets):
Compatibility
Tested with:
Additional Information
MoENeuronConfig(notNeuronConfig),moe_tp_degree=8,moe_ep_degree=1,router.dtype=float32,router.act_fn=softmax,normalize_top_k_affinities=True._call_shard_hidden_kernelLNC=2 path is not present in shippedneuronx-cc 2.26.6360; current default (shard_on_intermediate_dynamic_while) is 13% faster than torch fallback at short prompts but scales linearly with prompt length. A future SDK drop should close this.moe_ep_degree > 1works for prefill but NxDI'sExpertMLPsV2raisesNotImplementedErrorfor selective-loading decode with EP whentop_k/num_experts < threshold.SharedExpertsonly sums without per-token gating, so we implement it directly inQwen35MoEBlock.Related Issues
Related to PR #173 (Qwen3.5-27B DeltaNet + GQA dense contribution).
vLLM Integration
vLLM integration is future work.
By submitting this PR, I confirm that: