The Verify Kernel Goes Live: A Milestone in Custom sm_120 Attention for Blackwell GPUs

In the high-stakes world of large language model inference, every microsecond of attention computation matters. When existing high-performance libraries like FlashMLA, cutlass-MLA, and flashinfer-MLA all lack support for a particular GPU architecture, the only path forward is to build from scratch. Message [msg 12290] captures a pivotal moment in such an effort: the moment a custom CUDA verify attention kernel for the sm_120 architecture (RTX PRO 6000 Blackwell consumer GPUs) goes live in a production SGLang service, serving real Kimi K2.6 speculative decoding traffic. This message is not merely a status update—it is a carefully reasoned transition from correctness validation to performance optimization, marking the completion of one phase and the clear-eyed identification of the next.

The Moment of Truth: From Validation to Production

The message opens with quiet confidence: "Great, the ON mode is working as expected." This understated sentence belies the enormous effort that preceded it. For multiple prior messages, the assistant had been wrestling with a kernel that produced outputs completely uncorrelated with Triton's—relative errors around 1.0, meaning the attention computation was essentially random. The root cause had been a subtle marshaling mismatch: SGLang's extend attention structure places prefix keys in the KV pool (indexed via kv_indices) while draft tokens are passed as separate k/v tensors and written to out_cache_loc. The assistant's kernel had incorrectly assumed all keys lived in the pool, leading to the wrong keys being attended to. Fixing this required building combined indices that concatenated prefix slot locations with draft token cache locations, and correctly extracting the tree visibility mask using the full sequence length stride rather than just the prefix length.

Now, after that fix, the outputs match Triton's "almost perfectly across different prompts." The assistant provides concrete evidence: the "capital of France" prompt produces identical completions, the "list three primes" prompt is identical, and the "ocean" sentence diverges only in a minor late-token word choice ("life forms" versus "life"). This divergence is correctly attributed to bf16 floating-point non-determinism at temperature zero—a well-documented phenomenon where different kernel implementations of the same mathematical operation produce slightly different rounding results, which can cascade into different token selections in later positions. The assistant explicitly notes this is "not a bug," demonstrating a mature understanding of numerical precision in ML inference.

The Reasoning: A Clear-Eyed Assessment of What Comes Next

The most intellectually substantial portion of the message is the assistant's reasoning about the next steps. Having achieved correctness, the assistant immediately identifies the critical blocker: "the service runs in eager mode because the patch breaks CUDA graph capture." CUDA graphs are a performance optimization that allows the GPU to replay a sequence of kernel launches without CPU involvement, eliminating launch overhead and enabling the GPU to run at full utilization. Without graph capture, the service must pay the cost of Python-level marshaling—building indices, extracting mask strides, and launching kernels—on every decode step.

The assistant's reasoning here is exemplary in its structure. It first identifies the problem (no CUDA graphs), then enumerates the specific capture-safety issues: "dynamic allocations in the marshaling loop, the .item() calls for host syncs, and the data-dependent shapes." Each of these is a known pattern that prevents CUDA graph capture. Dynamic allocations mean the graph's memory footprint isn't fixed. .item() calls require a CPU-GPU synchronization that breaks the purely GPU-execution model of a CUDA graph. Data-dependent shapes mean the kernel launch parameters themselves change between invocations.

Rather than despairing, the assistant formulates a clean experimental strategy: "To isolate the kernel improvement fairly, I need to compare eager-mine versus eager-triton (both without cuda graphs), which removes the eager-mode overhead as a confound." This is a textbook example of controlling for confounding variables in performance measurement. By comparing both kernels in eager mode, any difference can be attributed to the kernel itself rather than to the graph-capture infrastructure. The assistant also correctly notes that at long context lengths, "the attention should dominate and show the kernel benefit," because attention computation scales quadratically with sequence length while marshaling overhead is roughly constant per step.

The Benchmark: First Performance Data

The message concludes with the first performance measurement of the custom kernel in its integrated form. The benchmark runs at three context lengths—4,096 tokens (5,702 prompt tokens), 16,384 tokens (22,779 prompt), and 65,536 tokens (91,027 prompt)—measuring decode throughput in tokens per second:

| Context | Prompt Tokens | Decode tok/s | Decode Wall (s) | |---------|---------------|--------------|-----------------| | 4,096 | 5,702 | 35.2 | 0.43 | | 16,384 | 22,779 | 13.6 | 1.10 | | 65,536 | 91,027 | 3.6 | 4.11 |

These numbers tell a story of context-dependent performance degradation. At 4K context, the kernel achieves 35.2 tokens per second. At 16K, throughput drops to 13.6 tok/s (a 2.6× slowdown for 4× more context). At 65K, it falls to 3.6 tok/s (a 9.8× slowdown for 16× more context). This sub-linear scaling is expected for attention mechanisms, but the magnitude of the drop hints at underlying inefficiencies that the assistant will need to address.

What makes these numbers meaningful is the comparison that follows in the next message ([msg 12291]). The Triton eager baseline at the same context lengths yields 38.0, 6.7, and 1.8 tok/s respectively. The custom kernel is slightly slower at 4K (35.2 vs 38.0, about 7% slower) but dramatically faster at longer contexts: 2× faster at 16K (13.6 vs 6.7) and 2× faster at 65K (3.6 vs 1.8). This pattern is exactly what the assistant hypothesized—the kernel's advantage grows with context length because attention computation dominates at longer sequences, and the custom kernel's flash-decode design with KV splitting and vectorized loads extracts more bandwidth from the sm_120 architecture than Triton's generic MLA implementation.

Assumptions, Decisions, and Knowledge

The message reveals several key assumptions and decisions. The assistant assumes that comparing both kernels in eager mode is a fair test, which is correct for isolating kernel performance but leaves open the question of whether Triton's graph-mode performance would be even better. The decision to benchmark at three context lengths rather than a single point is deliberate—it allows the assistant to observe the scaling trend and confirm that the kernel's benefit grows with context length.

The input knowledge required to understand this message is substantial. One must understand SGLang's extend attention architecture (the split between prefix KV pool and extend KV tensors), CUDA graph capture semantics (what makes a kernel capture-safe), the structure of MLA (Multi-head Latent Attention) with its absorbed query representation, and the sm_120 instruction set architecture (which lacks the Hopper/Blackwell-DC features like wgmma and TMA that other optimized libraries rely on). The assistant draws on all of this knowledge to interpret the benchmark results and plan the next steps.

The output knowledge created by this message is equally significant. It establishes that a custom sm_120 verify attention kernel can achieve numerical parity with Triton on real K2.6 serving tensors with TP8 and paged KV. It provides the first performance data point for such a kernel in an integrated SGLang service. And it identifies CUDA graph capture as the next critical optimization target, with a clear list of what needs to change (no dynamic allocations, no .item() calls, fixed shapes).

The Thinking Process: Instrumented Decision-Making

What makes this message particularly valuable as a case study is the transparency of the assistant's thinking process. The reasoning section reads like a well-structured engineering log: state the current status, identify the next blocker, enumerate the options, formulate a measurement strategy, execute the measurement, and document the results for comparison. This disciplined approach to instrumentation—replacing assumptions with evidence before making decisions—is a hallmark of effective systems engineering.

The assistant explicitly considers and rejects the temptation to immediately fix CUDA graph capture, instead choosing to first measure eager-mode performance to establish a baseline. This prevents wasted effort: if the kernel weren't faster even in eager mode, there would be no point in making it capture-safe. The benchmark results validate the approach, showing a clear 2× advantage at longer contexts that justifies the investment in graph capture.

Conclusion: A Foundation for Optimization

Message [msg 12290] represents the successful completion of a challenging engineering task—building a custom CUDA attention kernel for an unsupported GPU architecture and integrating it into a complex serving stack—while simultaneously laying the groundwork for the next phase of optimization. The custom sm_120 verify kernel is live, correct, and measurably faster than the Triton baseline at the context lengths that matter most for speculative decoding with large language models. The path forward is clear: make the kernel CUDA-graph capture-safe to unlock the full performance potential of the Blackwell GPUs, and continue optimizing the kernel's tile sizing and reduction strategy for the sm_120 architecture's unique characteristics. This message is a testament to the power of systematic reasoning, careful measurement, and the willingness to build from first principles when existing tools fall short.