The Investigation That Flipped: How Empirical Profiling Overturned the MoE-Centric Performance Model
Introduction
In the high-stakes world of large language model inference optimization, few experiences are as valuable—or as humbling—as the moment when empirical data decisively contradicts a well-reasoned hypothesis. This article chronicles one such moment: a deep-dive investigation into the performance characteristics of the DeepSeek-V4-Flash-NVFP4 model running on NVIDIA RTX PRO 6000 Blackwell GPUs, where a methodical, multi-phase analysis ultimately revealed that the investigation's core premise was inverted.
The session began with a clear mandate: investigate the MoE (Mixture of Experts) grouped-GEMM kernel to find tractable occupancy and throughput levers for high-concurrency decode (batch sizes 60–96). The model in question was a 43-layer transformer with 256 routed experts (top-6 active per token), using NVFP4 quantization for expert weights, deployed on SGLang with tensor parallelism across 8 Blackwell GPUs. The performance signature was puzzling: 97% SM utilization but only 57% power draw and 27% memory-controller utilization—a classic symptom of a latency-bound kernel that is neither compute-saturated nor memory-bandwidth-saturated.
What followed was a masterclass in systematic performance investigation: code tracing from configuration flags to actual GPU kernels, theoretical modeling of kernel behavior, careful experimental design for high-batch profiling, and—most crucially—a willingness to let the data overturn the operating hypothesis. The journey from "MoE is the bottleneck" to "attention is the bottleneck" is a case study in why empirical measurement must always trump assumption.
Phase One: Tracing the Code Path from Configuration to Kernel
The investigation began with a puzzle. The system was configured with --moe-runner-backend triton, suggesting that Triton-based kernels handled the MoE computation. Yet profile traces showed kernel names like cutlass_fp4_group_mm—a CUTLASS kernel, not a Triton one. This discrepancy demanded resolution before any meaningful optimization could proceed.
The assistant systematically traced the code path through the SGLang codebase, starting with the server argument parsing in server_args.py ([msg 5]), then following the dispatch logic through modelopt_quant.py where the ModelOptNvFp4FusedMoEMethod class was identified ([msg 8]). The critical discovery came when the assistant read the actual implementation: the NVFP4 quantization path unconditionally called cutlass_moe_fp4 directly, bypassing the MoeRunner object entirely ([msg 14]). The "triton" backend was a misnomer—for NVFP4 on Blackwell, the computation always routed through CUTLASS.
This finding had profound implications. Any performance optimization strategy based on Triton kernel characteristics—adjusting num_warps, num_stages, or Triton-specific autotuning parameters—would have zero effect. The actual kernels were CUTLASS templates, compiled with specific tile shapes determined at JIT time, governed by CUTLASS's resource allocation logic rather than Triton's parameters.
The assistant then drilled deeper, reading the actual CUTLASS kernel configuration for SM120 in nvfp4_blockwise_moe.cuh ([msg 17]). The kernel parameters were revealed: a thread block shape of 128×128×128 using TMA warp-specialized ping-pong persistent scheduling with grouped pointer arrays, no multi-CTA clustering, and automatic stage counting. This configuration immediately raised red flags for the decode use case.
Phase Two: Building the Theoretical Performance Model
With the actual kernel identified, the assistant constructed a detailed theoretical model of its behavior under decode conditions ([msg 18]). The core insight was devastatingly simple: at decode time, each expert receives only 0.75–2.25 tokens (calculated as batch_size × top_k / num_experts). The CUTLASS kernel's tile has an M-dimension of 128, meaning each tile iteration utilizes less than 2% of its rows. Yet the kernel still loads the full 128×4096 weight slice (256KB of FP4 data) for each output tile.
This is the classic signature of a weight-load-bound kernel: the arithmetic intensity is catastrophically low because the weight bytes per tile are fixed regardless of how many valid rows M carries. The assistant calculated that at batch size 96 with approximately 229 active experts and 32 N-tiles per expert, the system processes roughly 7,330 tile-iterations total, distributed across 188 CTAs handling 23–39 tiles each. The weight traffic alone—loading weights for ~229 active experts across 43 MoE layers—amounts to roughly 29.5 GB of reads per decode step per rank.
Yet the memory controller was only 27% utilized while SMs were 97% active. This suggested the real constraint was not bandwidth but latency: the persistent kernel's CTAs finish their tiny-M tiles quickly and then stall waiting for the next tile's data to arrive from HBM. The GPU's SMs are active but not doing useful work—they are stalled on memory latency, not saturating bandwidth.
The theoretical model was compelling, but it needed empirical validation. The assistant had earlier profiling data showing MoE consuming only 1.9–3.4% of step time at batch size 32, but this was from a period when a dominant indexer kernel was distorting the profile. After the indexer was fixed, MoE's fraction jumped to 28%. The question was: what would a fresh high-batch profile show?
Phase Three: Designing and Executing the High-Batch Profile
Capturing a clean high-batch profile required careful experimental design ([msg 22]). The assistant needed to start a sustained load, wait for the decode batch to fill, trigger the profiler at the right moment, and capture representative traces—all without disrupting the production system.
The assistant verified the decode server's environment configuration ([msg 21]), confirming critical flags: SGLANG_DSV4_FP4_EXPERTS=1 (FP4 expert computation enabled), SGLANG_SM120_MMA_FLASHMLA=1 (FlashMLA-style MMA operations active), SGLANG_SM120_TRITON_INDEXER=1 (Triton-based indexer for SM120), and SGLANG_TORCH_PROFILER_DIR=/root/prof (profiler output directory). The profiler infrastructure was confirmed operational.
The experiment was launched with 80 concurrent connections, 800 total requests, and 512 max tokens per request ([msg 23]). The concurrency of 80 was deliberately chosen to be below the cuda-graph-max-bs of 96 but high enough to stress the system into the regime where nearly all experts would be active. The metrics confirmed stable operation: num_running_reqs held steady at 80.0, and gen_throughput reached approximately 802 tokens/second.
With the system under sustained load, the assistant triggered the PyTorch profiler on the decode worker for 20 steps with shape recording enabled ([msg 24]). The profiler captured the trace files successfully, and the assistant then killed the benchmark to prevent interference.
Phase Four: The Revelation—Parsing the High-Batch Trace
The initial parse of the high-batch trace ([msg 27]) delivered a startling result. The raw breakdown showed:
_mma_sparse_decode_split_kernel(attention): 55.7% of GPU kernel time (964ms across 1596 calls)cutlass_fp4_group_mm(MoE grouped-GEMM): 17.4% (301ms across 1634 calls)ncclDevKernel_AllReduce_Sum_bf16_RING_LL(NCCL all-reduce): 7.2% This was the first hint that something was fundamentally wrong with the investigation's premise. The attention kernel consumed over half of all GPU kernel time, while the MoE grouped-GEMM—the ostensible subject of the deep-dive—accounted for less than a fifth. The assistant's reasoning in [msg 28] shows it grappling with this finding: "This is a significant finding: at batch size 80, attention is actually the dominant bottleneck, not MoE. The performance plan predicted MoE would be 28-35% and the primary driver, but the actual profile shows the opposite." But the assistant did not stop at the headline number. It immediately engaged in causal reasoning to explain why the profile looked this way: "The key insight is that attention work scales linearly with batch size—each token attends to a fixed top-512 KV set, so total attention compute grows directly with B. Meanwhile, MoE work depends on the number of active experts, which saturates as batch increases, causing it to scale sublinearly relative to attention." To validate this hypothesis, the assistant needed to compare against a baseline at lower batch size. It wrote a custom analyzer script (analyze2.py) that computed per-step timing for each kernel category and produced a delta table comparing batch size 32 against batch size 80 ([msg 30]).
Phase Five: The Decisive Cross-Batch Comparison
The output of the cross-batch comparison ([msg 31]) was the single most important piece of evidence in the entire investigation. The per-step delta table told a stark story:
| Category | bs32 (ms/step) | bs80 (ms/step) | Delta | |---|---|---|---| | ATTENTION | 14.22 | 51.18 | +36.96 | | MOE_GROUPGEMM | 15.55 | 17.36 | +1.81 | | NCCL | 6.00 | 8.15 | +2.15 | | MOE_QUANT | 4.40 | 4.69 | +0.29 | | MOE_ROUTING | 3.35 | 2.19 | -1.16 | | MOE_ACT | 3.20 | 2.09 | -1.11 | | GLUE | 8.57 | 7.83 | -0.74 | | OTHER | 2.74 | 1.48 | -1.26 | | TOTAL | 66.37 | 93.83 | +36.95 |
The conclusion was unambiguous. Of the total +36.95ms increase in per-step wall time when batch size grew from 32 to 80, attention alone contributed +36.96ms—more than 100% of the marginal cost. MoE grouped-GEMM contributed only +1.81ms. The remaining categories were flat or even decreased slightly (likely due to better overlap at higher occupancy).
The assistant's reasoning in [msg 32] crystallized the finding: "This flips the premise: MoE is a large fixed cost (~15-16ms/step, batch-independent) that won't improve the C60→C90 slope, but attention is the real lever for raising that ceiling."
Phase Six: Explaining Why MoE Is Batch-Independent
Having made the empirical finding, the assistant then explained why MoE is batch-independent at high concurrency. The analysis was rooted in GPU architecture, memory hierarchy, and kernel design.
The token-to-expert mapping problem: At batch size 80 with top-6 routing, approximately 213 experts are active per step. Each active expert processes roughly 2 tokens on average. The GEMM operation for each expert has an M dimension of approximately 2, while the tile size of the CUTLASS kernel is 128×128×128. This is the fundamental mismatch: the kernel is designed for large M, but at decode time M is tiny.
Latency-bound, not bandwidth-bound: The memory controller utilization of 24-27% and power draw of 337W (56% of 600W TDP) confirmed that the kernel is latency-bound, not bandwidth-bound. When a kernel is bandwidth-bound, the memory controllers are saturated and the GPU is typically power-limited. When latency-bound, the compute units finish their work quickly and then stall waiting for data.
The HBM bandwidth sanity check: The assistant calculated a theoretical lower bound for MoE computation if it were purely HBM-bandwidth-limited: 28 GB per step at 1.5 TB/s yields 18.7 ms. But the measured MoE time was only 15.85 ms—faster than the pure HBM-streaming bound. This forced a self-correction: "Wait, the measured time is actually faster than the pure HBM-streaming lower bound, which means L2 is already catching a meaningful fraction of the weights through reuse across tokens." The headroom for bandwidth-oriented optimization was smaller than initially estimated.
The persistent kernel design: The MoE grouped-GEMM uses a persistent CUTLASS kernel where thread blocks are launched once and remain resident on the SMs, processing tiles from a work queue. However, the persistent scheduler's ability to hide latency is limited when each individual tile is so small that it completes before the next tile's data arrives. The bottleneck is per-tile latency, not pipeline depth.
The Attention Problem: Where Optimization Effort Should Really Go
Having thoroughly analyzed MoE and concluded that its optimization headroom was limited and would not address the C60→C90 scaling problem, the assistant implicitly redirected attention to the real bottleneck: the sparse decode attention kernel.
The attention kernel (_mma_sparse_decode_split_kernel) consumed 56.2% of GPU time at batch size 80, up from 24.5% at batch size 32. The marginal cost of +36.96 ms per step dwarfed everything else. The per-launch timing revealed a heavily skewed distribution: most launches completed in ~86 microseconds, but some extended to 2090 microseconds, with a mean of 604 microseconds. This skew suggested that the split-K adaptivity in the attention kernel was responding to variable context lengths across different decode positions. As requests progressed through their sequences, the accumulated KV cache grew, making later attention operations significantly more costly than early ones.
The assistant also noted a critical nuance: the attention launches per step increased from 158.7 at bs32 to 168 at bs80—less than proportional to the batch size increase. This suggested that the attention kernel's scaling was influenced by context length accumulation, not just batch size. The heavy tail in per-launch timing meant that a small fraction of attention operations accounted for a large fraction of total attention time.
The Verification Step: Grounding Analysis in Configuration
The assistant concluded by verifying several assumptions empirically ([msg 32]). It checked the model configuration, confirming 43 hidden layers, 256 routed experts, 6 experts per token, intermediate size 2048, and hidden size 4096. It searched for SM100/SM120 gating in the CUTLASS MoE path and found that the code only checks for SM100 support (for MXFP8), with no SM120-specific optimization path. It searched for autotune configuration files for FP4 MoE and found nothing—confirming that there were no pre-tuned configurations available for this kernel path.
This verification step closed the loop. The assistant had made a strong empirical case that MoE optimization was not the right lever, and the configuration check confirmed that there were no hidden knobs or alternative implementations that might change the conclusion. The CUTLASS FP4 grouped-GEMM path was what it was, and on SM120, there was no better option available without significant kernel engineering effort.
The Broader Implications
The findings in this investigation have implications beyond this specific model and deployment:
MoE scaling properties: The observation that MoE cost is batch-independent while attention cost scales linearly with batch size is likely a general property of MoE transformer architectures. As batch sizes grow, attention optimization becomes increasingly important relative to MoE optimization.
The small-M problem: The M≈2 problem (few tokens per expert) is inherent to MoE decode with top-k routing at moderate batch sizes. This is a fundamental challenge that cannot be solved by better kernel implementations alone—it requires architectural changes such as expert sharing, larger batch sizes, or different routing strategies.
Latency-bound MoE: The finding that MoE is latency-bound rather than bandwidth-bound at high batch sizes challenges the common assumption that MoE inference is primarily bandwidth-constrained. This may be specific to FP4 quantization and SM120 architecture, but it suggests that bandwidth-oriented optimizations may have diminishing returns once a certain threshold is reached.
The importance of comparative profiling: The key methodological insight is that comparing two batch sizes reveals the marginal cost structure in a way that single-point profiling cannot. Without the bs32 baseline, the assistant would have seen that MoE consumes 17.4% of GPU time and attention consumes 56.2%, but would not have known which one is the scaling bottleneck.
Conclusion: When the Data Flipped the Premise
This investigation is a cautionary tale about the dangers of task framing. The subagent was asked to investigate "MoE grouped-GEMM occupancy deep-dive + high-batch trace," which implicitly assumed that MoE was the problem. Without the assistant's willingness to challenge this premise and follow the data where it led, the investigation could have wasted significant effort optimizing a kernel that was not the bottleneck.
The most important optimization insight sometimes comes not from solving the problem you were asked to solve, but from discovering that you were asking the wrong question. The assistant's systematic methodology—tracing the actual code path, building a theoretical model, designing a careful experiment, and—most crucially—comparing across operating points to separate fixed costs from marginal costs—transformed what could have been a narrow MoE kernel analysis into a fundamental re-examination of where GPU cycles go during MoE model inference at high batch sizes.
The journey from "MoE is the bottleneck" to "attention is the bottleneck" was not a failure of the investigation. It was its greatest success.