NVFP4 inference on NVIDIA DGX Spark GB10 — finally faster than AWQ

The first open-source vLLM image to unlock full NVFP4 performance on NVIDIA DGX Spark. NVFP4 is ~20% faster than AWQ INT4 — and with MTP speculative decoding, it peaks at 111.9 tok/s from an 80B-parameter model on a single desktop GPU.

Read the full benchmark write-up: NVFP4_BREAKTHROUGH_DGX_SPARK.md

• Quick Start
• Why This Exists
• Performance
• Architecture
• Configuration
• File Reference
• Build History
• License

docker pull avarok/dgx-vllm-nvfp4-kernel:v22

docker run -d --name vllm-nvfp4 \
  --network host --gpus all --ipc=host \
  -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=nvidia/Qwen3-Next-80B-A3B-Instruct-NVFP4 \
  -e PORT=8888 -e GPU_MEMORY_UTIL=0.90 \
  -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

To enable MTP speculative decoding, add to VLLM_EXTRA_ARGS:

--speculative-config '{"method":"mtp","num_speculative_tokens":2}' --no-enable-chunked-prefill

Note: MTP requires fix_mtp_nvfp4_exclusion.py to be run inside the container before serving. Mount it and prepend to the entrypoint, or apply it at build time.

The Docker image takes 30–60 minutes to build. It compiles vLLM, CUTLASS kernels, and all SM121 patches from source.

git clone https://github.com/Avarok-Cybersecurity/dgx-vllm.git
cd dgx-vllm
docker build -t dgx-vllm:v22 .

The sparse_fp4_kernel/ directory contains a custom CUDA GEMV kernel that exploits natural 2:4 sparsity in NVFP4 weights, reading only the 2 non-zero values per group of 4. This cuts MoE memory traffic by ~25% and frees 9 GiB for MTP speculative decoding.

To build and install the kernel inside a running container:

# Start a container with the repo mounted
docker run -it --gpus all --ipc=host \
  -v $(pwd)/sparse_fp4_kernel:/workspace/sparse_fp4_kernel \
  dgx-vllm:v22 bash

# Inside the container:
cd /workspace/sparse_fp4_kernel
python setup_v7.py build_ext --inplace
cp sparse_fp4_v7*.so $(python -c "import site; print(site.getsitepackages()[0])")/
cp sparse_v7_moe_patch.py $(python -c "import site; print(site.getsitepackages()[0])")/

Then serve with the sparse kernel enabled:

# Inside the container:
cd /workspace/sparse_fp4_kernel
python vllm_serve_v7.py serve nvidia/Qwen3-Next-80B-A3B-Instruct-NVFP4 \
  --host 0.0.0.0 --port 8888 --max-model-len 4096 \
  --gpu-memory-utilization 0.90 --enforce-eager \
  --attention-backend flashinfer --kv-cache-dtype fp8

Or set the environment variable VLLM_NVFP4_SPARSE_V7=1 and import sparse_v7_moe_patch before model loading — the patch hooks into vLLM's weight-loading path automatically.

# Wait ~10 min for startup (model load + torch.compile + CUDA graphs)
curl http://localhost:8888/v1/models

curl -s http://localhost:8888/v1/chat/completions \
  -H "Content-Type: application/json" \
  -d '{"model":"nvidia/Qwen3-Next-80B-A3B-Instruct-NVFP4",
       "messages":[{"role":"user","content":"Hello!"}],
       "max_tokens":100}' | jq -r '.choices[0].message.content'

The DGX Spark's GB10 GPU has native FP4 tensor cores — but its SM 12.1 compute capability is missing a critical PTX instruction (cvt.rn.satfinite.e2m1x2.f32) required for NVFP4 quantization. Without it:

                        NVIDIA DGX Spark GB10
                Grace Blackwell Superchip (SM_121)
           119.7 GB Unified LPDDR5X @ 273 GB/s bandwidth

    FP4 Tensor Cores    mma.sync.m16n8k64.e2m1.e2m1     [WORKS]
    FP8 Tensor Cores    mma.sync.m16n8k32.e4m3.e4m3     [WORKS]
    FP4 Convert         cvt.rn.satfinite.e2m1x2.f32     [MISSING]

• Vanilla vLLM falls back to Python — 1.1 tok/s (unusable)
• NVIDIA's official image uses flashinfer-cutlass — ~36 tok/s (same as AWQ, no advantage)
• TensorRT-LLM hits a hard ceiling — 29.6 tok/s (CUTLASS cooperative-only scheduling)

We wrote a 15-line software E2M1 conversion function, enabled the Marlin MoE backend, and patched SM 12.1 capability routing. NVFP4 now runs at ~42 tok/s without speculative decoding — 20% faster than AWQ. With MTP, it reaches ~67 tok/s average and peaks at 111.9 tok/s.

MTP exceeds the single-token bandwidth ceiling by generating ~2.4 tokens per forward pass (63-89% draft acceptance).

See OPTIMIZATIONS.md for detailed analysis including failed experiments.

This image bridges GB10's hardware FP4 tensor cores and vLLM's inference engine through 7 layers of patches and integrations:

 Layer 7  Model        Qwen3-Next-80B-A3B (MoE, 512 experts, NVFP4)
            |
 Layer 6  vLLM V1      CUDA graphs, chunked prefill, FlashInfer attention,
            |           MoE routing, MTP speculative decoding
            |
 Layer 5  Patches      Qwen3Next prefix fix, EMULATION backend fix
            |           Capability 121 -> SM_120 routing, FlashInfer JIT
            |
 Layer 4  CUTLASS      FP4 MoE GEMM (BlockScaled, Cooperative, 4 tiles)
            |           FP4/FP8 scaled_mm (SM120 kernels)
            |
 Layer 3  Software     patch_nvfp4_utils_sw_e2m1.py
            |  E2M1     15-line device function replacing missing PTX
            |
 Layer 2  CUDA 13.0    nv_fp4_dummy.h (FP4 type definitions)
            |           CCCL + FlashInfer header patches
            |
 Layer 1  Base Image   nvidia/cuda:13.0.2-cudnn-devel-ubuntu24.04
            |           PyTorch 2.10+cu130, Triton 3.6.0
            |
 Layer 0  Hardware     GB10 GPU: SM_121, 119.7 GB LPDDR5X, ARM64 Grace

#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ == 1210
__device__ __forceinline__ uint8_t _sw_float_to_e2m1(float x) {
  uint8_t sign = (uint8_t)((__float_as_uint(x) >> 28) & 8u);
  float ax = fabsf(x);
  uint8_t mag;
  if      (ax <= 0.25f)  mag = 0;  // 0.0
  else if (ax <  0.75f)  mag = 1;  // 0.5
  else if (ax <= 1.25f)  mag = 2;  // 1.0
  else if (ax <  1.75f)  mag = 3;  // 1.5
  else if (ax <= 2.5f)   mag = 4;  // 2.0
  else if (ax <  3.5f)   mag = 5;  // 3.0
  else if (ax <= 5.0f)   mag = 6;  // 4.0
  else                    mag = 7;  // 6.0 (satfinite)
  return sign | mag;
}
#endif

IEEE 754 round-to-nearest-even for E2M1, matching hardware behavior exactly. Applied to all 5 NVFP4 kernel files via patch_nvfp4_utils_sw_e2m1.py.

docker run ... avarok/dgx-vllm-nvfp4-kernel:v22 serve        # vLLM API server (default)
docker run ... avarok/dgx-vllm-nvfp4-kernel:v22 ray-head     # Ray head node
docker run ... avarok/dgx-vllm-nvfp4-kernel:v22 ray-worker   # Ray worker node
docker run ... avarok/dgx-vllm-nvfp4-kernel:v22 bash         # Interactive shell

Built on vLLM (Apache 2.0) and NVIDIA CUDA containers.

Open source at github.com/Avarok-Cybersecurity/dgx-vllm.

Built by Avarok with Claude Code.