Last updated on

vLLM on the NVIDIA DGX Spark (GB10 / sm_121): A Complete Technical Guide

A companion, published after the field notes. I published the field-notes post, vLLM on the DGX Spark, the Hard Way, first: that’s the operational story, what broke, what we run in production, and the measured numbers. This post is the deep background behind it, the kernel, CUTLASS/FlashInfer, quantization, and backend-oracle detail I had to work through to understand why those field observations happen at all. Researching it was most of the actual work, and it seemed worth putting in one place for anyone hitting the same walls, so I’m releasing it as a companion. Read the field notes for what to do today; read this for why it works underneath. One caveat: this is a dated implementation audit (2026-06-02 against vLLM main), so treat the version-specific operational bits as a snapshot and defer to the field notes’ timeline for the current state.

Why this report exists. On 2026-06-01 the vLLM team published its DGX Spark deep-dive. We implemented its guidance on our production fleet the next day, and then ran it for seven weeks of real traffic. This report is what that produced: a commit-level audit of the implementation details the high-level recipe leaves out (SM121 arch targeting, backend oracle behavior, dense-vs-MoE split routing, FlashInfer cubin/JIT behavior, CUDA graph escape hatches), plus an honest scorecard of what worked, what didn’t survive production, and what we had to discover ourselves. Where the official blog and repository behavior diverge, the repository behavior is treated as authoritative for the audited commit. The body is a dated snapshot of that audit; the update timeline records what has changed between the audit and this post’s publication.

Scope. This is an all-in-one technical reference for running and understanding vLLM on the NVIDIA DGX Spark. It covers the hardware, why the platform has been hard for vLLM, the kernel libraries involved (CUTLASS, FlashInfer, Marlin, QuTLASS), how quantization (NVFP4 / MXFP4 / FP8) is wired up, how vLLM’s fused-MoE machinery selects backends, what --enforce-eager is and why it has been a crutch, the development history, and concrete guidance on prebuilt containers vs. source builds.

Provenance. Code citations are against the vLLM main checkout at commit 689e... (2026-06-02). Release-timing and the operational recommendations are reconciled with the official blog “vLLM on NVIDIA DGX Spark” (https://vllm.ai/blog/2026-06-01-vllm-dgx-spark). Where the repo and field experience diverge, the repo wins and the divergence is called out.

⏱ Freshness & maintenance: implementation audit dated 2026-06-02 (vLLM main @ 689e…). v0.22.0 is treated throughout as the first audited stable baseline containing the Spark fixes discussed, not as the latest release as of this post’s date (stable releases have continued since; see the update timeline). Re-audit backend selection and dependency pins against the exact image or tag you deploy. sm_121 support is a fast-moving target (vLLM stable cadence ≈ every 2–3 weeks), so claims carry inline date tags: [verified YYYY-MM-DD] = checked against the repo/blog on that date; [since vX.Y.Z] = release a fix first shipped in. Time-volatile sections to re-check every vLLM release: §4.4 (the b12x-is-opt-in exclusion, the single highest-churn fact here; it flips when flashinfer-ai/flashinfer#3134 lands upstream), §7.2–7.3 (backend auto-selection order), §9 (release→fix mapping), §10 (recommended image + dependency pins). The architectural material (§1, §3, §5–§6, §8) is structural and ages slowly. When you re-verify, bump the date tag even if nothing changed: a current date is the signal the claim still holds.

What’s changed since the audit

The body below is the 2026-06-02 audit snapshot. Between that audit and this post’s publication (body assembled 2026-07-20; this timeline last extended 2026-08-05), the following moved. Each entry is either verified in our own production fleet (details and receipts in the GB10 field notes) or checked against the public release record on the date shown.

  • 2026-06-05 → 2026-07-05: MTP under real concurrency [fleet-verified]: in-checkpoint MTP validated at 86.8% acceptance, then hit CUDA error: an illegal memory access every few hours in production, only under concurrent load (upstream #41190/#42084).
  • 2026-07-11/12: native FP4 confirmed working in production [fleet-verified]: a genuine W4A4 checkpoint (4-bit activations, lm_head unquantized) put dense and MoE layers on the FLASHINFER_CUTLASS native-FP4 path, faster than forced Marlin on our boxes, and it eliminated the MTP crash (8-way soak, 400/400, zero IMA). §6’s checkpoint-format emphasis aged well; any lingering “Marlin is mandatory on GB10” reading of §5 did not.
  • 2026-07-14: vLLM v0.25.1 released [PyPI, checked 2026-07-20]: multiple stable releases now postdate the audited v0.22.0 baseline.
  • 2026-07-19: official cu129 stable still ships no sm_121 cubins [fleet-verified]: smoke-testing the v0.25.1 stable candidates reproduced the report’s central lesson: wheel/cubin coverage, not version number, decides whether a path works. Our fleet runs digest-pinned nightlies for exactly this reason.
  • 2026-07-27: v0.26.0 released; cu129 aarch64 stable now HAS sm_121 NVFP4 cubins [fleet scratch-tested 2026-08-05]: auto selects FLASHINFER_CUTLASS for both dense and MoE with no no kernel image is available, which breaks the 2026-07-19 lesson above (the 0.22.0 → 0.25.1 cu129 cubin drought is over). It still could not serve our unsloth W4A4 checkpoint, though: EngineCore init dies on a CUTLASS FP8 W8A8 GEMM (cutlass_gemm_caller.cuh:62, Error Internal), a path the same checkpoint clears on our cu13-line nightly. The lesson refines rather than breaks: the failure moved from a missing NVFP4 kernel image to an FP8 CUTLASS error, still wheel/CUDA-line-specific, not version-specific.
  • 2026-07-29: #41190 (MTP illegal-memory-access) fixed on main [PR #48245/#49620]: root cause turned out to be a kernel-agnostic scheduler bug (async-decode output-placeholder underflow under KV-pressure preemption), not the drafter kernel. Not yet in a tagged stable; upstream marks it a probable “may fix.” Stable-wheel operators stay exposed until a nightly carries it.
  • 2026-08-05: #42084 (prefix-caching + spec-decode OOB) appears resolved on a v0.26.1-dev nightly [fleet scratch-tested]: prefix caching and MTP now run together via an experimental Mamba/SSM cache-align mode (it aligns the attention block size to the linear-attention page size). Validated with a shared-prefix + MTP soak (40% cache hit, 86.9% KV, zero OOB or asserts). This lifts the prefix-caching-vs-spec-decode either/or on recent nightlies, with the caveat that vLLM marks the align path experimental. See the field notes’ 2026-08-05 update for the run detail.
  • 2026-08-05: #31424 (NGC container lag) closed not_planned: the vendor-container lag is now permanent-by-neglect.
  • Re-verified 2026-08-05: §4.4’s b12x auto-selection exclusion still holds, checked in _POSSIBLE_NVFP4_KERNELS (dense) and the MoE oracle at v0.26.0 and main; b12x is excluded from auto on both paths and the flashinfer-ai/flashinfer#3134 guard has not landed. Still not re-audited, re-check before relying on: §7.2–7.3 (backend auto-selection order) and §10 (image and dependency pins).

The scorecard: what worked, what didn’t, what we found

We adopted the blog’s recipe in one commit (2026-06-02, one day after publication) and have run it under 24/7 agent traffic since. The verdict, with receipts:

What worked

  • The cu130-nightly image track: adopted as a gitsha-pinned nightly (0.22.1rc1.dev26, CUDA 13.0.2), per the blog’s own “pin a digest” clause. It unlocked native FP4 immediately: --moe-backend auto → FLASHINFER_CUTLASS, ~55.8–58.5 tok/s vs ~50 on forced Marlin (2026-06-02, batch-1 decode). Still the fleet’s serving basis today.
  • Trusting auto backend selection, on that image track. The blog’s claim that “recent FlashInfer CUTLASS paths are much stronger than older Spark guidance suggests” measured out on our boxes the same day.
  • CUDA graphs enabled. We retired --enforce-eager fleet-wide the same commit (it had been a v0.20.0 Mamba graph-capture workaround, §8) and never looked back.
  • Built-in --reasoning-parser nemotron_v3 over NVIDIA’s HF-card parser plugin: simpler, coherent, no plugin file.
  • The JIT pre-warm ping (max_tokens: 3 at boot): we run the same trick the blog describes.

What didn’t, or didn’t survive production

  • auto on the images that predated the blog. On cu129/v0.20.0 images, auto crashed; our cards from that era read “cu129 stock: auto crashes; marlin correct.” The blog’s advice was true of the image track it tested and wrong for the installed base that preceded it. Reconciling that (blog vs. repo vs. our logs) is how we learned it was a version gap, not a dispute, and it’s why every claim in this report carries a commit and date.
  • The demo-shaped numeric defaults. --gpu-memory-utilization 0.85 matches our own single-model figure, but co-residency forced 0.3–0.4, and taught us the unified-memory rule the hard way: the GPU reservation is unswappable, so the reservation itself is the memory pressure (0.6 ≈ 71 GiB got our 35B OOM-killed on 2026-07-03). --max-num-seqs 4 and --max-model-len 131072 became 4/8/16 and 32K/262K per box role. None of these are errors in the blog; they’re demo-shaped numbers that don’t transfer to a multi-model production box.
  • MTP as “evaluate.” We productionized --speculative-config MTP and paid the tuition: 86.8% acceptance in validation, then CUDA illegal memory access every few hours only under concurrent load (upstream #41190/#42084). The configuration that avoided it, a genuine W4A4 checkpoint putting experts and drafter on the CUTLASS path, is the report’s biggest self-discovered lesson (below). The blog’s “validate before relying on it” clause was fully vindicated; batch-1 validation was not enough.
  • The fp8 KV-cache caution: inverted by measurement. The blog frames --kv-cache-dtype fp8 as a memory-pressure concession (“may affect predictability… noticeable performance cost… avoid unless memory pressure requires it and quality checks pass”). We ran the quality check it prescribes (2026-07-20, A/B on the production 35B, temp-0 task-scored extraction + greedy probes): byte-identical extraction output across arms 10/10 docs, fp8 more deterministic intra-arm (10/10 vs 9/10), and ~27% faster decode (96.6 vs 76.3 tok/s, halved KV traffic on a bandwidth-bound box, with MTP re-reading KV every verification step), plus double the 262K-context capacity in the same pool (~3.2 vs ~1.6 sequences). Not an error in the blog (its escape clause is exactly what we executed), but on decode-bound long-context serving the cost-benefit framing runs the other way: fp8 KV is a default with a per-checkpoint validation step, not a concession.

What we found that the blog never mentions

  • The checkpoint’s quant format decides the FP4 backend, not your flags (W4A4 → native CUTLASS; W4A16 → Marlin fallback). The load-bearing rule of this entire report (§6), and the actual fix for the MTP crashes.
  • b12x is opt-in, not auto at the audited commit (§4.4): auto cannot select the Spark-native MoE path the blog’s era shipped.
  • Backend flags the recipe needs but the blog never names: --attention-backend (TRITON_ATTN / flashinfer per model), --mamba-ssm-cache-dtype float32 for hybrid-Mamba, --max-num-batched-tokens prefill chunking, and --async-scheduling, including the discovery that spec-decode configurations can silently disable async scheduling and tax prefill (§the n-gram lesson in our field notes).
  • Coordinator interactions: behind a llamactl-style proxy, --served-model-name 404s every request; the blog’s standalone recipe breaks unmodified behind a gateway.
  • Unified-memory operational arithmetic: KV-pool math from the boot log (per-token cost, boot floor ≈ 0.24 at 262K ctx, ≥15 GiB host headroom) that turns --gpu-memory-utilization from folklore into subtraction.

Table of contents

  1. The DGX Spark architecture
  2. Why DGX Spark has been hard for vLLM
  3. CUTLASS: the templated GEMM engine
  4. FlashInfer: the kernel-serving layer
  5. Marlin: the universal fallback
  6. NVFP4 / MXFP4 / FP8 quantization on Spark
  7. How vLLM does MoE, and how it lands on Spark
  8. CUDA graphs and --enforce-eager
  9. The history: how Spark support was built
  10. Prebuilt containers vs. rolling your own
  11. Appendix: glossary and file index

1. The DGX Spark architecture

1.1 What the box is

The DGX Spark is built on the NVIDIA GB10 Grace Blackwell superchip, a Grace ARM64/aarch64 CPU and a Blackwell GPU fused onto one package with a 128 GB unified LPDDR5X memory pool shared between CPU and GPU. This single fact drives most of the operational quirks: there is no separate “VRAM”; weights, KV cache, the OS page cache, and the container runtime all draw from the same 128 GB. Memory bandwidth (~273 GB/s class) is far below a datacenter card, which is why decode throughput is the binding constraint on Spark, not compute.

vLLM recognizes the unified-memory nature of the platform explicitly:

  • vllm/utils/mem_utils.py:110: “On UMA (Unified Memory Architecture) platforms where CPU and GPU share physical memory (e.g. GH200, DGX Spark, Jetson Orin)…”
  • vllm/platforms/interface.py:799: is_integrated_gpu() documents DGX Spark as a UMA device.

1.2 The compute capability that matters: sm_121

The GB10 GPU reports compute capability 12.1, i.e. sm_121 (architecture-specific target sm_121a). This is the single most important number in this entire document. It places Spark in the “SM12x family” of consumer/workstation Blackwell parts:

Part Compute capability nvcc target Class
RTX 5090 12.0 sm_120a consumer Blackwell
RTX PRO 6000 12.0 sm_120a workstation Blackwell
DGX Spark (GB10) 12.1 sm_121a workstation Blackwell
B200 / GB200 10.0 sm_100a datacenter Blackwell

vLLM’s own source documents the mapping in two kernel generators:

csrc/quantization/marlin/generate_kernels.py:18 and csrc/moe/marlin_moe_wna16/generate_kernels.py:18: # SM89 and the SM12x family (SM120 RTX 5090, SM121 DGX Spark GB10)

1.3 The trap: sm_121 is not sm_100

Datacenter Blackwell (sm_100) and workstation Blackwell (sm_120/sm_121) are different architectures with different tensor-core instruction sets, even though both are marketed as “Blackwell.” A kernel compiled and tuned for sm_100 will not run on sm_121. In CUTLASS terms they require entirely separate kernel instantiations: different BlockwiseScaleConfig, different warp-specialized kernel schedules, and different cutlass::arch::Sm1xx tags (see §3). The practical consequence: almost all generic “Blackwell NVFP4” guidance on the internet is written for sm_100 and does not transfer to Spark. This is the root cause of nearly every Spark gotcha.

1.4 What the tensor cores natively support

The Marlin generator records the key hardware fact about the SM12x family:

csrc/quantization/marlin/generate_kernels.py:18-23:

# SM89 and the SM12x family (SM120 RTX 5090, SM121 DGX Spark GB10)
# fully support mma.sync.aligned.m16n8k32.row.col.f32.e4m3.e4m3.f32.
# SM90 and SM100 can use this PTX, but it's simulated
# with FP16 MMA, so it cannot achieve any acceleration.
if arch == 89 or arch // 10 == 12:
    SUPPORT_FP8 = True

So sm_121 has native FP8 (e4m3) tensor-core MMA, and (via OpClassBlockScaledTensorOp) native block-scaled FP4 tensor cores. The hardware is capable; the historical problems were entirely about software coverage for this specific architecture.

2. Why DGX Spark has been hard for vLLM

Four compounding reasons:

1. Architecture fragmentation + arch-specific cubins. Modern quantized GEMM kernels (CUTLASS, FlashInfer’s CuteDSL kernels) are compiled into architecture-specific cubins. A wheel that ships sm_100a and sm_120a cubins but not sm_121a will fail at runtime on Spark, not with a graceful fallback, but with a hard error or, worse, silent garbage. The canonical symptom:

csrc/.../fp4/nvfp4_scaled_mm_entry.cu:66-68 (runtime SM dispatch):

STD_TORCH_CHECK_NOT_IMPLEMENTED(
    false, "No compiled nvfp4 mm kernel for SM ", sm,
    ". Recompile with CUDA >= 12.8 and CC >= 100.");

2. sm_121 is the newest, least-covered target. sm_120 (RTX 5090) support landed in vLLM in mid-2025; sm_121 is a strict latecomer that had to be threaded through every kernel’s arch list one at a time (see §9). For a long window, sm_121 was partially covered: some kernels had it, others silently excluded it, producing inconsistent behavior.

3. aarch64. Spark is ARM64. Historically the CUDA Python ecosystem was x86-first; aarch64 CUDA wheels were patchy. This is now largely solved (PyTorch ≥2.11.0 ships aarch64 CUDA wheels on PyPI, and vLLM publishes aarch64 wheels), but it added friction during the early Spark period and still bites anyone on older torch.

4. CUDA toolkit floor. sm_121 kernels require CUDA ≥ 12.8 to even compile; CUDA 13.0 is preferred (it introduces family targets, see §3.4). Building against an older toolkit silently drops Spark from the arch list:

CMakeLists.txt:99-110:

if(CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL 13.0)
  set(CUDA_SUPPORTED_ARCHS "7.5;8.0;8.6;8.7;8.9;9.0;10.0;11.0;12.0")   # family targets
elseif(CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL 12.8)
  set(CUDA_SUPPORTED_ARCHS "7.5;8.0;8.6;8.7;8.9;9.0;10.0;10.1;10.3;12.0;12.1")
else()
  set(CUDA_SUPPORTED_ARCHS "7.0;7.5;8.0;8.6;8.7;8.9;9.0")               # no sm_121 at all
endif()

The rest of this guide explains the moving parts that these four problems touch.

3. CUTLASS: the templated GEMM engine

3.1 What CUTLASS is and why it matters

CUTLASS (CUDA Templates for Linear Algebra Subroutines) is NVIDIA’s open-source C++ template library for building high-performance GEMMs directly on the tensor cores. vLLM uses CUTLASS 3.x (“c3x”) to implement its first-choice quantized matmuls: FP8 W8A8, NVFP4 W4A4, and the blockwise-scaled variants. When CUTLASS works for your arch + dtype, it is typically the fastest path, which is why the backend oracles try it before falling back to Marlin.

Key files:

  • csrc/quantization/w8a8/cutlass/scaled_mm_c3x_sm120.cu: FP8/INT8 scaled-MM entry for SM120/121
  • csrc/quantization/w8a8/cutlass/c3x/scaled_mm_sm120_fp8_dispatch.cuh: per-tensor FP8, with M-shape-based tile selection
  • csrc/quantization/w8a8/cutlass/c3x/scaled_mm_blockwise_sm120_fp8_dispatch.cuh: blockwise FP8
  • csrc/quantization/fp4/nvfp4_scaled_mm_sm120_kernels.cu: NVFP4 GEMM
  • csrc/quantization/fp4/nvfp4_blockwise_moe_kernel.cu: NVFP4 blockwise MoE
  • csrc/quantization/fp4/nvfp4_scaled_mm_entry.cu: runtime SM dispatch (the source of the “No compiled nvfp4 mm kernel for SM …” error)

3.2 Why CUTLASS needs separate kernels for sm_100 vs sm_120/sm_121

CUTLASS kernels are parameterized by an architecture tag that selects the MMA atoms and the warp-specialized scheduling strategy. Datacenter and workstation Blackwell get different tags and different schedules:

Aspect sm_100 (datacenter) sm_120/sm_121 (workstation)
Arch tag cutlass::arch::Sm100 cutlass::arch::Sm120
Blockwise scale config Sm100BlockwiseScaleConfig Sm120BlockwiseScaleConfig
Kernel schedule KernelTmaWarpSpecializedBlockwise{1,2}SmSm100 KernelTmaWarpSpecializedBlockwise{Pingpong,Cooperative}Sm120

These are distinct compiled kernels; there is no shared binary. That is exactly why “it works on my B200” tells you nothing about Spark.

3.3 The enable_sm120_family guard: one binary for SM120 and SM121

The mechanism that lets a single compiled kernel serve both RTX-class (sm_120) and Spark (sm_121) is a thin device-side wrapper:

csrc/cutlass_extensions/common.hpp (≈ lines 128-142):

// SM12x family includes SM120 (RTX 5090) and SM121 (DGX Spark GB10)
template <typename Kernel>
struct enable_sm120_family : Kernel {
  template <typename... Args>
  CUTLASS_DEVICE void operator()(Args&&... args) {
#if defined __CUDA_ARCH__
  #if (__CUDA_ARCH__ >= 1200 && __CUDA_ARCH__ < 1300)
    Kernel::operator()(std::forward<Args>(args)...);
  #else
    printf("This kernel only supports sm120f.\n");
    asm("trap;");
  #endif
#endif
  }
};

The range check __CUDA_ARCH__ >= 1200 && < 1300 admits 1200 (sm_120) and 1210 (sm_121), the whole SM12x family. Several Spark bugs over the past months were precisely cases where a kernel used a narrower sm_120-only guard and thus trapped or fell through on sm_121; the fixes consisted of routing those kernels through enable_sm120_family (e.g. the per-tensor FP8 dispatch fix, PR #41215).

3.4 The a vs f suffix, and cuda_archs_loose_intersection

nvcc arch targets carry a suffix:

  • 12.1a (architecture-specific): one cubin for exactly sm_121. To cover the family you must list every member: 12.0a;12.1a.
  • 12.0f (family): one cubin that runs across the whole SM12x family: sm_120, sm_121, and future SM12x. Requires CUDA ≥ 13.0. (See NVIDIA’s family-specific-architecture docs.)

vLLM’s CMake picks the right form by toolkit version. For example, scaled-MM:

CMakeLists.txt:755-758:

if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
  cuda_archs_loose_intersection(SCALED_MM_ARCHS "12.0f" "${CUDA_ARCHS}")
else()
  cuda_archs_loose_intersection(SCALED_MM_ARCHS "12.0a;12.1a" "${CUDA_ARCHS}")
endif()

The same pattern appears for FP4_ARCHS, MLA_ARCHS, MARLIN_*_ARCHS, etc. The helper cuda_archs_loose_intersection (cmake/utils.cmake:368-416) is what makes the family form match a user’s specific arch: given source 12.0f and a target list containing 12.1a, it strips suffixes, compares the major version (both 12), and matches, so a single 12.0f build covers Spark. This family-level matching (added in PR #38126, “Fix DGX Spark logic”) is why a CUDA-13 build with just 12.0f correctly serves sm_121.

Operator takeaway: if you build from source on CUDA 12.8/12.9 you must list 12.1a (or 12.0a;12.1a); on CUDA 13.0+ 12.0f suffices. The in-repo Docker Grace-Blackwell example uses '9.0 10.0+PTX'. That is for GH200/GB200 and is wrong for GB10; you must override it.

4. FlashInfer: the kernel-serving layer

4.1 What FlashInfer is in vLLM

FlashInfer is a kernel library that vLLM leans on heavily for inference-optimized kernels. vLLM’s wrapper is vllm/utils/flashinfer.py (~1000 lines). Surface area spans:

  • Attention: paged/ragged batch decode & prefill, TRT-LLM attention backends, MLA (DeepSeek), FP4 KV cache.
  • Fused MoE: TRT-LLM BF16/FP8/FP4 block-scale MoE, CUTLASS fused MoE, CuteDSL NVFP4 MoE, and the SM12x b12x fused MoE.
  • Quantized GEMM: FP4/MXFP4 quantize, NVFP4 scaled MM, block-scale FP8.
  • RoPE and collective all-reduce (MNNVL/NVLink).

4.2 How FlashInfer ships kernels: prebuilt cubins vs. JIT

This is the crux of FlashInfer’s relationship with Spark. FlashInfer kernels arrive one of two ways (vllm/utils/flashinfer.py:27-63):

  • Prebuilt cubins: the flashinfer-cubin package (pinned flashinfer-cubin==0.6.11.post2 in requirements/cuda.txt:13) downloads precompiled cubins from an NVIDIA artifact repository. Fast, but only covers the arches NVIDIA chose to build.
  • JIT: if cubins are missing, FlashInfer compiles kernels at first request, which requires nvcc to be present in the container: 
    if not has_flashinfer_cubin() and shutil.which("nvcc") is None:
    # FlashInfer unavailable: no cubins and no nvcc to JIT
    return False

Why Spark broke: for a long while the prebuilt FlashInfer wheels shipped zero sm_120/sm_121 cubins, and stock inference images don’t carry nvcc (it’s in the devel image, not runtime). So FlashInfer’s FP4 path either errored (Failed to run cutlass FP4 gemm on sm120/sm121 … Error Internal) or (the dangerous case) produced incoherent tokens with no error. This is the original sin behind the field rule “force Marlin on Spark.”

4.3 b12x: the modern SM12x MoE / FP4 path

The newer, Spark-native path is b12x (FlashInfer’s SM12x block-scaled GEMM, built on CuteDSL, the CUTLASS C++ tensor-layout DSL). It exists in two forms:

  • MoE expert: vllm/model_executor/layers/fused_moe/experts/flashinfer_b12x_moe.py:

    “FlashInfer CuteDSL fused MoE expert for SM12x (SM120/SM121, RTX Pro 6000 / DGX Spark). Uses b12x_fused_moe … which fuses token dispatch, two GEMMs, SwiGLU activation, and topk-weight reduction into a single kernel call. Input quantization (BF16→FP4) is performed inside the kernel … Only NVFP4 … is supported.”

    Notably, it takes BF16 hidden states directly and quantizes inside the kernel; scale factors are pre-converted to MMA layout once in process_weights_after_loading.

  • Dense linear: vllm/model_executor/kernels/linear/nvfp4/flashinfer.py:221:

    class FlashInferB12xNvFp4LinearKernel(NvFp4LinearKernel): """NVFP4 GEMM via FlashInfer's b12x CuTe DSL warp-level MMA kernel (SM120+)."""

Both require CC ≥ 12.0. b12x landed in vLLM in PR #40082 (first released in v0.22.0).

4.4 The critical nuance: b12x is opt-in, not auto [verified 2026-06-02: highest-churn fact; re-check each release]

This is the single most under-appreciated fact for Spark operators, and it refines the common “just leave backends on auto” advice. The official blog recommends leaving backend selection on auto, but at the audited vLLM main commit (689e…, 2026-06-02), b12x is deliberately excluded from auto-selection until an upstream CUTLASS sm_121 MMA-guard issue (flashinfer-ai/flashinfer#3134) is resolved:

vllm/model_executor/layers/fused_moe/oracle/nvfp4.py:170-173:

# NOTE: the kernels are selected in the following order.
# FLASHINFER_B12X is intentionally excluded from auto-selection until
# the upstream CUTLASS SM121 MMA op guard is resolved; use
# moe_backend="flashinfer_b12x" to opt in explicitly.

and the matching note for dense linears in vllm/model_executor/kernels/linear/__init__.py:383-395:

# FlashInferB12xNvFp4LinearKernel excluded from auto-selection until
# upstream CUTLASS SM121 MMA op guard is resolved; use
# VLLM_NVFP4_GEMM_BACKEND=flashinfer-b12x to opt in explicitly.

So on that code path, --moe-backend auto will not pick the Spark-native b12x kernel; it walks the list below (§7) and lands on whatever else supports sm_121. To actually use b12x you pass --moe-backend flashinfer_b12x (and/or --linear-backend flashinfer-b12x). This is why the audited-commit Spark recipe is either “force Marlin,” “explicitly opt into b12x and validate,” or, as our production fleet later confirmed, get onto the FLASHINFER_CUTLASS native path with a genuine W4A4 checkpoint. Re-check this against the exact release you deploy: it is expected to change once the upstream guard lands.

Update [verified 2026-08-10]: the exclusion still holds at v0.26.0 and main, but two details in the quoted source comments have since moved — treat the comments above as a 2026-06-02 snapshot:

  • The env-var opt-in is gone. VLLM_NVFP4_GEMM_BACKEND was deprecated in v0.23 (deprecated_env(..., "v0.23", "Use --linear-backend.")) and removed in v0.25. On any current image, VLLM_NVFP4_GEMM_BACKEND=flashinfer-b12x is a silent no-op — stale recipes still hand it out. The only surviving opt-in is the CLI value: --linear-backend flashinfer_b12x / --moe-backend flashinfer_b12x.
  • The exclusion is global, not sm_121-scoped. FLASHINFER_B12X is absent from AVAILABLE_BACKENDS on every GPU — an sm_120 RTX PRO 6000 also falls to Marlin under auto. sm_121 is the reason cited in the comment, not the scope. The still-open PR #47577 would re-add auto-selection for sm_120 only and deliberately keep sm_121 on the explicit opt-in — auto keeps skipping b12x and falling through to whatever the image + checkpoint actually support (FLASHINFER_CUTLASS on a genuine W4A4 checkpoint with sm_121 cubins, Marlin on a W4A16 checkpoint — see §4.4), not an unconditional Marlin default. So b12x remains opt-in-only, never auto-selected on a Spark, in every released version — and forcing it currently carries a cluster of open bugs (activation memory roughly doubling → ~55% less KV cache, #49079; an Xid 31 MMU fault under concurrent chunked prefill, #50189).

4.5 flashinfer_cutlass_moe vs flashinfer_b12x_moe

flashinfer_cutlass_moe (older) flashinfer_b12x_moe (newer)
Dispatch per-block GEMM single fused kernel (dispatch + 2 GEMMs + SwiGLU + reduce)
Quant FP8 block, MXFP4, static NVFP4 NVFP4 only
Input pre-quantized activations BF16 (quantize inside kernel)
SM target SM90 / SM100(+) SM120+
Auto-selected? yes (in the list) no, opt-in

PR #39825 (“Disable FlashInfer CUTLASS MoE on SM121”) was the stopgap that kept the older path from misbehaving on Spark before b12x existed.

5. Marlin: the universal fallback

5.1 What Marlin is

Marlin (originating from IST-DASLab, integrated via Neural Magic; csrc/quantization/marlin/marlin.cu) is a mixed-precision, weight-only quantized GEMM: 16-bit activations (FP16/BF16) multiplied by low-bit weights. The “wna16” naming = weight n-bit, activation 16-bit. It supports a broad matrix of weight formats (csrc/quantization/marlin/generate_kernels.py:29-100): AWQ-INT4, GPTQ-INT4/INT8, FP8 (kFE4M3fn), NVFP4 (kFE2M1f weights, kFE4M3fn scales), MXFP4, and MXFP8.

5.2 Why Marlin is the robust fallback on Spark

Two structural reasons:

1. It is template-generated, not arch-pinned. Marlin emits kernels from Jinja2 templates parameterized by thread/tile/group config, so the same source covers sm_80, sm_89, sm_120, sm_121 without bespoke per-arch kernel files the way CUTLASS does. Its CMake arch lists explicitly include the SM12x family:

CMakeLists.txt: MARLIN_FP8_ARCHS = "8.9;12.0a;12.1a" (CUDA<13) / "8.9;12.0f" (CUDA≥13); MARLIN_MOE_FP8_ARCHS = "8.9;12.0;12.1".

2. It covers both dense linears and MoE. Dense: csrc/quantization/marlin/. MoE: csrc/moe/marlin_moe_wna16/. So forcing Marlin fixes the whole model in one move, which is exactly why it became the Spark workaround when CUTLASS/FlashInfer FP4 paths were broken for both the dense and MoE halves of a model.

5.3 Marlin for NVFP4

Marlin has a real FP4 path. The NVFP4 linear kernel is MarlinNvFp4LinearKernel (“NVFP4 weight-only GEMM via Marlin (W4A16)”, vllm/model_executor/kernels/linear/nvfp4/marlin.py:18). It performs a non-trivial scale repacking, converting the FP8-E4M3 block scales into Marlin’s internal “S0E5M3” dequant layout (vllm/.../utils/marlin_utils_fp4.py:53-114). You route to it explicitly with VLLM_NVFP4_GEMM_BACKEND=marlin (env, deprecated per the blog) or, preferably, the CLI --linear-backend marlin / --moe-backend marlin.

5.4 Where Marlin sits in selection, and #40923

In the NVFP4 MoE oracle (§7) Marlin sits near the end of the auto list, a deliberate “always works, decent speed” backstop after the CUTLASS/FlashInfer attempts. PR #40923 (“Marlin MoE: include SM 12.x in default arch list”) added the SM12x family to the default Marlin-MoE build arch list so stock builds carry Spark Marlin-MoE kernels out of the box (currently only in v0.22.1rc0).

Cost of Marlin: as a weight-only W*A16 kernel it dequantizes weights to 16-bit and runs the matmul in 16-bit, so it does not exploit Spark’s native FP4/FP8 tensor cores as fully as a true block-scaled CUTLASS/b12x kernel can. On a bandwidth-bound box like Spark this is often acceptable, but it’s the reason Marlin is the fallback, not the goal.

6. NVFP4 / MXFP4 / FP8 quantization on Spark

6.1 NVFP4 format

NVFP4 is a 4-bit float with two-level block scaling:

  • Elements: e2m1 (2-bit exponent, 1-bit mantissa; max 6.0), packed 2-per-byte.
  • Per-block scale: FP8 e4m3, one per 16 contiguous elements along K.
  • Global scale: an fp32 scalar (one for weights, one for activations); the kernel combines them as alpha = 1 / (input_global_scale * weight_global_scale).

So the effective value is element × fp8_block_scale × fp32_global_scale. (See vllm/.../quantization/modelopt.py around the NVFP4 method, and csrc/.../fp4/nvfp4_utils.cuh.) This block structure is why NVFP4 needs block-scaled tensor-core MMA (OpClassBlockScaledTensorOp) to run fast, which only sm_100/sm_120/sm_121 provide, in arch-specific form.

6.2 The dense-vs-MoE split: the gotcha that wasted everyone’s time

NVFP4 weights appear in two independent code paths:

  • Dense linear layers → ModelOptNvFp4LinearMethod, served by the linear kernel oracle (cutlass_scaled_fp4_mm, FlashInfer, or Marlin).
  • MoE expert weights → the fused-MoE oracle (§7).

These are selected independently. The historical Spark trap: the model-card env vars VLLM_USE_FLASHINFER_MOE_FP4=1 / VLLM_FLASHINFER_MOE_BACKEND=… only reroute the MoE half. The dense NVFP4 linears keep going through the broken CUTLASS FP4 path and crash or emit garbage. The only single-switch fix that covers both halves is Marlin (--linear-backend marlin --moe-backend marlin), or, on current images, explicitly opting both into b12x.

6.3 The kernels that consume NVFP4, and the runtime SM dispatch

The CUTLASS FP4 entry point dispatches on the runtime SM and fails loudly if the matching cubin wasn’t compiled in:

csrc/.../fp4/nvfp4_scaled_mm_entry.cu (≈ 48-68):

const int32_t sm = get_sm_version_num();
#if defined(ENABLE_NVFP4_SM100) ...
  if (sm >= 100 && sm < 120) { cutlass_scaled_fp4_mm_sm100a(...); return; }
#endif
#if defined(ENABLE_NVFP4_SM120) ...
  if (sm >= 120 && sm < 130) { cutlass_scaled_fp4_mm_sm120a(...); return; }   // SM120/121
#endif
  STD_TORCH_CHECK_NOT_IMPLEMENTED(false, "No compiled nvfp4 mm kernel for SM ", sm, ...);

This device guard + runtime dispatch was added/hardened in PR #29711. The lesson: a wheel built without ENABLE_NVFP4_SM120 / the 12.1a arch simply has no Spark FP4 kernel, and you hit the NOT_IMPLEMENTED at first FP4 matmul.

6.4 NVFP4 vs MXFP4 (and who uses which)

NVFP4 MXFP4
Element e2m1 (4-bit) e2m1 (4-bit)
Block size (K) 16 32
Block scale FP8 e4m3 e8m0 (exponent-only)
Typical models Qwen, Nemotron-3 GPT-OSS (MoE)

MXFP4 (vllm/.../quantization/mxfp4.py) is what the GPT-OSS family ships in. The block-scaled FP4/MXFP4 CUTLASS kernels for sm_120/sm_121 are partly provided through QuTLASS (cmake/external_projects/qutlass.cmake, fetched from IST-DASLab; requires CUDA 12.8+ and CC 10.0f/12.0f).

6.5 FP8: the path of least resistance on Spark

Because sm_121 has native e4m3 FP8 MMA (§1.4) and FP8 is covered by both CUTLASS (scaled_mm_sm120_fp8_dispatch) and Marlin, FP8 is usually far less fragile than NVFP4 on Spark; on the audited paths it needed no Marlin dance. Treat that as image- and model-dependent rather than a law (Spark FP8 failures have surfaced in specific nightlies and models); smoke-test the exact model and container before relying on auto. On a 128 GB unified box the memory saving of FP4 over FP8 is often irrelevant for the models people actually run, so a practical rule is: prefer FP8 (or BF16) unless a model is only published in NVFP4/MXFP4.

7. How vLLM does MoE, and how it lands on Spark

7.1 The modular fused-MoE architecture

vLLM decomposes a fused-MoE layer into composable stages (vllm/model_executor/layers/fused_moe/modular_kernel.py:45-79):

[Router] → [Quantize-Dispatch] → [Permute-Experts-Unpermute] → [Combine]

implemented via three classes: FusedMoEPrepareAndFinalizeModular (prep + finalize), FusedMoEExpertsModular (the expert GEMMs + activation + optional quant), and FusedMoEModularKernel (the glue). Each concrete backend is an “experts” implementation under vllm/model_executor/layers/fused_moe/experts/ (e.g. flashinfer_b12x_moe.py, marlin_moe.py, triton_moe.py, flashinfer_cutlass_moe.py).

7.2 The oracle: how a backend is chosen [verified 2026-06-02]

For NVFP4 MoE, select_nvfp4_moe_backend() (oracle/nvfp4.py:160-333) walks a priority list and asks each backend whether it supports the current config/device:

AVAILABLE_BACKENDS = [
    NvFp4MoeBackend.FLASHINFER_TRTLLM,
    NvFp4MoeBackend.FLASHINFER_CUTEDSL,
    NvFp4MoeBackend.FLASHINFER_CUTEDSL_BATCHED,
    NvFp4MoeBackend.FLASHINFER_CUTLASS,
    NvFp4MoeBackend.VLLM_CUTLASS,
    NvFp4MoeBackend.MARLIN,
    NvFp4MoeBackend.EMULATION,
]
# FLASHINFER_B12X intentionally excluded from auto-selection (opt in explicitly)

The --moe-backend CLI value maps onto this enum in map_nvfp4_backend() (oracle/nvfp4.py:141-157), and the available options + their docstrings (including “flashinfer_b12x: … for SM12x (RTX Pro 6000 / DGX Spark)”) live in vllm/config/kernel.py:171-190.

7.3 What this means on Spark

  • --moe-backend auto on sm_121 will skip b12x (opt-in only, at the audited commit, §4.4) and try TRT-LLM / CuteDSL / CUTLASS paths, then VLLM_CUTLASS, then Marlin, then emulation. Whether the earlier entries actually run on sm_121 depends on the image’s cubins; this is why auto behavior on Spark has historically been a coin-flip and why people pinned Marlin.
  • For a predictable Spark deployment you have two solid choices:
    1. Force Marlin for both linear and MoE: guaranteed-correct, slightly slower.
    2. Explicitly opt into b12x (--moe-backend flashinfer_b12x, --linear-backend flashinfer-b12x) on a v0.22.0+ image (the fast native path) and validate output coherence before trusting it (the b12x exclusion-from-auto exists precisely because the upstream guard isn’t fully settled).

8. CUDA graphs and --enforce-eager

8.1 What CUDA graphs do, and what enforce_eager turns off

vLLM captures the decode forward pass into a CUDA graph so the GPU replays a recorded sequence of kernel launches with near-zero per-launch CPU overhead. enforce_eager disables this entirely:

vllm/config/model.py: enforce_eager: bool = False: “If True, we will disable CUDA graph and always execute the model in eager mode. If False, we will use CUDA graph and eager execution in hybrid for maximal performance and flexibility.”

When set, vLLM forces cudagraph_mode = NONE and clears the capture-size list.

8.2 Piecewise vs. full capture (and the V1 default)

vllm/config/compilation.py:599-634 defines the modes:

  • NONE: no capture.
  • PIECEWISE: capture everything except graph-incompatible ops (e.g. some attention), which run outside the graph for flexibility.
  • FULL / FULL_DECODE_ONLY.
  • FULL_AND_PIECEWISE: full graphs for decode batches, piecewise for prefill/mixed. This is the V1 default and the most performant for most models.

Graphs are captured at a discrete set of batch sizes (cudagraph_capture_sizes), auto-generated roughly as [1,2,4] + range(8,256,8) + range(256, max, 16), with max_cudagraph_capture_size defaulting to min(max_num_seqs*2, 512). This is the mechanism behind the field trick of capping --max-num-seqs to bound the capture set.

8.3 Why Mamba / hybrid (nemotron_h) models need eager on Spark

First, the scope: the official blog’s benchmark path kept CUDA graphs enabled; this section covers a separate field failure mode, not a blanket Spark rule. Hybrid SSM/Mamba models maintain recurrent state and use Triton kernels with persistent in-place state updates. Graph capture freezes kernel-launch semantics and pointer bindings, which collides with that dynamic state; hence Mamba models implement special copy_inputs_before_cuda_graphs / get_seqlen_agnostic_capture_inputs hooks (vllm/model_executor/models/mamba.py:270-274, mamba2.py, jamba.py, nemotron_h.py). On sm_121 specifically, graph capture of the Mamba-2 Triton path has surfaced an “illegal instruction” during capture in our deployments, so the operational fix has been --enforce-eager, applied only when a specific model/image combination shows graph-capture instability, and documented with the tested model, image, CUDA version, and error.

8.4 Why eager is a crutch, not a destination

enforce_eager pays for itself in launch overhead on every decode step; field measurements on the Nemotron hybrid models put the cost around ~37% throughput on Spark, and the broader range is roughly 5–30% depending on batch size and model depth. It’s an escape hatch for a kernel/capture incompatibility, not a tuning choice. The right long-term fix is upstream: make the Mamba-2 Triton kernels capture-safe on sm_121 (or provide a CUDA-graph-compatible SSM path), at which point Spark hybrid models can drop --enforce-eager and recover that throughput. Until then it remains model-specific: apply it to nemotron_h/Mamba hybrids, not to dense or standard-MoE models that capture fine.

9. The history: how Spark support was built

Support arrived in two waves: RTX-class sm_120 first, then Spark sm_121 threaded in behind it. Release attributions below are verified with git tag --contains [verified 2026-06-02].

Wave 1, sm_120 (RTX Blackwell) groundwork, mid-2025:

  • #17280 CUTLASS W8A8 FP8 for sm_120; #21309 NVFP4 W4A4 for sm_120; #22131 block FP8; #29242 NVFP4 MoE CUTLASS; #29711 device guard + runtime SM dispatch for cutlass_scaled_fp4_mm.

Wave 2, sm_121 (DGX Spark) era, 2026:

PR What First release
#33517 (2026-02-06) enable_sm120_family template: first explicit DGX Spark commit; one binary for SM120+SM121 early 2026
#38126 (2026-03-27) “Fix DGX Spark logic”: CMake family matching (12.0f↔12.1a), arch lists gain 12.1 v0.19.0
#38423 (2026-03-30) NVFP4 bugfixes for DGX Spark + RTX50 (post-v0.19)
#39825 (2026-04-14) Disable FlashInfer CUTLASS MoE on SM121 (stopgap) v0.20.0
#35568 (2026-05-15) Fix incorrect SM121 exclusion from Marlin/CUTLASS FP8 paths v0.22.0
#40082 (2026-05-20) Integrate FlashInfer b12x MoE + FP4 GEMM for SM120/121 v0.22.0
#41215 (2026-05-20) Route per-tensor FP8 through enable_sm120_family on SM12.1 v0.22.0
#40923 (2026-05-28) Marlin MoE: include SM12.x in default arch list v0.22.1rc0

The key takeaway: the cluster of fixes that make Spark coherent end-to-end, the SM121 FP8 exclusion fix (#35568) and the b12x integration (#40082), first ship together in v0.22.0 (2026-05-29). Anything older (including the popular v0.20.0) lacks b12x and the exclusion fix; it can be made to work only by forcing Marlin. Release cadence is ~2–3 weeks.

10. Prebuilt containers vs. rolling your own

10.1 Decision summary [verified 2026-06-02: image tags & version floors drift; re-check each release]

You are… Do this
Running stock models via container (most people) Use a pinned digest of an image you have smoke-tested on Spark. v0.22.0 is the first stable baseline in this audit containing b12x (#40082) + the SM121 FP8 exclusion fix (#35568), a historical minimum, not the July 2026 recommendation; newer stable or nightly images may be preferable but must be re-audited for backend selection and cubin coverage (see the update timeline: the cu129 v0.25.1 stable still shipped no sm_121 cubins on 2026-07-19).
Tracking the newest kernels / want #40923 Marlin-MoE arch vllm/vllm-openai:v0.22.1rc0 or the blog’s compatibility track cu130-nightly (pin a digest; it’s CUDA-13).
Modifying C++/CUDA, or need a kernel no image ships Build from source with TORCH_CUDA_ARCH_LIST=12.1a, CUDA ≥ 12.8 (13.0 preferred).
On an old image (≤ v0.21) you can’t change Force Marlin (--moe-backend marlin --linear-backend marlin); expect dense+MoE to need it.

You almost never need a from-scratch native compile on the Spark itself unless you’re editing kernels; the published aarch64 images/wheels exist precisely so you don’t. From-source is the path for kernel work or when a needed fix is on main but not yet released.

10.2 If you build your own: the arch traps

  • List the right arch. CUDA 12.8/12.9 → 12.1a (or 12.0a;12.1a). CUDA 13.0+ → 12.0f (family target covers Spark). Below 12.8, sm_121 cannot be built at all (CMakeLists.txt:99-110).
  • Do not copy the in-repo Grace-Blackwell Docker example’s torch_cuda_arch_list='9.0 10.0+PTX'. That targets GH200/GB200, not GB10. Override it with 12.1a/12.0f.
  • Match the dependency pins to the revision you build. Current main (requirements/cuda.txt) pins flashinfer-python==0.6.11.post2, flashinfer-cubin==0.6.11.post2, and nvidia-cutlass-dsl[cu13]==4.5.2 [verified 2026-06-02: these pins move; read the pin in the exact tag you build]. (Older field notes citing cutlass-dsl 4.4.2 are stale.) Update [verified 2026-08-10]: this pin has since diverged by build path. The upstream main pin is not the whole story — Luke Alonso’s b12x kernel library (the one eugr’s experimental fork image builds against) now requires nvidia-cutlass-dsl==4.6.0, while 4.5.0 is known to emit PTX that ptxas rejects on sm_121/sm_120 (NVIDIA/cutlass#3227, fixed in 4.5.1/4.5.2). So “4.4.2 is stale” is now itself stale: read the cutlass-dsl pin in the exact tag and build path you deploy.
  • Include nvcc if you rely on FlashInfer JIT; otherwise FlashInfer disables itself when prebuilt cubins are missing for your arch (§4.2).

10.3 The non-negotiable: validate on the image you deploy

Because Spark correctness hinges on whether the specific image carries sm_121 cubins and which backend auto lands on, always smoke-test before trusting a config, watching for both the hard error and silent garbage:

  1. Serve one NVFP4 MoE model and one NVFP4-heavy dense model on the candidate image.
  2. Watch logs for Failed to run cutlass FP4 gemm on sm120/sm121 … Error Internal or No compiled nvfp4 mm kernel for SM 121.
  3. Inspect actual generated text for coherence (the silent-garbage failure mode produces no error).
  4. Try, in order: (a) auto, (b) explicit --moe-backend flashinfer_b12x --linear-backend flashinfer-b12x, (c) forced Marlin. Adopt the fastest that is coherent and stable.
  5. Fire a warm-up request (max_tokens=3, same model + chat template) at startup: the first request JIT-compiles Inductor/FlashInfer kernels (~25 s) and you don’t want a user to eat it.

10.4 Backend selection cheat-sheet for Spark

  • FP8 / BF16 model: auto is usually fine; FP8 uses native sm_121 tensor cores, so no Marlin dance in the common case. Treat it as image- and model-dependent, though: the v0.26.0 cu129 stable actually died on a CUTLASS FP8 W8A8 GEMM for our checkpoint (see the update timeline), so smoke-test the exact model and container before trusting auto.
  • NVFP4 model, want guaranteed-correct: --moe-backend marlin --linear-backend marlin.
  • NVFP4 model, want fastest native path: --moe-backend flashinfer_b12x --linear-backend flashinfer-b12x on v0.22.0+, then validate coherence.
  • Hybrid Mamba (nemotron_h): add --enforce-eager if your model/image combination shows the §8.3 graph-capture failure; it is a documented per-model workaround, not a Spark default.
  • Prefer CLI flags over the VLLM_* env vars: VLLM_NVFP4_GEMM_BACKEND was deprecated in v0.23 and removed in v0.25 (a no-op on current images), and VLLM_USE_FLASHINFER_MOE_FP4 is deprecated. Use --moe-backend / --linear-backend.

11. Appendix: glossary and file index

Glossary

  • GB10 / Grace Blackwell: the DGX Spark superchip (Grace ARM CPU + Blackwell GPU, unified memory).
  • sm_121 / CC 12.1: DGX Spark’s compute capability; workstation-Blackwell SM12x family, distinct from datacenter sm_100.
  • a vs f arch suffix: 12.1a = arch-specific cubin (one SM); 12.0f = family cubin (whole SM12x, needs CUDA 13).
  • CUTLASS: NVIDIA C++ template GEMM library; vLLM’s first-choice quantized matmul; needs arch-specific kernels.
  • FlashInfer: inference kernel library (attention, MoE, FP4/FP8 GEMM); ships prebuilt cubins or JITs at runtime.
  • CuteDSL: CUTLASS’s tensor-layout DSL; basis of the b12x SM12x kernels.
  • b12x: FlashInfer’s SM12x block-scaled fused-MoE / FP4-GEMM path; fast on Spark but opt-in (not auto, at the audited commit, §4.4).
  • Marlin / wna16: template-generated weight-only (W4/W8 ·A16) GEMM; broad arch coverage; Spark’s robust fallback for both dense and MoE.
  • QuTLASS: IST-DASLab CUTLASS extension providing FP4/MXFP4 kernels for sm_100/sm_120.
  • NVFP4: 4-bit float (e2m1), 16-element blocks, FP8-e4m3 block scales + FP32 global scale.
  • MXFP4: 4-bit float, 32-element blocks, e8m0 block scales; GPT-OSS format.
  • enforce-eager: disables CUDA-graph capture; needed for Mamba/hybrid on Spark; ~5–37% throughput cost.

File index (paths under the vLLM repo root)

Topic Path
SM12x family guard csrc/cutlass_extensions/common.hpp (enable_sm120_family)
CUTLASS FP8 dispatch (SM120) csrc/quantization/w8a8/cutlass/c3x/scaled_mm_sm120_fp8_dispatch.cuh
CUTLASS FP4 entry + SM dispatch csrc/quantization/fp4/nvfp4_scaled_mm_entry.cu
CUTLASS FP4 SM120 kernel csrc/quantization/fp4/nvfp4_scaled_mm_sm120_kernels.cu
Arch lists / family logic CMakeLists.txt (SCALED_MM_ARCHS, FP4_ARCHS, MARLIN_*_ARCHS), cmake/utils.cmake (cuda_archs_loose_intersection)
FlashInfer wrapper / JIT vllm/utils/flashinfer.py
b12x MoE expert vllm/model_executor/layers/fused_moe/experts/flashinfer_b12x_moe.py
b12x linear kernel vllm/model_executor/kernels/linear/nvfp4/flashinfer.py
Older FlashInfer CUTLASS MoE vllm/model_executor/layers/fused_moe/experts/flashinfer_cutlass_moe.py
NVFP4 MoE oracle (auto order, b12x exclusion) vllm/model_executor/layers/fused_moe/oracle/nvfp4.py
NVFP4 linear oracle vllm/model_executor/kernels/linear/__init__.py
MoE/linear backend options vllm/config/kernel.py
Modular fused-MoE vllm/model_executor/layers/fused_moe/modular_kernel.py
Marlin (dense) csrc/quantization/marlin/{marlin.cu,generate_kernels.py}
Marlin MoE csrc/moe/marlin_moe_wna16/
Marlin NVFP4 linear + scales vllm/model_executor/kernels/linear/nvfp4/marlin.py, vllm/.../quantization/utils/marlin_utils_fp4.py
NVFP4 / ModelOpt quant method vllm/model_executor/layers/quantization/modelopt.py
MXFP4 quant method vllm/model_executor/layers/quantization/mxfp4.py
QuTLASS build cmake/external_projects/qutlass.cmake
CUDA graph / eager config vllm/config/compilation.py, vllm/config/model.py
Mamba graph hooks vllm/model_executor/models/{mamba,mamba2,jamba,nemotron_h}.py
UMA / integrated-GPU detection vllm/utils/mem_utils.py, vllm/platforms/interface.py
Install / aarch64 / Docker docs/getting_started/installation/gpu.cuda.inc.md, docker/Dockerfile

Compiled from the vLLM main source tree (2026-06-02) and the official vLLM DGX Spark blog (2026-06-01). Backend selection details, especially the b12x opt-in, reflect current main; re-check the oracle comments in oracle/nvfp4.py against the exact image you deploy, since the upstream CUTLASS sm_121 MMA-guard fix may flip b12x into auto-selection in a future release.