- What this post is: a step-by-step tour of a from-scratch CUDA inference runtime for Qwen2.5-Coder-7B-Instruct on NVIDIA H100 (`sm_90`) — every hot path annotated for why, not just what.
- What the runtime does today: greedy graph-tier token match passes on H100 at ≤4096 context, steady-state decode lands around \~16.7 ms/token (\~60 tok/s), and first-token latency is \~128 ms for a 512-token prompt.
- The one number that made everything else make sense: eager decode was \~119 ms/token. Wrapping decode in a captured CUDA graph collapsed that to \~17 ms/token. A 7× win — same kernels, just less driver.
- Three bugs did most of the teaching:
__syncthreads()inside a warp-specialized branch is undefined behaviour (and corrupts softmax at KV page tails), CUDA graphs are not optional, andprmt.b32beat__dp4aon Hopper forgate/up/downGEMV shapes — the INT8-activation path was tried and removed. llama.cppis the reference implementation, not the rival: at commitb3040with Q4_K_M on the same GPU class, the samepp512/tg128protocol posts \~43 ms TTFT and \~4.95 ms/token decode. It is the yardstick to check whether your own build is behaving reasonably.
TL;DR up front, so you can leave with the point: if you have ever wanted to actually build an LLM inference runtime yourself — pack your own weights, own every barrier, capture your own CUDA graphs — this is what that journey looks like on an NVIDIA H100 (Hopper, sm_90) for Qwen2.5-Coder-7B. You will emerge with strong opinions about warp specialization, TMA, __syncthreads(), and the fact that a CUDA graph is not a bonus, it is a necessity. This article is a tour of a small LLM runtime called annotated-llm-runtime — roughly two dozen CUDA/C++ source files, every hot path commented for why, not just what — and the three war (figurative) stories that produced most of the annotations.
Github Repo (will be used heavily later throughout the post): https://github.com/AnubhabBanerjee/Annotated-LLM-Runtime
1. Why build your own LLM runtime at all?
Most production LLM inference already routes through a very well-tested path: GGUF → llama.cpp → done. It works, it is fast, it has an army of contributors, and it is the right tool for shipping. Let me be clear right now: nothing in this post is trying to challenge that, and I recommend you keep using it in production, as I do.
Ownership of the decode stack is where the value is. If you ever need to change the quantization format, add a custom sampler, plug in a novel attention variant, or ship on a new architecture, you cannot really do it from inside a black box. You need to be able to open every kernel and know what will break — and what has already broken and been fixed.
I think, at some point of time, we have already or we will ask ourselves the following question:
Can I write decode myself, understand every launch, own every barrier, and still get the correct tokens out?
The short answer is yes. If you have ever been curious about how a real decoder is actually glued together — packed weight files, INT4 dequant, paged KV, warp specialization, CUDA graphs, TMA — but you have been too intimidated to read all fifty-thousand lines of GGML to find out, this repo (and this post) is the same journey compressed into roughly two dozen CUDA/C++ source files with dense comments. Every hot path in the code carries the why. Every barrier says which threads arrive on it and which threads must not. Every kernel says which shape it was tuned for and which shape it silently regressed on. It is the version of the codebase I wish had existed when I started.
Some of the annotations exist because I had to talk myself through fixing the code. Others exist because I had to talk myself out of a bad idea I had already implemented. This is the second kind of comment, in the wild:
// prmt.b32 path exercises Hopper integer permute pipes — baseline before fused scale broadcast.
// ValueShuffle lane map is frozen to match offline packer — runtime cannot reinterpret nibbles.
That comment is the epitaph of a wasted afternoon.
2. Where the stack currently sits
Before we walk through the internals, the honest current state of the runtime. This is measured on the same H100 class, at the benchmark_e2e protocol locked in the config (batch\=1, pp512/tg128, --warmup 3 --prompt-tokens 512 --gen-tokens 128):
| Metric | This Solution | llama.cpp Q4_K_M |
| --- | --- | --- |
| TTFT (512 token prompt) | 128ms | 43ms |
| Decode ITL (steady-state) | 16.7 ms/token | 4.95 ms/token |
| Decode throughput | 60 tok/s | 200 tok/s |
If you are rebuilding this stack yourself, the llama.cpp column is a reference implementation for the same protocol — pinned at commit b3040 with the standard offload flags (-ngl 99 -c 8192 -b 512 -cb) and the official Qwen/Qwen2.5-Coder-7B-Instruct-GGUF weights. Same GPU class, same model, same greedy decode, different amount of tuning behind them. It is the yardstick to check that your own build is behaving reasonably, not a rival to chase this early.
For the CUDA-flavored readers: the single biggest lever inside this stack was CUDA graph decode. Eager per-token decode on the same kernels came in around 119 ms/token. Wrapping steady-state decode in a captured graph and replaying it dropped that to \~17 ms/token. Same kernels, just submitted differently. Whole story on that in Section 6.
3. The stack in one page
If you close this tab after this section you will still have gotten the shape. Every input token walks the following path, once per generated token, on the H100:
A little bit of grounding, because the numbers matter for what comes later. From csrc/common/qwen25_model_dims.h:
static constexpr int kHiddenSize = kNumAttentionHeads * kHeadDim;
static constexpr int kIntermediateSize = 18944;
static constexpr int kKvProjectionDim = kNumKvHeads * kHeadDim;
// ...
static constexpr int kNumHiddenLayers = 28;
static constexpr int kVocabSize = 152064;
28 decoder layers, hidden size of 3584, MLP width 18944 and vocab 152064. Standard Qwen2.5-Coder-7B-Instruct dimensions are captured inside constexpr so the C++ compiler owns them and never has to parse a YAML at runtime.
For grouped-query attention (GQA), from csrc/common/paged_kv_layout.h:
// GQA: 28 query heads share 4 KV heads — ratio 7:1 on decode attention CTAs.
static constexpr int kNumAttentionHeads = 28;
static constexpr int kNumKvHeads = 4;
static constexpr int kHeadDim = 128;
28 query heads, 4 KV heads — a 7:1 GQA ratio. That single number determines a lot: it says the KV footprint per token is small (4 heads × 128 × 2 bytes \= 1 KB per token), which is why paging KV at 16 tokens per page in FP16 lands at a clean 4 KiB per KV head slice and a full 16 KiB per physical page. Every one of those numbers is chosen to align to Hopper’s 128-byte L2 cache line, and the file will lecture you about it in the comments if you forget:
static_assert(kKvBytesPerPage % 128 == 0, "KV page must align to 128-byte L2 lines");
Weights are symmetric group-wise INT4, group size 128, one FP16 scale per group per output row, scale = absmax / 7.0, zero-point fixed to 0. That covers exactly seven per-layer projections — q_proj, k_proj, v_proj, o_proj, gate_proj, up_proj, down_proj. Everything else, including embed_tokens, all rms_norm, and lm_head, stay FP16 raw. This is a deliberate design by choice. When a numerical bug shows up (and trust me, at least one will), you want the search space to be small.
INT4 is packed into uint32 words, eight nibbles per word, in a specific “ValueShuffle” layout that will make an entire section of this article’s fault when we get to Bug #3. For now, know only this: the whole engine treats a .nanoqwen file as a mmap-able binary blob with a 256-byte header, a magic string, and packed INT4 payloads aligned to 128-byte L2 boundaries. No YAML on the hot path. Ever.
4. The five-line weight file
The offline packer produces a .nanoqwen file. The runtime opens that file, checks eight bytes, and refuses to talk to it if those eight bytes are wrong.
// NANOQWEN magic bytes: 0x4E414E4F5157454E ("NANOQWEN") asserted before any VRAM weight upload.
// First 8 bytes of every .nanoqwen file — corrupt files fail fast before large cudaMemcpy.
static constexpr unsigned char kNanoqwenMagicBytes[8] = {
0x4E, 0x41, 0x4E, 0x4F, 0x51, 0x57, 0x45, 0x4E,
};
The file simply starts with the word “NANOQWEN”. That is the entire safety check. If a file doesn’t start with those exact letters, the system refuses to load six gigabytes of unknown data into the GPU. It isn’t a glamorous design, but it has already saved me from two major headaches (including one messy code merge).
The rest of the file is organized into a strict, predictable grid. Since the spacing is hardcoded, both the Python and C++ code know exactly where to look for data without needing to read any extra setup instructions. Every piece of the project—the data preparation tool, the Python tests, and the C++ engine—uses this exact same map. If they don’t, they simply refuse to communicate.
The rest of the config (from config/nanoqwen.yaml) mirrors into csrc/common/nanoqwen_constants.h at compile time. Group size, ValueShuffle layout id, L2 cache line width, storage-type tags for INT4 vs FP16 tensors — all static constexpr int. The runtime hot path never sees a YAML parser, which is the single simplest way to make sure the hot path is not accidentally allocating on the heap somewhere.
Everything past this point is just the story of what happens when a uint32 full of eight signed 4-bit integers hits the arithmetic pipes on an H100. And the story is unfortunately more interesting than it should be.
5. Bug #1: the __syncthreads() that ate the softmax at KV page tails
This is where the abstraction drops and shows the actual work.
The paged attention kernel decodes attention against a KV cache stored in 16-token pages, per KV head. On the H100 the plan is fashionable and warp-specialized: one producer warp issues Hopper’s bulk TMA copies (cp.async.bulk) to slurp K and V pages out of HBM into SMEM, and six consumer warps wait, then compute the online softmax + weighted-V accumulation in FP32 registers. Triple-buffered SMEM stages give the producer somewhere to write while the consumers are still chewing the previous tile.
If you have never seen this shape drawn out, this is the picture the kernel is trying to be:
Now here is the bug in one sentence: an earlier version of this kernel had a __syncthreads() sitting inside the if (warp_id == 0) branch — the producer warp — and nowhere else. That is undefined behaviour in CUDA. __syncthreads() is a block-wide barrier. If only one warp arrives, the other six warps skip it entirely and cheerfully march ahead. 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.
Full pages of K/V looked fine — the copy usually finished before consumers got there anyway. But the moment the KV cache ended on a partial tail page (say, sequence length 49 → the last block holds one token, not sixteen), the timing shifted, consumers read whatever stale garbage was still sitting in stage-0 SMEM, and softmax produced confident, coherent, wrong tokens.
The current kernel does the safe thing. Warp-scope work is fenced with __syncwarp() inside the producer branch. The block-wide fence — the one every warp must reach — is placed outside any warp-id branch, immediately before consumers touch the tile:
if (warp_id == 0) {
// ... issue TMA for K page tile into smem_k_tile ...
__syncwarp();
// ... issue TMA for V page tile into smem_v_tile ...
__syncwarp();
}
// Block-wide fence: consumer warps must see producer SMEM fills before dot/softmax reads.
__syncthreads();
That comment is not decoration. It is the scar tissue from a painful debugging session.
The __syncwarp() calls work because they only need the 32 threads of a single warp to agree. The __syncthreads() call works because every warp in the block hits it. But if you place __syncthreads() inside an if (warp_id == 0) check, the GPU hangs because the other warps never reach that barrier. That mistake cost me three hours of staring at a KV tail.
To make things safer, the synchronization between the producer and consumers uses mbarriers—Hopper’s transaction-aware barrier objects—instead of __syncthreads(). The producer tells the mbarrier exactly how many bytes to expect. The consumers wait on that same mbarrier and only wake up when the hardware confirms the full data transfer is completely finished.
The core lesson: in a warp-specialized kernel, putting __syncthreads() inside any warp-id branch is a guaranteed bug, not an optimization. The softmax computation at KV tails is exactly where this race condition will crash your code first.
To a career HPC engineer, this mistake is so stupid it borders on offensive. But as a telecom architect transitioning to bare-metal GPUs, I am used to networks where dropped packets eventually just figure out a handshake and find their way home. It turns out CUDA threads are far less forgiving—they do not retransmit; they just sit there silently turning your expensive Hopper GPU into a very inefficient space heater.
6. Bug #2: 119 → 17 ms per token, from a submission trick
If the first bug was about correctness, the second one is about why you should never trust a per-token latency number that was measured with eager launches.
Here is the original timeline. Before implementing CUDA graphs, calculating a single token meant firing off about 10 kernel launches per layer across 28 layers, plus the lm_head and the argmax operations. That is over 280 individual kernel launches for just one token. Every single launch forces a round-trip to the CPU driver. Those microsecond delays stack up into milliseconds, and over thousands of tokens, the total runtime becomes an absolute embarrassment.
Running eager decode on this setup clocked in at a painful 119 ms/token. When I looked at the ITL breakdown (Inter-Token Latency), the actual GPU execution time barely registered. The scheduler sat idle 90% of the time because the hardware was starving for instructions. The bottleneck was never the math; it was the massive host CPU overhead caused by cudaLaunchKernel.
The solution is captured graphs. Once your decode step is completely deterministic—using the exact same kernels and shapes with zero Python-level branching—you can capture the entire sequence into a single cudaGraphExec_t. You simply hand the CUDA driver one command to replay the whole structure, entirely bypassing the per-kernel launch overhead. Implementing this dropped my steady-state decode from \~119 ms/token down to \~17 ms/token. That is a 7x speedup just from submitting the exact same kernels more efficiently.
There is one crucial catch. The paged-KV decode process has a hidden branching point depending on whether the current token crosses a memory page boundary (position % 16 == 0). Since you cannot put two different shapes into one static graph, the runtime solves this by pre-capturing both graph variants during setup and dynamically selecting the correct one for each step.
constexpr int kPageBoundaryCapturePosition = 512;
constexpr int kNoPageCapturePosition = 513;
Two positions, one on a page boundary and one not, captured back-to-back at warmup:
cudaError_t page_error = capture_one_graph(
&state->decode_graphs.exec_with_page_boundary,
state,
kPageBoundaryCapturePosition,
true,
stream);
cudaError_t no_page_error = capture_one_graph(
&state->decode_graphs.exec_no_page_boundary,
state,
kNoPageCapturePosition,
false,
stream);
And at replay time, the picker is a one-liner:
const bool page_boundary = (position_offset % kPagedKvTokensPerBlock) == 0;
cudaGraphExec_t exec = page_boundary ? state->decode_graphs.exec_with_page_boundary
: state->decode_graphs.exec_no_page_boundary;
Two graphs, chosen by a modulo. That is the whole trick. cudaGraphLaunch(exec, stream) and you are done. Every scalar the kernels need — current position offset, running sequence length, the current token id — is a device pointer that gets patched with a small cudaMemcpyAsync before the launch, so the graph does not have to be re-instantiated per step.
There is a per-token ITL breakdown structure in csrc/runtime/decode_itl_breakdown.h that tracks eleven regions — qkv_gemv, append_tail_kv, paged_attention, o_proj, gate_up_gemv, swiglu_down, and so on — plus a kernel_gap region that specifically measures the host overhead around cudaGraphLaunch. That last one exists because “kernel gap” is exactly what CUDA graphs eliminate, and if the number ever goes back up, someone accidentally reintroduced eager decode somewhere.
Learned Lesson: if a decode step contains hundreds of kernel launches and each of them touches the driver, you are not really benchmarking your kernels. You are benchmarking cudaLaunchKernel. CUDA graphs are not a performance tweak. They are the difference between measuring your GPU and measuring your driver.
7. Bug #3: prmt.b32 beat __dp4a, and INT8 activations lost anyway
Every one of the seven quantised projections is INT4 weights × FP16 activations. Which sounds simple until you realise the two competing ways of doing this on Hopper look very similar on paper and very different in a profiler.
Path A was __dp4a. Hopper still has fast INT8 dot-product-plus-accumulate: pack four signed INT8 operands into a uint32, feed two of those to __dp4a, and you get one INT32 accumulator update. The catch is that both sides of that dot product have to be INT8. So Path A required an extra kernel per decode step to quantise activations from FP16 to INT8 before every GEMV, plus a per-group activation scale to un-do that quantisation inside the accumulator.
Path B kept activations in FP16 and used a prmt.b32-flavoured INT4 unpack — Hopper’s byte-permute PTX instruction — with an FP32 FMA multiply-and-accumulate. The unpack itself is short:
__device__ __forceinline__ unsigned int prmt_b32(
unsigned int source_low,
unsigned int source_high,
unsigned int selector) {
unsigned int result = 0;
asm volatile("prmt.b32 %0, %1, %2, %3;"
: "=r"(result)
: "r"(source_low), "r"(source_high), "r"(selector));
return result;
}
Followed by a sign-extension helper for one signed nibble, because INT4 in symmetric packing lives in [-8, 7] and does not natively fit into a uint32 shift:
__device__ __forceinline__ int sign_extend_int4_nibble(unsigned int nibble_unsigned) {
unsigned int nibble = nibble_unsigned & 0xFu;
if (nibble & 0x8u) {
return static_cast<int>(nibble | 0xFFFFFFF0u);
}
return static_cast<int>(nibble);
}
That is the whole decode contract for one INT4 weight: mask to four bits, sign-extend if bit 3 is set, multiply by the group’s FP16 scale, use as an operand in the row’s FP32 accumulator.
On paper, Path A should have won on gate_proj / up_proj / down_proj — the giant MLP GEMV shapes where activation bandwidth from HBM is the wall. INT8 activations are half the bandwidth of FP16 activations, so, easy win, right?
Not on Hopper. Not for these shapes. The prmt.b32 FP16-activation path — Path B — beat __dp4a INT8-activations — Path A — on the exact gate/up/down shapes it was supposed to lose on. So the INT8 activation path was tried, benchmarked, and then removed. The code carries the tombstone: the __dp4a helper is still present in csrc/kernels/fused_gemv.cu as b1_accumulate_packed_word_dp4a, but the shipping launch path never turns it on.
There is a related detail hiding in the packing layout. The eight INT4 nibbles inside one packed uint32 do not live in a naive [0..7] order:
// Matches config/nanoqwen.yaml packed_layout [w0,w2,w4,w6,w1,w3,w5,w7].
// Lane order interleaves even/odd nibbles so prmt-friendly byte patterns emerge on Hopper.
switch (lane_in_octet) {
case 0: return 0;
case 1: return 16;
case 2: return 4;
case 3: return 20;
case 4: return 8;
case 5: return 24;
case 6: return 12;
case 7: return 28;
default: return 0;
}
Even lanes go into the low 16 bits, odd lanes into the high 16 bits. That is not aesthetic. That is the offline packer laying nibbles out so a future variant of this kernel can permute them into a register-file layout that pairs nicely with Hopper’s byte-permute datapath. The layout is called ValueShuffle and it lives at layout id 1; if the runtime ever sees a .nanoqwen file with a different layout id, it refuses to load, because the nibble map inside the C++ dequant is compiled in — you cannot silently reinterpret packed weights and hope for the best.
Learned Lesson: on Hopper, “INT8 is fewer bytes than FP16, therefore INT8 activations win on memory-bound GEMV” is not a theorem. It depends on the exact tile shape, the exact SM occupancy, and whether the INT4 unpack you replace is cheap enough that the extra activation-quant kernel becomes the new bottleneck. In this project it did.
8. The honest scoreboard
At this point, you might be wondering how I actually prove that this custom engine works. The rules are written directly into the config/nanoqwen.yaml file, and the runtime enforces them through a strict validation ladder:
- Primary goal: Achieve over 80% Memory-Bandwidth Utilization (MBU) on the core math operations (the fused INT4 GEMV). The H100 SXM GPU has a massive peak memory bandwidth of 3.35 TB/s; MBU is the actual percentage of that speed we manage to pull during execution.
- Secondary goal: Match the coding benchmark (HumanEval pass@1) of
llama.cpp‘sQ4_K_Mformat exactly. Same greedy decode, same prompts, same answers. - Tertiary goal: Achieve a 99.5% or better token match against a standard PyTorch INT4 simulation.
The tests are broken down into three stages that must be passed in order: unit tests (testing single math operations), layer tests (testing one full decoder block against PyTorch), and graph tests (generating tokens for 100 prompts to match the PyTorch simulation). Crucially, I intentionally do not try to force a perfect bit-for-bit match with the llama.cpp layout. Their Q4_K_M format and my symmetric INT4 weights are fundamentally different at the math level, and pretending they can align perfectly is how a project like this starts lying to itself.
When the README says “correctness: yes,” it means the engine successfully passes that final graph-tier test on a real H100. The speed benchmarks—TTFT, ITL, and tokens/s—are the receipts I showed earlier. Hitting that primary 80% MBU target is the ongoing chase.
There is one more crucial rule for honesty. The configuration explicitly refuses to lock the GPU clocks (lock_clocks: false). Locking the clocks on an H100 requires sudo nvidia-smi (root access). My philosophy here is simple: if you need root privileges to reproduce a benchmark, your benchmark is not reproducible. The speeds you see are based on default clock behavior on real, working silicon—not a artificially stable, locked-clock lab experiment.
9. Wrapping Up: What I actually Built and Learnt
I wrote a custom CUDA inference engine from scratch for Qwen2.5-Coder-7B on the NVIDIA H100 which completely bypasses standard frameworks to use a custom memory-mapped file format (.nanoqwen), handles its own INT4 mathematical operations, manages a paged KV cache, and runs the entire steady-state token generation using just two captured CUDA graphs.
Right now, the engine hits a Time-To-First-Token (TTFT) of \~128 ms and generates tokens at \~16.7 ms/token (\~60 tokens/second). For context, the absolute gold standard, llama.cpp, runs the exact same test at \~43 ms TTFT and \~4.95 ms/token. My runtime is not trying to beat llama.cpp today—those numbers are a yardstick to show exactly how much extreme optimization goes into a world-class, production-grade engine.
Building this taught me three massive lessons about bare-metal AI:
- Barrier Bugs are Silent Killers: Putting
__syncthreads()inside a warp-specialized branch is a fatal error. It will silently corrupt your math (specifically the softmax at the KV page tails). You must use__syncwarp()for branches, or better yet, rely on Hopper’smbarriersto handle the timing safely. - CUDA Graphs are Mandatory, Not Optional: Running this engine with standard per-kernel launches resulted in a terrible 119 ms/token. Capturing the exact same operations into CUDA graphs collapsed the latency to 17 ms/token. That is a 7x speedup purely from deleting the CPU submission overhead, with zero changes to the underlying math.
- Math Formats Matter: On Hopper architecture, using specific INT4 unpacking tricks with FP16 activations simply outperformed trying to force INT8 activations. I built the INT8 path, tested it, and ultimately ripped it out.
Ultimately, this project is designed to be read, not just run. It is a step-by-step guide of exactly what happens between your AI model and the physical GPU memory bus. If that turns into the beginning of your own decode stack, it is an even better outcome.
Disclaimer: The illustrations in this article were generated using AI (Claude Opus 4.8). They are illustrative, not photographic, and any labels visible inside the images are stylized rather than authoritative — refer to the article body and the code itself for precise function names, metric values, and architecture details.