Add Qwen3.5-2B (support vision and text input)#179
Open
qingzwang wants to merge 9 commits into
Open
Conversation
Qwen3.5-2B is a 2B-parameter image-text-to-text decoder with a hybrid attention stack: 24 layers arranged as [3 gated DeltaNet + 1 full GQA] x 6. Adapts the modeling code contributed by PR aws-neuron#173 (Qwen3.6-27B sibling, which shares model_type=qwen3_5) to the 2B variant. Key deltas vs the 27B code: - Add update_state_dict_for_tied_weights (2B ties embed_tokens<->lm_head; 27B does not). - Fallback shims for cancel/finish/prepare_hybrid_apc_request so the module loads on stock NxDI SDK 2.29 (those symbols are only exported by the unmerged PR aws-neuron#173 branch; hybrid APC is disabled by default anyway). - VL generate output parsing: read tokens off CausalLMOutputWithPast.tokens when on_device_sampling is enabled (logits is None in that mode). - Pad vision_embeddings / vision_mask / mRoPE position_ids to the CTE bucket size instead of the raw input length. Validated on trn2.48xlarge (TP=8, bf16, seq_len=512): - TTFT ~17.6 ms, TPOT ~4.0 ms (250 tok/s) across prompt lengths 16-256. - 3/5 prompts fully match HF transformers 5.13 CPU greedy (16/16 tokens); 66% aggregate token-match rate. - End-to-end VL path runs on a real image via CPU-side ViT + Neuron text decoder. Runs on stock /opt/aws_neuronx_venv_pytorch_2_9_nxd_inference/ with no library modifications. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
The VL orchestrator copied from PR aws-neuron#173 is not verified upstream. When wired end-to-end here it produces degenerate output (repeated token) on real images even though each component works in isolation: - CPU vision encoder: cosine 0.99 vs HF Qwen3_5VisionModel on a real image - get_rope_index: 100% token-position match vs HF compute_3d_position_ids - Text decoder compiles cleanly with use_text_only_cte_inputs=False The remaining gap is in the vision-scatter path inside the traced graph. Root cause is not yet isolated — likely candidates: the has_real_vision_inputs = shape[1] != seq_length gate in NeuronQwen35Model.get_model_output (~L5747), the padded_seq_len-1 fill for vision_mask pad slots (which overwrites a real text token position with vision garbage), or an interaction with the DeltaNet deltanet_padding_mask. README updated to prominently flag VL as WIP and list the debug pointers. Also drops one now-inaccurate comment in run_vl_smoke's pad section. No changes to text-only path (still: TTFT ~17.6 ms, TPOT ~4.0 ms, 3/5 prompts fully match HF greedy). Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Vision path is now working: - Image → HF AutoProcessor → CPU vision encoder (Qwen3-VL ViT, cos=0.99 vs HF) → vision embeddings scatter into text decoder → coherent text output that identifies image content. - Verified on a real cat image: model output correctly describes "a fluffy, light-colored cat with a thick, woolly coat". Root cause of prior _quantity garbage: the DEFAULT fused-multihead DeltaNet NKI kernel is numerically unstable on structured vision embeddings (per-dim mean up to 2.88 despite std 0.17). The legacy_direct kernel is stable. Changes: - Replace has_real_vision_inputs = shape[1] != seq_length gate with a static config-flag gate (traced_with_vision = not use_text_only_cte_inputs and vision inputs are 3-D with seq_length matching bucket). This makes the scatter unconditionally trace into the graph when the model is compiled with vision support, matching upstream qwen3_vl behavior. - Change vision pad convention in _prepare_vision_args_for_padded_seq to match upstream qwen3_vl: pad rows are zeros, pad-slot targets are padded_seq_len-1 (a guaranteed attention_mask==0 slot). - Simplify encode_vision_to_input to match upstream scatter_by_index_put: a single index_put_(accumulate=False), no masked scatter, no clone. Same correctness contract as qwen3_vl. - run_vl_smoke.py now sets QWEN36_DELTANET_CTE_IMPL=legacy_direct and QWEN36_DELTANET_MULTIHEAD_CTE=0 before importing the model so compile produces the numerically stable graph. - run_vl_smoke.py aligns its own pad convention with upstream. - README updated to reflect working VL, with a note on the required kernel flag and remaining follow-ups (ViT to Neuron, root-cause fix for fused kernel). No changes to text-only path; TTFT / TPOT and greedy-match numbers are unchanged from the previous commit. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Adds run_vl_benchmark.py that compiles once with NxDI CTE bucketing enabled
(buckets [512, 1024, 2048, 4096, 8192]) and measures TTFT / TPOT / output
on 512×512, 1024×1024, 2048×2048 resizes of the same image.
Numbers (trn2.48xlarge, TP=8, bf16, CPU vision + Neuron text):
size vision tokens TTFT (ms) TPOT (ms) tok/s
512×512 256 169 4.1 242
1024×1024 1,024 1,091 4.3 232
2048×2048 4,096 5,025 3.9 197
TTFT is dominated by the CPU vision encoder; the Neuron text-CTE alone would
be far below 200 ms at 4k tokens. Tracing the vision encoder to Neuron is
the biggest remaining latency win.
Enabling supports:
- Qwen35ModelWrapper.pad_inputs now pre-pads / truncates vision_embeddings
and vision_mask to the target CTE bucket BEFORE calling super().pad_inputs,
so upstream ModelWrapper.pad_inputs (image_to_text base) does not replace
them with dummies when the caller-provided shape ≠ pad_length. Fixes the
"Vision mask shape [1, X, 1] does not match padded sequence length Y" path.
- pad_inputs mrope handling adds a truncate branch for shape[-1] > padded_seq_len
(formerly threw integer-overflow because pad_size was negative).
- modeling_qwen35_vision.CPUVisionModel uses F.scaled_dot_product_attention
when available. This gets 2048×2048 CPU vision (4k vision tokens, 24 blocks)
from >6 minutes to ~5 seconds; matches HF reference qualitatively.
Accuracy note: HF greedy identifies "Pallas's cat" at all 3 sizes. Neuron
matches at 1024×1024, mis-identifies as "Pangolin" at 512×512 (a 2B-model
capacity issue reproducible in HF at reduced quality), hedges to
"wildcat/lynx" at 2048×2048. Text quality is coherent at all 3 sizes.
README updated with the size sweep table and accuracy note.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Adds compile_vision_encoder.py that torch_neuronx.trace's the CPUVisionModel
per patch-token bucket. Wrapper (NeuronQwen35VisionModelWrapper.load_compiled)
already supported multi-bucket loading; the fix here is that when the bucket
selection exceeds all compiled buckets, we now transparently fall back to
the CPU vision encoder instead of raising RuntimeError.
Compiled buckets: 1024 (512×512 image), 4096 (1024×1024 image).
The 16384 bucket (2048×2048 image → 16k patch tokens → 512 MB bf16 attention
mask) OOMs on trn2 single-core HBM; not compiled. 2048×2048 requests fall
back to CPU vision transparently.
Benchmark (trn2.48xlarge, TP=8, bf16, text_bucket [512..8192]):
image vision_tokens TTFT (CPU VE) TTFT (Neuron VE) Δ
512×512 256 169 ms 86 ms 2.0× ↓
1024×1024 1,024 1,091 ms 526 ms 2.1× ↓
2048×2048 4,096 5,025 ms 4,823 ms ~1.0× (CPU fallback)
TPOT unchanged at ~4 ms across all sizes (text decode is Neuron-only
regardless). Accuracy unchanged vs the CPU-vision path.
run_vl_benchmark.py gains a --vision-compiled-dir flag; when passed, loads
compiled buckets plus the CPU model for oversized-input fallback.
README updated with the CPU vs Neuron VE table and compile command.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Adds NeuronQwen35VisionModelWrapper._tiled_forward. When pixel_values
exceeds all compiled Neuron VE buckets AND its (H, W) patch grid is
divisible by 2*spatial_merge_size AND (H*W/4) fits in some compiled
bucket, split into 2×2 spatial tiles of (H/2, W/2) patches each, encode
each tile through the small-bucket kernel, and re-interleave the merged
outputs into full-image row-major order.
Bench (trn2.48xlarge, TP=8, bf16, vision buckets [1024, 4096]):
image CPU VE Neuron VE Neuron VE + tile
512×512 169 ms 86 ms 86 ms
1024×1024 1,091 ms 526 ms 528 ms
2048×2048 5,025 ms 4,823 ms* 2,090 ms ← 2.4× ↓ vs CPU
(* CPU fallback because 16384 bucket OOMs)
2048×2048 accuracy trade-off: tile boundaries lose cross-tile attention
(known LLaVA-NeXT-style limitation). On the cat test image output drops
from "wildcat/lynx" (CPU) to "wolf/canid" (tiled). Text quality remains
coherent; a follow-up would be to TP the vision encoder or compile a
windowed-attention 16384-bucket kernel.
TPOT (~4 ms) unchanged at all sizes — text decode is Neuron-only.
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>
Adds compile_vision_encoder_tp.py which uses NxD's parallel_model_trace to shard the ViT's ColumnParallelLinear / RowParallelLinear ops across N cores. NeuronQwen35VisionModelWrapper.load_compiled now auto-detects TP subdirs (bucket_<N>/tp_*.pt) alongside legacy vision_encoder_<N>.pt in the same directory. Measured: TP=4 vision on bucket 4096 drops standalone latency 308→101 ms (3.05×) and 1024×1024 E2E TTFT 526→318 ms. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Text-only (seq_len=1024): TTFT ~42ms flat across prompt lengths, TPOT ~4.75ms (1.76× / 1.19× slower than TP=8). VL: 512×512 TTFT 126ms, 1024×1024 488ms (both ~1.5× slower than TP=8). 2048×2048 requires CTE bucket 8192 not compiled in this run; marked n/a with extrapolation note. 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-2B contrib model — NeuronX Distributed Inference port of
Qwen/Qwen3.5-2B, a 2B-parameter vision-language decoder with a hybrid attention stack: 24 transformer layers arranged as [3 gated DeltaNet (linear attention) + 1 full GQA attention] × 6.This is the first hybrid linear-attention + softmax-attention model in the contrib folder. It supports both text-only and vision-language inference end-to-end on trn2.48xlarge.
Key implementation highlights:
head_dim=256,partial_rotary_factor=0.25)[1024, 4096]), 2×2 spatial tiling for 2048×2048 images,get_rope_index3-D mRoPE, andindex_put_scatter at image-token positions. Also supports CPU-only VE fallback.Model Information
Model Name: Qwen3.5-2B
Model Architecture: Hybrid decoder — [3 gated DeltaNet (linear attention) + 1 GQA (softmax attention)] × 6, with ViT vision encoder
Purpose: Text generation and vision-language (image-to-text) inference
Checklist
Required Components
Accuracy Test (
test/integration/test_model.py)run_text_smoke.py,run_benchmark.py,run_accuracy_check.py,run_hf_reference.py,run_vl_smoke.py,run_vl_benchmark.py,compile_vision_encoder.py,compile_vision_encoder_tp.pyREADME.md with the following sections:
Qwen/Qwen3.5-2B(~4.5 GB bf16)Source Code (
src/)modeling_qwen35.py— text decoder (DeltaNet + GQA + MRoPE, ~324 KB)modeling_qwen35_vision.py— ViT encoder wrapper (~40 KB)modeling_qwen35_vl.py— VL orchestrator (~28 KB)hybrid_apc.py— Automatic Prefix Cache (~70 KB, unused by default)nki_kernels/— 5 NKI kernel files for DeltaNetOptional Components
test/unit/directory exists)Folder Structure
Testing
How did you test this change?
Tested on trn2.48xlarge with:
transformers==5.13.0CPU bf16 greedyTest Results:
Text-only performance (TP=8, bf16, seq_len=512, batch=1):
TP=4 comparison (seq_len=1024):
Vision-language (TP=8, Neuron VE with CTE bucketing):
Optional TP-sharded vision encoder (
compile_vision_encoder_tp.py) further parallelizes the ViT across multiple NeuronCores viaparallel_model_trace(tp_degree=N).Accuracy vs. HF reference: 53/80 = 66% exact-token greedy match (3/5 prompts fully match). Qualitatively coherent on all prompts.
Compatibility
Tested with:
Additional Information
/opt/aws_neuronx_venv_pytorch_2_9_nxd_inference/DLAMI venv.[1024, 4096]), 2× TTFT speedup vs CPU baseline. Optional TP-sharded mode (compile_vision_encoder_tp.py) distributes QKV/MLP across cores. 2048×2048 images use transparent 2×2 spatial tiling through the 4096-bucket kernel.run_vl_smoke.pycan fall back to CPU vision (cosine similarity 0.99 vs. HF reference) for debugging.Qwen3.5-27B).qwen3_5architecture identifier was introduced in transformers 5.x (the Neuron runtime only needs 4.57+).Related Issues
Related to PR #173 (Qwen3.5-27B DeltaNet + GQA contribution).
vLLM Integration
vLLM integration is future work.
By submitting this PR, I confirm that: