The Last Backend Standing: Systematically Eliminating GPU Kernel Options for Blackwell SM120
In the high-stakes world of deploying large language models on cutting-edge hardware, there is rarely a single "correct" answer — only a series of increasingly refined approximations. Message [msg 5972] captures one such moment of convergence: the launch of an SGLang inference server for the 397-billion-parameter Qwen3.5-A17B-NVFP4 model on eight NVIDIA RTX PRO 6000 Blackwell GPUs, using a carefully chosen combination of backend kernels. On its surface, the message is a single bash command — a nohup launch of sglang.launch_server with a dozen flags. But beneath that surface lies the culmination of an exhaustive debugging session that tested four different GPU kernel backends across two separate computational paths (Mixture-of-Experts routing and dense FP4 matrix multiplication), eliminating each in turn until only one viable configuration remained.
The Message: A Launch Command Under the Microscope
The command itself is dense with meaning:
ssh root@10.1.230.174 'export CUDA_HOME=/usr/local/cuda-13.0 && \
export PATH="/usr/local/cuda-13.0/bin:/root/ml-env/bin:$PATH" && \
export LD_LIBRARY_PATH=/usr/local/cuda-13.0/lib64 && \
export OMP_NUM_THREADS=16 && \
nohup /root/ml-env/bin/python3 -m sglang.launch_server \
--model-path /data/models/Qwen3.5-397B-A17B-NVFP4 \
--served-model-name qwen3.5-397b \
--tp 8 \
--trust-remote-code \
--quantization modelopt_fp4 \
--host 0.0.0.0 \
--attention-backend triton \
--moe-runner-backend flashinfer_cutlass \
--fp4-gemm-backend flashinfer_cutlass \
--reasoning-parser qwen3 \
--tool-call-parser qwen3_coder \
--disable-custom-all-reduce \
> /tmp/sglang_allcutlass_test.log 2>&1 &'
Every flag tells a story. The environment variables establish a CUDA 13.0 toolchain — a recent upgrade that was itself the subject of extensive effort earlier in the session (see [msg 5936]). The --tp 8 flag spreads the model across all eight GPUs using tensor parallelism. The --quantization modelopt_fp4 selects NVIDIA's ModelOpt FP4 quantization format, which stores weights in 4-bit floating point for dramatically reduced memory footprint. But the two critical flags are --moe-runner-backend flashinfer_cutlass and --fp4-gemm-backend flashinfer_cutlass — both set to the same cutlass backend, representing the assistant's final hypothesis after a systematic elimination process.
The Reasoning: Why This Command Was Written
This message was written because the assistant had spent the preceding hour testing backend after backend, and each had failed in a different way. The context messages leading up to [msg 5972] reveal a rigorous experimental methodology:
flashinfer_trtllmfor MoE ([msg 5960]): Crashed immediately with a process kill. The TRT-LLM fused MoE kernel attempted to compile for SM100 compute capability (major version 10), but the RTX PRO 6000 Blackwell GPUs are SM120 (major version 12). The kernel binary was incompatible.flashinfer_cutedslfor MoE ([msg 5963]): Loaded successfully and began CUDA graph capture — a promising sign — but produced garbage output. The smoke test detected repeated!characters, a hallmark of NaN or corrupted computation in the MoE pathway. The CUTE DSL JIT compiler could generate code for SM120, but the generated kernels produced incorrect results.flashinfer_trtllmfor dense FP4 GEMM ([msg 5969]): Also crashed immediately. The TRT-LLM dense FP4 kernels, like their MoE counterparts, were compiled only for SM100 and lacked SM120 support.flashinfer_cudnnfor dense FP4 GEMM (implicitly tested earlier): Worked correctly, but the assistant was exploring whethercutlassmight offer better performance. Each failure was informative. The TRT-LLM failures revealed that NVIDIA's hand-tuned TensorRT-LLM kernels are Blackwell-specific in a narrow sense — they target the B100/B200 (SM100) architecture but not the RTX PRO 6000 (SM120), despite both being "Blackwell" family. The CUTE DSL failure was more subtle: the kernels compiled and ran, but produced incorrect results, suggesting either a bug in the JIT code generation for SM120 or an incompatibility with the FP4 weight layout used by the ModelOpt checkpoint.
Assumptions and Their Validity
The assistant made several assumptions in crafting this command, each carrying its own risk:
Assumption 1: flashinfer_cutlass supports SM120 for both MoE and dense FP4. This was a reasonable inference from the elimination process. The cutlass backend uses NVIDIA's CUTLASS library, which is architecture-generic and tends to have broader compatibility than TRT-LLM's hand-tuned kernels. Earlier tests in the session had confirmed that flashinfer_cutlass for MoE produced correct output on SM120 (see [msg 5953]), unlike cutedsl which produced garbage. The open question was whether the dense FP4 GEMM path would also work with cutlass — the assistant had been using cudnn for that path and was now testing cutlass for both.
Assumption 2: The environment setup (CUDA 13.0, nightly PyTorch 2.12.0, flashinfer 0.6.5) is sufficient. This was validated by earlier successful launches with other backend combinations. The stack had been painstakingly assembled over the course of the session, including building sgl-kernel from source with SM120 patches applied by community contributor catid.
Assumption 3: --disable-custom-all-reduce is necessary. This flag disables SGLang's custom all-reduce implementation, which was known to have compatibility issues with Blackwell GPUs in the PCIe configuration used here (eight GPUs connected via PCIe rather than NVLink). Earlier testing had shown that custom all-reduce caused hangs or crashes on this hardware.
Input Knowledge Required
To fully understand this message, one needs familiarity with several domains:
- GPU architecture tiers: The distinction between SM100 (B100/B200 datacenter GPUs) and SM120 (RTX PRO 6000 workstation GPUs) is critical. Both are "Blackwell" but have different compute capabilities and kernel compatibility.
- SGLang server architecture: The separation between the MoE runner backend (handling the sparse Mixture-of-Experts computation) and the FP4 GEMM backend (handling dense matrix multiplications with FP4 weights) reflects SGLang's modular design. Each can be configured independently.
- FP4 quantization: The
modelopt_fp4format stores weights in 4-bit floating point using NVIDIA's ModelOpt framework, requiring specialized kernel support for dequantization and computation. - The experimental method: The assistant's approach of testing one variable at a time — changing only the backend flags between launches — is a textbook application of the scientific method to systems debugging.
Output Knowledge Created
This message produced several forms of knowledge:
- A testable hypothesis: That
flashinfer_cutlassfor both backends would produce correct output on SM120. The server launched without immediate crash (the bash tool timed out after 120 seconds, but the process was running). - A benchmarkable configuration: Once the server was healthy, the assistant would run smoke tests and throughput benchmarks to compare this configuration against the previous
flashinfer_cutlass(MoE) +flashinfer_cudnn(dense FP4) combination. - Documentation of a working backend matrix: The systematic testing established which backends work on SM120: - MoE:
flashinfer_cutlass✅,flashinfer_trtllm❌,flashinfer_cutedsl❌ (garbage) - Dense FP4:flashinfer_cutlass✅,flashinfer_cudnn✅,flashinfer_trtllm❌
The Thinking Process
The assistant's reasoning, visible across the preceding messages, follows a clear pattern of hypothesis generation and falsification. After discovering that the default auto backend produced garbage output on SM120, the assistant methodically enumerated the available backends by inspecting FlashInfer's source code and documentation ([msg 5945]). Each backend was tested in isolation, with the assistant carefully cleaning GPU memory between tests (fuser -k /dev/nvidia*) to avoid interference.
The pivot to flashinfer_cutlass for both backends represents a conservative choice — the assistant had already confirmed that cutlass works for MoE and that cudnn works for dense FP4, but wanted to test whether using cutlass for both would yield better performance through kernel reuse or reduced context switching. This is the kind of micro-optimization that only becomes possible after the correctness constraints are satisfied.
The message also reveals the assistant's operational discipline: every launch is logged to a timestamped file (/tmp/sglang_allcutlass_test.log), environment variables are explicitly set rather than relying on system defaults, and the process is launched with nohup to survive shell termination. These practices reflect experience with remote server management where connection drops are common and debugging requires persistent logs.
Conclusion
Message [msg 5972] is a single data point in a larger experimental arc, but it encapsulates the essence of systems engineering for modern AI infrastructure. The command is simple; the reasoning behind it is not. Every flag represents a lesson learned from a previous failure, and the configuration as a whole represents the intersection of what is known to work, what is hypothesized to be optimal, and what the hardware constraints allow. In the messages that follow, the assistant would confirm that this configuration produces correct output and benchmark its throughput — but the decision to launch with flashinfer_cutlass for both backends was the critical juncture where the search space collapsed to a single viable path.