The SM120 GEMM Optimization Campaign: How Systematic Research Reframed a Structural Bottleneck

Introduction

In the high-stakes world of large language model inference on cutting-edge hardware, the difference between a system that delivers usable throughput and one that collapses under latency constraints often comes down to a single question: where is the bottleneck? For engineers deploying Mixture-of-Experts (MoE) models like Kimi-K2.5 (1 trillion parameters, 384 experts) and GLM-5-NVFP4 (744B parameters, 256 experts) on NVIDIA's Blackwell (SM120) GPUs, that question has a sharp edge. Each decode step launches hundreds of tiny matrix multiplications (GEMMs) that barely scratch the surface of the GPU's advertised 1,850 FP4 TFLOPS. The measured utilization hovers below 1%. Something is fundamentally broken in the software stack.

This article examines a focused subagent session—a deep-dive research campaign into SM120 GEMM optimization—that unfolded across nine messages within a larger coding session. The subagent was given an open-ended brief: "Research SM120 GEMM optimization (agent: explore)." Over the course of its investigation, it read approximately 30 files spanning model configurations, service definitions, benchmark scripts, profiling analyses, kernel source code, and improvement proposals. The result was a comprehensive synthesis that fundamentally reframed the optimization problem—from "how do we make the GEMM faster?" to "how do we change the parallelism strategy to create larger GEMM dimensions?"

This article traces the arc of that investigation, examining the subagent's methodology, its key findings, the assumptions that shaped its analysis, and the structural bottleneck diagnosis that emerged. It is a case study in how deep system understanding emerges from the synthesis of many data sources, and how the framing of a problem determines the range of possible solutions.

The Mission: An Open-Ended Exploration Brief

The subagent's mission originated from a parent session that had been conducting a comprehensive profiling campaign of the Kimi-K2.5 INT4 model on eight RTX PRO 6000 Blackwell GPUs. The profiling had identified AllReduce as the dominant bottleneck at 51.5% of decode time, but the team suspected that the real story was more nuanced—that the tiny per-expert GEMM dimensions (N=256 after TP=8 sharding) were the deeper structural problem.

The user's brief (message 0) was remarkably thorough [2]. It asked for four categories of information: (1) the current kernel paths for MoE in vLLM, (2) a search of local project files for any previous analysis, (3) the exact GEMM dimensions for Kimi-K2.5 INT4 and MiniMax-M2.5 under TP=8, and (4) the quantization formats in use. The closing instruction was telling: "Return all findings in detail — I need to understand exactly what GEMM kernels are being called, what dimensions they operate on, and what the theoretical vs actual throughput is."

This was not a casual question—it was a targeted intelligence-gathering operation. The user had reached a decision point. The profiling data told them what was slow (AllReduce), but not why the compute portion wasn't faster. The hypothesis was clear: tiny GEMMs cannot saturate the GPU. But the hypothesis needed evidence. The user needed to know: are we calling the right kernels? Are the dimensions as small as we think? Is the quantization format forcing suboptimal kernel selection? Only with this information could they decide whether to invest in kernel engineering, architectural changes, or higher-level algorithmic optimizations like speculative decoding.

The Methodology: Systematic Information Gathering

The subagent's approach to this investigation reveals a sophisticated, multi-phase research methodology that unfolded across messages 1 through 8.

Phase 1: Broad-First Exploration (Messages 1-2)

The subagent began by casting a wide net [7]. In message 1, it issued ten parallel searches across the local project directory at /home/theuser/glm-kimi-sm120-rtx6000bw/, covering file patterns for findings, improvements, GEMM analysis, MoE configurations, profiling scripts, kernel references, markdown documents, text configs, and Python scripts. This was not random—each pattern was deliberately chosen to cover a specific category of performance engineering knowledge.

The glob results revealed a rich ecosystem of prior work [3]: a master findings document (FINDINGS.md), a comprehensive GLM-5 findings document (glm5findings.md), twelve numbered improvement proposals (glb5improvement-01 through glb5improvement-12), multiple benchmark scripts, MoE configuration JSONs, service definition files, and logs from previous experiments. The directory structure itself told a story of a systematic, multi-month optimization campaign.

In message 2, the subagent selected ten files to read in parallel [3]. The selection logic revealed a sophisticated understanding of information hierarchy. The subagent read the executive summaries first (FINDINGS.md, glm5findings.md) to establish the big picture. It then read three strategically chosen improvement proposals: number 4 (Expert Parallelism, status "READY TO TEST," expected impact 30-100%), number 7 (Persistent Grouped GEMM, status "NEEDS SIGNIFICANT ENGINEERING," expected impact 2-5x), and number 8 (FP4 Structured Sparsity, status "REQUIRES MODEL MODIFICATION," expected impact up to 2x). These three covered the spectrum from immediately actionable to ambitious research. Finally, it read the analytical tools (decode_latency_breakdown.py, theoretical_max_analysis.py), benchmark scripts (bench_fp4_cublaslt.py), and configuration artifacts (MoE JSON configs).

Phase 2: Targeted Code Search (Messages 3-5)

Having established the high-level landscape, the subagent pivoted to targeted investigation. In message 3, it read ten more files [1]—the second wave of reads that drilled from high-level findings into specific implementation details. These included the FP4 benchmark scripts (bench_fp4_cublaslt_v2.py, bench_fp4_cublaslt_v3.py), the cuBLASLt FP4 summary script (cublaslt_fp4_summary.py), additional improvement proposals (numbers 10, 11, 6), the memory bandwidth optimization research document, the MoE configs list, and the service definition files for both Kimi-K2.5 and MiniMax-M2.5.

The service files were particularly revealing. They showed that the system uses PCIe-only GPU interconnect (NCCL_P2P_LEVEL=SYS), which means no NVLink or NVSwitch—a critical constraint that shapes all communication optimization strategies. The Kimi-K2.5 service used INT4 quantization with the compressed-tensors format, while MiniMax-M2.5 used FP8.

In message 4, the subagent continued filling gaps [8], reading the patches directory (which revealed that only an attention patch existed—no MoE-specific patches), the logs directory (12 log files documenting experiments with different configurations), the source directory (custom attention implementations, MXFP4 quantization code, SM120 detection utilities), the system improvement document, and improvement proposal 12 (Opportunistic Expert Activation, the only implemented optimization at +5.7% throughput).

Message 5 marked a critical transition [6]: the subagent conducted two targeted grep searches for kernel-specific terms. The first search targeted fused_moe|CompressedTensors|Marlin|marlin|cutlass_scaled_mm|WNA16|W4A16|W8A8 and found references in analyze_profile.py (mentioning flashinfer's cutlass_fused_moe), decode_gap_analysis.py (importing SGLang's grouped_gemm_triton), and source/mxfp4.py. The second search expanded to include intermediate_size, hidden_size, and MiniMax, revealing that the improvement documents contained estimated values for these critical architectural parameters.

Phase 3: Primary Source Verification (Messages 6-7)

In message 6, the subagent read the primary sources [4]: the Kimi-K2.5 model configuration JSON (configs/kimi-k25-config.json), the MXFP4 quantization implementation (source/mxfp4.py), the decode gap analysis script (decode_gap_analysis.py), and the profiler analysis script (analyze_profile.py). This was the moment when the subagent stopped trusting summaries and started examining raw evidence.

The profiler analysis contained a bombshell finding: unrolled_elementwise_kernel (copy_/dtype cast) consumed 1.938 seconds out of 2.797 total CUDA time—69.3% of all GPU time was spent on data copying and type casting, not on actual matrix multiplication. The ncclDevKernel_AllReduce_Sum_bf16_RING_LL kernel accounted for only 0.2 seconds. This finding fundamentally challenged the earlier assumption that AllReduce was the primary bottleneck.

Message 7 was the final data-gathering step [5]: the subagent read the MoE tuning log (moe-tuning.log) to see exact GEMM dimensions used during tuning, and the AllReduce fusion improvement proposal (glb5improvement-05-allreduce-fusion.md). With a deliberate "Now I have a comprehensive picture," the subagent signaled the transition from data gathering to synthesis.

The Synthesis: Message 8's Comprehensive Analysis

Message 8 [9] is the culmination of the entire investigation—a detailed technical report spanning kernel paths, exact GEMM dimensions, quantization formats, profiling data, and theoretical throughput analysis. Its structure reveals the subagent's reasoning process.

Kernel Paths Mapped

The subagent identified three distinct model-quantification combinations, each with its own kernel path:

  1. Kimi-K2.5 INT4: Uses Marlin INT4 kernels via Triton's fused MoE layer. The quantization method is CompressedTensorsWNA16MarlinMoEMethod. There was a known SM120 issue: Triton kernels required 147 KB shared memory, but SM120 only has 101 KB. This was fixed by PR #16975, which set is_persistent=False, num_stages=1, and used StridedLayout instead of TMA block layout.
  2. GLM-5-NVFP4: Uses FlashInfer CUTLASS FP4 grouped GEMM. The MoE runner backend is flashinfer_cutlass. CUTLASS tile constraints on SM120 were documented: only tiles that fit in 100 KB shared memory work (128x128x128 and 128x128x256), while larger tiles fail. The FP4 MMA instruction on SM120 was identified as mma.sync.aligned.m16n8k64.row.col.kind::mxf4nvf4.block_scale.scale_vec::4X.f32.e2m1.e2m1.f32.ue4m3—notably, an Ampere-era instruction, not SM100's tcgen05.mma.
  3. GLM-4.7 FP8: Uses custom-tuned Triton configs for SM120, stored in a specific configuration file.

The GEMM Dimensions Reveal the Structural Problem

The subagent extracted the exact model architecture parameters from the Kimi-K2.5 configuration JSON and computed the per-expert GEMM dimensions after TP=8 sharding. For the down_proj of each expert:

The Dtype Cast Bottleneck

The profiling data from analyze_profile.py revealed the startling finding: 69.3% of CUDA time was spent in unrolled_elementwise_kernel (a dtype cast/copy operation). This corresponded to 64.6 ms per decode step, called 78 times per step—once per transformer layer.

The subagent connected this to the FlashInfer CUTLASS path, hypothesizing that it was internally casting all expert weights per layer. The implication was devastating: the GPU was spending more than two-thirds of its time moving data around and converting formats, not actually computing. The FP4 GEMM utilization numbers confirmed the severity: at 1024 concurrency, the total FP4 peak across 8 GPUs was 14,800 TFLOPS, but the achieved throughput was only 128.6 TFLOPS—0.87% of peak.

Theoretical vs. Actual Throughput

The subagent presented a comparison that quantified the gap. For single-stream GLM-5-NVFP4:

The Structural Bottleneck Diagnosis

The most important contribution of message 8 was the diagnosis of a structural bottleneck. This finding had profound implications for the optimization strategy.

If the bottleneck were computational—if the GEMM kernels were simply not fast enough—the solution would be to write better kernels. This is the kind of problem that kernel engineers love: optimize tile sizes, adjust warp scheduling, exploit new hardware instructions. It is difficult but well-understood work.

If the bottleneck were communication—if AllReduce were the dominant cost—the solution would be to reduce communication volume, fuse operations, or overlap communication with computation. This is also well-understood, with established techniques like AllReduce fusion and async communication.

But the bottleneck identified in message 8 was different. The problem was that the GEMM dimensions themselves were too small for the hardware to achieve reasonable utilization. This was not a kernel problem or a communication problem—it was a parallelism strategy problem. The choice of TP=8, combined with the model's architecture (384 experts, intermediate_size=2048), produced per-expert GEMMs that were inherently inefficient.

The solution space for structural bottlenecks is different. Options include:

The Thinking Process: How the Subagent Reasoned

The subagent's reasoning process is visible in the structure and content of its investigation. Several patterns emerge.

Bottom-Up Synthesis

The subagent did not start with a hypothesis and seek to confirm it. Instead, it read broadly across many files and let the patterns emerge from the data. The progression from kernel paths → GEMM dimensions → quantization details → profiling data → theoretical analysis mirrored the logical flow of understanding: first you must know what code is running, then what data it is processing, then how that data is represented, then how fast it actually runs, then how that compares to what is theoretically possible.

This bottom-up approach was appropriate for an exploration task. The subagent did not know in advance what it would find. It had to be open to surprises—and the 69.3% dtype cast finding was indeed a surprise that reshaped the entire analysis.

Connecting Disparate Sources

The subagent demonstrated the ability to connect information across files. The profiling data showing 69.3% dtype cast time came from analyze_profile.py. The kernel path information came from source/mxfp4.py, decode_gap_analysis.py, and the improvement proposal documents. The model architecture parameters came from configs/kimi-k25-config.json. The service configurations came from the .service files. The subagent wove these together into a coherent picture that no single file provided.

This synthesis was the core value of the investigation. The individual files contained raw data; message 8 contained understanding.

Identifying Structural vs. Tuning Problems

A key insight in the subagent's reasoning was the distinction between problems that can be fixed by tuning parameters and problems that are structural. The small GEMM dimensions after TP sharding are structural: they are a direct consequence of the decision to use TP=8 with a model that has 384 experts and a small intermediate size. No amount of kernel tuning can make a 256-wide GEMM achieve high utilization on a GPU designed for massive parallelism.

This distinction is crucial because it determines what kind of solutions are viable. Tuning problems can be addressed by adjusting kernel parameters, tile sizes, or launch configurations. Structural problems require changes to the parallelism strategy, the model architecture, or the hardware configuration.

Assumptions and Blind Spots

Every analysis rests on assumptions, and this investigation contained several that deserve examination.

Assumption that the profiling data is representative. The profiling data came from analyze_profile.py, which captured 30 decode steps of GLM-5-NVFP4. The subagent assumed that these 30 steps were representative of the model's behavior across different inputs, batch sizes, and system states. In practice, profiling data can vary significantly depending on GPU clock speed, thermal state, memory controller contention, and NCCL topology.

Assumption that the dtype cast is attributable to FlashInfer CUTLASS. The subagent attributed the 69.3% dtype cast time to the FlashInfer CUTLASS path, stating that it "corresponds to ~1.5GB of data being cast per call at HBM bandwidth, suggesting the flashinfer CUTLASS path is internally casting all expert weights per layer." This was a reasonable inference, but it was not proven. The dtype cast could also come from the MoE routing logic, the attention mechanism, or the residual connection handling.

Assumption that the theoretical minimum is achievable. The theoretical minimum TPOT of 3.24 ms was computed assuming perfect HBM bandwidth utilization and no overhead from kernel launches, synchronization, or data movement. In practice, no real system achieves this bound. The subagent acknowledged this implicitly by presenting the gap analysis, but the theoretical minimum served as an aspirational target rather than a realistic expectation.

Assumption that TP=8 is the right baseline. The analysis assumed TP=8 as the parallelism configuration and evaluated GEMM dimensions after TP=8 sharding. But TP=8 is a choice, not a given. The improvement proposal documents in the project directory discussed alternatives like EP8 (expert parallelism) and DP2+TP4. The subagent did not evaluate whether a different parallelism configuration would produce better GEMM dimensions—it took TP=8 as the baseline and analyzed its consequences.

The Broader Significance

This subagent investigation exemplifies how systematic research can reframe a performance optimization problem. The team had been focused on kernel tuning and AllReduce optimization, but the subagent's analysis revealed that the fundamental issue was structural: the parallelism strategy itself was producing GEMM dimensions that could never achieve high utilization on Blackwell GPUs.

The investigation also demonstrated the value of open-ended exploration. The subagent was not told what to find. It was told to explore, and it returned with a diagnosis that reframed the entire optimization problem. This is the power of giving an AI system the freedom to read, synthesize, and report without being constrained by a narrow task definition.

For anyone working on large-scale GPU inference optimization, this investigation offers a template for how to approach bottleneck analysis: start with the kernel paths, compute the actual dimensions, examine the quantization details, measure the real performance, compare to theoretical limits, and synthesize everything into a structural understanding. The answer is rarely in a single number or a single file. It emerges from the connections between them.

Conclusion

The SM120 GEMM optimization campaign was a masterclass in systematic technical investigation. Over nine messages, a subagent read approximately 30 files, conducted targeted code searches, verified primary sources, and synthesized everything into a comprehensive analysis that fundamentally reframed the optimization problem.

The key findings were stark: per-expert GEMM dimensions after TP=8 sharding were so small (K=256) that the GPU achieved less than 1% of its theoretical peak FP4 throughput. A dtype cast operation consumed 69.3% of CUDA time, suggesting that the kernel implementation was spending most of its time moving data rather than computing. The gap between theoretical and actual throughput was a factor of 30x.

But the most important contribution was the structural diagnosis. The problem was not that the kernels were slow—it was that the parallelism strategy (TP=8 with 384 experts and a small intermediate size) produced GEMM dimensions that were inherently incapable of saturating the GPU. This reframing opened up a new axis of optimization: changing the parallelism strategy through expert parallelism, data parallelism, or reduced TP degree, rather than continuing to optimize kernel parameters within a fundamentally broken configuration.

The investigation concluded with the team pivoting to explore speculative decoding as a software-only optimization path—a decision informed by the deep understanding that the structural bottleneck could not be fixed by kernel tuning alone. The SM120 GEMM optimization campaign stands as a testament to the power of systematic research in the complex, high-stakes world of large-scale ML inference.