Most production LLM inference routes through well-tested paths. It works, it is fast, and it is the right tool for shipping. But ownership of the decode stack is where the value is. If you ever need to change quantization format, add a custom sampler, or ship on a new architecture, you need to know what will break.
I wrote a custom CUDA inference engine from scratch for Qwen2.5-Coder-7B on the NVIDIA H100. This is the story of what I learned.
The Numbers
At commit with Q4_K_M quantization, the engine posts ~60 tokens/second decode and ~128 ms time-to-first-token for a 512-token prompt[reference:69]. For context, llama.cpp runs the same test at ~4.95 ms/token. My runtime is not trying to beat llama.cpp today. Those numbers are a yardstick to show how much extreme optimization goes into a production-grade engine[reference:70].
Bug #1: __syncthreads() Inside a Warp-Specialized Branch
The paged attention kernel decodes attention against a KV cache. On the H100, the plan is warp-specialized: one warp issues bulk TMA copies to slurp K and V pages from HBM into SMEM, then compute online softmax in FP32 registers.
An earlier version of this kernel had a __syncthreads() sitting inside the producer warp branch and nowhere else[reference:71]. That is undefined behaviour in CUDA. __syncthreads() is a barrier. If only one warp arrives, the other six warps skip it entirely. In this kernel, "march ahead" meant consumer warps started reading K and V rows out of SMEM before the producer's bulk TMA had actually landed the bytes[reference:72].
Full pages of K/V looked fine. But the moment the KV cache ended on a partial block, timing shifted. Consumers read stale garbage from SMEM, and softmax produced confident, coherent, wrong tokens[reference:73].
The fix: place the block-wide fence outside any warp-id branch, immediately before consumers touch the tile[reference:74]. In a warp-specialized kernel, putting __syncthreads() inside any warp-id branch is a guaranteed bug, not an optimization[reference:75].
Bug #2: Eager Decode and the CUDA Graph Necessity
Before implementing CUDA graphs, calculating a single token meant firing off about 10 kernel launches per layer across 28 layers. Over 280 individual kernel launches for just one token[reference:76]. Every launch forces a round-trip to the CPU driver. Microsecond delays stack into milliseconds.
Running eager decode clocked in at a painful 119 ms/token[reference:77]. The actual GPU execution time barely registered. The scheduler sat idle 90% of the time. The bottleneck was never the math. It was the CPU overhead[reference:78].
The solution: captured graphs. Once the decode step is completely deterministic, you can capture the entire sequence into a single graph[reference:79]. Hand the CUDA driver one command to replay the whole structure. This dropped steady-state decode from ~119 ms/token to ~17 ms/token[reference:80]. A 7x speedup from submitting the same kernels more efficiently[reference:81].
One catch: the decode process has a hidden branching point depending on whether the current token crosses a memory page boundary[reference:82]. The runtime pre-captures both graph variants during setup and dynamically selects the correct one[reference:83].
Bug #3: INT8 Activations Lost to INT4 Unpack
Every quantised projection is INT4 weights × FP16 activations. Two competing ways exist on Hopper[reference:84].
Path A: quantise activations from FP16 to INT8 before every GEMV, use INT8 dot-product, then un-do the quantisation. INT8 activations are half the bandwidth of FP16. On paper, easy win[reference:85].
Path B: keep activations in FP16, use INT4 unpack with FP32 FMA multiply-and-accumulate[reference:86].
On Hopper, for these shapes, Path B beat Path A on the exact shapes it was supposed to lose on[reference:87]. The INT8 activation path was tried, benchmarked, and removed[reference:88].
The lesson: "INT8 is fewer bytes, therefore INT8 activations win on memory-bound GEMV" is not a theorem. It depends on tile shape, SM occupancy, and whether the INT4 unpack you replace is cheap enough that the extra activation-quant kernel becomes the new bottleneck[reference:89].
What I Learned
Three lessons stand out[reference:90]:
- Putting
__syncthreads()inside a warp-specialized branch will silently corrupt your math. Use block-wide fences outside branches. - CUDA graphs are not a performance tweak. They are the difference between measuring your GPU and measuring your driver.
- On Hopper, INT4 unpack with FP16 activations beat INT8 activations. The counterintuitive result won.
This project is designed to be read, not just run. It is a step-by-step guide of what happens between your AI model and the physical GPU memory bus.[reference:91]