When CUTLASS Tiles Don't Fit: Diagnosing SM120 Shared Memory Failures in Expert Parallelism

Introduction

In the high-stakes world of large language model inference optimization, few moments are as revealing as the crash of a freshly launched server. On February 19, 2026, during an intensive optimization campaign targeting the GLM-5-NVFP4 model on 8 NVIDIA RTX PRO 6000 Blackwell GPUs, an assistant encountered exactly such a moment. Message [msg 1186] captures a pivotal diagnostic episode: the Expert Parallelism (EP8) server, launched with a carefully tuned memory-safe configuration, crashes during warmup with a CUTLASS MoE kernel failure. What follows is a masterclass in GPU kernel debugging — a rapid, precise diagnosis that traces a stack trace through shared memory limits, GEMM dimension changes, and architecture-specific tile constraints.

This article examines that single message in depth, unpacking the reasoning, assumptions, and knowledge that make it a compelling case study in modern ML systems engineering.

The Message in Full

The assistant writes:

EP8 is crashing in the CUTLASS MoE kernel — specifically, it's trying to use tile 128x256x128 and 256x128x128 which we know fail on SM120 due to shared memory overflow. With EP8, the GEMM dimensions change (N=2048 instead of N=256), so the autotune is trying tiles that don't fit.

>

The error is during warmup in dispatchMoeGemmSelectClusterShapeTmaWarpSpecialized with cute::tuple<cute::C<128>, cute::C<256>, cute::C<128>> — the 128x256x128 tile that exceeds 100KB SMEM.

>

This is a known issue. The FlashInfer TRT-LLM MoE backend tries all tiles and some fail on SM120. Let me check if there's a way to filter or if the server continues past this: [bash command to check log]

The message then executes a bash command to probe the server log for signs of life — counting "ready to roll" lines, error lines, and tailing the final log entries.

Why This Message Was Written: Context and Motivation

To understand why this message exists, one must understand the broader optimization campaign consuming the assistant's attention. The GLM-5-NVFP4 model — a Mixture-of-Experts architecture with 256 experts — was being deployed on 8 RTX PRO 6000 Blackwell GPUs (SM120 architecture). The assistant had already achieved impressive throughput gains: an sglang update alone yielded a 2x improvement at 256 concurrency, and Opportunistic Expert Activation (OEA) had been implemented and benchmarked.

The natural next step was Expert Parallelism (EP8), where each GPU hosts a subset of experts, reducing the allreduce overhead that plagued the TP8 (Tensor Parallelism) configuration. Earlier EP attempts had crashed due to memory exhaustion — the A2A (All-to-All) communication buffers consumed too much GPU memory. The assistant had responded by crafting a "memory-safe" configuration: --mem-fraction-static 0.75 --max-running-requests 512, which left approximately 34 GB free per GPU.

The message at index 1186 is the moment that memory-safe EP8 server starts, and the assistant discovers it's crashing — not from OOM this time, but from a far more interesting architectural constraint.

The Reasoning and Diagnostic Process

The assistant's reasoning in this message is a textbook example of GPU kernel debugging. It proceeds through several logical steps:

Step 1: Identify the crash signature. The stack trace points to dispatchMoeGemmSelectClusterShapeTmaWarpSpecialized — a CUTLASS MoE kernel dispatcher that tries various tile configurations during warmup. The specific failing tile is 128x256x128.

Step 2: Connect to prior knowledge. The assistant immediately recognizes this tile as one that "we know fail on SM120 due to shared memory overflow." This is not a new discovery — earlier in the optimization campaign, the team had identified that SM120 (Blackwell architecture) has a 100 KB shared memory limit per block, and certain CUTLASS tiles exceed this budget. The 128×256×128 tile requires more than 100 KB of shared memory for its accumulator and operand buffers.

Step 3: Explain why EP8 triggers this. The key insight is that Expert Parallelism changes the GEMM dimensions. In TP8, each GPU handles all 256 experts but with a smaller hidden dimension (N=256). In EP8, each GPU handles only 32 experts but with the full hidden dimension (N=2048). The CUTLASS autotuner, seeing different matrix dimensions, tries different tile configurations — and some of those configurations, while valid for the new dimensions, exceed SM120's shared memory budget.

Step 4: Classify as a known issue. The assistant notes that "The FlashInfer TRT-LLM MoE backend tries all tiles and some fail on SM120." This frames the crash not as a bug in the assistant's configuration but as a known limitation of the FlashInfer MoE backend on Blackwell hardware.

Step 5: Formulate the next question. Rather than declaring failure, the assistant asks: "Let me check if there's a way to filter or if the server continues past this." This is the crucial pivot — is the crash fatal, or does the CUTLASS autotuner gracefully fall back to working tiles?

Assumptions and Their Validity

The message rests on several key assumptions, some explicit and some implicit:

Assumption 1: The crash is caused by shared memory overflow. This is well-supported. The SM120 architecture has a documented 100 KB shared memory limit per block, and the 128×256×128 tile is known to exceed this. The stack trace confirms the tile dimensions. This assumption is correct.

Assumption 2: EP8 changes N from 256 to 2048. This is correct for the GLM-5-NVFP4 model's architecture. In TP8, the hidden dimension is sharded across 8 GPUs (2048/8 = 256). In EP8, each GPU has the full hidden dimension (2048) but fewer experts (256/8 = 32).

Assumption 3: The autotune tries all tiles regardless of SMEM budget. This appears to be correct based on the observed behavior — the FlashInfer TRT-LLM backend enumerates tile configurations and attempts each one, rather than pre-filtering based on shared memory requirements.

Assumption 4: The server might continue despite the error. This is the key uncertainty the assistant tests with the bash command. And indeed, the subsequent message ([msg 1187]) reveals that the server did recover — the 160 "errors" were non-fatal autotune failures, and the server is running. This assumption was validated.

Input Knowledge Required

To fully understand this message, a reader needs:

  1. CUTLASS tile geometry: Understanding that GEMM operations are decomposed into tiles (M×N×K dimensions) that must fit in GPU shared memory. The 128×256×128 tile means 128 rows of output, 256 columns of output, and a K-dimension of 128 for the reduction.
  2. SM120/Blackwell architecture: Knowledge that NVIDIA's Blackwell architecture has a 100 KB shared memory limit per block, which constrains which CUTLASS tiles can execute.
  3. Expert Parallelism vs Tensor Parallelism: Understanding that TP shards the hidden dimension across GPUs (reducing N), while EP shards the expert count across GPUs (reducing the number of experts per GPU but keeping the full hidden dimension).
  4. FlashInfer TRT-LLM MoE backend: Knowledge that this backend uses CUTLASS for FP4 GEMM operations and performs autotuning during warmup, trying multiple tile configurations.
  5. The GLM-5-NVFP4 model architecture: Specifically that it has 256 experts and a hidden dimension of 2048, which determines the GEMM dimensions in different parallelism configurations.
  6. The prior optimization history: Understanding that earlier EP8 attempts failed due to memory exhaustion, leading to the memory-safe configuration being tested here.

Output Knowledge Created

This message produces several valuable pieces of knowledge:

  1. EP8 on SM120 triggers different CUTLASS tiles than TP8: The change in GEMM dimensions (N=2048 vs N=256) causes the autotuner to explore a different tile space, including tiles that exceed SM120's shared memory.
  2. The FlashInfer TRT-LLM backend's autotune is resilient: Even though tiles fail during warmup, the backend can fall back to working tiles and continue serving. This is a critical reliability property.
  3. The specific failing tile signature: 128×256×128 in dispatchMoeGemmSelectClusterShapeTmaWarpSpecialized is a known failure mode on SM120.
  4. A diagnostic methodology: The assistant demonstrates how to check whether autotune failures are fatal by counting "ready to roll" vs error lines in the log, and by tailing the final log output for server status.

The Thinking Process

What makes this message particularly interesting is the compressed reasoning chain visible in just a few sentences. The assistant moves from observation ("EP8 is crashing") to root cause ("trying to use tile 128x256x128 which we know fail on SM120") to explanation ("With EP8, the GEMM dimensions change") to classification ("This is a known issue") to next-step formulation ("Let me check if the server continues past this").

This is not a novice debugging session where each hypothesis is tested independently. The assistant draws on accumulated knowledge from earlier in the optimization campaign — the SM120 shared memory limit had been discovered and documented previously, and the specific failing tiles had been encountered. The message demonstrates how institutional knowledge (even in an AI assistant's context window) accelerates debugging: the assistant doesn't need to Google the error or experiment with tile sizes; it recognizes the pattern immediately.

The bash command at the end is also revealing. Rather than assuming the crash is fatal, the assistant checks empirically. It counts both "ready to roll" (success indicator) and "RuntimeError|FAILED|Killed" (failure indicators), then tails the log for the most recent entries. This dual-signal approach — looking for both positive and negative indicators — is a robust diagnostic technique.

The Broader Significance

This message sits at an inflection point in the optimization campaign. The EP8 approach — which promised to eliminate allreduce overhead by partitioning experts across GPUs — is encountering fundamental hardware limitations. The CUTLASS tile failures on SM120 are not configuration bugs; they are architectural constraints that no amount of software tuning can fully circumvent. The assistant's recognition of this ("This is a known issue") signals a mature understanding of where the optimization frontier lies.

Yet the story doesn't end with this diagnosis. The server does recover and run EP8, as the next message reveals. And the benchmarks that follow show EP8 achieving 9.79 tok/s single-stream (vs 10.36 tok/s baseline) — close but not superior. The real bottleneck remains the small per-expert GEMMs on SM120, a theme that will dominate the remainder of the optimization campaign.

Conclusion

Message [msg 1186] is a compact but rich example of GPU kernel debugging in the context of large-scale ML inference optimization. It demonstrates how architectural knowledge (SM120 shared memory limits), model architecture understanding (GEMM dimension changes under different parallelism strategies), and empirical validation (checking if the server recovered) combine to diagnose a complex failure mode. The assistant's reasoning is precise, contextual, and forward-looking — it doesn't just identify the crash but explains why it happens and formulates the next diagnostic step. For anyone working on LLM inference optimization, particularly on NVIDIA's Blackwell architecture, this message offers a valuable case study in how CUTLASS tile selection interacts with parallelism strategies, and how to distinguish fatal errors from recoverable autotune failures.