From Empirical Limits to Research-Driven Optimization: The SM120 FP4 Kernel Research Campaign

Introduction

In the high-stakes world of large language model inference, optimization is rarely a straight line from problem to solution. More often, it is an iterative cycle of implementation, measurement, failure, research, and renewed implementation. This chunk of the opencode session captures one of the most critical phases of that cycle: the moment when empirical trial-and-error hits a fundamental hardware constraint, and the only way forward is through systematic research.

The session had been making impressive progress deploying the GLM-5-NVFP4 model—a massive Mixture-of-Experts (MoE) transformer using FP4 quantization with block scaling—on a machine with eight NVIDIA RTX PRO 6000 Blackwell GPUs. Updating sglang to the latest commit had yielded a 2x throughput improvement at 256 concurrent requests. Single-stream and dual-stream benchmarking showed excellent linear scaling (10.36 tok/s and 19.29 tok/s respectively). A comprehensive findings document (glm5findings.md) had been written, capturing over 500 lines of discoveries and lessons learned.

But then came the wall. The Opportunistic Expert Activation (OEA) technique, which attempted to skip low-probability experts at decode time, showed near-zero average throughput improvement on random data. More critically, the EP8 (Expert Parallelism with 8-way sharding) attempt failed with a CUTLASS tile configuration error: the 128×256×128 tile exceeded the SM120's 100KB shared memory limit. This was not a configuration mistake—it was a fundamental hardware constraint that empirical methods alone could not surmount.

It was at this inflection point that the assistant pivoted from empirical optimization to research-informed engineering. This chunk captures that pivot in its entirety: from the meticulously structured research task (message 0), through the systematic multi-round research process (messages 1–5), to the comprehensive synthesis report (message 6) that would inform the next wave of optimization work.

The Research Task: A Map of Unknown Territory

The research began with a detailed eight-point brief from the user ([msg 0]), which is analyzed in depth in [6]. This was not a vague request to "look into FP4 stuff." It was a precisely scoped research agenda that reflected deep understanding of the problem space:

"Research the latest developments in NVIDIA SM120 (RTX PRO 6000 / RTX 5090 / Blackwell consumer) FP4 kernel optimization. This GPU has 100KB shared memory per SM (same as Ampere), uses mma.sync instructions (not SM100's tcgen05), and has a theoretical 1.85 PFLOPS FP4 dense peak per GPU."

The brief's structure reveals a layered decomposition of the optimization problem. Items 1–4 target the existing software ecosystem (CUTLASS kernels, TensorRT-LLM, FlashInfer, community projects)—the most likely sources of pre-built solutions. Items 5–6 address algorithmic techniques (MoE optimization, expert batching, token routing)—the next line of inquiry if pre-built kernels don't exist. Items 7–8 cover advanced architectural approaches (warp specialization, CUTLASS version updates)—deeper, more speculative research that might yield longer-term solutions.

The explicit mention of architectural constraints—100KB shared memory, mma.sync instructions, 1.85 PFLOPS theoretical peak—serves a dual purpose. It constrains the search to avoid irrelevant results (anything requiring tcgen05.mma or more than 100KB shared memory is automatically filtered out), and it communicates to the research agent that the user already understands the hardware fundamentals. This is not a beginner asking for an introduction; this is a practitioner who has already seen OutOfResources: Required: 147456, Hardware limit: 101376 errors and needs to know every trick in the book to work around them.

The Research Process: Systematic Intelligence Gathering

The assistant's response to this task was a model of systematic research methodology, spanning messages 1 through 5. As analyzed in [3], [5], [4], [1], and [2] respectively, each message represents a distinct phase in a carefully orchestrated intelligence-gathering campaign.

Message 1 ([3]) opened with a broad-first strategy: five parallel Exa web searches covering CUTLASS, TensorRT-LLM, FlashInfer, warp specialization, and MoE inference optimization. The initial results returned a rich set of sources: GitHub issues about SM120 support in vLLM, TensorRT-LLM FP4 support questions, FlashInfer documentation for MXFP4 grouped GEMM, and a CUTLASS issue about FP4 instruction support on SM_101 architecture. This broad sweep established the landscape and identified the most promising veins to mine.

Message 2 ([5]) shifted to deep-dive mode. The assistant evaluated the initial results and fetched the most critical sources in detail: the CUTLASS Blackwell documentation (which contains the definitive SM120-specific section), the DeepWiki page on SM120 architecture, the HuggingFace blog post comparing FP4 MoE frameworks, and launched additional searches for CUTLASS 4.x updates and MoE caching papers. This phase was about separating signal from noise—identifying which sources contained authoritative information versus which were superficial.

Message 3 ([4]) continued the deep-dive strategy with targeted fetches of the CUTLASS SM100 GEMM documentation (which also contains SM120-specific sections), the CUTLASS changelog, and searches for SM120 tile size constraints and vLLM SM120 support. The changelog fetch was particularly strategic—it provided a complete version history of SM120 fixes and features, which would become the backbone of the final report's sections on CUTLASS capabilities.

Message 4 ([1]) narrowed the focus further, fetching the Colfax Research tutorial on sub-byte GEMM on Blackwell GPUs and searching for specific vLLM PRs related to SM120 NVFP4 MoE. The search for DeepGEMM SM120 support was a critical negative finding: DeepGEMM does not work on SM120 because it requires tcgen05.mma instructions that only exist on SM100. This finding—that CUTLASS is the only viable FP4 path on consumer Blackwell—would become one of the report's most important conclusions.

Message 5 ([2]) represented the final intelligence-gathering push, fetching the two most critical community sources: the NVIDIA Developer Forum thread about SM121 CUTLASS kernel optimization (BTankut's work achieving 356 TFLOPS dense NVFP4 on DGX Spark) and the DGX Spark FP4 benchmarking thread with real-world performance numbers. The assistant also fetched the Opportunistic Expert Activation paper from arXiv, completing the research coverage.

The overall research strategy is methodical and efficient: broad searches first to map the territory, then deep dives on the most promising sources, then targeted searches for specific questions that remain unanswered. The assistant never fetches the same source twice and never searches for something that has already been found.

The Research Report: A Comprehensive Synthesis

The culmination of this research effort was message 6 ([7]), a comprehensive research report that synthesized findings across all eight research directions into a coherent, actionable document. The report is structured into eight numbered sections plus a summary, following the user's original research questions almost exactly—but with a deliberate narrative architecture that makes it more than a mechanical response to a numbered list.

Section 1: CUTLASS Kernels — The Foundation

The report opens with CUTLASS because it is the foundational layer. Everything else—TensorRT-LLM, FlashInfer, vLLM—builds on CUTLASS kernels or wraps them. The architectural comparison table (SM100 vs SM120) is the single most important piece of information in the entire report:

| Feature | SM100 (B200) | SM120 (RTX 5090/PRO 6000) | |---|---|---| | MMA instruction | tcgen05.mma (5th gen) | mma.sync (Ampere-era) | | Tensor Memory (TMEM) | 256 KB | None | | Shared Memory | 228 KB | ~100 KB (101,376 bytes) | | 2-SM CTA Pairs | Yes | No | | TMA Multicast | Yes | No | | Cluster Size | Flexible | 1×1×1 only |

This table tells a stark story: SM120 has no TMEM, half the shared memory of SM100, no 2-SM CTA pairs, no TMA multicast, and only 1×1×1 clusters. Every subsequent finding in the report is a consequence of these hardware limitations.

The report identifies seven specific CUTLASS examples targeting SM120 FP4 (Examples 79a through 79d, 80, 87, and the CuTe DSL example), providing a curated list that saves a practitioner hours of searching through the CUTLASS repository. It also documents the PTX instruction used on SM120:

mma.sync.aligned.m16n8k64.row.col.kind::mxf4nvf4.block_scale.scale_vec::4X.f32.e2m1.e2m1.f32.ue4m3

This is an mma.sync instruction (NOT tcgen05.mma), SM120+ only, processing 64 FP4 elements per MMA operation.

Sections 2–4: Framework and Community Landscape

The TensorRT-LLM section confirms that NVIDIA's flagship inference framework does support SM120 for NVFP4, but with a critical caveat: a known bug (issue #11368) where B200-sized tile configs overflow SM120's shared memory. Even NVIDIA's own engineers sometimes forget that SM120 has only 100KB.

The FlashInfer section reveals that FP8 grouped GEMM doesn't support SM120 (only SM100 and SM103), closing off certain optimization paths. However, MXFP4 grouped GEMM does work with specific tile parameters (tile_m=128, tile_n=128, tile_k=128). Performance data from the HuggingFace blog post by @apsys establishes the relative characteristics: FlashInfer CuteDSL excels at large batch sizes with many experts but struggles at interactive batch sizes due to preprocessing overhead.

The Community Projects section is arguably the most valuable for a practitioner. BTankut's SM121 optimization results (356 TFLOPS dense NVFP4, 154 TFLOPS MoE grouped GEMM) provide a proven recipe for success on consumer Blackwell hardware. The DGX Spark community benchmarks offer real-world performance numbers for actual models (Qwen3-Coder-30B-A3B) running on actual hardware—far more useful than theoretical peak FLOPs. The critical finding that "Default MoE configs from SGLang/vLLM request ~147KB SMEM and crash on SM120/SM121" directly explains the EP8 failure the team had encountered.

Sections 5–6: Algorithmic Techniques

The report broadens from kernel engineering to system-level optimization, surveying papers on expert routing (Opportunistic Expert Activation), expert caching (BuddyMoE, Harvest), multi-batch scheduling (Klotski), and dynamic expert remapping (Lynx). Each paper offers a technique that could be applied on top of whatever kernel solution is chosen.

Section 6 distills these papers into seven actionable techniques for increasing effective batch size per expert: batch-aware routing, expert-first layout, token reordering with fused reduction, expert caching, multi-batch pipeline, DeepEP dispatch, and adaptive grid sizing. The inclusion of actual CUDA code for adaptive grid sizing grounds the report in implementation reality:

while (grid.x <= multiProcessorCount && block.x > 64) {
    grid.x *= 2;
    block.x = (block.x + 1) / 2;
}

Sections 7–8: Advanced Architecture and Version History

The warp specialization section confirms that SM120 supports both Pingpong and Cooperative warp specialization schedules within the KernelTmaWarpSpecialized1SmSm120 kernel schedule—important because it means advanced scheduling techniques developed for Hopper and SM100 can be adapted, albeit with the 1-SM limitation.

The CUTLASS version history traces SM120 support from 4.0.0 through 4.4.0, noting which features were added in each version. This is valuable for deciding which version to target and understanding the maturity of the support.

The Shared Memory Budget Analysis

The report's capstone analysis is the quantification of the 100KB shared memory budget. For a 128×128×128 FP4 block-scaled tile with 2 pipeline stages:

Key Findings and Their Impact

The report concludes with eight actionable findings that directly inform the GLM-5-NVFP4 deployment strategy:

  1. CUTLASS is the only viable FP4 path on SM120 — DeepGEMM requires tcgen05.mma and fails
  2. 100KB SMEM is the binding constraint — all tile/pipeline configs must be budgeted within 101,376 bytes
  3. Community-proven peak: 356 TFLOPS NVFP4 dense on SM121, 154 TFLOPS MoE grouped GEMM
  4. Default framework configs crash — custom tile tuning within SMEM budget is mandatory
  5. Kernel fusion matters more on SM120 than SM100 due to tighter memory bandwidth
  6. Batch-aware expert routing (Opportunistic Expert Activation) can give 39% MoE latency reduction
  7. CUTLASS 4.4.0 has critical SM120 fixes (pingpong memory fence, SMEM alignment for scale factors)
  8. FlashInfer CuteDSL excels with many experts (256+) at large batches; SGLang better for interactive (BS=1-16) These findings would directly shape the next wave of optimization work documented in the subsequent chunk: updating sglang to the latest commit, implementing OEA with proper sigmoid scoring, retrying EP8 with memory-safe configurations, and benchmarking with an understanding of what the hardware can actually deliver.

The Broader Significance

This chunk captures a universal pattern in engineering work at the frontier of AI infrastructure. When you are deploying a state-of-the-art model on newly released hardware, you cannot rely on pre-tuned configurations. The default settings—designed for datacenter GPUs with 228KB shared memory—will crash on consumer Blackwell hardware with 100KB. You must understand the hardware constraints at a deep level, survey the ecosystem for existing solutions, and synthesize findings into actionable recommendations.

The research campaign documented in this chunk is a masterclass in this kind of work. It demonstrates that effective optimization is not about randomly trying things until something works—it is about systematically gathering intelligence, building a mental model of the constraints, and using that understanding to make informed engineering decisions. The assistant acted as a knowledge broker, aggregating information from GitHub issues, NVIDIA documentation, HuggingFace blog posts, arXiv papers, and community forums, and synthesizing it into a coherent picture tailored to a specific deployment context.

In the broader arc of the opencode session, this chunk is a turning point. Before it, the session was characterized by empirical trial-and-error. After it, the session can pivot to informed optimization—choosing the right CUTLASS version, selecting appropriate tile configurations, implementing expert routing improvements, and benchmarking with an understanding of what the hardware can actually deliver. The research did not solve every problem, but it provided a map of the territory that made subsequent engineering work vastly more efficient.

Conclusion

The SM120 FP4 kernel research campaign represents a strategic pivot from empirical optimization to research-informed engineering. Faced with a fundamental hardware constraint—the 100KB shared memory limit on consumer Blackwell GPUs—the assistant executed a systematic research process spanning multiple rounds of broad searches, deep dives, and targeted fetches. The resulting comprehensive report synthesized findings across the entire software stack, from low-level CUTLASS kernel primitives through inference frameworks to algorithmic techniques for MoE optimization.

The report's central contribution is its focus on the 100KB shared memory constraint as the organizing principle of SM120 FP4 optimization. Every finding—every tile size recommendation, every framework limitation, every community result—is connected back to this hardware reality. This gives the report a conceptual unity that makes it more than the sum of its parts, transforming a collection of search results into a coherent knowledge artifact that directly informs engineering decisions.

For the practitioner deploying GLM-5-NVFP4 on RTX PRO 6000 GPUs, the research provides a roadmap: use CUTLASS 4.4.0, budget your shared memory carefully, expect default configurations to fail, fuse your kernels, and consider batch-aware expert routing. For the observer of this session, it provides a window into how expert-level technical research is conducted in the age of large language models—systematically, collaboratively, and with a clear focus on actionable outcomes.