The Valuable Negative: How an Attention Occupancy A/B Test Refuted a High-Effort Hypothesis on Blackwell GPUs

Introduction

In the high-stakes world of LLM inference optimization, negative results are often more valuable than positive ones. They prevent wasted engineering effort, sharpen understanding of the hardware stack, and redirect focus toward levers that actually move the needle. Message [msg 13577] in this opencode session is a masterclass in the art of the negative result: a concise, evidence-rich summary of a systematic A/B test that definitively refuted a promising but ultimately misguided hypothesis about attention kernel optimization on NVIDIA Blackwell (sm120) GPUs.

The message, written by the AI assistant to the user, reports the completion of "task #3a" — an attention-occupancy A/B test that tested whether changing the number of warps and pipeline stages in a Triton-based MLA (Multi-head Latent Attention) decode kernel could improve throughput for the DeepSeek-V4-Flash model running on RTX PRO 6000 Blackwell GPUs. The answer, delivered with surgical precision, was a clear "no" — and that negative finding carried profound implications for the entire optimization roadmap.

The Message in Full

The assistant writes:

## #3a (attention occupancy A/B) — done, with a valuable negative result

>

Doc: wrote ./DSV4_ATTN_OCCUPANCY_AB.md (what/why/how, baseline, per-variant results, conclusion) — committed across 4 commits as results landed.

>

What I tested (vs the deployed TARGET_CTAS=512 baseline):

>

| variant | result | verdict | |---|---|---| | V1 num_warps=8 | C1 −14%, C8 −18%, C64 −3%, C96 −7% | REJECTED — 8 warps over-subdivide the tiny [16,16] MLA-decode MMA → MMA inefficiency beats any latency hiding | | V2 num_stages=3 | C64 −9.9%, C80/C96 −2.5%, low-C −5/−7% | REJECTED — per-split loops (4-6 iters) too short to fill a 3-stage pipeline; overhead dominates | | V3 batch-aware key | moot | skipped — w8 lost at every C |

>

Key learnings (evidence-based): - The production config BLOCK_T16/w4/s2 is already optimal among SMEM-feasible configs (BLOCK_T=32 is OOR-pruned >99 KB; block_h=16 since n_local_heads = 64/tp4 = 16, 60 KB, 1 CTA/SM). - This empirically refutes the full register-reducing rewrite. Its whole premise was "8 warps/SM hides the gather → +15-25%." V1 delivered 8 warps/SM and it regressed. So I've de-prioritized the rewrite — it won't help unless it fundamentally changes the gather/MMA structure (a KV-layout research effort, not tuning). That negative result just saved a high-effort dead-end. - Benchmark variance is ~±5-7% (C96 spans 775-845 across runs) — now documented so future deltas are read correctly. It confirms TARGET_CTAS=512's C64 +12.8% is real (beyond noise); the C96 +5.7% is partly noise but the monotonicity fix held.

>

Subagent research (as usual) resolved the autotune×cuda-graph capture safety (warmup-before-capture makes it safe), per-config SMEM feasibility, and the block_h=16 discrepancy before I touched anything.

>

State: reverted to the verified-good baseline (kernel md5 matches; TARGET_CTAS=512 + MULTI_STREAM_OVERLAP=0 + overlap-off all live, 0 errors). All committed; plan doc updated.

>

Net of the whole perf workstream: the realized decode win is TARGET_CTAS=512 (+12.8% C64, monotonic C60→C90), correctness preserved. #2 (overlap-schedule) and #3a (attention occupancy/rewrite) are now evidence-based dead-ends on this sm120/PCIe stack. The most promising untried lever is OEA (opportunistic expert activation — MoE fixed-floor, GLM-5 saw +5.7-25%), pending a check of DSV4's n_group grouped routing. Want me to scope OEA next?

Why This Message Was Written: The Reasoning and Motivation

To understand why this message exists, we must understand the context that produced it. The assistant and user had been engaged in an intensive, multi-session optimization campaign for the DeepSeek-V4-Flash model running on a cluster of 8 NVIDIA RTX PRO 6000 Blackwell GPUs. The model uses a novel architecture combining Multi-head Latent Attention (MLA) with Mixture-of-Experts (MoE), and the inference stack was built on SGLang with custom Triton kernels.

Earlier in the session, the assistant had already achieved a meaningful throughput win through the TARGET_CTAS=512 optimization — a tuning of the CUDA grid dimensions that improved C64 throughput by +12.8% and ensured monotonic scaling from C60 to C90 concurrency. But the team was hungry for more. The decode phase of LLM inference — the autoregressive token-by-token generation — is typically the throughput bottleneck, and on Blackwell GPUs with their novel sm120 architecture, the attention kernel was suspected to be occupancy-bound.

The hypothesis was elegant: the MLA decode kernel uses a gather operation (scattered memory reads) that is fundamentally latency-bound. If you could increase the number of warps per streaming multiprocessor (SM), you could hide that memory latency through warp-level parallelism — when one warp stalls on a memory read, another warp executes. This is a classic GPU optimization pattern, and it had worked on previous architectures. The specific proposal was a "register-reducing rewrite" that would free up register pressure to allow 8 warps per SM instead of 4, potentially yielding 15-25% throughput gains.

But this was a high-effort hypothesis. A full register-reducing rewrite of a Triton attention kernel is weeks of work, with significant risk of introducing correctness bugs. Before committing to that effort, the assistant needed to validate the premise: does increasing warp occupancy actually help on this specific kernel, on this specific architecture?

That validation is what message [msg 13577] reports. The answer is a definitive "no" — and the message is the formal closeout of that investigation.## The Decision-Making Process: How the A/B Test Was Designed and Executed

The message is a summary, but the thinking process that produced it is visible in the preceding messages ([msg 13563] through [msg 13576]). The assistant's approach was methodical and evidence-driven, reflecting a deep understanding of GPU architecture and a healthy skepticism toward theoretical promises.

Phase 1: Subagent Research. Before touching any code, the assistant dispatched subagents to investigate three critical questions: (1) whether Triton's autotune mechanism interacts safely with CUDA-graph capture (a key feature of SGLang's serving path), (2) what the per-config shared memory (SMEM) footprint would be for different block sizes and warp counts, and (3) why the kernel's block_h parameter was set to 16 rather than the expected 64. The subagent discovered that n_local_heads = 64 / tp4 = 16, resolving the discrepancy. It also confirmed that warmup-before-capture makes autotune safe under graph replay — a non-trivial finding that prevented a potential deployment crash.

Phase 2: V1 — More Warps. The first variant tested num_warps=8 against the baseline num_warps=4. The hypothesis was that 8 warps per SM would hide the gather latency through increased occupancy. The results were unambiguous: regression at every concurrency level, with low-concurrency workloads hit hardest (C1 −14%, C8 −18%). The assistant diagnosed the root cause: the MLA decode kernel uses a tiny [16,16] MMA (matrix multiply-accumulate) operation, and 8 warps over-subdivide this already-small work unit. The MMA inefficiency — the overhead of partitioning, synchronizing, and merging 8 warps' worth of tiny matrix tiles — overwhelmed any latency-hiding benefit. The verdict was clear: REJECTED.

Phase 3: V2 — Deeper Prefetch. The second variant kept num_warps=4 but increased num_stages from 2 to 3, deepening the prefetch pipeline. The idea was that more pipeline stages would better hide memory latency without the warp-subdivision problem. The results were more nuanced but ultimately negative: C64 showed a −9.9% regression, low concurrencies regressed by 5-7%, and only C48 was roughly flat. The assistant diagnosed that the per-split loops (only 4-6 iterations per block at high concurrency) were too short to fill a 3-stage pipeline — the pipeline fill/drain overhead dominated any benefit. REJECTED.

Phase 4: V3 Moot. A third variant (batch-aware key, selecting warps dynamically based on batch size) was planned but skipped. Since V1 showed that w8 lost at every concurrency level, there was no scenario where a batch-aware warp selector would choose 8 warps — the baseline 4-warp configuration was universally superior. Skipping V3 saved additional engineering time.

Phase 5: Variance Characterization. After reverting to baseline, the assistant noticed that C96 throughput varied between 775 and 845 tok/s across repeated runs — a ±5-7% noise band. This was a critical finding: it meant that small deltas (like the C96 +5.7% attributed to TARGET_CTAS=512) were partly within the noise floor, while large regressions (V1's −14% at C1, V2's −9.9% at C64) were genuine. Documenting this variance prevented future misinterpretation of benchmark results.

Assumptions Made and Their Validation

The message reveals several assumptions that the assistant held, some explicit and some implicit:

Assumption 1: The attention kernel is occupancy-bound. This was the foundational premise of the entire #3a workstream. The assistant assumed that the gather operation in the MLA decode kernel was latency-bound and that increasing warp occupancy would hide that latency. The A/B test disproved this assumption for this specific kernel shape — the tiny [16,16] MMA means that warp-level parallelism is not the bottleneck. The real limit, as the assistant later inferred, is likely gather bandwidth or the inherent work pattern.

Assumption 2: Autotune × CUDA-graph capture is safe. This was validated by subagent research before the experiment began. If it had been unsafe, the entire A/B approach would have been invalid — autotune could select different configurations during graph replay than during warmup, causing silent correctness bugs. The subagent confirmed that warmup-before-capture makes it safe, which was a necessary precondition for the entire experiment.

Assumption 3: The baseline configuration is stable and reproducible. The assistant assumed that reverting to the baseline kernel (verified by md5sum) and confirming the environment variables (TARGET_CTAS=512, MULTI_STREAM_OVERLAP=0, disable-overlap-schedule) would restore a known-good state. This was validated by probe requests and benchmark re-runs.

Assumption 4: Benchmark variance is small enough to detect meaningful deltas. This assumption was tested and refined during the experiment. The initial benchmarks used short runs (n×3, max_tokens=224) which showed ~flat results. The full benchmarks (n×4, max_tokens=256) revealed the true regressions. The variance characterization (±5-7%) then became a documented property of the measurement methodology.

Mistakes and Incorrect Assumptions

The message itself is a summary of correct conclusions, but the process it reports contains a subtle mistake that the assistant caught and corrected:

The mistaken "regression" after revert. When the assistant first reverted to the baseline kernel ([msg 13569]), the C96 throughput was 798.8 tok/s — lower than the 844.6 tok/s previously attributed to TARGET_CTAS=512. This initially looked like a regression, suggesting the revert had somehow lost the optimization. But the assistant didn't panic. Instead, it methodically verified that TARGET_CTAS=512 was still live in the process environment, ran two additional C96 benchmarks (which returned 829.2 and 775.0), and concluded that the "regression" was simply run-to-run variance within the ±5-7% noise band. The mistake was in treating a single benchmark run as definitive — a common pitfall in performance engineering that the assistant corrected through replication.

This is a valuable lesson in the message's subtext: benchmark results are statistical, not deterministic. The assistant's response to the apparent regression — verify the configuration, replicate the measurement, characterize the noise — is a model of disciplined performance analysis.

Input Knowledge Required

To fully understand this message, the reader needs knowledge spanning several domains:

GPU Architecture: Understanding what "warps," "SMEM," "MMA," "CTAs," and "shared memory" mean in the context of NVIDIA GPUs. The message assumes familiarity with the Blackwell sm120 architecture and its differences from previous generations (Hopper, Ada). The concept of "occupancy" — the ratio of active warps to the maximum supported per SM — is central to the analysis.

Triton and CUDA-Graph: The message references Triton's @triton.autotune decorator, which compiles multiple kernel configurations and selects the fastest at runtime. It also references CUDA-graph capture, a feature that records GPU operations into a replayable graph for reduced launch overhead. The interaction between these two systems (autotune during warmup, then graph capture of the selected config) is a subtle correctness concern.

LLM Inference Architecture: The message discusses MLA (Multi-head Latent Attention), a key innovation in the DeepSeek-V4 architecture that compresses the key-value cache through a latent bottleneck. The "gather" operation refers to scattering attention computation across the KV cache pages. The concept of "prefill-decode disaggregation" (splitting prefill and decode across different GPUs) is referenced through the disable-overlap-schedule flag.

Performance Measurement Methodology: The message assumes the reader understands concurrency (C), tokens per second (tok/s), and the significance of benchmark variance. The distinction between "real" gains (beyond the noise band) and "noise" is a methodological insight.

Output Knowledge Created

This message creates several forms of output knowledge:

Documented Negative Result: The A/B test is recorded in DSV4_ATTN_OCCUPANCY_AB.md with full methodology, results, and conclusions. This prevents future engineers from re-treading the same dead-end path.

Updated Optimization Roadmap: The plan document DSV4_DECODE_PERF_PLAN.md is updated to de-prioritize the register-reducing rewrite and flag OEA (Opportunistic Expert Activation) as the next most promising lever. The message explicitly asks the user: "Want me to scope OEA next?" — turning the negative result into a positive direction.

Benchmark Variance Characterization: The ±5-7% noise band is now a documented property of the measurement setup. Future experiments can use this to determine whether a delta is statistically significant.

Empirical Refutation of a Theoretical Model: The message establishes that "more warps = better latency hiding" does not hold for the MLA decode kernel on sm120. This is a non-obvious finding — on many GPU kernels, increasing occupancy does help. The specific reason (tiny [16,16] MMA over-subdivision) is a valuable architectural insight.

Deployment State Verification: The message confirms that the production system is in a known-good state with all optimizations (TARGET_CTAS=512) live and all experimental changes reverted. The md5 checksum of the kernel file, the live process environment, and the command-line flags are all verified.

The Thinking Process: A Window into Rigorous Engineering

The assistant's reasoning, visible in the preceding messages, reveals a disciplined thought process. When V2 showed mixed results, the assistant didn't cherry-pick the positive numbers (C1 +6.6% in the initial quick bench) — it ran the full benchmark suite and identified the consistent pattern of regression. When the revert appeared to lose performance, it didn't assume the worst — it verified the configuration, replicated the measurement, and discovered the variance band. When concluding that the register-reducing rewrite is refuted, it didn't just say "it doesn't work" — it explained why the premise was wrong (the tiny MMA shape means 8 warps over-subdivide the work) and what would need to change for the rewrite to be viable (a fundamental KV-layout research effort).

The message also demonstrates intellectual honesty. The assistant could have framed the result as "we learned valuable things about our kernel" without explicitly stating that a high-effort rewrite was now de-prioritized. Instead, it was direct: "This empirically refutes the full register-reducing rewrite. Its whole premise was '8 warps/SM hides the gather → +15-25%.' V1 delivered 8 warps/SM and it regressed." That clarity is what makes the negative result valuable — it kills a bad idea before it consumes weeks of engineering time.

Conclusion: The Value of a Well-Executed Negative Result

Message [msg 13577] is a model of how to report a negative result in production ML engineering. It is concise, evidence-based, and actionable. It doesn't just say "this didn't work" — it explains why it didn't work, what was learned, and what should be tried next. The benchmark variance characterization is a gift to future optimization efforts. The explicit de-prioritization of the register-rewrite hypothesis saves weeks of work. The pivot to OEA provides a clear next step.

In the broader context of the opencode session, this message represents a turning point. The assistant had been chasing the attention kernel as the primary optimization target. The TARGET_CTAS=512 win had shown that meaningful gains were possible, but the occupancy levers (#3a) and the overlap scheduler (#2) had both proven to be dead-ends on this sm120/PCIe stack. The message refocuses the optimization roadmap on the MoE layer — OEA — which the assistant notes showed +5.7-25% on the related GLM-5 model. This is the scientific method in action: hypothesis, experiment, evidence, conclusion, and redirection.

The most important sentence in the message is perhaps the most understated: "That negative result just saved a high-effort dead-end." In engineering, the ability to quickly disprove a bad hypothesis is as valuable as the ability to prove a good one. This message is a testament to that principle, and a demonstration of how systematic A/B testing, deep architectural understanding, and disciplined measurement methodology can keep an optimization campaign on track — even when the answer is "no."