The GPU Utilization Probe That Reframed a Performance Mystery

Introduction

In the middle of a deep investigation into why long-context speculative decoding was running at a crawl, a single diagnostic message arrived that fundamentally shifted the team's understanding of the bottleneck. Message [msg 12186] contains the results of a GPU utilization probe run against a 64k-context decode workload on an RTX PRO 6000 Blackwell GPU. The numbers it returned were simultaneously illuminating and puzzling: 99.8% SM utilization during decode, yet only 1.2% memory controller utilization and a mere 133.6W of power draw. These three numbers told a story that contradicted the prevailing hypothesis and sent the investigation in a new direction.

The Context: A Slowdown That Defied Explanation

The session leading up to this message had been an exhaustive investigation into why the Kimi K2.6 model with DFlash speculative decoding was experiencing severe decode slowdowns at long context lengths. Earlier benchmarks had shown step times ballooning from 159ms at 32k tokens to a staggering 1,430ms at 185k tokens — a roughly 9× increase for a 5.8× increase in context length. The theoretical floor for reading the KV cache at 1.8 TB/s was about 5ms, making the observed 1,430ms a ~286× overhead.

The initial suspect was the page_size=1 configuration, which forces the KV cache to be paged one token per page. This creates a scatter-gather memory access pattern where the attention kernel must chase pointers across 128k individual pages rather than performing coalesced memory reads. The Triton-based MLA attention kernel was presumed to be thrashing on this fragmented access pattern, causing the memory subsystem to stall.

But before diving into a kernel rewrite, the assistant needed data. The team had already ruled out one major concern — that the drafter was failing at long context. Through a series of carefully designed probes using predictable "The cat sat on the mat" repetition text, the assistant confirmed that the drafter maintained a healthy commit length of 7–8 tokens at 1k, 16k, 64k, and 128k context lengths alike. The commit_len=1 observed in earlier benchmarks was purely an artifact of the synthetic "continue this discussion" prompt producing genuinely unpredictable novel prose, not a long-context bug.

With the drafter exonerated, attention turned to GPU utilization. The question was simple: what is the GPU actually doing during those painfully slow decode steps? Is it memory-bound, thrashing on page-table walks? Is it compute-bound, doing too much work per token? Or is it stalled on something else entirely?

The Message: A Probe Designed for Clarity

The message itself is a single bash command execution followed by its output. The assistant copies a GPU utilization probe script to the remote server and runs it with a carefully chosen workload: 64k tokens of novel prose, generating 80 tokens. The novel prose is deliberate — it forces the drafter to achieve commit_len=1, meaning every decode step produces exactly one new token. This eliminates the variable of variable-length draft acceptance and ensures a clean measurement of the per-step cost.

The probe script samples GPU metrics from nvidia-smi at intervals during both the prefill and decode phases. The results are striking:

prompt~64000 gen=80  prefill=87.5s decode=26.2s (GPU0, 211 samples)
  PREFILL: SM%= 99.4  memctrl%=  5.0  power_W= 301.7  sm_clk=  2415  n=162
  DECODE : SM%= 99.8  memctrl%=  1.2  power_W= 133.6  sm_clk=  2359  n=49

What the Numbers Actually Mean

To understand why these numbers are so important, we need to decode what each metric represents in the context of GPU architecture.

SM% (Streaming Multiprocessor utilization) — This is the percentage of time that at least one warp is active on a multiprocessor. At 99.8% during decode, the SMs are never idle. There is always at least one warp running. But crucially, this does not mean all warps are doing useful arithmetic — it only means the SMs are not completely empty.

memctrl% (Memory Controller utilization) — This measures how busy the memory controllers are servicing requests. At 1.2%, the memory subsystem is essentially idle. The GPU is barely touching its HBM memory during decode.

power_W — The GPU power draw. During prefill, the GPU draws 301.7W — near its thermal design power, indicating heavy compute activity. During decode, it draws only 133.6W — less than half, despite the SMs being "99.8% utilized."

The contradiction is immediately apparent: how can the SMs be 99.8% busy while the memory controllers are 1.2% utilized and the power draw is only 133.6W? The answer is that the SMs are occupancy-bound, not compute-bound. The warps are present on the SMs but stalled — waiting on instruction latencies, shared memory bank conflicts, synchronization barriers, or pipeline hazards. They consume scheduler resources and keep the SM "occupied" in the nvidia-smi sense, but they are not doing useful arithmetic or memory access.

This is a fundamentally different diagnosis from the page_size=1 scatter-gather hypothesis. If the bottleneck were memory-side — page walks, TLB misses, or non-coalesced HBM access — we would expect to see elevated memctrl% as the memory controllers struggle to service the scattered requests. Instead, we see 1.2%, meaning the memory subsystem is twiddling its thumbs. The bottleneck is entirely within the compute pipeline of the SM.

Reframing the Problem

This single message reframed the entire performance investigation. The earlier hypothesis — that page_size=1 was causing memory scatter-gather overhead — was not supported by the data. The memory controllers were barely active. Something else was causing the SMs to stall.

The low power draw is the most telling signal. During prefill, the GPU draws 301.7W because the SMs are doing intensive matrix multiply operations — tensor cores firing, arithmetic units busy, registers and shared memory in heavy use. During decode, the GPU draws only 133.6W despite 99.8% SM "utilization" because the warps are stalled. A stalled warp consumes almost no dynamic power — it's just sitting in a scheduler queue waiting for its dependencies to resolve. The SM reports as "occupied" because there's a warp assigned to it, but that warp isn't executing instructions.

This points to a kernel that is instruction-throughput-limited or shared-memory-bank-conflict-limited rather than memory-bandwidth-limited. The Triton-generated MLA attention kernel, when operating at long context with page_size=1, likely generates a code path where warps spend most of their time waiting for data from other warps, resolving shared memory bank conflicts, or synchronizing across thread blocks. The actual arithmetic and memory access are sparse, but the warp scheduling overhead is enormous.

For the RTX PRO 6000 Blackwell (sm_120 architecture), this is particularly relevant. The Blackwell consumer GPU (sm_120) has a different ISA from the datacenter Blackwell (sm_100/sm_120a), lacking the advanced tensor core instructions (wgmma, TMA, tcgen05) that make attention kernels efficient on Hopper and datacenter Blackwell. The assistant had already discovered that all optimized MLA kernels — FlashMLA, cutlass-MLA, flashinfer-MLA — were compiled only for sm_90a/sm_100a/sm_103a and none supported sm_120. The Triton kernel that was running was a fallback path, and the GPU utilization data now showed just how inefficient that fallback was.

The Thinking Process Visible in the Message

The assistant's reasoning leading up to this message shows a disciplined, hypothesis-driven investigation. The chain of reasoning went:

  1. Observe the symptom: Step time grows super-linearly with context length (159ms at 32k → 1,430ms at 185k).
  2. Form hypothesis A: The drafter is failing at long context, causing commit_len=1 and eliminating the benefit of speculative decoding. Test this by running predictable text at multiple context lengths. Result: drafter works fine (commit_len ~7-8 at all lengths). Hypothesis A rejected.
  3. Form hypothesis B: The page_size=1 KV cache is causing memory scatter-gather overhead. The Triton kernel is thrashing on page-table walks. This would manifest as high memory controller utilization.
  4. Test hypothesis B: Measure GPU utilization during decode. If memory controllers are busy, hypothesis B is supported. If not, we need a new hypothesis.
  5. Execute the test: Write a GPU utilization probe that samples nvidia-smi metrics during a long-context decode with novel prose (to force clean single-token steps).
  6. Analyze the result: memctrl=1.2%, power=133.6W. Hypothesis B is not supported. The bottleneck is not memory-side. This is textbook diagnostic methodology: form hypotheses, design experiments that can falsify them, execute, and let the data guide the next step. The assistant did not assume the answer — it went and measured.

Assumptions Made and Corrected

Several assumptions were implicitly or explicitly present in this investigation:

Assumption 1: The decode slowdown is caused by memory access patterns. This was the prevailing hypothesis before the probe. The page_size=1 configuration was the prime suspect. The probe data disproved this — memory controller utilization was negligible.

Assumption 2: High SM utilization means the GPU is busy doing useful work. The 99.8% SM utilization number could easily be misinterpreted as "the GPU is fully utilized." But the power draw and memory controller data told a different story: the SMs were occupied but stalled. This is a subtle but critical distinction that the assistant correctly interpreted.

Assumption 3: The Triton-generated MLA kernel is reasonably efficient for the Blackwell architecture. The probe data showed this was false — the kernel was producing massive warp stalls and consuming power at barely half the GPU's capacity.

Assumption 4: The prefill and decode phases would show similar GPU utilization patterns. They did not. Prefill was genuinely compute-bound (99.4% SM, 301.7W, 5% memctrl — a healthy profile). Decode was occupancy-bound (99.8% SM, 133.6W, 1.2% memctrl — a pathological profile). This asymmetry is a key insight: the same GPU, the same model, but fundamentally different bottlenecks depending on the phase of inference.

Input Knowledge Required

To fully understand this message, the reader needs:

Output Knowledge Created

This message produced several critical pieces of knowledge:

  1. The decode bottleneck is not memory-bound. Memory controller utilization at 1.2% rules out the page_size=1 scatter-gather hypothesis as the primary cause.
  2. The decode kernel is occupancy-bound with stalled warps. The combination of 99.8% SM utilization, 133.6W power, and 1.2% memctrl indicates warps are present on SMs but stalled on instruction latencies or synchronization.
  3. Prefill and decode have fundamentally different bottleneck profiles. Prefill is compute-bound (high power, high SM, moderate memctrl). Decode is occupancy-bound (low power, high SM, low memctrl). Any optimization strategy must account for this asymmetry.
  4. The Triton MLA kernel on sm_120 is highly inefficient for long-context decode. The kernel keeps SMs occupied but delivers minimal useful work per cycle, as evidenced by the low power draw.
  5. A custom kernel is likely necessary. The existing optimized MLA kernels (FlashMLA, cutlass-MLA, flashinfer-MLA) don't support sm_120, and the Triton fallback is demonstrably inefficient. This finding directly motivated the subsequent work on building a custom sm_120 verify attention kernel.

Conclusion

Message [msg 12186] is a masterclass in diagnostic discipline. Faced with a confusing performance regression, the assistant did not guess or rely on intuition — it designed a targeted experiment, executed it cleanly, and let the data speak. The data said something unexpected: the GPU was "99.8% utilized" yet drawing only 133.6W, with memory controllers at 1.2%. This contradiction revealed that the bottleneck was not where anyone expected it to be.

The probe results reframed the entire investigation. Instead of optimizing memory access patterns (which would have yielded minimal gains given the 1.2% memctrl), the team could now focus on the real problem: the Triton-generated MLA kernel was producing massive warp stalls on sm_120, keeping SMs occupied without delivering useful throughput. This insight directly led to the development of a custom sm_120 verify attention kernel — work that would eventually achieve a 3–6× decode speedup over the Triton baseline.

In the broader narrative of the coding session, this message represents the turning point where a vague suspicion ("something is slow at long context") crystallized into a specific, measurable, and actionable diagnosis ("the decode kernel is occupancy-bound on sm_120 due to warp stalls"). Without this probe, the team might have spent weeks optimizing memory access patterns that were never the real problem. With it, they had a clear target: build a kernel that keeps warps genuinely busy, not just occupied.