Diagnosing the FP4 GEMM Catastrophe: A Probe into SM120 Kernel Support

Introduction

In the middle of an intensive optimization session for the GLM-5-NVFP4 large language model running on eight NVIDIA RTX PRO 6000 Blackwell GPUs, the assistant sent a single bash command that would prove to be a pivotal diagnostic moment. Message [msg 865] is a deceptively simple Python script executed over SSH, designed to test whether FlashInfer's mm_fp4 function — the core FP4 matrix multiplication kernel — actually works on the SM120 microarchitecture. What it reveals is a layered failure: every available backend for FP4 GEMM computation fails on this GPU, yet the model somehow produces coherent text. This contradiction becomes the central mystery that drives the next phase of investigation.

The Context: A Performance Puzzle

The story leading up to this message is one of careful hypothesis testing. The assistant had been running the GLM-5-NVFP4 model — a Mixture-of-Experts architecture with 256 experts, quantized to NVIDIA's FP4 format — across eight RTX PRO 6000 Blackwell GPUs. Earlier in the session ([msg 855]–[msg 857]), the assistant benchmarked two different parallelization strategies: Tensor Parallelism across all 8 GPUs (TP8) and a hybrid TP4+PP2 configuration (4-way tensor parallelism with 2-way pipeline parallelism). The results were stark: TP4+PP2 achieved roughly half the throughput of TP8 at every concurrency level.

The user immediately identified the likely cause ([msg 858]): "Doesn't the halving of perf sort of imply we're compute bound -> kernels not tuned to the GPU?" The assistant agreed and elaborated ([msg 859]), reasoning that if the bottleneck were allreduce communication, reducing the allreduce participants from 8 to 4 while keeping them NUMA-local should have improved throughput. Instead, throughput halved, directly correlating with the doubling of per-GPU GEMM sizes (each GPU now processes N/4 instead of N/8 of each layer). This pointed squarely at compute-bound execution with poorly tuned kernels.

To investigate, the assistant launched three parallel research tasks ([msg 861]): one to survey SM120 FP4 GEMM kernel support in CUTLASS, one to check FlashInfer's autotune results, and one to research SM120 tensor core specifications. The findings were alarming. The research revealed that FlashInfer had a recently filed bug report ([msg 862]) claiming that the CUTLASS FP4 backend "returns all zeros" on SM120, that cuDNN was unsupported, and that the TRTLLM backend explicitly rejects SM120. If the attention projection FP4 GEMMs were silently returning zeros, the model's output would be garbage — yet the model was generating coherent text. This contradiction demanded direct verification.

The Message: A Diagnostic Probe

Message [msg 865] contains a single bash command that executes a Python script on the remote machine. The script constructs FP4 matrix multiplication inputs with carefully chosen dimensions (M=128, N=256, K=256) and iterates over all three available backends — cutlass, cudnn, and trtllm — printing the result or error for each.

The key technical detail is how the script handles the FP4 scale tensors. FP4 quantization uses block-wise scaling: every group of 16 elements shares a single FP8 scale factor. The previous attempt ([msg 864]) had used torch.randn with dtype=torch.float8_e4m3fn, which crashed because PyTorch's normal_kernel_cuda doesn't support that dtype. The subject message fixes this by creating uint8 tensors with torch.randint(1, 127, ...) and then viewing them as float8_e4m3fn via .view(torch.float8_e4m3fn). This is a clever workaround: it bypasses PyTorch's missing random-number support for FP8 while still producing valid FP8 bit patterns in the valid positive range.

The script also omits the alpha and out parameters that the mm_fp4 API signature ([msg 863]) showed as optional. This turns out to be significant.

The Results: A Triple Failure

The output is devastating:

cutlass   : ERROR: TypeError: Mismatched type on argument #4 when calling: `fp4_gemm(...)`. Expected `DLTensor*` but got `None`
cudnn     : ERROR: RuntimeError: detail::finalize(matmul_operation.get_raw_desc()) failed with message: ptrDesc->finalize(), and code: CUDNN_STATUS_BAD_PARAM
trtllm    : ERROR: BackendSuppor...

All three backends fail. The CUTLASS backend fails because argument #4 (the alpha tensor) is None — the function requires an alpha parameter but the script didn't provide one. The cuDNN backend fails with a bad parameter error. The TRTLLM backend fails with a support error, consistent with the research finding that it explicitly rejects SM120.

This is a critical moment. The research suggested CUTLASS might be returning zeros; instead, it's outright crashing. The model is somehow running despite the FP4 GEMM function being completely non-functional in this test. The contradiction deepens.

Assumptions and Their Consequences

The assistant made several assumptions in crafting this test. First, it assumed that the mm_fp4 function could be called with default parameters for alpha and out, based on the API signature showing them as optional. In practice, the CUTLASS backend requires these parameters — the "optional" designation in the type signature doesn't mean the backend will accept None. This assumption was corrected in the very next message ([msg 866]), where providing explicit alpha and out tensors made CUTLASS produce non-zero results.

Second, the assistant assumed that the test dimensions (M=128, N=256, K=256) were representative of the actual GEMM shapes encountered during model inference. In reality, the model's FP4 GEMMs operate on much larger dimensions — the hidden size is 6144, and the MoE intermediate size is 2048. The test dimensions were chosen for quick verification, not for performance profiling.

Third, the assistant implicitly assumed that a failure in the standalone mm_fp4 test would correspond to a failure in the model's inference path. The fact that the model works despite this test failing suggests that either: (a) the model uses a different code path for FP4 GEMMs (perhaps through SGLang's own kernel wrappers or through a different backend selection), (b) the FP4 GEMMs in the attention path are handled differently from those in the MoE path, or (c) the model is falling back to a different precision or kernel implementation at runtime.

Input Knowledge Required

To fully understand this message, one needs knowledge of several domains. The FP4 quantization format — where two 4-bit values are packed into each byte, with block-wise FP8 scale factors — is essential for interpreting the tensor shapes (K//2 for packed data, K//16 for scales with block_size=16). The SM120 architecture (NVIDIA's Blackwell desktop/workstation microarchitecture, compute capability 12.0) and its relationship to SM100 (the datacenter Blackwell) is crucial context: SM120 lacks the autonomous tensor units (tcgen05.mma) that SM100 has, and instead uses older mma.sync instructions with FP4 data types bolted on. The FlashInfer library's architecture — with its multiple backends (CUTLASS, cuDNN, TRTLLM) for FP4 GEMM — and the recently filed bug report about CUTLASS returning zeros on SM120 provide the immediate motivation for the test.

Output Knowledge Created

This message produces several important pieces of knowledge. It empirically confirms that the CUTLASS FP4 backend crashes (rather than returning zeros) when called without explicit alpha and out tensors on SM120. It confirms that cuDNN's FP4 support is non-functional on SM120 with the current cuDNN version (9.1). It confirms that the TRTLLM backend explicitly rejects SM120. Most importantly, it establishes a contradiction — the model works despite the FP4 GEMM function failing — that drives the next phase of investigation.

The message also implicitly validates the workaround for creating FP8 scale tensors (using uint8 with .view()), which becomes the standard pattern for subsequent FP4 benchmarking scripts.

The Thinking Process

The reasoning visible in this message is diagnostic and systematic. The assistant had received research results suggesting that CUTLASS FP4 "returns all zeros" on SM120, but the model was generating coherent text. The natural next step is to verify the claim directly. The assistant constructs a minimal test — small dimensions for speed, all three backends for completeness, and a simple pass/fail metric (checking all_zeros and printing abs_max/abs_mean).

The progression from [msg 862] (first test with wrong dimensions) to [msg 863] (checking API signature) to [msg 864] (correct dimensions but wrong dtype) to [msg 865] (correct dimensions and dtype but missing alpha/out) shows a classic debugging workflow: each failure narrows the problem space. The assistant is iterating toward a working test, and each error message provides information about the API contract.

The choice to test all three backends in a single loop reflects an understanding that the backend landscape is fragmented — different backends have different SM120 support levels, and knowing which (if any) works is essential for diagnosing the model's behavior.

The Broader Significance

This message sits at a critical juncture in the optimization effort. The assistant had spent hours setting up infrastructure, benchmarking configurations, and researching kernel support. The TP4+PP2 experiment had ruled out communication bottlenecks. The research tasks had raised alarm bells about FP4 kernel quality. This test is the first direct empirical check of those alarms. Its failure — combined with the model's apparent functionality — forces a reassessment: either the FP4 GEMM path in the model is different from the standalone mm_fp4 function, or the model is using a fallback that avoids FP4 GEMMs entirely.

The follow-up message ([msg 866]) resolves the immediate issue by providing alpha and out tensors, showing that CUTLASS does produce non-zero results. But the deeper question — why do all three backends fail in the default configuration, and what path does the model actually use — remains open. This tension between the test results and the model's behavior drives the subsequent investigation into cuBLASLt, SGLang's internal kernel wrappers, and the detailed FP4 GEMM benchmarking that reveals the true performance characteristics of SM120.

Conclusion

Message [msg 865] is a textbook example of a diagnostic probe in systems optimization: a focused, minimal test designed to verify a specific claim (FP4 GEMMs return zeros on SM120) that challenges an observed phenomenon (the model generates coherent text). Its triple failure is more informative than a success would have been — it reveals that the FP4 kernel ecosystem on SM120 is fragmented and unreliable, that the API contracts are not what they appear from type signatures alone, and that the model's inference path must be using a different mechanism than the one being tested. The message's true value lies not in the answers it provides but in the questions it forces: if the FP4 GEMMs don't work in isolation, how does the model actually compute? The answer to that question would lead to deeper investigations into kernel selection, fallback paths, and ultimately to the discovery that FP4 GEMMs do work when called correctly — but with far lower efficiency than theoretical peaks would suggest.