The Build Script Edit That Made a Custom CUDA Kernel Real

[edit] /home/theuser/glm-kimi-sm120-rtx6000bw/kdtree-engine/scripts/build_nvcc.sh Edit applied successfully.

At first glance, message [msg 12262] appears almost trivial: a single-line edit to a build script, acknowledged with a terse "Edit applied successfully." Yet this message sits at a pivotal juncture in a months-long engineering effort to deploy a custom speculative decoding kernel on NVIDIA RTX PRO 6000 Blackwell GPUs (sm_120 architecture). It is the moment where a carefully designed piece of CUDA code — the paged+bf16 verify attention kernel — transitions from being a source file on disk to being a compiled, loadable, testable component of a live inference stack. Without this edit, the kernel written in [msg 12259] would remain inert code, never linked into the shared library, never called by the SGLang backend, never delivering the 3–6× decode speedup that the team had been chasing for weeks.

The Context: A Custom Kernel for an Unsupported Architecture

To understand why this edit matters, one must appreciate the broader engineering crisis that preceded it. The team was deploying the GLM-5-NVFP4 model with DFlash speculative decoding on RTX PRO 6000 Blackwell GPUs — a consumer-class Blackwell variant with compute capability sm_120. The problem was that every optimized MLA (Multi-head Latent Attention) kernel available in the ecosystem — FlashMLA, cutlass-MLA, flashinfer-MLA — was compiled exclusively for sm_90a (Hopper datacenter), sm_100a, or sm_103a (Blackwell datacenter). None supported sm_120's Ada-like instruction set, which lacks the wgmma, TMA, and tcgen05 instructions that datacenter Blackwell relies on.

The consequence was severe: the verify attention step in speculative decoding was falling back to a Triton-based MLA kernel with page_size=1, causing scattered KV cache accesses that achieved only ~14 GB/s effective bandwidth — roughly 130× below the GPU's 1.8 TB/s peak. At long context lengths (185k tokens), decode throughput collapsed to 0.7 tokens per second.

The user and assistant made a deliberate decision: rather than wait for upstream support that might never arrive for consumer Blackwell, they would build an owned sm_120 kernel. This was a significant commitment — writing custom CUDA for a niche architecture, maintaining it against SGLang's evolving interface, and owning the correctness and performance characteristics entirely.

The Message: What Actually Happened

Message [msg 12262] is the third in a four-message sequence that wires a new kernel into the build pipeline:

  1. [msg 12259]: The assistant wrote verify_attn_flash_paged.cu — a paged, bfloat16-capable variant of the flash-decode verify kernel, designed to read KV via kv_indices (SGLang's paged layout) and accept a compact prefix-length + visibility-matrix representation instead of a dense attention mask.
  2. [msg 12260]: The assistant added the function declaration to the header file verify_attn.cuh, making the kernel callable from other compilation units.
  3. [msg 12261]: The assistant added the C-ABI binding in capi.cu, creating a C-compatible entry point that Python (via ctypes or torch's CUDA extension mechanism) could invoke.
  4. [msg 12262] (target): The assistant edited build_nvcc.sh to include the new source file in the compilation command, ensuring that verify_attn_flash_paged.cu is compiled and linked into the shared library alongside the existing kernels.
  5. [msg 12263]: The assistant ran the build, which succeeded: [3/5] shared C-ABI lib completed, confirming the kernel compiled and linked correctly. The edit to build_nvcc.sh is the enabling step — it connects the kernel implementation to the build system so that the entire pipeline produces a binary that includes the new code.

The Reasoning: Why This Edit Exists

The assistant's thinking, visible in the reasoning blocks of surrounding messages, reveals a careful cost-benefit analysis. In [msg 12259], the assistant considered two strategies for validating the paged kernel against SGLang's real tensor layouts:

Option A (Tensor Capture): Temporarily patch the running SGLang service to dump actual tensors from a verify call, then develop the kernel offline against those real inputs. This would allow fast iteration without restarting the service, but required building an elaborate capture mechanism.

Option B (Double-Compute Validation): Build the kernel with best-guess shapes and layouts, then implement a backend that runs both the custom kernel and Triton's in parallel, logs differences, and returns Triton's result until differences converge to zero. This keeps the service correct throughout validation but requires 6-minute rebuild-and-restart cycles.

The assistant chose Option B, reasoning that "a few iterations to verify correctness is more practical than engineering an elaborate capture mechanism." This decision shaped everything that followed: the kernel needed to be built, compiled, and deployed into the live service quickly, with the build script being the critical path item.

The edit to build_nvcc.sh is therefore not a routine maintenance task — it is the direct consequence of a deliberate validation strategy that prioritized practical iteration over theoretical perfection.

Assumptions and Their Implications

Several assumptions underpin this message and the surrounding work:

Assumption 1: The build script is the correct integration point. The assistant assumed that adding the new .cu file to build_nvcc.sh would be sufficient to compile and link it. This assumes that the build script uses a glob or explicit file list, that the CUDA compilation flags are compatible with the new kernel (which uses bf16 and paged memory access patterns), and that no linker conflicts arise between the new kernel's symbols and existing ones. The successful build in [msg 12263] validated this assumption.

Assumption 2: The kernel's interface matches SGLang's actual tensor layouts. The assistant deduced the MLA key buffer shape as [num_slots, 576] (512 absorbed + 64 rope dimensions), the uniform q_len across the verify batch, and the custom mask format. These were educated inferences from reading SGLang's source code, not confirmed by dumping real tensors. The double-compute strategy was designed precisely to catch mismatches in these assumptions.

Assumption 3: CT200's freedom from production traffic justifies rapid iteration. The user confirmed that "CT200 does not serve any meaningful traffic right now," which gave the assistant license to restart services freely and use full GPU memory. This assumption made the Option B validation strategy viable — without it, the 6-minute restart penalty would have been unacceptable.

Assumption 4: The sm_120 architecture is stable enough to justify a custom kernel investment. The assistant and user committed to building and maintaining a kernel for an architecture that may never see broad adoption. This assumes that the performance gains (ultimately 3–6× over Triton) justify the long-term maintenance cost, and that future NVIDIA driver or CUDA toolkit updates won't break the custom kernel.

Input Knowledge Required

To fully understand this message, one needs:

  1. CUDA build system mechanics: Knowledge that .cu files must be explicitly listed in compilation commands or build scripts, and that adding a new kernel requires updating the build configuration.
  2. SGLang's speculative decoding architecture: Understanding that the verify attention kernel is called during DFlash speculative decoding to compute attention between draft tokens and the KV cache, and that SGLang uses a paged KV cache with page_size=1 for MLA.
  3. The sm_120 architecture gap: Awareness that consumer Blackwell GPUs (sm_120) lack the specialized instructions that datacenter Blackwell (sm_100a/sm_103a) provides, forcing custom kernel development.
  4. The project's file structure: The kernel source lives in src/kernels/, headers in src/kernels/verify_attn.cuh, C-ABI in src/kernels/capi.cu, and the build script at scripts/build_nvcc.sh.
  5. The double-compute validation strategy: Understanding that the backend will run both the custom kernel and Triton's in parallel, comparing outputs before committing to the custom path.

Output Knowledge Created

This message creates:

  1. A buildable kernel binary: After this edit, running build_nvcc.sh produces a shared library containing the paged+bf16 verify kernel, ready for testing and deployment.
  2. A validated build pipeline: The successful compilation in [msg 12263] confirms that the new kernel integrates cleanly with existing code, with no symbol conflicts or compilation errors.
  3. A foundation for backend integration: With the kernel compiled, the next step (Phase 2c) can implement the SGLang backend that calls this kernel during verify attention, using the double-compute strategy to validate correctness against real service tensors.
  4. A reusable pattern: The edit demonstrates how to extend the build system for future custom kernels — a pattern that may be reused if additional sm_120-specific kernels are needed.

Thinking Process: From Architecture to Build Script

The reasoning visible across messages [msg 12259] through [msg 12263] shows a disciplined engineering thought process:

  1. Problem diagnosis: The Triton MLA kernel with page_size=1 causes scattered memory access, achieving only ~14 GB/s bandwidth. The root cause is architectural — no optimized MLA kernel supports sm_120.
  2. Solution design: Build a custom flash-decode verify kernel that avoids page_size=1 scattering by reading KV directly via kv_indices, using split-K reduction for occupancy, and accepting a compact visibility matrix instead of a dense mask.
  3. Validation strategy selection: Choose double-compute over tensor capture, prioritizing simplicity and correctness guarantees over iteration speed.
  4. Implementation sequence: Write kernel → add header declaration → add C-ABI → update build script → compile → test. Each step depends on the previous, and the build script edit is the gate that enables all downstream steps.
  5. Risk management: The double-compute strategy ensures the service never serves incorrect results — the custom kernel's output is compared against Triton's, and Triton's result is returned until parity is confirmed.

Conclusion

Message [msg 12262] is a study in how the most consequential engineering work often hides in the most mundane operations. An edit to a build script — four words of diff — represents the culmination of architectural analysis, kernel design, validation strategy, and risk management. It is the moment when a custom CUDA kernel for an unsupported GPU architecture transitions from concept to compiled reality, enabling the 3–6× decode speedup that the team had been pursuing across weeks of debugging, profiling, and tuning.

The message also illustrates a key principle of systems engineering: the build system is not infrastructure; it is the delivery mechanism for every optimization, every fix, every new capability. Neglecting it means the best kernel ever written remains a file on disk. Attending to it — with a single, deliberate edit — turns code into performance.