The 99KB Wall: Tracing a CUTLASS Shared Memory Limit on Blackwell GPUs

In the middle of a deep investigation into why certain CUTLASS tile configurations crash during MoE (Mixture-of-Experts) inference on NVIDIA Blackwell GPUs, the assistant issues a short but pivotal message. Message 919 reads in full:

There's a static_assert checking SharedStorageSize <= sm120_smem_capacity_bytes. Let me check that capacity constant:

`` [bash] ssh root@10.1.230.174 'grep -rn "sm120_smem_capacity_bytes\|sm120_smem" /root/ml-env/lib/python3.12/site-packages/flashinfer/data/cutlass/include/ 2>/dev/null' /root/ml-env/lib/python3.12/site-packages/flashinfer/data/cutlass/include/cutlass/arch/arch.h:45:constexpr int sm120_smem_capacity_bytes = 101376; /root/ml-env/lib/python3.12/site-packages/flashinfer/data/cutlass/include/cutlass/gemm/collective/builders/sm120_blockscaled_sparse_mma_builder.inl:307: detail::sm120_smem_capacity_bytes, sparse_elem<ElementAMma::sparsity, SmemAllocTypeA>, ``

This message is the culmination of a multi-step forensic trace through the CUTLASS kernel compilation and initialization pipeline. It represents the moment the investigation transitions from "what is failing?" to "why it cannot work." Understanding this message requires reconstructing the chain of reasoning that led to it and the conclusions that followed from it.

The Investigation Arc

The story begins several messages earlier ([msg 906]), when the user asks the assistant to "fix tiles" — specifically, two CUTLASS tile configurations for FP4 block-scaled GEMM on SM120 (Blackwell architecture): M128×N256×K128 and M256×N128×K128. These tiles crash during initializeMoeGroupedGemm with a generic "Error Internal" status, while the smaller 128×128×128 and 128×128×256 tiles work fine. The assistant sets out to understand why.

The assistant traces the error through the TRT-LLM grouped GEMM launcher code ([msg 907]–[msg 909]), finding that can_implement passes but gemm.initialize() returns kErrorInternal. This points to a runtime shared memory allocation failure — the kernel requests more dynamic shared memory than the device can provide, even after cudaFuncSetAttribute is called to request the opt-in maximum.

A critical discovery comes in [msg 912], where the assistant queries the GPU's shared memory limits:

| Parameter | Value | |---|---| | Default shared memory per block | 48 KB (49,152 bytes) | | Opt-in max per block | 99 KB (101,376 bytes) | | Total per SM | 100 KB (102,400 bytes) |

These numbers are hardware-imposed. The Blackwell SM120 architecture provides 100 KB of shared memory per SM, of which at most 99 KB can be allocated to a single block (the remaining 1 KB is reserved for other uses). Any kernel that needs more than 99 KB of shared memory simply cannot launch.

The SMEM Budget Calculation

The assistant then writes a Python script to estimate shared memory requirements for each tile configuration ([msg 913]). For FP4 block-scaled GEMM, each matrix operand is packed at 0.5 bytes per element (FP4), with FP8 scale factors every 16 elements. The estimates reveal:

128×128×128 (working): 50 KB with default epilogue, 82 KB with finalize — fits comfortably.
128×128×256 (working): 68 KB default, 100 KB finalize — barely fits, explaining why the finalize variant fails.
128×256×128 (failing): 91 KB default, 155 KB finalize — even the default epilogue barely fits, finalize is impossible.
256×128×128 (failing): Same as above. The problem is geometric: doubling the output tile dimension (from 128 to 256 in either M or N) doubles the epilogue storage needed for the output accumulator. The 128×256 output tile alone needs 65 KB for BF16 staging, leaving only 34 KB for the mainloop pipeline — not enough for even a single stage of the input matrices (27.6 KB).

Message 919: Confirming the Constant

This brings us to message 919. The assistant has just found a static_assert in the non-grouped CUTLASS kernel code ([msg 918]) that checks SharedStorageSize <= sm120_smem_capacity_bytes. This is a compile-time safety net: if the kernel's total shared storage exceeds the hardware capacity, compilation fails with a clear error message. The fact that this check exists confirms that CUTLASS developers anticipated SMEM overflow as a real concern on SM120.

The assistant's next logical step is to check the actual value of sm120_smem_capacity_bytes. The grep command searches the entire CUTLASS include tree for the constant definition and its usage. The results are revealing:

The constant is defined in cutlass/arch/arch.h at line 45 as constexpr int sm120_smem_capacity_bytes = 101376. This is the canonical hardware limit — 99 KB.
The constant is used in sm120_blockscaled_sparse_mma_builder.inl at line 307, where it's passed to sm100_compute_stage_count_or_override_blockscaled. This function calculates how many pipeline stages can fit within the SMEM budget. The assistant's reasoning here is precise: if the constant is 101376 (99 KB), and the failing tiles need 91 KB for the mainloop alone (with 1 stage) plus an epilogue that requires at least 65 KB for the output tile, then the total exceeds the budget regardless of how stages are reduced. The StageCountAutoCarveout mechanism — which subtracts the epilogue's shared storage from the available SMEM before computing the number of pipeline stages — would compute a negative or zero stage count, causing the initialization to fail.

The Critical Distinction: Grouped vs. Non-Grouped Kernels

A key insight that emerges from this message and its immediate successor ([msg 920]) is the difference between how the non-grouped and grouped (MoE) CUTLASS kernels handle SMEM overflow:

Non-grouped kernels have a static_assert at compile time that catches the overflow. The kernel simply won't compile if the tile configuration is too large.
Grouped GEMM kernels (used for MoE) go through the TRT-LLM launcher, which uses runtime initialization via gemm.initialize(). There is no static_assert, so the overflow manifests as a runtime kErrorInternal error during initializeMoeGroupedGemm. This is precisely the error the assistant has been chasing. This distinction matters because runtime errors are much harder to diagnose than compile-time errors. The assistant had to trace through multiple layers of code — from the TRT-LLM launcher, through the CUTLASS collective builder, to the architecture constants — to discover that the root cause is a fundamental hardware limitation, not a software bug or configuration mistake.

The Pivot

The consequence of message 919 is immediate and consequential. In [msg 922], the assistant updates its todo list:

Fix failing CUTLASS M128x256 and M256x128 tiles for SM120 — SMEM overflow is fundamental (99KB limit), tiles need 91KB+ mainloop + epilogue. Would require reducing epilogue SMEM or splitting epilogue into separate kernel. DEPRIORITIZED.

The tile fix is abandoned as a dead end. The assistant correctly recognizes that no amount of software tuning can overcome a hardware-imposed 99 KB shared memory limit. The only way to make these tiles work would be to redesign the epilogue to use less shared memory (perhaps by splitting it into a separate kernel launch), which is a major engineering effort far beyond the scope of the current session.

The assistant pivots to the other approaches the user suggested: increasing effective batch size per expert (via num_continuous_decode_steps and larger max_running_requests) and trying alternative kernel paths like cuBLASLt FP4. These turn out to be far more productive — the assistant later achieves a 28% throughput improvement through server parameter tuning alone ([chunk 7.0]).

Input Knowledge Required

To fully understand message 919, one needs:

CUTLASS architecture: Understanding that CUTLASS uses template metaprogramming to generate GPU kernels, with tile shapes, pipeline stages, and shared memory layouts all resolved at compile time or kernel initialization time.
SM120 shared memory hierarchy: Blackwell GPUs have 100 KB of shared memory per SM, with a 99 KB opt-in maximum per block. This is a decrease from Hopper (SM90) which offered 228 KB per SM, making SMEM pressure a new constraint for Blackwell kernel designers.
FP4 block-scaled GEMM: The model uses FP4 weights with FP8 block scaling factors. This means each matrix operand has both data (FP4, 0.5 bytes/element) and scales (FP8, 1 byte per 16 elements), doubling the shared memory pressure compared to a standard FP8 GEMM.
MoE grouped GEMM: Mixture-of-Experts models dispatch tokens to different experts, requiring a grouped GEMM where multiple small matrix multiplications are batched together. The TRT-LLM launcher handles this grouping, but uses a different initialization path than standard CUTLASS kernels.
The difference between can_implement and initialize: In CUTLASS, can_implement checks whether the kernel could work given the tile shapes and data types, while initialize actually allocates shared memory and prepares the kernel for launch. A tile can pass can_implement but fail initialize if SMEM is insufficient.

Output Knowledge Created

This message produces several concrete pieces of knowledge:

The exact SMEM capacity constant: sm120_smem_capacity_bytes = 101376 (99 KB), defined in cutlass/arch/arch.h:45.
The constant's usage in the builder: It's passed to sm100_compute_stage_count_or_override_blockscaled to auto-compute the number of pipeline stages that fit within the SMEM budget.
Confirmation that the grouped GEMM path lacks compile-time protection: The static_assert exists only in the non-grouped kernel, explaining why the MoE launcher fails at runtime rather than compile time.
A definitive diagnosis: The tile failure is not a bug or configuration error — it's a fundamental hardware limitation. The 128×256 and 256×128 output tiles require more shared memory than SM120 can provide.

The Broader Significance

Message 919 is a small but crucial node in a much larger investigation spanning multiple sessions and segments. The assistant has been working on deploying the GLM-5-NVFP4 model — a massive Mixture-of-Experts language model — across 8 NVIDIA RTX PRO 6000 Blackwell GPUs. Performance tuning has revealed that the GPUs draw only ~235W out of 600W TDP during inference, and the CUTLASS FP4 kernels achieve merely 0.8–55 TFLOPS for the small per-expert batch sizes typical during decode (see [chunk 7.0]).

The SMEM limitation discovered in this message is one piece of a larger puzzle: Blackwell's reduced shared memory (100 KB vs. 228 KB on Hopper) makes it harder to achieve high utilization for small-batch GEMM operations. The assistant's subsequent work focuses on compensating through other means — increasing batch sizes through server parameter tuning, exploring expert parallelism, and investigating alternative kernel paths.

This message exemplifies a pattern common in systems performance work: the most valuable outcome of an investigation is often not a fix, but a clear understanding of why something cannot be fixed within the current constraints. The 99 KB SMEM wall is real, and knowing it saves the assistant from pursuing a dead end.