The Build Script Edit: How a Single Line of Infrastructure Unlocked Custom CUDA Kernels on sm_120

The Message

[assistant] ## Agent Reasoning I need to update the build script to add verify_attn_flash to both the KOBJ list and the compilation loop that currently handles verify_attn, tree_accept, and capi. [edit] /home/theuser/glm-kimi-sm120-rtx6000bw/kdtree-engine/scripts/build_nvcc.sh Edit applied successfully.

At first glance, this message from the opencode coding session appears unremarkable—a routine build system update, barely a paragraph of reasoning followed by a single file edit. Yet this message represents the precise moment when weeks of architectural design, kernel implementation, and performance debugging crystallized into executable code. It is the infrastructural keystone that connected a newly written custom CUDA kernel to the compilation pipeline, the test harness, and ultimately the live SGLang inference service running on RTX PRO 6000 Blackwell GPUs. To understand why this message matters, one must trace the chain of reasoning that led to it, the architectural constraints that shaped it, and the performance breakthroughs it enabled.

Context: The Long Road to a Custom sm_120 Kernel

The subject message arrives near the climax of a sustained engineering effort to deploy the Kimi K2.6 model with DFlash speculative decoding on Blackwell GPUs. The project had already achieved a 3–6× decode speedup over Triton-based attention by tuning the existing verify kernel—increasing NSPLIT from 16 to 64 and adding 128-bit vectorized bf16 KV loads. But a deeper bottleneck remained: the verify attention kernel was still using Triton's MLA (Multi-head Latent Attention) implementation, which was compiled only for sm_90a (Hopper), sm_100a (Blackwell DC), and sm_103a architectures. The RTX PRO 6000 Blackwell consumer cards, with their sm_120 architecture and smaller 100KB shared memory limit, were left unsupported by every optimized MLA kernel library—FlashMLA, cutlass-MLA, and flashinfer-MLA alike.

This forced the assistant into a position that any serious ML engineer will recognize: the available tooling simply did not cover the target hardware. The only path forward was to write an owned kernel from scratch, tailored to sm_120's constraints. The assistant had produced a detailed per-phase plan (plans/0002-sm120-verify-kernel-defrag.md) and implemented Phase 1: a KV-split flash-decode MLA verify kernel in verify_attn_flash.cu. The kernel used a block-per-(batch, query-position, head-tile) decomposition, streaming KV tiles into shared memory with online softmax so that any prefix length could be handled within the ~70KB shared memory budget. The design reduced redundant global reads of the KV cache from 64× (naive) down to 8× per query, a 16× reduction when accounting for the dual read pattern of scoring and PV accumulation.

Why This Message Was Written: The Build System as the Last Mile

The assistant had already written the kernel source file, declared the function in the header (verify_attn.cuh), and wired it into the C-ABI shim (capi.cu). But none of that code would ever execute unless the build script knew to compile it. The subject message addresses exactly this gap: updating build_nvcc.sh to add verify_attn_flash to both the KOBJ list (the list of object files to link) and the compilation loop that iterates over source files.

This is the "last mile" problem in systems engineering. The intellectually demanding work—designing the flash attention algorithm, reasoning about shared memory budgets, choosing tile sizes (HT=8, TK=16) to stay under the 100KB limit, implementing online softmax for arbitrary-length prefixes—was already complete. But without this build script change, the kernel would remain an inert file on disk, never compiled, never tested, never deployed. The message embodies the principle that in production systems, infrastructure is not separate from innovation; it is the vehicle through which innovation reaches users.

The Reasoning Process: Efficiency Through Integration

The agent's reasoning reveals a deliberate strategy for minimizing round-trips and maximizing feedback velocity. The thinking text states: "I need to update the build script to add verify_attn_flash to both the KOBJ list and the compilation loop that currently handles verify_attn, tree_accept, and capi." This is not a random observation—it reflects a conscious decision made several messages earlier, where the assistant considered two approaches:

Approach A (cautious): Sync changes to the CT200 server first, run the existing baseline tests to confirm the environment works, then incrementally add the new kernel.

Approach B (aggressive): Write all the code—kernel, header, C-ABI, tests, benchmarks, and build script changes—locally in one batch, then sync and build everything together.

The assistant chose Approach B, reasoning that "the baseline tests still run as part of the full build, confirming the environment works while I add the new kernel code." This decision compressed what could have been multiple round-trips (sync, test, write kernel, sync, test, write tests, sync, test) into a single deployment. The build script edit in the subject message is the final piece of that integration strategy: it ensures that when the code lands on CT200 and the build runs, verify_attn_flash.cu is compiled alongside its siblings, and the test binary that compares both kernels against the golden reference will link correctly.

Input Knowledge Required

To understand this message, one must grasp several layers of context:

  1. The project structure: The kdtree-engine repository contains custom CUDA kernels for speculative decoding with DFlash and DDTree. The build system is a shell script (build_nvcc.sh) that compiles individual .cu files into object files and links them into test binaries and a shared library.
  2. The compilation model: Each kernel source file (e.g., verify_attn.cu, tree_accept.cu, capi.cu) is compiled separately by nvcc and the resulting .o files are linked together. Adding a new source file requires adding it to both the list of objects (KOBJ) and the compilation loop that emits the nvcc commands.
  3. The hardware constraint: The RTX PRO 6000 Blackwell GPUs use sm_120 architecture, which lacks support for Hopper/Blackwell-DC instructions (wgmma, TMA, tcgen05) used by all major MLA kernel libraries. This forced the custom kernel effort.
  4. The performance context: The existing Triton-based verify kernel was achieving ~14 GB/s effective bandwidth on long contexts—130× below the 1.8 TB/s peak—because of scattered KV access with page_size=1. The custom flash kernel aimed to close this gap through tiling, coalesced loads, and online softmax.
  5. The testing strategy: The test suite (test_verify_attn.cu) compares kernel outputs against a golden reference generated by a Python oracle. The build script must compile the new kernel so the test can A/B it against both the naive kernel and the reference.

Output Knowledge Created

This message produced a concrete, verifiable change: the build script at scripts/build_nvcc.sh was edited to include verify_attn_flash in the compilation pipeline. The consequences of this change ripple outward:

Assumptions and Potential Pitfalls

The message rests on several assumptions that are worth examining:

  1. The build script is correctly structured. The edit assumes that build_nvcc.sh has a well-defined KOBJ list and a compilation loop that iterates over source files. If the script used a different pattern (e.g., wildcard globbing, CMake integration, or a Makefile), this edit would fail. The assistant had previously read the script (in [msg 12223]), confirming its structure.
  2. The kernel compiles without errors. The edit assumes that verify_attn_flash.cu is syntactically correct and compatible with the nvcc flags used for the other kernels (architecture flags, include paths, etc.). If the kernel uses features not supported by the CUDA version on CT200, or if it has naming conflicts with existing symbols, the build will fail—but that failure will be caught in the next round when the assistant builds and tests.
  3. The compilation order doesn't matter. The edit assumes that adding the new file to the loop in any position is safe. If there are inter-file dependencies (e.g., the kernel calls functions defined in another .cu file that must be compiled first), the order could matter. However, the project uses separate compilation with linking, so order within the compilation loop is irrelevant.
  4. No other build targets need updating. The edit only touches the main build script. If there are secondary build targets (e.g., a Python package build, a Dockerfile, a CI configuration), they would need separate updates. The assistant implicitly assumes that the primary build_nvcc.sh is the sole compilation pathway.

The Broader Significance

This message, for all its brevity, captures a truth about systems engineering that is often overlooked: the most intellectually sophisticated algorithm is worthless if it cannot be compiled and deployed. The flash attention kernel with its carefully tuned tile sizes, its online softmax for arbitrary-length prefixes, and its 16× reduction in global memory reads—all of that engineering effort converges on this single edit to a shell script. Without it, the kernel is a theoretical exercise. With it, the kernel becomes a measurable intervention in the live inference pipeline.

The message also illustrates the opencode session's distinctive rhythm: extended periods of deep reasoning and design (the earlier messages where the assistant iterated through three design iterations for the kernel's thread mapping, weighing shared memory budgets and register tiling strategies) punctuated by brief, decisive actions (this edit). The assistant's thinking in the subject message is notably terse—just a single sentence of reasoning—because the decision space had already been fully explored in prior messages. The build script edit is not the product of new analysis; it is the execution of a plan already formed.

In the subsequent messages, the assistant extended the test suite to A/B both kernels against the golden reference and added flash rows to the benchmark. The build script edit was the prerequisite for all of that work. When the code was synced to CT200 and built, the flash kernel compiled successfully, passed token-exact correctness tests, and eventually achieved the 3–6× decode speedup that became the session's headline result. But none of that would have been possible without this quiet, infrastructural edit—the moment when a new kernel became real code.