The Moment of Repetition: Debugging a NaN Crash on Blackwell GPUs

The Message

Message 183 (Assistant):

>

Server crashed again. Let me check the log:

>

``bash ssh 10.1.230.175 'grep -i "error\|exception\|assert\|crash\|fatal\|sigquit" ~/sglang-glm5.log | tail -10' ``

>

`` CUDA kernel errors might be asynchronously reported at some other API call, so the stacktrace below might be incorrect. Compile with TORCH_USE_CUDA_DSA to enable device-side assertions. , scipy.spatial._distance_wrap, scipy.spatial.transform._rotation_cy, scipy.spatial.transform._rigid_transform_cy, scipy.optimize._direct, sklearn.__check_build._check_build, scipy.integrate._odepack, scipy.integrate._quadpack, scipy.integrate._vode, scipy.integrate._dop, scipy.interpolate._fitpack, scipy.inter... ``

Context: A Long Night of Debugging

To understand message 183, one must appreciate the debugging odyssey that precedes it. The conversation up to this point represents hours of iterative troubleshooting on a cutting-edge machine learning deployment: running the GLM-5-NVFP4 model — a 400B+ parameter Mixture-of-Experts (MoE) model quantized to FP4 — on eight NVIDIA RTX PRO 6000 Blackwell GPUs using the SGLang inference server. This is bleeding-edge infrastructure: Blackwell (SM120) architecture GPUs, a nightly build of SGLang from its main branch, and a model architecture (glm_moe_dsa) that requires Transformers 5.2.0, itself a very recent release with known compatibility issues.

The assistant had already cycled through multiple configurations. The first attempt used the recommended settings from the HuggingFace model card: --attention-backend flashinfer with --kv-cache-dtype fp8_e4m3. The server loaded the model, captured CUDA graphs, and appeared healthy — until the first real query triggered a device-side assert triggered error during decode. The assistant then tried switching to --attention-backend triton, which failed with a different assertion: the Triton backend does not support NSA (Native Sparse Attention) with FP8 KV cache, and GLM-5 uses DSA/NSA sparse attention. A third attempt used --nsa-decode-backend flashmla_sparse, which appeared to work: the server started, the warmup query completed, and the log showed "The server is fired up and..." ([msg 180]). But when the assistant sent a real chat completion request ([msg 181]), the response was empty. A verbose curl showed "Connection refused" ([msg 182]), confirming the server had crashed silently after the warmup.

The Message Itself

Message 183 is the assistant's reaction to this third failure. It is terse — almost resigned. "Server crashed again. Let me check the log:" — the word "again" carries the weight of the preceding hours. The assistant then runs a diagnostic grep command to extract error-related lines from the SGLang server log.

The output is revealing but incomplete. It shows the familiar CUDA kernel errors might be asynchronously reported warning, the suggestion to compile with TORCH_USE_CUDA_DSA for more precise diagnostics, and then — oddly — a fragment of a Python stack trace showing scipy and sklearn module imports. This fragment is a red herring: it's not the actual error but rather the tail end of a traceback that happened to contain those module names in its path resolution. The critical information — the actual assertion failure — was likely captured in earlier lines that the tail -10 truncated.

Why This Message Matters

At first glance, message 183 appears to be a simple status check, a routine diagnostic step. But it represents a pivotal moment in the debugging process. The assistant has now exhausted three different attention backend configurations, all producing the same fundamental failure mode: a CUDA device-side assert during decode. The pattern is clear: the server loads, captures graphs, serves the warmup query (which is a prefill-only operation), but crashes on the first real decode step.

The message is written because the assistant needs to confirm the failure mode is identical to previous attempts before formulating a new hypothesis. It is a data-gathering step, but it is also a moment of hypothesis consolidation. The assistant is implicitly asking: "Is this the same crash or a new one? Is the flashmla_sparse backend also broken on SM120, or is something else going on?"

Assumptions and Their Consequences

The assistant operates under several assumptions in this message. First, it assumes that grepping for error-related keywords will capture the root cause. This is a reasonable heuristic but has a blind spot: CUDA device-side asserts are asynchronous and may not produce clean Python tracebacks. The output confirms this — the grep captures the CUDA DSA warning but misses the actual assertion message, which was likely "probability tensor contains either inf, nan or element < 0" (as revealed in subsequent messages <msg id=185-186>).

Second, the assistant assumes the crash is still an attention-backend compatibility issue with SM120. This assumption is implicit in the message — the assistant doesn't yet question whether the root cause might be elsewhere. The subsequent messages show a pivot: after message 183, the assistant reads the full scheduler exception traceback ([msg 185]) and discovers the NaN/Inf probability tensor, then traces it to the DeepGemm scale format warning ([msg 186]). This pivot is the direct consequence of message 183's inconclusive output — the grep didn't provide a clear answer, forcing deeper investigation.

Input Knowledge Required

To understand message 183, the reader needs to know:

Output Knowledge Created

Message 183 produces a confirmation that the flashmla_sparse backend also crashes with the same device-side assert error. This negative result is valuable: it rules out flashmla_sparse as a solution and narrows the search space. The message also surfaces the truncated stack trace fragment, which — while not immediately useful — hints that the error path involves Python module loading, possibly during a fallback or error-handling codepath.

More importantly, message 183 sets the stage for the breakthrough in messages 185-187. The assistant, having confirmed that three different attention backends all produce the same crash, is forced to look beyond the attention backend hypothesis. The next message ([msg 184]) explicitly states: "Same device-side assert triggered even with flashmla_sparse. The issue isn't the attention backend choice." This realization is the critical pivot that leads to the DeepGemm scale format diagnosis.

The Thinking Process

The assistant's reasoning in message 183 is compressed but visible. The phrase "Server crashed again" acknowledges the pattern. The choice of grep pattern — error\|exception\|assert\|crash\|fatal\|sigquit — reflects a systematic approach to log analysis, covering multiple failure modes. The tail -10 limits output to the most recent occurrences, focusing on the current crash rather than historical ones.

The assistant does not, however, immediately analyze the output within this message. The message ends with the raw grep result. This is a deliberate choice: the assistant is collecting data before forming a conclusion. The analysis happens in the next round, when the assistant reads the full scheduler exception ([msg 185]) and connects the NaN probability tensor to the DeepGemm warning ([msg 186]).

A Broader Perspective

Message 183 exemplifies a common pattern in complex systems debugging: the moment when multiple hypotheses have been exhausted and the debugger must step back and reconsider fundamental assumptions. The assistant had been pursuing an attention-backend-centric hypothesis — trying different combinations of --attention-backend, --nsa-decode-backend, and --sampling-backend flags. Each attempt failed in the same way. Message 183 is the point where this hypothesis tree is fully explored and found barren.

The actual root cause — DeepGemm's automatic selection on Blackwell combined with an incompatible checkpoint scale format (ue8m0 vs. the checkpoint's actual format) — was a subtle interaction between the quantization format of the model weights and the GEMM backend's assumptions. The warning had been logged from the very first attempt ([msg 186] shows it appeared at 00:09:54, early in the server startup), but it was buried among other log messages and initially dismissed. The NaN/Inf probability tensor was a downstream consequence: DeepGemm was computing incorrect results due to the scale format mismatch, producing garbage logits that triggered the CUDA assert when the softmax normalization encountered non-finite values.

This is a cautionary tale about the dangers of "auto" configuration on new hardware. SGLang's --fp8-gemm-backend auto correctly detected Blackwell and selected DeepGemm, but the model checkpoint was not formatted for DeepGemm's expected scale representation. The fix — forcing --fp8-gemm-backend cutlass ([msg 187]) — bypassed the incompatible kernel path entirely.

Conclusion

Message 183 is a small but crucial node in the debugging graph. It does not contain the solution, but it is the message that forces the pivot. Without it — without the confirmation that yet another attention backend had failed — the assistant might have continued down the attention-backend rabbit hole, trying fa3, tilelang, aiter, or trtllm backends. Instead, the repeated failure pattern triggered a re-examination of the log warnings, leading to the DeepGemm insight.

In the broader narrative of this coding session, message 183 represents the transition from "which attention backend works on SM120?" to "what is fundamentally wrong with the computation on SM120?" — a shift from configuration debugging to numerical debugging. It is a reminder that sometimes the most productive debugging step is not finding the answer, but conclusively ruling out the wrong questions.