The Negative Result That Saved Months of Work: How Empirical Evidence Killed a GPU Kernel Optimization Hypothesis

Introduction

In the high-stakes world of GPU kernel optimization for large language model inference, the most valuable result is often the one you didn't expect: a definitive "no." Message [msg 13569] of this opencode session captures precisely such a moment—a systematic, evidence-driven conclusion that two well-motivated optimization levers for a Triton attention kernel were dead ends, and that the kernel was already operating at its optimal configuration. This message is a masterclass in disciplined experimental science applied to systems engineering, demonstrating how rigorous A/B testing, careful reasoning about GPU architecture, and the courage to accept negative results can save weeks or months of wasted engineering effort.

The message is the culmination of a focused investigation into the mma_sparse_decode_split kernel—a custom Triton kernel powering the DeepSeek-V4-Flash model's sparse attention mechanism on NVIDIA Blackwell RTX PRO 6000 GPUs (sm_120 architecture). The assistant had been chasing a hypothesis that the kernel's throughput (saturating at ~845 tok/s at high concurrency) could be improved by increasing GPU occupancy—either by using more warps per streaming multiprocessor (SM) or by deepening the software pipeline to better hide memory latency. What followed was a textbook example of hypothesis-driven optimization, complete with controlled experiments, careful measurement, and the intellectual honesty to abandon a cherished hypothesis when the data contradicted it.

The Context: Two Experiments, Two Failures

To understand the significance of [msg 13569], we must first understand what came before it. The assistant had formulated an "occupancy hypothesis": that the decode attention kernel was bottlenecked by poor GPU occupancy—the ratio of active warps to the maximum a given SM can support. The intuition was sound: if the kernel is latency-bound (spending most of its time waiting for memory loads rather than computing), then increasing the number of active warps per SM should help hide that latency by keeping the compute units busy while some warps wait for data.

The first experiment (V1, [msg 13559]–[msg 13562]) tested the "more warps" lever. The assistant modified the kernel's autotune configuration to force 8 warps per SM instead of the baseline 4, expecting improved latency hiding. The results were unambiguous and damning: performance regressed across every tested concurrency level, from −14% at C=1 to −3% at C=96. The root cause was clear upon reflection: the kernel's MMA (matrix multiply-accumulate) tiles are tiny—[32, 16] per split—and subdividing such small tiles across 8 warps creates severe inefficiency. Each warp gets only a sliver of computation, the MMA units are underutilized, and synchronization overhead dominates. The autotuner's original choice of 4 warps was correct.

The second experiment (V2, [msg 13564]–[msg 13568]) tested the "deeper prefetch" lever. Instead of increasing warps, the assistant increased num_stages from 2 to 3, keeping the efficient 4-warp configuration. This deepens the software pipeline, allowing more K-tiles to be prefetched before computation begins, theoretically hiding gather latency. The assistant first had to resolve the block_h parameter—the number of attention heads processed per block—which required tracing through the model code to compute n_local_heads = 64 // attn_tp_size(4) = 16, confirming block_h=16. This meant the shared memory budget was ~60KB for the baseline, and num_stages=3 would use ~81KB—fitting within the 99KB limit. The experiment was feasible.

V2's results, presented and analyzed in [msg 13569], were equally disappointing. At C=64, the three-stage pipeline showed a −9.9% regression compared to baseline. At C=96, a −2.6% regression. At low concurrency (C=1, C=8), regressions of −5 to −7%. Only at C=48 was performance approximately flat. The deeper pipeline was hurting, not helping.

The Reasoning: Why the Levers Failed

The assistant's reasoning in [msg 13569] reveals a deep understanding of GPU architecture and the specific characteristics of this kernel. The core insight is that the loop iterations are too short to benefit from a deeper pipeline. At high concurrency, each block processes only 4–6 K-tile iterations. With num_stages=3, the pipeline needs at least 3 iterations to fill before it reaches steady state, and then additional iterations to drain. When the total iteration count is only 4–6, the fill and drain phases dominate the computation, and the increased shared memory pressure (81KB vs 60KB) adds register pressure that further degrades performance.

This is a classic trade-off in software pipelining: deeper pipelines help when loop trip counts are large enough to amortize the fill/drain overhead. For short loops, the overhead dominates. The assistant correctly identified that the kernel's working set at high concurrency (where each split handles a small slice of the total KV cache) simply doesn't provide enough iterations for num_stages=3 to be beneficial.

The reasoning also connects V2's failure back to V1's failure, synthesizing a unified conclusion: both occupancy-increasing levers are counterproductive for this kernel shape. The tiny MMA tiles ([32, 16]) that make 8 warps inefficient also constrain the benefits of deeper pipelining. The kernel is not latency-hiding-bound in the way the assistant initially hypothesized—the bottleneck lies elsewhere, likely in the gather bandwidth from the KV cache or in the inherent work pattern of the sparse attention computation.

The Decision: Revert, Document, and Move On

The most important decision in [msg 13569] is the decision to stop. The assistant reverts the kernel to the known-good baseline, verifies the revert with an MD5 checksum, restarts the decode service, waits for it to become healthy, runs a quick benchmark to confirm baseline performance is restored (C=64 at 782.0 tok/s, consistent with historical data), and then prepares to document the negative result.

This decision carries significant weight because the assistant had been pursuing a larger vision: a "full register-reducing rewrite" of the attention kernel that would fundamentally restructure the computation to enable 8 warps per SM. The V1 result had already "downgraded the full-rewrite premise" ([msg 13562]), but V2's failure seals the case. The assistant explicitly states: "the full register-reducing rewrite is empirically refuted (its premise — 8 warps/SM hides the gather — made things worse)."

The discipline here is remarkable. It would have been easy to rationalize the negative results as measurement noise, or to propose yet another variant (V3 with batch-aware warp sizing, which the assistant explicitly dismisses as "moot since w8 lost at every batch"). Instead, the assistant accepts the data, reverts to the working configuration, and preserves the one optimization that did work: the TARGET_CTAS=512 wave-fill setting that delivered 6–13% gains at high concurrency. The message ends with the system back in a known-good state, all experimental changes rolled back, and a clear path forward.

Input Knowledge Required

To fully understand [msg 13569], several pieces of domain knowledge are essential:

GPU Architecture (NVIDIA Blackwell sm_120): The kernel targets the Blackwell architecture's sm_120 compute capability. Key concepts include streaming multiprocessors (SMs), warps (groups of 32 threads), shared memory (on-chip scratchpad, ~99KB per SM for this generation), and the Tensor Memory Accelerator (TMA) for asynchronous data movement. The num_warps and num_stages parameters control how the Triton compiler schedules work across these hardware resources.

Triton Compiler and Autotuning: The kernel uses Triton's @triton.autotune decorator, which compiles multiple candidate configurations and selects the fastest at runtime based on input shapes. The BLOCK_T parameter controls the tile size along the sequence dimension, num_warps controls thread-block size, and num_stages controls software pipeline depth. The autotuner's configs list defines the search space.

Attention Kernel Structure: The mma_sparse_decode_split kernel implements sparse multi-head latent attention (MLA) for the DeepSeek-V4 architecture. It uses "split-K" parallelization, where each CTA (cooperative thread array) processes a slice of the total KV cache, and results are combined via a cross-CTA reduction. The TARGET_CTAS parameter controls how many CTAs are launched to fill the GPU's SMs (wave quantization).

DeepSeek-V4 Model Architecture: The model uses 64 attention heads, sharded across 4 attention-tensor-parallelism (attn_tp) ranks, yielding 16 local heads per rank. This determines block_h=16—the number of heads processed per kernel block. The n_hg=1 (number of head groups) follows from the MLA architecture where query and key heads are grouped.

Benchmark Methodology: The bench_tput.py script measures aggregate throughput (tok/s), per-request throughput, latency percentiles, and error counts across varying concurrency levels (C). The n parameter controls total requests, and max_tokens limits generation length. The assistant consistently uses n=C*4 and max_tokens=256 for apples-to-apples comparison.

Output Knowledge Created

[msg 13569] produces several concrete outputs:

1. Empirical refutation of the occupancy hypothesis: Both tested levers (more warps, deeper pipeline) regress performance. The kernel is already optimal at BLOCK_T=16, num_warps=4, num_stages=2 among feasible configurations.
2. Confirmation of block_h=16: Through careful code tracing, the assistant resolves the ambiguity between two subagent estimates, confirming that n_local_heads = 64 // 4 = 16, which means block_h = min(32, 16) = 16. This is critical for shared memory budgeting and future kernel work.
3. A restored, verified baseline: The system is returned to a known-good state with the original 3-config autotune list, the TARGET_CTAS=512 optimization intact, and benchmark numbers matching historical data (C=64 at ~782 tok/s).
4. Documentation of the negative result: The assistant updates the DSV4_ATTN_OCCUPANCY_AB.md document with the V1 and V2 results, preserving the knowledge for future reference and preventing others from pursuing the same dead ends.
5. A clear go-forward direction: The message implicitly defines what not to work on next. The assistant's subsequent efforts ([msg 13570] onward) pivot to other optimization avenues, informed by the knowledge that occupancy tuning is exhausted.

Assumptions and Potential Mistakes

The assistant makes several assumptions that deserve scrutiny:

Assumption 1: The benchmark is representative of production workloads. The benchmark uses synthetic requests with fixed max_tokens=256 and no tool-calling or multi-turn interaction. Real agentic workloads involve variable-length generation, tool call formatting, and multi-turn context accumulation. It's possible that occupancy tuning could benefit different workload profiles, though the assistant's reasoning about short loop iterations applies broadly to decode-phase attention.

Assumption 2: The autotuner's baseline configs are optimal. The assistant treats the original 3-config autotune list as the ground truth, but this list was itself produced by an earlier autotuning process that may have had its own biases (e.g., limited search space, cold-start artifacts). The assistant never questions whether the baseline could be improved through a different mechanism—only that the specific occupancy levers tested are ineffective.

Assumption 3: The regression at C=64 (−9.9%) is caused by pipeline overhead. While the reasoning is plausible, the assistant doesn't profile the kernel to confirm. The regression could also be caused by increased register pressure spilling to local memory, or by changes in the compiler's instruction scheduling for the deeper pipeline. A roofline analysis or instruction-level profiling would strengthen the causal claim.

Assumption 4: The "full register-reducing rewrite" premise is fully refuted. The assistant concludes that because num_warps=8 regresses, any rewrite targeting 8 warps/SM is doomed. However, a rewrite could fundamentally change the computation structure (e.g., fusing the cross-CTA reduction, using different tile shapes) such that 8 warps become efficient. The assistant is correct that the current kernel shape doesn't benefit from 8 warps, but a different kernel shape might. The conclusion is specific to the current implementation, not a universal impossibility.

Potential mistake: Not testing intermediate concurrency values more granularly. The benchmark tests C=1, 8, 48, 64, 80, 96. The regression at C=64 is the most severe, but C=48 is flat. A more granular sweep between C=48 and C=64 might reveal a threshold where the pipeline overhead becomes dominant, providing deeper insight into the mechanism. However, given the consistent negative trend at both lower and higher values, this is unlikely to change the overall conclusion.

The Thinking Process: A Window into Debugging Methodology

The reasoning section of [msg 13569] reveals the assistant's cognitive process in remarkable detail. It begins with raw data analysis ("Looking at the performance benchmarks... I'm seeing mixed results"), then immediately moves to causal reasoning ("The deeper prefetch pipeline... isn't delivering the expected latency-hiding benefits, likely because the loop iterations are too short"). This is followed by a quantitative justification ("At C=64 specifically, the eight-way split with 64-byte chunks means the three-stage pipeline barely fills before draining"), and then a generalization ("This tells me that both occupancy-increasing levers... are actually regressing performance").

The assistant then synthesizes the V1 and V2 results into a unified conclusion: "the attention kernel isn't improvable through these autotune knobs." This is a strong claim, and the assistant immediately qualifies it by identifying what isn't the bottleneck: "The ~845 tok/s saturation point isn't a latency-hiding problem I can fix with occupancy adjustments—it's something more fundamental about the gather bandwidth or the inherent work pattern."

The final paragraph of the reasoning shows the assistant planning the reversion and documentation, demonstrating an engineering discipline that prioritizes system stability and knowledge preservation over the sunk cost of the investigation. The decision to revert is presented as non-negotiable: "I'm reverting V2 back to the baseline, verifying it's healthy, documenting this result and the overall conclusion that #3a is a dead end, and then committing."

Broader Significance

[msg 13569] is more than just a negative experimental result—it's a case study in how to do GPU kernel optimization correctly. The hallmarks of good methodology are all present:

1. Clear hypothesis formulation: The occupancy hypothesis is stated explicitly and connected to a specific mechanism (latency hiding).
2. Isolated variable changes: Each experiment changes exactly one parameter (num_warps or num_stages) while holding everything else constant.
3. Apples-to-apples comparison: Benchmarks use identical parameters (n=C*4, max_tokens=256) across all experiments.
4. Verification of system state: The assistant verifies the baseline is restored with MD5 checksums and benchmark confirmation.
5. Documentation of negative results: The findings are recorded in a living document, preserving the knowledge for the team.
6. Resistance to escalation bias: The assistant explicitly rejects the temptation to propose yet another variant (V3), accepting that the hypothesis is falsified. In a field where "just one more optimization" is a constant temptation, and where negative results are rarely published or shared, this message stands as a model of intellectual honesty and engineering rigor. The months of engineering effort that might have been wasted on a full kernel rewrite were saved by two focused experiments and the discipline to accept their verdict.

Conclusion

[msg 13569] is the denouement of a tightly-scoped investigation into GPU kernel occupancy optimization. Through two controlled experiments (V1: num_warps=8, V2: num_stages=3), the assistant definitively showed that the mma_sparse_decode_split kernel is already operating at its optimal configuration among the feasible autotune settings. The negative result is not a failure—it is a valuable piece of knowledge that prevents wasted effort, guides future optimization work toward more promising avenues (such as the already-deployed TARGET_CTAS=512 wave-fill optimization), and demonstrates the kind of disciplined, evidence-driven engineering that separates effective optimization from guesswork.

The message also reveals something about the nature of GPU kernel optimization: that the most intuitive levers (occupancy, pipeline depth) are not always the right ones, and that the specific characteristics of the computation—tiny MMA tiles, short loop iterations, sparse memory access patterns—can subvert general-purpose heuristics. The only way to know is to measure, and the only way to trust the measurement is to control the experiment carefully. This message does both, and the result is knowledge that will serve the project long after the specific benchmarks have faded from memory.