Diagnosing NaN Outputs in FP4 Inference: The Critical Backend Configuration for Qwen3.5-397B on Blackwell GPUs
In the high-stakes world of deploying large language models on cutting-edge hardware, the difference between a working system and a broken one often comes down to a handful of command-line flags. Message [msg 5854] captures exactly such a moment — a crisp diagnostic pivot where an assistant, armed with an external reference from a GitHub gist, identifies the root cause of numerical instability in a 397-billion-parameter MoE model running on NVIDIA Blackwell (SM120) GPUs and applies the fix with surgical precision.
The Context: A Fresh Deployment Gone Wrong
The narrative leading to this message is one of ambitious deployment. The assistant had just transitioned from a hardened production setup of the Kimi-K2.5 INT4 model to deploying a newer, more efficient architecture: nvidia/Qwen3.5-397B-A17B-NVFP4. This model is a 397-billion-parameter Mixture-of-Experts (MoE) transformer with 512 experts (10 active per token), 60 layers, and a massive 262,144-token context window, all compressed using NVIDIA's FP4 quantization (modelopt_fp4). The model had been downloaded (223 GB across 6 safetensor shards), the latest SGLang main branch had been built from source with CUDA 13 and SM120 patches applied, and a systemd service (sglang-qwen.service) had been created.
The initial launch attempt failed with an assertion error — the hybrid GDN architecture (mixing linear attention layers with full attention) required the triton or trtllm_mha attention backend on Blackwell, not flashinfer. The assistant fixed this by switching to --attention-backend triton and the server came up in about two minutes.
Then the real problem surfaced. The first test query — a simple request to write a prime-checking Python function — produced garbage: the model output nothing but repeated exclamation marks (!!!!!!!!!!!!...), captured as reasoning_content by the Qwen3 reasoning parser. A second test with higher temperature confirmed the worst: the response finished with "matched_stop": "NaN happened". The model was producing numerical garbage — not just wrong answers, but mathematically invalid floating-point computations that collapsed to NaN.
The Investigation: Tracing the NaN to Its Source
The assistant's differential diagnosis was methodical. The NaN output pointed squarely at the FP4 quantization kernels — the specialized CUDA kernels that perform matrix multiplications using 4-bit floating-point arithmetic on Blackwell's tensor cores. The FP4 GEMM (General Matrix Multiply) runner is the component responsible for executing these quantized computations, and the MoE runner handles the expert routing and sparse computation for the 512-expert mixture-of-experts layers.
The assistant had previously inspected the available FP4 GEMM backend choices in SGLang's server arguments (message [msg 5847]), finding four options: auto, flashinfer_cudnn, flashinfer_cutlass, and flashinfer_trtllm. The default was flashinfer_cutlass. The question was whether this default backend was compatible with Blackwell's SM120 architecture.
The user then provided a crucial external reference — a GitHub gist at https://gist.github.com/catid/87cca824963f17fe7479a0ed26221397 titled "Qwen 3.5 NVFP4 Setup SM120" (messages [msg 5850] and [msg 5852]). The assistant fetched this gist in message [msg 5853], though the raw fetch returned mostly GitHub page HTML rather than the gist's rendered content.
The Subject Message: Extracting the Signal from the Noise
Message [msg 5854] is where the assistant synthesizes the gist's actionable content and executes the fix. The message contains two distinct parts: a reasoning/analysis section and a bash command execution.
The reasoning section reads:
The key fix from the gist — the!spam / NaN is caused by the wrong MoE and FP4 GEMM backend. Need to pin: ---moe-runner-backend flashinfer_cutlass---fp4-gemm-runner-backend flashinfer_cudnnAlso notes we need to buildsgl-kernelfrom source with SM120 support. But let me first try just the backend flags since we already have sgl-kernel 0.3.21:
This is a masterclass in pragmatic debugging. The assistant identifies three pieces of information from the gist:
- The MoE runner backend must be explicitly set to
flashinfer_cutlass - The FP4 GEMM runner backend must be explicitly set to
flashinfer_cudnn(note: not the defaultflashinfer_cutlass) - The
sgl-kernelpackage may need to be rebuilt from source with SM120 support enabled But critically, the assistant makes a strategic decision: try the simpler fix first. Sincesgl-kernelversion 0.3.21 is already installed, there's a chance it already includes SM120 support. The backend flags alone might suffice. This is the classic debugging heuristic of "test the simplest hypothesis first" — changing two command-line flags is much faster and less risky than rebuilding a kernel package from source. The bash command that follows stops the running service, kills any processes holding NVIDIA GPU file handles (fuser -k /dev/nvidia*), and waits for the GPUs to be released. This is standard procedure when reconfiguring a GPU-bound server — the old process must be fully terminated before the new one can claim the devices.
The Reasoning Process: What the Assistant Understood
To appreciate the intelligence in this message, one must understand the technical landscape the assistant was navigating. The FP4 quantization format is NVIDIA's 4-bit floating-point format, part of the Blackwell architecture's native support for low-precision computation. The modelopt_fp4 quantization scheme uses NVIDIA's ModelOpt toolkit to compress weights to this format. However, executing FP4 matrix multiplications requires specialized GPU kernels that are architecture-specific.
The Blackwell SM120 architecture introduced new tensor core capabilities for FP4, but these require matching kernel implementations. The flashinfer_cutlass backend uses CUTLASS (CUDA Templates for Linear Algebra Subroutines) to generate FP4 kernels, while flashinfer_cudnn uses cuDNN (CUDA Deep Neural Network library). The key insight from the gist is that on SM120, the FP4 GEMM runner needs cuDNN (flashinfer_cudnn), not CUTLASS (flashinfer_cutlass), while the MoE runner conversely needs CUTLASS. This asymmetric backend configuration is a subtle but critical detail that would be easy to get wrong.
The assistant also understood that the MoE runner and the FP4 GEMM runner are separate components with different backend requirements. The MoE runner handles the expert routing, gating, and sparse computation — operations that benefit from CUTLASS's flexible kernel generation. The FP4 GEMM runner handles the dense matrix multiplications within each expert — operations where cuDNN's Blackwell-optimized kernels provide numerical stability.
Assumptions and Potential Pitfalls
The message reveals several assumptions worth examining. First, the assistant assumes that sgl-kernel version 0.3.21, which was installed via pip (likely from a pre-built wheel), already includes SM120 support. This is a reasonable assumption — the version number suggests a relatively recent release — but it's not guaranteed. NVIDIA's kernel libraries often require explicit architecture flags during compilation, and a generic wheel might not include SM120-specific code paths.
Second, the assistant assumes that the backend flags alone will resolve the NaN issue without requiring any changes to the attention backend or other configuration. This is supported by the fact that the attention backend was already set to triton (required for hybrid GDN models on Blackwell), and the NaN was specifically a GEMM/MoE computation issue, not an attention issue.
Third, the assistant implicitly assumes that the gist's instructions are correct and applicable to this exact setup. The gist was titled "Qwen 3.5 NVFP4 Setup SM120," which matches the model and hardware perfectly, but the assistant hadn't verified the gist's contents beyond the HTML fetch. The trust in the gist's authority is an assumption — albeit a reasonable one given that the user explicitly provided it as a reference.
The Broader Significance
This message represents a classic pattern in ML infrastructure debugging: the "wrong kernel backend" problem. As GPU architectures diversify and quantization schemes multiply, the combinatorial explosion of backend configurations becomes one of the hardest challenges in deployment. A model that works perfectly on one GPU architecture may produce NaN outputs on another simply because the default kernel backend isn't optimized for the new hardware.
The assistant's approach — isolate the symptom (NaN), identify the component (FP4 GEMM and MoE runners), consult an external reference (the gist), apply the minimal fix first (flags only, deferring kernel rebuild) — is a model of disciplined debugging. It also highlights the importance of community knowledge sharing through gists and documentation, especially for bleeding-edge hardware like Blackwell where official documentation may lag behind practical experience.
The message also demonstrates the value of understanding the system architecture at multiple levels. The assistant didn't just blindly apply flags from the gist — it understood why those flags mattered, which allowed it to make an informed decision about the order of operations. This depth of understanding is what separates a script-follower from a true engineer.
Conclusion
Message [msg 5854] is a compact but dense example of applied ML systems engineering. In just a few lines of reasoning and a single bash command, the assistant diagnosed a numerical stability issue in FP4 inference on Blackwell GPUs, identified the correct backend configuration from an external reference, and executed the fix with minimal disruption. The message captures the moment of diagnostic insight — the "aha" where scattered clues (garbage output, NaN finish reason, gist reference) coalesce into a coherent action plan. It's a reminder that in the world of large-scale model deployment, the most impactful work often happens not in writing new code, but in finding the right combination of flags to make existing code work correctly on new hardware.