The CUTLASS Tile That Wouldn't Fit: Debugging EP8 on Blackwell SM120

In the high-stakes world of large language model inference optimization, few moments are as revealing as the precise instant a carefully planned experiment fails. Message 1185 of this opencode session captures exactly such a moment: the assistant's attempt to launch an Expert Parallelism (EP8) configuration for the GLM-5-NVFP4 model on 8 NVIDIA RTX PRO 6000 Blackwell GPUs, only to be met with a CUTLASS tile failure that exposes a fundamental architectural constraint of the SM120 compute architecture.

The Context: A Methodical Optimization Campaign

This message sits within a larger narrative of systematic performance optimization. The assistant had been engaged in an intensive campaign to maximize inference throughput for the GLM-5-NVFP4 model — a Mixture-of-Experts (MoE) architecture with 256 experts and FP4 quantization. Earlier in the session, the assistant had already achieved impressive results: updating sglang to the latest commit yielded a 2x throughput improvement at 256 concurrency, and implementing Opportunistic Expert Activation (OEA) showed marginal peak-throughput gains. The assistant had also benchmarked single-stream and dual-stream performance, finding excellent linear scaling (10.36 tok/s single-stream, 19.29 tok/s dual-stream).

The natural next step was to test Expert Parallelism (EP), a technique that distributes MoE experts across GPUs to reduce the per-GPU computation burden. Previous EP8 attempts had crashed under load, likely due to memory exhaustion from all-to-all (A2A) communication buffers. The assistant's hypothesis was that reducing the memory fraction — from the default to --mem-fraction-static 0.75 — would leave sufficient headroom for these buffers to operate without OOM errors.

The Message: Optimism Meets Reality

The message begins with a cautiously optimistic status report:

EP8 loaded with 34.15 GB available and 59.52 GB used per GPU. With 0.75 mem fraction, that leaves ~34 GB free — much more room for A2A buffers now. Still loading. Let me wait:

The assistant had successfully loaded the model weights across 8 GPUs with expert parallelism. The memory numbers looked healthy: 59.52 GB used out of ~94 GB total (the RTX PRO 6000 has 96 GB VRAM), leaving 34 GB free per GPU. This seemed to confirm that the memory-safe configuration was working — the reduced memory fraction had prevented the OOM that plagued earlier attempts.

But the real test was whether the server would complete its initialization and reach the "ready to roll" state. The assistant set up a polling loop with a 90-iteration timeout, checking for either success ("ready to roll") or failure signals ("Traceback", "RuntimeError", "CUDA error"). The loop ran with 10-second intervals, meaning the assistant was prepared to wait up to 15 minutes for the server to initialize.

The result came back swiftly: "CRASH DETECTED".

The Crash: A Deep-Dive into CUTLASS Tile Geometry

The crash trace reveals the precise nature of the failure:

1       0x789a90580ca6 tensorrt_llm::common::throwRuntimeError(...)
2       0x789a90cf6341 /root/.cache/flashinfer/0.6.3/120a/cached_ops/fused_moe_120/fused_moe_120.so(+0x90d341)
3       0x789a9076d948 void tensorrt_llm::kernels::cutlass_kernels_oss::dispatchMoeGemmSelectClusterShapeTmaWarpSpecialized<cutlass::arch::Sm120, __nv_fp4_e2m1, __nv_fp4_e2m1, __nv_bf...

This is a CUTLASS MoE GEMM (General Matrix Multiply) kernel failure. The function name dispatchMoeGemmSelectClusterShapeTmaWarpSpecialized tells us exactly what's happening: the FlashInfer MoE backend (which uses TensorRT-LLM's CUTLASS kernels) is attempting to select a tile configuration for the SM120 architecture, and it's failing.

The root cause, which the assistant would analyze in the subsequent message ([msg 1186]), is that the autotuner is attempting to use tile dimensions of 128×256×128 and 256×128×128. These tiles require more than 100KB of shared memory — the fixed budget available on SM120 (Blackwell architecture). The 128×256×128 tile alone would need approximately 128 × 256 × 2 bytes (for FP4) × some overhead for accumulators and staging, easily exceeding the 100KB limit.

Why EP8 Changes the Tile Geometry

A crucial insight here is why EP8 triggers this failure when TP8 (tensor parallelism) did not. With expert parallelism, each GPU handles a subset of experts, but the hidden dimension (N) per expert changes. In the GLM-5-NVFP4 architecture, each expert processes an intermediate dimension. When experts are distributed across 8 GPUs via EP, the per-GPU GEMM dimensions shift — specifically, the N dimension (the output feature dimension) becomes 2048 instead of 256. This different geometry causes the CUTLASS autotuner to explore different tile configurations, including the problematic 128×256×128 tile that exceeds SM120's shared memory budget.

The TP8 configuration that had been working previously used different tile geometries that fit within the 100KB budget. The EP8 configuration, by changing the matrix dimensions, pushed the autotuner into a region of the tile configuration space that is invalid for SM120.

Assumptions and Their Consequences

The assistant made several assumptions in this message, most of which were reasonable but ultimately incorrect:

  1. Memory was the bottleneck: The assistant assumed that the previous EP8 crashes were due to OOM from A2A buffers. Reducing --mem-fraction-static from its default (typically 0.9 or higher) to 0.75 was intended to leave more room. While this did prevent OOM (the model loaded successfully), the actual crash was in the kernel compilation/warmup phase, not in memory allocation.
  2. The FlashInfer MoE backend would handle SM120 gracefully: The assistant knew that FlashInfer's TRT-LLM MoE backend had SM120 support (it was explicitly compiled for SM120, as shown by the fused_moe_120.so path). However, the assumption that the autotuner would only select valid tile configurations proved wrong — the autotuner attempted tiles that exceeded the shared memory budget.
  3. The server would either fully succeed or cleanly fail: The polling loop checked for "ready to roll" or crash signals. In reality, the server did eventually recover and start (as [msg 1186] shows: "The server is fi..." and a subsequent check found 1 match for "ready to roll" and 160 matches for errors). This mixed state — a server that starts despite kernel compilation failures — is a subtle failure mode that the assistant's monitoring didn't fully anticipate.

Input Knowledge Required

To fully understand this message, one needs:

Output Knowledge Created

This message produces several valuable pieces of knowledge:

  1. EP8 with --mem-fraction-static 0.75 loads successfully: The model weights fit comfortably with 34 GB free per GPU. This is a useful datapoint for future EP experiments.
  2. CUTLASS tile failure is the EP8 blocker, not OOM: The root cause is architectural (SM120 shared memory limit), not a configuration issue. This redirects future optimization efforts away from memory tuning and toward kernel-level solutions.
  3. The specific failing tile is 128×256×128: This precise geometry information is critical for anyone working on CUTLASS kernel configuration for Blackwell GPUs.
  4. The server can recover from warmup failures: Despite the crash during warmup, the server eventually started and served requests (as confirmed in [msg 1186]). This resilience is useful to know — kernel compilation failures during warmup don't necessarily kill the server.

The Thinking Process

The assistant's reasoning in this message is methodical and evidence-driven. The sequence of actions reveals a clear mental model:

  1. Verify the prerequisite: Check memory usage to confirm the model loaded with sufficient headroom. The assistant reports 34 GB free, which validates the memory-safe config.
  2. Set up monitoring: The polling loop is carefully designed with multiple failure detection patterns — checking for both explicit crash signals and the success condition. The 90-iteration timeout shows the assistant expects initialization to take up to 15 minutes.
  3. Interpret the crash: The assistant doesn't just report "CRASH DETECTED" — it captures the full stack trace, which provides the critical information about the CUTLASS tile failure.
  4. Connect to prior knowledge: In the following message ([msg 1186]), the assistant immediately recognizes the 128×256×128 tile as the known SM120 shared memory overflow issue, connecting this crash to earlier discoveries about Blackwell's 100KB limit.

Broader Implications

This message illustrates a fundamental challenge in ML inference optimization: the interaction between parallelism strategies and hardware constraints. Expert Parallelism seemed like a natural optimization path — distribute the 256 experts across 8 GPUs to reduce per-GPU computation. But the change in GEMM dimensions pushed the CUTLASS autotuner into invalid tile configurations, revealing a constraint that wasn't visible in the TP8 configuration.

The lesson is that parallelism strategies don't just change computation distribution — they change the fundamental shapes of the operations being performed, which can interact unpredictably with kernel-level optimizations. A configuration that works perfectly with TP may fail with EP, not because of memory or communication issues, but because the underlying GEMM operations have different dimensionalities that trigger different code paths in the kernel autotuner.

This is the kind of deep, architectural insight that only emerges from systematic experimentation — trying the obvious optimization, watching it fail, and understanding why it failed at the level of individual tile dimensions and shared memory budgets. The assistant's methodical approach — benchmarking, documenting, and analyzing each failure — transforms a seemingly simple crash into a rich understanding of the Blackwell architecture's constraints and capabilities.