The FP4 Frontier: Deconstructing a Deep Research Campaign into SM120 Tensor Core Optimization

Introduction

In the rapidly evolving landscape of GPU-accelerated machine learning inference, few challenges are as tantalizing—and as maddeningly complex—as maximizing the utilization of NVIDIA's newest tensor core architectures. This article synthesizes a remarkable multi-round research campaign conducted within an opencode coding session, where an AI assistant was tasked with investigating the most cutting-edge approaches to maximize FP4 (4-bit floating point) tensor core utilization on NVIDIA's SM120 architecture—the compute engine powering the RTX PRO 6000 Blackwell workstation GPUs.

The research campaign, spanning seven messages across a dedicated subagent session, represents a systematic deep dive into one of the most pressing optimization problems in contemporary ML inference: how to close the gap between theoretical peak throughput and achieved real-world performance for small-batch Mixture-of-Experts (MoE) decode operations. The assistant's journey—from initial web searches through source code analysis, documentation parsing, microbenchmark interpretation, and final synthesis—offers a masterclass in technical research methodology. This article examines each phase of that journey, the reasoning that drove it, the mistakes made along the way, and the actionable insights ultimately produced.

The Context: A Real-World Deployment Under Pressure

The research did not occur in a vacuum. It was embedded within a larger session (segment 7) where the assistant had already accomplished substantial work: benchmarking a TP4+PP2 (Tensor Parallelism 4, Pipeline Parallelism 2) configuration across 8 RTX PRO 6000 Blackwell GPUs, deeply analyzing FP4 GEMM kernel efficiency on SM120, achieving a 28% throughput improvement through server parameter tuning (raising --max-running-requests to 2048 and setting --num-continuous-decode-steps 8), and beginning to document optimization approaches as glb5improvement-xx.md files. The model being deployed was GLM-5-NVFP4, a large language model using NVIDIA's NVFP4 (4-bit floating point) weight format, served via SGLang.

The user had already established that the model was compute-bound rather than communication-bound—TP4+PP2 was 2× slower than TP8, ruling out allreduce latency as the primary bottleneck. This meant the path to better performance lay in optimizing the tensor core operations themselves. The measured numbers painted a stark picture: CUTLASS kernels peaked at approximately 1,300 TFLOPS out of a theoretical ~1,850 TFLOPS dense FP4 (about 70% utilization) for large GEMMs, but for the small-batch MoE decode operations that dominate real inference—batch sizes of 16–64 tokens per expert—utilization plummeted to a mere 15–55 TFLOPS, or roughly 1–3% of peak.

This was the problem the user wanted solved. And they knew enough to ask precisely the right questions.

The Opening Salvo: A Masterfully Framed Research Request

The user's opening message ([msg 0]) to the subagent was not a casual question. It was a meticulously crafted technical inquiry that revealed weeks of prior work and a sophisticated mental model of the problem [1]. The user provided precise architectural context: SM120 uses mma.sync instructions (the Ampere-era programming model) with FP4 data types bolted on, has only 100KB of shared memory, lacks TMEM (tensor memory), and cannot form 2-SM cooperative CTA pairs—all features that distinguish it from the datacenter SM100 Blackwell architecture.

The seven research questions the user posed formed a taxonomy of the solution space:

  1. CUTLASS Example 79 ("blackwell_geforce_gemm")—what tile sizes and scheduling does it use?
  2. SM120-specific FP4 kernels in TensorRT-LLM, FlashInfer, vLLM, and other frameworks
  3. NVIDIA's SM120-optimized libraries released since late 2025
  4. The SM120_16x8x64_TN_VS MMA instruction—its per-cycle throughput and kernel structuring implications
  5. Persistent kernel approaches to avoid launch overhead for tiny GEMMs
  6. FP4 structured sparsity (2:4) —applicability to dense NVFP4 weights
  7. Any projects achieving >80% FP4 utilization on SM120 The user's inclusion of specific batch sizes (16–64), TFLOPS ranges (15–55), and the exact CUTLASS example number (79) demonstrated that this was not a hypothetical inquiry but a grounded investigation driven by measured data [1]. The user had already done the hard work of characterizing the problem; the subagent's job was to find the solutions.

The Research Cascade: Five Rounds of Systematic Investigation

The assistant's response unfolded across five rounds of information gathering (messages 1–5), each building on the previous round's results. This cascade reveals a clear research methodology: broad search, targeted fetch, gap-fill, deep documentation, and final verification.

Round 1: Broad Web Searches

In message 1, the assistant launched four parallel Exa web searches targeting the CUTLASS example, the MMA instruction, general RTX 5090 FP4 optimization, and inference framework implementations [5]. The search results were a mixed bag: the CUTLASS query returned documentation for SM100 (datacenter Blackwell) rather than SM120, revealing an early challenge in distinguishing between the two architectures. A TensorRT-LLM GitHub discussion about RTX 5090 deployment offered a more promising lead, while an NVIDIA blog post about NVFP4 inference focused on datacenter GPUs rather than the consumer/workstation variant.

The assistant's choice of which topics to prioritize in this first round—and which to defer—reveals strategic thinking. It addressed the most concrete leads first (the known CUTLASS example, framework implementations) while deferring more speculative topics (persistent kernels, structured sparsity) to later rounds [5].

Round 2: Direct Source Code and Blog Fetching

Message 2 marked a pivot from search to direct retrieval [3]. The assistant fetched the raw CUTLASS Example 79 source code from GitHub (79b_blackwell_geforce_nvfp4_nvfp4_gemm.cu) and the HuggingFace blog post "TFLOPS Gap: Why FP4 MoE Kernel Engineering Matters on Blackwell," while simultaneously launching new searches for persistent kernel approaches and FP4 structured sparsity.

The decision to fetch raw source code rather than relying on documentation summaries was significant. Documentation might simplify or omit critical details, but the source code contains the exact template parameters, tile shapes, and scheduling policies. This is the difference between reading a map and walking the terrain [3].

However, the tool environment imposed limitations: the webfetch returned only the beginning of the CUTLASS source file (truncated due to file size), and the HuggingFace blog post returned only its HTML structure without the full content. The assistant had to work with partial information—a realistic constraint of web research.

Round 3: Targeted Gap-Fill

Message 3 addressed the remaining unanswered questions [4]. The assistant launched three targeted searches: one for the exact SM120_16x8x64_TN_VS instruction throughput, one for the FlashDMoE persistent kernel paper, and one for CUTLASS SM120 tile scheduling specifics.

The search for the MMA instruction was remarkably precise—SM120 mma.sync.aligned.m16n8k64 FP4 e2m1 instruction throughput cycles per warp—revealing that the assistant had already absorbed architectural knowledge from earlier rounds and knew the exact instruction mnemonic. The result was a GitHub issue (#2908) about FP4 instructions being unsupported on SM_101, which indirectly confirmed SM120 support but provided no throughput numbers.

The FlashDMoE search was more successful, returning the arXiv paper "FlashDMoE: Fast Distributed MoE in a Single Kernel" (2506.04667v2)—a directly relevant academic work addressing the user's question about persistent kernel approaches. The CUTLASS scheduler search returned only a PyPI package page, confirming that low-level scheduling internals are poorly documented in searchable form [4].

Round 4: Deep Documentation Fetching

Message 4 represented a shift to authoritative sources [2]. The assistant fetched three specific documents: the NVIDIA CUTLASS documentation for Blackwell, Lei Mao's benchmark of Tensor Core MMA instruction peak performances, and the DeepWiki page on SM120 Blackwell GeForce architecture.

Lei Mao's blog was particularly valuable. As a well-known figure in the CUDA optimization community, his microbenchmark results provide ground-truth measurements of instruction throughput that official documentation often lacks. The DeepWiki page offered community-indexed architectural details about Grid Dependency Control (GDC) and other SM120-specific features [2].

Round 5: The SM100 Trap

Message 5 contained a critical error that illustrates the challenges of navigating NVIDIA's fragmented architecture landscape [6]. The assistant attempted to fetch the "SM120-specific CUTLASS documentation page" but instead fetched a URL (blackwell_functionality.html) whose content was titled "Blackwell SM100 GEMMs"—documenting the datacenter architecture that uses tcgen05.mma instructions, has TMEM, and supports 2-SM CTA pairs.

This was fundamentally the wrong architecture. SM120 uses the older mma.sync instruction model, has no TMEM, and cannot form 2-SM cooperative thread arrays. The assistant's assumption that a page titled "Blackwell Functionality" in the CUTLASS documentation would cover both SM100 and SM120 was incorrect [6]. This mistake—while understandable given NVIDIA's ambiguous documentation naming—could have led the assistant to internalize optimization strategies that simply do not apply to the target architecture.

Fortunately, the second action in message 5 was more productive: a search for the Nunchaku (SVDQuant) project's SM120 FP4 kernels, which targeted the correct architecture and offered a promising line of investigation [6].

The Grand Synthesis: Message 6's Comprehensive Findings

Message 6, the culmination of the research campaign, is a remarkable piece of technical synthesis spanning all seven research questions [7]. The assistant compiled findings from CUTLASS source code, Lei Mao's microbenchmarks, the HuggingFace community's framework comparisons, academic papers, and GitHub issue discussions into a coherent analysis.

Finding 1: CUTLASS Example 79 Configuration

The assistant extracted the exact configuration of NVIDIA's reference SM120 FP4 GEMM implementation: Sm120 architecture tag, OpClassBlockScaledTensorOp operator class, 128×128×128 thread block shape (M, N, K), 1×1×1 cluster shape (mandatory—SM120 cannot use TMA multicast), and KernelTmaWarpSpecializedPingpong kernel schedule [7].

The assistant then constructed a comparison table between SM100 and SM120 that distilled the critical architectural differences:

| Feature | SM100 (Datacenter) | SM120 (Consumer) | |---|---|---| | MMA Instruction | tcgen05.mma (5th gen) | mma.sync (Ampere-era) | | Memory | TMEM (Tensor Memory) | SMEM only (100KB) | | CTA Pairs | 2-SM cooperative | 1-SM only | | Cluster Shape | Up to 16× | 1×1×1 only | | SMEM Budget | 228KB | 100KB |

This comparison explains why SM120 cannot simply reuse SM100 kernel configurations. The 100KB SMEM budget (versus 228KB on SM100) is a particularly severe constraint—it limits tile sizes, reduces software pipeline depth, and forces more frequent data movement from global memory [7].

Finding 2: The MMA Instruction's True Capability

Lei Mao's microbenchmarks on an RTX 5080 provided ground-truth measurements of the SM120_16x8x64_TN_VS instruction: 933.1 TFLOPS, exceeding the advertised peak of 900.4 TFLOPS (103.6% utilization) due to boost clock variance. The k=64 variant achieved approximately 4× the throughput of the k=32 variant, confirming that "you MUST use the k=64 block-scaled variant" [7].

This finding was crucial: it proved that the tensor cores themselves are not the bottleneck. At the instruction level, SM120 can achieve perfect FP4 utilization. The problem lies entirely in the kernel infrastructure surrounding the tensor core instructions.

Finding 3: Framework Support Survey

The assistant surveyed TensorRT-LLM, SGLang, FlashInfer, and Nunchaku for SM120 FP4 support. Key findings included a critical bug in TensorRT-LLM (GitHub #11368) where FP4 CUTLASS GEMM fails on SM121 due to shared memory overflow from B200-sized tile configurations—a direct consequence of the 100KB SMEM limit. SGLang was identified as "the best-performing FP4 MoE stack on Blackwell according to the HuggingFace benchmark study," with architecture-agnostic optimizations applicable to SM120 [7].

Finding 4: The Core Problem Quantified

The assistant's most important contribution was quantifying the fundamental problem. For batch sizes of 16–64 with hidden_dim=7168 and K=2048:

Finding 5: Persistent Kernel Recommendations

The assistant synthesized a five-step practical recommendation for SM120 MoE decode: fuse gate_up GEMM + SiLU + down GEMM into a single persistent kernel, use Grouped GEMM to batch all expert computations, pad expert token counts to 128-byte boundaries for TMA alignment, use adaptive grid sizing with smaller thread blocks for small batches, and chain layers using GDC [7].

Finding 6: Sparse FP4 Assessment

The assistant confirmed that SM120 has native sparse FP4 tensor core instructions (SM120_SPARSE_16x8x128_TN_VS achieving 1872 TFLOPS—exactly 2× the dense peak) but was appropriately cautious about applicability. Converting dense NVFP4 weights to 2:4 sparsity would reduce effective bit width from ~4.5 bits/weight to ~2.25 bits/weight, which is "aggressive." However, for MoE models specifically, expert sparsity offers a promising angle because different experts can be pruned independently [7].

Finding 7: The Definitive Answer

The assistant's answer to whether anyone has achieved >80% FP4 utilization on SM120 was nuanced and definitive: "At the instruction level, YES. At the kernel level for real workloads, NO." Lei Mao's microbenchmarks prove the tensor cores can achieve 103.6% of advertised peak, but for real inference workloads at small batch sizes, utilization plummets to 1–8% or even below 1%. The root cause: "The fundamental bottleneck for MoE decode (BS=16–64) is that the problem is memory-bandwidth-bound, not compute-bound" [7].

The Broader Significance

This research campaign illuminates a fundamental challenge in GPU-accelerated inference: the growing gap between theoretical peak throughput and achievable real-world performance. As GPU architectures become more specialized and tensor core throughput increases exponentially, the overhead of kernel launch, memory loading, and scheduling becomes proportionally more important. The finding that launch overhead dominates compute by 20–40× for small-batch operations is not specific to SM120 or FP4—it is a general problem affecting all GPU inference workloads at small batch sizes.

The campaign also highlights the tension between datacenter and consumer architectures. NVIDIA's Blackwell family is bifurcated: SM100 gets the full set of innovations (TMEM, 2-SM CTA pairs, large SMEM), while SM120 gets a stripped-down variant with Ampere-era tensor core instructions bolted onto a Hopper-style TMA infrastructure. This fragmentation creates real problems for kernel developers, as evidenced by the TensorRT-LLM SMEM overflow bug.

Conclusion

The FP4 research campaign examined in this article demonstrates the value of systematic technical research in an era of information abundance. The assistant did not have access to a single authoritative source that answered all seven research questions. Instead, it assembled answers from a dozen different sources, cross-referencing and synthesizing to produce a coherent picture. The journey included dead ends (the SM100 documentation trap), partial successes (truncated source code fetches), and genuine breakthroughs (Lei Mao's microbenchmarks proving 100% instruction-level utilization).

The final synthesis—that the problem is memory-bandwidth-bound, not compute-bound—reframes the optimization challenge entirely. It shifts the focus from "how do we make the tensor cores faster?" to "how do we feed the tensor cores more efficiently?" This insight, combined with the prioritized action plan of persistent grouped GEMM, kernel fusion, and GDC chaining, provides a concrete roadmap for the next phase of optimization work on the GLM-5-NVFP4 deployment.## References

[1] "The SM120 FP4 Optimization Inquiry: A Deep Dive into a Cutting-Edge Research Task" — Analysis of the user's opening message (msg 0) that launched the research campaign.

[2] "The Research Deep Dive: Fetching SM120 FP4 Kernel Documentation" — Analysis of message 4, where the assistant fetched authoritative documentation sources.

[3] "Drilling Deeper into SM120 FP4: The Data-Gathering Pivot" — Analysis of message 2, the transition from broad search to targeted deep-dives.

[4] "Closing the Loop: The Third Research Pass in SM120 FP4 Kernel Optimization" — Analysis of message 3, the gap-fill round targeting remaining unanswered questions.

[5] "The Opening Salvo: How a Single Message Launched a Deep Dive into SM120 FP4 Kernel Optimization" — Analysis of message 1, the assistant's first response with parallel web searches.

[6] "The SM100 Trap: A Research Assistant's Critical Misstep in the Hunt for SM120 FP4 Kernel Optimization" — Analysis of message 5, where the assistant fetched SM100 documentation instead of SM120.

[7] "The FP4 Frontier: Deconstructing a Deep Dive into SM120 Tensor Core Optimization" — Analysis of message 6, the comprehensive synthesis of all research findings.