The 2x Gap: Diagnosing Suboptimal FP4 GEMM Kernels on Blackwell SM120

In the high-stakes pursuit of optimizing GLM-5-NVFP4 inference across eight NVIDIA RTX PRO 6000 Blackwell GPUs, a single message from the assistant crystallized a pivotal realization: the FP4 tensor core kernels, the computational heart of the model's throughput, were delivering barely half their theoretical potential. Message [msg 870] is a compact but dense diagnostic turn in a much larger investigation—a moment where the assistant pivots from measuring raw FP4 performance to comparing it against a baseline (BF16), quantifying the efficiency gap, and then immediately searching for alternative kernel implementations that might close it.

The Investigation So Far

To understand why this message matters, we must first appreciate the context that led to it. The preceding messages ([msg 862] through [msg 869]) had been a deep, methodical autopsy of FP4 GEMM performance on SM120—the compute architecture of the RTX PRO 6000 Blackwell. The assistant had discovered a cascade of problems. FlashInfer's CUTLASS FP4 backend, the primary path for the model's matrix multiplications, was silently broken in certain configurations. The largest CUTLASS tile shapes (M128×N256 and M256×N128) failed to initialize because they overflowed SM120's 100 KB shared memory limit. The MoE autotune was losing its best tactics. And most damningly, the raw benchmark numbers told a sobering story: while the theoretical peak for FP4 on 188 SMs was around 2,000 TOPS, the CUTLASS kernels plateaued at roughly 1,200–1,300 TFLOPS—only 60–65% of peak.

But the really painful part was what happened at inference-serving batch sizes. During decode, each of the model's 256 Mixture-of-Experts (MoE) experts processes only a handful of tokens. With 512 concurrent requests and 8 experts activated per token, each expert gets roughly 16 tokens on average. At those tiny per-expert batch sizes (M=16–64), the FP4 GEMMs were achieving a mere 0.8–55 TFLOPS—between 0.02% and 3% of theoretical peak. The model was memory-bandwidth-bound, and the small GEMM sizes were catastrophic for utilization.

The Pivot: Comparing FP4 Against BF16

Message [msg 870] opens with a flash of insight that reframes the entire performance picture:

Interesting — BF16 torch.mm at batch=512 does 225 TFLOPS, while FP4 CUTLASS does 437 TFLOPS for the same dimensions. FP4 is only ~2x faster than BF16, not 4x as it should be theoretically. That confirms FP4 kernels aren't optimally tuned.

This comparison is deceptively simple but analytically powerful. The assistant has chosen the M=512, N=2048, K=6144 dimension from the earlier benchmark suite—a size that corresponds roughly to a MoE up-projection with a batch of 512 tokens under tensor parallelism across 8 GPUs. At this size, the FP4 CUTLASS kernel achieved 437 TFLOPS. By running the same dimensions through standard BF16 torch.mm (which uses cuBLAS internally), the assistant gets 225 TFLOPS.

The ratio is 437/225 ≈ 1.94x. FP4 is barely 2x faster than BF16.

Why should this be surprising? FP4 (4-bit floating point) stores data in half the bits of BF16 (16-bit), so in a memory-bandwidth-limited scenario, FP4 should load data twice as fast. Additionally, NVIDIA's tensor cores typically offer higher throughput for lower-precision formats. On the Blackwell architecture, FP4 tensor core throughput should theoretically be about 4x that of BF16—a 2x gain from the reduced data movement and another 2x from the tensor core's higher arithmetic density. Getting only 2x instead of 4x signals that something is fundamentally suboptimal in how the FP4 kernels are executing on SM120.

The Assumption Under the Microscope

The assistant's reasoning reveals a key assumption: that FP4 should deliver approximately 4x the throughput of BF16 on the same hardware. This assumption is grounded in NVIDIA's published specifications for Blackwell tensor cores, which advertise 4x FP4 over FP16/BF16. But the assumption carries nuance that the assistant implicitly acknowledges by calling the kernels "not optimally tuned."

On SM120 specifically—which uses Ampere-era mma.sync instructions retrofitted with FP4 data types, rather than the more advanced tcgen05.mma autonomous tensor unit found on SM100—the FP4 path may have architectural overhead that the simpler BF16 path avoids. The 2x gap might reflect not just suboptimal CUTLASS tile selection but also fundamental instruction-level inefficiencies in how FP4 is bolted onto the older programming model. The assistant doesn't yet know which factor dominates, but the comparison provides a clear, quantitative target: if the kernels were optimally tuned, the FP4 number should be closer to 900 TFLOPS (4× the BF16 baseline) rather than 437 TFLOPS.

Discovering the sgl_kernel Alternative

The second half of the message is a rapid, opportunistic pivot:

Also, sgl_kernel has cutlass_scaled_fp4_mm and cutlass_moe — let me benchmark those directly:

The assistant has discovered that SGLang's own kernel library (sgl_kernel) exposes FP4-specific functions that might be better tuned for the Blackwell architecture than FlashInfer's CUTLASS backend. The naming is suggestive: cutlass_scaled_fp4_mm appears to be a direct FP4 matrix multiply with scaling factors, and cutlass_moe might be a fused MoE kernel that handles the expert-scatter-gather pattern more efficiently.

The bash command that follows calls help(sgl_kernel.cutlass_scaled_fp4_mm) to inspect the API signature. The result reveals:

cutlass_scaled_fp4_mm(a: torch.Tensor, b: torch.Tensor, block_scale_a: torch.Tensor, 
                      block_scale_b: torch.Tensor, alpha: torch.Tensor, 
                      out_dtype: torch.dtype) -> torch.Tensor

This signature confirms the function takes the same fundamental inputs as FlashInfer's mm_fp4—two FP4-packed weight matrices, their per-block scaling factors, an alpha scaling coefficient, and an output data type. The similarity means the assistant can drop this function into the existing benchmark harness with minimal modification and directly compare throughput against the FlashInfer CUTLASS path.

The Deeper Significance

What makes this message compelling is not just the data it produces but the thinking process it reveals. The assistant is operating in a classic performance-engineering loop: measure, compare, hypothesize, test. The BF16 comparison provides a sanity check that quantifies the efficiency gap. The discovery of sgl_kernel functions opens a new branch of investigation—perhaps the SGLang team has already tuned these kernels for SM120, or perhaps they use different CUTLASS tile configurations that avoid the shared memory overflow that plagued FlashInfer's autotune.

The message also demonstrates a crucial skill in systems optimization: knowing when to zoom out. After spending several messages deep in the weeds of FlashInfer's FP4 implementation—testing backends, debugging scale tensor types, verifying non-zero outputs—the assistant zooms out to compare against a completely different implementation path (BF16 cuBLAS) and a completely different kernel library (sgl_kernel). This cross-validation is what separates superficial benchmarking from genuine performance understanding.

Output Knowledge and Next Steps

This message produces several concrete outputs. First, it establishes a quantitative baseline: BF16 achieves 225 TFLOPS for the M=512, N=2048, K=6144 case, giving a reference point for all future FP4 comparisons. Second, it confirms the hypothesis that FP4 kernels are underperforming—the 2x ratio versus the expected 4x is a clear signal. Third, it identifies sgl_kernel.cutlass_scaled_fp4_mm as a promising alternative path worth benchmarking.

The message ends mid-investigation, with the assistant having just discovered the API signature but not yet run the benchmark. The reader is left in suspense: will sgl_kernel's FP4 implementation be faster? Will it close the gap from 2x toward 4x? The answer will come in subsequent messages, but the framing has already shifted the investigation. The question is no longer "how fast is FlashInfer's FP4?" but "which FP4 implementation is closest to the theoretical peak, and what can we learn from the gap?"

For anyone following the optimization journey, message [msg 870] is the moment where the assistant stops taking FP4 performance at face value and starts asking the harder question: compared to what?