The Moment the Zeros Disappeared: A Pivotal FP4 Kernel Test on Blackwell SM120

In the middle of a deep investigation into why the GLM-5-NVFP4 model was performing poorly on eight NVIDIA RTX PRO 6000 Blackwell GPUs, the assistant executed a small Python test that quietly overturned a critical assumption. The message at <msg id=866> is deceptively simple: a script calling flashinfer.mm_fp4 with the cutlass backend, checking whether the output is all zeros. But behind this seemingly trivial test lies a moment of scientific correction — where a hypothesis formed from a GitHub issue filed just 24 hours earlier was put to the test and found to be incomplete.

The Investigation That Led Here

The story begins with a performance mystery. The assistant had been benchmarking the GLM-5-NVFP4 model — a Mixture-of-Experts model using NVIDIA's FP4 (4-bit floating point) quantization format — across different parallelism configurations. The model was deployed using SGLang with FlashInfer kernels on eight RTX PRO 6000 Blackwell GPUs (compute capability 12.0, architecture codenamed SM120).

The key finding was puzzling: switching from TP8 (tensor parallelism across all 8 GPUs) to TP4+PP2 (tensor parallelism across 4 GPUs with 2-stage pipeline parallelism) resulted in roughly half the throughput. The user astutely observed in <msg id=858> that this pattern suggested the model was compute-bound rather than communication-bound — because if the bottleneck were allreduce communication, halving the number of allreduce participants should have helped, not hurt. The fact that throughput dropped in proportion to the doubled per-GPU matrix sizes pointed squarely at kernel efficiency.

This launched a multi-pronged investigation. The assistant dispatched three research agents in <msg id=861> to study SM120 FP4 GEMM kernel support, CUTLASS autotune results, and SM120 tensor core specifications. The synthesized findings in <msg id=862> painted a dire picture: SM120's tensor cores were architecturally distinct from the datacenter SM100 (Blackwell B200/B100), using older Ampere-era mma.sync instructions instead of SM100's dedicated tcgen05.mma tensor units. But the most alarming finding was about FlashInfer specifically.

The "All Zeros" Hypothesis

The research in <msg id=862> uncovered FlashInfer bug #2577, filed just one day earlier (February 18, 2026). The bug report claimed that on SM120, the cutlass backend for mm_fp4 — the function responsible for FP4 matrix multiplication in attention projections — returns all zeros. The cudnn backend was reported as unsupported on SM120, and the trtllm backend explicitly rejected SM120.

This was a bombshell. If the attention FP4 GEMMs were silently producing zeros, the model's outputs would be garbage — or the model might be working through some fallback path that the assistant hadn't yet identified. The assistant declared in <msg id=862>: "The situation is worse than 'not tuned' — it's fundamentally broken in multiple ways."

But a hypothesis, no matter how well-sourced, demands experimental verification.

The Testing Journey

The assistant began testing in <msg id=862> with a straightforward call to flashinfer.mm_fp4, but immediately hit a "K dimension mismatch" error — the tensor shapes were wrong. After consulting the API documentation in <msg id=863>, the assistant corrected the shapes in <msg id=864> but encountered a NotImplementedError because Float8_e4m3fn scale tensors couldn't be directly created on CUDA. A workaround was attempted in <msg id=865> using uint8 tensors viewed as float8_e4m3fn, but this produced a cryptic TypeError about mismatched argument types — the cutlass backend expected a DLTensor* but received None.

Each failure taught the assistant something about the API. The function signature, revealed via help(flashinfer.mm_fp4) in <msg id=863>, showed optional parameters alpha and out that the earlier tests had omitted. The assistant had been calling the function with positional arguments that didn't align with what the underlying C++ kernel expected.

The Subject Message: A Corrected Test

In <msg id=866>, the assistant synthesized everything learned from the previous failures and wrote a corrected test. The key changes were:

  1. Providing the alpha parameter: A scalar tensor alpha = torch.tensor(1.0, ...) was explicitly created and passed. This scaling factor is used in the fused multiply-add operation that the kernel performs.
  2. Providing the out parameter: An output tensor out = torch.zeros(M, N, ...) was pre-allocated and passed. This allows the kernel to write results into a pre-existing buffer rather than allocating internally.
  3. Correct tensor shapes: The b tensor was shaped as (K//2, N) — matching the column-major convention where the second matrix has K rows and N columns — rather than the earlier incorrect (N, K//2).
  4. Proper scale tensor construction: The uint8-to-float8_e4m3fn view trick was retained from the previous attempt, but now the function was called with the complete set of required arguments. The test also checked the auto backend, which would attempt to select the best available backend automatically.

The Results: A Correction

The output was unambiguous:

Device: NVIDIA RTX PRO 6000 Blackwell Server Edition
Capability: (12, 0)
cutlass: shape=torch.Size([128, 256]), all_zeros=False, abs_max=18743296.0000, abs_mean=626688.000000
auto ERROR: RuntimeError: detail::finalize(matmul_operation.get_raw_desc()) failed with message: ptrDesc->finalize(), and code: CUDNN_STATUS_BAD_PARAM

The cutlass backend did NOT return all zeros. The all_zeros flag was False, and the output had a mean absolute value of 626,688 with a maximum of 18,743,296. These numbers are large — suspiciously large for what should be a normalized matrix multiplication — but that's likely an artifact of using random scale factors. The critical finding is that the kernel is alive and producing non-zero output.

The auto backend failed with a CUDNN_STATUS_BAD_PARAM error, confirming that the automatic backend selection tries cuDNN first and fails on SM120. This means that for the auto backend, the fallback chain doesn't reach cutlass — it simply errors out.

What This Means

This single test result has profound implications for the investigation:

The "all zeros" hypothesis is disproven for this specific configuration. The cutlass FP4 kernel on SM120 does produce meaningful output when called with the correct API signature. This means the model's attention projections are likely computing correctly — the performance problem is not that the kernels are silently broken, but that they are inefficient.

The auto backend is broken. Since auto tries cuDNN first and cuDNN rejects SM120, any code using backend="auto" on SM120 will fail. This is a critical finding for the SGLang deployment, which may be using the auto backend internally.

The API is fragile. The fact that omitting alpha and out caused a TypeError rather than a graceful fallback suggests that the Python wrapper has a bug or an undocumented requirement. The assistant discovered this through trial and error — a process that consumed several rounds of debugging.

The Thinking Process

What's visible in the assistant's reasoning is a methodical, hypothesis-driven approach. The initial hypothesis ("cutlass returns zeros on SM120") was treated as provisional — worth investigating but not accepted as fact. Each failed test narrowed down the problem: first shape mismatch, then type issues, then missing arguments. The assistant didn't give up after three consecutive failures; instead, each error message was parsed for clues about what the API actually required.

The assistant's comment "Interesting — cutlass needs more arguments" shows a shift in mental model. The earlier assumption was that the kernel was fundamentally broken. The new understanding is that the kernel works but has a fussy API. This is a much more tractable problem — it means optimization efforts can focus on performance rather than correctness.

Output Knowledge Created

This message produces several concrete pieces of knowledge:

  1. Empirical confirmation: The FlashInfer CUTLASS FP4 kernel on SM120 (RTX PRO 6000 Blackwell, compute 12.0) produces non-zero output when called correctly.
  2. API requirements documented: The mm_fp4 function requires explicit alpha and out tensors when using the cutlass backend on SM120 — omitting them causes a TypeError.
  3. Auto backend failure confirmed: The auto backend fails on SM120 because it attempts cuDNN first, which rejects the architecture.
  4. A reproducible test case: The exact Python script serves as a minimal verification test that can be reused by anyone deploying FP4 models on SM120 hardware.

The Broader Context

This message sits at a turning point in the investigation. Before it, the team was operating under the assumption that the FP4 kernels were fundamentally broken — a "disaster" as the assistant described it. After it, the focus shifts to performance tuning: the kernels work, but they're slow. The assistant goes on to achieve a 28% throughput improvement through server parameter tuning, and begins documenting optimization approaches in glb5improvement-xx.md files (as noted in the chunk summary for segment 7).

The lesson is a classic one in systems debugging: when a performance problem seems catastrophic, verify the catastrophic hypothesis first. A GitHub issue filed yesterday is not a substitute for running the test yourself. The zeros disappeared, and the real work of optimization could begin.