The SM120 FP4 Kernel Disaster: A Deep Dive into Blackwell's Broken Matrix Math

Introduction

In the high-stakes world of large language model inference, few moments are as simultaneously exhilarating and devastating as the one captured in message 862 of this opencode session. The assistant, having just spent hours benchmarking the GLM-5-NVFP4 model across 8 NVIDIA RTX PRO 6000 Blackwell GPUs, sits down to synthesize the findings from three parallel research agents. What emerges is not merely a performance analysis—it is a forensic investigation into a fundamental architectural betrayal. The message, titled internally as "Root Cause Analysis: SM120 FP4 Kernel Disaster," pulls back the curtain on why Blackwell's consumer-grade SM120 microarchitecture is catastrophically underperforming for FP4 matrix operations, and why the model's inference throughput is stuck at a fraction of its theoretical potential.

This article examines message 862 in isolation, exploring its reasoning, its assumptions, its mistakes, and the critical knowledge it produces. The message is a turning point in the session: the moment when vague suspicions of "untuned kernels" crystallize into a precise, multi-layered diagnosis of hardware and software failure.

Why This Message Was Written: The Context

To understand message 862, one must understand the chain of events that led to it. The session had been a long, grueling optimization campaign for the GLM-5-NVFP4 model, a massive Mixture-of-Experts language model quantized to NVIDIA's FP4 format. The assistant had been running on 8 RTX PRO 6000 Blackwell GPUs inside a Proxmox LXC container, and had achieved a respectable ~3,740 total tok/s throughput through careful tuning of FlashInfer's CUTLASS MoE autotune and server parameters.

But a nagging question remained: was the model communication-bound or compute-bound? The assistant designed an experiment to find out. If the bottleneck was the allreduce communication between GPUs during tensor parallelism (TP), then reducing the TP degree from 8 to 4 while adding pipeline parallelism (PP) of 2 should improve throughput—fewer GPUs to synchronize, less communication overhead. If the bottleneck was computation, the opposite would happen: each GPU would have to compute larger matrix multiplications (1/4 of each layer instead of 1/8), and throughput would drop.

The assistant benchmarked TP4+PP2 and found it was roughly 2× slower than TP8 across the board. At 256 concurrency, TP8 achieved 1,867 tok/s while TP4+PP2 managed only 727 tok/s. At 512 concurrency, the gap was 2,800 tok/s vs 1,318 tok/s. The user, observing these results, made a sharp comment in message 858: "Doesn't the halving of perf sort of imply we're compute bound -> kernels not tuned to the GPU?"

This was the spark. The assistant immediately recognized the correctness of this observation and, in message 859, laid out the mathematical reasoning: the ~2× slowdown directly correlates with the 2× larger GEMMs per GPU. If the bottleneck were communication, TP4 would be faster. The fact that it's slower points to compute-bound execution, which means the FP4 GEMM kernels are the problem.

In message 861, the assistant dispatched three parallel research tasks (subagents) to investigate:

  1. Research SM120 FP4 GEMM kernel support
  2. Check CUTLASS autotune results on the remote machine
  3. Research SM120 tensor core specifications Message 862 is the synthesis of these three research threads. It is the moment of diagnosis.

The Synthesis: Four Layers of Failure

The assistant opens message 862 with a tone of grim discovery: "This is a goldmine. Let me synthesize the critical findings." The word "goldmine" is telling—it conveys both the richness of the data and the severity of the problem. The assistant then presents a "Root Cause Analysis: SM120 FP4 Kernel Disaster" organized into four findings, each one peeling back a deeper layer of dysfunction.

Finding 1: SM120 Tensor Cores Are a Different Architecture Entirely

The first and most fundamental finding is that SM120 (used in the RTX PRO 6000 Blackwell and RTX 5090) is not the same architecture as SM100 (used in the B200/B100 datacenter GPUs), despite both being marketed as "Blackwell." The assistant explains that SM120 uses "Ampere-era mma.sync instructions" (warp-level, 32 threads), while SM100 uses the new tcgen05.mma autonomous tensor unit with dedicated tensor memory (TMEM). This is a critical architectural divergence: SM120 gets FP4 data types bolted onto an older programming model, while SM100 has dedicated hardware for the new formats. The result is that per-SM FP4 throughput is 2–4× lower on SM120 than on SM100.

This finding is significant because it challenges the assumption that "Blackwell" means a unified architecture. The assistant had been operating under the implicit assumption that the FP4 kernels tuned for B200 would work similarly on the RTX PRO 6000. This assumption was wrong, and the performance numbers were the evidence.

Finding 2: FlashInfer's CUTLASS FP4 Is Silently Returning Zeros

The second finding is more alarming. The assistant references a FlashInfer bug (#2577) filed just the day before (February 18, 2026). According to this bug report, the cutlass backend for mm_fp4 (the FP4 matrix multiplication used in attention projections) returns all zeros on SM120. The cudnn backend is not supported on SM120, and the trtllm backend explicitly rejects SM120.

If the attention FP4 GEMMs are returning zeros, the model's outputs should be garbage. Yet the model was producing coherent text in earlier benchmarks. The assistant notes this contradiction with concern: "If mm_fp4 (attention FP4 GEMMs) is returning zeros, the model might be working through some fallback or the MoE path is different, but this is extremely concerning." This is a moment of intellectual honesty—the assistant acknowledges that the data doesn't fully align with the hypothesis, and flags it for further investigation.

Finding 3: MoE Autotune Is Losing the Best Tactics

The third finding addresses the MoE (Mixture-of-Experts) path specifically. The assistant had previously enabled FlashInfer's CUTLASS MoE autotune, which searches for the best kernel tile configurations. The investigation reveals that the two largest tile shapes—M128×256×128 and M256×128×128—fail with "Error Internal" during autotune. The root cause is SM120's 100KB shared memory limit (actually 99KB usable), which is insufficient for these larger tiles. The surviving tactics use smaller tiles with worse latency hiding, meaning the MoE kernels are running suboptimally by necessity, not by choice.

Finding 4: Memory Bandwidth Is the Real Wall

The fourth finding reframes the entire problem. The RTX PRO 6000 has 1.6 TB/s of GDDR7 memory bandwidth, compared to the B200's 8 TB/s of HBM3e. With approximately 2,000 TOPS of FP4 compute but only 1.6 TB/s of bandwidth, the ops:byte ratio needed to be compute-bound is approximately 1,250. During decode with small batch sizes (per-expert batches of 16–64 tokens), the workload is almost always memory-bandwidth-bound, not compute-bound. This means that even if the kernel issues were fixed, the GPUs would still be limited by memory bandwidth for small-batch decode operations.

This finding is particularly insightful because it reframes the problem from "the kernels are untuned" to "the hardware has a fundamental bandwidth bottleneck that kernel tuning alone cannot fix." The assistant is careful to note that this applies to small batch sizes during decode—at larger batch sizes, the compute wall may still be reachable.

The Bash Command: Testing the Hypothesis

After presenting the four findings, the assistant outlines next steps: investigating whether cuBLASLt has properly tuned SM120 FP4 kernels (it reportedly achieves 99.6 TFLOPS on SM121), and verifying whether the attention FP4 path is actually working or silently corrupted.

To test this, the assistant runs a Python script via SSH that creates random FP4 inputs and calls flashinfer.mm_fp4 with the cutlass backend. The result is:

Device: NVIDIA RTX PRO 6000 Blackwell Server Edition
Capability: (12, 0)
CUTLASS error: K dimension mismatch in mm_fp4. got a.shape[1] = 128, b.shape[0] = 256

The test fails with a dimension mismatch error, not the zeros bug. This is because the assistant's test code incorrectly constructed the FP4 tensors—the packed FP4 format has specific dimension requirements (K must be divisible by the block size and the packing factor). The assistant created a_fp4 with shape (M, K//2) expecting K=256, but mm_fp4 expects the K dimension to be handled differently.

This is a subtle mistake. The assistant was trying to test whether mm_fp4 returns zeros, but the test itself had a bug in the tensor construction. The error message—"K dimension mismatch"—is a red herring that prevents the assistant from getting the answer it sought. The zeros bug remains unverified in this message.

Assumptions and Their Consequences

Message 862 reveals several assumptions, some correct and some incorrect:

Correct assumption: The TP4+PP2 benchmark conclusively demonstrates that the model is compute-bound. The 2× slowdown with larger per-GPU GEMMs is strong evidence that kernel efficiency, not communication, is the primary bottleneck.

Correct assumption: SM120 is architecturally different from SM100, and FP4 support on SM120 is immature. The research findings confirm this.

Incorrect assumption (implicit): The mm_fp4 test code correctly constructs FP4 tensors. The dimension mismatch error reveals that the assistant did not fully understand the FP4 packing format's dimension requirements. The tensors need to be constructed with specific block sizes and layout conventions that the test code didn't follow.

Incorrect assumption (implicit): The zeros bug from FlashInfer issue #2577 would be reproducible with a simple random tensor test. In reality, the zeros bug may depend on specific scale tensor values, matrix dimensions, or other conditions that the simple test didn't replicate.

Unverified assumption: cuBLASLt might have better SM120 FP4 kernels. The assistant proposes this as a next step but hasn't tested it yet. This assumption turns out to be partially correct in later messages (cuBLASLt is found to be no faster than CUTLASS).

Input Knowledge Required

To fully understand message 862, the reader needs knowledge of:

  1. NVIDIA GPU architecture generations: Understanding that SM120 (consumer Blackwell) and SM100 (datacenter Blackwell) are different microarchitectures despite sharing the "Blackwell" brand name.
  2. Tensor core programming models: The difference between mma.sync (warp-level matrix multiply-accumulate, used since Volta/Ampere) and tcgen05.mma (autonomous tensor unit with dedicated TMEM, new in SM100).
  3. FP4 quantization format: NVIDIA's FP4 format packs two 4-bit values into one byte, with per-block scaling factors. The packing format affects tensor shapes and memory layouts.
  4. CUTLASS tile configurations: The M×N×K notation for GEMM tile shapes, and how shared memory limits constrain tile sizes.
  5. FlashInfer's backend system: The cutlass, cudnn, and trtllm backends for FP4 GEMM operations, and their respective SM support levels.
  6. The ops:byte ratio concept: The ratio of arithmetic operations to memory bytes needed for a kernel to be compute-bound rather than memory-bound.
  7. The session's prior context: The TP4+PP2 benchmark results, the user's observation about compute-bound behavior, and the three parallel research tasks.

Output Knowledge Created

Message 862 produces several important pieces of knowledge:

  1. A confirmed diagnosis: The model is compute-bound, and the FP4 GEMM kernels on SM120 are fundamentally broken or suboptimal for multiple independent reasons.
  2. A prioritized list of issues: Architectural mismatch (SM120 vs SM100), the zeros bug in FlashInfer CUTLASS, the MoE autotune tile size failures, and the memory bandwidth wall.
  3. A test result: The mm_fp4 call on SM120 fails with a dimension mismatch, confirming that the FP4 path is not trivially working (though not confirming the zeros bug specifically).
  4. A research methodology: The message demonstrates how to combine web research (subagent tasks), local file inspection (autotune cache check), and live GPU testing (the Python script) to diagnose a complex performance issue.
  5. A roadmap for next steps: Investigate cuBLASLt, verify the attention FP4 path, and potentially explore alternative kernel backends.

The Thinking Process

The assistant's reasoning in message 862 is a model of systematic diagnostic thinking. The structure is:

  1. Synthesize the three research threads into a coherent picture
  2. Organize findings by severity and causality (architecture → kernel bug → autotune failure → fundamental bandwidth limit)
  3. Acknowledge contradictions (the model works despite the zeros bug)
  4. Propose concrete next steps (test cuBLASLt, verify attention FP4 path)
  5. Execute a test to validate or invalidate the hypothesis The thinking is notable for its intellectual honesty. When the assistant says "the situation is worse than 'not tuned' — it's fundamentally broken in multiple ways," it's acknowledging that the problem is deeper than simple optimization. The word "disaster" in the section title is not hyperbole—it's an accurate assessment of the situation. The bash command at the end shows the assistant moving from theory to practice. Rather than just reporting the research findings, it immediately tries to verify them with a live test. The test fails, but not in the way expected—and the assistant doesn't try to hide this or spin it. The raw error output is presented without commentary, letting the reader (and the user) draw their own conclusions.

Conclusion

Message 862 is a turning point in the GLM-5-NVFP4 optimization campaign. It transforms vague performance suspicions into a precise, multi-layered diagnosis of hardware and software failure. The assistant demonstrates how to combine web research, local investigation, and live testing to diagnose a complex performance issue, and it does so with intellectual honesty and systematic thinking.

The message also reveals the challenges of deploying cutting-edge AI models on new hardware architectures. The RTX PRO 6000 Blackwell GPUs, despite their impressive specifications, suffer from immature software support, architectural compromises, and fundamental bandwidth limitations that no amount of kernel tuning can fully overcome. The "SM120 FP4 Kernel Disaster" is a cautionary tale about the gap between hardware marketing and real-world performance—a gap that only careful, systematic investigation can reveal.