Thanks for the recipe — it runs. Two problems showed up in daily use on a single Spark (GB10, 121.6 GiB unified memory, TP1, general23 image), and both have fixes that don't need a rebuild. Full measurements are linked at the end.
1. --gpu-memory-utilization 0.90 leaves the host ~6 GB and the driver runs out of page-table memory
On GB10 the GPU pool is the OS RAM. 0.90 reserves 109.5 GiB for vLLM; the host (OS, docker, the vLLM processes' own RSS ~6 GB) needs ~10.5 GiB, so MemAvailable after init is ~6 GB. The first real request then allocates ~4.4 GB on the GPU outside the profiler's budget (JIT of mhc_pre_big_fuse_with_norm_tilelang and _prepare_dflash_inputs_kernel during inference, allocator growth), and the first prompt above 7168 tokens another ~2.5 GB. At that point NVRM logs NV_ERR_NO_MEMORY from _memdescAllocInternal and you get either Xid 31 FAULT_PDE or an engine that hangs while the host thrashes (SSH dead, power button). We hit both. Note that --memory=115g on the container does nothing here: GPU allocations are not charged to the cgroup, and earlyoom with its default swap condition never fires.
Measured: 0.90 → 6.0 GB free after init → first request pushes it to 2.2 GB with 4× NV_ERR_NO_MEMORY. 0.85/64k → 14.2 GB after init, plateau ~7.5 GB, zero errors.
2. Why 262k needs 16 GiB of KV in this recipe, and how to get it down to 3.9 GiB
The KV pool is one global pool of 7168-token blocks (MLA page must hold the KDA state) and every block id costs 51.8 MiB regardless of which group owns it. For a single 262k request the need is ceil(L/7168) + 4·(2+K) + 1 + drafter blocks. The drafter (DFlash2, sliding window 2048, block 64 padded to the MLA page) reserves ceil((2047 + 2·max_num_batched_tokens)/64)+1 = 257 blocks with async scheduling on — 13.0 GiB for 0.31 GiB of draft KV, 80% of the 16.34 GiB total. Context itself is 37 blocks (1.9 GiB).
Three changes, all verified on this box (decode and prefill unchanged, acceptance 3.3–3.6):
| change |
effect on KV needed for 262k |
--no-async-scheduling (drafter window 2047+1·MNBT → 145 blocks; bs=1 decode 25–27 tok/s before and after) |
16.34 → 10.68 GiB |
drafter compact block 64 → 1024 in the padded slot-share path (kv_cache_utils.py: compact_block = 1024 and the != 64 validation). FlashAttention reports MultipleOf(16), select_common_block_size returns the manager block, so kernel block == manager block and the strided view stays valid |
10.68 → 3.85 GiB |
GLM53_INDEXER_WORKSPACE=rightsize (your patch, off by default in start.sh) |
consumed −1.28 GiB |
Result at --gpu-memory-utilization 0.80, 262144 context: KV 10.1 GiB (2.6× concurrency), 20.5 GB host free after init, 13.6 GB minimum during a 196k-token cold prefill (~745 tok/s), decode 27 tok/s. 524288 also boots (1.7×), 100k/250k multi-needle 5/5.
3. Prefix-cache hits crash the engine (and below the threshold, silently corrupt the KDA state)
Any request whose cached prefix is ≥ 8 full blocks (57 344 tokens) dies immediately with an illegal memory access in the first KDA layer (kda.py chunk_kda_with_fused_gate → prepare_chunk_indices), Xid 31 in dmesg, at 14 GB host free. 1/4/6-block hits "work" but restore the wrong recurrent state.
Root cause is upstream vLLM: gpu/model_states/mamba_hybrid.py add_request seeds the mamba state slot with (num_computed_tokens-1) // cache_config.block_size, and engine/core.py has by then lowered cache_config.block_size to the smallest prefix-cacheable group — the drafter's 64 (or 1024) — while the mamba blocks and the block-table row width are in units of 7168. _copy_mamba_state_block then reads block_table[src_col] with src_col up to 111 in a 42-wide row. Your README's warm numbers were probably measured with --max-num-seqs 4 (the read lands in another row: no fault, wrong state) — and this may be what issue #1 ("endless loops after a while under load", no errors in logs) looks like from the outside.
One-line fix, applied with sed at container start: // self.cache_config.mamba_block_size. After it: 8/13/34/63-block hits complete, warm TTFT at 100k is 8.7 s (was 139 s cold), needle answers on warm requests match cold (5/5 at 100k and 250k). I'm filing this upstream as well: vllm-project/vllm#55600.
Env knobs I added to scripts/serve-one-spark.sh / start.sh (happy to send a PR)
ONE_SPARK_UTIL, ONE_SPARK_CTX, ONE_SPARK_SEQS (these existed but were never passed into the container by start.sh), ONE_SPARK_ASYNC=0|1, ONE_SPARK_MNBT, ONE_SPARK_DRAFT_BLOCK=1024, ONE_SPARK_APC=0|1, ONE_SPARK_MAMBA_SEED_FIX=1, and GLM53_INDEXER_WORKSPACE passthrough.
4. What the combined configuration survives (numbers)
- 3-hour soak at ctx 524288, util 0.80, prefix caching + seed fix, watchdog on: 133 requests (57 multi-turn conversation turns with prefix hits up to 184k tokens, 22 fresh prompts of 41k–354k tokens, short/JSON/3000-token generations), 133 OK, zero Xid, MemAvailable min 13.3 GB with no drift; warm TTFT median 13.8 s.
- Cold prefill stays linear to 443k tokens (~700–745 tok/s); 5/5 multi-needle recall at 100k, 250k, 350k and 450k.
- One caveat for your watchdog logic if you have any: vLLM logs no
Avg prompt throughput lines during a long prefill (only after), so a liveness check keyed on them kills a healthy engine mid-prefill; the per-forward exl3 fat-expert P0 lines are a usable liveness signal.
5. Quantization quality, measured against the full model (same harness, same items)
100 GSM8K + full HumanEval (164), temperature 0, max_tokens 4096, reference = z-ai/glm-5.3-flash via OpenRouter:
|
GSM8K |
HumanEval |
HumanEval excl. finish=length |
cut at 4096 |
| full model, thinking |
98% |
95.1% |
98.7% |
6 |
| EXL3 2.05 bpw, thinking |
97% |
83.5% |
97.9% |
24 |
| EXL3 2.05 bpw, non-thinking |
96% |
93.3% |
93.3% |
0 |
On tasks it finishes the quant is as accurate as the full model; the loss is that ~15% of coding tasks send its reasoning into non-convergent deliberation (still length at 12 000 tokens where the full model finishes in 0.8–2.7k). Non-thinking mode recovers almost everything at ~6× lower latency. This may be part of what issue #1 ("endless loops after a while under load") looks like, on top of the prefix-hit state corruption above.
Upstream vLLM issue for the seed bug: vllm-project/vllm#55600. Measurement scripts and logs available on request.
Thanks for the recipe — it runs. Two problems showed up in daily use on a single Spark (GB10, 121.6 GiB unified memory, TP1,
general23image), and both have fixes that don't need a rebuild. Full measurements are linked at the end.1.
--gpu-memory-utilization 0.90leaves the host ~6 GB and the driver runs out of page-table memoryOn GB10 the GPU pool is the OS RAM. 0.90 reserves 109.5 GiB for vLLM; the host (OS, docker, the vLLM processes' own RSS ~6 GB) needs ~10.5 GiB, so MemAvailable after init is ~6 GB. The first real request then allocates ~4.4 GB on the GPU outside the profiler's budget (JIT of
mhc_pre_big_fuse_with_norm_tilelangand_prepare_dflash_inputs_kernelduring inference, allocator growth), and the first prompt above 7168 tokens another ~2.5 GB. At that point NVRM logsNV_ERR_NO_MEMORYfrom_memdescAllocInternaland you get eitherXid 31 FAULT_PDEor an engine that hangs while the host thrashes (SSH dead, power button). We hit both. Note that--memory=115gon the container does nothing here: GPU allocations are not charged to the cgroup, and earlyoom with its default swap condition never fires.Measured: 0.90 → 6.0 GB free after init → first request pushes it to 2.2 GB with 4×
NV_ERR_NO_MEMORY. 0.85/64k → 14.2 GB after init, plateau ~7.5 GB, zero errors.2. Why 262k needs 16 GiB of KV in this recipe, and how to get it down to 3.9 GiB
The KV pool is one global pool of 7168-token blocks (MLA page must hold the KDA state) and every block id costs 51.8 MiB regardless of which group owns it. For a single 262k request the need is
ceil(L/7168) + 4·(2+K) + 1 + drafterblocks. The drafter (DFlash2, sliding window 2048, block 64 padded to the MLA page) reservesceil((2047 + 2·max_num_batched_tokens)/64)+1 = 257blocks with async scheduling on — 13.0 GiB for 0.31 GiB of draft KV, 80% of the 16.34 GiB total. Context itself is 37 blocks (1.9 GiB).Three changes, all verified on this box (decode and prefill unchanged, acceptance 3.3–3.6):
--no-async-scheduling(drafter window 2047+1·MNBT → 145 blocks; bs=1 decode 25–27 tok/s before and after)kv_cache_utils.py:compact_block = 1024and the!= 64validation). FlashAttention reportsMultipleOf(16),select_common_block_sizereturns the manager block, so kernel block == manager block and the strided view stays validGLM53_INDEXER_WORKSPACE=rightsize(your patch, off by default instart.sh)Result at
--gpu-memory-utilization 0.80, 262144 context: KV 10.1 GiB (2.6× concurrency), 20.5 GB host free after init, 13.6 GB minimum during a 196k-token cold prefill (~745 tok/s), decode 27 tok/s. 524288 also boots (1.7×), 100k/250k multi-needle 5/5.3. Prefix-cache hits crash the engine (and below the threshold, silently corrupt the KDA state)
Any request whose cached prefix is ≥ 8 full blocks (57 344 tokens) dies immediately with an illegal memory access in the first KDA layer (
kda.py chunk_kda_with_fused_gate→prepare_chunk_indices), Xid 31 in dmesg, at 14 GB host free. 1/4/6-block hits "work" but restore the wrong recurrent state.Root cause is upstream vLLM:
gpu/model_states/mamba_hybrid.py add_requestseeds the mamba state slot with(num_computed_tokens-1) // cache_config.block_size, andengine/core.pyhas by then loweredcache_config.block_sizeto the smallest prefix-cacheable group — the drafter's 64 (or 1024) — while the mamba blocks and the block-table row width are in units of 7168._copy_mamba_state_blockthen readsblock_table[src_col]withsrc_colup to 111 in a 42-wide row. Your README's warm numbers were probably measured with--max-num-seqs 4(the read lands in another row: no fault, wrong state) — and this may be what issue #1 ("endless loops after a while under load", no errors in logs) looks like from the outside.One-line fix, applied with
sedat container start:// self.cache_config.mamba_block_size. After it: 8/13/34/63-block hits complete, warm TTFT at 100k is 8.7 s (was 139 s cold), needle answers on warm requests match cold (5/5 at 100k and 250k). I'm filing this upstream as well: vllm-project/vllm#55600.Env knobs I added to
scripts/serve-one-spark.sh/start.sh(happy to send a PR)ONE_SPARK_UTIL,ONE_SPARK_CTX,ONE_SPARK_SEQS(these existed but were never passed into the container bystart.sh),ONE_SPARK_ASYNC=0|1,ONE_SPARK_MNBT,ONE_SPARK_DRAFT_BLOCK=1024,ONE_SPARK_APC=0|1,ONE_SPARK_MAMBA_SEED_FIX=1, andGLM53_INDEXER_WORKSPACEpassthrough.4. What the combined configuration survives (numbers)
Avg prompt throughputlines during a long prefill (only after), so a liveness check keyed on them kills a healthy engine mid-prefill; the per-forwardexl3 fat-expert P0lines are a usable liveness signal.5. Quantization quality, measured against the full model (same harness, same items)
100 GSM8K + full HumanEval (164), temperature 0, max_tokens 4096, reference =
z-ai/glm-5.3-flashvia OpenRouter:finish=lengthOn tasks it finishes the quant is as accurate as the full model; the loss is that ~15% of coding tasks send its reasoning into non-convergent deliberation (still
lengthat 12 000 tokens where the full model finishes in 0.8–2.7k). Non-thinking mode recovers almost everything at ~6× lower latency. This may be part of what issue #1 ("endless loops after a while under load") looks like, on top of the prefix-hit state corruption above.Upstream vLLM issue for the seed bug: vllm-project/vllm#55600. Measurement scripts and logs available on request.