The 99KB Barrier: Diagnosing CUTLASS Tile Failures on Blackwell's SM120 Architecture

In the high-stakes world of large language model inference optimization, performance gains often hinge on the smallest details of GPU kernel configuration. Message 920 of this opencode session captures a pivotal diagnostic moment: the assistant has been systematically investigating why certain CUTLASS tile configurations—specifically M128×N256 and M256×N128—fail during grouped GEMM initialization on NVIDIA's Blackwell SM120 architecture. The message is a concise synthesis of findings from a deep-dive investigation spanning multiple tool calls, and it marks the transition from symptom observation to root-cause understanding.

The Context: Chasing Throughput on Blackwell

The broader session involves deploying the GLM-5-NVFP4 model—a large language model using FP4 block-scaled quantization—across 8 NVIDIA RTX PRO 6000 Blackwell GPUs. The assistant has been systematically tuning server parameters and kernel configurations to improve inference throughput. Earlier benchmarking revealed that the model is compute-bound, with GPUs drawing only ~235W out of their 600W TDP during decode. The CUTLASS FP4 GEMM kernels, which handle the matrix multiplications in the Mixture-of-Experts (MoE) layers, plateau at ~1,300 TFLOPS (70% of dense peak) only for very large matrices. During actual inference, per-expert batch sizes of ~16–64 tokens achieve merely 0.8–55 TFLOPS—a tiny fraction of peak capability.

A critical clue emerged when the assistant attempted to use larger CUTLASS tile configurations (M128×N256 and M256×N128) to improve small-batch GEMM efficiency. These tiles failed during initialization with a cryptic "Error Internal" from the initializeMoeGroupedGemm function. The user had explicitly requested this investigation ([msg 905]: "Fix tiles, then try 2/3"), making it the highest-priority task.

Tracing the Error: From Symptom to Root Cause

The assistant's investigation in the preceding messages ([msg 907] through [msg 919]) traced the error through multiple layers of the GPU software stack. The journey began by locating the error source in FlashInfer's copy of the TRT-LLM (TensorRT-LLM) grouped GEMM launcher at moe_gemm_tma_ws_launcher.inl. The error originated from gemm.initialize(args, ...) returning kErrorInternal—a CUTLASS runtime error distinct from the can_implement check, which had passed successfully.

This distinction is crucial. The can_implement method performs a compile-time feasibility check on the kernel configuration, while initialize performs runtime setup including shared memory (SMEM) allocation. The fact that can_implement passed but initialize failed strongly suggested a runtime memory constraint violation—specifically, the kernel was requesting more dynamic shared memory than the device could provide, even after calling cudaFuncSetAttribute to opt into the maximum allocation.

The assistant then queried the GPU's shared memory capabilities ([msg 910]–[msg 912]), discovering the critical numbers:

The Core Discovery: A Missing Static Assert

Message 920 opens with the assistant's synthesis of what these numbers mean for the tile failure problem:

Good — sm120_smem_capacity_bytes = 101376 (99 KB). The static_assert in the CUTLASS non-grouped kernel catches this at compile time. But for the grouped GEMM (MoE), the code path is different — it goes through the TRT-LLM launcher which doesn't have this static_assert, so it fails at runtime during initialize().

This is the key insight. In the previous message ([msg 919]), the assistant discovered a static_assert in the non-grouped CUTLASS kernel header (sm120*.hpp) that checks SharedStorageSize <= sm120_smem_capacity_bytes at compile time. This assertion would prevent any tile configuration that exceeds the 99 KB SMEM limit from even compiling. But the grouped GEMM path—used for MoE layers where multiple small matrix multiplications are batched together—goes through a different code path: the TRT-LLM launcher, which doesn't include this static assertion. Consequently, oversized configurations compile successfully but fail at runtime when CUDA attempts to allocate the requested shared memory.

The assistant then articulates two hypotheses for why the auto-configuration selects an invalid setup:

The key problem is that the TRT-LLM grouped GEMM wrapper uses KernelScheduleAuto for SM120, and CUTLASS auto-selects a configuration that either: 1. Doesn't know about SM120's 99KB limit and picks too many stages 2. The epilogue is union'd with mainloop but the auto-carveout math doesn't account for it properly

These hypotheses reflect a deep understanding of CUTLASS's collective builder architecture. The KernelScheduleAuto mechanism is supposed to automatically select pipeline stages and tile configurations based on the target architecture's shared memory capacity. But the assistant suspects this auto-selection is either unaware of SM120's specific limit (perhaps because the grouped GEMM builder was written before SM120 support was finalized) or miscomputes the carveout—the portion of shared memory reserved for the epilogue (output accumulation) versus the mainloop (input data pipeline).

The SMEM Budget Analysis

The assistant's earlier SMEM estimation ([msg 913]–[msg 914]) had already revealed why the larger tiles fail. For the working 128×128×128 tile with a single pipeline stage, the mainloop requires 18,432 bytes for FP4 input data and scaling factors, plus 32,768 bytes for the default epilogue (BF16 output staging), totaling 51,200 bytes—well within the 99 KB limit. Even with the more expensive FP32 accumulator epilogue, the total reaches 83,968 bytes, still fitting comfortably.

But for the failing 128×256×128 tile, the numbers tell a different story. The output tile is twice as wide (N=256 instead of 128), so the epilogue alone requires 65,536 bytes for BF16 staging. With the mainloop consuming 27,648 bytes for a single stage, the total reaches 93,184 bytes. While this technically fits in 99 KB, it leaves almost no margin. The finalize epilogue (with FP32 accumulation) would need 155,648 bytes—catastrophically over budget.

The 256×128×128 tile faces the same problem in the M dimension instead of N. The auto-carveout mechanism should theoretically reduce to a single pipeline stage and use the minimum epilogue, but the assistant suspects the math doesn't account for SM120's constraints properly in the grouped GEMM path.

The Decision to Read Source

The most important decision in message 920 is the assistant's next investigative step:

Let me look at what the actual CollectiveBuilder resolves to for SM120 block-scaled grouped GEMM

Rather than continuing to speculate about SMEM allocation or attempting to patch the kernel at the binary level, the assistant chooses to read the CUTLASS collective builder source code directly. This is a strategic decision: the sm120_blockscaled_mma_builder.inl file is the authoritative source for how tile configurations are constructed for SM120. By reading it, the assistant can determine exactly what KernelScheduleAuto resolves to, whether the SM120 capacity constant is properly referenced, and where the carveout computation might go wrong.

This approach exemplifies good debugging methodology: when the behavior of a complex system is unexpected, go to the source of truth rather than continuing to infer from symptoms. The assistant could have attempted workarounds (manually setting stage counts, patching the binary, or using a different kernel path), but understanding the root cause in the builder logic is more likely to yield a correct and maintainable fix.

Assumptions and Potential Pitfalls

The assistant makes several assumptions in this message that deserve scrutiny. First, it assumes that the TRT-LLM launcher genuinely lacks the static assertion, which appears confirmed by the earlier grep results ([msg 918]–[msg 919]). However, the static assertion might exist in a different form—perhaps as a runtime check in the launcher's initialization path that the assistant hasn't found yet. The grep only searched for specific patterns; a differently named assertion could exist.

Second, the assistant assumes that KernelScheduleAuto is the root cause of the misconfiguration. While plausible, the issue could also be in the tile dispatch logic itself—the builder might be selecting these larger tiles based on a heuristic that assumes more SMEM is available than SM120 provides. The KernelScheduleAuto might be working correctly given the tile choice, but the tile choice itself might be wrong.

Third, the SMEM estimates in the earlier message were acknowledged as rough: "Actually wait — my epilogue estimate might be wrong. CUTLASS epilogues use different strategies." The assistant correctly recognizes that the actual CUTLASS epilogue implementation might use different memory layouts or sharing strategies that change the SMEM budget calculation. Reading the builder source will resolve this uncertainty.

Input Knowledge Required

To fully understand this message, one needs knowledge of several domains:

Output Knowledge Created

This message creates several valuable outputs:

  1. Root cause identification: The grouped GEMM path lacks the compile-time SMEM capacity check that protects the non-grouped path, allowing oversized configurations to compile but fail at runtime.
  2. Actionable next step: The decision to read the CollectiveBuilder source provides a clear path to understanding the exact mechanism of failure and potentially fixing it.
  3. Documented hypotheses: The two hypotheses about stage counting and epilogue carveout provide a framework for interpreting the builder source code once read.
  4. Boundary understanding: The 99 KB SMEM limit is now firmly established as the hard constraint that all tile configurations must respect, providing a clear design target for any kernel optimization work.

The Thinking Process Visible

The assistant's reasoning in this message reveals a methodical, layered approach to debugging. It starts with the highest-level symptom (tile initialization failure), traces to the specific error location (initialize() returning kErrorInternal), identifies the missing safety check (static_assert in non-grouped path), and then formulates hypotheses about why the auto-configuration fails. Each step builds on evidence gathered from the previous tool call, creating a chain of reasoning that narrows the possible causes.

The assistant also demonstrates intellectual honesty by acknowledging when its estimates might be wrong ("my epilogue estimate might be wrong"). This self-correction is essential in debugging—holding onto incorrect assumptions can lead to wasted effort pursuing the wrong fix.

Significance in the Larger Narrative

Message 920 represents a turning point in the optimization effort. The assistant has moved from observing that larger tiles fail to understanding why they fail. This understanding is necessary before any fix can be attempted—whether that fix involves patching the builder to properly account for SM120's limits, forcing a different tile selection, or accepting that these tiles simply cannot work on Blackwell and pursuing alternative optimization strategies.

The 99 KB SMEM barrier is a fundamental architectural constraint of SM120, and understanding it shapes all subsequent optimization work. It means that any kernel using larger output tiles (M or N > 128) must be extremely careful about SMEM budgeting, potentially using fewer pipeline stages or more aggressive memory sharing strategies. It also means that the path to higher throughput on Blackwell may lie not in larger tiles but in better utilization of the working 128×128 configurations—perhaps through batching strategies like num_continuous_decode_steps that increase the effective batch size per GEMM call.

This message exemplifies the kind of deep, systematic investigation that high-performance GPU kernel optimization demands: tracing errors through multiple software layers, understanding hardware constraints, and using source code as the ultimate authority for system behavior.