GLM-5.3-Flash on 2x NVIDIA DGX Spark 43.4 tok/s PEAK [Checkpoint]

Another Checkpoint and work in progress waiting on DFLASH2 feel free to submit issues and PRs

From Agent -

GLM-5.3-Flash serving on 2x DGX Spark — day-0, TP2, 262K context, FP8 KV + MTP — full recipe + the 8 GB10 bugs we fixed

Got GLM-5.3-Flash (LibertAIDAI NVFP4 quant, 182 GiB) serving across 2x GB10 the same day the model dropped. Sharing the full recipe plus everything that broke on sm_121, since the day-0 image (vllm/vllm-openai:glm53-flash-arm64-cu130) works on B200 but fails several distinct ways on Spark.

Numbers (TP2, vLLM, 200-token greedy streaming, 3-run medians):

  • 21.8 tok/s decode (peak 22.7) with the native MTP head, 4 draft tokens
  • TTFT 0.29 s
  • 262,144-token context, FP8 KV cache, 507K-token KV pool
  • Tool calling (glm47 parser) and reasoning parser working; thinking togglable per-request

What breaks on GB10 and the fixes (all patches + probe scripts in the repo):

  1. GLM-5.3 is NoPE MLA (qk_rope_head_dim=0) — the only stock sm_12x sparse-attention backend hard-requires DeepSeek’s fp8_ds_mla packed layout (pe_dim=64 assert in concat_and_cache_mla; same crash RTX PRO 6000 users hit). Fix: extend vLLM’s SM90 NoPE sparse-MLA backend to capability 12 on the FA2 path — kernel probed on-GPU with the model’s real geometry first.
  2. FlashInfer 0.6.17’s FA2 MLA kernel silently NaNs on 64–256-row batches on sm_121 — tiny and huge batches are clean, normal prompts land in the kill zone, output degenerates to one repeated token. Fixed in the 0.6.18 nightly.
  3. …but that nightly silently downgrades nvidia-nccl-cu13 to 2.29.7 (NCCL “internal error” on the fabric) and skews nvidia-cutlass-dsl (CuTeDSL compiler ICE). Re-pin both. Audit transitive deps after any pip install in these images.
  4. --block-size 2304 is mandatory: the hybrid block aligner picks a size whose kpool storage tiles by 32, but DeepGEMM’s arch-12 fp8 paged-MQA only accepts 64-entry pool pages.
  5. The kpool indexer’s top-k buffer is torch.empty and the kernels only fill min(k, valid) entries — uninitialized pool ids → bogus token indices → attention gathers unwritten KV → intermittent NaN. Init to -1 + clamp.
  6. vLLM enables PDL for capability ≥ 9 in the KDA recurrent-state kernels — unvalidated on sm_121; gated off.
  7. FP8 KV “requires SM90” is actually a two-line fix: the fa2 fp8 branch forces a CTA tile sized for Hopper’s 228 KB smem, over-requesting GB10’s ~101 KB opt-in max (cudaFuncSetAttribute → invalid argument). Cap the tile and fp8 KV runs clean (rel-err ~0.005 vs fp32). As far as we can tell, the first fp8 KV for a NoPE-MLA model on consumer Blackwell.
  8. The GB10 classic: NVRM: NV_ERR_NO_MEMORY from _memdescAllocInternal while tens of GB show “available.” The driver needs the KV slab as truly free pages and won’t reclaim page cache — which a 182 GiB shard read just filled. And vLLM’s “free memory” on integrated GPUs is MemAvailable (counts reclaimable cache), so it suggests KV sizes the driver can’t deliver. We mapped the exact ceiling with a six-boot controlled ladder + /proc/meminfo traces (docs/GB10-KV-MEMORY-LADDER.md), and ship the cache-flusher sidecar that runs during every load. Short version: take vLLM’s suggested --kv-cache-memory verbatim.

Bonus: TP4 across four Sparks dissolves the memory ceiling entirely (~50 GiB weights/rank) — 35.7 tok/s and a 1.26M-token FP8 KV pool. Separate repo.

Everything — Dockerfiles v1→v8, launchers, NCCL fabric env, probes, deploy report, upstream issue drafts:

Credits: zai-org (model), LibertAIDAI (quant — their sm_121 notes were spot-on), the vLLM PR #53906 authors (day-0 image), barrydeen (gmu reference). Happy to answer questions — receipts for all of it.

did you tried this quant as well? seems to be better quality