The Moment the Bottleneck Shifted: Interpreting GPU Utilization Heatmaps in Custom CUDA Kernel Development

In any deep optimization effort, there comes a pivotal moment when a hard-won improvement reveals the next, deeper constraint — a bottleneck that was always present but hidden behind a more immediate one. Message 12331 of this opencode session captures exactly such a moment. The assistant, having just delivered a dramatic 3–6× decode speedup over the Triton baseline by tuning a custom sm_120 verify attention kernel, receives two user messages containing GPU utilization screenshots. What follows is a masterclass in interpreting performance data, reasoning about architectural constraints, and making pragmatic decisions under uncertainty.

This article examines message 12331 in depth: the reasoning that produced it, the assumptions it tests, the knowledge it creates, and the pivotal role it plays in the broader optimization narrative. The message is notable not for a breakthrough implementation — the kernel edit it contains is a routine continuation of prior work — but for the conceptual synthesis it achieves. The assistant correctly identifies that the optimization frontier has shifted from attention latency to MoE expert imbalance, a structural ceiling that no amount of kernel tuning can address at batch size 1.

The Context: A Custom Kernel in Flight

To understand message 12331, we must first understand the trajectory that led to it. The session concerns deploying the Kimi K2.6 model — a Mixture-of-Experts (MoE) architecture with 384 experts — on an 8× RTX PRO 6000 Blackwell workstation using SGLang with speculative decoding (DDTree). The assistant had built a custom CUDA verify attention kernel for the sm_120 architecture (Blackwell consumer GPUs) because all optimized MLA kernels from FlashMLA, cutlass-MLA, and flashinfer-MLA target sm_90a/sm_100a/sm_103a — architectures with instructions (wgmma, TMA, tcgen05) that simply don't exist on sm_120's Ada-like ISA.

The custom kernel, verify_attn_flash_paged.cu, implements a KV-split flash-decode MLA verify attention with a partial+reduce design. The assistant had already made it CUDA-graph capture-safe and had just discovered a critical performance insight: in the TP8 regime (tensor parallelism across 8 GPUs), each rank has only 8 heads (H=8) instead of the full 64, which drastically reduces the number of GPU blocks the kernel launches. With NSPLIT=16, the grid was only 1×9×8×16 = 1152 blocks — far too few to saturate 188 SMs, leading to occupancy starvation and terrible bandwidth utilization (~2% DRAM). Bumping NSPLIT to 64 (a simple environment variable change, no rebuild required) delivered a 2.1–3.3× speedup across context lengths 4k–65k.

Buoyed by this success, the assistant was in the process of adding 128-bit vectorized bf16 KV loads to squeeze further performance from the paged kernel. The edit had already been applied to the source file, and the assistant was about to rebuild and test — when the user's messages arrived.

The User's Observation: A Heatmap and a Question

The user's two messages (12329 and 12330) arrived in quick succession. The first reads:

Seems we're now into expert compute bound territory, some gpus full on copute and on pcie bw while others pretty idle, maybe try TP8?

The second is a screenshot with a caption:

@2026-05-31-183115_3335x1907_scrot.png - top prefill?? bottom decode?

The user has been monitoring GPU utilization via nvidia-smi or a similar tool and has captured a heatmap showing SM compute utilization across all 8 GPUs over time. They've correctly identified two distinct phases: the top portion of the image shows high, uniform utilization across all GPUs (prefill), while the bottom portion shows lower, patchy utilization (decode). Their diagnosis — "expert compute bound territory, some gpus full on compute and on pcie bw while others pretty idle" — is remarkably accurate. Their suggestion, "maybe try TP8?", is where the interesting reasoning begins.

The Assistant's Reasoning: Interpreting the User

The assistant's Agent Reasoning section opens with a careful parsing of the user's intent:

The user is showing me a performance heatmap with questions about prefill versus decode utilization, and they're suggesting TP8 — but we're already using TP8, so they likely meant expert parallelism (EP) for the MoE model instead, or something else entirely to address the compute bottleneck where some GPUs are maxed out while others are idle.

This is a critical moment of reasoning. The user said "try TP8" — but the system is already configured with TP8 (tensor parallelism degree 8). The assistant recognizes this discrepancy and correctly infers that the user probably meant EP (expert parallelism), not TP. This is a natural confusion: both are parallelism strategies for distributing model computation across GPUs, but they work differently. TP splits individual layers' weight matrices across GPUs (each GPU holds a slice of every layer), while EP splits the experts themselves (each GPU holds a disjoint set of experts). For an MoE model like Kimi K2.6 with 384 experts, EP would mean each GPU holds 48 experts, and tokens are routed only to the GPUs whose experts they activate.

The assistant then performs a detailed analysis of the heatmap:

Looking at the per-SM compute heatmap, the top section is definitely prefill (high utilization across all GPUs, compute-bound) while the bottom is decode (lower and patchier, showing expert load imbalance).

This interpretation is grounded in the model's architecture. During prefill, the model processes a long prompt (thousands of tokens) in parallel. All tokens route to their respective experts, and with enough tokens, all 384 experts are likely activated, distributing work evenly across all 8 GPUs. The result is high, uniform compute utilization — exactly what the top of the heatmap shows.

During decode, the situation is radically different. With speculative decoding using DDTree, the verify step processes only ~9 tokens (the draft tree). Each token activates roughly 8 experts (the top-k routing in MoE), so we have about 72 expert activations total. With 384 experts spread across 8 GPUs (48 per GPU), the probability that any given GPU's subset of experts is hit depends on the routing pattern. With only 9 tokens and random expert selection, some GPUs will be hit multiple times while others are hit rarely or not at all. The result is the patchy utilization the user observed: some GPUs saturated on both compute and PCIe bandwidth, others sitting idle.

The MoE Imbalance: A Structural Ceiling

The assistant's reasoning then deepens into a quantitative analysis of why this imbalance is fundamental:

At bs=1 decode (verify q=9), only ~9 tokens route to ~8 experts each → up to ~72 expert activations across 384 experts on 8 GPUs. With TP, each GPU holds 384/8=48 experts. The 9 tokens' selected experts are scattered → some GPUs have hits, others don't → IMBALANCE + the idle GPUs wait (sync at all-reduce). And PCIe BW = the TP all-reduce / all-gather of the MoE output.

This is the key insight. Under tensor parallelism, every layer's computation requires an all-reduce or all-gather communication step to synchronize partial results across GPUs. When some GPUs have no work (their experts weren't activated) and others have work, the idle GPUs still must participate in the communication barrier. The result is that the busy GPUs are slowed down by the communication overhead, while the idle GPUs waste cycles waiting. The PCIe bandwidth between GPUs becomes a bottleneck because the all-reduce of MoE outputs must traverse the PCIe bus (these are PCIe-connected GPUs, not NVLink).

The assistant correctly identifies this as a fundamental limitation:

This is the fundamental low-batch MoE inefficiency: few tokens → few experts active → imbalanced → idle GPUs + comm overhead. The fix for THROUGHPUT is batching (more tokens → all experts active → balanced). For single-request LATENCY, it's inherent (can't balance 9 tokens across 8 GPUs' 384 experts).

This distinction between throughput and latency is crucial. Batching multiple requests together would increase the token count, activating more experts and improving balance — but it wouldn't help a single request's decode latency. For single-request latency, the MoE imbalance at low batch is a structural ceiling that no amount of kernel optimization can突破.

The Pragmatic Decision: Build First, Analyze Later

Despite this deep insight, the assistant makes a pragmatic decision:

I was mid-flight on the vectorized-load kernel edit and should finish the build+test, but the user's observation reframes things: further attention-kernel optimization has diminishing returns if MoE now dominates. Let me verify by building and testing the vectorized load to see if it still helps the attention portion.
My plan: build the vectorized-load kernel (already edited), rebuild, restart, measure — this tells me if attention still matters or MoE dominates. Then address the user's MoE-bound observation in my response, defrag Tier 0, record, and commit.

This is a textbook example of evidence-driven decision-making. Rather than assuming that MoE imbalance makes further attention optimization pointless, the assistant decides to measure whether the vectorized loads still produce a measurable improvement. If decode throughput barely moves, it confirms MoE dominance. If it improves, attention is still a meaningful contributor. Either way, the data will guide the next step.

The assistant also clarifies the "try TP8" comment in its response to the user:

The bottleneck has shifted. My kernel + NSPLIT fix removed the attention-latency stall, so decode is now MoE-expert-bound with load imbalance (some GPUs full on compute+PCIe, others idle) — exactly your read. This is the inherent low-batch MoE behavior: at bs=1/verify-q=9, only ~9 tokens route to ~8 experts each, so only some of each rank's 48 experts (384/8) get hit → idle ranks + a PCIe all-reduce sync. We're already on TP8; the imbalance is intrinsic to TP+low-batch, and even EP can't balance 9 tokens across 8 GPUs — the real levers are batching (throughput) or EP+EPLB for many concurrent requests, not single-request latency.

This response serves multiple purposes. It validates the user's observation (the heatmap interpretation is correct), provides a clear technical explanation for the observed behavior, corrects the misunderstanding about TP8 (we're already using it), and sets realistic expectations about what can and cannot be fixed at the kernel level.

The Knowledge Created by This Message

Message 12331 creates several important pieces of knowledge that shape the remainder of the session:

1. The bottleneck hierarchy is now understood. The original bottleneck was the Triton MLA verify attention with page_size=1, causing scattered KV access at ~14 GB/s effective bandwidth (130× below the 1.8 TB/s peak). The custom kernel + NSPLIT tuning fixed this, achieving 3–6× speedup. The new bottleneck is MoE expert imbalance under TP at low batch. This hierarchy is now empirically confirmed: attention → MoE → (next bottleneck).

2. The TP8 misunderstanding is resolved. The user suggested trying TP8, but the system was already on TP8. The assistant correctly inferred that the user likely meant EP (expert parallelism) and explained why even EP wouldn't help single-request latency at batch size 1. This prevents wasted effort on a misdirected optimization.

3. The prefill/decode utilization signature is characterized. Prefill: high, uniform compute utilization across all GPUs (compute-bound, healthy). Decode: lower, patchy utilization with some GPUs saturated and others idle (MoE-imbalance-bound). This signature becomes a diagnostic tool for future performance analysis.

4. The limits of kernel optimization are established. The assistant recognizes that no amount of attention kernel tuning can fix MoE imbalance at batch size 1. This is a structural ceiling imposed by the combination of MoE routing, tensor parallelism, and low token count. The levers for improvement are architectural (batching, EP with load balancing), not kernel-level.

5. A measurement methodology is reinforced. The assistant's decision to finish the vectorized-load build and measure its impact before drawing conclusions demonstrates a commitment to evidence over assumption. This methodology — instrument, measure, analyze, decide — is a recurring theme throughout the session.

Assumptions and Their Validation

The message rests on several assumptions, most of which are explicitly stated and testable:

Assumption 1: The heatmap shows prefill (top) and decode (bottom). This is a visual interpretation of GPU utilization data. The assistant's reasoning — prefill has many tokens and is compute-bound (high utilization), decode has few tokens and is MoE-bound (patchy utilization) — is consistent with the model's architecture and the observed behavior. This assumption could be validated by correlating the heatmap timestamps with request logs.

Assumption 2: The user meant EP when they said TP8. This is a charitable interpretation that assumes domain knowledge (the user knows the system is TP8 and is suggesting a different parallelism strategy). The alternative — that the user forgot the system was already TP8 — is less productive and would lead to a different response. The assistant's interpretation is both respectful and technically sound.

Assumption 3: MoE imbalance is the dominant bottleneck at batch size 1. This is supported by the heatmap data and the quantitative analysis (9 tokens × ~8 experts = ~72 activations across 384 experts on 8 GPUs). However, the assistant wisely decides to test this by completing the vectorized-load build: if decode throughput improves significantly, attention is still a meaningful contributor.

Assumption 4: PCIe bandwidth is a limiting factor. The GPUs are PCIe-connected (not NVLink), so the all-reduce communication for MoE outputs must traverse the PCIe bus. This is consistent with the user's observation that some GPUs are "full on compute and on pcie bw." The assistant correctly identifies this as a contributor to the overall bottleneck.

The Thinking Process: A Window Into Expert Debugging

The Agent Reasoning section of message 12331 provides a rare, detailed look into how an expert systems engineer thinks through a complex performance problem. Several patterns are worth highlighting:

Pattern 1: Interpret the user's intent. The assistant doesn't take "try TP8" literally. Instead, it considers what the user meant given the context (we're already TP8, the user sees MoE imbalance, they're suggesting a fix). This requires both technical knowledge (what TP and EP are) and social reasoning (the user is an expert who knows the system configuration).

Pattern 2: Connect observation to architecture. The heatmap shows patchy GPU utilization. The assistant connects this to the specific architectural feature that causes it: MoE routing with few tokens under TP. This isn't generic reasoning — it's specific to the Kimi K2.6 model (384 experts, top-8 routing), the parallelism strategy (TP8), and the batch size (bs=1 with verify q=9).

Pattern 3: Distinguish throughput from latency. The assistant explicitly separates the two optimization goals. Batching improves throughput by activating more experts and improving balance. But single-request latency at bs=1 is fundamentally limited by the few tokens' expert activations. This distinction prevents conflating two different optimization problems.

Pattern 4: Use measurement to resolve uncertainty. Rather than assuming MoE dominance, the assistant decides to finish the vectorized-load build and measure. This is a commitment to evidence that characterizes the entire session. The assistant repeatedly instruments code, profiles components, and lets data guide decisions.

Pattern 5: Communicate clearly with the user. The response to the user validates their observation, explains the technical reasoning, corrects the misunderstanding gently, and sets realistic expectations. It's a model of how to communicate complex performance analysis to a technical collaborator.

The Broader Significance

Message 12331 is significant not just for its immediate content but for what it represents in the optimization journey. The assistant has successfully:

  1. Diagnosed the original bottleneck (scattered KV access in Triton MLA)
  2. Built a custom kernel to address it (verify_attn_flash_paged.cu)
  3. Made the kernel production-ready (CUDA graph capture-safe)
  4. Tuned it for the target architecture (NSPLIT=64 for occupancy)
  5. Identified the next bottleneck (MoE imbalance)
  6. Communicated the findings to the user
  7. Made a pragmatic plan for next steps This is the arc of a successful optimization effort: identify, fix, measure, identify the next constraint, repeat. Each cycle reveals a deeper layer of the performance stack. The assistant has moved from the CUDA kernel level (attention) to the system architecture level (MoE parallelism), and the next steps will likely involve batching strategies or expert parallelism with load balancing. The message also demonstrates the importance of GPU utilization monitoring in performance debugging. The user's heatmap screenshot provided the key insight that shifted the optimization focus. Without it, the assistant might have continued optimizing the attention kernel well past the point of diminishing returns. The combination of automated profiling (the assistant's instrumentation) and human visual pattern recognition (the user's heatmap interpretation) proved more powerful than either alone.

Conclusion

Message 12331 captures a moment of synthesis in a complex optimization effort. The assistant, mid-flight on a kernel edit, receives new data from the user that reframes the entire problem. Rather than dismissing the user's observation or blindly following their suggestion, the assistant performs a careful analysis: interpreting the heatmap, connecting it to architectural constraints, correcting a misunderstanding about parallelism strategies, and making a pragmatic decision to gather more data before pivoting.

The message creates lasting knowledge: the bottleneck hierarchy is now understood (attention → MoE), the prefill/decode utilization signature is characterized, the limits of kernel optimization are established, and a measurement methodology is reinforced. It also demonstrates the value of evidence-driven reasoning, the importance of understanding architectural constraints, and the art of communicating complex technical findings to a collaborator.

The assistant's final action — rebuilding the vectorized-load kernel — is both a continuation of prior work and a test of the new hypothesis. If decode throughput improves, attention still matters. If it doesn't, MoE truly dominates. Either way, the data will speak, and the next optimization cycle will begin from a position of knowledge rather than assumption. This is the essence of disciplined performance engineering: let the machine tell you where the bottleneck is, and then fix it.