The Art of Reading the Manual: Debugging FP4 GEMM Kernels on Blackwell SM120

In the middle of a high-stakes optimization sprint for the GLM-5-NVFP4 model on 8× NVIDIA RTX PRO 6000 Blackwell GPUs, the assistant encounters a moment of quiet methodological discipline. Message <msg id=863> is deceptively simple — a single bash command that runs a Python help() call on a remote server. But this message represents a critical turning point in the investigation: the moment when the assistant stops guessing and starts reading the source of truth.

The Message

[bash] ssh root@10.1.230.174 'source /root/ml-env/bin/activate && python3 -c "
import torch
import flashinfer

# Test mm_fp4 on SM120 - check the correct API signature
help(flashinfer.mm_fp4)
" 2>&1 | head -40'

The output reveals the function signature:

mm_fp4(a: torch.Tensor, b: torch.Tensor, a_descale: torch.Tensor, b_descale: torch.Tensor, 
       alpha: Optional[torch.Tensor] = None, out_dtype: torch.dtype = torch.bfloat16, 
       out: Optional[torch.Tensor] = None, block_size: int = 16, 
       use_8x4_sf_layout: bool = False, 
       backend: Literal['cudnn', 'trtllm', 'cutlass', 'auto'] = 'auto', 
       use_nvfp4: bool = True) -> torch.Tensor

Why This Message Was Written

The immediate trigger is a failed experiment. In the preceding message <msg id=862>, the assistant attempted to test whether FlashInfer's mm_fp4 function was silently returning zeros on SM120 — a critical concern raised by a recently filed FlashInfer bug report (issue #2577). That test crashed with a cryptic error:

CUTLASS error: K dimension mismatch in mm_fp4. got a.shape[1] = 128, b.shape[0] = 256

The assistant had constructed the FP4 input tensors manually: a_fp4 of shape (M, K//2) and b_fp4 of shape (N, K//2), with scale tensors scale_a of shape (M, K//64) and scale_b of shape (N, K//64). But the API rejected these dimensions. Rather than continuing to guess at the correct tensor shapes — which could lead to endless trial-and-error or, worse, silently producing incorrect results that look plausible — the assistant chose to step back and consult the authoritative API documentation via Python's built-in help() function.

This decision reflects a deeper understanding of the debugging process. When a numerical library rejects your tensor shapes, there are at least three possible explanations: (1) you constructed the tensors incorrectly, (2) the library's dimension-checking logic has a bug, or (3) you're using the wrong API entirely. The fastest way to disambiguate these cases is to verify the exact parameter semantics from the source. The help() output provides the full function signature, parameter names, types, and defaults — everything needed to construct a valid call.

The Broader Context: A Crisis of Confidence in FP4 Kernels

To understand why this single API lookup matters, we need to appreciate the crisis that precipitated it. The assistant had spent the previous several hours benchmarking the GLM-5-NVFP4 model across different parallelism configurations. The results were alarming:

| Configuration | 256 Concurrency | 512 Concurrency | |--------------|----------------|-----------------| | TP8 | 1,867 tok/s | 2,800 tok/s | | TP4+PP2 | 727 tok/s | 1,318 tok/s |

The TP4+PP2 configuration was roughly 2× slower than TP8. The user astutely observed in <msg id=858> that this halving of performance implied the model was compute-bound, not communication-bound — because reducing the allreduce participants from 8 to 4 should have helped if communication were the bottleneck. Instead, throughput dropped proportionally to the increase in per-GPU matrix sizes (N/4 instead of N/8 per layer).

This observation triggered a deep investigation into FP4 GEMM kernel efficiency on SM120 (Blackwell's workstation microarchitecture). The research agents dispatched in <msg id=861> uncovered a cascade of bad news:

  1. SM120 tensor cores are fundamentally different from SM100 (B200/B100). Despite both being "Blackwell," SM120 uses Ampere-era mma.sync instructions rather than SM100's autonomous tcgen05.mma tensor units. FP4 is bolted onto an older programming model, yielding 2-4× lower per-SM throughput.
  2. FlashInfer bug #2577 (filed just the day before, on February 18, 2026) reported that the cutlass backend for mm_fp4 was silently returning all zeros on SM120. The cudnn backend was unsupported, and the trtllm backend explicitly rejected SM120.
  3. MoE autotune was losing its best tactics. The two largest CUTLASS tile configurations (M128×256×128 and M256×128×128) failed during autotune with "Error Internal" — they overflowed SM120's 100 KB shared memory limit (compared to SM100's larger shared memory). The surviving tactics used smaller tiles with worse latency hiding.
  4. The GPUs were power-starved. During inference, the RTX PRO 6000 GPUs drew only ~235 W out of their 600 W TDP, with SM occupancy at 98% but power at 55%. This classic sign of a memory-bound workload meant the SMs were stalled waiting on data, not doing useful FP4 math.
  5. Memory bandwidth was the ultimate ceiling. With 1.6 TB/s GDDR7 versus B200's 8 TB/s HBM3e, the ops:byte ratio needed for compute-bound execution was ~1,250 — far beyond what small-batch decode could achieve. The assistant's attempted mm_fp4 test in <msg id=862> was meant to verify whether the attention projection FP4 GEMMs were actually producing correct results. If they were silently returning zeros, the model's outputs might be degraded or entirely garbage — yet the benchmark numbers looked reasonable, suggesting either the MoE path used different kernels or there was a fallback mechanism. The dimension mismatch error cut that investigation short, leading directly to the API documentation lookup in <msg id=863>.

Assumptions and Their Consequences

The assistant made several assumptions in constructing the failed test. First, it assumed that mm_fp4 accepted packed FP4 tensors with the same layout as the model's checkpoint format — where each pair of FP4 values (two 4-bit numbers) is packed into a single uint8 byte, reducing the first dimension from K to K//2. This is a reasonable assumption given that NVFP4 (NVIDIA's 4-bit floating point format) commonly uses this packing scheme.

Second, the assistant assumed the scale tensors should have shape (M, K//64) based on a block size of 64 elements per shared scale factor — a common pattern in block-scaled quantization where each block of 64 values shares one scale. The API documentation reveals that the default block_size is 16, not 64, which would change the expected scale dimensions.

Third, the assistant assumed that mm_fp4 would accept the tensors in a particular order — first the packed data, then the scales. The actual API signature shows that a_descale and b_descale are separate named parameters, and the function also accepts an optional alpha parameter for fused scaling.

The key mistake was not checking the API signature before constructing the test tensors. A few minutes spent reading the documentation would have saved the failed experiment and provided the correct tensor shapes immediately. This is a classic debugging pitfall: when a library function fails, the natural instinct is to tweak the inputs and retry, rather than consulting the documentation to verify fundamental assumptions about the API contract.

Input and Output Knowledge

The input knowledge required to understand this message includes:

The Thinking Process Visible in the Message

The reasoning in this message is subtle but important. The assistant's comment — "Test mm_fp4 on SM120 - check the correct API signature" — reveals a deliberate methodological choice. Rather than continuing to debug the dimension mismatch error by trial and error (which could involve guessing different tensor shapes, reading error messages, and iterating), the assistant chooses to go to the source: Python's help() function, which displays the function's docstring.

This is particularly notable because the assistant is working remotely via SSH, adding friction to every command. The decision to run help() and pipe the output through head -40 shows an awareness of both the value of the information and the cost of obtaining it. The head -40 truncation is pragmatic — the full docstring might be hundreds of lines, but the function signature and parameter descriptions (typically the first 30-40 lines) contain the critical information needed.

The choice to use help() rather than reading the source code directly (e.g., via grep or by opening the Python file) is also telling. help() provides the documented API, which may differ from the implementation. For a debugging scenario where the question is "what shapes does this function expect?", the docstring is the authoritative answer — it represents the contract the function promises to uphold.

The Path Forward

With the correct API signature in hand, the assistant can now construct a valid mm_fp4 call. The key insights are: