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 (
glm47parser) and reasoning parser working; thinking togglable per-request
What breaks on GB10 and the fixes (all patches + probe scripts in the repo):
- GLM-5.3 is NoPE MLA (
qk_rope_head_dim=0) — the only stock sm_12x sparse-attention backend hard-requires DeepSeek’sfp8_ds_mlapacked layout (pe_dim=64 assert inconcat_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. - 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.
- …but that nightly silently downgrades
nvidia-nccl-cu13to 2.29.7 (NCCL “internal error” on the fabric) and skewsnvidia-cutlass-dsl(CuTeDSL compiler ICE). Re-pin both. Audit transitive deps after any pip install in these images. --block-size 2304is 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.- The kpool indexer’s top-k buffer is
torch.emptyand the kernels only fillmin(k, valid)entries — uninitialized pool ids → bogus token indices → attention gathers unwritten KV → intermittent NaN. Init to -1 + clamp. - vLLM enables PDL for capability ≥ 9 in the KDA recurrent-state kernels — unvalidated on sm_121; gated off.
- 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. - The GB10 classic:
NVRM: NV_ERR_NO_MEMORYfrom_memdescAllocInternalwhile 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 isMemAvailable(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/meminfotraces (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-memoryverbatim.
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:
- TP2: GitHub - tonyd2wild/GLM-5.3-Flash-NVFP4-262K-2x-DGX-Spark: GLM-5.3-Flash (NVFP4) on 2x NVIDIA DGX Spark - vLLM TP2, 262K context, MTP. World-first deploy recipe: 7 day-0 bugs found and fixed, patched sm121 image, probes and full report. · GitHub
- TP4: GitHub - tonyd2wild/GLM-5.3-Flash-NVFP4-1M-KV-4x-DGX-Spark: GLM-5.3-Flash (320B MoE) at TP4 across four DGX Sparks, same day as the model drop: 36 tok/s, 1.26M-token FP8 KV pool, 262K context, MTP spec decode. First TP4 glm5_next outside B200 hardware. Full patched-image recipe + the GB10 memory study. · GitHub
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.