The FP4 GEMM Kernel Reality Check: When Theoretical Peak Collides with SM120's Hardware Limits

In the high-stakes world of large language model inference, the difference between a model that "works" and one that performs can be measured in orders of magnitude. Message [msg 867] of this opencode session represents a critical inflection point — a moment when the assistant moves from high-level configuration benchmarking to a deep, systematic investigation of the actual computational efficiency of the underlying kernels. This message, delivered by the AI assistant during a complex ML deployment session on 8× NVIDIA RTX PRO 6000 Blackwell GPUs, is a masterclass in diagnostic reasoning: it confirms that the FP4 GEMM kernels are functional (dispelling a recently filed bug report), then immediately pivots to quantify exactly how far from peak performance they actually run.

The Context: A Performance Mystery Unfolds

To understand why message 867 was written, we must trace the thread of reasoning that led to it. The session had been deploying the GLM-5-NVFP4 model — a large Mixture-of-Experts (MoE) transformer using NVIDIA's NVFP4 quantization format — across eight RTX PRO 6000 Blackwell GPUs. The assistant had been experimenting with different parallelism configurations, comparing Tensor Parallelism with 8 GPUs (TP8) against a hybrid Tensor Parallelism with 4 GPUs plus Pipeline Parallelism with 2 stages (TP4+PP2).

The results were puzzling. At 256 concurrent requests, TP8 delivered 1,867 total tokens per second. TP4+PP2 delivered only 727 — roughly 2.4× slower. At 512 concurrency, TP8 hit 2,800 tok/s while TP4+PP2 managed just 1,318. The user, in message [msg 858], made a sharp observation: "Doesn't the halving of perf sort of imply we're compute bound -> kernels not tuned to the GPU?"

The assistant immediately recognized the insight. In TP4+PP2, each GPU computes 1/4th of each layer (instead of 1/8th in TP8), meaning the per-layer GEMM sizes double. If the bottleneck were communication (allreduce latency), reducing the allreduce participants from 8 to 4 — and keeping them NUMA-local — should have improved throughput. Instead, throughput roughly halved, directly correlating with the 2× larger GEMMs per GPU. This is the classic signature of a compute-bound system: larger matrix multiplications take proportionally longer, and communication overhead is not the dominant factor.

This realization triggered a multi-pronged investigation. The assistant launched three parallel research tasks (visible in message [msg 861]): one to research SM120 FP4 GEMM kernel support, one to check CUTLASS autotune results, and one to research SM120 tensor core specifications. The results, returned in subsequent messages, painted a grim picture: SM120 (the Blackwell workstation/desktop microarchitecture) uses fundamentally different tensor core instructions than SM100 (the datacenter Blackwell). SM120 relies on Ampere-era mma.sync warp-level instructions rather than SM100's autonomous tcgen05.mma tensor units, resulting in 2-4× lower per-SM FP4 throughput. A FlashInfer bug report (#2577, filed just the day before) claimed that the CUTLASS FP4 backend "returns all zeros" on SM120, while the cudnn backend is unsupported and trtllm explicitly rejects the architecture. The autotune was failing on the two largest tile shapes (M128 and M256) due to SM120's 99KB shared memory limit. And the memory bandwidth — 1.6 TB/s GDDR7 versus B200's 8 TB/s HBM3e — meant the ops:byte ratio was fundamentally constrained.

Message 867: The Kernel Reality Check

Message 867 opens with a critical correction. The assistant had been alarmed by the FlashInfer bug report suggesting CUTLASS FP4 returns zeros on SM120. But after fixing the test harness — providing proper alpha and out tensors — the CUTLASS backend produces non-zero results with reasonable magnitudes:

OK — so CUTLASS does produce non-zero results when given proper inputs with alpha and out tensors. The auto backend tries cudnn first and fails. The bug report #2577 may have been about a different setup or an older version.

This is an important moment of intellectual honesty. The assistant had just spent significant effort researching a potential show-stopping bug, only to discover the issue was likely a test configuration problem rather than a genuine kernel defect. Rather than dwelling on the dead end, the assistant immediately reframes the question: "The model IS generating coherent text, so the FP4 GEMMs are working to some degree. The question is whether they're running at peak efficiency."

This reframing is the core of message 867. The assistant decides to benchmark raw FP4 GEMM speed across a range of matrix sizes that are representative of the GLM-5 model's actual operations. The benchmark is meticulously designed:

Matrix size selection: The assistant identifies the key dimensions of GLM-5 — hidden_size=6144, moe_intermediate_size=2048 — and computes the per-GPU sizes under TP8: hidden/8=768, intermediate/8=256. This demonstrates a deep understanding of how model parallelism maps matrix dimensions across GPUs.

Operation coverage: The benchmark covers both attention projection GEMMs (M×768×6144) and MoE expert up-projection GEMMs (M×2048×6144), across batch sizes from 1 to 1024. This spans the full range from single-token decode (batch=1) to high-throughput serving (batch=1024).

Measurement methodology: The assistant uses 5 warmup iterations followed by 100 timed iterations with CUDA synchronization, measuring microseconds per call and computing achieved TFLOPS using the formula 2 * M * N * K (accounting for multiply-add operations).

The results are devastatingly clear:

| M | N | K | Time | TFLOPS | Description | |---|---|---|---|---|---| | 1 | 768 | 6144 | 31.9 μs | 0.30 | Single token attn proj | | 64 | 768 | 6144 | 38.9 μs | 15.53 | Batch=64 attn proj | | 256 | 768 | 6144 | 29.7 μs | 81.33 | Batch=256 attn proj | | 512 | 768 | 6144 | 30.0 μs | 160.92 | Batch=512 attn proj | | 1 | 2048 | 6144 | 30.5 μs | 0.82 | Single token MoE up |

At single-token decode — the most common operating point for interactive serving — the FP4 GEMM achieves a mere 0.30 TFLOPS. The theoretical peak for FP4 on SM120 is approximately 2,000 TFLOPS. That's 0.015% of theoretical peak. Even at batch=512, the attention projection reaches 160.92 TFLOPS — still only about 8% of peak.

The Reasoning Behind the Benchmark

The assistant's decision to run this specific benchmark reveals a sophisticated diagnostic strategy. Rather than continuing to tweak server parameters (like --max-running-requests or --num-continuous-decode-steps) and hoping for incremental improvements, the assistant recognized that the fundamental question was whether the kernels themselves were efficient. This is a textbook example of moving up the abstraction hierarchy: when configuration tuning hits diminishing returns, you must verify that the underlying computational primitives are performing as expected.

The benchmark design also reflects an understanding of the MoE architecture's unique characteristics. In a dense transformer, all GEMMs operate on the full hidden dimension. In an MoE transformer, each token is routed to a subset of experts, and each expert processes its assigned tokens independently. This means the effective batch size per expert is much smaller than the global batch size. With 64 experts in GLM-5 and 8-way tensor parallelism, a global batch of 256 tokens might result in per-expert batches of just 4-16 tokens. The benchmark shows that at these small sizes, the FP4 kernel achieves only 0.8-15 TFLOPS — a tiny fraction of the GPU's capability.

Assumptions and Their Validation

Several assumptions underpin this message, and the assistant implicitly validates them through the benchmark design:

Assumption 1: The model is compute-bound, not communication-bound. This was established in the previous message ([msg 859]) through the TP4+PP2 comparison. The benchmark results in message 867 provide the micro-level confirmation: if the kernels were running at high efficiency, the compute-bound hypothesis would be contradicted by the low power utilization (~235W out of 600W TDP) observed earlier. The fact that the kernels are running at 0.015-8% of peak explains both the low power draw and the poor throughput.

Assumption 2: The CUTLASS backend is the correct path to test. The assistant tests only the CUTLASS backend, not cuBLASLt or other alternatives. This is justified by the earlier research showing that cudnn is unsupported on SM120 and trtllm explicitly rejects it. CUTLASS is the only viable option.

Assumption 3: The benchmark matrix sizes are representative. The assistant assumes that the GLM-5 model's actual GEMM operations match the dimensions tested. This is a reasonable assumption given the model architecture, but it's worth noting that the actual operations may include additional transformations (e.g., reshaping, transposition) that could affect performance.

Assumption 4: The mm_fp4 function is the correct API for the model's FP4 GEMMs. The assistant tests flashinfer.mm_fp4 specifically, which handles NVFP4-formatted matrices. If the model uses a different FP4 variant or a different kernel path (e.g., through the MoE routing mechanism), the benchmark might not reflect actual inference performance.

Input Knowledge Required

To fully understand message 867, the reader needs:

  1. Understanding of tensor parallelism (TP): How splitting the hidden dimension across GPUs reduces per-GPU matrix sizes. TP8 means each GPU handles 1/8th of the hidden dimension; TP4 means 1/4th.
  2. Knowledge of FP4 quantization: NVFP4 is a 4-bit floating point format that packs two 4-bit values into one byte. The packed format explains why the input matrices have shape (M, K//2) instead of (M, K) — the K dimension is halved because each byte stores two FP4 values.
  3. Understanding of MoE architecture: Mixture-of-Experts models have multiple "expert" feed-forward networks, with each token routed to a subset of experts. This creates small per-expert batch sizes during decode.
  4. Familiarity with CUTLASS and FlashInfer: CUTLASS is NVIDIA's template library for GEMM kernels; FlashInfer is a library that wraps CUTLASS (and other backends) for inference workloads. The mm_fp4 function is FlashInfer's FP4 matrix multiplication interface.
  5. Knowledge of SM120 vs SM100 differences: The Blackwell architecture has two variants — SM100 (datacenter B200/B100) with autonomous tensor units, and SM120 (workstation/desktop RTX PRO 6000) with older-style tensor cores. This distinction is critical for understanding why FP4 performance differs.
  6. CUDA synchronization and benchmarking methodology: The use of torch.cuda.synchronize() ensures accurate timing by waiting for all GPU operations to complete before measuring elapsed time.

Output Knowledge Created

Message 867 produces several critical pieces of knowledge:

  1. CUTLASS FP4 works on SM120 (with caveats): Despite the bug report, the CUTLASS backend produces non-zero, reasonable-looking results when called with the correct API signature (including alpha and out tensors). This is a valuable data point for the FlashInfer community.
  2. Quantitative efficiency baseline: The benchmark establishes precise TFLOPS measurements for each matrix size, creating a baseline that can be compared against theoretical peak and against future optimizations.
  3. Confirmation of compute-bound hypothesis at the micro level: The extremely low TFLOPS at small batch sizes provides the microarchitectural explanation for the poor throughput observed at the system level.
  4. Identification of the critical operating region: The benchmark reveals that single-token decode (batch=1) is the worst case, achieving only 0.30 TFLOPS. This is the most important operating point for interactive applications.
  5. Methodology for kernel benchmarking: The assistant's approach — identifying representative matrix dimensions, running warmup iterations, synchronizing CUDA, and computing TFLOPS from measured time — provides a template for future kernel investigations.

The Thinking Process: A Diagnostic Masterclass

What makes message 867 exceptional is the clarity of its reasoning structure. The assistant follows a logical chain that any performance engineer would recognize:

  1. Hypothesis generation: "The model is compute-bound" (from the TP4+PP2 comparison)
  2. Hypothesis refinement: "The FP4 kernels might be running suboptimally" (from the low power utilization and autotune failures)
  3. Threat validation: "But first, let me check if the kernels even work correctly" (addressing the bug report)
  4. Hypothesis testing: "Let me benchmark the raw FP4 GEMM speed" (the core of message 867)
  5. Result interpretation: The benchmark confirms the hypothesis — kernels are running at 0.015-8% of peak This is the scientific method applied to systems engineering. The assistant doesn't jump to conclusions, doesn't assume the worst (the bug report), and doesn't stop at the system-level observation. It drills down to the microarchitectural level to find the root cause. The message also demonstrates intellectual humility. When the CUTLASS test initially failed, the assistant didn't declare the backend broken — it checked the API signature, fixed the test, and re-ran. When the bug report suggested a catastrophic failure, the assistant verified with actual hardware rather than accepting the report at face value. This willingness to question both external sources and one's own assumptions is a hallmark of expert troubleshooting.

The Broader Implications

Message 867 is not just a diagnostic step — it's a turning point that shapes the entire subsequent optimization effort documented in segment 7. The benchmark results directly motivate the exploration of:

Conclusion

Message 867 represents the moment when a performance investigation transitions from "what" to "why." The assistant had observed that TP4+PP2 was 2× slower than TP8. The user hypothesized compute-bound execution. The assistant confirmed this at the system level, then drilled down to the kernel level to measure exactly how inefficient the FP4 GEMMs were. The benchmark results — 0.30 TFLOPS at single-token decode versus a theoretical peak of ~2,000 TFLOPS — are stark evidence of a fundamental kernel efficiency problem on SM120.

This message is a textbook example of systematic performance debugging: start with a clear hypothesis, validate it at multiple levels of abstraction, measure precisely, and let the data guide the next steps. For anyone working on LLM inference optimization, message 867 demonstrates the critical importance of understanding not just whether your kernels work, but how efficiently they run — and the sometimes shocking gap between theoretical peak and practical reality.