The cuBLASLt FP4 Investigation: A Systematic Evaluation of Blackwell's Matrix Math Frontier

When the promise of a faster matrix multiplication path appears on the horizon, the temptation to adopt it immediately can be strong. The cuBLASLt library's FP4 support on NVIDIA's Blackwell architecture reportedly achieved an eye-popping 99.6 TFLOPS on GB10 SM121 — a figure that naturally sparked intense interest. Could this be the key to unlocking dramatically better inference performance for the GLM-5-NVFP4 model running on the RTX PRO 6000 Blackwell Server Edition?

The answer, as this investigation revealed, was far more nuanced. This article chronicles a deep, systematic evaluation of cuBLASLt FP4 on SM120 hardware, covering everything from initial discovery through rigorous benchmarking to a final verdict that surprised even the researchers conducting it. For a detailed account of the initial probing that kicked off this work, see [21] and [11], which document the earliest attempts to understand cuBLASLt's FP4 capabilities on Blackwell.

The Discovery: cuBLASLt FP4 Is Accessible on SM120

The investigation began with a straightforward question: is cuBLASLt FP4 even available on this system? The answer came quickly. PyTorch 2.9.1+cu128, installed on a machine with two (later eight) RTX PRO 6000 Blackwell GPUs, exposed torch._scaled_mm — the gateway to cuBLASLt's scaled matrix multiplication — and, crucially, supported the torch.float4_e2m1fn_x2 data type for packed FP4 values.

The first successful call was a landmark moment:

out = torch._scaled_mm(
    a_fp4,                          # (M, K//2) as float4_e2m1fn_x2
    b_fp4.t(),                       # (N, K//2) as float4_e2m1fn_x2, transposed
    scale_a=scale_a,                 # (M, K//16) as float8_e4m3fn
    scale_b=scale_b,                 # (N, K//16) as float8_e4m3fn
    out_dtype=torch.bfloat16         # Required!
)

This single line of code confirmed that CUDA_R_4F_E2M1 (enum value 33, defined in CUDA 12.8's library_types.h) was recognized by cuBLASLt's heuristic engine on SM120. The cublasLtMatmulAlgoGetHeuristic call successfully routed to FP4 kernels. This breakthrough moment is detailed in [16], which describes the exact sequence of discoveries that led from initial failure to working FP4 matmul.

But the discovery came with immediate constraints. The out_dtype parameter was mandatory — omitting it triggered a CUBLAS_STATUS_NOT_SUPPORTED error. The use_fast_accum flag was explicitly rejected for FP4 inputs. And most importantly, the M dimension had to be padded to a multiple of 128, a requirement that would prove costly for small-batch inference.

Verifying Correctness: FP4 Encoding and Block Scaling

Before any performance benchmarking could begin, the team needed to verify that cuBLASLt was interpreting FP4 values correctly. NVFP4 E2M1 encoding packs two 4-bit values into each byte, using a sign-exponent-mantissa format (1-2-1) with a bias of 1. The representable values are: 0, ±0.5, ±1.0, ±1.5, ±2.0, ±3.0, ±4.0, and ±6.0.

A series of targeted tests confirmed the encoding:

| Test Pattern | Packed Byte | Expected | Got | |---|---|---|---| | All 1.0 values | 0x22 | 64.0 | 64.0 | | All 0.5 values | 0x11 | 32.0 | 32.0 | | All -1.0 values | 0xAA | -64.0 | -64.0 | | All 6.0 values | 0x77 | 384.0 | 384.0 |

The block scaling mechanism was also validated. With 1x16 blocking (one scale per 16 elements), the scales are stored as float8_e4m3fn values in a plain row-major layout. A mixed-scale test — where the first half of the K dimension used scale 1.0 and the second half used scale 2.0 — produced exactly the expected output of 96.0 for K=64, confirming that cuBLASLt correctly applies per-block scales.

This was a critical finding: cuBLASLt expects plain row-major scales, not the swizzled (interleaved) layout produced by existing FP4 quantizers like flashinfer.fp4_quantize and sgl_kernel.scaled_fp4_quant. The swizzled layout, designed for CUTLASS's tcgen05 MMA scale factor format, is fundamentally incompatible with cuBLASLt's expectations. This scale format mismatch became a central theme of the investigation, explored in depth across multiple articles including [1], [2], and [25], which dissected the swizzle pattern and its implications for integration.

The Performance Benchmark: A Tale of Two Regimes

With correctness verified, the investigation moved to performance. A comprehensive benchmark suite was built, testing cuBLASLt FP4 against FlashInfer's CUTLASS backend and SGLang's own CUTLASS kernel across a range of configurations representative of the GLM-5-NVFP4 model's workload.

The benchmark used CUDA events for precise timing and covered dimensions from small decode batches (M=1) through large prefill configurations (M=2048), with K and N values matching the model's hidden size (7168), intermediate size (4096), and expert dimensions (2048).

The results revealed a clear performance dichotomy:

Small Batch Regime (M=1 to M=64)

For decode-time inference where batch sizes are small, cuBLASLt appeared to dominate — but the victory was illusory. Because cuBLASLt requires M to be padded to 128, a request with M=1 would compute 128× more work than needed. The benchmark reported 426.8 TFLOPS for M=1, but this was achieved by computing 128 rows instead of 1. The actual useful throughput was a fraction of the reported figure.

FlashInfer CUTLASS, which handles arbitrary M dimensions without padding, reported a more honest 2.7 TFLOPS for the same configuration — reflecting the true cost of processing a single token.

Large Batch Regime (M=128 to M=2048)

At larger batch sizes, the picture shifted dramatically. With M ≥ 128, cuBLASLt's padding requirement no longer imposed a penalty, and the comparison became fair:

| Configuration | cuBLASLt (TFLOPS) | FI-CUTLASS (TFLOPS) | sgl-CUTLASS (TFLOPS) | |---|---|---|---| | Attn M=256 | 796.8 | 835.0 | 831.3 | | Attn M=1024 | 1053.3 | 1102.8 | 894.9 | | MLP-up M=2048 | 1186.2 | 1256.0 | 1029.1 |

FlashInfer CUTLASS was consistently 5-8% faster than cuBLASLt at large batch sizes. The peak observed throughput was approximately 1,256 TFLOPS for FI-CUTLASS versus 1,186 TFLOPS for cuBLASLt at M=2048. The vaunted 99.6 TFLOPS claim from GB10 SM121 did not translate to superior performance on SM120. The full benchmark results, including the initial syntax-error-plagued attempts and the pivot to file-based benchmarking, are documented in [19], [28], and [35].

The Quantization Compatibility Wall

Perhaps the most significant obstacle to integration was the scale format incompatibility. Both flashinfer.fp4_quantize and sgl_kernel.scaled_fp4_quant produce scales in a swizzled layout optimized for CUTLASS's tcgen05 MMA instructions. cuBLASLt, by contrast, expects plain row-major scales.

The FP4 packed data itself was compatible — a simple .view(torch.float4_e2m1fn_x2) on the existing uint8 tensor sufficed. But the scales could not be reused directly. Attempting to reshape swizzled scales into the expected (M, K//16) shape produced incorrect results, as the numerical values in each scale position were permuted by the interleaving pattern.

Three potential solutions were identified:

  1. Write a new CUDA kernel for non-swizzled FP4 quantization
  2. Write a de-swizzle kernel to convert existing scales
  3. Compute scales separately using PyTorch operations Option 3 was prototyped and benchmarked. Computing plain scales from raw BF16 data added approximately 0.02ms per call — a 245% overhead on the quantization step itself. While a dedicated CUDA kernel would be faster, the overhead would still eat into any performance advantage cuBLASLt might offer.

The MoE Showstopper

The most definitive limitation emerged when investigating MoE (Mixture of Experts) support. The GLM-5-NVFP4 model uses a Mixture of Experts architecture where each token is routed to a subset of expert networks. Efficient inference requires grouped GEMM — computing multiple small matrix multiplications in a single kernel call.

PyTorch provides torch._scaled_grouped_mm for this purpose, but the function explicitly rejects SM120:

FP4 _scaled_grouped_mm error: torch._scaled_grouped_mm is only supported 
on CUDA devices with compute capability = [9.0, 10.0], or ROCm MI300+

This means cuBLASLt cannot be used for the MoE path — the dominant compute component in models like GLM-5-NVFP4 — on Blackwell GPUs. Each expert would need an individual _scaled_mm call, losing all the batching efficiency that grouped GEMM provides. The existing FlashInfer CUTLASS path, which supports grouped execution natively, remains the only viable option for MoE.

The Verdict: Stay the Course

After days of systematic investigation — spanning API discovery, correctness verification, performance benchmarking, quantization analysis, and integration planning — the verdict was clear:

Do not switch to cuBLASLt FP4 for SM120.

The reasoning rested on four pillars:

  1. Performance parity at best. FlashInfer CUTLASS is 5-8% faster at large batch sizes and comparable elsewhere. The SM121 performance claims do not transfer to SM120.
  2. MoE is blocked. The inability to use grouped GEMM on SM120 means the dominant compute path cannot benefit from cuBLASLt.
  3. Integration cost is prohibitive. A new non-swizzled quantizer, dual scale storage, alpha multiplication logic, and M-padding handling would require substantial engineering effort with no performance payoff.
  4. M=128 padding wastes compute. For the small-batch decode regime that dominates interactive inference, padding to 128 rows means computing 128× more work than needed. The recommendation was to stay with FlashInfer CUTLASS (--fp4-gemm-backend flashinfer_cutlass), which is already well-integrated, performs slightly better on SM120, and handles MoE efficiently through its native grouped execution path. This verdict is explored in greater detail in [43] and [36], which provide comprehensive analyses of the kernel selection decision and the scale format mismatch respectively.

References

  1. [1] "Deciphering the Scale Layout: How a Single Error Message Unlocked cuBLASLt FP4 Matmul on Blackwell GPUs"
  2. [2] "Decoding cuBLASLt's FP4 Scale Layout: A Detective Story in Kernel Integration"
  3. [11] "Peering into the cuBLASLt FP4 Abyss: A Technical Deep-Dive into GPU Kernel Routing"
  4. [16] "The Breakthrough and the Stumble: Discovering cuBLASLt FP4 on Blackwell SM120"
  5. [19] "The Benchmark That Wasn't: A Syntax Error in the Pursuit of FP4 Performance"
  6. [21] "Probing the Frontier: Investigating cuBLASLt FP4 GEMM Support for Blackwell GPUs"
  7. [25] "Decoding the Swizzle: A Deep Dive into cuBLASLt FP4 Scale Layout Compatibility on Blackwell GPUs"
  8. [28] "The Pivot to File-Based Benchmarking: A Case Study in Debugging Strategy"
  9. [35] "The Pivot to File-Based Benchmarking: A Methodological Turning Point in FP4 GEMM Research"
  10. [36] "The cuBLASLt FP4 Integration Investigation: A Deep Dive into Scale Format Mismatches and Kernel Design Decisions"
  11. [43] "The Verdict on cuBLASLt FP4: A Deep Dive into GPU Kernel Selection for Blackwell"

Epilogue: The Value of Systematic Investigation

This investigation stands as a case study in the importance of rigorous, systematic evaluation before adopting new technology. The initial allure of cuBLASLt FP4 — backed by impressive performance claims and the promise of a "better" kernel path — gave way to a nuanced understanding of real-world constraints. The M-padding requirement, the scale format incompatibility, the MoE limitation, and the ultimate performance parity all emerged only through careful, methodical testing.

The 44-message conversation that produced this analysis, along with the dozens of subagent sessions and benchmark scripts, represents a significant investment of computational and human effort. But that investment paid off: it prevented a costly integration that would have delivered no performance benefit while adding substantial complexity.

In the end, the best path forward was the one already in place — a reminder that sometimes the most valuable outcome of research is the confidence to stay the course.