Is there a reason why you use this flag? It will slow down prompt processing.
If anything, AWQ version runs very well and the recipe for it is included in the repo already :)
[ vllm-marlin-sm12x/test_e2m1_ptx.cu at main · flash7777/vllm-marlin-sm12x · GitHub ] proofs …
The F-Suffix Problem
SM120/SM121 have FP4 Tensor Cores, but the PTX instruction cvt.rn.satfinite.e2m1x2.f32 is only available with the F-suffix:
| Gencode | Compiles? | Runs on SM121? |
|---|---|---|
compute_120, sm_120 |
❌ ptxas error | — |
compute_120a, sm_120a |
❌ ptxas error | — |
compute_120f, sm_120f |
✅ | ✅ 8/8 tests pass |
compute_121f, sm_121f |
✅ | ✅ 8/8 tests pass |
The compiler (ptxas) behaves correctly — with the wrong suffix it aborts rather than inserting NOOPs. With the correct suffix all test cases execute correctly (including saturation on overflow).
Proof: test_e2m1_ptx.cu — Commit 88f6f21 in github.com/flash7777/vllm-marlin-sm12x
CUTLASS Guard Bug (NVIDIA Bug)
CUTLASS incorrectly checks for the A-variants instead of the F-variants:
// WRONG in CUTLASS float_subbyte.h:
defined(CUTLASS_ARCH_MMA_SM120A_ENABLED) // sm_120a does NOT have cvt.e2m1
defined(CUTLASS_ARCH_MMA_SM121A_ENABLED) // sm_121a does NOT have cvt.e2m1
// CORRECT would be:
defined(CUTLASS_ARCH_MMA_SM120F_ENABLED) // sm_120f DOES have it
defined(CUTLASS_ARCH_MMA_SM121F_ENABLED) // sm_121f DOES have it
Consequence: Building CUTLASS code with -gencode arch=compute_120a,code=sm_120a that uses NumericConverter<float_e2m1_t, float> results in a build error — the guard enables PTX code that ptxas rejects.
NVIDIA’s own vLLM CMake works around the bug by using "12.0f" for NVFP4 kernels. FlashInfer/Avarok hits the same bug because JIT compiles with sm_121a.
Sunday Thesis
The suffix system (a vs. f vs. none) is fragile and error-prone. Even NVIDIA’s own CUTLASS team got the guard wrong. Correct NVFP4 usage on Blackwell requires the F-suffix. Very conFusing?
If you are confused too - it seems that Avarok has solved a problem that does not exist anymore (software fallback for e2m1), only exists in older versions of CUTLASS and is “solved” in newer ones by just using cpu instruction sets that the compiler does not want you to use if you havent said the magic word “family”.
vllm-marlin-sm12x on github flash7777 could work on 590.x.y driver and updated nv ubuntu. keep in mind that I drink a lot, rather espresso yirgasheffe.
and yes, its a deep deep rabbithole and no way out.
This entire compute capability thing is messed up indeed.
According to NVIDIA documentation, if you specify 12.1a as the compilation target, it should include all family related features + architecture-specific features (it acts as a superset).
In vLLM itself it’s the case now, all architecture codes get normalized to family with minor version stripped, so 12.1f/a will be treated the same way as 12.0 in CMakeFiles.txt.
However, in addition to compile-time guards, there are runtime checks, and those often don’t include 12.1 variants (like in flashinfer).
I tried to compile flashinfer and vLLM with 12.0f, 12.1a, 12.1f - there was no difference at all - all models behaved the same way. I settled on 12.1a, as it’s a “proper” code, but I’m keeping an eye on others’ findings too.
Oh, and there is a known bug, that you should let the builds use the system PTXAS to avoid this error - that’s why I have ENV TRITON_PTXAS_PATH=/usr/local/cuda/bin/ptxas in the Docker.
I believe, they only recently fixed it in Flashinfer, but not in other places.
This looks awesome, thanks for for this @tbraun!
I just got a DGX Spark unit a week ago and starting to play around this weekend, so the timing was excellent.
I just want to leave the setup I used with the avarok/dgx-vllm-nvfp4-kernel:v22 for others new into the field as well. Somewhat redundant from the github quick guide, a setup below with some extra steps for newbies.
In order to get 68-70+ tps the speculative decoding needs to be enabled.
For this, the patch fix_mtp_nvfp4_exclusion.py needs to be applied in the docker environment before the llm model is launched.
This can be done manually in an interactive way.
On the other hand, in order to automate the launch with an entrypoint, the following bash script start_vllm.sh may be useful:
#!/bin/bash
#1. Clear the potentially corrupted FlashInfer cache
rm -rf /root/.cache/flashinfer
#2. Add these environment variables for Marlin compilation
export VLLM_USE_FLASHINFER_MOE_FP4=0
export VLLM_TEST_FORCE_FP8_MARLIN=1
export VLLM_NVFP4_GEMM_BACKEND=marlin
#3. Hard-force Offline Mode to run locally
export HF_HUB_OFFLINE=1
export TRANSFORMERS_OFFLINE=1
export PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True
#4. Apply the Python MTP patch
python3 /tmp/fix_mtp.py
MODEL=“nvidia/Qwen3-Next-80B-A3B-Instruct-NVFP4”
echo “[INFO] Launching model from: $MODEL_PATH”
#5. Start vLLM with the explicit path
python3 -m vllm.entrypoints.openai.api_server
-–model “$MODEL”
-–quantization nvfp4
-–host 0.0.0.0
-–port 8888
-–gpu-memory-utilization 0.90
-–max-model-len 32768
-–attention-backend flashinfer
-–kv-cache-dtype fp8
-–no-enable-chunked-prefill
-–speculative-config ‘{“method”:“mtp”,“num_speculative_tokens”:2}’
-–trust-remote-code
in the directory where the script is you may run the command:
sudo docker run --rm -it --name vllm_node \
-–network host --gpus all --ipc=host \
-–entrypoint /bin/bash \
-v $HOME/.cache/huggingface:/root/.cache/huggingface \
-v $(pwd)/fix_mtp_nvfp4_exclusion.py:/tmp/fix_mtp.py \
-v $(pwd)/start_vllm.sh:/tmp/start_vllm.sh \
-e HF_HUB_OFFLINE=1 \
-e TRANSFORMERS_OFFLINE=1 \
avarok/dgx-vllm-nvfp4-kernel:v22 \
/tmp/start_vllm.sh
moreover, if you want to further simplify the execution command, the script can be also be used in a .yalm docker compose file, for example docker-compose.yml with content below.
services:
vllm-qwen3:
image: avarok/dgx-vllm-nvfp4-kernel:v22
container_name: vllm_node
network_mode: host
ipc: host
deploy:
resources:
reservations:
devices:
- driver: nvidia
count: all
capabilities: [gpu]
volumes:
- /home/[user_name]/.cache/huggingface:/root/.cache/huggingface
- ./fix_mtp_nvfp4_exclusion.py:/tmp/fix_mtp.py
- ./start_vllm.sh:/tmp/start_vllm.sh
environment:
- HF_HUB_OFFLINE=1
- TRANSFORMERS_OFFLINE=1
- VLLM_USE_FLASHINFER_MOE_FP4=0
- VLLM_TEST_FORCE_FP8_MARLIN=1
- VLLM_NVFP4_GEMM_BACKEND=marlin
- PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True
- MODEL_NAME=nvidia/Qwen3-Next-80B-A3B-Instruct-NVFP4
entrypoint: /bin/bash
command: /tmp/start_vllm.sh
then run sudo docker compose up and you may have the docker with the MTP fix for enabling the speculative decoding and the avarok/dgx-vllm-nvfp4-kernel:v22 launch without extra manual steps.
Did it somewhat quickly, some aspects most likely may be improved. Hope it helps for a quick launch and broader testing.
Okay. I created a working version on my branch. Here is the guide: dgx-vllm/MINIMAX_M25_GUIDE.md at optimize/cutlass-sparse-gemm · Avarok-Cybersecurity/dgx-vllm · GitHub
Please let me know if the NVFP4 image works or not. You will need 2 sparks, since the NVFP4 REAP version does NOT work (model was ruined during training). I am getting consistently 17-18 tok/s, which is suboptimal, but, we are also talking about a relatively large model.
without MTP:
docker run --rm --name dgx-vllm-nvfp4 \
--network host --gpus all --ipc=host \
-v "${HOME}/.cache/huggingface:/root/.cache/huggingface" \
-v /home/edison/Downloads/vllm/models:/models \
-e VLLM_USE_FLASHINFER_MOE_FP4=0 \
-e VLLM_TEST_FORCE_FP8_MARLIN=1 \
-e VLLM_NVFP4_GEMM_BACKEND=marlin \
-e PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True \
-e MODEL=/models/Qwen3-Next-80B-A3B-Instruct-NVFP4 \
-e PORT=8000 -e GPU_MEMORY_UTIL=0.70 \
-e MAX_MODEL_LEN=65536 -e MAX_NUM_SEQS=128 \
-e VLLM_EXTRA_ARGS="--attention-backend flashinfer --kv-cache-dtype fp8" \
avarok/dgx-vllm-nvfp4-kernel:v22 serve
llama-benchy --base-url "http://0.0.0.0:8888/v1" --model "/models/Qwen3-Next-80B-A3B-Instruct-NVFP4" --tokenizer "/home/edison/Downloads/vllm/models/Qwen3-Next-80B-A3B-Instruct-NVFP4" --pp 512 2048 8192 --tg 32 128 --runs 5
| model | test | t/s | peak t/s | ttfr (ms) | est_ppt (ms) | e2e_ttft (ms) |
|:------------------------------------------|-------:|-----------------:|-------------:|----------------:|----------------:|----------------:|
| /models/Qwen3-Next-80B-A3B-Instruct-NVFP4 | pp512 | 1421.66 ± 355.59 | | 403.34 ± 159.32 | 399.95 ± 159.32 | 403.41 ± 159.33 |
| /models/Qwen3-Next-80B-A3B-Instruct-NVFP4 | tg32 | 29.27 ± 1.67 | 29.93 ± 1.86 | | | |
| /models/Qwen3-Next-80B-A3B-Instruct-NVFP4 | pp512 | 1560.70 ± 43.31 | | 331.69 ± 9.14 | 328.31 ± 9.14 | 331.75 ± 9.14 |
| /models/Qwen3-Next-80B-A3B-Instruct-NVFP4 | tg128 | 28.42 ± 0.04 | 29.00 ± 0.00 | | | |
| /models/Qwen3-Next-80B-A3B-Instruct-NVFP4 | pp2048 | 2053.03 ± 9.73 | | 1000.86 ± 4.64 | 997.48 ± 4.64 | 1000.92 ± 4.63 |
| /models/Qwen3-Next-80B-A3B-Instruct-NVFP4 | tg32 | 28.32 ± 0.08 | 29.00 ± 0.00 | | | |
| /models/Qwen3-Next-80B-A3B-Instruct-NVFP4 | pp2048 | 2036.31 ± 13.92 | | 1009.17 ± 6.86 | 1005.79 ± 6.86 | 1009.24 ± 6.86 |
| /models/Qwen3-Next-80B-A3B-Instruct-NVFP4 | tg128 | 28.16 ± 0.09 | 29.00 ± 0.00 | | | |
| /models/Qwen3-Next-80B-A3B-Instruct-NVFP4 | pp8192 | 1922.55 ± 2.58 | | 4264.40 ± 5.70 | 4261.01 ± 5.70 | 4264.47 ± 5.70 |
| /models/Qwen3-Next-80B-A3B-Instruct-NVFP4 | tg32 | 27.71 ± 0.04 | 28.00 ± 0.00 | | | |
| /models/Qwen3-Next-80B-A3B-Instruct-NVFP4 | pp8192 | 1922.68 ± 5.77 | | 4264.15 ± 12.76 | 4260.77 ± 12.76 | 4264.22 ± 12.77 |
| /models/Qwen3-Next-80B-A3B-Instruct-NVFP4 | tg128 | 27.57 ± 0.09 | 28.00 ± 0.00 | | | |
with MTP:
docker run --rm --gpus all \
--name dgx-vllm-nvfp4 \
--network host \
--ipc=host \
-v /home/edison/Downloads/vllm/models:/models \
-v "$(pwd)/fix_mtp_nvfp4_exclusion.py:/tmp/fix_mtp.py" \
-v "${HOME}/.cache/huggingface:/root/.cache/huggingface" \
-e VLLM_USE_FLASHINFER_MOE_FP4=0 \
-e VLLM_TEST_FORCE_FP8_MARLIN=1 \
-e VLLM_NVFP4_GEMM_BACKEND=marlin \
-e PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True \
-e MODEL=/models/Qwen3-Next-80B-A3B-Instruct-NVFP4 \
-e PORT=8000 \
-e GPU_MEMORY_UTIL=0.70 \
-e MAX_MODEL_LEN=65536 \
-e MAX_NUM_SEQS=128 \
-e VLLM_EXTRA_ARGS="--attention-backend flashinfer --kv-cache-dtype fp8 --speculative-config {\"method\":\"mtp\",\"num_speculative_tokens\":2} --no-enable-chunked-prefill" \
--entrypoint bash avarok/dgx-vllm-nvfp4-kernel:v22 \
-c "python3 /tmp/fix_mtp.py && exec vllm serve \$MODEL --host 0.0.0.0 --port \$PORT --max-model-len \$MAX_MODEL_LEN --gpu-memory-utilization \$GPU_MEMORY_UTIL --max-num-seqs \$MAX_NUM_SEQS \$VLLM_EXTRA_ARGS"
llama-benchy --base-url "http://0.0.0.0:8888/v1" --model "/models/Qwen3-Next-80B-A3B-Instruct-NVFP4" --tokenizer "/home/edison/Downloads/vllm/models/Qwen3-Next-80B-A3B-Instruct-NVFP4" --pp 512 2048 8192 --tg 32 128 --runs 5
| model | test | t/s | peak t/s | ttfr (ms) | est_ppt (ms) | e2e_ttft (ms) |
|:------------------------------------------|-------:|-----------------:|-------------:|------------------:|------------------:|------------------:|
| /models/Qwen3-Next-80B-A3B-Instruct-NVFP4 | pp512 | 1069.15 ± 462.31 | | 1024.34 ± 1255.70 | 1021.68 ± 1255.70 | 1024.42 ± 1255.70 |
| /models/Qwen3-Next-80B-A3B-Instruct-NVFP4 | tg32 | 17.02 ± 1.46 | 18.22 ± 1.48 | | | |
| /models/Qwen3-Next-80B-A3B-Instruct-NVFP4 | pp512 | 1285.44 ± 9.84 | | 400.99 ± 3.06 | 398.33 ± 3.06 | 401.07 ± 3.05 |
| /models/Qwen3-Next-80B-A3B-Instruct-NVFP4 | tg128 | 16.29 ± 0.08 | 17.40 ± 0.49 | | | |
| /models/Qwen3-Next-80B-A3B-Instruct-NVFP4 | pp2048 | 1886.24 ± 7.93 | | 1088.44 ± 4.56 | 1085.77 ± 4.56 | 1088.49 ± 4.56 |
| /models/Qwen3-Next-80B-A3B-Instruct-NVFP4 | tg32 | 16.36 ± 0.50 | 17.07 ± 0.65 | | | |
| /models/Qwen3-Next-80B-A3B-Instruct-NVFP4 | pp2048 | 1875.47 ± 1.63 | | 1094.66 ± 0.95 | 1091.99 ± 0.95 | 1094.72 ± 0.95 |
| /models/Qwen3-Next-80B-A3B-Instruct-NVFP4 | tg128 | 16.10 ± 0.14 | 17.00 ± 0.00 | | | |
| /models/Qwen3-Next-80B-A3B-Instruct-NVFP4 | pp8192 | 1988.55 ± 5.71 | | 4122.29 ± 11.82 | 4119.63 ± 11.82 | 4122.35 ± 11.83 |
| /models/Qwen3-Next-80B-A3B-Instruct-NVFP4 | tg32 | 16.32 ± 0.52 | 17.19 ± 0.83 | | | |
| /models/Qwen3-Next-80B-A3B-Instruct-NVFP4 | pp8192 | 1991.25 ± 3.91 | | 4116.67 ± 8.07 | 4114.01 ± 8.07 | 4116.73 ± 8.07 |
| /models/Qwen3-Next-80B-A3B-Instruct-NVFP4 | tg128 | 15.90 ± 0.15 | 16.60 ± 0.49 | | | |
Oh, that was probably part of my trying to troubleshoot it. I’ll remove it.
Thanks for that efford. I was basically waitng for that to order my Spark. :-)
Just in case you haven’t seen the other thread, this doesn’t require any special builds and works with any recent vLLM build, including our community Docker out of the box. Thanks @tbraun96 for discovering the obscure VLLM flags, but it was already part of vLLM.
./launch-cluster.sh -t vllm-node-20260221 \
--solo \
-e VLLM_NVFP4_GEMM_BACKEND=marlin \
-e VLLM_TEST_FORCE_FP8_MARLIN=1 \
exec vllm serve nvidia/Qwen3-Next-80B-A3B-Instruct-NVFP4 \
--gpu-memory-utilization 0.7 \
--host 0.0.0.0 --port 8888 \
--max-model-len 128000 \
--load-format fastsafetensors
| model | test | t/s | peak t/s | ttfr (ms) | est_ppt (ms) | e2e_ttft (ms) |
|---|---|---|---|---|---|---|
| nvidia/Qwen3-Next-80B-A3B-Instruct-NVFP4 | pp2048 | 3317.78 ± 387.70 | 633.90 ± 78.21 | 626.41 ± 78.21 | 634.05 ± 78.31 | |
| nvidia/Qwen3-Next-80B-A3B-Instruct-NVFP4 | tg32 | 44.27 ± 1.96 | 45.71 ± 2.03 |
llama-benchy (0.3.2.dev5+g0dd75fc4b.d20260221)
date: 2026-02-22 10:13:03 | latency mode: api
Thanks for pointing and testing this out. The next challenge appears to leverage gb10+nvfp4 w/ flashinfer?
Well, I didn’t dig too deep into this but basically current high-performance FP4 implementation in flashinfer is restricted to sm10x where it can use tcgen05. GB10 has tensor cores too, just lacks dedicated tensor memory (TMEM). And shared on-chip memory is smaller too.
./launch-cluster.sh -t vllm-node \
--solo \
-e VLLM_NVFP4_GEMM_BACKEND=marlin \
-e VLLM_TEST_FORCE_FP8_MARLIN=1 \
exec vllm serve nvidia/Qwen3-Next-80B-A3B-Instruct-NVFP4 \
--gpu-memory-utilization 0.7 \
--host 0.0.0.0 --port 8888 \
--max-model-len 128000 \
--load-format fastsafetensors
llama-benchy --base-url "http://0.0.0.0:8888/v1" --model nvidia/Qwen3-Next-80B-A3B-Instruct-NVFP4 --pp 2048 --tg 32 --runs 5
llama-benchy (0.3.2.dev1+g17b42667a)
Date: 2026-02-23 10:37:59
Benchmarking model: nvidia/Qwen3-Next-80B-A3B-Instruct-NVFP4 at http://0.0.0.0:8888/v1
Concurrency levels: [1]
Loading text from cache: /home/edison/.cache/llama-benchy/cc6a0b5782734ee3b9069aa3b64cc62c.txt
Total tokens available in text corpus: 141518
Warming up...
Warmup (User only) complete. Delta: 8 tokens (Server: 29, Local: 21)
Warmup (System+Empty) complete. Delta: 13 tokens (Server: 34, Local: 21)
Measuring latency using mode: api...
Average latency (api): 3.00 ms
Running test: pp=2048, tg=32, depth=0, concurrency=1
Run 1/5 (batch size 1)...
Run 2/5 (batch size 1)...
Run 3/5 (batch size 1)...
Run 4/5 (batch size 1)...
Run 5/5 (batch size 1)...
Printing results in MD format:
| model | test | t/s | peak t/s | ttfr (ms) | est_ppt (ms) | e2e_ttft (ms) |
|:-----------------------------------------|-------:|-----------------:|-------------:|---------------:|---------------:|----------------:|
| nvidia/Qwen3-Next-80B-A3B-Instruct-NVFP4 | pp2048 | 2182.13 ± 144.04 | | 945.43 ± 60.56 | 942.44 ± 60.56 | 945.54 ± 60.52 |
| nvidia/Qwen3-Next-80B-A3B-Instruct-NVFP4 | tg32 | 30.95 ± 2.34 | 31.74 ± 2.69 | | | |
llama-benchy (0.3.2.dev1+g17b42667a)
date: 2026-02-23 10:37:59 | latency mode: api
I got about 30-32 T/s only.
When was the last time you built the container?
There was a commit that slowed down all Qwen3-Next models considerably (including FP8). Also, make sure you pull all the latest changes from the repo as building process changed a little bit last week (shouldn’t affect the speed though).
This is vLLM version I’m running currently: version 0.16.0rc2.dev371+gb71fbd06e.d20260222
We expect to see this, right?
(EngineCore_DP0 pid=205) WARNING 02-23 02:45:32 [marlin_utils_fp4.py:150] Your GPU does not have native support for FP4 computation but FP4 quantization is being used. Weight-only FP4 compression will be used leveraging the Marlin kernel. This may degrade performance for compute-heavy workloads.
Yes, that’s expected with Marlin kernel.
I ran git pull about 3 hours ago, and executed ./build-and-copy.sh. The tests above were done using the latest image built with ./build-and-copy.sh. When starting vLLM, it shows the version as: version 0.16.0rc2.dev385+g944ffb596.d20260223.
OK, let me try to rebuild with the most recent version and see what’s going on. There can be another regression.