Skip to content

Add Qwen3.5-35B-A3B (support vision and text input)#180

Open
qingzwang wants to merge 5 commits into
aws-neuron:mainfrom
qingzwang:qwen3.5-35b-a3b-hybrid-deltanet-moe
Open

Add Qwen3.5-35B-A3B (support vision and text input)#180
qingzwang wants to merge 5 commits into
aws-neuron:mainfrom
qingzwang:qwen3.5-35b-a3b-hybrid-deltanet-moe

Conversation

@qingzwang

Copy link
Copy Markdown

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 (from moe_v2) via a new Qwen35MoEBlock. Both text-only and vision-language inference are validated end-to-end on trn2.48xlarge.

Key implementation highlights:

  • Qwen35MoEBlock wraps NxDI's initialize_moe_module for 256 routed experts + a per-token sigmoid-gated shared expert (NxDI's built-in SharedExperts lacks per-token gating)
  • Weight converter handles HF stacked 3D expert weights → NxDI's ExpertMLPsV2 transposed layout
  • Three MoE blockwise-matmul backends evaluated: shard_on_intermediate_dynamic_while (default, 13% faster at short prompts), use_torch_block_wise, and monkey-patched _call_shard_hidden_kernel
  • DeltaNet NKI kernels — unchanged from dense siblings (fused chunked forward CTE)
  • Neuron VL — Neuron-compiled vision encoder (buckets 1024/4096) + text model with CTE bucketing, 2×2 spatial tiling for 2048×2048 images
  • No NxDI library modifications required — runs on stock DLAMI venv

Model 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_text

Purpose: Text generation and vision-language inference

Checklist

Required Components

  • Accuracy Test (test/integration/test_model.py)

    • All 5 prompts produce coherent, factually correct Neuron output
    • HF greedy comparison deferred (67 GB model too large for CPU bf16 in reasonable time)
  • README.md with the following sections:

    • Usage Example: Full CLI commands for text-only and VL
    • Compatibility Matrix: Detailed component version table
    • Example Checkpoints: Link to Qwen/Qwen3.5-35B-A3B (~67 GB bf16)
    • Testing Instructions: Commands for text smoke, benchmark, VL compile + benchmark
  • Source Code (src/)

    • modeling_qwen35.py — DeltaNet + GQA + NEW Qwen35MoEBlock (~336 KB)
    • modeling_qwen35_vision.py — ViT encoder wrapper
    • modeling_qwen35_vl.py — VL orchestrator
    • hybrid_apc.py — Automatic Prefix Cache
    • nki_kernels/ — DeltaNet NKI kernels (unchanged from dense siblings)

Optional Components

  • Unit Tests (test/unit/ directory)

Folder Structure

/contrib/models/Qwen3.5-35B-A3B/
  README.md
  /src
    __init__.py
    modeling_qwen35.py
    modeling_qwen35_vision.py
    modeling_qwen35_vl.py
    hybrid_apc.py
    /nki_kernels
      __init__.py
      nki_deltanet.py
      nki_deltanet_chunked.py
      nki_deltanet_fused.py
      nki_deltanet_fused_legacy.py
      qwen_qk_norm_rope.py
  /test
    __init__.py
    /unit
    /integration

Testing

How did you test this change?

Tested on trn2.48xlarge (TP=8, bf16) with:

  • Text-only: compile + generate (5 diverse prompts, all coherent + factually correct)
  • Benchmark: TTFT/TPOT sweep across prompt lengths with 3 MoE backends
  • VL: Neuron vision encoder (buckets 1024/4096) + text model, tested across 512×512, 1024×1024, 2048×2048 images
  • Correctly identifies Pallas's cat (manul) in test images

Test Results:

Text-only performance (TP=8, bf16, seq_len=512, shard_on_intermediate default):

Prompt tokens TTFT (ms) TPOT (ms) Throughput (tok/s)
16 480.0 ~7.7 ~130
64 575.5 ~7.7 ~130
256 667.1 ~7.7 ~130

TPOT of ~7.7 ms matches the ~3B activated parameter budget (sits between 9B at 6.9ms and 4B at 5.7ms from dense siblings).

Vision-language (TP=8, Neuron VE, CTE buckets):

Image TTFT (ms) TPOT (ms) tok/s
512×512 1,037 8.2 122
1024×1024 2,004 8.1 124
2048×2048 (2×2 tiled) 5,490 8.1 123

Compatibility

Tested with:

  • neuronx-cc: 2.26.6360.0
  • NKI: 0.5.0
  • Instance Type: trn2.48xlarge (TP=8)
  • NxDI: 0.10.18399
  • NxD: 0.19.28492
  • PyTorch: 2.9.1 (torch-neuronx 2.9.0.2)
  • Python: 3.12
  • transformers: 4.57.6 (HF CPU reference needs ≥ 5.13)

Additional Information

  • First DeltaNet + MoE model on Neuron — combines hybrid linear-attention architecture with sparse MoE routing on NxDI.
  • MoE config: MoENeuronConfig (not NeuronConfig), moe_tp_degree=8, moe_ep_degree=1, router.dtype=float32, router.act_fn=softmax, normalize_top_k_affinities=True.
  • Shipped SDK MoE kernel gap: The preferred _call_shard_hidden_kernel LNC=2 path is not present in shipped neuronx-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.
  • Expert Parallelism (EP=1): moe_ep_degree > 1 works for prefill but NxDI's ExpertMLPsV2 raises NotImplementedError for selective-loading decode with EP when top_k/num_experts < threshold.
  • Per-token sigmoid-gated shared expert: Qwen3.5-MoE specific — NxDI's built-in SharedExperts only sums without per-token gating, so we implement it directly in Qwen35MoEBlock.
  • HF greedy comparison deferred — 67 GB model too large for CPU bf16 in reasonable time; all Neuron outputs are qualitatively correct.
  • VL uses legacy-direct DeltaNet kernel — same as dense siblings for numerical stability.

Related Issues

Related to PR #173 (Qwen3.5-27B DeltaNet + GQA dense contribution).

vLLM Integration

  • This model/feature is intended for use with vLLM
  • Documentation includes vLLM registration instructions

vLLM integration is future work.


By submitting this PR, I confirm that:

  • I have read and followed the contributing guidelines
  • This is a community contribution and may have limited testing compared to officially-supported models
  • The code follows best practices and is well-documented
  • All required components listed above are included

Ubuntu and others added 5 commits July 13, 2026 07:16
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>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant