This is the English companion to the lab’s Russian long-read on Habr: getting a 35B-class MoE to serve a 256K context window on a single DGX Spark, and what we found broken along the way. That article is the primary source for every figure here — the numbers below are quotes from it, not lab registry claims.
Read the version pin first. Everything on this page describes one moment in a moving target: vLLM 0.15-era mainline with PyTorch 2.10, FlashInfer 0.6.8, driver 580.142, CUDA 13.0, DGX OS 7.5.0, measured in May 2026. Mainline moves weekly. What was broken then may be fixed by the time you read this, and what worked then may have regressed — treat this page as dated evidence with its pins attached, and check the linked upstream issues before repeating our conclusions.
The machine and the goal
One DGX Spark: GB10, Blackwell SM_121, 128 GiB of unified LPDDR5x at 273 GB/s . The goal was a genuinely long window — 262,144 tokens — for a Qwen3.6-35B-A3B-class MoE, on hardware whose compute architecture most open-source kernels had never been compiled for.
The configuration hunt
What refused to run, in the order we hit it:
- FlashInfer FP8 attention rejected the chip outright — the kernel supports sm120, and SM_121 got illegal synchronization errors instead of a fallback.
- CUTLASS/FA3 FP4 paths carry a heuristic that requires a minor compute-capability of zero, which excludes GB10 by construction.
- Triton fallback ran and stayed stable, at a cost of roughly 10–30% against FlashInfer .
- MoE backend flags needed hand-setting:
VLLM_USE_FLASHINFER_MOE_FP4=0because the autotuner skipped ten of seventeen tactics as unsupported , whileVLLM_USE_FLASHINFER_MOE_FP8=1worked on SM_121 with FlashInfer 0.6.x. - Full CUDA-graph capture hung on vLLM 0.15-era builds;
cudagraph_mode=piecewiseis the workaround. - Driver 580.142 is the pin. The 590.48.01 driver leaked unified memory
until 80 of 128 GB were consumed with nothing visible in AnonPages or Slab ,
and 595.58.03 broke NVFP4 with a
cuTensorMapEncodeTiledillegal memory access. If your Grafana looks fine while this happens, see why GB10 monitoring lies.
Why NVFP4 was broken in the mainline build we tested
The pain query that probably brought you here. In the mainline build we
pinned, NVFP4 kernels were compiled for compute_120f, while the native
NVFP4 instructions exist only in compute_120a and compute_121a. On
SM_121 that meant quantized weights were unpacked in software, bit
manipulation in the shader with the tensor cores idle. The log fingerprint
was [AutoTuner]: Skipped 10 unsupported tactic(s) for trtllm::fused_moe::gemm2.
The measured consequence: the quantization with half the memory footprint decoded slower than FP8 — 40.9 against 51.2 tok/s single-stream at the 260K configuration . Half the bytes, more time: the unpacking overhead ate the bandwidth win.
To be precise about tense: NVFP4 was broken in that build, on that chip, at that date. This is a compilation-target gap, not a law of the hardware — the upstream tracker was flashinfer issue #3170, the “DGX Spark SM_121 Audit”, with seventeen items open when the article shipped . Verify against current mainline before deciding anything.
What worked
The configuration that held 260K and went fastest was nothing like stock:
vLLM built from source with seven local patches, FlashInfer 0.6.8 plus PRs
#2520 and #2702, the community AEON-7/Qwen3.6-35B-A3B-heretic-NVFP4
checkpoint, and the DFlash block-diffusion drafter at fifteen speculative
tokens:
--max-model-len 262144
--max-num-seqs 128
--max-num-batched-tokens 65536
--gpu-memory-utilization 0.85
--enable-chunked-prefill --enable-prefix-caching
--speculative-config '{"method":"dflash","num_speculative_tokens":15}'
--attention-backend flash_attn
The trap in the fast path: the heretic fine-tune breaks tool calling (its
issue #4, unfixed at the time). For agent and MCP serving the article’s
recommendation was vanilla Qwen/Qwen3.6-35B-A3B-FP8 on stock vLLM —
slower, but function calls come back intact.
Measured behavior at 256K
Quoted from the article’s single-box runs :
| Configuration | c=1 decode | ×32 aggregate | Weights in memory | Context |
|---|---|---|---|---|
| Gemma 4 26B-A4B BF16 | 23.3 tok/s | 297.8 tok/s | 49 GiB | 252K |
| Qwen3.6 FP8 | 51.2 tok/s | 349.2 tok/s | 35 GiB | 260K |
| NVFP4, mainline | 40.9 tok/s | 448 tok/s | 22 GiB | 260K |
| NVFP4 + MTP | 54.1 tok/s | 498.6 tok/s | 22 GiB | 131K |
| AEON-7 + DFlash | 69.7 tok/s | 350.1 tok/s | 22 GiB | 260K |
Warm time-to-first-token was around 100 ms including drafter compilation .
The most useful surprise is how cheap the window itself was: growing
max-model-len from 65K to 260K cost about two and a half percent of single-stream decode ,
while a single user’s active context rarely exceeded 30K of the 260K window .
The speculative numbers swing hard by workload — DFlash acceptance was 17.8% overall but 78% at position zero on math and code ,
which spread decode between 117 tok/s on math and 45–57 tok/s on business chat .
And the ceiling is memory, not marketing: at 273 GB/s the article’s roofline puts effective speculative decode near 320 tok/s ,
regardless of what the FP4 petaflop figure on the box implies.
Honesty section
- Everything here is
lab_single_run: one lab, one box, single passes. None of these figures are registry claims — they are quotes from our own article, carrying its caveats. - The version pin at the top is the whole point. A “broken” verdict without its build fingerprint is noise; with it, it is a reproducible bug report.
- The two-node serving recipe with its own traps is a separate page, and why published Spark numbers disagree with each other is its own investigation.