The Milestone Commit: How a Git Checkpoint Captured the Birth of a Custom Blackwell GPU Kernel

Introduction

In the middle of a high-stakes optimization session for the GLM-5-NVFP4 model on NVIDIA RTX PRO 6000 Blackwell GPUs, a single git commit marks a quiet but significant transition. Message [msg 12249] is not a dramatic code change or a breakthrough benchmark—it is a checkpoint. The assistant runs git add -A && git commit -q -m "Phase 1: owned sm_120 flash-decode verify kernel..." and then displays the resulting commit hash. On the surface, this is a routine operation: stage all changes, write a descriptive message, and record the state. But in the context of the surrounding work, this commit represents the formal closure of a demanding engineering phase and the deliberate documentation of what was accomplished, what was learned, and what remains.

This article examines that single message—why it was written, what knowledge it encapsulates, the assumptions it makes, and the thinking process it reveals. To understand this commit, one must understand the journey that led to it: a multi-day effort to build a custom CUDA kernel for speculative decoding on a novel GPU architecture, where every microsecond of latency mattered and where off-the-shelf solutions had failed.

Context: The Problem That Demanded a Custom Kernel

The broader session (segment 66 of the conversation) was focused on deploying the GLM-5-NVFP4 model with 200k-token context length on CT200, a server equipped with RTX PRO 6000 Blackwell GPUs (compute capability sm_120). The assistant had already diagnosed a severe decode bottleneck: the existing Triton-based verify attention kernel was achieving only ~14 GB/s effective memory bandwidth—roughly 130× below the GPU's 1.8 TB/s peak. The root cause was that Triton's MLA (Multi-head Latent Attention) implementation used page_size=1 for its KV cache, causing scattered, non-contiguous memory accesses that crippled throughput.

The user directed the assistant to build an owned sm_120 kernel and implement K/V defragmentation. The assistant responded with a detailed plan (plans/0002-sm120-verify-kernel-defrag.md) and proceeded to implement Phase 1: a custom flash-decode MLA verify kernel. This was not a trivial port of existing code. The assistant discovered that all optimized MLA kernels (FlashMLA, cutlass-MLA, flashinfer-MLA) were compiled only for sm_90a, sm_100a, and sm_103a—none supported sm_120. The RTX PRO 6000 Blackwell consumer GPU uses an Ada-like ISA without the wgmma, TMA, or tcgen05 instructions found on Hopper and Blackwell-DC architectures. The assistant had to write a kernel from scratch for this specific target.

What the Commit Captures

The commit message in [msg 12249] is structured as a mini-manifesto of Phase 1 deliverables. Let us examine each element:

The kernel itself: src/kernels/verify_attn_flash.cu is the centerpiece. The commit message claims it is "token-exact vs naive oracle" (meaning it produces bitwise-identical attention outputs to a straightforward implementation) and achieves "1.7-4.2x faster on full-H bench." The "full-H" qualifier is important—it means the benchmark uses all 64 heads without tensor parallelism. In production with TP8, each rank would handle only 8 heads, making the per-head design even more efficient. The kernel also "handles arbitrary prefix the naive kernel can't fit in smem"—a critical capability for the 200k-token context length target.

The interface and test infrastructure: "C-ABI kdtree_verify_attn_flash; test --flash A/B; bench naive-vs-flash to 65k." The assistant built a C-compatible ABI for the kernel, enabling integration into the existing kdtree-engine codebase. The --flash A/B testing mode allows direct comparison between the naive oracle and the flash kernel, and the benchmark suite covers sequence lengths up to 65,536 tokens.

The planning and documentation: "plans/0002 spec; bench_results_ct200/verify_attn_flash.txt; context_scaling root-cause." The plan document that guided the work, the benchmark results file, and a root-cause analysis of the context-scaling problem are all included in the commit. This is not just code—it is the intellectual trail of the investigation.

Diagnostic tools: "commit_probe.py, gpu_util_probe.py (long-ctx GPU-idle finding)." These are diagnostic scripts that helped identify the original bottleneck. The "long-ctx GPU-idle finding" refers to the discovery that GPU tensor core utilization was at only ~3% during decode despite 99.8% SM occupancy—a finding that directly motivated the custom kernel work.

The Thinking Process Revealed

The commit message reveals several layers of the assistant's reasoning. First, there is the deliberate choice of what to include. The assistant stages all changes (git add -A), not just the kernel source. This signals that the commit is intended as a complete snapshot of Phase 1—not just the code, but the tests, benchmarks, plans, and diagnostics that give the code its meaning.

Second, the bullet-point structure shows what the assistant considers the key claims worth documenting: correctness ("token-exact"), performance ("1.7-4.2x faster"), capability ("handles arbitrary prefix"), and integration readiness ("C-ABI"). These are the dimensions that matter for the next phase of work.

Third, the commit message title—"Phase 1: owned sm_120 flash-decode verify kernel (split-K, warp-per-key, online softmax)"—encapsulates the three key architectural decisions that made the kernel work. "Split-K" refers to splitting the KV cache across multiple thread blocks to improve occupancy. "Warp-per-key" means each warp of 32 threads cooperates on computing the attention score for one key, keeping all threads active during the score phase. "Online softmax" is a numerically stable algorithm that computes softmax incrementally without storing all scores. These three techniques together resolved the occupancy bottleneck that had limited earlier designs.

Assumptions and Their Implications

The commit makes several assumptions worth examining. It assumes that Phase 1 is genuinely complete—that the kernel is correct, fast enough, and ready for the next phase. But the assistant's earlier reasoning (visible in [msg 12244] and [msg 12247]) shows awareness that further optimizations are possible: double-buffering, tensor core utilization, and bf16 support are all mentioned as deferred improvements. The commit is a pragmatic checkpoint, not a declaration of perfection.

The commit also assumes that the benchmark results (1.7-4.2x) are representative of production performance. The assistant explicitly notes that the benchmark uses all 64 heads, while production with TP8 would use only 8 heads per rank, making the per-head design even more advantageous. This is a reasonable assumption, but it remains untested until Phase 2 integration.

Another assumption is that the C-ABI interface is sufficient for SGLang integration. The assistant has built the kernel as a standalone library with a C-compatible ABI, but integrating it into SGLang's Python-based serving stack requires additional work (Phase 2 in the plan). The commit assumes this interface is a solid foundation, which is reasonable but not yet validated.

Input Knowledge Required

To fully understand this commit, one needs knowledge of several domains. First, the GPU architecture context: sm_120 is the compute capability of the RTX PRO 6000 Blackwell consumer GPU, which has a different ISA from the data-center Blackwell (sm_100a/sm_103a). Understanding why existing MLA kernels don't work on this architecture is essential to appreciating why a custom kernel was necessary.

Second, the speculative decoding context: the verify attention kernel is part of a DFlash DDTree (Draft Tree) speculative decoding system. The kernel computes attention scores for draft tokens against the full KV cache, verifying whether the draft predictions are correct. This is a memory-bandwidth-bound operation in the decode phase, which is why the Triton kernel's scattered access pattern was so damaging.

Third, the software stack: the kdtree-engine project, the SGLang serving framework, the CUDA toolkit, and the build system. The commit references specific file paths, test commands, and benchmark outputs that assume familiarity with this stack.

Output Knowledge Created

This commit creates several forms of knowledge. Most concretely, it creates a permanent, versioned record of the Phase 1 kernel and all associated artifacts. Anyone who checks out this commit can reproduce the benchmarks, run the tests, and understand the kernel's capabilities and limitations.

The commit message itself is a form of documentation. It summarizes the key results in a way that is accessible to someone who wasn't present during the development session. The bullet points serve as a quick reference for what Phase 1 delivered.

The commit also creates a boundary between phases. By marking Phase 1 as complete, it enables Phase 2 (SGLang integration) to begin with a clear baseline. If something goes wrong in Phase 2, the team can always return to this commit and know that the kernel was correct and performing at this level.

The Deeper Significance

Beyond its technical content, this commit represents a deliberate pause in a fast-moving optimization session. The assistant could have continued iterating on the kernel—adding bf16 support, trying tensor core instructions, or optimizing the reduction kernel further. Instead, it chose to checkpoint, document, and commit. This decision reflects an understanding that engineering progress is not just about raw performance gains but about creating a reliable foundation for the next steps.

The commit hash—7a8dca3—is now a fixed point in the project's history. It marks the moment when a custom CUDA kernel for an unsupported GPU architecture went from an experiment to a deliverable. The 1.7-4.2x speedup over the naive oracle, the ability to handle 65k-token prefixes that the naive kernel couldn't fit, and the suite of diagnostic tools that revealed the original bottleneck are all captured in this single revision.

For the reader of this conversation, message [msg 12249] is a milestone marker. It tells us that Phase 1 is done, the kernel works, and the assistant is ready to move on to the harder problem of integrating this custom kernel into a live production service. The commit is both a conclusion and a beginning—the end of one engineering chapter and the start of the next.