cuTensorMapEncodeTiled illegal memory access

What you see

Two fingerprints from the same NVFP4 path on GB10. On driver 595.58.03 NVFP4 broke with a cuTensorMapEncodeTiled illegal memory access. The report records it in those words, not as a captured log line; the exact runtime message is not on the page.

On the pinned driver NVFP4 loaded and served, but the log carried this line and decode came out slower than FP8 at the same context. Half the bytes, more time:

[AutoTuner]: Skipped 10 unsupported tactic(s) for trtllm::fused_moe::gemm2

Where we saw it

• Hardware: one DGX Spark, GB10, Blackwell SM_121.
• Runtime: vLLM 0.15-era mainline with PyTorch 2.10, FlashInfer 0.6.8, CUDA 13.0, DGX OS 7.5.0. Driver 580.142 is the pin; 595.58.03 is where the illegal memory access came from.
• Model: a Qwen3.6-35B-A3B-class MoE in NVFP4, at the report’s long-context configuration.
• When: measured in May 2026; the report is dated 2026-08-21.

Cause

The report does not root-cause the illegal memory access; it correlates it with the driver version. It records what each newer driver did: 590.48.01 leaked unified memory with nothing visible in AnonPages or Slab, 595.58.03 broke NVFP4 with the cuTensorMapEncodeTiled error. Grafana can look fine while it happens.

The slow path is a compilation-target gap. In the mainline build we pinned, the NVFP4 kernels were compiled for compute_120f, while the native NVFP4 instructions exist only in compute_120a and compute_121a. On SM_121 the quantized weights were unpacked in software, bit manipulation in the shader with the tensor cores idle, and the autotuner skipped most of the fused-MoE GEMM tactics as unsupported.

Separately, the report lists the CUTLASS/FA3 FP4 paths as refusing the chip outright: they carry a heuristic that requires a minor compute-capability of zero, which excludes GB10 by construction.

Fix

Driver: pin 580.142. That is the whole fix on our box; neither newer driver the report tried survived: 590.48.01 leaked unified memory, 595.58.03 broke NVFP4.

Kernel gap: the report documents a workaround, not a fix. Hand-set the MoE backend flags:

VLLM_USE_FLASHINFER_MOE_FP4=0
VLLM_USE_FLASHINFER_MOE_FP8=1

FP4 off because the autotuner skipped its tactics; FP8 worked on SM_121 with FlashInfer 0.6.x.

The configuration that held the long window with an NVFP4 checkpoint was vLLM built from source with seven local patches, FlashInfer 0.6.8 plus PRs #2520 and #2702, the AEON-7/Qwen3.6-35B-A3B-heretic-NVFP4 checkpoint and the DFlash drafter; launch flags are in the report. Nothing there says those patches recompile the kernels for compute_121a, so do not read the fast path as a fix. The fast path has its own trap: the heretic fine-tune breaks tool calling. For agent and MCP serving, the article the report quotes recommended vanilla Qwen/Qwen3.6-35B-A3B-FP8 on stock vLLM.

Still open upstream?

No ticket is named for the driver error; it is recorded as a pin. flashinfer issue #3170, the DGX Spark SM_121 support audit, tracks the kernel gap behind the AutoTuner line, with items still open when the article shipped; verify against current mainline first.

Evidence

• Source page: vLLM at 256K context on DGX Spark, evidence level lab_single_run.
• The article it quotes: the lab’s long-read on Habr.
• Upstream tracker named by the page, for the kernel gap only: flashinfer issue #3170.
• Why the box can look fine meanwhile: why GB10 monitoring lies.