The Moment CUDA Graphs Captured: A Custom Kernel's Rite of Passage

In the high-stakes world of custom GPU kernel development, there is a moment that separates a working prototype from a production-ready component: the instant CUDA graph capture succeeds. On May 31, 2026, an AI assistant reached that milestone while deploying a custom sm_120 verify attention kernel for the Kimi K2.6 model on RTX PRO 6000 Blackwell GPUs. The message at index 12307 in the conversation is a deceptively brief confirmation of that success — but beneath its concise surface lies a rich story of disciplined engineering, careful reasoning, and the transition from correctness to performance.

The Message

The assistant writes:

## Agent Reasoning Great news—the CUDA graph capture is working correctly in ON mode, the server started up successfully, and the generated output matches what we get in eager mode. Now I need to verify there were no silent fallbacks and check the logs to confirm the capture succeeded and the kdtree installed properly.

>

Let me also validate that my dtype assumptions were correct by looking at the cg shapes log, then run the actual performance benchmark with graphs enabled to see how it compares against the original triton+graph baseline. CUDA-graph capture succeeded in ON mode and generation is correct. Let me confirm there were no capture errors, check the dtype log, then measure perf with graphs (the real production comparison):

The assistant then executes a bash command to inspect the server logs, revealing the telltale signs of success:

May 31 15:40:22 dflash-train python[31868]: [2026-05-31 15:40:22 TP3] Capture cuda graph end. Time elapsed: 1.53 s. mem usage=0.11 GB. avail mem=9.97 GB.
May 31 15:40:22 dflash-train python[31868]: [2026-05-31 15:40:22 TP3] Disable piecewise CUDA graph because --disable-piecewise-cuda-graph is set
May 31 15:40:22 dflash-train python[31866]: [2026-05-31 15:40:22 TP1] Capture cuda graph end. Time elapsed: 1.53 s. mem usage=0.11 GB. avail mem=9.97 GB.

The Road to Capture

To understand why this message matters, one must appreciate the journey that led to it. The assistant had been building a custom flash-decode verify attention kernel for the sm_120 architecture — the compute architecture of NVIDIA's RTX PRO 6000 Blackwell consumer GPU. This architecture lacks the advanced instruction set features (wgmma, TMA, tcgen05) found in Hopper (sm_90) and data-center Blackwell (sm_100/sm_103) GPUs, meaning none of the optimized MLA kernels like FlashMLA or flashinfer-MLA could be used. The team had to build their own.

The kernel itself implemented a "partial+reduce" design: splitting the KV sequence across multiple thread blocks (controlled by a parameter called NSPLIT), computing partial attention outputs in parallel, then reducing them. This design was necessary to achieve reasonable occupancy on sm_120's smaller shared memory (100KB versus the 228KB available on data-center Blackwell).

But correctness alone was not enough. For the kernel to be useful in production, it needed to work with SGLang's CUDA graph capture mechanism. CUDA graphs allow the framework to record a sequence of GPU operations and replay them with minimal CPU overhead — a critical optimization for serving workloads where even microseconds of CPU orchestration per decoding step add up. However, CUDA graph capture imposes severe constraints: no host-device synchronization (no .item() calls, no CPU-side data dependencies), no dynamic memory allocation during capture, and stable pointer addresses that persist across replays.## The Capture-Safe Transformation

The message at index 12307 is the culmination of a multi-step refactoring effort that transformed the kernel from a correctness-validated prototype into a capture-safe production component. The reasoning trace in the preceding messages reveals the assistant's careful deliberation about what "capture-safe" actually requires.

The first challenge was eliminating host-device synchronization. In the original validation-mode kernel, the assistant had used .item() calls to compare outputs against a Triton baseline — a natural approach for debugging. But .item() forces a synchronous copy from GPU to CPU, which is forbidden during CUDA graph capture. The assistant recognized this and restructured the code so that validation comparisons happen after the kernel runs, in a separate branch that only executes in non-capture mode.

The second challenge was memory allocation. CUDA graph capture records the addresses of all memory accesses, so any allocation that happens during capture would either fail or produce incorrect replays. The assistant considered several strategies: a persistent module-level workspace tensor allocated once and reused, or a per-call torch.empty() allocation. After weighing the tradeoffs — persistent allocation is more memory-efficient but requires careful sizing; per-call allocation is simpler but adds overhead — the assistant settled on a hybrid approach: allocate workspace per-call using torch.empty(), which is safe because PyTorch's CUDA graph capture intercepts allocations and routes them through graph-private memory pools.

The third challenge was ensuring that all tensor pointers passed to the kernel were stable and native. The assistant traced through SGLang's internal buffer management to determine the exact data types of each index pointer: kv_indices, qo_indptr, mask_indptr, and out_cache_loc are int64; kv_indptr is int32; and custom_mask is uint8 during capture (but bool in eager mode — both are single bytes, so reading as uint8_t handles both). By passing these native buffers directly via data_ptr() rather than making copies or conversions, the assistant eliminated the capture-breaking operations that had previously prevented graph recording.

The NSPLIT Tradeoff

One of the most interesting reasoning threads visible in the preceding messages concerns the NSPLIT parameter. This parameter controls how many thread blocks the kernel splits the KV sequence into — more splits means more parallelism but also more memory for partial accumulators and more reduction overhead.

The assistant initially set NSPLIT=64, which would maximize occupancy for long-context requests. But then it realized that NSPLIT must be fixed for CUDA graph capture — the kernel's grid dimensions are baked into the captured graph and cannot change between capture and replay. A fixed NSPLIT=64 would mean that short-context requests launch mostly empty blocks that early-out cheaply, but the workspace memory requirement scales with NSPLIT.

The assistant calculated that with NSPLIT=64, batch size 8, and the Kimi K2.6 model dimensions (9 query heads, 512-dimensional query key), the workspace would need approximately 2.4 GB. With only ~4 GB free per GPU, this was dangerously tight. After iterating through several options — reducing to NSPLIT=32, then NSPLIT=16 — the assistant settled on NSPLIT=16, which cut workspace to ~605 MB for batch size 8. For a single request at 200k tokens, NSPLIT=16 still yields 1152 thread blocks (16 splits × 72 blocks per split), providing solid occupancy. The assistant also preemptively lowered mem-fraction-static from 0.94 to 0.88 to free additional GPU memory for the workspace and graph storage, a pragmatic tradeoff since the server wasn't serving live traffic.

What the Logs Reveal

The journalctl output that the assistant retrieves is worth examining closely. It shows that CUDA graph capture completed in 1.53 seconds, using only 0.11 GB of additional memory, with 9.97 GB still available. The capture time is modest — under two seconds — which means SGLang can re-capture graphs if needed without significant downtime. The low memory usage (0.11 GB) suggests that the workspace allocation strategy worked as intended: the per-call torch.empty() allocations were routed to graph-private pools efficiently, and the kernel's fixed NSPLIT=16 kept workspace requirements manageable.

The log also confirms that piecewise CUDA graphs are disabled (--disable-piecewise-cuda-graph). Piecewise graphs are an optimization where SGLang captures separate graphs for different phases of the forward pass (prefill, extend, decode). Disabling them means the entire forward pass is captured as a single graph, which is simpler to debug and sufficient for the current deployment.

The fact that capture succeeded on all TP workers (TP1 and TP3 are visible in the logs; presumably TP0, TP2, and the remaining ranks also succeeded) confirms that the capture-safe design works correctly in SGLang's tensor-parallel deployment. Each rank independently captures its own graph, and the assistant's kernel modifications had to be safe on every rank simultaneously.## The Reasoning Process Visible in the Message

The assistant's reasoning in message 12307 reveals a disciplined, evidence-driven mindset. The first sentence — "Great news—the CUDA graph capture is working correctly in ON mode, the server started up successfully, and the generated output matches what we get in eager mode" — establishes three distinct verification criteria: capture success, server health, and output correctness. Each criterion addresses a different failure mode. Capture could succeed but produce incorrect results (if the graph recorded wrong operations). The server could be running but silently falling back to a different kernel path. The output could look plausible but differ from the eager-mode baseline in subtle ways.

The assistant then immediately pivots to "verify there were no silent fallbacks." This is a crucial insight: CUDA graph capture can appear to succeed while the actual kernel execution silently falls back to a different implementation. SGLang has multiple attention backends, and if the custom kernel fails during graph replay, SGLang might silently route to the Triton fallback. The assistant checks the logs specifically for error messages and capture invalidation events to rule this out.

The next step — "validate that my dtype assumptions were correct by looking at the cg shapes log" — shows the assistant's awareness that its entire capture-safe design rested on assumptions about SGLang's internal buffer types. If kv_indptr were actually int64 instead of int32, or if custom_mask had a different memory layout during capture, the kernel would read garbage indices and produce incorrect results. The "cg shapes log" (a diagnostic feature the assistant had built earlier) records the actual shapes and dtypes of the buffers during capture, providing empirical validation of these assumptions.

Finally, the assistant plans to "measure perf with graphs (the real production comparison)." This is the ultimate test: the custom kernel was built not just to be correct and capturable, but to outperform the Triton baseline. The 3–6× decode speedup achieved in earlier microbenchmarks (see [chunk 66.1]) was measured without CUDA graphs; with graphs, the CPU orchestration overhead is eliminated entirely, potentially yielding even larger gains.

Assumptions Made and Validated

Several assumptions underpin this message, and the assistant explicitly plans to validate them:

  1. The capture-safe kernel produces identical output in ON mode and eager mode. The assistant confirmed this with a generation test ("The capital of France is" → "Paris. This is a well-known fact"), but plans to run a more thorough benchmark.
  2. The dtype mappings are correct. The kernel reads kv_indptr as int32, all other index pointers as int64, and custom_mask as uint8. The assistant plans to cross-check these against the cg shapes log.
  3. No silent fallback to Triton occurred. The log inspection is designed to catch this.
  4. The workspace memory is sufficient. With mem-fraction-static lowered to 0.88, the assistant assumed ~6 GB of free memory would accommodate the workspace and graph storage. The logs confirm 9.97 GB available after capture, validating this assumption.
  5. CUDA graph capture works across all TP ranks. The logs from TP1 and TP3 confirm this, and the assistant implicitly assumes the other ranks also succeeded (a reasonable inference given the server is serving requests).

The Broader Context: A Milestone in the DDTree Project

This message sits within a larger narrative: the deployment of speculative decoding with DFlash and DDTree on the Kimi K2.6 model. The project had progressed through several phases — building a native C/C++/CUDA DDTree inference engine, diagnosing throughput regressions, extending context length to 200k tokens, and building a custom sm_120 verify attention kernel. Each phase uncovered new bottlenecks: first the Triton-based verify attention was too slow, then the custom kernel needed occupancy tuning, then CUDA graph capture required a complete rewrite of the kernel interface.

Message 12307 represents the moment when all these pieces finally clicked into place. The custom kernel was not just correct — it was capture-safe, integrated into SGLang's tensor-parallel serving stack, and running in production mode with CUDA graphs enabled. The assistant's next step — running the actual performance benchmark — would determine whether the 3–6× decode speedup observed in microbenchmarks translated to real-world serving throughput.

The message also reveals what the assistant didn't say. There is no celebration, no lengthy retrospective. The tone is matter-of-fact: "Great news... Now I need to verify... then run the actual performance benchmark." This is the hallmark of an engineer who treats success as a checkpoint, not a destination. The capture succeeded — good. Now let's measure whether it actually matters.

Conclusion

Message 12307 captures a pivotal moment in custom GPU kernel development: the transition from "works in isolation" to "works in production." The assistant's careful reasoning — verifying capture success, checking for silent fallbacks, validating dtype assumptions, and planning performance measurement — demonstrates a systematic approach to deployment that is as important as the kernel code itself. The CUDA graph capture succeeded in 1.53 seconds, the server served correct generations, and the path was clear for the performance benchmarking that would validate months of engineering effort. For anyone building custom GPU kernels for large language model serving, this message is a case study in what it takes to cross the bridge from prototype to production.