Qwen3.5-122B-A10B on DGX Spark: 28.3 → 52 tok/s (+82%)
Optimizations for Qwen3.5-122B-A10B inference on a single NVIDIA DGX Spark from 28.3 to 52 tok/s (+82%), with 256K context support, no quality degradation. Headline 52 = round of 51.6 tok/s cross-prompt average measured 2026-05-09 on v2.4 (autotune + PR #38325, both default-on); LongCode peak 54.9 tok/s.
Results
| Configuration |
tok/s |
Improvement |
Build |
| Baseline (vLLM 0.19 + AutoRound INT4 + FlashInfer) |
28.3 |
– |
– |
| + Hybrid INT4+FP8 Dense Layers |
30.8 |
+8.8% |
step 1 |
| + MTP-2 Speculative Decoding |
38.4 |
+35.7% |
step 2 |
| v2 (+ INT8 LM Head v2) |
51 |
+80% |
Dockerfile.v2 |
v2.4 (+ @triton.autotune on LM Head + vLLM PR #38325 swapAB FP8 SM120, both default-on since 2026-05-09) |
52 (54.9 LongCode peak) |
+82% |
./install.sh (default) |
| v2.6 (stable CUDA 13.0 dependency set + reliability fixes) |
52 (unchanged) |
+82% |
./install.sh (default) |
| v2.7 (INT8 LM Head patch applies to vLLM 0.27.x as well) |
52 (unchanged) |
+82% |
./install.sh (default) |
| v2-tq (+ TurboQuant KV Cache) |
39 |
+38% |
Dockerfile.v2-tq |
The same optimizations also work with Qwen3.5-35B-A3B (same architecture, smaller): 112 tok/s.
v2.7 release
v2.7 makes the INT8 LM Head patch apply to vLLM 0.27.x and records what the
optimizations are worth there. The 0.19 build path behaves as before: the
in-image dist-packages path is still the one tried first, and the only change
to the file the patch writes on 0.19.1 is the head_dtype guard below, which
is inert on a release that has no head_dtype to begin with.
patch_int8_lmhead.py recognises the 0.27 form of _get_logits, which
routes through _apply_head, alongside the 0.19 form that calls the quant
method directly. An unrecognised form aborts without editing the file.- The target file is located through the importable vLLM when the in-image
path is absent, so a virtualenv install is covered as well.
- The INT8 path emits float16, so it declines the conversion when a
head_dtype other than the model dtype is requested and leaves those runs
to _apply_head.
- A file still carrying the v1 patch is recognised by release rather than
assumed to be 0.19.
What the optimizations measure on stock 0.27.1 with the plain Intel AutoRound
checkpoint is documented in On newer vLLM (0.27.x).
Those figures do not replace the headline numbers, which are the full v2 stack
on 0.19.1.
vLLM 0.27 support and the 0.27.1 measurements were contributed by
moep90.
Upgrading from v2.6 rebuilds only the ~1 second final layer, since the patch is
baked into the image at build time. No inference kernels or model arithmetic
changed in this release.
v2.6 maintenance release
v2.6 keeps the tested vLLM and performance patches frozen while making fresh
builds reproducible again:
- replaces expired nightly PyTorch wheels with a pinned stable CUDA 13.0 set;
- pins FlashInfer 0.6.11 from checksummed official release wheels instead of
the incompatible rolling upstream alias;
- freezes the known-compatible FastAPI routing stack used by vLLM 0.19;
- prevents CPU-only Torch replacement and verifies
libtorch_cuda.so plus
vllm._C before an image is accepted;
- backports the warning-only FLA fix for valid short prompts;
- makes the required MTP registration step explicit for manual hybrid builds;
- documents thinking and non-thinking API requests separately.
No inference kernels or model arithmetic changed in this maintenance release.
About the v2.4 numbers (5 runs × 2 sub-runs = 10 sub-runs, n=50 prompt measurements, bench_qwen35.sh 2026-05-09):
| Prompt |
Mean |
Median |
Std |
| Q&A 256 |
51.3 |
51.3 |
0.70 |
| Code 512 |
52.8 |
52.9 |
0.23 |
| JSON 1024 |
51.1 |
51.2 |
0.78 |
| Math 64 |
47.8 |
48.1 |
0.81 |
| LongCode 2048 |
54.9 |
55.0 |
0.31 |
Cross-prompt mean of per-sub-run averages: 51.58 ± 0.30 tok/s. Headline 52 = round(51.58); LongCode 54.9 reflects the most decode-bound prompt (long sustained generation) where LM Head + shared_expert FP8 paths dominate.
Composition: wonderwork v2 (canonical 51) + @triton.autotune (+1.2% A/B) + vLLM PR #38325 swapAB FP8 SM120 (+0.76% marginal A/B). The +1.1% headline gain (51 → 51.6) is at ~1.5σ of bench noise; the per-prompt deltas (especially LongCode 54.2 → 54.9 and Code 52.0 → 52.8) are above noise individually. See Optimization 4 for autotune details and Optimization 5 for PR #38325.
256K Context Support
v2 supports 256K context out of the box (355K token KV cache). No TurboQuant needed for single-user 256K.
| Config |
KV Cache |
Concurrent Users @ 256K |
| v2 (standard) |
355K tokens |
1 |
| v2-tq (TurboQuant) |
1.4M tokens |
5 |
Quick Start
All optimizations are independent — pick what you need:
| Path |
Steps |
tok/s |
What you get |
| MTP only (easiest) |
0 → 2 → 3 → 4 → 5 |
~44 |
MTP-2 + INT8 LM Head, no hybrid |
| Full v2.6 (recommended, default) |
0 → 1 → 2 → 3 → 4 → 5 |
52 (54.9 LongCode peak) |
v2.4 performance stack plus stable CUDA 13.0 dependencies and v2.6 reliability fixes. |
First-time builds take ~30-60 min because vLLM has no prebuilt SM121 wheels — Step 3 NVCC-compiles the entire vLLM C extension. PR #38325 is baked into that build at zero extra time cost. Subsequent re-runs of install.sh skip Step 3 if the image is cached (~3 min total).
Existing users (cached vllm-sm121:latest from before PR #38325 became default): the cache check skips Step 3, so you keep your old base and miss the +0.76%. To pick up PR #38325, pass --no-cache (full rebuild, ~30-60 min) or run docker rmi vllm-sm121:latest and re-run install.sh.
Need vanilla (no PR #38325)? Pass --no-pr38325 — useful if the patch breaks your build, or you want to keep a pristine vllm-sm121:latest cache. You’ll only lose the marginal +0.76% (autotune still applies).
Automated install (TL;DR)
If you want everything done for you, just run:
./install.sh
This walks through Steps 0-4 automatically with progress bars, elapsed time, and a final prompt to launch the container. It is idempotent — re-running skips steps whose outputs already exist.
The script never invokes sudo itself: if a prerequisite is missing (python3-venv, docker daemon access, etc.) it prints the exact sudo command you should run, then exits non-zero so you can fix it and re-run.
Useful flags:
./install.sh --launch # build, then auto-launch container (no prompt)
./install.sh --no-launch # build only, never prompt for launch
./install.sh --no-cache # nuke existing images + BuildKit cache and rebuild
# from scratch. Required for existing users whose
# cached vllm-sm121:latest predates PR #38325 default.
./install.sh --no-pr38325 # SKIP the PR #38325 swapAB FP8 cherry-pick (default
# IS to apply it). Use only if the patch breaks your
# build or you want to reuse a pristine vllm-sm121
# cache without the ~30-60 min recompile. Costs the
# +0.76% marginal contribution from PR #38325.
./install.sh --build-images-only # skip 122B model preparation and verify only
# the images; use a smaller model for runtime QA
./install.sh --help # full flag reference
For an isolated validation build that does not replace existing image tags:
SM121_IMAGE=vllm-sm121-v26-test \
FINAL_IMAGE=vllm-qwen35-v26-test \
./install.sh --build-images-only
If you prefer to understand the workflow, follow Steps 0-4 below. The automated
installer adds fail-closed dependency and runtime validation around these same
steps and is the supported path for a v2.6 build.
Out of scope for install.sh: the TurboQuant variant (see “Optional: TurboQuant KV Cache Compression” later) and the runtime Step 6 benchmark.
Host-side Python environment
Steps 0-2 run on the host (not inside Docker) and need a small set of Python packages. Pick whichever install style you prefer — both produce the same result.
Option A — virtualenv (recommended, clean):
python3 -m venv .venv
source .venv/bin/activate
pip install -U pip
pip install torch numpy safetensors huggingface_hub
Keep this venv active for the rest of Steps 0-2. Step 3 onward runs inside Docker and does not use the host venv.
Option B — system-wide (quick & dirty):
pip install --break-system-packages torch numpy safetensors huggingface_hub
On recent Ubuntu (24.04+), a plain pip install is blocked by PEP 668, hence the --break-system-packages flag. This works fine for a one-shot model prep but pollutes your global Python — if you ever use this machine for anything else, prefer Option A.
Note: torch and numpy are only needed by the Step 1 hybrid-checkpoint script (safetensors falls back to numpy when saving non-contiguous tensors). Steps 0 and 2 use just huggingface_hub and the stdlib. If you’re taking the MTP-only path and skipping Step 1, you can drop torch and numpy from the install list.
Step 0: Download the model
hf download Intel/Qwen3.5-122B-A10B-int4-AutoRound
INTEL_DIR=$(find ~/.cache/huggingface/hub/models--Intel--Qwen3.5-122B-A10B-int4-AutoRound/snapshots -maxdepth 1 -mindepth 1 -type d)
Note: huggingface_hub 1.x renamed the CLI from huggingface-cli to hf. If you installed an older version (0.x), use huggingface-cli download ... instead.
Step 1: Build hybrid checkpoint (optional, +9%)
Replaces BF16 shared expert weights with FP8 from the official Qwen checkpoint. Skip this if you just want MTP — the Docker image works with both hybrid and non-hybrid checkpoints (the FP8 dispatch simply doesn’t activate if no FP8 weights are present).
python patches/01-hybrid-int4-fp8/build-hybrid-checkpoint.py \
--gptq-dir "$INTEL_DIR" \
--fp8-repo Qwen/Qwen3.5-122B-A10B-FP8 \
--output ~/models/qwen35-122b-hybrid-int4fp8 \
--force
Takes ~20 minutes. Output: ~71 GB. If you skip this step, use $INTEL_DIR as your model path in step 2 and 4.
MTP is not included by this converter alone. If
model_extra_tensors.safetensors is present, the script prints the exact
add-mtp-weights.py command at startup and completion. Continue with Step 2
before launching speculative decoding. install.sh performs Step 2
automatically.
Step 2: Add MTP weights
Intel AutoRound ships model_extra_tensors.safetensors (4.8 GB, 785 MTP tensors) as a separate shard. The hybrid converter rebuilds its target index from the main model shards, so this extra shard must be copied and registered afterward. This script copies the file (if needed) and adds the 785 tensor mappings to the target index, so vLLM loads them for speculative decoding.
# Target = hybrid checkpoint (step 1) or original Intel dir (if skipping step 1)
MODEL_DIR=~/models/qwen35-122b-hybrid-int4fp8 # or $INTEL_DIR
python patches/02-mtp-speculative/add-mtp-weights.py \
--source "$INTEL_DIR" \
--target "$MODEL_DIR"
Step 3: Build base vLLM image for SM121
DGX Spark requires vLLM compiled for SM121 (Blackwell). Pre-built wheels from PyPI don’t support this architecture. Use eugr/spark-vllm-docker at the exact commit we tested against:
PROJECT_DIR=$(pwd) # path to this repo (DGX_Spark_Qwen3.5-122B-A10B-AR-INT4)
git clone https://github.com/eugr/spark-vllm-docker.git
cd spark-vllm-docker
git checkout 49d6d9fefd7cd05e63af8b28e4b514e9d30d249f
# Remove two "TEMPORARY PATCH" RUN blocks that curl-apply vLLM PRs 35568
# and 38919. Both target main-branch bugs that don't affect v0.19.0, and
# both PRs were force-pushed after 2026-04-04, so their current .diff no
# longer applies to v0.19.0 at all. Our reference image was built without
# them (verified by MD5 comparison against pristine v0.19.0).
sed -i '/# TEMPORARY PATCH for broken FP8 kernels/,/&& rm pr35568.diff/d' Dockerfile
sed -i '/# TEMPORARY PATCH for broken compilation/,/&& rm pr38919.diff/d' Dockerfile
# Pin stable PyTorch CUDA versions in BOTH stages. The upstream Dockerfile runs
# `uv pip install torch torchvision torchaudio triton ...` twice (builder
# stage ~L50 and runner stage ~L311). Without a pin, those two invocations
# resolve independently and can bake an ABI mismatch into the image. Symptom:
# `ImportError: undefined symbol: _ZN2at4cuda24getCurrentCUDABlasHandleEv`
# at startup (PyTorch changed the signature of at::cuda::getCurrentCUDABlasHandle
# between builds, so vllm/_C.abi3.so and libtorch_cuda.so disagree). The
# longer timeout covers large aarch64 wheel downloads without host networking.
sed -i 's|uv pip install torch torchvision torchaudio triton --index-url https://download.pytorch.org/whl/nightly/cu130|UV_HTTP_TIMEOUT=300 uv pip install torch==2.12.0+cu130 torchvision==0.27.0+cu130 torchaudio==2.11.0+cu130 triton==3.7.0 --index-url https://download.pytorch.org/whl/cu130|g' Dockerfile
# Cherry-pick vLLM PR #38325 (swapAB SM120 CUTLASS blockwise FP8 GEMM)
# — adds ~+0.76% throughput on shared_expert decode. SKIP this if you want
# vanilla vLLM (run the docker build with --build-arg VLLM_BASE=vllm-sm121:latest
# in Step 4 in either case; the same final tag is used).
cp "${PROJECT_DIR}/patches/05-pr38325-swapab/pr38325-swapab-fp8-sm120.diff" local-pr38325.diff
python3 -c "
import re
txt = open('Dockerfile').read()
inject = '\nCOPY local-pr38325.diff /tmp/local-pr38325.diff\nRUN git apply -v /tmp/local-pr38325.diff && rm /tmp/local-pr38325.diff\n'
new_txt = re.sub(r'(RUN if \[ -n \"\\\$VLLM_PRS\" \]; then.*? fi\n)', r'\1' + inject, txt, count=1, flags=re.DOTALL)
open('Dockerfile', 'w').write(new_txt)
"
./build-and-copy.sh -t vllm-sm121 --vllm-ref v0.19.0 --tf5
cd ..
This takes 30-60 minutes (compiles PyTorch + FlashInfer + Triton for SM121).
Why build-and-copy.sh and not docker build directly: the upstream Dockerfile does COPY build-metadata.yaml, and that file is generated at build time by build-and-copy.sh (and removed after). A plain docker build will fail with "/build-metadata.yaml": not found.
Why --tf5 and --vllm-ref v0.19.0: our image was built with transformers_5: true and vllm_ref: v0.19.0 (read from /workspace/build-metadata.yaml inside the image). The script’s defaults are different, and an image built with defaults will not be binary-compatible with our patches.
Why the two “TEMPORARY PATCH” sed commands: the upstream Dockerfile at 49d6d9f has two hardcoded RUN curl ... .diff | git apply blocks that pull vLLM PRs 35568 and 38919 live from GitHub. Both PRs target bugs in main, not in v0.19.0, and both were force-pushed after our original 2026-04-04 build, so their current diffs no longer apply to v0.19.0 (they reference files moved under csrc/libtorch_stable/ which didn’t exist yet at that path in v0.19.0). We verified via MD5 comparison that marlin_utils.py inside our reference image is byte-identical to pristine v0.19.0 — i.e. PR 35568 was never actually applied to our image in the first place. Removing the two blocks reproduces the original build.
Why the torch pin sed command: the upstream Dockerfile installs the
dependency set twice — once where vLLM’s C extension is compiled and again
in the runner image. v2.6 uses the exact stable CUDA 13.0 package set in both
stages. Unlike date-stamped nightly artifacts, these wheels are published on
the retained stable index. install.sh also prevents later dependency
resolution from replacing CUDA Torch with a CPU wheel or upgrading past the
compatible FastAPI route implementation, then verifies
libtorch_cuda.so and imports vllm._C before accepting the image.
Version pinning: Tested and verified with vLLM 0.19.1 (exact build:
0.19.1.dev0+g2a69949bd.d20260404, commit 2a69949bd,
spark-vllm-docker commit 49d6d9f) against PyTorch 2.12.0+cu130,
torchvision 0.27.0+cu130, torchaudio 2.11.0+cu130, and
Triton 3.7.0, and FlashInfer 0.6.11. Patches are version-specific — do not use with other
vLLM versions without re-testing.
A direct manual Dockerfile build must also preserve CUDA Torch and the known
compatible FastAPI stack during later wheel and Ray installation steps.
install.sh applies those resolver overrides and verifies the resulting
runtime automatically; use it instead of the abbreviated manual sequence for
a release-equivalent image.
Step 4: Build v2 image
docker build -t vllm-qwen35-v2 -f docker/Dockerfile.v2 .
This is a ~1-second thin layer on top of vllm-sm121:latest that COPYs patches/01-hybrid-int4-fp8/inc.py into the image and runs patches/03-int8-lm-head/patch_int8_lmhead.py to text-replace vllm/model_executor/layers/logits_processor.py with the INT8 LM Head v2 + autotune kernel. To layer on a different base (e.g. a vanilla vllm-sm121:latest you built without PR #38325), pass --build-arg VLLM_BASE=<image-name>:<tag>.
Step 5: Launch
docker run -d --name vllm-qwen35 \
--gpus all --net=host --ipc=host \
-v ~/models:/models \
vllm-qwen35-v2 \
serve /models/qwen35-122b-hybrid-int4fp8 \
--served-model-name qwen \
--port 8000 \
--max-model-len 262144 \
--gpu-memory-utilization 0.90 \
--reasoning-parser qwen3 \
--attention-backend FLASHINFER \
--speculative-config '{"method":"mtp","num_speculative_tokens":2}'
If you skipped step 1, replace the model path with your Intel AutoRound directory.
Wait ~13 minutes for loading + warmup (weights load ~10 min + compile/warmup ~2.5 min + graph capture ~30s; cached re-launches drop to 5-7 min). Then:
curl localhost:8000/health
curl http://localhost:8000/v1/chat/completions \
-H "Content-Type: application/json" \
-d '{"model":"qwen","messages":[{"role":"user","content":"Hello!"}],"max_tokens":256}'
Thinking and non-thinking requests
Qwen3.5 enables thinking by default. --reasoning-parser qwen3 does not turn
thinking on or off; it separates the model’s <think>...</think> output into
the API reasoning field. If the parser is omitted while thinking remains
enabled, raw reasoning text and the closing </think> tag can appear in
content.
Keep the parser enabled for thinking requests:
curl http://localhost:8000/v1/chat/completions \
-H "Content-Type: application/json" \
-d '{
"model":"qwen",
"messages":[{"role":"user","content":"Explain TCP and UDP."}],
"max_tokens":512
}'
Disable thinking per request when only the final answer is wanted:
curl http://localhost:8000/v1/chat/completions \
-H "Content-Type: application/json" \
-d '{
"model":"qwen",
"messages":[{"role":"user","content":"Explain TCP and UDP."}],
"chat_template_kwargs":{"enable_thinking":false},
"max_tokens":512
}'
Per-request control preserves both use cases and avoids changing the model’s
global default for every client.
Step 6: Benchmark
./bench_qwen35.sh "v2"
Expected: ~51 tok/s with hybrid, ~44 tok/s without hybrid (Run 2; Run 1 is JIT warmup).
Running in Production
The launch command in Step 5 (and install.sh) is a minimal smoke test. For daily use, you’ll likely want a custom launch script tailored to your setup. Here’s a community example from @whpthomas that demonstrates common production flags:
run-qwen.sh — start the server:
#!/bin/bash
docker rm vllm-qwen35
sudo sh -c 'sync; echo 3 > /proc/sys/vm/drop_caches'
docker run -it --name vllm-qwen35 \
--gpus all --net=host --ipc=host \
-v ~/models:/models \
vllm-qwen35-v2 \
serve /models/qwen35-122b-hybrid-int4fp8 \
--served-model-name qwen/qwen3.5 \
--max-model-len 196608 \
--max-num-batched-tokens 32768 \
--gpu-memory-utilization 0.88 \
--port 8000 \
--host 0.0.0.0 \
--load-format fastsafetensors \
--attention-backend FLASHINFER \
--reasoning-parser qwen3 \
--speculative-config '{"method":"mtp","num_speculative_tokens":2}' \
--enable-chunked-prefill \
--enable-auto-tool-choice \
--tool-call-parser qwen3_coder \
--generation-config auto \
--override-generation-config '{"temperature": 0.7, "top_p": 0.8, "top_k": 20, "presence_penalty": 0.0, "repetition_penalty": 1.0}'
stop-qwen.sh — stop the server:
#!/bin/bash
docker stop $(docker ps -q --filter "name=vllm")
Notable flags in this example:
| Flag |
Purpose |
sudo sh -c 'sync; echo 3 > /proc/sys/vm/drop_caches' |
Flush page cache before launch for consistent memory |
--host 0.0.0.0 |
Listen on all interfaces (LAN access) |
--load-format fastsafetensors |
Faster checkpoint loading |
--enable-chunked-prefill |
Better TTFT on long prompts |
--reasoning-parser qwen3 |
Routes thinking tokens to the API reasoning field; does not enable or disable thinking |
--enable-auto-tool-choice --tool-call-parser ... |
Agent/tool-calling support (see note below) |
--override-generation-config '{...}' |
Server-side sampling defaults |
Adapt to your needs. These are suggestions, not requirements. The only flags critical for this project’s optimizations are --attention-backend FLASHINFER and --speculative-config. Everything else depends on your use case. See vLLM documentation for the full flag reference.
Note: --enable-prefix-caching is intentionally omitted — it crashes on Qwen3.5 due to DeltaNet hybrid attention (see Troubleshooting).
Tool-call parser options: vLLM ships two parsers for Qwen models. Pick the one that matches your model:
| Parser |
Format |
Models |
qwen3_xml |
<tool_call>{"name": "fn", "arguments": {...}}</tool_call> (JSON in XML tags) |
Qwen3.5-*, Qwen3-*-Instruct |
qwen3_coder |
<tool_call><function=fn><parameter=x>val</parameter></function></tool_call> (custom XML) |
Qwen3-Coder-* |
For Qwen3.5-122B (this project), use --tool-call-parser qwen3_xml. The example above uses qwen3_coder which also works but is designed for Coder-series models.
Known issue (vLLM 0.19): When --reasoning-parser qwen3 and --tool-call-parser are both active, tool calls emitted inside <think> blocks may be silently dropped in non-streaming mode (vllm#39056).
Troubleshooting
| Symptom |
Fix |
health returns nothing |
Wait. It takes ~13 min on first launch (weights load ~10 min + compile/warmup ~2.5 min + graph capture ~30s). Cached re-launches: ~5-7 min. |
| Garbage output |
Ensure patched image, not vanilla vLLM |
| OOM at startup |
Lower --gpu-memory-utilization to 0.85 |
content: null with a populated reasoning field |
The model is still thinking or exhausted max_tokens before its final answer. Increase the token limit, or send "chat_template_kwargs":{"enable_thinking":false} for a non-thinking response. |
Raw reasoning or a trailing </think> appears in content |
Thinking is enabled but --reasoning-parser qwen3 was omitted. Add the parser, or disable thinking per request with chat_template_kwargs.enable_thinking=false. |
Input tensor shape suggests potential format mismatch |
Rebuild the final image with v2.6 or later. Earlier images contain a false-positive FLA warning for valid short prompts; tensor computation is unaffected. |
ImportError: libtorch_cuda.so |
Rebuild with the v2.6 or later stable CUDA wheels. The installer pins CUDA Torch through later dependency resolution and verifies libtorch_cuda.so plus vllm._C before accepting the image. |
Resolver reports flashinfer-python and vllm are incompatible |
A rolling upstream FlashInfer asset moved beyond vLLM 0.19’s cuDNN frontend range. Use the v2.6 or later installer; it downloads the checksummed official FlashInfer 0.6.11 wheel set. |
| Only ~38 tok/s |
Check --speculative-config has num_speculative_tokens:2 |
| Stale Triton cache after rebuild |
docker exec <name> rm -rf /root/.cache/triton and restart |
| MTP doesn’t work with PyTorch backend |
MTP requires FlashInfer backend (--attention-backend FLASHINFER). PyTorch backend is not supported. Also check vLLM version — MTP had bugs in pre-0.19 versions (#36843, #36917). |
| Multi-node / Ray cluster issues |
This project is tested on a single DGX Spark only. Multi-node setups (Ray, 2x Spark) have different requirements and are not covered here. Community reports suggest up to 56 tok/s on 2x Spark with Ray, but cluster configuration is outside our scope. |
| 408 “unexpected unmatched FP8 tensor” warnings during checkpoint build |
Normal. These are DeltaNet linear_attn projections (36 of 48 layers) plus some attention norms/gates. They exist in the Qwen FP8 checkpoint but have no matching counterparts in Intel AutoRound INT4 (different naming conventions). The script only replaces shared_expert dense layers (144 tensors) with FP8 — everything else stays in its original format. Use --force to proceed. |
Hardware
- System: NVIDIA DGX Spark (ASUS Ascent GX10)
- GPU: NVIDIA GB10 (Blackwell, SM121)
- Memory: 128 GB unified CPU-GPU (LPDDR5x, 273 GB/s)
- CUDA: 13.0
- Architecture: aarch64 (ARM Grace CPU)
Tested Environment
These exact versions were used for all benchmarks. Mismatched versions may cause errors.
| Component |
Version |
| vLLM |
0.19.1.dev0+g2a69949bd.d20260404 (commit 2a69949bd) |
| PyTorch |
2.12.0+cu130 |
| torchvision |
0.27.0+cu130 |
| torchaudio |
2.11.0+cu130 |
| CUDA Toolkit |
13.2 (V13.2.51) |
| CUDA (torch) |
13.0 |
| FlashInfer |
0.6.11 (flashinfer-jit-cache 0.6.11+cu130) |
| Triton |
3.7.0 |
| FastAPI |
0.136.1 |
| Starlette |
0.52.1 |
| prometheus-fastapi-instrumentator |
7.1.0 |
| Python |
3.12.3 |
| OS |
Ubuntu 24.04.4 LTS (aarch64) |
| Build flags |
TORCH_CUDA_ARCH_LIST=12.1a FLASHINFER_CUDA_ARCH_LIST=12.1a |
| Release |
v2.7 INT8 LM Head patch applies to vLLM 0.19.1 and 0.27.x; v2.6 dependency and reliability set and the v2.4 performance stack unchanged |
| Anchors back-out |
--no-pr38325 rebuilds without PR #38325 (loses +0.76%, keeps autotune); ad-hoc reverts of bd16c23 are possible but not exposed as a flag (autotune is byte-identical arithmetic, no quality risk) |
Why the complete package set is pinned: the base image resolves PyTorch,
torchvision, torchaudio, and Triton independently in the builder and runner
stages. Exact stable CUDA 13.0 versions keep the two stages ABI-identical and
avoid expiry of nightly artifacts. A manual build must reproduce the Step 3
replacement.
Prerequisites
For the actual build commands and rationale, see Step 3 above. This section just summarizes the SM121-specific build settings and one community alternative, both of which Step 3 implicitly covers.
The critical compile-time env vars are:
TORCH_CUDA_ARCH_LIST="12.1a" # SM121 (Blackwell consumer)
FLASHINFER_CUDA_ARCH_LIST="12.1a" # FlashInfer kernels for SM121
CUDA_HOME=/usr/local/cuda-13.0 # or 13.2
The base CUDA image should be nvidia/cuda:13.2.0-devel-ubuntu24.04 (aarch64). Build takes 30-60 minutes on DGX Spark. Pre-built vLLM wheels from PyPI do not support SM121 — compile from source.
For spark-vllm-docker users: The hybrid INC patch is also available as a community mod (enable-hybrid-int4fp8). Apply it with --apply-mod in the launch script. Note that the inc.py patch is tied to a specific vLLM version — if you update the community Docker, the patch may need adjusting (the internal inc.py API changes frequently). PR #38325 cherry-pick is NOT available as a community mod; use Step 3 if you want it.
Optimization Details
Optimization 1: FlashInfer Attention Backend
Effect: 24.0 → 28.3 tok/s (+16%)
vLLM defaults to FLASH_ATTN on SM121. FlashInfer has optimized kernels that better utilize the Blackwell memory hierarchy. One flag, +16% free.
--attention-backend FLASHINFER
Optimization 2: Hybrid INT4+FP8 Dense Layers
Effect: 28.3 → 30.8 tok/s (+8.8%)
MoE expert weights stay in INT4 (Marlin). Shared expert MLP weights replaced with FP8 from the official Qwen FP8 checkpoint, using native SM121 CUTLASS FP8 block-128 kernels.
The patch (patches/01-hybrid-int4-fp8/inc.py) fixes a bug where shared expert layers (marked as 16-bit by AutoRound) loaded FP8 weights without scale tensors.
Optimization 3: MTP-2 Speculative Decoding
Effect: 30.8 → 38.4 tok/s (+25%)
Qwen3.5-122B ships with a native MTP head (785 tensors, 4.8 GB BF16). MTP-2 (num_speculative_tokens:2) predicts 2 additional tokens per step with ~80% acceptance rate on position 2.
Why MTP-2, not MTP-1: MTP-1 was the initial v1 configuration (38.4 tok/s). MTP-2 provides an additional +10% at no quality cost, with ~80% acceptance rate on position 2.
The MTP weights live in model_extra_tensors.safetensors in the Intel AutoRound checkpoint but are missing from the model index. The script add-mtp-weights.py registers all 785 tensors.
FAQ: Can I use MTP without the hybrid checkpoint? Yes. add-mtp-weights.py only copies model_extra_tensors.safetensors and updates the index — it works on any Qwen3.5 checkpoint (INT4, FP8, or hybrid). The hybrid patch (step 1) and MTP (step 2) are independent optimizations.
What happens if I add --speculative-config without running add-mtp-weights.py? vLLM won’t find the MTP tensors and will either error out or silently disable speculative decoding (no speedup, no error — just no effect).
Optimization 4: INT8 LM Head v2
Effect: 38.4 → 51 tok/s (+33%)
The LM Head matrix (vocab 248320 × hidden 3072 ≈ 762M params; 1.5 GB BF16, 729 MB after INT8 quant) is the single largest weight read in the decode step — it has to be re-read from memory for every token. The v2 shared-weight Triton GEMV kernel:
- Quantizes BF16 → INT8 at runtime (per-channel, no calibration needed). One-shot at first request, costs ~5-10s of startup time, saves ~770 MB of bandwidth on every subsequent token.
- Single kernel launch reads the 729 MB weight matrix ONCE per batch, regardless of batch size (the v1 kernel launched N times for N tokens).
- Triton INT8 GEMV achieves ~84% memory-bandwidth utilization on a single SM121 LM-head call — vs ~24% for the default BF16 matmul, which is GEMV-shape-unfriendly (M=1) and falls back to a non-tiled CUTLASS path.
No quality degradation: INT8 per-channel quantization of the output layer preserves top-k token rankings.
On vLLM 0.27.x the ~24% BF16 figure no longer holds — upstream reaches 64% of bandwidth at batch 1 and 86% above it, so the INT8 kernel is a 2.0-2.8x replacement there rather than 3.5x. See On newer vLLM (0.27.x) for the full measurements.
v2.4: @triton.autotune (default since 2026-05-09, commit bd16c23)
The v2 kernel originally hardcoded BLOCK_M=128, BLOCK_K=256 with Triton-default num_warps=4, num_stages=2. A community contributor on the NVIDIA DGX Spark forum noticed the docstring claimed autotune but the decorator was missing — wrapping the same byte-for-byte-identical kernel in @triton.autotune over 8 configs (BLOCK_M ∈ {64,128,256}, BLOCK_K ∈ {128,256,512}, num_warps ∈ {4,8}, num_stages ∈ {2,3}) lets Triton pick the best for each (M, K, NUM_BATCH) shape. Measured on Qwen3.5-122B/Spark:
| Test |
Hardcoded |
Autotuned |
Δ |
| Q&A 256 |
50.3 |
50.5 |
+0.4% |
| Code 512 |
52.0 |
52.6 |
+1.1% |
| JSON 1024 |
51.0 |
51.3 |
+0.5% |
| Math 64 |
46.5 |
47.7 |
+2.6% |
| LongCode 2048 |
54.2 |
54.85 |
+1.2% |
Average +1.2% with no quality cost. One-time autotune cost ~6s (1.6s × 4 unique NUM_BATCH ∈ {1…4}) on first request after container start. The contributor measured +8.5% on their vLLM 0.20.1 stack with a 36 tok/s baseline (LM head dominant); the smaller win here reflects our already-optimized stack where LM head is one of several balanced bottlenecks.
Optimization 5: vLLM PR #38325 swapAB SM120 FP8 GEMM
Effect: +0.76% marginal on top of autotune / +2.0% cumulative over the wonderwork v2 baseline. Default-on since 2026-05-09. Skip with --no-pr38325.
vLLM upstream PR #38325 adds a “swapAB” CUTLASS dispatch (B-major weight) for SM120 family blockwise FP8 GEMM. The runtime auto-selects swapAB whenever M ≤ 64 || M % 4 != 0 — exactly the decode shape our shared_expert FP8 layers hit at batch 1-4. SM121 inherits the path via enable_sm120_family<...>.
Measured marginal contribution on top of autotune:
| Test |
autotune-only |
+ PR #38325 |
Δ |
| Q&A 256 |
50.5 |
50.75 |
+0.5% |
| Code 512 |
52.6 |
53.4 |
+1.5% |
| JSON 1024 |
51.3 |
51.6 |
+0.6% |
| Math 64 |
47.7 |
48.1 |
+0.8% |
| LongCode 2048 |
54.85 |
55.05 |
+0.4% |
Why default-on, not opt-in. The patch costs a full vllm-sm121 base image rebuild (~30-60 min). On a fresh install that cost is paid anyway — vLLM has no prebuilt SM121 wheels, so Step 3 NVCC-compiles vLLM regardless. Adding PR #38325 to that build is effectively free time-wise. Existing users with cached vllm-sm121:latest from before this default flip will skip Step 3 (cached-image check) and miss the +0.76% — they need --no-cache to pick it up.
Why an opt-out exists. The single .cuh diff is conservative (auto-active on small M only, no other hot paths affected), but if it ever breaks a build for a future eugr/spark-vllm-docker pin or future torch nightly, --no-pr38325 is the escape hatch. Cost of using it: lose the marginal +0.76%, keep the autotune +1.2%.
The diff lives in patches/05-pr38325-swapab/pr38325-swapab-fp8-sm120.diff, rewritten for v0.19.0 source paths (the upstream PR was authored against a later tree where csrc/libtorch_stable/... and torch::stable::Tensor exist; v0.19.0 still has csrc/quantization/... and torch::Tensor). When vLLM ≥0.20 becomes our build target, this whole opt-out goes away — PR #38325 was merged into upstream 0.20.x.
On newer vLLM (0.27.x)
patch_int8_lmhead.py now recognises the 0.27 form of
LogitsProcessor._get_logits as well as the 0.19 one, and locates
logits_processor.py through the importable vLLM instead of only the in-image
path, so it also applies to a virtualenv install. Everything below is what the
optimizations are worth there. It does not replace the headline numbers: those
are the full v2 stack, and this is stock upstream vLLM with the plain Intel
AutoRound checkpoint.
Setup. DGX Spark (GB10, SM121), vLLM 0.27.1 from PyPI, no custom image,
Intel/Qwen3.5-122B-A10B-int4-AutoRound without the hybrid conversion,
--kv-cache-dtype fp8, --max-model-len 262144, --max-num-seqs 2. Throughput
figures are vllm bench serve output throughput over 4 requests of 1024 in /
512 out at concurrency 1, so they include TTFT and are not comparable to
bench_qwen35.sh decode rates. TPOT is quoted where the comparison matters.
The bf16 baseline moved. Optimization 4 rests on the LM head GEMV reaching
only ~24% of memory bandwidth. On 0.27.1 it no longer does. Measuring
F.linear against the real head shape, 248320 x 3072 bf16, 1.53 GB, on 273
GB/s:
| batch |
bf16 |
of peak |
INT8 kernel |
of peak |
speedup |
| 1 |
8.80 ms |
64% |
3.17 ms |
88% |
2.77x |
| 2 |
6.46 ms |
86% |
3.19 ms |
88% |
2.03x |
| 4 |
6.47 ms |
86% |
3.28 ms |
85% |
1.97x |
The kernel still does what it claims, at 85-88% of bandwidth. What changed is
the thing it is compared against: upstream recovered most of the gap on its
own, and only batch 1 still picks a noticeably worse path.
End to end it is worth less than the kernel suggests. Same server, INT8 LM
head on and off:
| scenario |
off |
on |
delta |
| 1024 in / 512 out, c1 |
28.54 tok/s |
30.48 tok/s |
+6.8% |
| 1024 in / 512 out, c4 |
48.03 tok/s |
48.55 tok/s |
+1.1% |
| 16K in / 1024 out, c1 |
24.37 tok/s |
25.55 tok/s |
+4.8% |
TPOT at c1 goes 32.41 -> 30.80 ms. That is 1.6 ms saved where the isolated
kernel predicts 5.6 ms. The gap is not explained here. A plausible candidate is
CUDA graph batch padding: capture sizes are 1, 2, 4, 8, and if the head runs at
the padded shape then the bf16 path was already at 86% and the honest
comparison is 6.46 -> 3.19 ms rather than 8.80 -> 3.17 ms. Anyone rebuilding
this on 0.27.x should measure rather than assume the +33%.
Other optimizations, same setup.
|
this project on 0.19 |
measured on 0.27.1 |
| FlashInfer backend |
+16% |
+8.5% (26.3 -> 28.5 tok/s) |
| MTP-2 |
+25% |
+2.4% (28.5 -> 29.2 tok/s) |
MTP acceptance is not the problem: 8532 accepted tokens over 5798 draft steps
is 1.47 of 2, or 73.6%, close to the ~80% reported here. The overhead is. The
MTP head ships BF16 at 4.8 GB against an INT4 target model, so one draft step
reads enough weight to cost a large fraction of a real decode step at 273 GB/s.
On the hybrid + INT8 LM head stack the ratio between draft and target cost is a
different one, which is consistent with +25% being real there. MTP also cost
KV capacity at 256K, 10.58 concurrent sequences down to 7.77.
Prefix caching is not broken on 0.27.1. vLLM does not disable it for this
architecture; it selects Mamba cache mode align and logs that the mode is
experimental. A cold prefill and a full cache hit of the same 15.5K token
prompt returned byte-identical greedy output, and so did a partial prefix hit.
See the note under “What Didn’t Work”.
Quality of the INT8 head, and what was not checked. Against the bf16
reference at the shapes above, argmax matched in every case. The top-5 set
reordered in 2 of 4 cases at batch 4, at a relative logit error around 9e-3.
Greedy decoding is unaffected by a reorder below the argmax; sampling sees a
small perturbation. Full continuations were not diffed between the two heads,
so this says nothing about how a divergence would compound over a long
generation.
Optional: TurboQuant KV Cache Compression
TurboQuant (Google, ICLR 2026) compresses the KV cache from bf16 to ~3.5 bits per element, giving 4x more KV cache capacity (1.4M tokens vs 355K) at the cost of -22% generation speed (39 vs 51 tok/s).
How it works
Each KV vector (head_size=256, bf16 = 512 bytes) is compressed to 120 bytes using two techniques:
- MSE quantization: Vectors are rotated via structured Hadamard transform, then each coordinate is quantized to the nearest Lloyd-Max centroid (3-bit for outlier dims, 2-bit for regular).
- QJL residual: The quantization residual is encoded as 1-bit sign projections through a random matrix, preserving inner product estimates.
Packed format per KV position (120 bytes total):
Group 0 (128 outlier dims): 48B MSE indices + 16B QJL signs + 4B norms = 68 bytes
Group 1 (128 regular dims): 32B MSE indices + 16B QJL signs + 4B norms = 52 bytes
Compression: 512 bytes (bf16) → 120 bytes = 4.27x
The attention score is computed directly on compressed data without decompression:
score(q, k) = Σ_group [ vector_norm × (
dot(codebook[mse_indices], F @ q) + // MSE term
residual_norm × dot(qjl_signs, G @ q × scale) // QJL term
)]
where F and G are structured Hadamard rotation matrices, and the codebook has only 8 entries (3-bit) or 4 entries (2-bit) — fits entirely in GPU registers.
Why -22% speed is unavoidable
The performance penalty comes from two architectural changes:
-
Triton replaces FlashInfer (-15%): TQ requires a custom attention kernel that reads packed data. FlashInfer doesn’t support custom KV formats, so TQ uses a Triton-based attention backend. Triton generates less optimized GPU code than FlashInfer’s hand-tuned CUDA kernels.
-
PIECEWISE instead of FULL CUDA graphs (-7%): The Triton TQ kernel has JIT compilation behavior that’s incompatible with FULL CUDA graph capture. PIECEWISE graphs add CPU launch overhead between kernel segments.
We also developed and benchmarked a custom CUDA fused kernel that computes attention directly on TQ packed data (in patches/04-turboquant/cuda_tq_fused/). After 8 optimization iterations (238x speedup from v1 to v8), the CUDA kernel achieves 38 GB/s effective bandwidth — but FlashInfer achieves 221 GB/s on the same hardware. The gap is fundamental: bit unpacking + codebook lookup requires ~1.5x more integer ALU operations per KV position than a simple bf16 dot product. With 4.3x less data but 1.5x more compute per byte, the net result is always slower than bf16 FlashInfer on bandwidth-bound SM121 hardware.
Bottom line: TQ is a memory-for-speed trade-off. The -22% penalty is the cost of 4x compression. There is no free lunch.
When to use TQ
- High-throughput serving: 4x more concurrent users at 256K context (5 vs 1)
- Not needed for 256K context on single user: the standard v2 image already fits 256K with 355K token cache
Important notes
- Qwen3.5 was tested by its developers on context lengths up to 256K only. Longer contexts are not validated.
- The current INT4 quantization (Intel AutoRound) was calibrated for standard context lengths. For deeper contexts (>256K), a custom AutoRound calibration would be needed.
- The 1.4M token KV cache theoretically supports ~1M context, but model quality beyond 256K is unverified.
TurboQuant Recipes
The generate_tq_metadata.py script supports multiple recipes. Choose based on your trade-off between memory, speed, and quality:
| Recipe |
K Storage |
V Storage |
Memory vs fp16 |
Outlier Dims |
Best For |
turboquant35 |
TQ35 |
TQ35 |
4× (128B/head) |
50% K+V same |
Default, balanced |
turboquant25 |
TQ25 |
TQ25 |
5× (108B/head) |
25% K+V same |
Maximum compression |
turboquant_asym |
TQ35 |
TQ35 |
4× (128B/head) |
50% K≠V disjoint |
Better needle retrieval |
turboquant_q8k_tq35v |
int8 |
TQ35 |
~2× (258B/head) |
K=full, V=50% |
Best quality |
turboquant_q8k_tq25v |
int8 |
TQ25 |
~2× (258B/head) |
K=full, V=25% |
Quality + V compression |
Recipe details:
-
turboquant35 (default): Symmetric 3.5-bit encoding. First 50% of dimensions are high-precision outliers for both K and V. Best balance of compression and quality.
-
turboquant25: More aggressive 2.5-bit encoding. Only 25% of dimensions are high-precision outliers. Higher compression ratio but more reconstruction error.
-
turboquant_asym (Asymmetric): Same storage as TQ35 (128B/head), but K uses the first 50% of dimensions while V uses the last 50%. This disjoint selection decorrelates K and V index sets, improving reconstruction quality for long-context needle-in-haystack tasks without extra memory cost.
-
turboquant_q8k_tq35v (True Asymmetric): K stored as int8 + fp16 scale (258 bytes/head), V stored as TQ35. ~2× memory overhead vs fp16 (vs 4× for symmetric TQ). K’s full int8 precision preserves attention score quality, while V’s TQ compression reduces the overall footprint. Recovers needle-in-haystack quality to 3/3 at 256K context (vs 0/3 with symmetric TQ35).
-
turboquant_q8k_tq25v: Same as above but V uses TQ25 (25% outlier dims). Same memory as q8k-tq35v, more aggressive V compression.
Build & Run (TQ variant)
# Step 1: Generate TQ metadata for each recipe you want to use.
# Each recipe produces its own metadata file — they are NOT interchangeable.
# --kv-cache-dtype at runtime must match the recipe used here.
# TQ35 — default, balanced (4x memory reduction)
python patches/04-turboquant/generate_tq_metadata.py \
--model-dir ~/models/qwen35-122b-hybrid-int4fp8 \
--output-path ~/models/qwen35-122b-hybrid-int4fp8/turboquant_kv_tq35.json
# TQ25 — maximum compression (5x memory reduction, lower quality)
python patches/04-turboquant/generate_tq_metadata.py \
--model-dir ~/models/qwen35-122b-hybrid-int4fp8 \
--recipe turboquant25 \
--output-path ~/models/qwen35-122b-hybrid-int4fp8/turboquant_kv_tq25.json
# turboquant_asym — same memory as TQ35, better long-context needle retrieval
python patches/04-turboquant/generate_tq_metadata.py \
--model-dir ~/models/qwen35-122b-hybrid-int4fp8 \
--recipe turboquant_asym \
--output-path ~/models/qwen35-122b-hybrid-int4fp8/turboquant_kv_asym.json
# turboquant_q8k_tq35v — K=int8, V=TQ35 (best quality, ~2x memory)
python patches/04-turboquant/generate_tq_metadata.py \
--model-dir ~/models/qwen35-122b-hybrid-int4fp8 \
--recipe turboquant_q8k_tq35v \
--output-path ~/models/qwen35-122b-hybrid-int4fp8/turboquant_kv_q8k_tq35v.json
# turboquant_q8k_tq25v — K=int8, V=TQ25 (quality + more V compression, ~2x memory)
python patches/04-turboquant/generate_tq_metadata.py \
--model-dir ~/models/qwen35-122b-hybrid-int4fp8 \
--recipe turboquant_q8k_tq25v \
--output-path ~/models/qwen35-122b-hybrid-int4fp8/turboquant_kv_q8k_tq25v.json
# Step 2: Build TQ image
docker build -t vllm-qwen35-v2-tq -f docker/Dockerfile.v2-tq .
# Step 3: Run — --kv-cache-dtype must match the recipe used in Step 1.
# Each recipe has its own metadata file; using the wrong file causes a startup error.
# TurboQuant auto-selects TRITON_ATTN backend regardless of other flags.
# TQ35 (default)
docker run -d --name vllm-qwen35-tq \
--gpus all --net=host -v ~/models:/models \
vllm-qwen35-v2-tq \
serve /models/qwen35-122b-hybrid-int4fp8 \
--served-model-name qwen --port 8000 \
--max-model-len 262144 --gpu-memory-utilization 0.90 \
--reasoning-parser qwen3 \
--kv-cache-dtype turboquant35 --enable-turboquant \
--turboquant-metadata-path /models/qwen35-122b-hybrid-int4fp8/turboquant_kv_tq35.json \
--speculative-config '{"method":"mtp","num_speculative_tokens":2}'
# TQ25 (maximum compression)
docker run -d --name vllm-qwen35-tq \
--gpus all --net=host -v ~/models:/models \
vllm-qwen35-v2-tq \
serve /models/qwen35-122b-hybrid-int4fp8 \
--served-model-name qwen --port 8000 \
--max-model-len 262144 --gpu-memory-utilization 0.90 \
--reasoning-parser qwen3 \
--kv-cache-dtype turboquant25 --enable-turboquant \
--turboquant-metadata-path /models/qwen35-122b-hybrid-int4fp8/turboquant_kv_tq25.json \
--speculative-config '{"method":"mtp","num_speculative_tokens":2}'
# turboquant_asym (disjoint K/V outlier dims)
docker run -d --name vllm-qwen35-tq \
--gpus all --net=host -v ~/models:/models \
vllm-qwen35-v2-tq \
serve /models/qwen35-122b-hybrid-int4fp8 \
--served-model-name qwen --port 8000 \
--max-model-len 262144 --gpu-memory-utilization 0.90 \
--reasoning-parser qwen3 \
--kv-cache-dtype turboquant_asym --enable-turboquant \
--turboquant-metadata-path /models/qwen35-122b-hybrid-int4fp8/turboquant_kv_asym.json \
--speculative-config '{"method":"mtp","num_speculative_tokens":2}'
# turboquant_q8k_tq35v (K=int8, V=TQ35 — best quality)
docker run -d --name vllm-qwen35-tq \
--gpus all --net=host -v ~/models:/models \
vllm-qwen35-v2-tq \
serve /models/qwen35-122b-hybrid-int4fp8 \
--served-model-name qwen --port 8000 \
--max-model-len 262144 --gpu-memory-utilization 0.90 \
--reasoning-parser qwen3 \
--kv-cache-dtype turboquant_q8k_tq35v --enable-turboquant \
--turboquant-metadata-path /models/qwen35-122b-hybrid-int4fp8/turboquant_kv_q8k_tq35v.json \
--speculative-config '{"method":"mtp","num_speculative_tokens":2}'
# turboquant_q8k_tq25v (K=int8, V=TQ25 — quality + more V compression)
docker run -d --name vllm-qwen35-tq \
--gpus all --net=host -v ~/models:/models \
vllm-qwen35-v2-tq \
serve /models/qwen35-122b-hybrid-int4fp8 \
--served-model-name qwen --port 8000 \
--max-model-len 262144 --gpu-memory-utilization 0.90 \
--reasoning-parser qwen3 \
--kv-cache-dtype turboquant_q8k_tq25v --enable-turboquant \
--turboquant-metadata-path /models/qwen35-122b-hybrid-int4fp8/turboquant_kv_q8k_tq25v.json \
--speculative-config '{"method":"mtp","num_speculative_tokens":2}'
First launch is slow (~15-20 min vs ~10 min for v2). TQ patches modify vLLM internals at startup, and the Triton decode kernels are JIT-compiled on first run. Subsequent launches with cached Triton kernels are faster.
Important: --kv-cache-dtype must match the recipe embedded in the metadata file. Each recipe produces its own file (e.g. turboquant_kv_q8k_tq35v.json vs turboquant_kv_q8k_tq25v.json). Passing the wrong file causes a ValueError at startup.
Attention Backend: TurboQuant automatically uses TRITON_ATTN. The selector switches to TRITON_ATTN whenever any turboquant* dtype is detected — do not pass --attention-backend FLASHINFER with TurboQuant, it will fail.
TQ Benchmarks
| Config |
tok/s |
KV Cache |
Concurrent @ 256K |
Memory/head |
Notes |
| v2 (standard) |
51 |
355K |
1 |
512B (fp16) |
Baseline |
| v2-tq (TQ35) |
39 (-22%) |
1.4M |
5 |
128B |
Best memory efficiency |
| v2-tq-asym |
~39 |
1.4M |
5 |
128B |
Better long-context retrieval |
| v2-tq-q8k-tq35v |
~35 (-31%) |
~700K |
~3 |
258B |
Best quality, 2× memory |
What Didn’t Work
We tested 20+ optimization approaches across speculative decoding, quantization, sparsity, kernel optimization, and system tuning. 5 worked (above), the rest didn’t. Here’s the full graveyard so you don’t repeat our experiments.
Speculative Decoding
EAGLE-3 — +10%, but loses to MTP-2. EAGLE-3 requires downloading and storing separate draft model weights (~5 GB). MTP-2 uses the built-in MTP head (already in the checkpoint) and is both simpler and faster. Two patches were created for EAGLE-3 integration but ultimately abandoned since MTP-2 wins on every metric.
KnapSpec Self-Speculative — Skipped after analysis. KnapSpec uses the model itself as a draft by skipping layers. Math showed it’s bandwidth-inferior to MTP: the draft forward pass reads ~75% of full model weights, while the MTP head reads only 4.4%. On a bandwidth-bound system like DGX Spark, this makes KnapSpec slower than MTP by design.
MTP-1 → MTP-2 — MTP-1 replaced, not failed. MTP-1 (num_speculative_tokens:1) gave 38.4 tok/s. MTP-2 (num_speculative_tokens:2) gives 51 tok/s with ~80% acceptance rate on position 2. Strictly better — more tokens per step with no quality cost.
Expert-Level Optimizations
SERE Expert Re-routing — 0% improvement. SERE re-routes tokens from underutilized experts to similar popular ones. But Qwen3.5’s 256 experts are extremely specialized — pairwise similarity is only ~0.02 (max 0.08). SERE needs coarser MoE architectures (8-16 experts) where redundancy exists.
MoE Expert Pruning (MoE-Spec) — Skipped. Since SERE proved all 256 experts are unique, pruning any of them would degrade quality. No intra-expert redundancy to exploit.
2:4 Structured Sparsity — Skipped. Expert FFN hidden size is only 512. Too small for structured pruning patterns (2:4 needs meaningful redundancy within weight matrices). Same conclusion as SERE/MoE-Spec: quality is sacred.
LM Head / Output Layer
AdaptiveSoftmax — 0% improvement. Mathematically correct idea (only compute logits for likely tokens), but the scattered memory reads for frequent-token subsets achieve only 2.5% bandwidth utilization. The INT8 GEMV approach (full matrix, 84% BW utilization) wins decisively by reading contiguously.
INT8 Shared Expert — GIBBERISH output. The shared expert weights are already in FP8 (from hybrid patch) with carefully calibrated per-tensor scales. Naive FP8→INT8 re-quantization with per-channel scales destroys those calibrations — produced garbage output. Never re-quantize calibrated FP8 weights to INT8.
Quantization
NVFP4 (RedHatAI) — -42% slower (16.6 tok/s). SM121 doesn’t have working FP4 CUTLASS kernels in vLLM yet, so it falls back to Marlin SM80 PTX which handles FP4 poorly. Waiting for vLLM PRs #38957 and #31607.
FP8 KV Cache — +0.2 tok/s (negligible). --kv-cache-dtype fp8 adds almost nothing. SM121 lacks native FP8 attention kernels; the dtype conversion overhead eats any bandwidth savings. Not worth the risk of subtle accuracy issues.
Abliterated (Uncensored) Model — OOM. The abliterated variant needs 244 GB in BF16. DGX Spark has 128 GB + 56 GB swap = 184 GB. Doesn’t fit.
Kernel Optimizations
Triton Native SM121 MoE Kernels — 0% improvement. We forced vLLM to use Triton-compiled native SM121 kernels instead of Marlin SM80 PTX for MoE expert GEMM. Exactly the same speed. The bottleneck is LPDDR5x memory bandwidth (273 GB/s), not compute. Both kernel implementations achieve the same memory throughput.
Native SM121 FP4 CUTLASS — Not possible. SM121 (consumer Blackwell) does NOT have WGMMA or tcgen05.mma tensor core instructions — those are datacenter-only (SM100/SM103). SM121 uses the same mma.sync as SM80 Ampere. The “3.65x speedup” reported on NVIDIA forums was on datacenter Blackwell, not DGX Spark.
MARLIN_USE_ATOMIC_ADD — 0% on single GPU. VLLM_MARLIN_USE_ATOMIC_ADD=1 is designed for multi-GPU tensor parallel. No effect on single-GPU SM121.
System / Runtime
CPU/JIT Overhead Reduction — Noise (+0.4%). torch.profiler showed 183% CPU utilization (2 of 20 cores), GIL-bound. Pinning to big ARM cores gave +0.4% — within noise. The OS already schedules correctly on Grace CPU.
Prefix Caching — Broken on Qwen3.5. DeltaNet layers maintain recurrent state that conflicts with KV prefix caching. Enabling --enable-prefix-caching produces incorrect outputs. vLLM correctly disables it automatically for hybrid attention architectures. Experimental ‘align’ mode gave -2% (slightly slower). On vLLM 0.27.1 this no longer holds: align is selected automatically for this architecture instead of prefix caching being disabled, and cold prefill against full cache hit returned byte-identical greedy output on a 15.5K token prompt. See “On newer vLLM (0.27.x)”.
vLLM PR Cherry-picks — 0% improvement. PR #38990 (shared expert overlap): not applicable to v0.19.1, the bug was introduced in a later refactor. PR #37700 (FLA/TMA SM12x fix): applied cleanly, 0% speedup — DeltaNet layers aren’t the decode bottleneck.
SM121 Architecture Notes
- ISA: Same
mma.sync as SM80 (Ampere). No datacenter-only tensor core instructions.
- No native FP4: FP4 tensor core ops are SM100/SM103 only (datacenter Blackwell).
- Memory-bound at batch=1: 273 GB/s LPDDR5x is the ceiling.
- FlashInfer wins: +16% over FlashAttention2. Always use
--attention-backend FLASHINFER.
Competitive Landscape (April 2026)
| Setup |
tok/s |
vs Ours |
| This work (v2) |
51 |
– |
| Previous best (v1, MTP-1 + Hybrid) |
38.4 |
-25% |
| Intel AutoRound INT4 (vLLM, FlashInfer) |
28.3 |
-45% |
| llama.cpp GGUF Q5_K |
23.0 |
-55% |
| NVFP4 RedHatAI (vLLM) |
16.6 |
-67% |
| Official Qwen GPTQ-Int4 (vLLM) |
15.0 |
-71% |
File Structure
.
├── README.md
├── bench_qwen35.sh # Benchmark script (5 tests × 2 runs)
├── LICENSE # Apache 2.0
├── patches/
│ ├── 01-hybrid-int4-fp8/
│ │ ├── inc.py # Pre-patched vLLM INC module
│ │ ├── inc.py.patch # Diff for reference
│ │ └── build-hybrid-checkpoint.py # Hybrid checkpoint builder (~20 min)
│ ├── 02-mtp-speculative/
│ │ └── add-mtp-weights.py # Register MTP weights in model index
│ ├── 03-int8-lm-head/
│ │ └── patch_int8_lmhead.py # INT8 LM Head v2 (baked into image at build time)
│ ├── 04-turboquant/ # Optional: TurboQuant KV cache
│ │ ├── generate_tq_metadata.py # Generate turboquant_kv.json
│ │ ├── kv_cache_interface.py # TQ-aware KV cache interface
│ │ ├── patch_turboquant_v2.py # Main TQ patch script
│ │ ├── turboquant_kv_cache.py # TQ layout, pack/unpack, codebooks
│ │ ├── turboquant_metadata.py # Per-layer TQ config + JSON
│ │ ├── triton_turboquant_decode.py # Triton decode attention kernel
│ │ ├── triton_turboquant_kv_update.py # Triton encode (compress) kernel
│ │ ├── selector.py # TQ-aware attention backend selector
│ │ ├── triton_attn.py # Modified Triton attention backend
│ │ └── cuda_tq_fused/ # Experimental: CUDA fused kernel
│ │ ├── tq_fused_decode.cu # CUDA kernel (v8, SM121)
│ │ ├── tq_fused_decode.py # Python wrapper + Hadamard transforms
│ │ ├── setup.py # CUDAExtension build
│ │ ├── lab_2_debug.py # Correctness verification
│ │ ├── lab_2_perf.py # Performance benchmark
│ │ └── lab_2_tq_fused_bench.py # Full benchmark suite
│ └── 06-fla-shape-warning/
│ └── patch_fla_shape_warning.py # Remove false short-prompt layout warnings
├── tests/
│ ├── test_release_v26.py # Maintenance release regressions
│ ├── test_release_v27.py # Release version contract (installer / image labels)
│ └── test_int8_lmhead_versions.py # INT8 LM Head patch vs vLLM 0.19 / 0.27 forms
├── docker/
│ ├── Dockerfile.v2 # Main: vLLM + hybrid + INT8 LM Head (baked in)
│ └── Dockerfile.v2-tq # Optional: + TurboQuant (baked in)
└── configs/
├── launch-baseline.sh # 28.3 tok/s (reference)
├── launch-hybrid.sh # 30.8 tok/s (hybrid only)
├── launch-v2.sh # 51 tok/s (production)
└── launch-v2-tq.sh # 39 tok/s (TQ variant)
Acknowledgments
License
Apache 2.0, following the license of the original model.
评论