The SM120 FP4 GEMM Kernel Investigation: Uncovering the Shared Memory Crisis on Blackwell Workstation GPUs

Introduction

In the high-stakes world of large language model inference, performance optimization often descends through layers of abstraction — from server configuration to communication topology to kernel microarchitecture. This article examines a concentrated research effort within a broader coding session, where an AI assistant was tasked with diagnosing why FP4 GEMM (General Matrix Multiply) kernels were performing suboptimally on NVIDIA's RTX PRO 6000 Blackwell workstation GPUs. What began as a targeted research question rapidly uncovered a systemic crisis: the software ecosystem for Blackwell's workstation microarchitecture (SM120) was fundamentally broken for the very quantization format the model depended on.

The chunk under analysis (segment 7, chunk 0) captures a subagent session — session ID ses_389 — consisting of seven messages (indices 0–6) in which the assistant systematically investigated FP4 GEMM kernel support on SM120. The messages span the full arc of a research investigation: from the initial user request ([msg 0]), through parallel web searches ([msg 1]), deep dives into documentation and bug reports (<msgs id=2–5>), and culminating in a comprehensive synthesis of findings ([msg 6]). The articles [1]–[7] each analyze individual messages from this session, and this article synthesizes the entire investigation into a coherent narrative.

The Context: A Performance Anomaly on Cutting-Edge Hardware

To understand why this research was necessary, we must first understand the situation that prompted it. The broader session (segments 0–7 of the root session) describes an extensive ML environment setup on Ubuntu 24.04. The team had successfully installed NVIDIA drivers (590.48.01), CUDA Toolkit 13.1, and verified multiple RTX PRO 6000 Blackwell GPUs. After resolving complex build issues for flash-attn and related dependencies, the GLM-5-NVFP4 model — a large language model using NVIDIA's NVFP4 quantization format (4-bit floating-point weights with BF16 activations) — was deployed using a nightly build of SGLang.

By segment 7, the assistant had already confirmed through benchmarking that the model was compute-bound rather than communication-bound. A TP4+PP2 configuration (tensor parallelism across 4 GPUs with pipeline parallelism across 2 groups) proved to be 2× slower than TP8, ruling out allreduce latency as the primary bottleneck. Yet a stubborn anomaly remained: the GPUs were drawing only approximately 235–330 watts out of a 600-watt thermal design power (TDP) budget, even at 98% reported utilization. This power-utilization mismatch is a classic diagnostic signature of memory-bound execution — the GPU cores appear busy but are actually stalled waiting for data, unable to sustain the arithmetic intensity needed to reach peak power draw.

The user's research request at [msg 0] captured this anomaly and proposed a hypothesis: the FP4 GEMM kernels — the fundamental building blocks of the model's matrix multiplications — were poorly tuned for the SM120 architecture. The user provided critical context: SM120 has only 100KB of shared memory per streaming multiprocessor (SM), compared to 228KB on SM100 (the datacenter Blackwell variant used in B200/B100 GPUs). The model used NVFP4 quantization with BF16 activations, and the FlashInfer CUTLASS backend was being used for both Mixture-of-Experts (MoE) and attention FP4 GEMMs. The user's five specific research questions — covering CUTLASS SM120 support, FP4 tensor core throughput comparisons, known issues, cuDNN alternatives, and TensorRT-LLM support — defined the scope of the investigation.

The Research Campaign: Systematic Parallel Investigation

The assistant's response at [msg 1] launched the investigation with four parallel web searches, each targeting a distinct dimension of the problem. As analyzed in [6], this was a carefully orchestrated first move. The searches covered CUTLASS FP4 SM120 support, NVIDIA's official hardware specifications, shared memory tile configuration constraints, and TensorRT-LLM kernel support. Notably, the assistant omitted a dedicated cuDNN search in this first round — a deliberate triage decision, prioritizing the most likely sources of information.

The search results were remarkably on-point. The first result returned a vLLM GitHub issue titled "Add SM120 (RTX 6000/5000 Blackwell) support for native NVFP4 MoE kernels" — immediately confirming that SM120 NVFP4 support was an open problem, not a solved one. The third result returned a TensorRT-LLM bug report titled "[Bug] FP4 CUTLASS GEMM fails on GB10 (SM121)" — SM121 being an even more constrained variant of the Blackwell workstation architecture. This bug report described FP4 CUTLASS GEMM failures due to shared memory overflow, directly validating the assistant's hypothesis. The fourth result returned another TensorRT-LLM bug about kernel occupancy warnings on SM120, adding a second compounding factor.

These initial results set the stage for deeper investigation. The assistant had confirmed that the problem was real, that it affected the broader Blackwell workstation family, and that shared memory overflow was the likely root cause.

Deep Dives: Uncovering the Shared Memory Crisis

Messages 2 through 5 represent the assistant's deep-dive phase, where it fetched and analyzed the full content of the discovered issues, NVIDIA's CUTLASS documentation, and architecture whitepapers. As analyzed in [4] and [7], this phase revealed the full extent of the crisis.

The CUTLASS documentation, fetched at [msg 2], confirmed a critical architectural divergence. NVIDIA's Blackwell architecture encompasses two distinct microarchitectures: SM100 (used in datacenter B200/B300 GPUs) and SM120 (used in workstation RTX PRO 6000/RTX 5090 GPUs). While both support the same tcgen05.mma tensor core instructions for FP4 operations, they differ dramatically in resource budgets:

The Breaking Point: FlashInfer #2577

Perhaps the most alarming discovery came from FlashInfer issue #2577, fetched at [msg 4]. Filed just one day before the investigation, this bug report documented that ALL THREE NVFP4 mm_fp4 GEMM backends were broken on SM120:

| Backend | Result | |---|---| | cutlass | Returns all zeros (silent failure!) | | cudnn | cudnnGraphNotSupportedError: No execution plans support the graph | | trtllm | BackendSupportedError: mm_fp4 does not support backend &#39;trtllm&#39; with capability 120 |

The CUTLASS backend silently returning all zeros was particularly dangerous. As the assistant noted, the kernel launches and completes without error, but computes nothing — consistent with shared memory overflow causing silent corruption or early exit. For a production model serving pipeline, this could lead to silently degraded output quality without any obvious indication of failure. The cuDNN backend's error confirmed that NVIDIA's own deep learning library had not yet added SM120 FP4 support, eliminating a potential alternative path. The TRT-LLM backend's error was equally definitive: TensorRT-LLM explicitly did not support FP4 on SM120 through its CUTLASS path.

The Comprehensive Synthesis: Message 6

The investigation culminated in [msg 6], a comprehensive research report that synthesized all findings into a coherent diagnosis. As analyzed in [5], this message is a masterful piece of technical synthesis, organizing the findings into five sections plus a summary.

The message's central finding was the shared memory bottleneck. The assistant wrote:

"This is the root cause of your suboptimal performance. SM120 has 100KB shared memory vs SM100's 228KB. FP4 block-scaled GEMMs require storing: A tile of FP4 data, B tile of FP4 data, Scale factor tiles for both A and B, Pipeline staging buffers (multiple stages for latency hiding). The SM100 tile configs (e.g., 128×128×256B, 256×128×128B) were designed for 228KB SMEM and overflow on SM120."

The message then connected this root cause back to the original symptom:

"This directly explains the 330W out of 600W TDP — the GPU is memory-bound because the GEMM kernels cannot use large enough tiles to be compute-bound."

The architectural comparison table in the message quantified the gap between SM100 and SM120 across six dimensions. The 2-SM CTA pair limitation was identified as particularly significant: SM100 can fuse two SMs to work on a single 256×256 tile, but SM120 cannot, halving the maximum effective tile size. The memory bandwidth gap (8 TB/s vs 1.8 TB/s) means SM120 GEMMs become memory-bound at much larger batch sizes than SM100.

The message also identified one working path forward: cuBLASLt. Benchmarks from TensorRT-LLM #11368 showed cuBLASLt achieving 99.6 TFLOPS on GB10 (SM121), suggesting it had proper SM120-specific tile configurations that fit within the 100KB shared memory budget. The assistant recommended checking cuBLASLt availability and considering TensorRT-LLM's cuBLASLt path (nvfp4_gemm_cublaslt) as an alternative.

The Broader Significance

This investigation illuminates several broader truths about the state of GPU software ecosystems for cutting-edge hardware.

The Datacenter-to-Workstation Gap

NVIDIA's Blackwell architecture is not a monolithic design. The SM100 (datacenter) and SM120 (workstation) variants share the same tensor core ISA but have fundamentally different resource budgets. Software optimized for one does not automatically work on the other, and the ecosystem is still catching up. This gap is likely to persist as NVIDIA continues to differentiate its product lines.

The Silent Failure Problem

The FlashInfer CUTLASS backend returning all zeros without an error message is a stark reminder that GPU kernel failures can be silent. A user could deploy a model, observe seemingly normal operation, and receive completely wrong results. This is particularly dangerous in production environments where output quality monitoring may not catch systematic numerical errors.

The Fragmentation of Inference Frameworks

The investigation revealed a fragmented ecosystem where different frameworks have different levels of support for the same hardware:

The Value of Cross-Reference Validation

The assistant's methodology — gathering evidence from multiple independent sources and cross-referencing them — is a model for technical investigation. No single source provided the complete picture, but together they converged on a consistent diagnosis. The CUTLASS documentation established the architectural constraints. The TensorRT-LLM bug report provided concrete evidence of SMEM overflow. The FlashInfer issues documented the practical consequences. The NVIDIA whitepapers confirmed the hardware specifications. Each source independently pointed to the same root cause.

Conclusion

The SM120 FP4 GEMM kernel investigation is a case study in systematic performance diagnosis. Starting from a single puzzling symptom — high GPU utilization with low power draw — the assistant traced the problem through multiple layers of software and hardware to a fundamental architectural constraint: the shared memory capacity of Blackwell workstation GPUs is less than half that of their datacenter counterparts, and the software ecosystem has not yet adapted.

The investigation revealed that FP4 GEMM support on SM120 was not merely suboptimal but fundamentally broken across all major backends. The CUTLASS backend silently returned zeros. cuDNN had no SM120 FP4 execution plans. TensorRT-LLM explicitly rejected SM120 FP4 operations. Only cuBLASLt showed working support, but it was not directly accessible through the FlashInfer backend abstraction used by the deployment.

This investigation directly informed the subsequent optimization work in the broader session. With the root cause understood, the team could focus on practical mitigation strategies: tuning server parameters (achieving a 28% throughput improvement), exploring alternative parallelism strategies, and documenting optimization approaches. But the fundamental lesson remained: when deploying cutting-edge hardware, the software ecosystem may not be ready, and deep kernel-level investigation is often necessary to understand why performance falls short of expectations.## References

[1] "Deconstructing a Research Request: The SM120 FP4 GEMM Kernel Investigation" — Analysis of the initial user message ([msg 0]) that launched the research campaign.

[2] "Gathering the Last Pieces: Deep-Diving into SM120 FP4 Kernel Research" — Analysis of the assistant's follow-up research gathering additional evidence.

[3] "The Pivot Point: How One Message Uncovered the Root Cause of FP4 GEMM Failure on Blackwell Workstation GPUs" — Analysis of the message that identified the shared memory overflow as the root cause.

[4] "The Deep Dive: How One Assistant Message Uncovered the Shared Memory Crisis on Blackwell SM120" — Analysis of the deep-dive phase examining CUTLASS documentation and bug reports.

[5] "Decoding the Blackwell FP4 Puzzle: A Deep Dive into SM120 GEMM Kernel Analysis" — Analysis of the comprehensive synthesis message ([msg 6]) that compiled all findings.

[6] "The Moment of Discovery: How a Research Subagent Launched the Hunt for FP4 GEMM Kernel Bottlenecks on Blackwell SM120" — Analysis of the first assistant response ([msg 1]) initiating parallel web searches.

[7] "The Pivot Point: How One Assistant Message Uncovered the SM120 Shared Memory Crisis in FP4 GEMM Research" — Analysis of the message that crystallized the shared memory crisis diagnosis.