Cracking the SM120 Shared Memory Ceiling: A Deep Dive into CUTLASS Tile Configuration Debugging

In the high-stakes world of large language model inference optimization, few things are as frustrating as a kernel that fails to initialize with a cryptic "Error Internal" message. This is exactly the situation the assistant faced in message 914 of this opencode session, where the goal was to maximize throughput of the GLM-5-NVFP4 model on eight NVIDIA RTX PRO 6000 Blackwell GPUs. The assistant had already established that the model was compute-bound rather than communication-bound, and that the FP4 GEMM kernels were the critical performance bottleneck. Now it needed to understand why certain CUTLASS tile configurations—specifically M128×N256 and M256×N128—were crashing during initialization, while smaller tiles like 128×128×128 and 128×128×256 worked fine.

The Context: A Compute-Bound Model on Blackwell GPUs

The broader session (Segment 7) had already produced several key findings. The assistant had benchmarked TP4+PP2 versus TP8 configurations, confirming the model was compute-bound. GPU power draw during inference was only ~235W out of 600W TDP—just 39%. The CUTLASS FP4 kernels achieved at most ~1,300 TFLOPS (70% of dense peak) for large matrices, but during actual decode, per-expert batch sizes of ~16–64 tokens yielded a paltry 0.8–55 TFLOPS (0.02–3% of peak). The 99KB shared memory limit on SM120 was preventing larger CUTLASS tile configurations from initializing. The user had given a clear directive in message 905: "Fix tiles, then try 2/3"—meaning first fix the failing tile configurations, then try increasing effective batch size and exploring the cuBLASLt FP4 path.

Tracing the Failure

The assistant's investigation in messages 906–913 had traced the failure to the initializeMoeGroupedGemm function in the CUTLASS MoE GEMM launcher. The can_implement check passed, but initialize returned kErrorInternal. This pointed to a runtime shared memory allocation failure—the kernel was requesting more dynamic shared memory than the device could provide, even after cudaFuncSetAttribute had been called to opt into the maximum.

The critical numbers, extracted in message 912, were:

• Default shared memory per block: 48 KB
• Opt-in max per block: 99 KB (101,376 bytes)
• Total per SM: 100 KB (102,400 bytes) So a CUTLASS kernel could use at most 99 KB of shared memory if it properly requested the opt-in maximum. The question was: how much did each tile configuration actually need?

The Shared Memory Calculation

Message 914 is where the assistant synthesizes all this information into a crystal-clear diagnosis. The message begins with "Now it's crystal clear:" and presents a detailed analysis of shared memory requirements for each tile configuration.

The calculation considers:

• A tile data: M × K × 0.5 bytes (FP4 packed, 4 bits per element)
• B tile data: K × N × 0.5 bytes (FP4 packed)
• A scales: M × (K/16) × 1 byte (FP8 scales, one per 16 elements)
• B scales: (K/16) × N × 1 byte
• Epilogue default: M × N × 2 bytes (BF16 output staging)
• Epilogue finalize: M × N × 4 bytes (FP32 accumulator + BF16 output) The results are illuminating:
• 128×128×128 stages=1: 50KB (default epilogue) or 82KB (finalize) — comfortably fits in 99KB
• 128×128×256 stages=1: 68KB (default) or 100KB (finalize) — the finalize variant is right at the boundary, explaining why it might be in the "failing group2"
• 128×256×128 stages=1: 91KB (default) or 155KB (finalize) — default barely fits, finalize is far too large
• 256×128×128 stages=1: same as above The key insight is that the epilogue output tile (M×N×sizeof) is the real bottleneck, not the mainloop pipeline stages. For 128×256 or 256×128 output tiles, even the default epilogue needs 65,536 bytes (64 KB) just for output staging, leaving only about 34 KB for the mainloop—not enough for even a single stage of the input matrices (which need about 27.6 KB).## The "Crystal Clear" Moment What makes message 914 particularly interesting is the reasoning process it reveals. The assistant had spent several prior messages gathering data—checking the CUTLASS launcher source code, running Python scripts to query GPU properties via PyTorch and the CUDA driver API, and writing a shared memory estimation script. But message 914 is where all that data coalesces into a coherent explanation. The assistant notes a crucial subtlety: "The StageCountAutoCarveout should auto-reduce to 1 stage, but the epilogue output tile (M×N×sizeof) is the real problem." This is a sophisticated observation. StageCountAutoCarveout is a CUTLASS mechanism that automatically determines how many pipeline stages fit within the available shared memory budget, subtracting the epilogue storage requirement. The assistant correctly identifies that even with auto-carveout reducing to a single stage, the epilogue itself is too large for the larger tiles. There's also an important moment of intellectual honesty: "Actually wait — my epilogue estimate might be wrong." The assistant acknowledges that its initial calculation might be using the wrong epilogue strategy and immediately pivots to check the actual CUTLASS source code. This self-correction is a hallmark of effective debugging—recognizing when your model of the system might be incomplete.

Assumptions and Their Validity

The assistant makes several assumptions in this analysis:

1. The epilogue storage estimate: The assistant assumes BF16 output staging (2 bytes per element) for the default epilogue and FP32 accumulator + BF16 output (4 bytes per element) for the finalize epilogue. These are reasonable estimates based on typical CUTLASS epilogue patterns, but the assistant wisely hedges by noting they might be wrong and proceeding to check the actual source.
2. That StageCountAutoCarveout is working correctly: The assistant assumes the auto-carveout mechanism is functioning as designed and reducing to the minimum number of stages. If there were a bug in the carveout calculation itself, the diagnosis might be different.
3. That the failure is purely SMEM-related: Given that can_implement passes but initialize fails, the SMEM hypothesis is the most likely explanation. However, there could be other reasons for kErrorInternal—for instance, register pressure exceeding the architectural limit, or unsupported tensor memory accelerator (TMA) descriptors for certain tile dimensions.

Input Knowledge Required

To fully understand this message, one needs:

• Understanding of CUTLASS tile configurations: The M×N×K notation for GEMM tiles, and how tile dimensions affect shared memory usage and computational efficiency.
• Knowledge of FP4 block-scaled quantization: The model uses FP4 (4-bit floating point) with block scaling, meaning every 16 elements share an FP8 scale factor. This affects how data is packed and how much shared memory each tile consumes.
• Familiarity with SM120 architecture: The Blackwell GPU's shared memory hierarchy—48 KB default, 99 KB opt-in, 100 KB total per SM—and how cudaFuncSetAttribute is used to request additional shared memory.
• Understanding of CUTLASS pipeline stages: The StageCountAutoCarveout mechanism and how it interacts with epilogue shared storage to determine the final kernel configuration.
• Context from prior messages: The assistant had already established that the model was compute-bound, that GPU power was only 235W, and that the failing tiles were 128×256 and 256×128.

Output Knowledge Created

This message produces several valuable outputs:

1. A precise diagnosis of the tile failure: The larger tiles fail because the epilogue output tile (M×N) requires more shared memory than is available after accounting for even a single pipeline stage of the mainloop.
2. A quantitative model of shared memory usage: The assistant provides a reusable formula for estimating SMEM requirements for any tile configuration, which can be applied to other tile sizes or architectures.
3. A clear direction for the fix: The problem is not with the mainloop pipeline stages but with the epilogue. This suggests that solutions might involve using a different epilogue strategy (e.g., epilogue in registers or via TMA store), reducing the output tile size, or finding a way to partition the output across multiple kernel invocations.
4. An actionable next step: The assistant immediately follows the analysis by looking at the actual CUTLASS source code to verify the epilogue estimates and understand what StageCountAutoCarveout actually does, demonstrating a systematic approach to verification.