The 99KB Wall: Tracing a CUTLASS Tile Failure on Blackwell SM120

Introduction

In the middle of an intensive optimization session for the GLM-5-NVFP4 model running on eight NVIDIA RTX PRO 6000 Blackwell GPUs, the assistant encountered a puzzling failure: certain CUTLASS tile configurations for FP4 block-scaled matrix multiplication crashed during kernel initialization with a cryptic "Error Internal" message. Message [msg 921] captures the precise moment when the assistant traced this failure to its root cause—a shared memory capacity constraint on the Blackwell SM120 architecture that manifests differently in the grouped GEMM (MoE) kernel path than in standalone CUTLASS kernels.

This message is a turning point in the optimization narrative. It represents the culmination of a deep debugging dive into NVIDIA's CUTLASS library, the TensorRT-LLM launcher infrastructure, and the physical memory limits of the Blackwell GPU architecture. More importantly, it demonstrates the moment when the assistant correctly diagnosed a fundamental hardware-imposed constraint and made the strategic decision to pivot toward more promising optimization avenues.

The Debugging Context

To understand message [msg 921], we must first appreciate the chain of investigation that led to it. The assistant had been systematically benchmarking the GLM-5-NVFP4 model—a Mixture-of-Experts (MoE) architecture using FP4 block-scaled quantization—and discovered that the GPUs were drawing only ~235W out of their 600W TDP during inference. This suggested severe underutilization of the compute hardware.

The core problem lay in the FP4 GEMM (General Matrix Multiply) kernels used during the decode phase. During inference, each expert in the MoE layer processes a relatively small batch of tokens—approximately 16–64 tokens per expert at 1024 concurrency. The CUTLASS kernels that perform these matrix multiplications achieve their peak efficiency (~1,300 TFLOPS, or 70% of the dense theoretical peak of ~1,850 TFLOPS) only for very large matrices. For the small per-expert batches typical during decode, the kernels achieve a meager 0.8–55 TFLOPS—between 0.02% and 3% of peak.

The assistant had identified that larger CUTLASS tile configurations—specifically M128×N256 and M256×N128—could potentially improve small-batch GEMM efficiency by processing more output elements per threadblock. However, these tiles failed to initialize during initializeMoeGroupedGemm, throwing an "Error Internal" status. The investigation in messages [msg 907] through [msg 920] systematically narrowed down the cause: the TRT-LLM grouped GEMM launcher, which wraps CUTLASS for MoE computation, was failing during kernel initialization because the requested tile configurations exceeded available shared memory.

The Critical Discovery

Message [msg 921] opens with the assistant zeroing in on the critical line of code in the CUTLASS collective builder for SM120:

static constexpr int PipelineStages = cutlass::gemm::collective::detail::sm100_compute_stage_count_or_override_blockscaled<
    detail::sm120_smem_capacity_bytes, SmemAllocTypeA, SmemAllocTypeB, TileShape_MNK, SmemLayoutAtomSFA, SmemLayoutAtomSFB>(StageCountType{});

This line is the heart of the pipeline stage calculation for the CUTLASS kernel. It uses sm120_smem_capacity_bytes—a constant set to 101,376 bytes (99 KB)—to compute how many pipeline stages can fit within the available shared memory. The function sm100_compute_stage_count_or_override_blockscaled (note: named for SM100 but reused for SM120) takes the SMEM capacity, the memory allocation types for operands A and B, the tile shape, and the shared memory layout atoms, and returns the maximum number of pipeline stages that can be sustained.

The assistant immediately recognizes that "the builder itself is correct." The CUTLASS collective builder properly accounts for SM120's 99 KB shared memory limit. The problem lies in how the TRT-LLM launcher invokes this builder. The launcher passes StageCountAutoCarveout as the stage count type, which instructs CUTLASS to automatically compute the stage count while carving out space for the epilogue's shared storage. The epilogue—the portion of the kernel that handles output computation, bias addition, activation functions, and scatter operations—requires its own shared memory allocation.

The assistant's key insight is that "if that calculation results in 0 or negative stages, the kernel can't be initialized." This is precisely what happens for the larger tile configurations. The epilogue shared storage for the FINALIZE variant (which includes router scales, bias, and scatter operations) consumes so much shared memory that the carveout leaves insufficient room for even a single pipeline stage of the mainloop computation.

Shared Memory Arithmetic on Blackwell

To understand why this happens, we need to examine the shared memory budget on SM120. Earlier measurements in the conversation (message [msg 912]) established the following constraints:

The TRT-LLM Launcher Discrepancy

The assistant's deeper insight in message [msg 921] is that the failure is specific to the grouped GEMM path used for MoE computation. In standalone CUTLASS kernels (the non-grouped path), a static_assert at compile time catches SMEM overflow:

static_assert(SharedStorageSize <= cutlass::arch::sm120_smem_capacity_bytes, 
              "SMEM usage exceeded capacity.");

This assertion, discovered in message [msg 919], ensures that any kernel that exceeds the 99 KB limit fails at compile time with a clear error message. However, the TRT-LLM grouped GEMM launcher bypasses this check because it uses a different kernel template instantiation path—one that doesn't include the static_assert. Instead, the failure occurs at runtime during gemm.initialize(args, ...), which returns cutlass::Status::kErrorInternal when the SMEM allocation fails.

The assistant hypothesizes that the grouped GEMM kernel through TRT-LLM uses a different epilogue than the standalone CUTLASS builder expects. The StageCountAutoCarveout mechanism subtracts sizeof(CollectiveEpilogue::SharedStorage) from the available SMEM budget, and if the epilogue is larger than what the builder's stage count calculation assumes, the carveout can result in zero or negative remaining stages.

The Strategic Pivot

Message [msg 921] ends with the assistant preparing to examine the TRT-LLM launcher's epilogue construction in detail, issuing a bash command to read lines 395–445 of the launcher file. However, the very next message ([msg 922]) reveals a critical realization: the SMEM overflow is fundamental. The 99 KB limit on SM120 is a hardware constraint that cannot be circumvented without either:

  1. Reducing the epilogue's shared memory footprint (e.g., splitting the epilogue into a separate kernel)
  2. Using smaller tile configurations that fit within the budget The assistant correctly judges that attempting to fix the large tiles would require significant engineering effort—modifying the TRT-LLM launcher to use a different epilogue strategy or splitting the fused epilogue into a separate kernel pass—with uncertain payoff. Instead, the assistant deprioritizes the tile fix and pivots to more impactful optimization directions: increasing per-expert batch sizes through num_continuous_decode_steps, trying cuBLASLt's FP4 path for small GEMMs, and tuning server parameters. This pivot proved successful. In subsequent messages, the assistant achieved a 28% throughput improvement by raising --max-running-requests to 2048 and setting --num-continuous-decode-steps 8, reaching 2,095 output tok/s and 4,151 total tok/s at 2048 concurrency.

Input Knowledge Required

To fully appreciate message [msg 921], the reader needs substantial background knowledge:

  1. CUTLASS architecture: Understanding of how CUTLASS decomposes GEMM operations into threadblocks, warps, and pipeline stages, and how the collective builder selects tile configurations
  2. Shared memory hierarchy on NVIDIA GPUs: The distinction between per-block and per-SM shared memory, the opt-in mechanism for exceeding the default 48 KB limit, and how cudaFuncSetAttribute works
  3. Blackwell SM120 architecture: The specific shared memory capacity (99 KB opt-in, 100 KB total per SM) and how it differs from Hopper (SM90) and other architectures
  4. MoE grouped GEMM: How Mixture-of-Experts inference uses grouped matrix multiplications where each expert's weight matrix is multiplied by a different subset of tokens, requiring specialized kernel launchers
  5. TRT-LLM integration: How TensorRT-LLM wraps CUTLASS kernels for production inference, including the TmaWarpSpecializedGroupedGemm launcher and its epilogue variants
  6. FP4 block-scaled quantization: The memory layout of FP4 data (0.5 bytes per element) with FP8 scale factors (1 byte per 16 elements), which affects SMEM requirements

Output Knowledge Created

Message [msg 921] produces several valuable pieces of knowledge:

  1. Root cause identification: The failing tiles (M128×N256, M256×N128) fail because the TRT-LLM launcher's StageCountAutoCarveout subtracts epilogue shared storage from the 99 KB SMEM budget, leaving insufficient room for even one pipeline stage
  2. Architectural constraint confirmation: The 99 KB SMEM limit on SM120 is a hard barrier that cannot be exceeded, and the larger tile configurations' epilogue requirements fundamentally exceed this budget
  3. Diagnostic methodology: The assistant demonstrates a systematic approach to tracing kernel initialization failures—from runtime error to source code to hardware specification—that serves as a template for similar debugging efforts
  4. Strategic decision framework: The message implicitly establishes a cost-benefit analysis for optimization work: fixing the tiles would require modifying the TRT-LLM launcher's epilogue construction (high engineering cost, uncertain payoff), while other approaches offer more immediate gains

The Thinking Process

The assistant's reasoning in message [msg 921] reveals a sophisticated debugging methodology. Rather than treating the "Error Internal" message as a black-box failure, the assistant traces through multiple layers of abstraction:

  1. From error to code: The error originates in gemm.initialize(args, ...) returning kErrorInternal, which the assistant traces to the CUTLASS collective builder
  2. From code to constraint: The builder correctly uses sm120_smem_capacity_bytes (99 KB), so the builder itself isn't buggy
  3. From constraint to mismatch: The discrepancy must lie in how the TRT-LLM launcher interacts with the builder—specifically, the StageCountAutoCarveout mechanism that subtracts epilogue storage
  4. From mismatch to hypothesis: The epilogue in the grouped GEMM path is likely larger than what the standalone builder expects, causing the carveout to leave zero or negative stages This chain of reasoning demonstrates a deep understanding of the software stack, from the high-level Python benchmarking scripts down to the C++ template metaprogramming in CUTLASS. The assistant doesn't just find a bug—it identifies a fundamental architectural tension between the hardware's shared memory limits and the software's epilogue requirements.

Conclusion

Message [msg 921] is a masterclass in GPU kernel debugging. It captures the moment when a seemingly mysterious runtime error is traced to its physical root: the 99 KB shared memory wall on Blackwell SM120. More importantly, it demonstrates the wisdom of knowing when to stop digging. Rather than attempting an expensive workaround to force the large tiles to fit, the assistant correctly judges that the constraint is fundamental and pivots to more productive optimization strategies. This decision ultimately led to a 28% throughput improvement, validating the assistant's engineering judgment.

The message also serves as a valuable case study in the complexities of production ML inference optimization. The interaction between CUTLASS's template-based kernel generation, TRT-LLM's production launcher infrastructure, and the physical hardware constraints of the Blackwell architecture creates a rich landscape where failures can propagate through multiple abstraction layers before manifesting as cryptic runtime errors. Tracing these failures requires not just knowledge of each layer, but an understanding of how they interact—exactly the kind of systems thinking that message [msg 921] exemplifies.