The Turning Point: How Switching to flashmla_sparse Brought GLM-5-NVFP4 Past CUDA Graph Capture on Blackwell GPUs

In the high-stakes world of deploying cutting-edge large language models on novel hardware, success often hinges on a single configuration flag. Message [msg 179] captures precisely such a moment — a quiet but decisive turning point in a marathon debugging session. After hours of iterative troubleshooting involving OOM crashes, device-side assert failures, and incompatible attention backends, the assistant finally sees the GLM-5-NVFP4 server survive its CUDA graph capture phase on 8 RTX PRO 6000 Blackwell GPUs. The message itself is deceptively simple — a bash command that sleeps for 180 seconds and then tails a log file — but the output it returns tells a story of systematic debugging, architectural knowledge, and the razor-thin margins of compatibility that define state-of-the-art ML infrastructure.

The Message

The assistant writes:

[bash] sleep 180 && ssh 10.1.230.175 'tail -15 ~/sglang-glm5.log'
[2026-02-19 00:12:45 TP3] Capture cuda graph end. Time elapsed: 7.70 s. mem usage=1.62 GB. avail mem=6.60 GB.
[2026-02-19 00:12:45 TP5] Capture cuda graph end. Time elapsed: 7.74 s. mem usage=1.62 GB. avail mem=6.60 GB.
[2026-02-19 00:12:45 TP6] Capture cuda graph end. Time elapsed: 7.74 s. mem usage=1.62 GB. avail mem=6.60 GB.
[2026-02-19 00:12:45 TP4] Capture cuda graph end. Time elapsed: 7.77 s. mem usage=1.62 GB. avail mem=6.60 GB.
[2026-02-19 00:12:45 TP2] Capture cuda graph end. Time elaps...

On its surface, this is a routine status check. But the log lines it reveals — "Capture cuda graph end" appearing across multiple tensor parallelism ranks (TP2 through TP6) — represent a breakthrough. The server has successfully completed CUDA graph capture, the final and most fragile step of initialization, without crashing. This is the first time in the session that the server has gotten this far without a fatal error.

The Debugging Journey That Led Here

To understand why this message matters, we must trace the path that led to it. The assistant had been attempting to deploy the GLM-5-NVFP4 model — a massive 400-billion-parameter Mixture-of-Experts architecture quantized to FP4 — on a server equipped with 8 NVIDIA RTX PRO 6000 Blackwell GPUs, each with 96 GB of VRAM. This is bleeding-edge hardware (SM120 architecture) running a bleeding-edge model with a bleeding-edge serving framework (sglang built from the main branch).

The journey was fraught. The first attempt crashed with an OOM error because --mem-fraction-static 0.95 was too aggressive, leaving no room for CUDA graph buffers ([msg 150]). After lowering it to 0.88, the server loaded successfully and captured CUDA graphs, but then crashed during the very first decode operation with a cryptic device-side assert triggered error ([msg 164]). This was not a memory issue — it was a kernel crash, the kind that signals fundamental incompatibility between the CUDA kernels being invoked and the GPU architecture.

The assistant traced the crash to the decode phase of inference, specifically in the attention mechanism. The GLM-5 model uses a "Dynamic Sparse Attention" (DSA) architecture, which in sglang is handled by the NSA (Native Sparse Attention) backend system. The server logs revealed that sglang had auto-selected nsa_decode_backend='flashmla_kv' for the decode attention path ([msg 167]). This backend, it turned out, was crashing on SM120 silicon.

The Decision-Making Process

The assistant's response to this crash demonstrates a methodical approach to debugging kernel-level compatibility issues. Rather than guessing randomly, the assistant first enumerated all available NSA decode backends by querying sglang's help text ([msg 177]). The options were: flashmla_sparse, flashmla_kv, flashmla_auto, fa3, tilelang, aiter, and trtllm.

Each candidate was evaluated against known constraints:

Why This Message Is a Turning Point

Message [msg 179] is the verification step for that hypothesis. The assistant waits 180 seconds — a generous buffer to account for model loading (83 shards of safetensors), KV cache allocation, and CUDA graph capture — and then checks the log. The output shows that CUDA graph capture completed successfully across all tensor parallelism ranks. Each rank reports approximately 7.7 seconds of capture time, 1.62 GB of memory used for the CUDA graphs, and 6.60 GB of available memory remaining.

Several details in this output are worth highlighting:

The memory numbers tell a coherent story. With --mem-fraction-static 0.88, each GPU has approximately 84.5 GB allocated for model weights, KV cache, and runtime buffers. The CUDA graphs consume 1.62 GB, leaving 6.60 GB available — a healthy margin that prevents the OOM crashes seen in earlier attempts. This confirms that the memory configuration is stable.

The capture times are consistent across ranks. TP3 through TP6 all show capture times between 7.70 and 7.77 seconds, with identical memory usage. This uniformity suggests that the CUDA graph capture is well-balanced across the 8 GPUs, with no stragglers or imbalanced memory pressure.

The server did not crash. This is the most significant signal. Every previous attempt that reached the decode phase ended with a device-side assert or assertion error. The fact that CUDA graph capture — which compiles and caches the CUDA kernels for all supported batch sizes — completed without error strongly suggests that the flashmla_sparse backend is SM120-compatible where flashmla_kv was not.

Assumptions and Their Validation

The assistant made several assumptions in reaching this point. The first was that the crash was specifically in the NSA decode backend, not in prefill or elsewhere. This was validated by the log output showing successful prefill completion before the decode crash in earlier attempts ([msg 166]). The second assumption was that flashmla_sparse would be compatible with Blackwell GPUs — this was an educated guess based on the naming convention and the fact that it wasn't ruled out by architectural constraints like FA3 was. The third assumption was that 180 seconds would be sufficient for the full initialization sequence, which turned out to be generous — CUDA graph capture completed in under 8 seconds after model loading.

Input Knowledge Required

Understanding this message requires familiarity with several technical domains. One must know that CUDA graph capture is the process of recording and optimizing CUDA kernel launches for repeated execution, and that it is a critical initialization step for high-performance inference serving. One must understand tensor parallelism (TP) and why log lines from TP2 through TP6 (but not all 8 ranks) appear — the log may be truncated, or some ranks may have completed slightly earlier or later. One must know what NSA backends are and why the DSA architecture of GLM-5 requires them. And one must appreciate the significance of SM120 as a new GPU architecture (Blackwell) where kernel compatibility is not guaranteed.

Output Knowledge Created

This message creates several pieces of actionable knowledge. It confirms that --nsa-decode-backend flashmla_sparse is a viable configuration for GLM-5-NVFP4 on Blackwell GPUs, at least through the initialization phase. It provides empirical data on CUDA graph capture memory usage (1.62 GB) and available headroom (6.60 GB) at --mem-fraction-static 0.88. It establishes that the flashinfer attention backend combined with flashmla_sparse for NSA decode does not trigger the device-side assert that plagued the flashmla_kv backend. And it sets the stage for the next critical test: whether the server can actually complete a full decode pass and return a response to a user query.

The Broader Significance

This message exemplifies a pattern that recurs throughout ML infrastructure work: the difference between success and failure is often a single flag, discovered through systematic elimination. The assistant did not randomly change parameters — it read error messages, consulted documentation, enumerated options, reasoned about architectural constraints, and made targeted changes. The flashmla_sparse backend was not an obvious choice; it was the survivor of a淘汰 process that ruled out flashmla_kv (crashes), triton (incompatible with NSA+FP8), and fa3 (wrong architecture).

The subsequent messages confirm the breakthrough. In [msg 180], the server is up and the warmup completes without crashing. In [msg 181], the assistant sends a real query and gets a response. The turning point, however, is here in [msg 179] — the moment when the assistant sees those "Capture cuda graph end" lines and knows that the configuration is finally stable enough to proceed. It is a small victory in a long battle, but in the world of deploying 400-billion-parameter models on brand-new hardware, small victories are how progress is made.