The Autotune That Wouldn't Quit: When 160 "Errors" Mean Success
In the high-stakes world of large language model inference optimization, few moments are as disorienting as watching a server you know should be crashing... serve requests anyway. This is precisely the scenario that unfolds in message 1187 of an extensive optimization session targeting the GLM-5-NVFP4 model on eight NVIDIA RTX PRO 6000 Blackwell GPUs. The message captures a single, pivotal benchmark run that resolves a tense cliffhanger from the preceding messages and, in doing so, reveals something profound about how modern GPU kernel autotuners handle architectural constraints.
The Setup: A Long Campaign Against SM120's Limits
To understand message 1187, one must appreciate the months-long (or at least session-long) optimization campaign that preceded it. The assistant had been systematically working through a prioritized list of optimization techniques for the GLM-5-NVFP4 model, a Mixture-of-Experts (MoE) architecture with 256 experts running on the novel SM120 architecture of Blackwell GPUs. Earlier in the session, the assistant had benchmarked a baseline TP8 (Tensor Parallelism, 8 GPUs) configuration, achieving 10.36 tok/s single-stream and 19.29 tok/s dual-stream — excellent linear scaling that confirmed the model was compute-bound at low concurrency.
The next big idea was Expert Parallelism (EP8), which distributes experts across GPUs rather than sharding individual tensors. EP promised better scaling for MoE models by keeping each expert's weights intact and only routing tokens to the GPU hosting the needed expert. But earlier attempts at EP8 had ended in disaster: the server crashed under moderate load, running out of memory for all-reduce buffers. The assistant designed a "memory-safe" configuration using --mem-fraction-static 0.75 --max-running-requests 512, reserving more headroom. The launch script (run_tp8_ep8_memsafe.sh) was carefully crafted with NCCL tuning flags (NCCL_P2P_LEVEL=5, NCCL_MIN_NCHANNELS=8) and the correct CUDA 12.8 toolkit path.
The Cliffhanger: "CRASH DETECTED"
In message 1185, the assistant polled the EP8 server's log for readiness and was met with an alarming response:
CRASH DETECTED
1 0x789a90580ca6 tensorrt_llm::common::throwRuntimeError(...)
2 0x789a90cf6341 .../fused_moe_120.so(+0x90d341)
3 0x789a9076d948 void tensorrt_llm::kernels::cutlass_kernels_oss::dispatchMoeGemmSelectClusterShapeTmaWarpSpecialized<cutlass::arch::Sm120, __nv_fp4_e2m1, ...>The stack trace pointed directly at the CUTLASS MoE kernel's tile selection logic failing on SM120. The assistant's investigation in message 1186 identified the root cause: the autotuner was attempting a 128×256×128 tile configuration that exceeded SM120's 100KB shared memory limit. With EP8, the GEMM dimensions change (the N dimension grows from 256 to 2048), causing the autotuner to explore larger tiles that simply don't fit in Blackwell's shared memory budget.
But then came the twist. In message 1186, the assistant ran a more careful check:
grep -c "ready to roll" /root/sglang-server-ep8-memsafe.log → 1
grep -c "RuntimeError\|FAILED\|Killed" /root/sglang-server-ep8-memsafe.log → 160One "ready to roll" and 160 error lines. The server had logged 160 CUTLASS autotune failures... and then started serving requests anyway. The final lines showed Uvicorn running, a successful model info request, and a completed generate call. The server was alive.
The Subject Message: A Reversal and a Benchmark
Message 1187 opens with a triumphant declaration that carries the weight of the preceding uncertainty:
EP8 is running! The 160 "errors" were just CUTLASS autotune failures for tiles that don't fit — the autotune handles this gracefully by falling back to working tiles. The server is ready.
This single sentence reframes the entire preceding drama. What looked like a catastrophic failure was actually the autotuner doing exactly what it was designed to do: probing the space of possible tile configurations, failing on the ones that exceed hardware limits, and silently selecting a viable alternative. The 160 logged exceptions were not crashes — they were the sound of the autotuner learning the boundaries of SM120's shared memory capacity.
The assistant then executes the benchmark suite with a single bash command, running sglang.bench_serving at concurrency 1 and 2:
=== EP8, max-concurrency=1 ===
Output token throughput (tok/s): 9.79
Peak output token throughput (tok/s): 11.00
Total token throughput (tok/s): 20.96
Mean TTFT (ms): 263.91
Mean TPOT (ms): 99.02
Mean ITL (ms): 99.25
=== EP8, max-concurrency=2 ===
Output token throughput (tok/s): 18.28
Peak output token throughput (tok/s): 22.00
Total token throughput (tok/s): 39.08
Mean TTFT (ms): 326.70
Mean TPOT (ms): 103.34
Mean ITL (ms): 103.33 Interpreting the Numbers
The results tell a nuanced story. At single-stream (concurrency 1), EP8 achieves 9.79 tok/s output throughput with a TPOT (Time Per Output Token) of 99.02 ms. Compare this to the baseline TP8 single-stream benchmark from earlier in the session: 10.36 tok/s with 95.14 ms TPOT. EP8 is slightly slower — about 5.5% lower throughput and 4% higher latency.
At dual-stream (concurrency 2), EP8 delivers 18.28 tok/s output throughput with 103.34 ms TPOT. The baseline TP8 achieved 19.29 tok/s with 99.21 ms TPOT. Again, EP8 trails by about 5.2%.
The scaling efficiency is similar: EP8 achieves 1.87x throughput going from 1 to 2 streams (18.28 / 9.79), while TP8 achieved 1.86x (19.29 / 10.36). Both configurations scale nearly linearly, confirming that at low concurrency, the model is compute-bound rather than communication-bound.
The TTFT (Time To First Token) is notably higher for EP8: 263.91 ms vs the baseline's ~200 ms at concurrency 1, and 326.70 ms vs ~250 ms at concurrency 2. This makes sense — EP8 adds an all-to-all communication step before the first decode, routing tokens to the correct expert GPUs.
Why EP8 Underperforms TP8 at Low Concurrency
The fact that EP8 is slightly worse than TP8 at low concurrency is not surprising, and the assistant's analysis throughout the session explains why. The GLM-5-NVFP4 model has 256 experts, each relatively small. With EP8, each GPU hosts 32 experts (256/8). The per-expert GEMM operations are tiny, and the CUTLASS autotuner is forced to use suboptimal tile configurations because the larger tiles that would be efficient don't fit in SM120's 100KB shared memory. The assistant had previously identified this as the core bottleneck: "small per-expert GEMMs on SM120."
Moreover, EP8 introduces all-to-all communication overhead for every decode step. At low concurrency (1-2 streams), the batch size is tiny, so the computation is already fast and the communication overhead becomes a significant fraction of total latency. EP8 typically shines at high concurrency where large batches amortize the communication cost and expert load balancing matters more.
The Deeper Lesson: Graceful Degradation in GPU Autotuners
The most valuable insight from message 1187 is not the benchmark numbers themselves but the behavior of the CUTLASS autotuner. The FlashInfer TRT-LLM MoE backend, when faced with SM120's constrained shared memory (100KB vs 228KB on Hopper), does not abort. Instead, it iterates through possible tile configurations, catches the ones that fail with throwRuntimeError, and continues to the next candidate. This is a design choice worth examining.
The autotuner's search space for MoE GEMMs includes tiles like 128×256×128, 256×128×128, 128×128×256, and many others. On SM120, the 128×256×128 tile requires more than 100KB of shared memory per threadblock — it simply doesn't fit. But rather than failing with a hard error, the autotuner logs the failure and tries the next configuration. This is possible because CUTLASS's dispatchMoeGemmSelectClusterShapeTmaWarpSpecialized is designed as a selector: it probes tiles at kernel launch time and picks the first one that compiles and fits.
The 160 logged errors represent the autotuner's exploration of the tile space. Each "error" is a data point: "this tile doesn't work on this hardware." The autotuner accumulates these data points until it finds a working configuration, then proceeds silently. This is a beautiful example of what software engineers call "fail-fast with fallback" — the system tries aggressively, fails informatively, and recovers gracefully.
Assumptions and Knowledge Required
To fully understand message 1187, the reader needs several pieces of background knowledge. First, the concept of CUDA threadblock tiles: GPU kernels process matrices in small rectangular blocks (tiles), and the tile dimensions (M×N×K) determine register pressure, shared memory usage, and computational efficiency. Second, the SM120 architecture's 100KB shared memory limit — this is a known constraint of Blackwell GPUs that differs from Hopper's 228KB, making large tiles impossible. Third, the MoE GEMM dimension change under EP: when experts are distributed across GPUs, the N dimension (output features per expert) grows because each GPU handles a subset of experts with full weight matrices rather than sharded ones. Fourth, the FlashInfer MoE backend's architecture, which wraps TensorRT-LLM's CUTLASS kernels and performs autotuning at warmup time.
The assistant also assumes that the benchmark methodology is sound: using --max-concurrency to precisely control the number of concurrent requests, --request-rate 999 to ensure the server is fully loaded at the target concurrency, and --num-prompts 16 to get statistically stable measurements. The random input/output length of 128 tokens provides a consistent workload.
Output Knowledge Created
Message 1187 produces several concrete outputs. First, it establishes that EP8 on SM120 is viable but not beneficial at low concurrency for this model — a negative result that is valuable for guiding future optimization efforts. Second, it confirms that the CUTLASS autotuner gracefully handles SM120's shared memory constraints, which means future EP attempts can focus on performance rather than correctness. Third, the benchmark numbers provide a baseline for any future EP8 improvements (e.g., CUTLASS 4.4.0 with SM120 fixes, or the flashinfer_cutedsl backend that may have better tile selection for 256-expert models).
Most importantly, the message demonstrates a methodology: when an optimization technique fails, investigate how it fails. The 160 errors were not a wall — they were a map of the hardware's boundaries. By reading that map correctly, the assistant turned a "CRASH DETECTED" into a successful benchmark run.