The Kernel Archaeology of a Hard Ceiling: Planning the sm_120 Sparse-Decode Rewrite for Blackwell

Introduction

In the high-stakes world of deploying large language models on novel hardware, performance optimization is rarely a matter of flipping a single configuration flag. More often, it resembles archaeological excavation: you dig through layers of abstraction—CUDA graphs, NCCL tuning, expert parallelism, fused kernels—only to hit a bedrock of architectural reality. Message 12448 in this opencode session represents precisely that moment of archaeological reckoning. Having systematically eliminated every configurable lever and proven that the bottleneck was a single Triton fallback kernel consuming 63% of GPU time, the assistant now pivots from diagnosis to treatment. The user's directive—"Expand the findings with a detailed plan to build the missing kernels tuned for RTX PRO 6000 Blackwell"—triggers a message that is less a plan itself than the genesis of a plan: a reasoning-rich dispatch that launches two parallel subagent research tasks to map the kernel internals, assesses the reuse potential of prior work, and establishes the strategic framework for a multi-kernel rewrite campaign.

This article examines that message in depth: its reasoning, its decision-making, its assumptions, and the knowledge it both consumes and produces. It is a case study in how an experienced AI systems engineer approaches the problem of hardware-bound performance optimization when the easy knobs have all been turned.

The Context: A Bottleneck Laid Bare

To understand message 12448, one must first understand what preceded it. The session had been engaged in a prolonged effort to deploy DeepSeek-V4-Flash—a Mixture-of-Experts (MoE) model with sparse Multi-head Latent Attention (MLA)—on a machine equipped with 8× NVIDIA RTX PRO 6000 Blackwell GPUs (compute capability sm_120). The deployment had achieved prefill-decode disaggregation, a sophisticated orchestration where prefill runs on GPU 0–3 (NUMA 0) and decode runs on GPU 4–7 (NUMA 1), with KV cache transfer via NIXL/UCX. The infrastructure worked correctly.

The performance, however, was catastrophic relative to expectations. At batch size 1, throughput hovered around 10 tokens per second. At concurrency 16, it reached approximately 23–25 tok/s. The user's target was roughly 1000 tok/s—a gap of over 40×.

The assistant had methodically tried every configuration lever in its repertoire: NCCL LL+Ring protocol tuning, CUDA graphs (already enabled), tilelang indexer fusion (which failed to JIT-compile on sm_120 with CUDA 13), non-Marlin MoE backends (invalid for FP4 expert weights), and expert parallelism EP4 (which made performance worse due to PCIe all-to-all overhead). None moved the needle.

The definitive breakthrough came from GPU profiling ([msg 12442]). A trace of 40 decode steps revealed the total GPU kernel time was 3,830 ms, and a single kernel—_tiled_sparse_decode_kernel—accounted for 2,412 ms (63%). The MoE GEMV kernel consumed 10%, FP8 attention projections 6%, and NCCL all-reduce a mere 2%. The assistant had found its culprit.

Further investigation ([msg 12444]) confirmed the root cause at the grid level. The kernel's launch configuration was grid = (B, H)—at batch size 1, that meant only 64 blocks (1 batch × 64 query heads) on a GPU with approximately 170 streaming multiprocessors. Each block serially iterated through all 512 top-k tokens in tiles of 16–32. This was the same low-occupancy pathology that had plagued the earlier K2.6 verify kernel, which the team had fixed with a split-K strategy yielding 3–6× speedup.

The assistant's analysis concluded with an honest, quantified ceiling: even if the sparse-decode kernel were made instant, the remaining 37% of GPU time (MXFP4 MoE gemv, FP8 matmul, elementwise ops) would limit total speedup to ~2.7×, yielding roughly 60 tok/s at C=16. The 1000 tok/s target was unreachable on sm_120 via configuration alone—it required the fused DSA/MoE stack (DeepGEMM, trtllm-gen, FP4 indexer) that is arch-gated to SM100 (B200) hardware.

The Message: Planning Mode Engaged

Message 12448 is the assistant's response to the user's request for a detailed kernel-build plan. The message begins with an extended reasoning block that reveals the assistant's strategic thinking:

Agent Reasoning: I'm in planning mode to build a comprehensive strategy for optimizing the sparse-decode kernel for RTX PRO 6000 Blackwell architecture. Since this is the primary bottleneck at 63%, I need to thoroughly understand its algorithm, data flow, memory patterns, and how to best leverage the new hardware capabilities.

I'm also identifying the secondary optimization targets: the MXFP4 MoE gemv kernel and the FP8 attention matmul, where there's a quick win by tuning the missing config for this GPU. The key insight is that the repo already has a proven split-K flash-decode implementation with online softmax merging that I can use as a reference architecture, which significantly de-risks the plan. Realistically that gets us to 2.5-3x speedup and around 60-70 tokens/sec at C=16, but going further requires tackling the long tail of FP8 projections and fused elementwise ops. The core missing kernels to prioritize are the sparse-MLA decode attention with split-K (biggest win), an optimized FP4 tensor-core MoE for sm120a, tuned FP8 block-scaled GEMM, a fused kernel for the indexer logits, and the mHC. For implementation, I'm weighing CUDA with split-K reusing the kdtree verify_attn_flash kernel against improving the existing Triton kernels, and I can patch directly into the editable sglang source or integrate via a compiled .so.

I'm setting up targeted exploration to map out the sparse-MLA decode path in detail—the kernel inputs, grid layout, math, and how deepseek_v4_backend calls it with the actual tensor shapes and dtypes—then do the same for the MXFP4 MoE path including weight layout and the missing tuning config mechanism. I'll also pull the K2.6 kdtree CUDA kernels to assess what can be reused and grab the exact sm_120 hardware specs from the box.

The message then launches two task tool calls—subagent sessions that run in parallel to perform read-only analysis of the source code:

"Map sparse-MLA decode kernel": A task to read the Triton kernel implementation (flash_mla_sm120_triton.py), the Python wrapper (flash_mla_sm120.py), and the backend integration (deepseek_v4_backend.py), producing a precise technical specification for a split-K rewrite.
"Map MXFP4 MoE + FP8 GEMM kernels": A task to analyze the MXFP4 MoE decode kernel and FP8 block GEMM, identifying the missing tuning configuration mechanism and weight layout details. Both tasks are read-only; the assistant explicitly states "Do NOT modify anything; this is read-only analysis." The results of these tasks (truncated in the conversation data) return detailed technical specs that will form the foundation of the actual kernel implementation.

Why This Message Was Written: The Reasoning and Motivation

The message serves multiple strategic purposes, each reflecting a layer of the assistant's reasoning.

First, it establishes a research-first methodology. The assistant has already diagnosed the bottleneck at the grid-occupancy level, but it recognizes that writing a correct, performant split-K kernel requires precise knowledge of the data layout, tensor shapes, memory access patterns, and calling conventions. The existing Triton kernel is a black box at the algorithmic level—the assistant knows what it does (sparse decode attention with online softmax over top-k tokens) but needs to understand exactly how the inputs are structured, how the KV cache is indexed, and how the output feeds into the subsequent MoE layers. Launching subagent tasks to read and analyze the source code is a deliberate choice to ground the plan in concrete implementation details rather than abstract reasoning.

Second, it manages scope and expectations. The assistant explicitly acknowledges that the split-K rewrite of the sparse-decode kernel is the "biggest win" but also identifies four other missing kernels: an optimized FP4 tensor-core MoE for sm_120a, tuned FP8 block-scaled GEMM, a fused indexer logits kernel, and an mHC kernel. By enumerating the full set of required kernels, the assistant provides the user with a realistic roadmap. The phrase "Realistically that gets us to 2.5-3x speedup and around 60-70 tokens/sec at C=16" is a carefully calibrated expectation-setter—it says "here's what we can achieve with a focused effort" while implicitly acknowledging that the 1000 tok/s target remains out of reach without the SM100 fused stack.

Third, it evaluates implementation strategies. The assistant weighs two approaches: writing custom CUDA kernels (reusing the K2.6 kdtree verify_attn_flash kernel as a reference) versus improving the existing Triton kernels. It also considers integration strategies—patching directly into the editable SGLang source versus compiling a standalone .so that can be loaded dynamically. This is not a trivial decision: CUDA offers finer control over occupancy and memory coalescing but requires more development time, while Triton offers faster iteration but may be constrained by its autotuning framework. The assistant's reasoning shows an awareness of the trade-offs without prematurely committing to one path.

How Decisions Were Made

Several implicit and explicit decisions shape the message's direction.

Decision 1: Prioritize the sparse-decode kernel. This is the most obvious and defensible decision. At 63% of GPU time, the _tiled_sparse_decode_kernel is the single largest contributor to latency. The assistant's own analysis shows that even making this kernel instant yields only 2.7× overall speedup, but that is still the highest-leverage intervention available. The decision follows the Pareto principle: fix the one thing that dominates, then assess the remaining tail.

Decision 2: Use split-K as the primary optimization strategy. This decision is grounded in precedent. The assistant explicitly references the K2.6 kdtree verify kernel, which had the same low-occupancy pathology (grid = (B, H) with serial inner loop over the KV dimension) and was fixed by introducing a split-K dimension that distributed the workload across more blocks. The repo already contains a proven split-K flash-decode implementation with online softmax merging. Reusing this architectural pattern de-risks the implementation significantly—the assistant is not inventing a new algorithm but adapting an existing one to a different kernel.

Decision 3: Research before building. The assistant could have jumped directly into writing CUDA code, given that the bottleneck is well-understood. Instead, it chooses to invest in detailed source-code analysis via subagent tasks. This reflects an understanding that kernel implementation is unforgiving of guesswork: incorrect assumptions about tensor strides, memory layouts, or calling conventions can produce kernels that silently compute wrong results or fail to compile. The research phase is an insurance policy against wasted implementation effort.

Decision 4: Address the MoE and FP8 GEMM kernels as secondary targets. The assistant identifies the MXFP4 MoE gemv kernel (10% of GPU time) and FP8 attention matmul (6%) as secondary optimization targets. The reasoning mentions "a quick win by tuning the missing config for this GPU," suggesting that the assistant suspects these kernels have a configuration mechanism (similar to autotuning) that simply lacks entries for sm_120. This is a lower-effort, lower-reward intervention that can be pursued in parallel with the sparse-decode rewrite.

Assumptions Made

The message rests on several assumptions, some explicit and some implicit.

Assumption 1: The split-K approach from K2.6 is transferable. The assistant assumes that the architectural pattern used for the K2.6 verify kernel—splitting the KV/top-k dimension across multiple blocks and combining results with online softmax—will work for the DeepSeek-V4-Flash sparse-decode kernel. This is a reasonable assumption given that both kernels perform attention over a selected subset of KV pairs, but it is not guaranteed. The sparse-decode kernel may have additional constraints (e.g., the top-k selection is dynamic per layer, or the KV cache layout differs) that complicate the split-K adaptation. The subagent research task is designed to validate this assumption.

Assumption 2: The Triton kernel's grid structure is the sole cause of low occupancy. The assistant attributes the kernel's poor performance entirely to the (B, H) grid with serial inner loop. While the profiling data strongly supports this—64 blocks on ~170 SMs is undeniably low occupancy—there may be additional factors: inefficient memory access patterns, register pressure, or suboptimal tile sizes that the autotuning didn't explore. The split-K fix addresses occupancy but may not address all performance pathologies.

Assumption 3: The user is willing to invest in a multi-kernel rewrite. The assistant's plan implicitly assumes that the user has the time and tolerance for a development effort spanning multiple kernels. The message mentions "multi-week scope" for the full owned-sm_120-kernel route. The assistant is careful to offer bounded expectations ("2.5-3x speedup") but does not explicitly ask the user to confirm the investment level before proceeding with research. This is a minor risk: if the user expected a quick config fix, the kernel-rewrite plan may exceed their patience.

Assumption 4: The K2.6 kdtree CUDA kernels are available and reusable. The assistant plans to "pull the K2.6 kdtree CUDA kernels to assess what can be reused." This assumes that those kernels exist in the repository, are compatible with the sm_120 architecture, and have a license or ownership structure that permits reuse. If the K2.6 kernels were written for a different GPU architecture (e.g., sm_90 for Hopper) or are not available in the current workspace, this assumption would need revision.

Input Knowledge Required

To fully understand message 12448, one needs knowledge spanning several domains:

Deep learning inference architecture: Understanding what sparse-MLA attention is, how it differs from standard attention, the role of top-k token selection, and the decode-phase memory access patterns. The reader must know that during autoregressive decoding, each token attends to a subset of previous tokens (selected by a separate mechanism), and that this sparse attention reduces compute but introduces irregular memory access.

GPU compute architecture: Knowledge of streaming multiprocessors (SMs), occupancy, grid/block launch configurations, and the difference between CUDA cores and tensor cores. The sm_120 architecture (Blackwell) has specific capabilities (FP4 tensor cores, FP8 GEMM) that differ from sm_90 (Hopper) and sm_100 (B200). Understanding why the fused DSA stack is "arch-gated to SM100" requires knowing that certain hardware instructions (e.g., tcgen05 for tensor-core generation) are only available on newer architectures.

The SGLang inference engine: Familiarity with SGLang's architecture—how it dispatches attention kernels, how the backend integration works (deepseek_v4_backend.py), and how custom kernels can be injected into the serving path. The assistant references patching "directly into the editable sglang source" or integrating "via a compiled .so," which implies knowledge of SGLang's plugin mechanism.

Prior session history: The reader must know about the K2.6 verify kernel effort, which is referenced repeatedly as a precedent. The K2.6 work involved building a custom CUDA kernel with split-K that achieved 3–6× speedup over the Triton baseline. The assistant treats this as a validated template for the current problem.

CUDA and Triton programming models: The distinction between writing kernels in CUDA C++ (with explicit grid/block management) versus Triton (a Python-based DSL that auto-generates CUDA code) is central to the implementation strategy discussion. The assistant's weighing of these approaches assumes familiarity with their respective strengths and limitations.

Output Knowledge Created

Message 12448 produces several forms of knowledge:

A prioritized kernel development roadmap: The assistant enumerates five missing kernels in priority order: (1) sparse-MLA decode attention with split-K, (2) FP4 tensor-core MoE for sm_120a, (3) tuned FP8 block-scaled GEMM, (4) fused indexer logits kernel, (5) mHC kernel. This roadmap provides a structured path forward, preventing the effort from becoming unfocused.

An architectural reference point: By identifying the K2.6 kdtree verify kernel as a reusable reference, the assistant creates a bridge between prior work and the current problem. This reduces the cognitive load of designing a split-K kernel from scratch—the team can study the existing implementation, understand its patterns (how it splits the KV dimension, how it combines partial results with online softmax), and adapt them to the sparse-decode context.

A realistic performance ceiling: The assistant's estimate of "2.5-3x speedup and around 60-70 tokens/sec at C=16" serves as a crucial expectation-management tool. It prevents the user from expecting the split-K rewrite to close the entire 40× gap to 1000 tok/s. The message explicitly acknowledges that the remaining 37% of GPU time (MXFP4 MoE gemv, FP8 matmul, elementwise ops) will limit total speedup, and that reaching the user's target would require "a multi-kernel, multi-week effort" including owned sm_120 kernels for the entire pipeline.

Technical specifications (via subagent tasks): The two task tool calls produce detailed technical specs that document the exact kernel inputs, tensor shapes, dtypes, memory layouts, and calling conventions. These specs are the raw material from which the actual kernel implementation will be built. They represent a form of knowledge extraction—converting implicit information scattered across multiple source files into a structured, actionable document.

An evaluation of integration strategies: The assistant's consideration of patching SGLang source versus compiling a standalone .so creates knowledge about the engineering trade-offs. Patching is simpler but creates maintenance burden when SGLang is updated; a .so is more modular but requires a stable API boundary. This analysis helps the user make an informed decision about the integration approach.

The Thinking Process: A Window into Systems Engineering

The reasoning block in message 12448 is unusually rich, offering a rare window into how an experienced AI systems engineer approaches a complex optimization problem. Let me examine its structure.

The reasoning begins with a situation assessment: "I'm in planning mode to build a comprehensive strategy for optimizing the sparse-decode kernel for RTX PRO 6000 Blackwell architecture." This frames the message not as a final answer but as a preparatory step. The assistant is signaling that it needs to gather information before producing the deliverable.

Next comes a bottleneck reaffirmation: "Since this is the primary bottleneck at 63%, I need to thoroughly understand its algorithm, data flow, memory patterns, and how to best leverage the new hardware capabilities." The assistant anchors its reasoning in the empirical data from the profile, ensuring that the plan remains grounded in measured reality rather than speculation.

Then a secondary target identification: "I'm also identifying the secondary optimization targets: the MXFP4 MoE gemv kernel and the FP8 attention matmul, where there's a quick win by tuning the missing config for this GPU." This shows the assistant thinking in terms of a portfolio of interventions—not just the big win but also the low-effort, moderate-reward opportunities.

The reasoning then performs a risk assessment: "The key insight is that the repo already has a proven split-K flash-decode implementation with online softmax merging that I can use as a reference architecture, which significantly de-risks the plan." This is a critical strategic insight: the assistant recognizes that prior work in the same repository provides a template, reducing the uncertainty that typically accompanies novel kernel development.

Next, a realistic projection: "Realistically that gets us to 2.5-3x speedup and around 60-70 tokens/sec at C=16, but going further requires tackling the long tail of FP8 projections and fused elementwise ops." The assistant is careful to bound expectations, preventing the user from assuming that the split-K rewrite alone will solve the performance problem.

The reasoning then enumerates the full kernel portfolio: "The core missing kernels to prioritize are the sparse-MLA decode attention with split-K (biggest win), an optimized FP4 tensor-core MoE for sm120a, tuned FP8 block-scaled GEMM, a fused kernel for the indexer logits, and the mHC." This list serves as a scope definition—it tells the user exactly what needs to be built.

Finally, an implementation strategy evaluation: "For implementation, I'm weighing CUDA with split-K reusing the kdtree verify_attn_flash kernel against improving the existing Triton kernels, and I can patch directly into the editable sglang source or integrate via a compiled .so." The assistant is thinking ahead to the engineering logistics, considering both the programming model and the integration mechanism.

The reasoning concludes with a research plan: "I'm setting up targeted exploration to map out the sparse-MLA decode path in detail... then do the same for the MXFP4 MoE path... I'll also pull the K2.6 kdtree CUDA kernels to assess what can be reused and grab the exact sm_120 hardware specs from the box." This transitions from strategic thinking to tactical execution: the assistant launches the subagent tasks that will produce the technical specifications needed for implementation.

Mistakes and Incorrect Assumptions

While the message is generally sound, there are a few points worth examining critically.

Potential over-reliance on the K2.6 precedent. The assistant repeatedly references the K2.6 verify kernel as a template, but the DeepSeek-V4-Flash sparse-decode kernel may differ in significant ways. The K2.6 kernel operated on a tree-structured speculative decoding draft (the "kdtree"), where the KV set was the set of candidate tokens in the tree. The DSV4 sparse-decode kernel operates on a flat top-512 selection from a potentially much larger KV cache. The memory access patterns, indexing logic, and online softmax combine may differ enough that the K2.6 approach requires substantial adaptation. The assistant's subagent research task should surface these differences, but the reasoning does not explicitly acknowledge the risk that the template may not transfer cleanly.

The "quick win" for MXFP4 MoE may be optimistic. The assistant suggests that tuning the missing config for the MXFP4 MoE kernel on sm_120 could be a quick win. However, the profiling data shows this kernel consumes only 10% of GPU time. Even a 2× improvement would yield only ~5% overall speedup. The assistant may be overestimating the impact of this intervention relative to the effort required to identify and add the missing configuration entries.

No explicit fallback plan. The message does not discuss what happens if the split-K approach fails to achieve the expected speedup—for example, if memory bandwidth becomes the limiting factor rather than occupancy, or if the online softmax combine introduces numerical issues. A robust plan would include diagnostic criteria for success and contingency strategies if the initial approach underperforms.

Conclusion

Message 12448 represents a pivotal transition in the optimization campaign for DeepSeek-V4-Flash on Blackwell GPUs. Having exhausted every configuration lever and identified the root cause of the performance bottleneck with surgical precision, the assistant now pivots from diagnosis to treatment. The message is a masterclass in structured problem-solving: it reaffirms the bottleneck, identifies secondary targets, assesses risk through precedent, sets realistic expectations, enumerates the full scope of required work, evaluates implementation strategies, and launches the research phase that will produce the technical specifications for kernel development.

The message's true value lies not in the code it writes—it writes none—but in the thinking it surfaces. The reasoning block reveals a systems engineer operating at full capacity: weighing trade-offs, managing uncertainty, bounding expectations, and sequencing work for maximum impact. The two subagent tasks it spawns will produce the detailed technical knowledge needed to transform the strategic plan into working CUDA or Triton code.

For anyone interested in the art of GPU kernel optimization, this message is a case study in how to approach a hard ceiling: measure precisely, diagnose definitively, plan systematically, and research thoroughly before writing a single line of CUDA. The ceiling may not be broken in this message, but the foundation for breaking it is laid with care and precision.