This document is relevant for: Trn2, Trn3
Tutorial: Deploy dense Qwen3-32B with vLLM Neuron#
This tutorial prepares a dense Qwen3-32B checkpoint, launches an OpenAI-compatible server, and checks its responses. BF16 targets Trn2 and Trn3; native MXFP8 targets Trn3.
Overview#
Start with a TP8 single server using BF16 weights and BF16 KV. On Trn3, convert the same pinned BF16 source to MXFP8 and repeat the smoke procedure. MXFP8 is validated with both BF16 KV and packed FP8 KV. The DI deployment uses two Trn3 nodes, packed FP8 KV, TP4/DP16/LNC2 prefill, and TP8/DP16/LNC1 decode. Check its responses through the proxy.
For hardware, checkpoint formats, and KV-cache restrictions, see the model recipe. This tutorial uses dense Qwen3-32B.
Prerequisites#
A Trn2 instance for BF16, or Trn3 for BF16 or MXFP8, with at least eight available logical NeuronCores at LNC2.
For DI, two Trn3 nodes with NIXL/LIBFABRIC connectivity between them.
A working environment from the setup guide and familiarity with the online serving quickstart.
A checkout matching the plugin in the chosen image, including
examples/vllm_neuron/models/qwen3/.Access to
Qwen/Qwen3-32Band local storage for the source and converted checkpoints. Conversion also needs host memory for the source model.For conversion,
llmcompressor,transformers, andsafetensorsin a separate conversion environment. Freeze their versions with the converter commit.
Step 1: Record the environment and prepare checkpoints#
Use a Neuron 2.33 serving environment from the setup guide. Record its image digest and the installed versions before deployment:
git rev-parse HEAD
python -m pip show vllm vllm-neuron torch neuronx-cc
neuron-ls
Download an immutable source revision. Replace BF16_REVISION with the agreed
Hugging Face commit SHA before running:
export BF16_REVISION=REPLACE_WITH_CHECKPOINT_COMMIT
export BF16_MODEL=/models/Qwen3-32B-bf16
export MXFP8_MODEL=/models/Qwen3-32B-mxfp8
hf download Qwen/Qwen3-32B --revision "$BF16_REVISION" --local-dir "$BF16_MODEL"
For BF16 serving, proceed to Step 2. For MXFP8, run the converter in the conversion environment, using the downloaded directory so the source revision cannot change during conversion:
python examples/vllm_neuron/models/qwen3/build_mxfp8_checkpoint.py \
--model-checkpoint "$BF16_MODEL" \
--output-dir "$MXFP8_MODEL" \
--torch-dtype bfloat16 \
--ignore-modules lm_head
python -m pip freeze > /models/qwen3-conversion-requirements.txt
Record the converter’s Git commit and complete command. Hash every checkpoint file, including configuration and tokenizer files, after conversion:
python - <<'PY_HASH'
import hashlib
from pathlib import Path
import os
for name in ("BF16_MODEL", "MXFP8_MODEL"):
root = Path(os.environ[name])
if not root.is_dir():
continue
output = root.parent / (root.name + ".sha256")
with output.open("w") as manifest:
for path in sorted(root.rglob("*")):
if not path.is_file() or ".cache" in path.relative_to(root).parts:
continue
digest = hashlib.sha256()
with path.open("rb") as stream:
for chunk in iter(lambda: stream.read(8 * 1024 * 1024), b""):
digest.update(chunk)
manifest.write(f"{digest.hexdigest()} {path.relative_to(root)}\n")
PY_HASH
Copy the complete checkpoint directory to the serving environment and verify its checksums there. MXFP8 uses the compressed-tensors checkpoint produced by this builder; a similarly named repository is insufficient provenance.
Step 2: Launch a single server#
Run one weight format at a time. Select the backend being validated explicitly:
export VLLM_NEURON_BACKEND=neuron_native
export NEURON_LOGICAL_NC_CONFIG=2
export VLLM_NEURON_COMPILATION_TIMEOUT=1200
export NEURON_LIBTORCH_COMPILATION_TIMEOUT=1200
export VLLM_EXECUTE_MODEL_TIMEOUT_SECONDS=1200
export MODEL_PATH="$BF16_MODEL"
Use neuron_native for the baseline. Set VLLM_NEURON_BACKEND=vllm_neuron to
select the XLA backend. Repeat the response checks after a backend change.
The BF16 baseline uses BF16 KV. Prefix caching is on by default and needs
segmented prefill, so the maximum model length (10,240) must exceed the
8192-token batch, and kv_segment_size_buckets sets the segment size. The
repository examples use the same initial smoke approach.
vllm serve "$MODEL_PATH" \
--served-model-name qwen3-32b \
--dtype bfloat16 \
--optimization-level 2 \
--tensor-parallel-size 8 \
--max-model-len 10240 \
--max-num-batched-tokens 8192 \
--max-num-seqs 4 \
--additional-config '{
"neuron_config": {
"on_device_sampling_config": {},
"kv_segment_size_buckets": [8192],
"num_batched_tokens_buckets": [8192],
"num_seqs_buckets": [4]
}
}'
For MXFP8 on Trn3, set MODEL_PATH="$MXFP8_MODEL". MXFP8 is validated with both
BF16 KV and packed FP8 KV. Set "quantization": "mxfp8" in both cases; platform
validation rejects the compressed-tensors checkpoint without it.
For BF16 KV, use the Step 2 command with this additional configuration:
{
"neuron_config": {
"quantization": "mxfp8",
"on_device_sampling_config": {},
"kv_segment_size_buckets": [8192],
"num_batched_tokens_buckets": [8192],
"num_seqs_buckets": [4]
}
}
For packed FP8 KV, also add --kv-cache-dtype fp8 --block-size 128 and
"fp8_packed_kv": true:
{
"neuron_config": {
"quantization": "mxfp8",
"fp8_packed_kv": true,
"on_device_sampling_config": {},
"kv_segment_size_buckets": [8192],
"num_batched_tokens_buckets": [8192],
"num_seqs_buckets": [4]
}
}
Keep the server’s stdout and stderr in the validation record. Cold launch compiles model graphs; wait for startup completion before sending traffic. The batch limit and compiled sequence bucket must agree. The maximum batched tokens must match the largest compiled token bucket.
Step 3: Validate responses and recovery#
In another shell, check health and a deterministic completion:
curl --fail-with-body http://localhost:8000/health
curl --fail-with-body http://localhost:8000/v1/completions \
-H 'Content-Type: application/json' \
-d '{"model":"qwen3-32b","prompt":"The capital of France is","max_tokens":32,"temperature":0}'
Require HTTP 200 and a coherent choices[0].text that gives Paris as the answer.
Write the prompt without a trailing space. A trailing space makes the model
continue with an unrelated sentence. Preserve the request and
response. A qualitative smoke response does not replace logit or task accuracy.
Send concurrent requests, checking every command’s exit status and response:
python - <<'PY_TRAFFIC'
from concurrent.futures import ThreadPoolExecutor
import json
from urllib.request import Request, urlopen
def complete(index):
body = json.dumps({"model": "qwen3-32b", "prompt": "Count: 1, 2, 3,",
"max_tokens": 32, "temperature": 0}).encode()
request = Request("http://localhost:8000/v1/completions", data=body,
headers={"Content-Type": "application/json"})
with urlopen(request, timeout=1200) as response:
result = json.load(response)
assert result["choices"][0]["text"], result
return index, result
with ThreadPoolExecutor(max_workers=4) as pool:
for index, result in pool.map(complete, range(16)):
print(json.dumps({"request": index, "response": result}))
PY_TRAFFIC
Stop the server normally, restart with the same image, checkpoint, and command, and repeat health and completion checks. Record cold and warm startup time and any unexpected OOM, compile timeout, or runtime error.
Prefix caching and long context#
Validate prefix caching and long context as a separate configuration. Use the
single-server command and additional configuration from Step 2 for your weight
format and KV format, and change only --max-model-len to 32768. Keep prefix
caching on: do not add --no-enable-prefix-caching, which only the DI command in
Step 4 uses.
Send the same long prefix twice, verify output consistency, and inspect prefix-cache metrics. Use tokenizer-measured prompts beyond 8192 tokens to exercise segmented prefill and near the configured context limit while leaving room for generation. Record token counts, TTFT, and errors. Do not infer 32k validation from the 10,240-token baseline smoke.
LongBench-v2 at 32k is validated on a single TP8 server at batch 8 with BF16 KV,
for both BF16 and MXFP8 weights. That configuration uses --max-model-len 32768,
--max-num-batched-tokens 8192, --max-num-seqs 8, and "num_seqs_buckets": [8],
with prefix caching on by default. See the
model recipe for its acceptance floor.
Step 4: Deploy disaggregated inference#
The MXFP8 DI configuration runs 1P1D on two Trn3 nodes, and each role uses a whole node:
Setting |
Prefill node |
Decode node |
|---|---|---|
Parallelism |
TP4 / DP16 / LNC2 |
TP8 / DP16 / LNC1 |
Optimization level |
2 |
3 |
Maximum batch per replica |
1 |
64 |
API servers |
16 |
16 |
KV cache |
Packed FP8, block size 128 |
Packed FP8, block size 128, 5232 blocks |
Both roles use a 10,240-token context, an 8192-token batched-token bucket, async scheduling, and a 300 s KV lease. The DI configuration is validated with prefix caching off. Follow the DI tutorial for NIXL, LIBFABRIC, network setup, and proxy operation. Use the same image and checkpoint on both nodes.
Note: This DI configuration supports a maximum model length of up to 10,240 tokens. For 32k context, including LongBench evaluation, use the validated single-server configuration in Prefix caching and long context.
Launch the target roles and proxy#
Open a shell in the same serving image on each node. In both shells, set MXFP8_MODEL,
and set PREFILL_IP and DECODE_IP to the node addresses reachable by NIXL/LIBFABRIC.
Follow the network checks in the DI tutorial.
On the prefill node:
export PREFILL_IP=REPLACE_WITH_PREFILL_NODE_IP
export DECODE_IP=REPLACE_WITH_DECODE_NODE_IP
export NEURON_LOGICAL_NC_CONFIG=2
export DI_TP=4
export DI_OPT_LEVEL=2
export DI_MAX_SEQS=1
DI_KV_ARGS=()
export DI_ROLE_CONFIG='"kv_segment_size_buckets": [8192],'
export VLLM_NIXL_SIDE_CHANNEL_PORT=5559
export DI_ROLE=kv_producer
export DI_KV_RANK=0
export DI_ENGINE_ID=qwen3-prefill
export DI_PORT=8100
On the decode node:
export PREFILL_IP=REPLACE_WITH_PREFILL_NODE_IP
export DECODE_IP=REPLACE_WITH_DECODE_NODE_IP
export NEURON_LOGICAL_NC_CONFIG=1
export DI_TP=8
export DI_OPT_LEVEL=3
export DI_MAX_SEQS=64
DI_KV_ARGS=(--num-gpu-blocks-override 5232)
export NKILIB_EXPERIMENTAL_ATTN_TKG_NO_SWAP=0
export VLLM_NEURON_ATTN_MASK_TORCH=1
export VLLM_NEURON_ATTN_PV_SWAP=1
export VLLM_NIXL_SIDE_CHANNEL_PORT=5659
export DI_ROLE=kv_consumer
export DI_KV_RANK=1
export DI_ENGINE_ID=qwen3-decode
export DI_PORT=8200
Then run this command on each node. Each role sets its own LNC. Sixteen data-parallel replicas fill the node: 64 logical cores at LNC2 for prefill and 128 at LNC1 for decode. Each replica of the decode node admits up to 64 sequences, so the node holds 1,024 concurrent requests.
export VLLM_NEURON_BACKEND=neuron_native
export VLLM_NIXL_SIDE_CHANNEL_HOST=0.0.0.0
export VLLM_NEURON_NIXL_PROGRESS_THREAD_DELAY_US=200
export NEURON_RT_EXEC_TIMEOUT=600
export VLLM_ENGINE_READY_TIMEOUT_S=3600
export VLLM_NEURON_COMPILATION_TIMEOUT=1800
export VLLM_EXECUTE_MODEL_TIMEOUT_SECONDS=1200
vllm serve "$MXFP8_MODEL" \
--served-model-name qwen3-32b \
--dtype bfloat16 \
--optimization-level "$DI_OPT_LEVEL" \
--tensor-parallel-size "$DI_TP" \
--data-parallel-size 16 \
--api-server-count 16 \
--kv-cache-dtype fp8 \
--block-size 128 \
"${DI_KV_ARGS[@]}" \
--max-model-len 10240 \
--max-num-batched-tokens 8192 \
--max-num-seqs "$DI_MAX_SEQS" \
--no-enable-prefix-caching \
--async-scheduling \
--port "$DI_PORT" \
--kv-transfer-config "{
\"kv_connector\": \"NixlConnector\",
\"kv_role\": \"$DI_ROLE\",
\"kv_buffer_device\": \"cuda\",
\"engine_id\": \"$DI_ENGINE_ID\",
\"kv_ip\": \"$PREFILL_IP\",
\"kv_rank\": $DI_KV_RANK,
\"kv_parallel_size\": 2,
\"kv_connector_extra_config\": {
\"backends\": [\"LIBFABRIC\"],
\"kv_lease_duration\": 300,
\"prefill_ips\": [\"$PREFILL_IP\"],
\"decode_ips\": [\"$DECODE_IP\"]
}
}" \
--additional-config "{
\"neuron_config\": {
$DI_ROLE_CONFIG
\"quantization\": \"mxfp8\",
\"fp8_packed_kv\": true,
\"on_device_sampling_config\": {},
\"num_batched_tokens_buckets\": [8192],
\"num_seqs_buckets\": [$DI_MAX_SEQS]
}
}"
Each server log must report Started 16 API server processes. During traffic,
the log must not contain Releasing expired KV blocks.
Stop the single-server process on port 8000 first. Once both roles report startup complete, launch the proxy in a second shell on the prefill node:
curl --fail-with-body "http://$PREFILL_IP:8100/health"
curl --fail-with-body "http://$DECODE_IP:8200/health"
python examples/vllm_neuron/vllm/disaggregated_inference/toy_proxy_server.py \
--port 8000 \
--prefiller-hosts "$PREFILL_IP" --prefiller-ports 8100 \
--decoder-hosts "$DECODE_IP" --decoder-ports 8200
Repeat Step 3’s completion request through port 8000 and retain transfer logs from both roles. The toy proxy is a validation example; use the DI tutorial to configure a serving deployment.
KV cache size#
The decode engine pins 5232 KV blocks to avoid automatic cache sizing on shared LNC1 HBM.
The allocation is (10240 / 128) × (64 + 1) + 32 blocks.
It provides a 10,240-token context for the 64 sequences of each replica, plus headroom.
Keep the QK-swap environment settings on the decode role.
The model recipe includes logit and GSM8K results for this mixed-LNC configuration. Use the same role settings when comparing output or performance.
Add EAGLE3 speculative decoding#
For single-server BF16 or MXFP8 serving, use the matching
RedHatAI/Qwen3-32B-speculator.eagle3 draft. Pin its revision with the target checkpoint.
Add this option to the Step 2 command for the chosen weight format:
--speculative-config '{
"method": "eagle3",
"model": "RedHatAI/Qwen3-32B-speculator.eagle3",
"num_speculative_tokens": 5
}'
Note: EAGLE3 supports a batch size of up to 8 for an MXFP8 target on Trn3 and up to 4 for a BF16 target on Trn2, with TP8 decode at LNC2. Keep
--max-num-seqsandnum_seqs_bucketswithin these limits.
EAGLE3 needs more device memory outside the KV cache than the default headroom gives. Set the headroom before you launch the server:
# MXFP8 target on Trn3
export VLLM_NEURON_RESERVED_HEADROOM_GIB=17.2
# BF16 target on Trn2
export VLLM_NEURON_RESERVED_HEADROOM_GIB=9.4
With the default headroom of 4 GiB, an MXFP8 target on Trn3 starts and serves one request. The first concurrent batch then fails with a device memory allocation error, and the engine stops. Each additional GiB of headroom removes 1 GiB from the KV cache budget.
Target |
Hardware |
Validated configuration |
|---|---|---|
MXFP8, packed FP8 KV |
Trn3 |
TP8. Batch 8, 32k context, 8192-token bucket, |
BF16, BF16 KV |
Trn2 |
TP8, 10,240-token context, headroom 9.4 GiB. Batch 4, |
Each headroom value is validated only for its configuration.
For an MXFP8 target with packed FP8 KV, keep --kv-cache-dtype fp8 --block-size 128
from Step 2. Use num_speculative_tokens 5. Do not use 3 at batch 4: it gives
incorrect target logits on this path.
This example uses TP8. A BF16 target uses BF16 KV. With EAGLE3, the MXFP8 target
is validated with packed FP8 KV only. The target and draft compile during startup.
Repeat the completion and concurrency checks. Read the proposed and accepted
draft-token counters from /metrics:
curl --fail-with-body http://localhost:8000/metrics | \
grep -E '^vllm:spec_decode_num_(draft|accepted)_tokens_total'
Use counter deltas for the workload. Compute acceptance rate as accepted draft
tokens divided by proposed draft tokens. Accepted draft tokens show that the
speculative path is active. The mean accepted length is
1 + accepted / (proposed / num_speculative_tokens). Over a long generation
(for example GSM8K), a value below 1.5 shows that speculation does not work.
Short outputs give lower values. Compare throughput and latency with the same prompts,
concurrency, and non-speculative command. Keep the target’s accuracy checks in the comparison.
The mixed-LNC DI procedure above uses standard decoding.
Step 5: Benchmark the validated workload#
Benchmark the same image, checkpoint, flags, and endpoint used for accuracy and smoke validation. Start with the GPT-OSS benchmarking procedure and the repository’s Qwen3 performance targets. Record input/output token counts, request distribution, concurrency, per-replica batch, warmup, measurement duration, TP/DP, LNC, KV dtype, optimization level, and context buckets.
Report throughput, TTFT, inter-token latency, cold/warm startup, and memory use. For DI, benchmark through the proxy. For example, a workload that fills the 16 decode replicas at batch 64:
vllm bench serve \
--base-url http://localhost:8000 \
--model qwen3-32b \
--tokenizer "$MXFP8_MODEL" \
--dataset-name random \
--seed 42 \
--random-input-len 8192 \
--random-output-len 1024 \
--random-range-ratio 0 \
--num-prompts 51200 \
--max-concurrency 1024 \
--num-warmups 1024 \
--temperature 0 \
--percentile-metrics ttft,tpot,itl \
--ignore-eos
The concurrency of 1,024 fills the 16 decode replicas at batch 64. The time to fill and drain the replicas lowers the reported throughput of short runs, so send many prompts per unit of concurrency. Compare runs only at the same prompt count.
Compare identical workloads against the baseline and verify batch admission before reporting a high-batch result. Measured performance alone does not establish model accuracy.
Confirm your work#
You have completed this tutorial. For each weight format you deployed, you see the following:
GET /healthreturns HTTP 200, and the completion request returns Paris as the answer, both on a cold start and after a restart with the same image, checkpoint, and command.The 16 concurrent requests all return a non-empty
choices[0].text, with no OOM, compile timeout, or runtime error in the server log.For DI, both role servers report
Started 16 API server processes, the completion request succeeds through the proxy on port 8000, and the logs contain noReleasing expired KV blocks.For EAGLE3,
vllm:spec_decode_num_accepted_tokens_totalincreases under traffic, and the server stays up through the concurrent batch.
Keep the image digest, checkpoint record, and exact commands with these results. For the logit and task-score accuracy methods, see the model recipe.
Common issues#
Symptom |
Action |
|---|---|
Compile timeout |
Check the compile logs, timeout environment variables, and cache setup in the setup guide. Record cold and warm startup separately. |
OOM or incomplete batch admission |
Reduce context or concurrency for diagnosis, keep buckets aligned, and report the configuration that actually ran. Revalidate the intended release point before claiming it. |
MXFP8 fails to load |
Verify compressed-tensors metadata, builder commit, conversion options, checksums, and Trn3 hardware. Keep |
DI health succeeds but proxy requests fail |
Check side-channel addresses, NIXL/LIBFABRIC connectivity, core allocation, matching checkpoints, and producer/consumer transfer logs. |
EAGLE3 compile or numerical failure |
Use the baseline non-speculative command while investigating; repeat compile and accuracy checks for the EAGLE3 configuration. |
EAGLE3 server returns HTTP 500 after the first concurrent batch, and the log shows |
Set |