Understanding Inference for Systems Engineers

By Dmitri Melikyan | Sep 11, 2026

This article is AI-generated, human-curated, and based on real inference engines and their measurements.

Contents

1. The one number that governs everything
2. From text to tokens
3. The forward pass, one layer at a time
4. Quantised weights and the kernel that streams them
5. Prefill: processing many tokens at once
6. Assembling a step: launches, graphs, and megakernels
7. The agent loop: sessions, speculation, and grammars
8. Correctness: what “the engine is right” means
9. Profiling an inference engine
10. Building an engine for one model, one GPU and one workload
11. Glossary

Who this is for, and what we will use as the running example

You know what a GPU is, you have written or at least read a CUDA kernel, you know that memory bandwidth and occupancy exist, and you can read a profiler trace. You know that a neural network has a forward pass and that training also has a backward pass. You have heard the words “KV cache”, “prefill” and “decode” and have a rough idea of what they mean. What you do not have is a mental model of how all of these fit together when a model actually serves requests, and the documentation you find is either a paper written for ML researchers or a serving engine’s source tree written for people who already understand it.

This article is the missing middle. It walks through inference from the first byte of a prompt to the last generated token, in the order the machine performs the work, and it stops at every term that would otherwise send you to a search engine. Every mechanism is illustrated with one concrete engine: a purpose-built, standalone CUDA/C++ inference engine that was designed, verified and measured for a single model on a single GPU, and compared against two widely used open-source serving engines (vLLM and SGLang) running the same weights on the same machine. Parts 1 to 9 use that engine only as a source of examples and measurements; Part 10 describes how it was built and how the comparison was run. Where a number appears, it was either measured on that machine or is arithmetic on the model’s configuration file, and the text says which.

The running example:

The model was chosen because it contains almost every mechanism a modern inference engine has to deal with: two different kinds of sequence-mixing layer, a quantised weight format, a built-in draft head for speculative decoding, a chat template with structured tool-call markup, and a context window of 262,144 tokens. If you understand how this one is served, you can read the design of most engines.

A note on scope. Everything here is about inference: running a trained model forward to produce tokens. Backpropagation never happens. That single fact shapes the whole problem, because without a backward pass nothing needs the intermediate activations of every layer; the only state retained is the per-sequence attention cache and recurrent state, and the entire cost structure collapses into one question: how fast can you move the weights through the GPU?

Part 1. The one number that governs everything

1.1 Decode is weight streaming

Start with the simplest possible description of what a language model does at generation time. It takes the sequence of tokens so far, runs a fixed computation over it, and produces a probability distribution over the next token. You pick one token from that distribution, append it, and run the computation again. That loop is called decode, and each iteration is a decode step. The step is dominated by large matrix products of the form

y = W · x

where W is a weight matrix stored in GPU memory and x is a small vector of activations computed from the current token. At batch size 1, x is a single vector, so each weight element is read from memory once and used in exactly one multiply-add. A modern GPU can perform vastly more floating-point operations per second than it can read bytes per second, so a computation that does about one operation per byte is not limited by arithmetic at all. It is limited by how fast bytes come out of memory.

For our model, the arithmetic on the configuration file says:

The gap between what this GPU can compute and what it can read is about three orders of magnitude: at 1.9 FLOP/B the tensor cores are essentially idle, waiting on memory. This is what people mean when they say decode is memory-bound or bandwidth-bound.

The consequence is a hard floor. The GPU’s sustainable read bandwidth, measured with a plain streaming kernel:

No kernel can run one forward pass of this model faster than about 17.5 ms on this GPU, because the weights have to be read. The scope of that statement matters: it is an idealised weight-streaming bound for one non-speculative decode step at batch 1 at the measured bandwidth. Part 7.4 shows how speculative decoding amortises one weight pass over several emitted tokens and brings the per-token cost well below it. For the plain decode step, every performance discussion is a discussion of what fraction of the measured bandwidth you reach. The two production engines measured at 21.3 ms and 21.4 ms per step, so the floor is about 82 % of their time, or equivalently they reach about 82 % of the measured bandwidth. The purpose-built engine’s discrete-kernel path reached 20.9 ms, about 83.5 %. Those percentages, not the raw milliseconds, are the numbers that tell you how good a kernel is.

This way of thinking is called a roofline: you compute the minimum time implied by the bytes that must move (the “memory roof”) and the minimum time implied by the arithmetic that must happen (the “compute roof”), and the larger of the two is the bound. For decode the memory roof is so much higher than the compute roof that the compute roof can be ignored.

1.2 Where the bytes go

If the token costs 26.9 GB of weight reads, it matters which weights those are, because each kind of layer gets a different kernel. The breakdown for our model, arithmetic on the configuration and confirmed against the checkpoint files byte for byte:

Two facts follow. First, almost two thirds of every token is the feed-forward network, the structurally simplest part of a transformer. Whatever you do to attention, you cannot speed up decode by more than about a third unless you also speed up the feed-forward path. Second, the vocabulary projection is one tensor that is almost a tenth of the whole model. Both facts drive design decisions later.

Notice what is not in the table: the token embedding table. It is another 2.54 GB, but decode reads one row of it (10 KB) per token, not the whole thing. Tables you index into are cheap; matrices you multiply by are expensive.

1.3 What else moves per token

Weights are not the only bytes. Two other kinds of state are read and written every step:

At short contexts these are rounding errors on top of the weights. But the key/value cache grows with context, at 64 KiB per token for this model, and the recurrent state does not grow at all. At the maximum context of 262,144 tokens the cache read is 17 GB per step, 64 % of the weight stream. Model-wide, the cache never overtakes the weights inside the context limit, but an individual attention layer’s own weights (105 MB) are overtaken by that layer’s cache read at about 25,000 tokens, which is why the attention kernel has to be designed for a bandwidth target of its own rather than as an afterthought.

Part 2. From text to tokens

2.1 What a tokenizer does

A language model never sees text. It sees a sequence of integers called token ids, each in the range [0, V) where V is the vocabulary size. A tokenizer converts bytes to ids and back. For this model V is 248,320, though only 248,077 ids are actually defined; the remaining 243 rows are padding so the vocabulary is a round multiple of a hardware-friendly number. The engine has to know this, because an id in the padding range has no text and must never be emitted.

The tokenizer is a byte-level BPE (byte pair encoding). The intuition: start with 256 tokens, one per byte value, then repeatedly find the most common adjacent pair in a training corpus and give that pair its own token, until you have the vocabulary size you want. The result is a merge table. Encoding text means splitting it into bytes and greedily applying merges in priority order. Common words become one token, rare words become several, and any byte sequence is representable, which is why it is called byte-level.

In practice there are more steps than that, and each one has to be reproduced exactly or the engine will feed the model a different sequence than the one it was trained on:

1. Special-token scan. The raw bytes are searched for a fixed list of added tokens (things like <|im_start|>, <tool_call>, <think>). These become ids directly and split the text into independent segments.
2. Unicode normalisation (NFC): different byte sequences for the same visible character are canonicalised.
3. Pre-tokenisation. A regular expression splits the text into chunks (roughly: words, numbers, punctuation runs, whitespace runs) so that a merge never crosses a chunk boundary. This regex has to be matched with the same alternation semantics as the reference implementation. For this model the tokenizer’s own configuration file declares one regex and the reference library silently uses another; the two disagree on 85 of 17,693 test strings. An engine has to follow the reference implementation the model is validated against, not the configuration file.
4. Merges within each chunk.

Decoding is the reverse: concatenate the byte strings of each id, then decode UTF-8, taking care to emit only complete characters when streaming (a multi-byte character can span two tokens).

The engine’s tokenizer was verified against the reference implementation on 17,693 test strings and a 2.2-million-string sweep over Unicode code points, encoding to identical ids and decoding byte-identically. This level of verification is necessary: a tokenizer defect produces a model that still generates fluent text, only slightly different text, and nothing downstream will detect it.

2.2 Generated id sequences are not always canonical

A given piece of text has one canonical tokenisation, the one the encoder produces. But an id sequence the model generated need not be canonical. In the agent workload of Part 7, when a grammar forces the model to emit the single character = and the model then freely generates start, the recorded ids are = followed by start. The canonical encoding of the text =start is a single, different token. The bytes are identical; the ids are not.

This matters for engineering because it means you cannot round-trip a session through text. If you detokenise a conversation and re-tokenise it, you get a different id sequence, the model’s cached state no longer corresponds to the input, and results silently change. The rule to adopt is simple: recorded ids are the truth, and nothing ever re-tokenises text that already has ids.

2.3 The chat template

A chat model is trained on conversations rendered into a specific plain-text format. The chat template is the recipe, and for this model it is a Jinja template shipped with the checkpoint. A rendered conversation looks roughly like:

<|im_start|>system
...instructions... # Tools ... <tools> {json schema of each tool} </tools> ...
<|im_end|>
<|im_start|>user
What files are in the current directory?<|im_end|>
<|im_start|>assistant
<think>
...the model's reasoning...
</think>

<tool_call>
<function=list_dir>
<parameter=path>
.
</parameter>
</function>
</tool_call><|im_end|>
<|im_start|>user
<tool_response>
README.md src/ ...
</tool_response><|im_end|>
<|im_start|>assistant
<think>

Every one of those angle-bracket markers is a token or a few tokens. <|im_start|>, <|im_end|>, <tool_call>, </tool_call>, <tool_response>, </tool_response>, <think> and </think> are each a single added token. <function= and <parameter= are three ordinary tokens each. The two ids <|im_end|> and <|endoftext|> are stop tokens: when the model emits either, generation of the current turn ends.

For the agent loop the important property is that the rendered conversation is append-only. Rendering the full message list after each tool result produces exactly the bytes you would get by concatenating the previous render, the model’s generation, and the new tool-response turn. The engine verified this at all 82 step boundaries of the workload of Part 7. It holds only because the template’s trimming rules happen to be no-ops on what this model emits (its reasoning traces end with exactly one newline before </think>). It is the property that lets an engine keep a conversation’s state on the GPU across steps rather than re-processing the whole history, and it is worth checking rather than assuming for any new model.

Part 3. The forward pass, one layer at a time

Now the actual computation. We describe the decode form, one token in, one probability distribution out, because it is the simplest and because it is where most of the time goes. Part 5 describes how the same computation is done for many tokens at once.

3.1 The residual stream

The model maintains a vector h of 5,120 numbers (the hidden size) that represents everything it knows about the current position. It starts as the embedding of the input token, a 5,120-element row gathered from the embedding table. Then 64 decoder layers each read h, compute something from it, and add the result back:

h = h + mixer(norm(h))        # mixer = attention or linear attention
h = h + ffn(norm(h))          # feed-forward network

The vector h flowing through all the layers with these additions is called the residual stream. The additions are why the network can be 64 layers deep: each layer only has to learn a correction. After the last layer, a final normalisation and one big matrix product turn h into a score for every vocabulary entry.

The two additions above, over 64 layers, are 128 sites where the reference implementation rounds to 16-bit. Matching them exactly is part of what “correct” means, as Part 8 discusses.

3.2 Numbers: bf16, fp32, and why both appear everywhere

Before the layers, the number formats, because every op description below says which one it uses.

bf16 (“brain float 16”) is a 16-bit floating-point format with 8 exponent bits and 7 mantissa bits. It has the same range as 32-bit float but only about three decimal digits of precision. Neural network weights and activations are stored in bf16 because halving the bytes halves the memory traffic, and models are tolerant of the precision loss.

fp32 is ordinary 32-bit float. Accumulating a long sum in bf16 would lose most of the terms, so every dot product accumulates in fp32 and rounds to bf16 once at the end. Normalisations, softmax and the linear-attention recurrent state also run in fp32.

A ulp (“unit in the last place”) is the gap between two adjacent representable numbers. For bf16 values between 16 and 32 the ulp is 0.125. This becomes important when comparing two implementations: a difference of one ulp in a score can flip which token wins.

The reference implementation makes specific choices about where to round from fp32 to bf16, and those choices are part of the model’s definition as far as reproducing its outputs is concerned. The op descriptions below note them.

3.3 RMSNorm

Every layer begins by normalising h. RMSNorm (root-mean-square normalisation) divides the vector by its root-mean-square magnitude and multiplies by a learned per-element weight:

xf   = fp32(x)                                   # widen
rstd = 1 / sqrt( mean(xf²) + 1e-6 )              # 1e-6 = epsilon, guards against zero
y    = xf * rstd * (1 + w)                       # w is the learned weight, stored zero-centred
out  = bf16(y)                                   # one rounding

It controls the scale of the input to each attention and feed-forward block, which is what makes training stable; in this pre-norm architecture the residual additions themselves are not normalised. The (1 + w) form is this model family’s convention: the stored weight is an offset from one, and the multiply is done in fp32 after normalisation. A model has several norm widths (5,120 for the residual stream, 256 for the per-head query/key norms in attention, 128 for a gated norm inside linear attention) and this model has two distinct norm shapes: the ordinary one above and a gated variant that rounds to bf16 before applying a plain (not 1 + w) weight and then multiplies by an activation of a second input. The engine implements exactly these three forms as device functions and nothing else. A decode step calls a norm 209 times, which is why they are inlined into neighbouring kernels rather than launched separately.

One instance from this engine’s verification: a norm implementation carried over from an earlier version computed the mean with a division instead of a multiply by a precomputed reciprocal. Mathematically identical; numerically not. It was caught only because the component was verified against the reference op with an instrument that reports the fp32 statistic before rounding, where a whole-tensor tolerance would have passed. The general rule: when exactness matters, test the intermediate quantity, not only the output.

3.4 Two kinds of layer

This model is a hybrid. Of its 64 layers, 48 are Gated DeltaNet layers (a form of linear attention, explained in 3.6) and 16 are ordinary full attention layers, in a repeating pattern of three linear layers then one attention layer. Each layer also has a feed-forward block (3.7).

Why mix them? Full attention lets every position look at every earlier position, which is powerful but costs a cache that grows with context and a per-step read of that whole cache. Linear attention compresses the entire history into a fixed-size state, so its cost per token is constant regardless of context. Interleaving them gets most of the modelling quality of full attention at a fraction of the long-context cost. For the engine it means two completely different mixer kernels, two completely different kinds of per-session state, and a cache-management problem that neither a pure-transformer nor a pure-recurrent design has.

3.5 The full attention layer

Attention answers the question “which earlier positions are relevant to the current one, and what do they say?” For each position the model computes a query vector q, a key vector k, and a value vector v. The current query is compared to every stored key by a dot product, the comparison scores are turned into weights that sum to one by a softmax, and the output is the weighted sum of the stored values. Keys and values from earlier positions are kept in the KV cache so they are not recomputed, and each new token appends its own k and v.

The model does this with several heads in parallel, each a smaller independent attention over a slice of the vector, so different heads can attend to different things. This model uses 24 query heads of dimension 256 but only 4 key/value heads; each key/value head is shared by 6 query heads. That sharing is grouped-query attention (GQA), and its purpose is purely to shrink the KV cache: 4 heads of keys and values instead of 24 means 6 times less cache per token. The engine’s kernel exploits it directly: one block handles a key/value head and all 6 of its query heads together, so the cache is read once per group rather than once per query head.

Step by step, for one new token at position pos, in this model:

1. Projections. Three matrix products turn the normalised h (5,120) into q (24 heads × 256, plus a same-sized output gate, see step 7), k (4 × 256) and v (4 × 256). The engine fuses them into one matrix product of 14,336 output rows so the weights stream once.
2. Query/key normalisation. Each q head and each k head is RMSNormed over its 256 dimensions. This is a stabilisation technique for training that inference must reproduce.
3. Rotary position embedding (RoPE). Attention as described so far has no idea where a token is. RoPE encodes position by rotating pairs of coordinates of q and k by an angle proportional to the position, with a different frequency per pair. Because a rotation by pos_q composed with a rotation by −pos_k depends only on the difference, the dot product q·k ends up depending on relative position, which is what you want. In this model only the first 64 of the 256 dimensions are rotated (“partial RoPE”), the pairs are (i, i+32) rather than adjacent, and the base frequency is 10 million. The multimodal variant of this model uses a three-axis positional scheme for images; for text all three axes carry the same position and it reduces exactly to ordinary RoPE, which the engine proved numerically (zero difference) before relying on it.
4. KV append. The rotated k and the raw v for this position are written into the cache: 4 heads × 256 × 2 bytes × 2 (k and v) = 4 KiB per token per layer, 64 KiB per token across the 16 attention layers.
5. Scores. s_j = (q · k_j) × 256^-0.5 for every cached position j ≤ pos. The scale factor keeps the dot products from growing with head dimension; here it is exactly 1/16, which is representable in every float format, so it introduces no rounding.
6. Softmax and weighted sum. p_j = exp(s_j − max s) / Σ exp(...) in fp32; the reference rounds p to bf16, then o = Σ p_j v_j. The subtraction of the maximum is numerical hygiene against overflow.
7. Output gate. This model multiplies the attention output element-wise by sigmoid(gate), where gate came out of the query projection in step 1. A gate in neural-network vocabulary is a learned multiplier that controls how much of a signal passes; here it is sigmoid(x) = 1/(1+e^-x), which squashes any number into (0, 1). Not every gate is bounded that way: the SiLU-shaped gate in the linear-attention layer (3.6) can exceed 1 and go slightly negative. This particular gate lets the model suppress attention output when it is not useful.
8. Output projection. One more matrix product (6,144 → 5,120), then add to the residual stream.

Steps 5 and 6 are the ones that read the cache and scale with context. Over a 4,096-token context that is 4 × 4,096 × 256 × 2 B × 2 = 16.8 MB per layer.

How the engine’s decode attention kernel is laid out. The context is split into P partitions and each block handles one (key/value group, partition) pair. With 4 groups and P = 47 that is exactly 188 blocks, one per SM. Each block computes the scores over its slice and the slice’s softmax statistics (its maximum m and its sum l = Σ exp(s − m)); the statistics of all partitions are then merged into the global M and L for each query head, which is possible because a partial sum taken against m can be rescaled to any other maximum by exp(m − M). This is the standard split-KV or flash-decoding structure. The ordering matters for exactness. The reference rounds each probability p_j = exp(s_j − M)/L to bf16 before multiplying by v_j, with M and L the global statistics. The engine’s kernel is therefore a two-pass design: scores and per-partition statistics first, then the merge into (M, L), and only then a second pass that forms each p_j under the global normaliser, rounds it to bf16 and accumulates p_j × v_j. A one-pass “online softmax” that rescales already-rounded or already-accumulated partial results afterwards is faster but produces different bits from this reference. The kernel reaches 71 % of the GPU’s read bandwidth at 32,768 tokens of context.

One subtle consequence of the split-KV structure: the fp32 merge of P partial results is not invariant to how the context is partitioned. Two runs with different P produce results that differ in the last bit. The engine fixes P = 47 at every position so that its two decode paths are bit-identical to each other; this kind of decision is what makes “bit-identical” claims possible and it has to be made deliberately.

3.6 The Gated DeltaNet layer

Now the layer that makes this model interesting. Linear attention replaces the softmax-weighted lookup over all past positions with a fixed-size recurrent state: a matrix S that is updated once per token and read once per token. Instead of storing every past key and value and re-reading them all, the model folds each new key/value pair into S and reads S with the query. The cost per token is the size of S, independent of how many tokens came before.

Gated DeltaNet (GDN) is a specific, recent linear-attention design. The two words in its name are two additions to plain linear attention:

• Delta rule. Plain linear attention adds k ⊗ v (an outer product) to S every step and S just accumulates. The delta rule instead first predicts what the state already thinks the value for this key is (S^T k), and only writes the difference between the actual value and that prediction, scaled by a learned amount. This is an error-correcting update, borrowed from classical online learning, and it lets the state overwrite stale associations instead of piling new ones on top.
• Gated. A learned per-head decay multiplies S by a factor between 0 and 1 every step, so old information fades. A second learned per-head write strength controls how much of the delta is applied.

Concretely, per token, per layer, in this model:

1. Projections. From the normalised h (5,120): a combined q|k|v vector of 10,240 elements (16 query heads of 128, 16 key heads of 128, 48 value heads of 128); a separate z vector of 6,144 that will gate the output; and two tiny vectors b and a of 48 elements each, one per value head, that become the gates.

2. Short causal convolution. Each of the 10,240 channels of q|k|v is passed through its own 4-tap filter over the last 4 positions (a depthwise causal conv1d), then through SiLU(x) = x × sigmoid(x), a smooth activation. The convolution gives each position a little local context before it enters the recurrence. To compute it at decode time the engine keeps the previous 3 inputs of every channel: the conv state, 10,240 × 3 × 2 bytes = 60 KiB per layer, kept as a small ring buffer indexed by position modulo 4.

3. Normalise q and k. Each 128-element head is scaled to unit length (l2norm). The reference does this in bf16 with five specific rounding points, and the engine reproduces all five, because an all-fp32 version differs on 28 % of rows. The normalised q is then multiplied by 1/√128 in fp32; the scaled vector is what the recurrence below reads the state with.

4. Gates. β = sigmoid(b), one per value head, is the write strength. g = −exp(A_log) × softplus(a + dt_bias), one per value head, is the log of the decay; exp(g) is between 0 and 1. A_log and dt_bias are small learned per-head constants; softplus(x) = log(1 + e^x) is a smooth version of max(0, x). This parameterisation comes from the state-space-model literature and guarantees the decay is a valid forgetting factor.

5. The 48-over-16 head mapping. There are 48 value heads but only 16 query/key heads. Value head j uses query/key head j // 3. So the state has 48 slices, each [128 keys × 128 values], and every three slices share one (q, k) pair. This is grouped attention again, in recurrent form.

6. The recurrence, in fp32, per value head:

   S      = S × exp(g)             # forget a little
   kv_mem = S^T k                  # what the state currently predicts for this key   (128 values)
   Δ      = (v − kv_mem) × β       # the correction, scaled by write strength
   S     += k ⊗ Δ                  # rank-1 update: S[i][j] += k[i] × Δ[j]
   o      = S^T q                  # read the state with the scaled query (q × 1/√128)  (128 values)

   S for one head is 128 × 128 fp32 = 64 KiB; for 48 heads, 3 MiB per layer; for 48 layers, 144 MiB per session. It is read and written once per token, 6 MiB per layer, and that is the entire per-token cost of the mixer regardless of context length.

7. Gated RMSNorm. The 128-element output of each head is RMSNormed (the gated form: round to bf16, multiply by a plain weight shared across heads) and multiplied by SiLU(z) for that head’s slice of z. This is where the z gate from step 1 acts: it lets the model scale each head’s contribution per token.

8. Output projection (6,144 → 5,120) and add to the residual stream.

Compare the two mixers as a systems engineer would:

The last two rows are the ones that make hybrid models hard to serve, and Parts 5 and 6 are largely about them.

A note on the state’s precision. The checkpoint’s configuration asks for the recurrent state in fp32. The reference implementation ignores that field and, through an accident of dtype propagation, keeps the state in fp32 during prompt processing but in bf16 during generation. Both production engines honour the config and use fp32. The purpose-built engine uses fp32 too, but had to add a bf16 mode purely so it could be compared against the reference during generation. A dtype toggle with no performance justification in an engine is usually explained by a requirement of this kind.

What the engine’s GDN kernel looks like. One block per (value head, 32 of the 128 value columns), 512 threads, 192 blocks. The state is laid out so that one value column’s 128 key entries are a contiguous 512-byte run, which makes both matrix-vector products in the recurrence warp-local reductions with no shared memory and no block-wide synchronisation. The conv state, the normalisation, the gates and the recurrence are all fused into one kernel, and the two 48-element gate projections (b and a) are computed inside it rather than as a separate matrix product, because a 96-row matrix product is far too small to fill 188 SMs and ran at 12 % of bandwidth as a standalone kernel. Fusing it recovered 0.23 ms per token, more than 1 % of the whole step. The kernel itself moves the state at 99.7 % of read bandwidth, but a fixed cost of about 5 µs per launch, 48 times a token, keeps the layer at 48 % of bandwidth overall. Per-launch fixed cost is the characteristic problem of small kernels; Part 6 covers the remedies.

3.7 The feed-forward network

After the mixer, every layer runs a feed-forward network (FFN), sometimes called the MLP. It has three weight matrices: two projections up to a wider dimension whose outputs are combined through a nonlinearity, then one projection back down, applied independently to each position:

gate = W_gate · x          # 5,120 → 17,408
up   = W_up   · x          # 5,120 → 17,408
act  = SiLU(gate) × up     # element-wise
y    = W_down · act        # 17,408 → 5,120

This particular shape, where one branch is passed through SiLU and multiplies the other, is called SwiGLU (Swish-gated linear unit; Swish and SiLU are the same function). The middle dimension, 17,408, is 3.4 times the hidden size, which is why the FFN is 267 MB per layer and 63.6 % of the whole token. The FFN is where a transformer stores most of its factual knowledge, and its size is most of what “27 billion parameters” means.

For the engine the FFN is three matrix-vector products, commonly executed as two kernel calls, and the design questions are how to fuse them. W_gate and W_up are stored concatenated along the row axis as one 34,816-row matrix, so one kernel streams both. Within that kernel a warp takes row j from the gate half and row j + 17,408 from the up half, so gate_j and up_j land in the same lanes and SiLU(gate_j) × up_j is computed in the epilogue as a register operation, never materialising the two 17,408-element intermediates. W_down then reads the 17,408-element activation and adds its output straight into the residual stream in its epilogue. Neither production engine fuses the SiLU-multiply into either FFN kernel on the fp8 path; both launch it separately. Fusing it saves a launch and a 17,408-element round trip per layer.

3.8 The head: from hidden state to token

After 64 layers, one final RMSNorm, then the lm_head (“language-model head”): a matrix product from the 5,120-element hidden state to a 248,320-element vector of logits, one score per vocabulary entry. In this model the head is a separate 2.54 GB bf16 matrix (some models tie it to the embedding table to save memory; this one does not, hence “untied”). The logits are rounded to bf16 because that is what the reference does, and that rounding is why ties between candidate tokens are common: 4.9 % of positions on average, and up to 15.6 % on some prompts, have two logits within one bf16 ulp of each other.

Then the token is chosen:

• Greedy decoding takes the argmax. On a tie the reference takes the lowest id. Verification always uses greedy, because it is deterministic.
• Sampling draws from the distribution softmax(logits / temperature) after optionally restricting to the top-k highest logits and the smallest set of tokens whose cumulative probability exceeds top-p. Each of these has an exact tie rule in the reference (top-k is a value threshold, so the kept set can be larger than k; top-p removes an ascending prefix and the highest tied id survives), and the engine measured each rule against the reference rather than assuming it.

The engine’s head kernel fuses the 248,320-row matrix product with the argmax or the top-k selection: each block produces a partial result (its maximum, the running log-sum-exp normaliser, and a small candidate list) and one small kernel merges them. The logits vector is never written to memory unless a debugging flag asks for it. At 2.54 GB against a 128 MiB L2 cache, the head is the one tensor in the model that can never be cache-resident, so every read of it is a DRAM read; the kernel reaches 98.5 % of bandwidth, the best of any kernel in the engine, because 248,320 rows give it far more parallel work than any other shape.

3.9 The whole step

Putting it together, one decoded token on this model is:

embed[token]                                   1 row gather
for each of 64 layers:
    norm → mixer (GDN or attention) → add residual
    norm → gate/up GEMV + SiLU·mul → down GEMV → add residual
final norm → lm_head GEMV → argmax / sample → next token

In the engine’s discrete-kernel form this is 677 kernel launches captured into one CUDA graph (Part 6 explains why a graph). Each launch’s configuration was chosen by a measured sweep. The table below is the component list; the GEMV kernel that appears on most rows is the subject of Part 4.

Part 4. Quantised weights and the kernel that streams them

4.1 fp8 with block scales

The checkpoint stores most weights in fp8, specifically the e4m3 format: 1 sign bit, 4 exponent bits, 3 mantissa bits. It has 256 encodings, of which two are zeros and two are NaN patterns, with a maximum magnitude of 448 and about one decimal digit of precision. That is far too coarse to store a weight matrix directly, so it is paired with block scales: the matrix is divided into 128 × 128 blocks, each block has one scale factor (stored in bf16 in this checkpoint, promoted to fp32 when loaded), and the actual weight is

W[n, k] = fp32(w_e4m3[n, k]) × scale[n / 128, k / 128]

The scale brings each block’s values into e4m3’s range, and the per-block granularity means one unusually large weight only costs precision for its own block. For a 5,120 × 17,408 matrix the scale grid is 40 × 136 = 5,440 numbers, about 0.01 % of the payload. The engine proved this rule (multiply, not divide; non-transposed index; per-block, not per-row) by falsification against the alternatives, and corroborated it by observing that every block’s maximum e4m3 magnitude is exactly 448, which is what a scale-to-fit quantiser produces.

Not everything is quantised. The lm_head, the embedding table, every norm weight, and the small GDN vectors stay bf16; their traffic is 2.6 GB of the 26.9 GB. The fp8 payload is 24.3 GB, half what it would be in bf16, which is the entire point: fp8 weights make the pure-bandwidth floor 17.5 ms instead of about 33 ms.

There is a companion question on the activation side. The checkpoint says activation_scheme: dynamic, meaning the reference implementation also quantises the input vector x to fp8 (one scale per 128 elements, computed on the fly) so that both operands of the matrix product are fp8 and a native fp8 × fp8 tensor-core instruction can be used. Both production engines do this; it costs them one small extra kernel per matrix product. The purpose-built engine does not: it keeps x in bf16 and widens the fp8 weights to bf16 in registers, which is lossless because every one of the 256 e4m3 values is exactly representable in bf16. This choice was made on precision grounds, not speed. When the engine compared the reference’s native fp8 path against a version with weights dequantised to bf16 up front, the dequantised version was 7 to 24 times closer to exact fp32 products per projection and 3.4 times more self-consistent between prompt processing and generation. The dynamic activation quantisation is the dominant error source in the reference. An engine that skips it is more accurate than the reference, which complicates verification: every disagreement with the reference then has to be explained rather than assumed to be the engine’s error.

4.2 DeepGEMM, CUTLASS, and why they are not the answer at batch 1

Two library families come up whenever fp8 block-scaled matrix products are discussed, and it helps to know what they are and why a batch-1 engine on this GPU uses neither directly.

DeepGEMM is an open-source library of fp8 block-scaled GEMM kernels written for datacenter Hopper and Blackwell GPUs. Its scale format is exactly what this checkpoint stores, so its layout conventions are reusable. Its kernels are not: they are built on the warp-group matrix instructions (wgmma) of Hopper and the tensor-memory instructions (tcgen05) of datacenter Blackwell, and they assume about 227 KB of shared memory per SM. The RTX PRO 6000 is a Blackwell workstation part (compute capability 12.0): it has the older warp-level mma.sync tensor-core instructions, no tensor memory, and 100 KB of shared memory per SM. DeepGEMM’s dispatcher does not even have a branch for it. One of the production engines has an explicit rule that disables DeepGEMM for this exact model type on Blackwell; the other resolves its backend to CUTLASS before DeepGEMM is ever considered. Either way, neither baseline runs DeepGEMM here, and a reader who sees “DeepGEMM” in a design note should read it as “the reference for the scale layout”, not “the kernel in use”.

CUTLASS is NVIDIA’s template library for tensor-core GEMMs, and it has a blockwise fp8 kernel for compute capability 12.0. Both production engines use it for this model. It is a GEMM kernel: it tiles the output matrix into rectangles (128 × 128 by default) and streams tiles of both operands through shared memory into the tensor cores. At batch 1 the “M” dimension of the problem (the number of input rows, i.e. tokens) is 1, and the smallest tile the kernel supports is 32 rows, so 31 of every 32 rows of tensor-core work are padding. That waste is free in a bandwidth-bound kernel: the tensor cores were idle anyway. What is not free is that the kernel is designed around large tiles and shared-memory staging, both of which exist to serve reuse that a single-row problem does not have. The production engines’ fp8 GEMM reaches about 82 % of read bandwidth at batch 1. One of them autotuned its kernel specifically on this model’s decode shapes, so this is a strong baseline.

4.3 Why the batch-1 GEMV uses tensor cores

The naive kernel for y = W·x at batch 1 is a GEMV (general matrix-vector product): each thread reads some weights, converts them from fp8, multiplies by the matching x elements, and accumulates. The engine measured that kernel first: it reached 1,164 GB/s, 75.6 % of bandwidth, and could not be pushed higher by any load-tuning.

The cause of that ceiling was not established. An instruction-issue budget was proposed and appeared to predict it: about 1.56 operations per weight byte against a per-SM budget of 4 instructions per cycle at 3.37 bytes per SM per cycle. That arithmetic compares thread-level operations with a budget of warp instructions, and a warp instruction covers 32 lanes, so the demand is about 0.16 warp instructions per cycle, far below the ceiling; corrected, issue throughput is not the limit, and the agreement with 75.6 % was a coincidence. The scalar path is more plausibly limited by how many bytes it keeps in flight per instruction, but that was not measured. What is established is the comparison: the scalar kernel reached 75.6 % of bandwidth and the tensor-core kernel below reached 88.7 %. One mma.sync instruction multiplies a 16 × 16 tile by a 16 × 8 tile and accumulates into 16 × 8 outputs, replacing 2,048 scalar multiply-adds, and it takes its operands from registers loaded in 16-byte units, which is the property the design exploits.

The engine’s GEMV therefore is a tensor-core kernel, but with the roles arranged for a single input vector:

• The A operand of the mma holds 16 output rows of W by 16 reduction elements. The B operand holds the activation x broadcast across all 8 of its columns. Column 0 of the result holds the 16 dot products; the other 7 columns are duplicates and are discarded. This wastes 7/8 of the tensor-core arithmetic, which is idle in a bandwidth-bound kernel anyway. Critically, the waste is on the n axis of the instruction, not on the row axis of a tile: no weight byte is ever loaded twice.
• At batch 1 the reduction axis can be permuted. Permuting matching weights and activations preserves the mathematical dot product but changes floating-point rounding; this kernel uses one fixed permutation, so its result is reproducible run to run, and its difference from the reference’s summation order is verified within one bf16 ulp per element. The kernel uses this freedom to choose which reduction slot each weight byte occupies so that every lane’s global load is 16 physically consecutive bytes. No shared memory for weights, no transpose, no ldmatrix.
• All 16 rows a warp owns fall in one 128-row scale block, and with a 128-wide reduction tile the whole 16 × 128 × 1 byte = 2 KB tile shares one scale. The scale multiply is applied once per 2 KB, to the fp32 accumulator, exactly as a blockwise GEMM’s main loop does it.
• The weights are read straight from the checkpoint’s own row-major byte order. No re-layout at load time.

Result: the whole family of decode-time projections runs in 19.7 ms per token, 88.7 % of read bandwidth. That beats the production engines’ whole-token figure of 82 %, but the relevant comparison is narrower: when one production engine’s own GEMM time was attributed from a profile it came to 19.5 ms. The purpose-built kernel is at parity with a well-tuned library kernel, within 1 %. That engine’s remaining 1.6 ms per token is spent elsewhere (small kernels, quantisation kernels, zero-fills), not in its GEMM.

4.4 Split-K and the row-count problem

A GEMV over a 5,120-row matrix has 5,120 outputs. On 188 SMs, with 16 rows per warp, that is only 320 row groups, and the kernel’s throughput turns out to sort by row-group count: the 5,120-row shapes reach 82 to 88 % of bandwidth while the lm_head’s 15,520 row groups reach 98.6 %. Fewer row groups means fewer independent memory streams in flight, and the kernel becomes latency-bound (waiting on individual loads) rather than bandwidth-bound.

The standard fix is split-K: divide the reduction dimension into S slices, have S times as many blocks each compute a partial dot product over its slice, and add the partials in a small second kernel. The engine chose S per shape by measurement (2, 4 or 8 in the table of 3.9). The reduce kernel adds partials in a fixed order so results are deterministic, which matters for the bit-identity claims of Part 8; an atomicAdd-based reduction would be faster to write and non-deterministic.

The trade-off is worth internalising: split-K costs an extra pass over the partials (small) and a launch (small under a graph), and buys memory-level parallelism. It is the same lever as “more threads in flight” but applied along the reduction axis when the output axis is too short.

Part 5. Prefill: processing many tokens at once

5.1 Prefill is a GEMM problem

Everything so far processed one token. When a prompt arrives, all of its tokens have to go through the model before the first output token can be produced. That is prefill. It dominates the time to first token (TTFT), which as measured at a client also includes request handling, the first vocabulary projection and sampling, and transport. In an agent loop there are two kinds of prefill: a cold one once per trajectory (the whole rendered prompt with tool schemas, often thousands of tokens) and an incremental one at each later step for the new tool result (typically a few hundred tokens).

Prefill runs the same layers, but with T rows of activations instead of one. Every matrix product becomes [T, 5120] × [5120, N], a real GEMM in which each weight is reused across the T rows. How many times a weight tile is actually fetched depends on the tiling and on L2 reuse, and activation traffic grows with T, so arithmetic intensity rises with T but sublinearly at large T, as the table shows:

On this GPU the crossover is somewhere between 16 and 64 tokens. A 200-token tool result is firmly compute-bound, and the consequence is direct: appending a tool result must not be implemented as 200 decode steps. It needs a GEMM kernel with real tiles, and the tensor-core waste that was free at batch 1 becomes the whole problem.

The fp8 question returns here. A native fp8 × fp8 mma instruction exists on this GPU and the engine verified it works and is 4 to 20 % faster; but it requires fp8 activations, which means quantising x, which means prefill and decode would compute different functions of the same weights. The engine kept bf16 activations in both, so that a token processed during prefill and the same token processed during decode agree to the limits of summation order.

The engine’s prefill GEMM was its weakest kernel. Whole-model GEMM time at 59 tokens was 57.9 ms against a 17.5 ms weight-streaming floor, about 30 % of bandwidth overall, and the three 5,120-row output shapes reached only 14 to 16 %, because a 59-row problem with 128-row tiles has neither enough rows to use the tensor cores well nor enough blocks to fill 188 SMs. A split-K version built later improved it (Part 10).

5.2 Prefill attention

Attention over T new tokens against past + T keys is a batched, causal problem: query i may only attend to keys at positions ≤ i. The cost grows as T × (past + T). This is the operation that flash-attention-style kernels exist for: tile the queries and the keys, keep the running softmax statistics in registers, never materialise the T × (past + T) score matrix.

The engine initially did not have such a kernel. It reused its decode attention kernel four queries at a time, which re-reads the entire KV prefix once per four queries. For a 4,096-token cold prefill that fallback was 43 % of the time; at 32,768 tokens it was 67 %, and a cold prefill took 37 seconds where the production engines took under 4. A two-pass tiled kernel built later cut cold prefill by 1.7 to 2×; Part 10 describes why it was nevertheless not adopted.

5.3 Chunked prefill for the recurrent layers: the WY form

Here is the part that has no analogue in a pure transformer. The GDN recurrence of 3.6 is inherently sequential: S after token t depends on S after token t − 1. Running it 4,096 times in a row for a prompt would be 4,096 dependent steps per layer, each a small matrix-vector product, and would take longer than the prompt is worth.

The solution is the chunked form: process tokens in chunks of 64, and within a chunk replace the 64 sequential rank-1 updates with a handful of small dense matrix products. The mathematical tool is the WY representation, borrowed from numerical linear algebra where it is used to represent a product of Householder reflections compactly. Intuitively: a sequence of rank-1 updates S ← S·decay + k_i ⊗ Δ_i can be rewritten so that all 64 Δ_i of a chunk are computed at once by solving a 64 × 64 lower-triangular linear system built from the pairwise key dot products and the decays, and then the chunk’s total effect on S is one [128 × 64] × [64 × 128] matrix product. Within a chunk the work is parallel; across chunks it is still sequential, but 64 times shorter.

In this model the chunked path per chunk is: compute the cumulative decays; build the 64 × 64 matrix of decayed key-key dot products scaled by β; invert (I − M) by forward substitution over the 64 rows; apply it to get the corrected values u and the corrected keys w; then run the inter-chunk scan (v_new = u − w·S, output, state update). The engine’s kernel decomposes this so that each block owns a slice of the 128 value columns and runs the entire sequential chunk scan for its columns with no cross-block synchronisation, because every line of the scan is element-wise in the value dimension.

Two names you will see next to this: FLA is flash-linear-attention, an open-source library of Triton kernels implementing exactly these chunked forms for a family of linear-attention models, and it is what both production engines run. The reference implementation has a pure-PyTorch fallback that is used when FLA’s fast path is unavailable, and that fallback is what the engine was verified against. The two use different conventions in places (FLA computes decays in base 2 with exp2; the reference uses natural exp), and the engine found by measurement that the exp2 form differs from the reference on 42 % of decay-matrix elements while the natural form is bit-exact. Same math, different bits; when your correctness criterion is the reference’s bits, you have to pick the reference’s convention.

5.4 The same math, different rounding: why prefill and decode disagree

The chunked form and the sequential recurrence compute the same real-number function. In fp32 they do not compute the same bits, because the operations are grouped differently. The reference implementation itself disagrees between its prompt-processing path (chunked) and its generation path (sequential). The engine measured that disagreement on 432 positions: the two paths of the reference produce top-1 tokens that differ at 2 of 432 positions, with logit differences up to 4.8. That is the noise floor of the reference. No implementation can be “more exact” than that against a reference that disagrees with itself, and any correctness criterion has to be calibrated to it.

There is a second, subtler consequence. If you process a prompt in two calls (say 300 tokens, then 200 more), the chunk boundaries fall at 0, 64, 128, … within each call. If you process all 500 at once, they fall at 0, 64, …, 448. Different chunking, different fp32 rounding, different state. The engine measured this: appending one token to a 319-token session and comparing against a 320-token single pass, 82 % of state elements differ (by about 4 parts per million). So a session built incrementally is not bit-identical to the same session built in one go, unless you do something about it.

The engine’s solution is worth knowing because the production engines face the same problem and solve it differently. The session keeps a checkpoint of the recurrent state and conv state at the last position that is a multiple of 64. An append of n tokens restores that checkpoint and re-runs from it: the up to 63 already-committed tokens since the checkpoint plus the n new ones, with chunks anchored at the sequence start. Those chunks are exactly the chunks a single pass would have used, so the result is bit-identical to a single pass over the whole stream. The price is an average of 31.5 extra tokens of prefill per append plus one 144 MiB state copy when a 64-boundary is crossed (0.2 ms). The engine proved the equivalence: zero differing bits across every tested combination of prompt length, append length and chunk size. The limit of the guarantee: an agent step is append, then prefill, then decode, and decode uses the sequential recurrence, so the state after a step is never bit-identical to a single chunked pass. The guarantee is about prefill against prefill.

Notice how much of this section is about numerics rather than speed. That is representative. In a hybrid model, “does incremental processing give the same answer as batch processing?” is a real engineering question with a measurable answer, and the answer defaults to “no”.

Part 6. Assembling a step: launches, graphs, and megakernels

6.1 The launch problem

A decode step in the engine’s discrete form is 677 kernels. In one of the production engines, measured with a profiler with CUDA graphs disabled, it is 1,559 kernels, of which 512 are zero-fills of padding rows created by the fp8 GEMM’s 4-row alignment requirement and 256 are the activation-quantisation kernels. Each launch costs CPU time to enqueue and GPU time to ramp up and drain.

How much? A micro-benchmark of back-to-back empty launches says 1.8 µs each. A real server measured graphs-off against graphs-on says otherwise: the production engine’s step went from 21.3 ms to 36.0 ms with graphs off, and 14 ms over 1,559 launches is 9.1 µs per launch, five times the micro-benchmark, because a real launch stream contends with itself rather than pipelining perfectly. The other production engine’s graphs-off penalty was 2.7 to 4.6×, confounded with the loss of its compiler fusions.

6.2 CUDA graphs

A CUDA graph records a sequence of kernel launches and their dependencies once, then replays the whole sequence with a single API call. The GPU’s front end walks the graph itself; the CPU is out of the loop. The engine measured the cost of one graph node at about 0.5 µs, so 677 nodes cost about 0.35 ms per token, 1.7 % of the step. Both production engines capture a graph for the batch-1 decode step, which is why their graphs-off penalty is a curiosity rather than their actual performance. The profiler shows the consequence: with graphs on, 93 % of the production engine’s wall time per token is GPU time inside graph replays, and only about 7 % is anything else.

The catch with a graph is that its kernel arguments are captured once. They can be updated between replays through the graph API, at a per-node cost each iteration, but the position, the input token and the sampling step change every step, so the engine takes the smaller route and makes those device-side: a small control block in GPU memory holds the step counter, the position, the current token and the stop flag, every kernel that needs them reads them from there, and a tiny finish kernel at the end of the step advances them. The graph is captured once at startup and replayed unchanged for every token. As an extra saving, the host keeps one replay in flight ahead of the token it is waiting for, so the GPU never idles between steps.

6.3 Getting the token back to the host

Each step produces one integer that the host needs (to detokenise and stream). The naive path, cudaMemcpy after each step, costs a synchronisation. The engine measured four alternatives on this discrete GPU:

The first won and is the default. Two points generalise beyond this engine. First, the whole term is 0.03 % of a 21 ms token; the hand-off is not where time goes at this model size, though it was a real lever on a smaller model where the token was 18× shorter. Second, the engine found that 188 blocks each polling a host-memory flag costs 145 µs per poll interval, versus 93 ns if one block reads it and mirrors it into device memory for the others. Host-mapped memory is for one reader, never for a grid.

6.4 The megakernel

If launch overhead and per-launch fixed cost are the problem, the extreme answer is one kernel: a persistent, cooperatively launched kernel with one block per SM that stays resident for the entire step (or the entire generation), executing the layers as a program of “instructions” and synchronising through counters in global memory instead of kernel boundaries. This is the megakernel design, and several research systems have demonstrated large wins with it, mostly on small models or on GPUs where the per-launch cost is a large fraction of a short step.

The engine built one: 188 blocks × 512 threads, executing a template of 355 instruction words (an embedding gather, 64 layers of projections, mixers and FFNs, the final norm, the head, the argmax, the hand-off), each block with its own slice of rows per instruction, every instruction a thin wrapper over the same device functions the discrete kernels use. Dependencies between instructions are counters in global memory incremented with integer red.add and waited on with acquire loads. It carries a multi-token dimension so that verifying several speculative drafts (Part 7.4) costs one weight pass. It was verified bit-identical to the discrete kernels at every stage of bring-up.

The measured result was parity: 21.30 ms per token against 21.22 ms for the same kernels under a CUDA graph, and in the composed engine 1.7 % slower. Its GEMV bodies ran at exactly the discrete kernel’s bandwidth. What it saved in launch cost it spent on its own synchronisation: counter waits, the serialisation of the GDN cores, and the skew at the end of the gate/up instruction where 188 producers wait for the slowest. At four speculative tokens per pass it was 17.7 % slower, because the multi-query attention instantiation spilled registers.

The conclusion is not that megakernels are ineffective. It is that the megakernel’s value has three sources, launch elimination, keeping loads in flight across instruction boundaries, and the in-kernel token loop, and on a 27 GB-per-token model on a GPU with cheap graph nodes, the first is already collected by the graph, the third is 0.03 % of the step, and the second has to beat the cost of 131 grid-wide joins per token (one per RMSNorm, whose mean square needs every SM’s slice of the residual). This outcome was anticipated as a possibility before the kernel was built, and the measurement confirmed it. An engineer reading a megakernel paper should ask: how long is one token on this hardware, and how much of it is launch overhead that a graph does not already remove?

6.5 Memory layout and the “nothing allocates after startup” rule

One more systems-level decision is worth naming. The engine allocates every device byte at startup: the weights, one fixed arena of per-step activations and scratch, and a session pool whose slots hold each conversation’s recurrent state and conv state, with a paged KV cache (256-token pages) drawn from one global free list. Session create, fork and evict are free-list operations, never allocations. The engine verified this with an interposer counting driver calls: zero allocations, zero frees, zero ioctls in a 1,000-token loop. The production engines also pre-allocate, but their allocators serve many concurrent requests and their page tables exist to serve batched attention kernels; at batch 1 with a fixed capacity most of that machinery reduces to constants.

Paged KV, for readers who have only heard the term. Attention needs to read every past key and value for the current sequence. Storing each sequence’s cache contiguously wastes memory (you must reserve the maximum length) and fragments it. Paging stores the cache in fixed-size pages indexed by a per-sequence page table, exactly like virtual memory, so pages are allocated as the sequence grows and freed when it ends. The attention kernel walks the page table. The engine keeps paging even at batch 1 because sessions are long-lived and have very different lengths.

Part 7. The agent loop: sessions, speculation, and grammars

Everything so far applies to any serving workload. This part is about the three mechanisms that make an agent workload different, and it is where the purpose-built engine concentrated its design effort. All three are also present in the production engines; the engine’s approach was to apply each of them at a finer granularity.

7.1 What an agent step costs

An agent step is: append the previous tool result to a conversation that already exists, run the new tokens through the model (incremental prefill), then decode until the model emits a stop token. A trajectory is the sequence of steps for one task. The metric is wall-clock per trajectory, which decomposes as, per step, T_prefill (submit to first token) plus T_decode (first token to stop token).

The workload used throughout this article is a recorded set of 13 agent trajectories: 82 steps, 11,239 emitted tokens, contexts from 1,500 to 34,000 tokens, reasoning effort from off to very high (Part 10 describes how it was recorded). Several facts about it shape everything below:

• 86.5 % of emitted tokens are unconstrained free text (reasoning traces, final answers, and string-typed tool-call arguments, which this model’s tool-call format renders as raw text).
• At every position where the grammar of a tool call actually restricts the next token, it restricts it to at most 382 candidates. There is nothing in between: the distribution of “how many tokens are allowed here” is bimodal with an empty middle.
• 9.1 % of tokens are forced: the grammar admits exactly one continuation, so the model’s opinion is irrelevant. No forced run is longer than 4 tokens.

7.2 Session-resident state and prefix caching

Both production engines keep a conversation’s state across turns through prefix caching: they hash blocks of the token stream, and when a new request shares a prefix with a cached one, they reuse the cached KV blocks instead of recomputing them. For a pure transformer that is a clean win at block granularity.

For a hybrid model it is not, because of the property from 3.6: a recurrent state cannot be sliced. You cannot take the state after 1,000 tokens and derive the state after 800. You can only resume from a state you saved. So both engines cache recurrent-state snapshots at boundaries and re-run the suffix. One of them, whose block size for this model is forced to about 800 tokens by the need to make attention pages and state pages the same size, resumes from that grid and re-prefills up to a full block per step; with speculative decoding enabled it backs off a further block, so a history under 1,600 tokens hits nothing. The other snapshots every 64 tokens during prefill and every 256 during decode and re-prefills from the nearest snapshot, at the cost of a 147 MiB state copy per hit.

The purpose-built engine keeps the live state of every session resident and checkpoints per agent step, so the only history re-processed per step is the tail since the last 64-token checkpoint, at most 63 tokens (5.4), plus the new tool-result tokens. Against the production engines’ own re-prefill, the saving is an estimate: pricing the tokens each baseline re-prefills on this workload at its own measured marginal prefill rate gives 1.0 to 1.2 % of trajectory time against the coarser-grained engine and 0.2 % against the finer one. Small, because re-prefill is cheap relative to decode and the workload has few steps per trajectory. Against not keeping sessions at all (re-prefilling the entire history each step, which is what a stateless server does) the effect was measured directly on a four-trajectory subset: with residency disabled the trajectories took 2.4× as long (229 s against 95 s), so residency removes about 59 % of the stateless time. The mechanism is real; the production engines already capture most of it.

Two more numbers systems engineers will want. A session’s fixed cost is 148 MiB (144 MiB recurrent state, 3.75 MiB conv state, page table), which equals the KV cache of about 2,200 tokens; below that context a session’s memory is mostly state, above it mostly KV. With 66 GiB free after weights, this GPU holds about 161 sessions of 4,096 tokens or 29 of 32,768. Forking a session (for a branch of an agent tree) costs one 148 MiB device-to-device copy, 0.23 ms, at any context, because the KV page table is shared copy-on-write. Restoring a parked session from host memory is estimated to beat re-prefilling it for any session longer than about 26 to 30 tokens; the restore cost is measured, while the re-prefill side of that estimate uses the production engines’ marginal prefill rates because this engine’s own rate was not yet available.

7.3 Where the time went

Before the two remaining mechanisms, the measured comparison of the three engines on this workload, because it frames why they matter. The method (Part 10): every engine is asked for exactly the recorded number of tokens per step; prefix caching, speculative decoding and grammar enforcement are enabled on all three. Two qualifications apply to every row. First, with grammar enforcement on, the production engines end a step when the grammar reaches a stop token, and on 10 of the 13 trajectories that happened before the recorded count, so only 3 trajectories have matching emitted-token counts across all three engines; the generated histories differ everywhere, since each engine’s own tokens are fed back. In a control run with grammar enforcement off, where all three emit exactly the recorded counts, the production engines total 121.1 s and 153.4 s against the same 255.6 s. Second, the purpose-built engine emits the recorded count regardless of where its grammar terminates, so its output is not always a complete, valid tool call: an independent grammar checker flagged 35 of its 82 steps on 11 trajectories in this mode, 7 that continued past the grammar’s end and 28 that ended before it. These figures are therefore fixed-count replay performance; they measure the cost of producing a given number of tokens under the grammar, not performance on equally valid, completed tool calls. Summed over the 13 trajectories:

The engine won decode by a quarter to more than a third and lost the trajectory by 2× because its prefill was 5 to 9 times slower. Everything in Parts 3, 4 and 6 about decode worked. Part 5’s two weak kernels (the prefill GEMM at 14 to 16 % of bandwidth and the borrowed attention fallback) dominated. A later prefill attention kernel and split-K GEMM brought the total to 178 s; Part 10 explains why that build was not adopted.

This is the central conclusion of the article: decode is one term in a sum. In a multi-turn workload the other terms are prefill and re-prefill, and they are compute-bound GEMM and attention problems with completely different kernel design constraints from the bandwidth-bound decode problem. An engine that is excellent at one and poor at the other loses.

7.4 Speculative decoding with the MTP head

This is the mechanism responsible for the engine’s decode advantage. Recall the floor: one token costs one pass over 26.9 GB of weights, because the tensor cores are idle. Speculative decoding exploits the idleness. If you could guess the next k tokens cheaply, you could run the model over all k + 1 positions in one pass, reading the weights once, and check which guesses were right. Every correct guess is a token you got for free. This is only possible because the model is deterministic given its inputs: the model’s output at position t + 3 given the guessed tokens at t + 1 and t + 2 is exactly what it would have produced sequentially, so verifying a chain of guesses is exact, not approximate.

Where do the guesses come from? This checkpoint ships its own drafter: a multi-token prediction head (MTP). It is one extra decoder layer, trained alongside the model, that takes the model’s final hidden state at position t and the embedding of the token at t + 1 and predicts the token at t + 2. To draft k tokens you run it k times in a chain, feeding each output back in. Its cost per draft step, arithmetic on the configuration, is 3.0 GB of weight reads, 11 % of a target step. But 84 % of that is the shared 2.54 GB lm_head, because the drafter has to produce a distribution over the full vocabulary too. The drafter’s own layer is 1.8 % of a target step.

The reference implementation, incidentally, does not implement the MTP head at all; it silently drops those tensors on load. Both production engines implement it. The engine had to write its own reference for the drafter from the two engines’ code and validate it by measuring acceptance: with the drafter wired as the production engines wire it, 85 % of its top-1 predictions agree with the target model’s; with the two halves of its input concatenated in the other order, 0 of 233. A wiring error in a drafter does not crash; it produces a drafter that is never right.

The engine’s speculation design, in five decisions:

1. Prune the draft head. Since 84 % of the draft cost is reading the head, read only the rows that matter. The chosen working vocabulary (this engine’s term for a pruned draft vocabulary) is the union of every token id that has appeared in the session so far and a static core of the 4,096 most useful tokens, rebuilt whenever it grows. A draft step over that head costs about 0.46 ms instead of 2.08 ms with the full head, and loses only 0.13 accepted tokens per round.
2. Verify k drafts in one pass. The target model runs over k + 1 rows (the current token plus k drafts) with the same weights read once: the GEMV kernel from Part 4 has an M ≤ 8 form whose per-row arithmetic is bit-identical to the M = 1 form, and the GDN kernel has a multi-step form that keeps the state in registers across the k + 1 recurrence steps, reading and writing the 3 MiB state once per layer instead of k + 1 times. Measured, verifying 8 rows costs 26.2 ms against 21.5 ms for one, a 22 % premium for 8× the work.
3. Accept greedily. The verify pass runs over the rows [current, d₁, …, dₖ]; row 0 predicts the token after the current one, row 1 the token after d₁, and so on. Draft dᵢ is therefore compared against the argmax of row i − 1. Accept drafts from the front as long as each matches; on the first mismatch at dᵢ, emit row i − 1’s own argmax (the “correction”) and stop. If all k match, row k supplies a free bonus token. For sampled decoding there is an exactness-preserving acceptance rule (accept draft d with probability min(1, p(d)/q(d)) where p is the target’s distribution and q the drafter’s, and on rejection sample from the normalised positive part of p − q), which the engine verified reproduces the target distribution to statistical precision.
4. Roll back the state exactly. This is the hybrid-model complication. The verify pass advanced the recurrent state by k + 1 steps; if only j drafts were accepted, the state must be the one after j + 1 steps. For the KV cache you move a cursor. For the recurrent state you cannot. Three solutions exist in the wild: write the state after every one of the k + 1 positions into its own slot and read the accepted one next step (cost: (k+1) × 147 MiB per in-flight request); snapshot to scratch and commit the accepted one (same order of memory); or keep a small replay ring (this engine’s term; state rollback by replay) of the raw per-step inputs (k, v, β, g per layer, about 1.5 MiB per step) and, on commit, re-run the recurrence from the frozen checkpoint over the accepted prefix. The engine chose the ring and fold, at 0.26 ms per commit with no per-session memory, and proved it bit-identical to sequential decode: zero differing bytes across all 48 states, the conv ring, 17 layers of KV and the position, for every accept count, at eight positions including page and partition boundaries. The drafter is a full-attention layer and never touches the recurrent state, which is the fact that makes all of this tractable.
5. Choose k by measurement. Time per emitted token is (T_verify(k) + k × T_draft) / tokens_per_round(k). Acceptance rises with k but with diminishing returns; verify cost rises slowly. The measured optimum was k = 7: 3.7 drafts accepted per round on average, 4.7 tokens per round, 6.5 ms per emitted token against 21.3 ms without speculation, a 3.3× decode speed-up. The production engines default to a chain of 3 to 4 drafts and accept 3.0 to 3.9 tokens per verify pass.

Acceptance depends on content. Inside a tool call the drafter is almost always right (4.5 of 8 accepted, 99 % at least one); in reasoning text it is weakest (3.7 of 8). A speculative engine’s throughput is therefore a property of the workload, and any benchmark that reports one number for it is hiding a distribution.

7.5 Grammar-constrained decoding

When the model emits a tool call, its output must parse. Constrained decoding guarantees it: a grammar (here derived from the tool schemas and the model’s tool-call format) is compiled into a matcher, and before each token the matcher produces the set of token ids that can legally come next. The sampler is restricted to that set. Both production engines use the same open-source grammar library and represent the allowed set as a bitmask: 248,320 bits, 31 KB, built on the CPU each step, copied to the GPU, and applied by writing −∞ into every disallowed logit before the argmax.

One observation about that pipeline: the mask is never used to reduce work. The 2.54 GB head is read in full, the −∞ scatter touches all 248,320 logits, and the argmax scans all of them, even when the grammar admits three tokens. Three full-vocabulary passes per token, for a decision that was already made.

The engine treats the grammar as a bandwidth optimisation with three separable mechanisms:

1. The mask as a filter. Unavoidable, and a cost: building the mask on the CPU (measured at 2.7 µs per position in C++, far under the decode step it overlaps) and honouring it inside the head kernel (13 µs, 0.07 % of a token).
2. The masked head (this engine’s term for a restricted vocabulary projection). When the admissible set is small, build a sorted row list from the bitmask on the GPU (a 7,760-word popcount scan, negligible) and run the head kernel over only those rows of the lm_head. At 382 rows that is 3.9 MB instead of 2.54 GB, 90× cheaper. The head kernel was designed from the start over row sets rather than row ranges for exactly this.
3. Forced spans. The grammar library exposes a “jump-forward” query: the longest byte string every legal continuation must begin with. Neither production engine calls it, so every forced token costs them a full decode step. The engine tokenises the forced bytes canonically and ingests them as known rows of the next verify pass (the same multi-row pass speculation uses), so a 4-token forced span plus the next real token costs one weight pass with no head read and no sampling for the forced positions. The tokens still traverse all 64 layers to update the state; this is a prefill of known tokens, not free.

Two correctness subtleties, both measured rather than reasoned. First, the forced-span invariant is byte identity, not token identity: ingesting the canonical tokens of = followed by later generating start yields different ids from the merged token the model might have produced for =start, and that is inherent. The engine’s check is that the bytes match, that each ingested token is accepted by an independent matcher, and that the tokens after the span are unchanged; the last of these held on all 359 recorded spans, but it is an observed result on this workload, not a guarantee. A different tokenisation of the same bytes changes the model’s input ids and can change its subsequent logits. Second, the last token of a forced span may merge with the model’s next free token. The results reported here use the recording’s rule and ingest all canonical tokens of the span (keep_last); the conservative alternative that ingests all but the last token and lets the model generate it, and a variant that skips one-token spans, were built and measured at the component level but are not the benchmarked configuration.

Measured at the component level on the token stream, without speculation, the three mechanisms remove 1.0 to 1.1 % of trajectory time and 52 % of head traffic inside a tool call. In the integrated engine with speculation on, the result is different, and it is the one that counts: switching the whole grammar lever off made the 13 trajectories 1.3 % faster (252.4 s against 255.6 s). The restricted vocabulary projection alone still saved 0.3 %, but forced spans cost 1.0 %, because a forced span interrupts a speculative round and the pass that ingests it emits fewer tokens than a full round would have. Two conclusions follow. The head-traffic saving is real and bounded by the bimodal statistic from 7.1: the grammar binds on only 4.4 % of tokens, and nothing about a grammar reduces the 63.6 % of every token that is the feed-forward network. And a mechanism that wins in isolation can lose in composition; Part 10 discusses why this one was measured both ways.

Part 8. Correctness: what “the engine is right” means

Every kernel above was admitted into the engine only after a correctness result from the same build, under one rule: a faster wrong kernel is a rejected kernel. This part is about what “right” means in practice, which is less straightforward than it appears.

8.1 The oracle

The oracle or reference is the implementation whose outputs define correctness: here, the model’s published PyTorch code run in the standard transformers library, with weights dequantised to bf16 up front (the more self-consistent of the two ways to load the checkpoint, as 4.1 explained). Goldens are recorded from it: for a set of prompts, the top-32 logits and their ids at every position, the full 248,320-wide logit vector at three positions per prompt, and the recurrent state, conv state and KV cache at chosen positions.

Two things about the oracle you need to know before comparing anything against it:

• It disagrees with itself. Its prompt-processing path and its generation path are different numerical paths through the same weights (5.4). Measured over 3,072 positions, they disagree on the top-1 token at 31 (98.99 % agreement) and differ in logits by up to 7.25. A uniform criterion tighter than that cannot be required of an implementation whose summation order differs from the oracle’s, which is every implementation not built on the oracle’s own kernels; tighter bounds can be applied at positions where the implementation is insensitive to order.
• Its logits are bf16, so it has ties, and 4.9 % of positions are within one ulp of one. An engine that differs in accumulation order will flip some of those ties. Whether a flip is a bug or a tie is decidable only if you look at the margin.

8.2 Teacher forcing

The primary comparison technique is teacher forcing: feed the engine the golden token at every position regardless of what it would have chosen, and compare its logits and its argmax at every position against the golden. This decouples the comparison from divergence: without it, a single early tie flip sends the engine down a different text and every later position is incomparable. Free-running greedy generation is compared too, but scored per prompt, with divergence allowed only at positions the oracle’s own margin explains.

The tolerance envelope is calibrated on the oracle’s own inconsistency: per-position logit differences bounded by what the oracle’s two paths disagree by, top-1 agreement at least the oracle’s own 98.99 %, per-prompt percentile and mean bounds, and a hard requirement of exact token match on a sanity prompt. The engine’s discrete-kernel path was evaluated at 4,546 positions over 24 prompts and scored 99.19 % raw top-1 agreement on the 2,607 same-path positions that the criterion applies to. That figure, and every oracle gate in this article, was run with the engine in its bf16 recurrent-state compatibility mode (3.6), because the oracle keeps its state in bf16 during generation and the two cannot otherwise be compared position by position. The fp32-state configuration that the performance numbers use was validated differently: its recurrent state is compared against the oracle’s after rounding to bf16, and its decode, prefill and speculation paths are held bit-identical to each other by the checks of 8.3. Agreement between the engine’s own paths in fp32 does not transfer the bf16 oracle-gate statistics to that configuration; it shows that the fp32 configuration is internally consistent and that its state stays within the oracle’s envelope. Where the engine’s own two prefill compositions disagree with each other by more than the bound at a position, that position is judged against the engine’s measured self-variation rather than the fixed bound, and the number of such positions left unexplained is reported (zero for the shipped configuration). It is calibrated on prompts of the workload’s length, because the envelope widens three to six times between a 50-token prompt and a 420-token tool-calling prompt, and an engine certified on short prompts would be certified against the wrong tolerance.

8.3 Bit identity where it is possible

Against the oracle, exactness is not achievable without reproducing its summation order, which this engine does not attempt. But between two paths of the same engine, exactness is a choice, and the engine makes it everywhere it can, because a bit-identity check is unambiguous where a tolerance is a judgement:

Each row required a design decision earlier in the article: the fixed attention partition count, the fixed-order split-K reduce, the checkpoint-anchored chunking, the register-resident multi-step recurrence, the monotonic counters. None of them happens by accident.

8.4 Negative controls

A test that cannot fail proves nothing. The engine keeps a build with a deliberately wrong RoPE base frequency and requires the correctness checker to reject it. It does, with nine argmax failures. When a checker is calibrated with a tolerance derived from the oracle’s own noise, a negative control is the minimum evidence that the tolerance rejects a gross error; it does not show that every bound is as tight as it could be.

Part 9. Profiling an inference engine

A short, practical section on what to measure, since the audience knows how to run a profiler and the question is what to look at.

Bandwidth utilisation per kernel, not FLOP/s. For every decode kernel the question is bytes moved divided by kernel time, as a fraction of measured (not theoretical) read bandwidth. The engine’s per-kernel figures: lm_head 98.5 %, the GEMV family 88.7 %, attention 71 % at 32K context, the GDN step 48 %. Those numbers immediately rank where the remaining time is.

Fixed cost per call. When a kernel’s bytes-per-second looks bad, check whether it is latency (too few independent memory streams; fix with split-K or more blocks) or a fixed per-launch cost (fix by fusing or by processing more work per launch). The GDN kernel’s marginal byte moved at 99.7 % of bandwidth; the 48 % came from about 5 µs of fixed cost per launch times 48 launches.

Occupancy as a threshold, not a gradient. On this GPU, every configuration with at least 512 resident threads per SM landed within 0.3 % of each other; 256 threads lost 8 %; 128 threads collapsed to 55 %. The binding constraint was memory-level parallelism (enough loads in flight to cover DRAM latency), and it saturates. Sweep occupancy once per GPU, then stop.

Launch count, with graphs off. The profiling path used here (the CUDA activity API) reports a graph replay as one unit rather than per kernel; node-level graph tracing exists in Nsight Systems but was not used. The practical method is to measure the launch inventory with graphs disabled and the wall time with graphs enabled. The gap between the two is the launch term the graph is saving you, and it tells you how much a megakernel could possibly recover.

The GPU-time fraction. Wall time per token minus GPU time per token is the host, the scheduler, the hand-off, and gaps. For the production engines at batch 1 it was 7 %. For a batch-1 engine anything above a few percent is a bug or a design flaw.

Cache effects in benchmarks. With a 128 MiB L2, a benchmark that calls one kernel repeatedly on the same operands measures L2, not DRAM. Every kernel benchmark in this engine rotates through disjoint operand copies larger than twice the L2, except the lm_head, whose 2.54 GB evicts itself. A “flush” buffer equal to the L2 size does not flush; the first bandwidth probe for this engine made exactly that mistake.

Published claims are hypotheses. Two published notes about this GPU’s architecture (that a particular cache-bypass load modifier bypasses L2, and that floating-point atomic adds are cheap counters) were tested and both were wrong on this hardware: the modifier did nothing, and float atomics under contention were 77× slower than integer ones because integer adds aggregate within a warp and float adds cannot. Measure before relying on such claims.

Part 10. Building an engine for one model, one GPU and one workload

Parts 1 to 9 used a purpose-built engine as their source of examples. This part describes what that engine is, how it was produced, how it was compared against the production engines, and what the process taught. Readers interested only in how inference works can stop at Part 9.

10.1 What a specific engine is

The production engines are general: one runtime serves hundreds of model architectures on a dozen GPU generations, with dispatch on dtype, architecture and hardware at every layer. The engine in this article is the opposite: a standalone CUDA/C++ binary with no Python runtime that contains only the code paths this one model executes on this one GPU for this one workload. There is one quantisation format, one attention configuration, one GDN configuration, one norm implementation per norm shape, and no dispatch on any of them. Everything the model does not execute is absent.

The premise is that specialisation buys two things. First, decisions that a general engine makes at run time (which kernel, which tile, which cache layout) are made once, by measurement, and fixed. Second, mechanisms that a general engine cannot afford because they only pay off for one workload (per-step session residency, a pruned draft vocabulary, a masked vocabulary projection, forced spans) can be built in.

10.2 How it was produced

The engine was generated in stages, each producing a document with its evidence before the next stage began:

1. Model specification. Every op, shape, dtype and rounding point derived from the checkpoint’s configuration and the reference PyTorch code (Parts 3 and 4.1); the byte budget per token (Part 1); the reference’s self-consistency measured to set a tolerance (Part 8).
2. Hardware specification. Bandwidth, launch and graph-node costs, barrier costs, L2 behaviour and occupancy thresholds measured on the target GPU (Parts 1, 6, 9), with published architectural claims tested rather than assumed.
3. Reference-engine analysis. The production engines’ source read to find which kernels they dispatch for this model on this GPU, how they manage hybrid state and prefix caches, how they speculate and roll back, and how they enforce grammars (Parts 4.2, 7.2, 7.4, 7.5), and a launch inventory per token derived from it (Part 6).
4. Workload recording. Thirteen agent trajectories were recorded by running the model against a fixed set of deterministic tools, with every tool result frozen so that a replay is reproducible. Each step boundary was checked to be byte-identical to a full re-render of the chat template (Part 2.3), and the grammar’s admissible-token set was recorded at every position (Part 7.1).
5. Component synthesis. One kernel family at a time (Parts 3 to 5), each with a specification, a reference op, a verification against it, a measured sweep of launch configurations, and a record of every variant that lost and by how much.
6. Composition. The components assembled into the discrete-kernel engine under a CUDA graph, the megakernel variant, and the agent engine with its levers (Parts 6 and 7); every path verified against the oracle (Part 8) before any number was recorded.
7. Benchmark and optimisation. A like-for-like comparison against the production engines, then a profile-driven loop: rank kernel time, change one thing, verify, measure, keep or revert.

10.3 How the comparison was run

The comparison in Part 7.3 follows a fixed method. Both production engines ran with the features that make them fast on this workload enabled: prefix caching, speculative decoding with the model’s own MTP head, and grammar-constrained output. They were fed token ids rather than text, because the recorded id stream is not always the canonical tokenisation of its own text (Part 2.2) and re-tokenising would change the input. Every engine was asked for exactly the recorded number of tokens per step; with grammar enforcement on, the production engines stopped early at grammar stop tokens on 10 of 13 trajectories, while the purpose-built engine ran to the recorded count whether or not its grammar had terminated (35 of 82 steps flagged by an independent checker), so a grammar-off control with exact counts was run alongside and the results are reported as fixed-count replay performance (Part 7.3). Each trajectory was replayed three times from a flushed cache and the median taken. Clocks were taken at the API boundary: submit time, first-token arrival and stop-token arrival per step. The two production engines are strong baselines: one of them autotuned its fp8 GEMM on this model’s decode shapes, and both reached about 82 % of the measured bandwidth on plain decode before speculation.

10.4 What the build taught

Three findings from the process generalise beyond this engine:

• Constants that survive a port. Kernels carried from an earlier engine for a smaller model of the same family contained ten literal 4s that meant “the number of query heads per KV head” (now 6), a norm prologue hard-wired to a 1,024-wide hidden size (now 5,120), and a GDN channel-offset computation that read the wrong key head for 47 of 48 value heads while staying in bounds and producing plausible output. All compiled. All would have passed a shape check. The response was a rule: components are never copied between engines; a prior engine is studied and its decisions re-derived.

• Verification bugs look like kernel bugs. In the speculation component, more of the stoppages on the way to a passing result were in the checker than in the kernels (a stale page-table entry written before the page was allocated; a reference that compared the wrong intermediate). When a check fails, suspect it with the same weight as the thing it checks.

• The faster kernel that could not ship. The prefill attention kernel and split-K GEMM of Part 5 brought the trajectory total from 256 s to 178 s, but the build containing them produced one unexplained token divergence on a 4,096-token prompt and one state element outside tolerance in the correctness gate. Under the rule that a change is adopted only if the gate passes, it was not, and Part 7.3 reports the slower, verified configuration. Whether that rule is too strict is a legitimate question; that it was applied is what makes the numbers trustworthy.

• Component measurements do not compose. The grammar mechanisms of Part 7.5 saved head traffic in isolation and cost time in the integrated engine, because forced spans interrupt speculative rounds. Every mechanism was therefore measured twice: alone against its component reference, and switched off one at a time in the composed engine with everything else on.

The single most useful habit of the process is to state, for every number, whether it was measured (and where), derived (and from what), or estimated (and to flag it as such). The second is to record rejected variants with their numbers, because a variant rejected for a reason that no longer holds on new hardware is free performance for the next engine. The third is that correctness is a first-class deliverable with its own artefacts, not a checkbox.

Glossary

Quick reference for the terms expanded in the text, in roughly the order they appeared.