Peeling the Kernel Onion: How a Single Benchmark Revealed the FP4 GEMM Bottleneck on Blackwell SM120

Introduction

In the long arc of optimizing large language model inference, there are moments when a single experiment crystallizes an entire performance problem into a single, undeniable number. Message 868 of this opencode session is one such moment. The assistant, having already established that the GLM-5-NVFP4 model is compute-bound rather than communication-bound, reaches for a deeper understanding of why the GPUs are underperforming. The message contains a single bash command that dispatches a Python script to the remote server, benchmarking FlashInfer's FP4 GEMM (general matrix-matrix multiply) kernel across a sweep of matrix sizes. The output, partially visible at the bottom of the message, tells a devastating story: the CUTLASS FP4 kernels on SM120 (NVIDIA RTX PRO 6000 Blackwell) plateau at approximately 1,200–1,300 TFLOPS — only 60–66% of the theoretical peak — and only reach that plateau for impractically large matrices. For the small matrix sizes that dominate real Mixture-of-Experts (MoE) inference, the kernels achieve barely 1–3% of peak throughput. This message is the moment the assistant quantifies the gap between hardware potential and actual delivered performance, and it sets the stage for every optimization that follows.

The Chain of Reasoning That Led to This Benchmark

To understand why message 868 exists, one must trace the reasoning chain that preceded it. The session had been investigating why the GLM-5-NVFP4 model — a 256-expert MoE model running on 8× RTX PRO 6000 Blackwell GPUs — was not achieving expected throughput. Earlier in the segment, the assistant benchmarked the model in two configurations: TP8 (tensor parallelism across all 8 GPUs) and TP4+PP2 (4-way tensor parallelism with 2-way pipeline parallelism). The results were stark: TP4+PP2 was roughly 2× slower than TP8 ([msg 857]). The user immediately recognized the implication ([msg 858]): if communication were the bottleneck, halving the allreduce participants should have improved throughput, not halved it. The fact that throughput dropped proportionally to the increase in per-GPU compute load pointed squarely at compute-bound execution.

The assistant seized on this insight ([msg 859]) and launched a multi-pronged investigation. Three research subagents were dispatched to investigate SM120 FP4 kernel support, CUTLASS autotune results, and SM120 tensor core specifications ([msg 861]). The findings were alarming: FlashInfer's CUTLASS FP4 backend was reported to return zeros on SM120 (though the assistant later disproved this with correct API usage), the two largest CUTLASS tile configurations failed during autotune due to shared memory overflow, and SM120's tensor cores turned out to be fundamentally different from SM100's — using Ampere-era mma.sync instructions rather than the new tcgen05.mma autonomous tensor units ([msg 862]).

A series of quick experiments followed ([msg 863]–[msg 867]) to verify that the FP4 GEMMs were actually producing correct results. The assistant discovered that the CUTLASS backend did produce non-zero output when given proper arguments (alpha and out tensors), and proceeded to benchmark the kernel at sizes representative of GLM-5 inference with TP8: attention projection matrices of shape (batch, 768, 6144) and MoE up-projection matrices of shape (batch, 2048, 6144). The results from that first benchmark (message 867) showed that at batch=1, the kernel achieved a paltry 0.30 TFLOPS for attention projections and 0.82 TFLOPS for MoE projections. Even at batch=512, it reached only 161 TFLOPS for attention and 437 TFLOPS for MoE — still far below the estimated ~2,000 TOPS peak.

This is the precise moment message 868 enters. The assistant has seen that small-batch performance is terrible, but doesn't yet know where the kernel saturates. "This is very revealing," the assistant writes. "Let me add larger sizes to see where we saturate."

The Benchmark Design: Probing the Saturation Point

The Python script in message 868 is carefully designed to answer a specific question: at what matrix size does the FP4 GEMM kernel stop being memory-bound and become compute-bound? The assistant selects sizes that span from moderately large (2048×2048×6144) to enormous (16384×8192×8192), covering both attention-like shapes (where N is large, matching the hidden dimension) and MoE-like shapes (where N is the intermediate size).

The choice of sizes reveals the assistant's mental model. The GLM-5 model has hidden_size=6144 and moe_intermediate_size=2048. With TP8, each GPU handles 1/8th of these dimensions, so the effective per-GPU GEMM sizes are (batch, 768, 6144) for attention projections and (batch, 256, 6144) for MoE projections — but the assistant's earlier benchmark already covered those. Now the assistant wants to understand the peak capability of the kernel, so it tests sizes that are 2×, 4×, and 8× larger than anything the model would actually use. This is a diagnostic benchmark, not a realistic one.

The script also computes arithmetic intensity (ops/byte), which is a critical metric for understanding whether a kernel is compute-bound or memory-bandwidth-bound. The formula bytes_loaded = (M * K // 2) + (K * N // 2) + (M * K // 16) + (K * N // 16) + (M * N * 2) accounts for the FP4 input data (packed as uint8, hence the //2), the FP8 scale factors (one per 16 elements, hence //16), and the BF16 output. With the RTX PRO 6000's 1,597 GB/s memory bandwidth and ~2,000 TOPS compute peak, the ops:byte ratio needed for compute-bound execution is approximately 1,250 — anything below that is memory-bandwidth-bound.

The benchmark uses 50 iterations per size with 3 warmup iterations, calls torch.cuda.synchronize() to ensure accurate timing, and reports both TFLOPS and arithmetic intensity. This is a well-structured microbenchmark, reflecting the assistant's understanding that kernel-level measurements require careful synchronization and warmup to avoid cold-start effects.

The Results: A Tale of Two Regimes

The output visible at the bottom of message 868 reveals the kernel's behavior across the size sweep:

| M | N | K | TFLOPS | AI | Description | |---|---|---|---|---|---| | 2,048 | 2,048 | 6,144 | 840.2 | 2,286 | MoE up 2K batch | | 4,096 | 2,048 | 6,144 | 1,188.9 | 2,712 | MoE up 4K batch | | 8,192 | 2,048 | 6,144 | 1,221.2 | 2,990 | MoE up 8K batch | | 1,024 | 6,144 | 6,144 | 904.3 | 2,070 | Attn QKV 1K batch | | 2,048 | 6,144 | 6,144 | 1,121.1 | 2,891 | Attn QKV 2K batch | | 4,096 | 6,144 | 6,144 | 1,227.2 | 3,608 | Attn QKV 4K batch |

The pattern is clear. At M=2,048, the kernel achieves 840 TFLOPS (42% of peak). At M=4,096, it jumps to 1,189 TFLOPS (59%). At M=8,192, it plateaus at 1,221 TFLOPS (61%). The attention-like shapes follow a similar trajectory: M=1,024 reaches 904 TFLOPS, M=2,048 reaches 1,121 TFLOPS, and M=4,096 reaches 1,227 TFLOPS.

The plateau at approximately 1,200–1,300 TFLOPS — about 60–66% of the estimated 2,000 TOPS peak — is the key finding. The kernel is not achieving full utilization even at the largest sizes. This could be due to several factors: the reduced shared memory on SM120 (100KB vs 228KB on SM100) forcing smaller tile sizes, the use of mma.sync instructions instead of tcgen05.mma, or simply that the CUTLASS kernel templates haven't been fully tuned for SM120.

But the truly devastating insight comes from the arithmetic intensity numbers. All the tested sizes have AI values well above 2,000 ops/byte — far above the ~1,250 ops/byte breakeven point. This means these kernels are firmly in the compute-bound regime. The problem is not memory bandwidth; it's that the compute units themselves are not being fully utilized.

The Critical Connection to Real Inference

The full impact of these numbers only becomes clear when connected to the actual inference workload. In the subsequent message ([msg 869]), the assistant performs this connection explicitly. During MoE inference, each 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 approximately 512 × 8 / 256 = 16 tokens on average. This puts the effective M dimension at 16 — squarely in the regime where the kernel achieves 0.3–55 TFLOPS (0.02–3% of peak).

This is the fundamental architectural tension at the heart of the GLM-5 optimization problem. The model has 256 experts to achieve strong sparsity and high model quality, but the price is that tokens are scattered so thinly across experts that the GEMM kernels can never reach their efficient operating regime. The GPUs are starving for work — not because they lack compute capacity, but because the work arrives in portions too small to feed the compute units.

Assumptions Embedded in the Benchmark

The assistant makes several assumptions in this benchmark that deserve scrutiny. First, the theoretical peak of ~2,000 TOPS for FP4 on SM120 is an estimate. The actual peak depends on the specific tensor core configuration, clock speeds, and whether the kernel can exploit second-order sparsity (which SM120 supports). The assistant's research subagent had estimated this value based on the 188 SMs and Blackwell's FP4 throughput characteristics, but the true peak may be higher or lower.

Second, the benchmark assumes that the CUTLASS backend is representative of the best available FP4 implementation. The assistant later checks cuBLASLt and sgl_kernel to see if alternatives exist ([msg 869]), finding that sgl_kernel does offer cutlass_fp4_group_mm and cutlass_scaled_fp4_mm functions, but the CUTLASS path through FlashInfer is the one actually used by the SGLang runtime.

Third, the benchmark measures isolated GEMM operations with random data, not the full computational graph of a transformer layer. In real inference, FP4 GEMMs are interleaved with other operations (attention, activation functions, allreduce communication), and the GPU may be able to overlap computation with communication. The benchmark captures only the raw kernel throughput, not the end-to-end pipeline efficiency.

Fourth, the benchmark uses a single GPU (cuda:0) and does not account for the effects of tensor parallelism. In TP8 mode, each GPU processes smaller matrices, which would further reduce the effective M dimension and worsen the utilization problem. The benchmark's sizes are chosen to represent per-GPU work, but the interaction between parallelism strategy and kernel efficiency is complex.

Knowledge Created and Its Impact

Message 868 creates concrete, quantitative knowledge that fundamentally reshapes the optimization strategy. Before this benchmark, the team knew the model was compute-bound but didn't know why. The benchmark reveals that the root cause is not a simple matter of kernel tuning — it's a structural mismatch between the MoE architecture's small per-expert batch sizes and the FP4 GEMM kernel's need for large matrices to achieve high utilization.

This knowledge directly informs the optimization approaches explored later in the segment. The assistant investigates expert parallelism (which would increase per-expert batch sizes by grouping experts on fewer GPUs), piecewise CUDA graphs (to reduce kernel launch overhead for small GEMMs), persistent grouped GEMM kernels (to amortize overhead across multiple small matrix multiplications), and FP4 structured sparsity (to reduce the computational demand). Each of these approaches can be understood as a response to the fundamental problem revealed in message 868: the kernels need larger batches, and the optimization strategies are all different ways of achieving that.

The benchmark also establishes a baseline for measuring improvement. When the assistant later achieves a 28% throughput improvement through server parameter tuning ([msg 912]), that improvement is measured against a baseline that includes the kernel inefficiency quantified here. Any future kernel-level optimization — whether through better CUTLASS tactics, custom CUDA kernels, or cuBLASLt integration — can be evaluated against the 1,200–1,300 TFLOPS plateau identified in this message.

Conclusion

Message 868 is a masterclass in diagnostic benchmarking. It takes a complex performance question — why are 8× RTX PRO 6000 Blackwell GPUs underperforming on GLM-5 inference? — and reduces it to a single, measurable phenomenon: the FP4 GEMM kernel achieves only 60–66% of theoretical peak even at large sizes, and real-world batch sizes are far too small to reach even that plateau. The message is the hinge point of the entire segment: before it, the investigation was about identifying the bottleneck; after it, the investigation becomes about working around it.

The assistant's approach in this message exemplifies the scientific method applied to systems optimization: form a hypothesis (the kernel is not saturating), design an experiment (sweep matrix sizes to find the saturation point), execute it carefully (with warmup, synchronization, and multiple iterations), and interpret the results in the context of the real workload. The output may look like a simple table of numbers, but those numbers encode a deep understanding of the GPU architecture, the MoE model structure, and the fundamental economics of matrix multiplication on modern tensor cores.