How To Construct Your Personal LLM Runtime From Scratch

• What this publish is: a step-by-step tour of a from-scratch CUDA inference runtime for Qwen2.5-Coder-7B-Instruct on NVIDIA H100 (`sm_90`) — each sizzling path annotated for why, not simply what.
• What the runtime does at this time: grasping graph-tier token match passes on H100 at ≤4096 context, steady-state decode lands round ~16.7 ms/token (~60 tok/s), and first-token latency is ~128 ms for a 512-token immediate.
• The one quantity that made all the things else make sense: keen decode was ~119 ms/token. Wrapping decode in a captured CUDA graph collapsed that to ~17 ms/token. A 7× win — identical kernels, simply much less driver.
• Three bugs did a lot of the instructing: __syncthreads() inside a warp-specialized department is undefined behaviour (and corrupts softmax at KV web page tails), CUDA graphs will not be non-compulsory, and prmt.b32 beat __dp4a on Hopper for gate/up/down GEMV shapes — the INT8-activation path was tried and eliminated.
• llama.cpp is the reference implementation, not the rival: at commit b3040 with Q4_K_M on the identical GPU class, the identical pp512/tg128 protocol posts ~43 ms TTFT and ~4.95 ms/token decode. It’s the yardstick to examine whether or not your personal construct is behaving moderately.

TL;DR up entrance, so you possibly can go away with the purpose: when you have ever wished to really construct an LLM inference runtime your self — pack your personal weights, personal each barrier, seize your personal CUDA graphs — that is what that journey appears to be like like on an NVIDIA H100 (Hopper, sm_90) for Qwen2.5-Coder-7B. You’ll emerge with sturdy opinions about warp specialization, TMA, __syncthreads(), and the truth that a CUDA graph just isn’t a bonus, it’s a necessity. This text is a tour of a small LLM runtime known as annotated-llm-runtime — roughly two dozen CUDA/C++ supply information, each sizzling path commented for why, not simply what — and the three struggle (figurative) tales that produced a lot of the annotations.

Github Repo (will likely be used closely later all through the publish): https://github.com/AnubhabBanerjee/Annotated-LLM-Runtime

---

1. Why construct your personal LLM runtime in any respect?

Most manufacturing LLM inference already routes by way of a really well-tested path: GGUF → llama.cpp → achieved. It really works, it’s quick, it has a military of contributors, and it’s the proper software for delivery. Let me be clear proper now: nothing on this publish is making an attempt to problem that, and I like to recommend you retain utilizing it in manufacturing, as I do.

Possession of the decode stack is the place the worth is. Should you ever want to vary the quantization format, add a customized sampler, plug in a novel consideration variant, or ship on a brand new structure, you can not actually do it from inside a black field. You want to have the ability to open each kernel and know what is going to break — and what has already damaged and been fastened.

I feel, all through time, we’ve got already or we are going to ask ourselves the next query:

Can I write decode myself, perceive each launch, personal each barrier, and nonetheless get the proper tokens out?

The quick reply is sure. In case you have ever been inquisitive about how an actual decoder is definitely glued collectively — packed weight information, INT4 dequant, paged KV, warp specialization, CUDA graphs, TMA — however you may have been too intimidated to learn all fifty-thousand traces of GGML to search out out, this repo (and this publish) is identical journey compressed into roughly two dozen CUDA/C++ supply information with dense feedback. Each sizzling path within the code carries the why. Each barrier says which threads arrive on it and which threads should not. Each kernel says which form it was tuned for and which form it silently regressed on. It’s the model of the codebase I want had existed once I began.

A number of the annotations exist as a result of I needed to discuss myself by way of fixing the code. Others exist as a result of I needed to discuss myself out of a foul concept I had already applied. That is the second form of remark, within the wild:

// prmt.b32 path workout routines Hopper integer permute pipes — baseline earlier than fused scale broadcast.
// ValueShuffle lane map is frozen to match offline packer — runtime can't reinterpret nibbles.

That remark is the epitaph of a wasted afternoon.

2. The place the stack at the moment sits

Earlier than we stroll by way of the internals, the sincere present state of the runtime. That is measured on the identical H100 class, on the benchmark_e2e protocol locked within the config (batch=1, pp512/tg128, --warmup 3 --prompt-tokens 512 --gen-tokens 128):

In case you are rebuilding this stack your self, the llama.cpp column is a reference implementation for a similar protocol — pinned at commit b3040 with the usual offload flags (-ngl 99 -c 8192 -b 512 -cb) and the official Qwen/Qwen2.5-Coder-7B-Instruct-GGUF weights. Similar GPU class, identical mannequin, identical grasping decode, totally different quantity of tuning behind them. It’s the yardstick to examine that your personal construct is behaving moderately, not a rival to chase this early.

For the CUDA-flavored readers: the one largest lever inside this stack was CUDA graph decode. Keen per-token decode on the identical kernels got here in round 119 ms/token. Wrapping steady-state decode in a captured graph and replaying it dropped that to ~17 ms/token. Similar kernels, simply submitted in another way. Entire story on that in Part 6.

3. The stack in a single web page

Should you shut this tab after this part you’ll nonetheless have gotten the form. Each enter token walks the next path, as soon as per generated token, on the H100:

A bit little bit of grounding, as a result of the numbers matter for what comes later. From csrc/frequent/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 dimension of 3584, MLP width 18944 and vocab 152064. Customary Qwen2.5-Coder-7B-Instruct dimensions are captured inside constexpr so the C++ compiler owns them and by no means has to parse a YAML at runtime.

For grouped-query consideration (GQA), from csrc/frequent/paged_kv_layout.h:

// GQA: 28 question heads share 4 KV heads — ratio 7:1 on decode consideration CTAs.
static constexpr int kNumAttentionHeads = 28;
static constexpr int kNumKvHeads = 4;
static constexpr int kHeadDim = 128;

28 question heads, 4 KV heads — a 7:1 GQA ratio. That single quantity determines quite a bit: 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 web page in FP16 lands at a clear 4 KiB per KV head slice and a full 16 KiB per bodily web page. Each a type of numbers is chosen to align to Hopper’s 128-byte L2 cache line, and the file will lecture you about it within the feedback if you happen to overlook:

static_assert(kKvBytesPerPage % 128 == 0, "KV web page should align to 128-byte L2 traces");

Weights are symmetric group-wise INT4, group dimension 128, one FP16 scale per group per output row, scale = absmax / 7.0, zero-point fastened to 0. That covers precisely seven per-layer projections — q_proj, k_proj, v_proj, o_proj, gate_proj, up_proj, down_proj. All the things else, together with embed_tokens, all rms_norm, and lm_head, keep FP16 uncooked. It is a deliberate design by selection. When a numerical bug reveals up (and belief me, not less than one will), you need the search house to be small.

INT4 is packed into uint32 phrases, eight nibbles per phrase, in a selected “ValueShuffle” format that may make a complete part of this text’s fault once we get to Bug #3. For now, know solely this: the entire 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 recent 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 speak to it if these eight bytes are flawed.

// NANOQWEN magic bytes: 0x4E414E4F5157454E ("NANOQWEN") asserted earlier than any VRAM weight add.
// First 8 bytes of each .nanoqwen file — corrupt information fail quick earlier than giant cudaMemcpy.
static constexpr unsigned char kNanoqwenMagicBytes[8] = {
    0x4E, 0x41, 0x4E, 0x4F, 0x51, 0x57, 0x45, 0x4E,
};

The file merely begins with the phrase “NANOQWEN”. That’s the total security examine. If a file doesn’t begin with these actual letters, the system refuses to load six gigabytes of unknown information into the GPU. It isn’t a glamorous design, but it surely has already saved me from two main complications (together with one messy code merge).

The remainder of the file is organized right into a strict, predictable grid. Because the spacing is hardcoded, each the Python and C++ code know precisely the place to search for information while not having to learn any additional setup directions. Every bit of the venture—the info preparation software, the Python exams, and the C++ engine—makes use of this very same map. In the event that they don’t, they merely refuse to speak.

The remainder of the config (from config/nanoqwen.yaml) mirrors into csrc/frequent/nanoqwen_constants.h at compile time. Group dimension, ValueShuffle format id, L2 cache line width, storage-type tags for INT4 vs FP16 tensors — all static constexpr int. The runtime sizzling path by no means sees a YAML parser, which is the one easiest way to verify the recent path just isn’t by accident allocating on the heap someplace.

All the things previous this level is simply the story of what occurs when a uint32 stuffed with eight signed 4-bit integers hits the arithmetic pipes on an H100. And the story is sadly extra attention-grabbing than it ought to be.

5. Bug #1: the __syncthreads() that ate the softmax at KV web page tails

That is the place the abstraction drops and reveals the precise work.

The paged consideration kernel decodes consideration in opposition to a KV cache saved in 16-token pages, per KV head. On the H100 the plan is trendy and warp-specialized: one producer warp points Hopper’s bulk TMA copies (cp.async.bulk) to slurp Okay and V pages out of HBM into SMEM, and six client warps wait, then compute the web softmax + weighted-V accumulation in FP32 registers. Triple-buffered SMEM levels give the producer someplace to put in writing whereas the customers are nonetheless chewing the earlier tile.

In case you have by no means seen this form drawn out, that is the image the kernel is making an attempt to be:

Now right here is the bug in a single sentence: an earlier model of this kernel had a __syncthreads() sitting contained in the if (warp_id == 0) department — the producer warp — and nowhere else. That’s undefined behaviour in CUDA. __syncthreads() is a block-wide barrier. If just one warp arrives, the opposite six warps skip it completely and cheerfully march forward. On this kernel, “march forward” meant: client warps began studying Okay and V rows out of SMEM earlier than the producer’s bulk TMA had truly landed the bytes.

Full pages of Okay/V appeared high-quality — the copy normally completed earlier than customers acquired there anyway. However the second the KV cache ended on a partial tail web page (say, sequence size 49 → the final block holds one token, not sixteen), the timing shifted, customers learn no matter stale rubbish was nonetheless sitting in stage-0 SMEM, and softmax produced assured, coherent, flawed tokens.

The present kernel does the secure factor. Warp-scope work is fenced with __syncwarp() inside the producer department. The block-wide fence — the one each warp should attain — is positioned exterior any warp-id department, instantly earlier than customers contact the tile:

if (warp_id == 0) {
    // ... challenge TMA for Okay web page tile into smem_k_tile ...
    __syncwarp();
    // ... challenge TMA for V web page tile into smem_v_tile ...
    __syncwarp();
}
// Block-wide fence: client warps should see producer SMEM fills earlier than dot/softmax reads.
__syncthreads();

That remark just isn’t ornament. It’s the scar tissue from a painful debugging session.

The __syncwarp() calls work as a result of they solely want the 32 threads of a single warp to agree. The __syncthreads() name works as a result of each warp within the block hits it. However if you happen to place __syncthreads() inside an if (warp_id == 0) examine, the GPU hangs as a result of the opposite warps by no means attain that barrier. That mistake value me three hours of watching a KV tail.

To make issues safer, the synchronization between the producer and customers makes use of mbarriers—Hopper’s transaction-aware barrier objects—as an alternative of __syncthreads(). The producer tells the mbarrier precisely what number of bytes to count on. The customers wait on that very same mbarrier and solely get up when the {hardware} confirms the complete information switch is totally completed.

The core lesson: in a warp-specialized kernel, placing __syncthreads() inside any warp-id department is a assured bug, not an optimization. The softmax computation at KV tails is precisely the place this race situation will crash your code first.

To a profession HPC engineer, this error is so silly it borders on offensive. However as a telecom architect transitioning to bare-metal GPUs, I’m used to networks the place dropped packets ultimately simply work out a handshake and discover their manner house. It seems CUDA threads are far much less forgiving—they don’t retransmit; they simply sit there silently turning your costly Hopper GPU into a really inefficient house heater.

6. Bug #2: 119 → 17 ms per token, from a submission trick

If the primary bug was about correctness, the second is about why you need to by no means belief a per-token latency quantity that was measured with keen launches.

Right here is the unique timeline. Earlier than implementing CUDA graphs, calculating a single token meant firing off about 10 kernel launches per layer throughout 28 layers, plus the lm_head and the argmax operations. That’s over 280 particular person kernel launches for only one token. Each single launch forces a round-trip to the CPU driver. These microsecond delays stack up into milliseconds, and over 1000’s of tokens, the overall runtime turns into an absolute embarrassment.

Working keen decode on this setup clocked in at a painful 119 ms/token. After I appeared on the ITL breakdown (Inter-Token Latency), the precise GPU execution time barely registered. The scheduler sat idle 90% of the time as a result of the {hardware} was ravenous for directions. The bottleneck was by no means the mathematics; it was the large host CPU overhead brought on by cudaLaunchKernel.

The answer is captured graphs. As soon as your decode step is totally deterministic—utilizing the very same kernels and shapes with zero Python-level branching—you possibly can seize the complete sequence right into a single cudaGraphExec_t. You merely hand the CUDA driver one command to replay the entire construction, completely bypassing the per-kernel launch overhead. Implementing this dropped my steady-state decode from ~119 ms/token all the way down to ~17 ms/token. That could be a 7x speedup simply from submitting the very same kernels extra effectively.

There’s one essential catch. The paged-KV decode course of has a hidden branching level relying on whether or not the present token crosses a reminiscence web page boundary (place % 16 == 0). Since you can not put two totally different shapes into one static graph, the runtime solves this by pre-capturing each graph variants throughout setup and dynamically deciding on the proper one for every step.

constexpr int kPageBoundaryCapturePosition = 512;
constexpr int kNoPageCapturePosition = 513;

Two positions, one on a web 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’s the entire trick. cudaGraphLaunch(exec, stream) and you might be achieved. Each scalar the kernels want — present place offset, operating sequence size, the present token id — is a tool pointer that will get patched with a small cudaMemcpyAsync earlier than the launch, so the graph doesn’t need to be re-instantiated per step.

There’s a per-token ITL breakdown construction in csrc/runtime/decode_itl_breakdown.h that tracks eleven areas — qkv_gemv, append_tail_kv, paged_attention, o_proj, gate_up_gemv, swiglu_down, and so forth — plus a kernel_gap area that particularly measures the host overhead round cudaGraphLaunch. That final one exists as a result of “kernel hole” is precisely what CUDA graphs eradicate, and if the quantity ever goes again up, somebody by accident reintroduced keen decode someplace.

Discovered Lesson: if a decode step incorporates a whole bunch of kernel launches and every of them touches the driving force, you aren’t actually benchmarking your kernels. You might be benchmarking cudaLaunchKernel. CUDA graphs will not be a efficiency tweak. They’re the distinction between measuring your GPU and measuring your driver.

7. Bug #3: prmt.b32 beat __dp4a, and INT8 activations misplaced anyway

Each one of many seven quantised projections is INT4 weights × FP16 activations. Which sounds easy till you realise the 2 competing methods of doing this on Hopper look very related on paper and really totally different in a profiler.

Path A was __dp4a. Hopper nonetheless has quick INT8 dot-product-plus-accumulate: pack 4 signed INT8 operands right into a uint32, feed two of these to __dp4a, and also you get one INT32 accumulator replace. The catch is that either side of that dot product need to be INT8. So Path A required an additional kernel per decode step to quantise activations from FP16 to INT8 earlier than each GEMV, plus a per-group activation scale to un-do that quantisation contained in the accumulator.

Path B stored 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 brief:

__device__ __forceinline__ unsigned int prmt_b32(
    unsigned int source_low,
    unsigned int source_high,
    unsigned int selector) {
    unsigned int end result = 0;
    asm unstable("prmt.b32 %0, %1, %2, %3;"
                 : "=r"(end result)
                 : "r"(source_low), "r"(source_high), "r"(selector));
    return end result;
}

Adopted by a sign-extension helper for one signed nibble, as a result of INT4 in symmetric packing lives in [-8, 7] and doesn’t natively match right into a uint32 shift:

__device__ __forceinline__ int sign_extend_int4_nibble(unsigned int nibble_unsigned) {
    unsigned int nibble = nibble_unsigned & 0xFu;
    if (nibble & 0x8u)  0xFFFFFFF0u);
    
    return static_cast(nibble);
}

That’s the entire decode contract for one INT4 weight: masks to 4 bits, sign-extend if bit 3 is about, multiply by the group’s FP16 scale, use as an operand within the row’s FP32 accumulator.

On paper, Path A ought to have gained on gate_proj / up_proj / down_proj — the enormous MLP GEMV shapes the place activation bandwidth from HBM is the wall. INT8 activations are half the bandwidth of FP16 activations, so, straightforward win, proper?

Not on Hopper. Not for these shapes. The prmt.b32 FP16-activation path — Path B — beat __dp4a INT8-activations — Path A — on the precise gate/up/down shapes it was imagined to lose on. So the INT8 activation path was tried, benchmarked, after which eliminated. The code carries the tombstone: the __dp4a helper continues to be current in csrc/kernels/fused_gemv.cu as b1_accumulate_packed_word_dp4a, however the delivery launch path by no means turns it on.

There’s a associated element hiding within the packing format. The eight INT4 nibbles inside one packed uint32 do not stay 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.
change (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 excessive 16 bits. That isn’t aesthetic. That’s the offline packer laying nibbles out so a future variant of this kernel can permute them right into a register-file format that pairs properly with Hopper’s byte-permute datapath. The format is known as ValueShuffle and it lives at format id 1; if the runtime ever sees a .nanoqwen file with a special format id, it refuses to load, as a result of the nibble map contained in the C++ dequant is compiled in — you can not silently reinterpret packed weights and hope for the perfect.

Discovered Lesson: on Hopper, “INT8 is fewer bytes than FP16, due to this fact INT8 activations win on memory-bound GEMV” just isn’t a theorem. It is determined by the precise tile form, the precise SM occupancy, and whether or not the INT4 unpack you change is reasonable sufficient that the additional activation-quant kernel turns into the brand new bottleneck. On this venture it did.

8. The sincere scoreboard

At this level, you is likely to be questioning how I truly show that this practice engine works. The principles are written immediately into the config/nanoqwen.yaml file, and the runtime enforces them by way of a strict validation ladder:

• Major purpose: Obtain over 80% Reminiscence-Bandwidth Utilization (MBU) on the core math operations (the fused INT4 GEMV). The H100 SXM GPU has an enormous peak reminiscence bandwidth of three.35 TB/s; MBU is the precise proportion of that pace we handle to tug throughout execution.
• Secondary purpose: Match the coding benchmark (HumanEval go@1) of llama.cpp‘s Q4_K_M format precisely. Similar grasping decode, identical prompts, identical solutions.
• Tertiary purpose: Obtain a 99.5% or higher token match in opposition to a regular PyTorch INT4 simulation.

The exams are damaged down into three levels that should be handed so as: unit exams (testing single math operations), layer exams (testing one full decoder block in opposition to PyTorch), and graph exams (producing tokens for 100 prompts to match the PyTorch simulation). Crucially, I deliberately don’t attempt to power an ideal bit-for-bit match with the llama.cpp format. Their Q4_K_M format and my symmetric INT4 weights are essentially totally different on the math stage, and pretending they will align completely is how a venture like this begins mendacity to itself.

When the README says “correctness: sure,” it means the engine efficiently passes that closing graph-tier take a look at on an actual H100. The pace benchmarks—TTFT, ITL, and tokens/s—are the receipts I confirmed earlier. Hitting that main 80% MBU goal is the continued chase.

There’s another essential 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 entry). My philosophy right here is easy: if you happen to want root privileges to breed a benchmark, your benchmark just isn’t reproducible. The speeds you see are primarily based on default clock habits on actual, working silicon—not a artificially secure, locked-clock lab experiment.

9. Wrapping Up: What I truly Constructed and Learnt

I wrote a customized CUDA inference engine from scratch for Qwen2.5-Coder-7B on the NVIDIA H100 which utterly bypasses commonplace frameworks to make use of a customized memory-mapped file format (.nanoqwen), handles its personal INT4 mathematical operations, manages a paged KV cache, and runs the complete steady-state token technology utilizing simply two captured CUDA graphs.

Proper 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, absolutely the gold commonplace, llama.cpp, runs the very same take a look at at ~43 ms TTFT and ~4.95 ms/token. My runtime just isn’t making an attempt to beat llama.cpp at this time—these numbers are a yardstick to point out precisely how a lot excessive optimization goes right into a world-class, production-grade engine.

Constructing this taught me three huge classes about bare-metal AI:

• Barrier Bugs are Silent Killers: Placing __syncthreads() inside a warp-specialized department is a deadly error. It would silently corrupt your math (particularly the softmax on the KV web page tails). It’s essential to use __syncwarp() for branches, or higher but, depend on Hopper’s mbarriers to deal with the timing safely.
• CUDA Graphs are Necessary, Not Non-compulsory: Working this engine with commonplace per-kernel launches resulted in a horrible 119 ms/token. Capturing the very same operations into CUDA graphs collapsed the latency to 17 ms/token. That could be a 7x speedup purely from deleting the CPU submission overhead, with zero modifications to the underlying math.
• Math Codecs Matter: On Hopper structure, utilizing particular INT4 unpacking tips with FP16 activations merely outperformed making an attempt to power INT8 activations. I constructed the INT8 path, examined it, and finally ripped it out.

Finally, this venture is designed to be learn, not simply run. It’s a step-by-step information of precisely what occurs between your AI mannequin and the bodily GPU reminiscence bus. If that turns into the start of your personal decode stack, it’s a good higher end result.

Disclaimer: The illustrations on this article had been generated utilizing AI (Claude Opus 4.8). They’re illustrative, not photographic, and any labels seen inside the pictures are stylized reasonably than authoritative — confer with the article physique and the code itself for exact perform names, metric values, and structure particulars.