Peering into the CUTLASS Kernel: A Single Bash Command That Uncovered Why Blackwell GPUs Stall

The Message

In the middle of an intensive performance debugging session for the GLM-5-NVFP4 model running on eight NVIDIA RTX PRO 6000 Blackwell GPUs, the assistant issued the following command:

ssh root@10.1.230.174 'sed -n "420,470p" /root/ml-env/lib/python3.12/site-packages/flashinfer/data/csrc/nv_internal/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/launchers/moe_gemm_tma_ws_launcher.inl'

The output was a single line from a C++ template instantiation:

      using CollectiveEpilogueFinalize = typename cutlass::epilogue::collective::                                                                                                                                                                                                                                       \
          CollectiveBuilder<                                        /**/                                                                                                                  ...

At first glance, this appears to be nothing more than a routine file read — a developer peeking at source code. But in the context of the session, this single command represents a critical turning point in a deep investigation into why Blackwell GPUs were dramatically underperforming during inference. It is the moment the assistant stopped treating the GPU as a black box and started reverse-engineering the very kernels that drive it.

Context: The Performance Mystery

To understand why this message matters, one must appreciate the problem that preceded it. The assistant had been deploying the GLM-5-NVFP4 model, a massive Mixture-of-Experts (MoE) transformer with 256 experts, quantized to NVFP4 (NVIDIA's 4-bit floating point format). The hardware was formidable: eight RTX PRO 6000 Blackwell GPUs, each rated at 3,700 TFLOPS for FP4 sparse compute (1,850 TFLOPS dense), with 1,597 GB/s of memory bandwidth. The theoretical system peak was nearly 30 PFLOPS.

Yet during inference, the GPUs were delivering a paltry 0.8 TFLOPS per expert — just 0.04% of their rated peak. Even in the best-case micro-benchmark, the CUTLASS FP4 GEMM kernels plateaued at roughly 1,300 TFLOPS (70% of dense peak), but only for impractically large matrices. During actual decode, where each expert saw only 16–64 tokens per batch, the arithmetic intensity was 72× below what was needed to be compute-bound.

The assistant had already established that the model was compute-bound rather than communication-bound (TP4+PP2 was 2× slower than TP8, ruling out allreduce latency). The bottleneck was squarely in the FP4 GEMM kernels themselves. But why were the kernels so inefficient on Blackwell?

Why This Message Was Written

Message [msg 896] was written as a direct follow-up to a discovery made in the preceding minutes. In [msg 891], the assistant had examined the autotuned kernel files generated by FlashInfer and found that only a subset of tile configurations were working on SM120 (Blackwell's compute architecture). Specifically:

| Tile (M×N×K) | Status | |---|---| | 128×128×128 | Working | | 128×128×256 | Working | | 128×256×128 | FAILS | | 256×128×128 | FAILS |

The two larger tiles — M128×N256 and M256×N128 — were failing to initialize. These were precisely the configurations that would improve performance by processing larger blocks of data, increasing data reuse and compute density. The assistant hypothesized that shared memory limitations on SM120 (99 KB per block) were preventing these larger tiles from being compiled.

In [msg 895], the assistant had searched the launcher file for shared memory related terms (smem, shared, SharedMemory, etc.) but found no direct references — only a line at position 437 referencing sizeof(typename CollectiveEpilogue::SharedStorage). This was a dead end with grep. The assistant needed to see the actual template code to understand how shared memory was being calculated and why the larger tiles failed.

Thus, message [msg 896] was born: a targeted sed command to read lines 420–470 of the 734-line launcher template file, specifically the region around the CollectiveEpilogue::SharedStorage reference. This was not a random read — it was a surgical probe into the exact section of code that governed the shared memory budget for the epilogue phase of the GEMM kernel.## The Reasoning Process Visible in This Message

What makes this message remarkable is the chain of reasoning that led to it. The assistant had been systematically working through a diagnostic tree:

  1. Macro level: Benchmark TP4+PP2 vs TP8 → confirmed compute-bound, not communication-bound.
  2. Micro level: Micro-benchmark FP4 GEMM at various batch sizes → found catastrophic underutilization at small batch sizes.
  3. Kernel level: Examine autotuned kernel files → discovered only 128×128 tiles work, larger tiles fail.
  4. Source level: Search for shared memory references → found a single sizeof(SharedStorage) at line 437.
  5. Template level: Read the launcher template around that line → this message. Each level peeled back another layer of abstraction. The assistant was not content to stop at "the kernels are slow" — it wanted to know why the kernels were slow at the CUTLASS template metaprogramming level. This is the difference between a superficial investigation and a deep one. The sed command itself reveals the assistant's mental model. It didn't read the entire file (734 lines) — it read exactly 50 lines centered on line 437, where the shared storage calculation lived. The assistant already knew from [msg 892] that the difference between working and failing tiles was the FINALIZE epilogue mode (vs NONE). The question was: what about the FINALIZE epilogue required more shared memory than the SM120's 99 KB budget allowed?

Assumptions and Their Consequences

The assistant made several assumptions in crafting this message:

Assumption 1: The shared memory calculation for the epilogue was the root cause of the tile initialization failure. This was a reasonable hypothesis — larger tiles require more shared memory for accumulator storage, and SM120 has a fixed 99 KB per block. However, as the subsequent message ([msg 897]) would reveal, the actual issue was more nuanced: SM120 uses KernelScheduleAuto (CUTLASS's automatic kernel scheduler) rather than the explicit warp-specialized schedules available for SM100. The auto scheduler likely picks suboptimal configurations, and the StageCountAutoCarveout mechanism auto-computes how many pipeline stages fit in the available SMEM. The assistant's assumption was partially correct — shared memory was involved — but the deeper issue was the absence of SM120-optimized kernel schedules in the CUTLASS codebase.

Assumption 2: The answer would be found in the launcher template file itself. The assistant assumed that the shared memory constraint was encoded in the template metaprogramming, visible by reading the source. In reality, the constraint emerges from the interaction between the template parameters (tile sizes, epilogue mode, pipeline stages) and the hardware limits (SMEM per block), which are resolved at compile time by the CUDA compiler. Reading the source could reveal the mechanism but not the specific values that cause the failure.

Assumption 3: The sed output would be complete and readable. In practice, the output was truncated — the long C++ template instantiation line was cut off with .... The assistant got only a partial view of the code. This is a limitation of running commands over SSH with line truncation. The assistant would need to follow up with more targeted reads in subsequent messages.

Input Knowledge Required

To understand this message, the reader needs:

  1. CUTLASS architecture: Knowledge that CUTLASS uses template metaprogramming to generate GPU kernels, with CollectiveBuilder composing epilogue and mainloop strategies. The CollectiveEpilogue handles the output computation (e.g., bias addition, activation, quantization).
  2. FlashInfer's integration: Understanding that FlashInfer bundles CUTLASS kernels from TensorRT-LLM, and the autotuner generates concrete instantiations from these templates. The file path (flashinfer/data/csrc/nv_internal/tensorrt_llm/...) reveals the provenance.
  3. SM120 shared memory limits: Blackwell GPUs (compute capability 12.0) have 99 KB of shared memory per thread block when using TMA (Tensor Memory Accelerator), down from 228 KB on Hopper (SM90). This is a critical constraint for tile-based GEMM kernels.
  4. MoE grouped GEMM semantics: The INSTANTIATE_TMA_WARP_SPECIALIZED_MOE_GEMM macro generates kernels that process multiple expert GEMMs in a single kernel launch, using warp-specialized programming models with TMA for data movement.
  5. The FINALIZE vs NONE epilogue distinction: FINALIZE epilogue performs the final output write (potentially with quantization or activation), while NONE skips it. FINALIZE requires additional shared memory for the epilogue's tile accumulator.

Output Knowledge Created

This message produced:

  1. A confirmed code location: The CollectiveEpilogueFinalize type alias at the top of the template instantiation block, confirming that the FINALIZE epilogue uses cutlass::epilogue::collective::CollectiveBuilder — the standard CUTLASS epilogue composition mechanism.
  2. A negative result: The single line of output was insufficient to determine the exact shared memory calculation. The assistant would need to read more of the file (which it did in [msg 897], finding the KernelScheduleAuto and StageCountAutoCarveout references).
  3. A direction for further investigation: The truncated output told the assistant that the template was deeply nested and would require more careful reading. This led to the subsequent discovery about SM120 using KernelScheduleAuto instead of SM100's explicit warp-specialized schedules.

Mistakes and Incorrect Assumptions

The most significant mistake was the assumption that the shared memory calculation was the sole reason for the tile failures. While shared memory is a constraint, the subsequent investigation revealed that the real issue was more architectural: CUTLASS simply lacks SM120-optimized kernel schedules. The KernelScheduleAuto fallback for SM120 means the compiler must auto-detect the best schedule, which may not handle larger tiles correctly. On SM100 (Hopper), explicit schedules like KernelPtrArrayTmaWarpSpecialized2SmBlockScaledMxNvf4UltraVs16Sm103 are available; on SM120, the auto scheduler is used instead.

Additionally, the assistant assumed that reading lines 420–470 would show the complete shared memory logic. In reality, the template metaprogramming spans multiple files and involves compile-time constant expressions that are not visible in a simple sed of one file. The sizeof(typename CollectiveEpilogue::SharedStorage) at line 437 is just the usage of the shared storage size — the actual calculation is in the CUTLASS CollectiveBuilder and the individual epilogue component templates.

The Broader Significance

This message exemplifies a pattern that recurs throughout the session: the assistant repeatedly drills down from high-level benchmarks to low-level kernel source code, treating no layer as opaque. When the GPU doesn't perform as expected, the assistant doesn't just try different flags or configurations — it reads the actual kernel templates to understand why.

The message also illustrates the challenges of debugging on cutting-edge hardware. Blackwell (SM120) is so new that CUTLASS doesn't have optimized schedules for it. The assistant is effectively doing architecture bring-up work, discovering SM120-specific limitations that the CUTLASS developers haven't yet addressed. The 99 KB shared memory limit, the auto-scheduler fallback, and the failing tile configurations are all symptoms of a GPU architecture that is ahead of its software ecosystem.

In the end, this single sed command — reading 50 lines of a C++ template file over SSH — represents the moment when the investigation shifted from "how do we work around the slow kernels" to "why are the kernels slow at the hardware level." It is the pivot point in the segment, after which the assistant begins documenting concrete optimization approaches (the glb5improvement-xx.md files) informed by a deep understanding of the SM120 architecture's constraints.