The FP4 GEMM Bottleneck: When 256 Experts Starve Your GPUs

In the high-stakes world of large language model inference, few moments are as illuminating as the one where raw benchmark data crystallizes into a clear understanding of the bottleneck. Message [msg 869] in this opencode session represents exactly such a turning point. The assistant, having spent multiple rounds investigating FP4 GEMM kernel performance on NVIDIA's new Blackwell SM120 architecture (the RTX PRO 6000), synthesizes a comprehensive set of microbenchmarks into a single table that tells a devastating story: the model's Mixture-of-Experts architecture, combined with the realities of small-batch decode, is starving the GPUs of work.

The Investigation So Far

To understand why message [msg 869] is so significant, we must first understand the journey that led to it. The session's overarching goal was deploying the GLM-5-NVFP4 model—a massive Mixture-of-Experts language model using FP4 (4-bit floating point) quantization—across eight RTX PRO 6000 Blackwell GPUs. Earlier segments had already resolved numerous infrastructure challenges: NaN decode crashes, virtualization-induced PCIe bottlenecks, and CUDA initialization failures in LXC containers. By segment 7, the assistant had achieved a respectable ~3,740 tok/s throughput, but was pushing for more.

The critical question was: why aren't the GPUs fully utilized? Power draw was hovering around 235W per GPU out of a 600W TDP. Something was holding them back. The assistant hypothesized that the bottleneck might be communication-related (allreduce latency in tensor parallelism) or compute-related (inefficient FP4 GEMM kernels).

A TP4+PP2 benchmark (tensor parallelism 4 + pipeline parallelism 2) had already confirmed the model was compute-bound rather than communication-bound—it was 2× slower than TP8, ruling out allreduce latency as the primary issue. This focused the investigation squarely on the FP4 matrix multiplication kernels running on SM120.

The SM120 Disappointment

Through a series of research subagents and web searches (see [msg 861]), the assistant had uncovered a troubling reality: SM120 (Blackwell desktop/workstation architecture) does not have the same tensor core capabilities as SM100 (Blackwell server architecture). Despite both being branded "Blackwell," SM120 uses older Ampere-era mma.sync instructions with FP4 data types bolted on, while SM100 has the new tcgen05.mma autonomous tensor unit with dedicated TMEM memory. Per-SM FP4 throughput on SM120 is estimated at 2–4× lower than SM100.

Even worse, FlashInfer's CUTLASS FP4 backend was found to have a bug (issue #2577, filed just the day before this message) where it silently returns zeros on SM120 under certain configurations. The cudnn backend doesn't support SM120 at all, and the trtllm backend explicitly rejects it. The MoE autotune was also failing on the two largest tile shapes (M128×256×128 and M256×128×128) because they overflow SM120's 100KB shared memory limit.

The Benchmark That Changed Everything

Message [msg 869] opens with a series of microbenchmarks the assistant ran on the actual hardware. These benchmarks measured FP4 GEMM performance across a range of matrix sizes, from tiny single-token decode operations up to large 16K×8K matrices designed to saturate the GPU. The results are presented in a clean table:

| Size | TFLOPS | % of ~2000 TOPS Peak | Notes | |------|--------|---------------------|-------| | Batch=1 decode | 0.3-0.8 | <0.1% | Pure memory-bound (launch overhead) | | Batch=64 | 15-55 | 1-3% | Still launch-overhead dominated | | Batch=256 | 81-219 | 4-11% | Starting to compute | | Batch=512 | 161-437 | 8-22% | Getting reasonable | | Batch=1024 | 849 | 42% | Good | | Batch=4K+ | 1,190-1,313 | 60-66% | Approaching saturation at ~1,300 TFLOPS |

The table reveals two critical facts. First, the CUTLASS FP4 kernels do work on SM120 and can achieve respectable throughput—up to ~1,300 TFLOPS for large matrices, which is about 65% of the estimated 2,000 TOPS theoretical peak. This is not terrible for a first-generation CUTLASS implementation on a new architecture, and it disproves the fear that the kernels were silently returning garbage.

But the second fact is the killer: at the batch sizes that matter for decode, performance is catastrophic. A single token achieves just 0.3–0.8 TFLOPS—less than 0.1% of peak. Even at batch 64, the throughput is only 1–3% of peak. The GPUs are spending virtually all their time on kernel launch overhead and memory latency, not actual computation.

The MoE Expert Scattering Problem

The assistant then connects these microbenchmark results to the real inference scenario with a crucial piece of analysis:

The critical insight for inference performance: during decode, each MoE expert processes at most batch_size / num_experts_per_token tokens. With 256 experts, 8 activated per token, and 512 concurrent requests, each expert gets ~512*8/256 = 16 tokens on average. That puts us in the 15-55 TFLOPS range — only 1-3% of peak! The small GEMM sizes are catastrophic for utilization.

This is the heart of the bottleneck. GLM-5-NVFP4 has 256 experts, with 8 activated per token via a routing mechanism. Even with 512 concurrent requests in flight, the tokens are scattered across the experts such that each expert processes only about 16 tokens per invocation. A 16×2048×6144 GEMM (the typical shape for an MoE expert's up-projection) achieves only about 55 TFLOPS—barely 3% of the GPU's theoretical capability.

The problem is structural: the MoE architecture with a large number of experts inherently creates tiny per-expert batch sizes during decode, and those tiny batches cannot utilize the tensor cores efficiently. The GPUs are not being lazy; they are being starved of work by the scattering of tokens across experts.

Assumptions and Their Validity

This analysis rests on several assumptions worth examining. The assistant assumes that the CUTLASS FP4 backend is the one actually being used during inference—a reasonable assumption given that the auto backend falls back to CUTLASS when cudnn fails and trtllm rejects SM120. The assistant also assumes the estimated 2,000 TOPS peak is accurate for SM120, which is based on extrapolation from published specifications rather than direct measurement. The benchmark results themselves assume that the microbenchmark GEMMs are representative of the actual model's computation patterns, which is reasonable given the assistant has the model's architecture parameters (hidden_size=6144, moe_intermediate_size=2048).

One potential blind spot is the assumption that the FP4 GEMM is the only significant compute operation. In reality, attention mechanisms, layer normalization, activation functions, and routing all consume GPU cycles too. However, for a MoE transformer, the FP4 GEMMs in the expert layers dominate the compute budget, so this assumption is likely safe for the purpose of identifying the primary bottleneck.

Knowledge Created

This message creates several important pieces of knowledge. First, it establishes a quantitative understanding of FP4 GEMM performance on SM120 across the full range of operation sizes, from 0.3 TFLOPS for single tokens to 1,300 TFLOPS for large matrices. This is valuable data that didn't exist before—the assistant had to run these benchmarks because no published benchmarks for SM120 FP4 GEMMs were available.

Second, it creates the analytical framework connecting microbenchmark performance to real inference throughput. The insight that per-expert batch size = batch_size × experts_per_token / total_experts is simple but powerful. It explains why simply increasing the number of concurrent requests yields diminishing returns: even at 2048 concurrent requests, each expert only gets ~64 tokens, still in the 4–11% utilization range.

Third, it definitively confirms that the model is compute-bound (or more precisely, memory-bandwidth-bound for small GEMMs) rather than communication-bound. This rules out a whole class of potential optimizations (better allreduce algorithms, NVLink tuning, etc.) and focuses attention on the actual problem: making small FP4 GEMMs run faster.

The Thinking Process

The assistant's reasoning in this message is a model of systematic data-driven analysis. It begins with the raw benchmark numbers, organizes them into a clear table, and then extracts the key insight about plateau behavior (~1,300 TFLOPS). Rather than stopping there, it immediately connects this to the real inference scenario by computing the per-expert batch size. The arithmetic is straightforward but revealing: 512 * 8 / 256 = 16. This transforms the abstract benchmark table into a concrete explanation for the observed GPU underutilization.

The assistant also shows intellectual honesty by acknowledging that the kernels "plateau around 1,200-1,300 TFLOPS — about 65% of the estimated 2,000 TOPS peak. That's not terrible for CUTLASS on a new architecture, but there's still ~35% headroom." This balanced assessment avoids both panic (the kernels are broken!) and complacency (everything is fine!). The kernels work, they're not optimally tuned, but the real problem is structural rather than kernel-level.

The Follow-Up: cuBLASLt Investigation

The message concludes with a bash command to check whether cuBLASLt (NVIDIA's proprietary BLAS library) offers better FP4 performance than CUTLASS. This is a logical next step: if NVIDIA's own library has better-tuned kernels for SM120, routing the GEMMs through cuBLASLt could recover some of the 35% headroom. The command also checks for sgl_kernel, which is SGLang's custom kernel library, and discovers it has several FP4-related functions including cutlass_fp4_group_mm and cutlass_scaled_fp4_mm.

The results show that BF16 torch.mm achieves 224.7 TFLOPS for a 512×2048×6144 GEMM—impressively high, but this is a dense matrix multiplication with no expert scattering. The FP4 version of the same operation would need to be compared directly. The discovery of sgl_kernel's FP4 functions opens another avenue for optimization: perhaps SGLang's custom kernels are better tuned for SM120 than the generic FlashInfer CUTLASS path.

Broader Implications

This message has implications beyond this specific deployment. The FP4 GEMM performance cliff at small batch sizes is a fundamental challenge for MoE models with many experts. It suggests that the trend toward larger numbers of experts (Mixtral uses 8, GLM-5 uses 256, and some models are exploring thousands) comes with a hidden cost: the per-expert batch size shrinks proportionally, making efficient GPU utilization increasingly difficult.

For the GLM-5-NVFP4 deployment specifically, this analysis points toward several optimization strategies: increasing the number of concurrent requests to push per-expert batch sizes higher, exploring expert parallelism to aggregate tokens across GPUs before routing to experts, and investigating kernel-level optimizations for small GEMMs (such as persistent kernels that avoid launch overhead). The assistant would go on to explore all of these in subsequent messages, documenting them as glb5improvement-xx.md files.

Conclusion

Message [msg 869] is a masterclass in diagnostic reasoning. It takes raw benchmark data, organizes it into a coherent picture, connects it to the real system architecture, and produces a clear, actionable understanding of the bottleneck. The assistant doesn't just report numbers—it interprets them, computes derived quantities (per-expert batch size), and draws the critical conclusion that the problem is structural rather than kernel-level. This understanding would guide all subsequent optimization efforts, ensuring they target the right bottleneck rather than chasing red herrings.

The message also demonstrates the value of systematic microbenchmarking in ML systems optimization. Without the FP4 GEMM benchmark data, the team might have continued chasing communication optimizations or kernel tuning, making marginal gains while missing the fundamental structural issue. By measuring the actual kernel performance across the relevant size range and connecting it to the model's architecture, the assistant identified the true bottleneck and set the stage for the optimization work that followed.