The 99KB Barrier: Tracing a CUTLASS Shared Memory Overflow on Blackwell SM120

Message Snapshot

[assistant] This is the mainloop only. The total SMEM is mainloop + epilogue. Let me directly check what CUTLASS computes for the failing configs by writing a small CUDA program, or better yet — look at the GemmKernel::SharedStorageSize: ``bash ssh root@10.1.230.174 'grep -rn "SharedStorageSize\|sizeof.*SharedStorage\|kSmemSize\|get_smem_size\|smem_size" /root/ml-env/lib/python3.12/site-packages/flashinfer/data/cutlass/include/cutlass/gemm/kernel/sm120*.hpp | head -20' ``

>

Output: `` 180: static constexpr int SharedStorageSize = sizeof(SharedStorage); 181: static_assert(SharedStorageSize <= cutlass::arch::sm120_smem_capacity_bytes, "SMEM usage exceeded capacity."); ``

At first glance, this message from the assistant (global index 918) appears to be a routine debugging step: a grep command followed by its output. But in the context of the larger investigation — a multi-session effort to squeeze maximum inference throughput from the GLM-5-NVFP4 model on eight NVIDIA RTX PRO 6000 Blackwell GPUs — this message represents a critical inflection point. It is the moment when the assistant transitions from estimating the root cause of a crashing CUTLASS kernel to verifying it against the compiler's own static analysis, and in doing so, uncovers a subtle mismatch between compile-time guarantees and runtime behavior that reshapes the entire optimization strategy.

The Investigation Leading Up to This Message

To understand why this message was written, one must appreciate the path that led to it. The assistant had been systematically benchmarking the GLM-5-NVFP4 model — a Mixture-of-Experts architecture using FP4 block-scaled quantization — across eight RTX PRO 6000 GPUs. Throughput had plateaued at roughly 1,640 output tokens per second at 2048 concurrency ([msg 902]), and GPU power draw was a puzzling ~235W per card — barely 39% of the 600W TDP ([msg 903]). The GPUs were clearly not being fed enough compute work.

The bottleneck traced back to the FlashInfer MoE grouped GEMM kernels, which dispatch per-expert matrix multiplications during the decode phase. These kernels use CUTLASS 3.x tile configurations tailored for SM120 (Blackwell's compute architecture). The assistant discovered that two promising tile configurations — M128×N256×K128 and M256×N128×K128 — were crashing during initializeMoeGroupedGemm with a kErrorInternal status ([msg 907]). Only the smaller 128×128×128 and 128×128×256 tiles worked.

The working hypothesis was that the larger output tiles (128×256 or 256×128) required more shared memory per block than the device could provide. The assistant had already gathered the critical device parameters: the RTX PRO 6000 Blackwell offers 100 KB of shared memory per SM, with a maximum opt-in per block of 99 KB (101,376 bytes) ([msg 912]). Manual calculations suggested that even a single pipeline stage for the 128×256 output tile would demand approximately 91 KB for the mainloop plus epilogue, leaving almost no headroom ([msg 913]). But these were estimates — rough calculations based on data type sizes and assumed epilogue strategies.

The Pivot: From Estimation to Compile-Time Verification

Message 918 marks the pivot from estimation to verification. The assistant explicitly acknowledges the limitation of the manual analysis: "This is the mainloop only. The total SMEM is mainloop + epilogue." This self-correction is significant — the earlier SMEM calculations had focused on the mainloop pipeline stages but had not fully accounted for the epilogue's shared storage requirements. The assistant then proposes two approaches: writing a small CUDA program to probe the failing configurations, or — more elegantly — examining what CUTLASS itself computes at compile time via GemmKernel::SharedStorageSize.

The decision to look at the CUTLASS source rather than writing a CUDA program reveals the assistant's engineering judgment. Writing a CUDA program would be time-consuming: it would require compiling a test harness with the exact tile template parameters, linking against CUTLASS, and running it on the target machine. The source code, by contrast, is already present in the FlashInfer package installed on the server. A simple grep can reveal whether CUTLASS has a compile-time shared memory check, and if so, what the computed size is relative to the device capacity. This is a classic "inspect the compiler's work" strategy — faster, more direct, and less error-prone than constructing a separate test.

The grep command targets sm120*.hpp files in the CUTLASS kernel directory, searching for symbols related to shared memory size computation: SharedStorageSize, sizeof.*SharedStorage, kSmemSize, get_smem_size, and smem_size. The output reveals two critical lines:

static constexpr int SharedStorageSize = sizeof(SharedStorage);
static_assert(SharedStorageSize <= cutlass::arch::sm120_smem_capacity_bytes, "SMEM usage exceeded capacity.");

What This Discovery Means

This discovery carries profound implications for the investigation. The static_assert means that CUTLASS performs a compile-time check ensuring that the total SharedStorage structure — which includes both mainloop pipeline storage and epilogue storage — fits within the device's SMEM capacity (sm120_smem_capacity_bytes). If a tile configuration violated this constraint, the kernel would fail to compile at all, not merely crash at runtime.

The fact that the 128×256 and 256×128 tiles compile successfully but crash at runtime means the SMEM size is within the compile-time limit but something else is going wrong. This invalidates the assistant's working hypothesis that the tiles simply require too much shared memory. The static_assert would have caught that. The runtime failure must have a different cause — perhaps a dynamic shared memory allocation that exceeds the per-block limit after cudaFuncSetAttribute, or a TMA (Tensor Memory Accelerator) descriptor alignment issue, or a warp-specialized pipeline deadlock specific to the larger tile shapes.

This is a classic debugging pattern: the obvious hypothesis (SMEM overflow) is ruled out by a compiler guarantee, forcing the investigator to look deeper. The assistant does not yet articulate this conclusion in message 918 — the message is purely about gathering the evidence — but the reader can sense the shift. The grep output is presented without commentary, but its significance is immediately clear to anyone familiar with CUTLASS's design: the static_assert is the compiler's promise that the kernel should work, and the runtime crash represents a broken promise that must be investigated at a different layer.

Assumptions and Their Validity

The assistant makes several assumptions in this message. First, it assumes that the SharedStorageSize constant is the definitive measure of SMEM usage — that sizeof(SharedStorage) accurately reflects the runtime shared memory footprint. This is generally correct for CUTLASS 3.x kernels, which use a union-based SharedStorage structure that encompasses all pipeline stages, tensor maps, and epilogue storage. However, there is a subtlety: sizeof(SharedStorage) captures the static layout, but CUTLASS may also request dynamic shared memory at kernel launch via cudaFuncSetAttribute. The static_assert checks only the static portion; the dynamic allocation is set separately and could theoretically exceed the device limit even if the static portion fits. This nuance is not addressed in the message.

Second, the assistant assumes that the sm120_smem_capacity_bytes constant accurately reflects the usable shared memory limit. This is a reasonable assumption — CUTLASS is a NVIDIA-maintained library and its architecture constants are typically correct. But the value of this constant is not shown in the grep output, and the assistant does not verify it against the device properties gathered earlier (99 KB opt-in per block). A discrepancy between sm120_smem_capacity_bytes and the actual cudaDevAttrMaxSharedMemoryPerBlockOptin could explain the runtime failure: CUTLASS might think it has more SMEM available than the driver actually grants.

Third, the assistant assumes that the failing tiles are compiled with the same SharedStorage structure as the working tiles. This is a safe assumption given that they share the same kernel template, differing only in tile dimension parameters. However, CUTLASS uses partial specialization and SFINAE (Substitution Failure Is Not An Error) extensively; it is possible that different tile dimensions trigger different template instantiations with different SharedStorage layouts.

Input Knowledge Required

To fully understand this message, the reader needs several layers of context. First, they need to know that CUTLASS is NVIDIA's CUDA template library for linear algebra subroutines, and that its SM120 backend implements warp-specialized kernels using TMA (Tensor Memory Accelerator) for data movement. The SharedStorage structure is a union of all temporary buffers needed by the kernel — mainloop pipeline stages, epilogue accumulators, tensor map descriptors — and its size directly determines the shared memory footprint per thread block.

Second, the reader must understand the MoE grouped GEMM architecture: during inference, each token is routed to a subset of experts (typically 2 out of 64 in GLM-5), and the per-expert batch size is small (roughly 32 tokens at 1024 concurrency). This is why larger CUTLASS tiles are desirable — they allow each thread block to process more work, improving arithmetic intensity. The failing 128×256 and 256×128 tiles would have doubled the output dimension compared to the working 128×128 tile, potentially improving throughput for the small-batch regime.

Third, the reader needs to know the Blackwell SM120 architecture specifics: 100 KB shared memory per SM, 99 KB opt-in per block, 188 SMs on the RTX PRO 6000, and the FP4 block-scaled data format where each 2-bit element is packed with an FP8 scale factor shared across 16 elements. These details were established in earlier messages (<msg id=910-913>) and form the quantitative foundation for the SMEM analysis.

Fourth, the reader must understand the difference between compile-time and runtime checks in CUDA. The static_assert is evaluated by the host compiler during kernel template instantiation. If it passes, the kernel is compiled into device code. But the compiled kernel can still fail at launch time if the requested dynamic shared memory exceeds the per-block limit, or if the TMA descriptors are misaligned, or if the warp scheduler encounters a resource conflict. The assistant's grep targets compile-time checks, but the failure mode is runtime — a distinction that becomes crucial in the subsequent investigation.

Output Knowledge Created

This message produces several valuable outputs. First and most concretely, it confirms that CUTLASS has a compile-time SMEM capacity check via static_assert. This is a design insight that the assistant can leverage: any tile configuration that compiles successfully is guaranteed by the library to fit within SMEM, at least according to CUTLASS's own calculations. This rules out the simplest explanation for the crash.

Second, the message identifies the exact file and line numbers where the SMEM check occurs: sm120_blockscaled_mma_array_tma.hpp at lines 180-181. This gives the assistant a precise location to examine the SharedStorage structure and understand its layout. The next logical step — which the assistant pursues in subsequent messages — is to examine the SharedStorage struct itself, compute its size for each tile configuration, and compare it against sm120_smem_capacity_bytes to see if there is a discrepancy.

Third, the message establishes a methodology: when investigating CUTLASS kernel failures, look at the compiler's own bounds checks before building custom test programs. This is a reusable debugging pattern that the assistant can apply to other kernel failures in the future.

Fourth, and perhaps most importantly for the narrative arc of the session, this message creates the knowledge that the SMEM hypothesis is incomplete. The static_assert passing means something else is causing the crash. This forces the investigation to broaden — from shared memory capacity to dynamic SMEM allocation, TMA descriptor alignment, warp specialization scheduling, or even a bug in the CUTLASS SM120 backend itself. The assistant does not yet know which direction to go, but it now knows that the obvious answer is wrong.

The Thinking Process Revealed

The assistant's reasoning in this message is visible in the structure of the grep command. The search pattern is carefully constructed: it includes not just SharedStorageSize (the exact symbol) but also sizeof.*SharedStorage (to catch any direct sizeof usage), kSmemSize (a common naming convention for SMEM constants), get_smem_size (a possible accessor function), and smem_size (a generic fallback). This breadth suggests the assistant is uncertain about the exact naming convention used in CUTLASS and is casting a wide net. The | head -20 limits output to a manageable size, indicating the assistant expects multiple matches and wants to see the most relevant ones first.

The decision to target sm120*.hpp specifically (rather than all CUTLASS kernel files) shows architectural awareness: the SM120 backend is the one relevant to Blackwell GPUs, and the assistant knows that different SM architectures have different kernel implementations. Searching only SM120 files avoids noise from SM90 (Hopper) or SM100 (Blackwell sparse) variants.

The message also reveals the assistant's debugging philosophy: prefer static analysis over dynamic testing. The throwaway mention of "writing a small CUDA program" is immediately superseded by the grep approach, which is faster, requires no compilation, and leverages the existing codebase. This is a pragmatic choice in a remote debugging session where compilation can be slow and error-prone.

The Broader Context

This message sits at a transition point in the optimization effort. The earlier messages (segments 0-6) had focused on environment setup, driver installation, and basic deployment. Segment 7, where this message appears, is the deep performance analysis phase. The assistant has moved from "can we run the model?" to "why isn't it faster?" The investigation into CUTLASS tile failures is a microcosm of this larger shift: the assistant is no longer satisfied with a working system; it wants to understand the fundamental limits of the hardware and software stack.

The discovery that the GPUs are drawing only 235W during inference ([msg 903]) was the catalyst for this deep dive. When hardware is operating at 39% of its thermal design power, there is clearly untapped performance. The CUTLASS tile investigation is one of several parallel optimization paths the assistant is pursuing — others include adjusting max-running-requests, tuning num-continuous-decode-steps, and exploring expert parallelism. The tile investigation is particularly important because it targets the core compute kernel, and fixing it could unlock gains across all configurations.

Conclusion

Message 918 is a small but pivotal step in a complex optimization journey. It represents the transition from hypothesis to verification, from estimation to measurement, and from surface-level debugging to deep architectural investigation. The grep command and its output may seem mundane, but they carry the weight of a disproven hypothesis and the seed of a new direction. The assistant's careful construction of the search, its architectural awareness in targeting SM120 files, and its preference for static analysis over dynamic testing all reveal a methodical engineering mind at work.

The static_assert discovery does not solve the problem — the tiles still crash — but it eliminates the most obvious explanation and forces the investigation to consider more subtle possibilities. In the subsequent messages, the assistant will examine the SharedStorage structure in detail, compute its size for each tile configuration, and eventually trace the crash to a dynamic shared memory allocation issue. But message 918 is where the path forks, and that makes it worthy of close study.