The Inference Engineering Course
Glossary
A compact reference to the course vocabulary, with chapter links for the full definitions and assumptions. This is a maintained glossary, not an automatically proven exhaustive term inventory.
Coverage
Organized across all fourteen parts; definitions are scoped to the stated model or source pin. Where a term means different things in the two engines, both names are given — the vocabulary divergence is itself worth knowing. Where a term names something that is absent — a capacity factor at inference, worker-level fault recovery — the row says so, because knowing what an engine deliberately does not have is as load-bearing as knowing what it does.
§1
The core loop and scheduling
| Term | One line | Chapter |
|---|---|---|
| Prefill | The forward pass over new prompt tokens, possibly chunked or extending cached context. It produces logits and reusable state; compute-bound behavior depends on shape, context, kernels and hardware. | 01-01 |
| Decode | Autoregressive continuation, conventionally one new input token per sequence per pass. Small-batch decode is often bandwidth- or launch-sensitive; speculation verifies multiple positions and large batches can be compute-bound. | 01-01 |
| Extend (SGLang) | SGLang's name for prefill — "continue a sequence whose prefix may already be cached", of which a cold prefill is the special case. | 00-02 |
| ForwardMode | SGLang's enum naming EXTEND / DECODE / MIXED / IDLE; the in-code name for the two regimes. | 00-02 |
| MIXED batch | One forward pass carrying both prefill-chunk tokens and one-token decode rows. | 01-01 |
| Uniform decode | vLLM's notion of a pure decode batch — every request scheduled exactly 1 + num_spec_tokens tokens; gates dense CUDA-graph dispatch. | 01-01 |
| Head-of-line blocking | One long prefill monopolising the GPU so every decoding stream stalls for its full duration. | 01-01 |
| Static batching | A fixed request set batched for its whole lifetime; runs until the slowest member finishes. | 01-03 |
| Dynamic batching | Forming a batch at request time with a size trigger and timeout; fixes batch-boundary padding only. | 01-03 |
| Continuous batching | Admitting and retiring requests at every iteration (Orca, OSDI 2022). Also called iteration-level scheduling. | 01-03 |
| Selective batching | Orca's split: batch per-token ops on a flattened token axis, run attention per sequence because KV lengths differ. | 01-03 |
| Persistent batch | vLLM's slot-indexed InputBatch, preallocated at max_num_seqs rows and mutated in place. | 01-03 |
condense() | vLLM's in-place compaction sliding the highest occupied slot into the lowest hole, keeping the batch a dense prefix. | 01-03 |
filter_batch / merge_batch | SGLang's gather-based removal and torch.cat-based admission on ScheduleBatch. | 01-03 |
| Straggler / padding tax | The row-steps a static batch spends on rows that have already finished. | 01-03 |
| Preemption (vLLM) / retraction (SGLang) | Evicting a running request's KV to reclaim capacity, returning it to the waiting queue. Both engines resume by recompute. | 01-04 |
| Watermark | The fraction of KV blocks kept free that new admissions — not running requests — must clear. | 01-04 |
| PrefillAdder | SGLang's stateful admission-control object; accumulates a budget and returns CONTINUE / NO_TOKEN / OTHER, distinguishing KV-capacity refusal from token-budget refusal. | 01-04 |
new_token_ratio | SGLang's feedback-controlled pessimism factor on how much of max_new_tokens each running request will actually generate. | 01-04 |
| KV survival fraction (f) | Share of a preemption victim's tokens still recoverable from the prefix cache on resume. | 01-04 |
| LPM | SGLang's longest-prefix-match admission ordering. Not the default — and it silently self-disables above 128 queued requests. | 02-04 |
| Token budget | The per-iteration cap on scheduled query positions that prefill and decode tokens both draw from. | 01-05 |
| Chunk | The slice of a prompt's remaining prefill admitted in one pass — the residue of the budget after decode tokens are placed. | 01-05 |
| Piggybacking | Decode tokens riding along with a prefill chunk, reusing its already-paid weight stream. | 01-05 |
chunked_req | SGLang's single-slot handle on the one in-flight partially-prefilled request. | 01-05 |
| P/D disaggregation | Running prefill and decode on separate machines and shipping the KV cache between them. | 01-06 |
| Bootstrap room | A 63-bit random correlation id the router injects into both the P and D copies of a request; the rendezvous key for the transfer. | 01-06 |
| Bootstrap server | Service on the prefill node publishing rank→(ip, port) so decode workers can find their peer. | 01-06 |
| KVPoll | SGLang's five-state transfer status: Failed, Bootstrapping, WaitingForInput, Transferring, Success. | 01-06 |
| KV connector | vLLM's plugin interface (scheduler half + worker half) carrying disaggregation, CPU offload, and remote KV stores through one registry. | 01-06 |
| Hetero-TP transfer | KV remapping when prefill and decode run at different tensor-parallel degrees. | 01-06 |
§2
Latency and measurement
| Term | One line | Chapter |
|---|---|---|
| TTFT | Time to first token: arrival to first streamed token, queueing included. | 01-02 |
| TPOT | (t_N − t_1)/(N−1); one scalar per request. A rival definition using E2E/N exists — say which you mean. | 01-02 |
| ITL | Inter-token latency, t_k − t_{k−1}. A distribution, not a scalar. vLLM samples per streamed chunk, SGLang per token — not comparable across harnesses. | 01-02 |
| E2EL | Arrival to last token (vLLM) or to stream close (SGLang — it includes the terminal SSE frame). | 01-02 |
| Goodput | Completed requests satisfying every SLO predicate, per second. A request served past its SLO counts as zero, not a fraction. The cited harness support is version-specific; the metric can be computed from request outcomes independently of engine. | 01-02 |
| Burstiness | vLLM's Gamma shape parameter on inter-arrival times; 1.0 is Poisson. SGLang offers exponential only. | 01-02 |
| Bucket floor | The lowest histogram edge, below which histogram_quantile carries no information. vLLM's ITL floor of 10 ms sits above a healthy 8B decode rate. | 01-02 |
| Batch invariance | The property that a prompt yields identical tokens regardless of what else is in the batch. Broken by split-K reduction order. | 10-04 |
§3
Memory and the KV cache
| Term | One line | Chapter |
|---|---|---|
| KV cache | Per-layer key/value state. For a growing autoregressive prefix, causal masking preserves earlier states under unchanged model, positions and inputs; fixed encoder/cross-attention state is another reusable case. | 00-02 |
| Cacheability precondition | Reuse requires the cached computation's inputs, weights, positions and mask semantics to remain unchanged. Extending an unrestricted bidirectional segment changes old states; an unchanged encoder input or cross-attention memory can still be cached. | 00-02 |
| Query length vs sequence length | The two independent quantities in every attention-metadata struct; equal in prefill, 1 : s in decode. | 00-02 |
| Cell size | Bytes of KV one token costs across all layers, 2·L·h_kv·d_h·b. SGLang's own term. | 02-01 |
| Page size (KV) | Bytes of one block of one layer; vLLM's AttentionSpec.page_size_bytes. | 02-01 |
| Resident tokens | Sum of all in-flight sequence lengths — the quantity KV capacity is actually linear in. | 02-01 |
| Latent cache | MLA's single per-token-per-layer vector of width kv_lora_rank + qk_rope_head_dim, replacing per-head K and V. | 02-01 |
| PagedAttention | Storing KV in fixed-size blocks addressed through a per-sequence block table. Now a memory-management design, not a kernel — vLLM deleted the CUDA kernel in PR #47361. | 02-02 |
| Block pool | The fixed set of physical KV blocks plus its free list; the allocator's arena. | 02-02 |
| Block table | Per-sequence array mapping logical block index to physical block id. | 02-02 |
| Slot mapping | Per-scheduled-token index into block-major KV storage saying where this step's K/V is written. | 02-02 |
kernel_block_size | The block size the attention kernel sees — a divisor-subdivision of the allocator's block size. | 03-03 |
| Internal fragmentation (KV) | Unwritten slots in a sequence's last block; bounded by block_size − 1 tokens per sequence. | 02-02 |
| Over-reservation | Capacity held because output length is unknown at admission — the dominant waste under contiguous allocation. | 02-02 |
| Null block | Block id 0, popped at construction and never cached or freed; a placeholder for skipped windows. | 02-02 |
| Copy-on-write | In vLLM V1, not sequence forking — the only CoW path fires on a partial prefix-cache hit whose boundary lands inside a block. | 02-02 |
| Prefix caching | Reusing already-computed KV for a prompt prefix a previous request processed. | 02-03 |
| Block hash chain | h_i = H(h_{i−1}, t_i, e_i) — a block's key fingerprints its whole prefix, not just its own tokens. | 02-03 |
NONE_HASH | The chain's initialisation vector; random per process for xxhash, fixed for cryptographic algorithms. | 02-03 |
| Extra keys | Non-token components of a block hash: LoRA name, multimodal id + offset, cache salt, prompt-embedding digest. The correctness boundary of the feature. | 02-03 |
cache_salt | Per-request string keyed into block 0; the chain propagates isolation to the whole request. | 02-03 |
| Alignment sensitivity | Boundary rounding loses fewer than one block from an otherwise identical prefix. Inserting or changing an early token instead changes chained keys for the entire suffix; that loss is not bounded by one block. | 02-03 |
| RadixAttention | SGLang's prefix cache: a compressed trie whose edges carry token sequences, giving longest-prefix matching and tree-aware eviction. | 02-04 |
| Edge splitting | Splitting one edge into parent + child when a new key diverges mid-edge. Clones the index tensor; never moves KV. | 02-04 |
lock_ref | Per-node refcount; locking a node locks the whole path to root and moves its tokens from evictable to protected. | 02-04 |
| Evictable leaf set | SGLang's structurally maintained set excluding interior and locked nodes. Leaf-first eviction is an invariant here and a convention in vLLM. | 02-04 |
| In-batch prefix caching | A throwaway simulated tree over the waiting queue that deprioritises requests colliding on an uncached prefix. | 02-04 |
UnifiedRadixCache | SGLang's default cache at this SHA. HiRadixCache is defined but never constructed; HiCache is enabled via init_hicache(). | 02-04 |
ChunkCache | The no-reuse cache implementation used when radix caching is off. | 02-04 |
| Crossover N* | Resident-token count where KV traffic per decode step equals weight traffic — where KV quantization becomes a latency lever. | 02-05 |
| Randomized Hadamard Transform | HDx: sign-flip then Hadamard, gaussianising a vector so a scalar quantizer becomes near-optimal and outliers vanish. | 02-05 |
| Lloyd-Max quantizer | MSE-optimal scalar quantizer for a known density; boundaries at centroid midpoints, centroids at cell conditional means. | 02-05 |
| Scale steganography | Hiding a 4-bit zero-point in the low mantissa bits of an fp32 scale. | 02-05 |
mem_fraction_static | SGLang's static-allocation fraction — the complement of a heuristic reserve, not vLLM's profiled cap. | 02-06 |
| Write-through / write-back | HiCache host-write policies, distinguished by a threshold and whether staging happens on the eviction path. | 02-06 |
| L1 / L2 / L3 | SGLang's names for the GPU pool, host pool, and storage backend. | 02-06 |
| Lazy offload | vLLM's policy of cursor-walking the GPU free queue and offloading only blocks near eviction. | 02-06 |
§4
GPU architecture and kernels
| Term | One line | Chapter |
|---|---|---|
| SM | Streaming multiprocessor — the GPU's independent scheduling unit. 132 on H100 SXM5. | 00-03 |
| Warp | 32 threads issuing one instruction together; the real unit of execution. | 00-03 |
| Occupancy | Resident warps over the maximum. A diagnostic, not a target. | 00-03 |
| Warp shuffle | Direct register exchange between lanes, with no memory traffic. | 00-03 |
| mma / wgmma | Tensor-core matrix-multiply-accumulate; Hopper's warp-group form reads operands straight from shared memory. | 00-03 |
| Kernel launch overhead | Per-launch CPU/driver cost, microsecond-scale, independent of the GPU work done. | 00-03 |
| Arithmetic intensity | FLOPs per byte of compulsory HBM traffic. | 00-04 |
| Compulsory traffic | Bytes that must cross the HBM boundary at least once; excludes cache re-reads, so it lower-bounds real traffic. | 00-04 |
| Roofline / ridge point | min(π, β·I), with the ridge at I* = π/β — 295 FLOP/byte on H100 SXM bf16. | 00-04 |
| Decode step floor | t ≳ (weight bytes + KV bytes)/β — 4.48 ms, 223 tok/s at batch 1 for Llama-3-8B bf16. | 00-04 |
| Online softmax | Softmax in one streaming pass via a running max and running sum, corrected whenever the max increases. | 03-01 |
| Rescaling factor (α) | exp(m_old − m_new), in (0, 1] for finite nonempty states in real arithmetic; initialization and underflow can yield zero; retro-normalises a partial state when a larger logit arrives. | 03-01 |
| Log-sum-exp (LSE) | m + log ℓ — the single fp32 scalar per (token, head) both engines pass between attention kernels. | 03-01 |
| Merge attention states | Combining partial attention results over disjoint key ranges. vLLM's merge_attn_states is a two-way merge; the N-way split-K reduction is Triton's reduce_segments. | 03-01 |
| Softcapping | c·tanh(x/c) applied to logits before softmax. | 03-01 |
| Tile (attention) | The B_r × d_h query slab and B_c × d_h KV slab held in SRAM for one loop iteration. | 03-02 |
| Split-Q vs split-K partitioning | Whether the query or key axis is divided across warps in a block. Split-Q (FA2) lets each warp own complete rows and removes a shared-memory reduction. | 03-02 |
| Warp specialisation | Dedicating whole warpgroups to roles — a producer issuing TMA copies, consumers running wgmma and softmax. | 03-02 |
| Ping-pong scheduling | Offsetting two consumer warpgroups by half a stage so one is in a GEMM while the other is in the SFU-bound softmax. | 03-02 |
| TMA | Hopper's hardware bulk-copy engine; one thread submits a descriptor and hardware handles addressing. | 03-02 |
| mbarrier | Shared-memory barrier counting bytes arrived rather than threads arrived; the producer/consumer handshake. | 03-02 |
| Incoherent processing | Multiplying Q and K by a random orthogonal matrix to spread FP8 outliers, leaving QKᵀ unchanged. | 03-02 |
varlen / cu_seqlens | Ragged-batch convention: requests concatenated into one axis with prefix-sum offsets, no padding. | 03-02 |
| Split-K (decode) | Partitioning the KV axis across thread blocks so each computes a partial softmax, combined by a second kernel. | 03-03 |
| FlashDecoding | The decode-shaped application of split-K: query length 1, KV split to fill the SMs. | 03-03 |
| Partial attention state | The triple (Õ, m, ℓ) — or (Õ, LSE) — a split emits before normalisation. | 03-03 |
| Attention backend | A class satisfying the engine's backend ABC. An adapter to an external kernel library, not a kernel. | 03-04 |
| Metadata builder | Per-attention-group object whose build() runs once per step, producing metadata every layer then reads. | 03-04 |
| Attention group | Layers sharing a backend, a KV spec, and a per-rank Q-head count — the unit of one metadata build. | 03-04 |
AttentionCGSupport | Four-level enum declaring how much CUDA-graph capture a builder tolerates. | 03-04 |
| plan/run split | FlashInfer's two-phase API: host-side layout and scheduling in plan(), launch in run(). | 03-04 |
| CuteDSL | CUTLASS's Python DSL; FlashAttention 4's implementation language, JIT-compiled rather than shipped compiled. | 03-04 |
| Group size (g) | Query heads per KV head, h/h_kv — the one knob that both divides KV bytes and multiplies decode intensity. | 03-05 |
| KV head replication | What max(1, h_kv // tp_size) does past TP = h_kv: extra ranks hold duplicates, so cluster KV capacity stops growing. | 03-05 |
| DP attention | Data-parallel attention with tensor-parallel FFN — the mode that makes MLA's replicated caches non-redundant. | 03-05 |
| Rotary dimension (r) | The leading slice of the head dimension RoPE rotates; the tail passes through untouched. | 03-06 |
| NeoX vs GPT-J pairing | The two conventions for which elements form a rotation pair; checkpoints permute QK weights to match. | 03-06 |
| Correction range | YaRN's band of pair indices between pure extrapolation and pure interpolation. | 03-06 |
| mscale (μ) | YaRN's attention-temperature factor 0.1·ln s + 1, scaling every logit by μ². | 03-06 |
| Tombstone node | An SGLang radix node whose sliding-window KV has been freed while its full-layer KV remains matchable. | 03-06 |
| Attention sink | The first few positions, which absorb disproportionate attention mass; pinning them is what makes windowed attention stable. | 03-06 |
§5
Numerics and quantization
| Term | One line | Chapter |
|---|---|---|
| Dynamic range | Ratio of largest to smallest representable magnitude; set by exponent bits alone. | 00-05 |
| Precision / machine epsilon | Gap between 1.0 and the next representable value, 2^−m; set by mantissa bits alone. | 00-05 |
| Accumulation dtype | The dtype used for partial sums, distinct from operand storage and output dtype. Many low-precision floating-point paths accumulate in FP32; integer and reduced-precision paths have different contracts. Inspect the selected instruction and kernel. | 00-05 |
| Affine quantization | x ≈ s(q − z); symmetric when z = 0. | 00-05 |
| Group shape | Scale-sharing granularity — per-tensor, per-token, per-channel, per-group. vLLM encodes it as a two-integer GroupShape. | 00-05 |
| E8M0 | 8-bit exponent-only shared scale used by OCP MX formats; a bare power of two. | 00-05 |
| MXFP4 / NVFP4 | Block-scaled 4-bit floats: block 32 with an E8M0 scale (4.25 bits/wt) versus block 16 with an E4M3 scale (4.5). | 00-05 |
| Saturating clamp | Out-of-range values pin to max rather than becoming inf — the reason FP8 failures are silent. | 00-05 |
| Effective bits per weight | Stored bits plus amortised scale metadata, b + b_s/g, with a further + b_z/g for an asymmetric scheme's zero-points. 4.125 for int4 at group 128. | 04-01 |
| QuantKey | vLLM's frozen dataclass identifying storage dtype, scale descriptor, second-level scale, and symmetry. | 04-01 |
| Static vs dynamic activation scale | Frozen at calibration versus recomputed per forward pass. | 04-01 |
| Online quantization | Round-to-nearest weight quantization at load time from a bf16 checkpoint. | 04-01 |
| Requantization | Converting an already-quantized checkpoint to a different format at load. | 04-01 |
§6
Decoding and constraints
| Term | One line | Chapter |
|---|---|---|
| Speculative decoding | Draft cheaply, verify in one target pass, accept a prefix — provably without changing the output distribution. | 06-02 |
| Acceptance rate / length | Fraction of drafted tokens accepted, and mean tokens emitted per target pass — the quantity deciding whether speculation pays. | 06-02 |
| Tree attention | Verifying a branching draft tree in one pass with a mask encoding the tree's ancestry. | 06-04 |
| EAGLE | Drafting from the target model's own hidden states with a small head, rather than a separate draft model. | 06-04 |
| Structured decoding | Masking logits each step so only grammar-permitted tokens can be sampled. | 06-05 |
| Jump-forward decoding | Emitting tokens the grammar forces without running the model at all. Dead code at both pinned SHAs — the SGLang scheduler driver was removed in commit 935cda944b (#4032) and the outlines backend now hard-codes jump_forward_map = None; vLLM never had it. Taught as mechanism, flagged as not running. | 06-05 |
| Neutral value encoding | Storing a disabled sampling parameter as the value that is a no-op — k = V, p = 1, ρ = 1, α = 0 — so one kernel runs unconditionally over every row with no branch on request identity. | 06-01 |
TOP_K_ALL | SGLang's 1 << 30 sentinel for "no top-k" — larger than any vocabulary, so it is neutral by arithmetic rather than by branch. | 06-01 |
| Gumbel-max / exponential race | argmax_v p_v/q_v with q_v ~ Exp(1) is distributed exactly as a draw from p. How both engines sample without torch.multinomial's host synchronisation. | 06-01 |
| Counter-based (hashed) noise | Deriving random draws from stable request coordinates rather than only a batch-global stream. Batch-independent noise requires a backend whose seed/counter mapping preserves those coordinates; exact speculative sampling relies on the verifier's acceptance and residual rule, not identical draft and target draws. | 06-01 |
| Min-p | Keep tokens with p_v ≥ μ · max_j p_j. Relative to the peak, so it adapts: sharp distributions keep few tokens, flat ones keep many. | 06-01 |
| Pivot-based truncation | vLLM's default Triton top-k: estimate a Gaussian cutoff from one sample tile, gather the outliers above it, ternary-search the pivot until exactly k survive. Nothing is sorted. | 06-01 |
| Repetition / frequency / presence penalty | Sign-aware multiplicative, count-scaled subtractive, fires-once subtractive. The only sampling operators that read the request's history rather than the current step. | 06-01 |
| Bonus token | The extra token sampled from the target's own distribution at the last verified position when every draft is accepted. The reason the exponent in E is κ+1 and not κ. | 06-02 |
| Ragged verify metadata | One flat draft-token tensor plus a prefix sum, so requests carrying different draft lengths — including zero — share a single verify pass. | 06-02 |
| Draft slots | KV slots the scheduler reserves per decode request per step for tokens the drafter will append, whether or not that request ends up speculating. Speculation costs KV capacity before it costs latency. | 06-02 |
threshold_acc / threshold_single | SGLang multipliers on the acceptance probability that raise α by biasing the output toward the draft. They void the exactness proof; both default to the identity. | 06-02 |
| Indifference curve (draft selection) | On the (α, c) plane, the locus of draft sources worth exactly as much as a given draft model. Where it crosses c = 0 is the acceptance a free matcher needs to tie. | 06-03 |
| Prompt lookup / n-gram proposer | Propose tokens following an earlier occurrence of the current suffix. No draft-model weights or draft GPU forward are needed, but host lookup, metadata and target verification still cost time. The two cited vLLM implementations choose the oldest occurrence of the longest match. | 06-03 |
| Suffix automaton (SAM) | A linear-space structure holding every substring of everything ingested, with per-state occurrence count and last-occurrence position. Gives ranked continuations at every achievable match length simultaneously. | 06-03 |
| Anchor (suffix decoding) | One node per achievable match length, collected by walking suffix links upward from the deepest match. The draft tree is built from all of them, with fan-out scaled by how long each anchor's match was. | 06-03 |
max_spec_factor | Suffix decoding's multiplier turning match length into draft length, which makes num_speculative_tokens a ceiling rather than a length. | 06-03 |
| Token-Level Intersection (TLI) | vLLM's use_heterogeneous_vocab mapping aligns shared token identities between tokenizers, whose integer ids may differ, and masks unsupported draft tokens. For a proposal supported on that intersection, exact one-step acceptance cannot exceed the target probability mass there. | 06-03 |
| External corpus | A tokenized corpus loaded into SGLang's suffix automaton at startup or over POST /add_external_corpus, so a draft can come from a domain corpus rather than from the prompt. | 06-03 |
| Feature-level autoregression | EAGLE's draft input, fc(cat(embed(x̃_j), f_{j−1})) — draft token j+1 is conditioned on j through both the embedding and the feature, so the tuple is a joint sample rather than a product of marginals. | 06-04 |
| Medusa | Multiple prediction heads on one target hidden state. Coordinatewise marginal argmaxes need not form a likely joint continuation; autoregressive drafting models conditional dependence but does not guarantee a good proposal. The cited implementation is vLLM-only. | 06-04 |
| MTP (multi-token prediction) | Multi-token prediction trains additional future-token prediction modules alongside the base model. The cited DeepSeek-style module separately normalizes token embeddings and prior features before concatenation; other MTP architectures differ. SGLang routes the cited implementation through NEXTN. | 06-04 |
d2t / draft vocabulary | An index map from EAGLE-3's reduced draft vocabulary to target ids. Vocabulary reduction can lower LM-head traffic; the two-thirds estimate belongs to the chapter's stated model shapes and is not a requirement of every drafting architecture. | 06-04 |
| Node budget (draft tree) | num_draft_tokens — how many tree nodes one target pass verifies. Not the same as the accept cap. | 06-04 |
max_tree_depth | spec_steps + 1, the width of every accepted-path array and the maximum tokens a tree can emit. Confusing it with num_draft_tokens is the commonest tree bug; it surfaces as an illegal memory access two kernels later. | 06-04 |
| Tree mask (ancestry mask) | The boolean mask making node i attend to the shared prefix, its own ancestors and itself and nothing else — so the logits at position i are what the target would have produced from that root-to-i path alone. | 06-04 |
FULL_MASK / QLEN_ONLY / bitpacked | Three encodings of that mask. A dense ancestry bit-matrix takes T² bits per request, but the tree structure can be encoded more compactly; FULL_MASK materialises (Σ seq_len)·T bytes purely to satisfy a kernel API. | 06-04 |
| First-child / next-sibling encoding | retrieve_next_token and retrieve_next_sibling: two T-wide int arrays representing an arbitrary tree with no pointers and no recursion. The verifier descends on the first and scans alternatives on the second. | 06-04 |
| Ragged verify | Verifying a different number of drafted tokens for each request in one batch. supports_ragged_verify() is true only for DSpark, but the capability is gated on SGLANG_RAGGED_VERIFY_MODE, which defaults to static — so it is off unless you ask for it. LoRA is compatible with DSpark on that default path; it is only the non-static modes (cap-accept, compact), whose per-request widths break LoRA's uniform-width segment layout, that are refused, and refused at startup. | 06-06 |
| Parallel drafting | Predicting a block in one draft forward rather than sequential draft forwards. The idealized draft-cost term changes from kappa*c to a measured block cost, which need not equal one sequential step. Missing sampled-token conditioning can reduce acceptance; the tradeoff must be measured. | 06-06 |
| Mask-token block | The [bs, γ] draft input whose column 0 is the last verified token and whose other columns are one repeated mask-token id. The concrete shape of a parallel draft. | 06-06 |
| Markov head | DSpark's sequential correction over that block: bias each position's logits by the token sampled at the previous one, with no further transformer forwards. A partial refund on parallel drafting's lost α. | 06-06 |
| Survival probability | DSpark's cumulative per-position confidence is nonincreasing along a request. Global top-k preserves contiguous prefixes only with an earlier-position-first tie rule or explicit closure repair; monotonicity alone is insufficient at equal scores. | 06-06 |
| Frozen-KV MTP | An assistant that owns no KV pool, reads the target's committed cache read-only at a rope phase that never advances, and therefore has no draft-extend pass. vLLM reached the same design under the name constant_draft_positions. | 06-06 |
| Block verification | Sun et al.'s rule uses joint prefix probabilities and residual mass. Under its assumptions, expected accepted length is at least that of the corresponding standard rule, with equality possible. Vocabulary reductions add work; the cited implementation uses vLLM's V2 runner. | 06-06 |
| Synthetic acceptance | An ablation harness that accepts according to a supplied curve instead of the proposal. At fixed draft width and mean accepted length, minimum variance does not change expected wasted slots; burstiness can still affect queueing, adaptive policies and nonlinear batch costs. | 06-06 |
| Adaptive speculation | SGLang's closed loop: an EMA of accepted drafts per request walks κ up or down a per-batch-size candidate ladder, with hysteresis, and the shipped default allows only κ = 0 at batch 64. | 06-06 |
| Adaptive step slot | One candidate κ in that ladder. Swapping between slots at runtime is six pointer assignments; the cost is paid at startup, where every candidate needs its own attention backends and CUDA-graph captures. | 06-06 |
| Multi-layer EAGLE | A distinct trained draft module per step rather than one applied recurrently. Total draft cost sums the costs of invoked modules; using kappa*c already counts kappa steps and must not multiply the per-step cost by kappa again. | 06-06 |
| STANDALONE | An independently trained draft LM with its own embeddings and LM head, run inside the target's process. Distinct from the genuinely out-of-process decoupled design, whose protocol exists but has no consumer at this SHA. | 06-06 |
| DFlash | The cited draft transformer uses target hidden states to construct its KV and shares target-side token/head machinery instead of owning the usual embedding and LM-head pair. The inspected repository excerpts cite no paper; that is not evidence that no paper exists. | 06-06 |
| DSpark | DFlash's parallel mask-token draft plus a Markov head and a confidence head. The only algorithm supporting ragged verify. | 06-06 |
| Bitmask (grammar) | One bit per vocabulary entry, 1 meaning allowed — 16,032 bytes per request per step at vocab 128,256. Identical grammar states under the same vocabulary can share the same allowed-token mask; different states require separate ownership and may produce different masks. | 06-05 |
accept_token / fill_bitmask / apply_token_bitmask | The three-operation grammar interface. In the described backend the mask application runs on the GPU. Advance/fill preserve order within each request; independently owned request states can be processed in parallel on the host. | 06-05 |
| Grammar barrier | SGLang's callback into the worker that resolves the previous batch's tokens and advances the FSM during the current forward. Where a speculative algorithm cannot support it, a grammar in the batch disables overlap scheduling for that step. | 06-05 |
| Structural tag | A constraint type alongside JSON Schema, regex and EBNF. Supported by xgrammar and llguidance only; outlines and lm-format-enforcer log "Skip unsupported" and return an invalid grammar. | 06-05 |
| Retokenisation hazard | Forced text must be re-tokenised whole, because BPE is not compositional across a boundary — the last cached token and the first forced characters may merge into a different token the KV cache has already committed to. | 06-05 |
| Reasoner grammar backend | A grammar backend that wraps another and suppresses the mask during the model's thinking block. SGLang composes it; vLLM handles the same concern inside its manager. | 06-05 |
§7
Parallelism and runtime
| Term | One line | Chapter |
|---|---|---|
| Tensor parallelism (TP) | Splitting individual weight matrices across GPUs. Two all-reduces per transformer block — one at the end of attention, one at the end of the MLP — so 160 per Llama-3-70B forward pass. | 05-01 |
| Pipeline parallelism (PP) | Splitting layers into stages. GPipe's micro-batch remedy inverts at decode, because stage time is set by weight traffic and does not shrink with the micro-batch; the engines pipeline independent batches instead, turning the bubble into a low-load phenomenon. | 05-02 |
| Expert parallelism (EP) | Placing MoE experts on different GPUs; turns routing into an all-to-all problem. | 05-03 |
| CUDA graph | A captured, replayable launch sequence, removing per-launch CPU cost from the decode step. | 08-01 |
| Column-then-row pairing | Split the first matmul by output columns and the second by input rows. Forced, not chosen: an elementwise nonlinearity commutes with a column split and not with a row split, and every other pairing puts a collective on the 3.5×-wider intermediate. | 05-01 |
| Residual stream | The [num_tokens, d] tensor every rank must agree on byte for byte at each block boundary. The all-reduce is the price of putting it back together. | 05-01 |
| Vocab-parallel embedding / LM head | Sharding along the vocabulary rather than the hidden dimension: the input masks out-of-range ids and all-reduces, the output all-gathers [n, V]. Vocabulary-sized logits can be much wider than hidden-state activations; some sampling paths avoid a full gather. | 05-01 |
num_kv_head_replicas | tp_size // total_num_kv_heads — the integer division in the QKV weight loader that physically creates KV-head duplication past TP = h_kv. | 05-01 |
| Vocab-parallel argmax | Gathering only per-rank (value, index) pairs instead of full logits: 56 bytes per token at TP=8 against 219 KiB. Used for greedy decoding and draft heads. | 05-01 |
| Sequence parallelism | Rewriting all-reduce → RMSNorm as reduce-scatter → local norm → all-gather, sharding the norm along the token axis at unchanged wire volume. Incompatible with FX-level piecewise splitting, because the residual changes shape. | 05-01 |
| Interconnect intensity (Inet) | FLOPs of compute per byte of TP traffic per rank. Compared against the machine's π/B_net it says where TP stops paying — past TP ≈ 25 on NVLink, past TP ≈ 2.35 on InfiniBand NDR. | 05-01 |
| Seam (PP) | A pipeline stage boundary. Its payload is two tensors — hidden states and the un-added residual, because both engines fuse the residual add into the next norm — so 2·T·d·b bytes, and both shard that send across the TP group. | 05-02 |
PPMissingLayer | torch.nn.Identity placeholders vLLM builds for the layers a PP rank does not own, keeping state-dict keys and layer indices global. They hold zero parameters and produce zero KV. | 05-02 |
| Bubble fraction | (S−1)/(S+M−1) — the idle share of device time in a GPipe schedule with S stages and M micro-batches. | 05-02 |
| Batch-level pipelining | Keeping S independent batches in flight from the continuous-batching queue rather than splitting one batch. The bubble stops being a per-step tax and becomes a load condition: below S in-flight batches, stages idle. | 05-02 |
max_concurrent_batches | vLLM's scheduler-side PP contract — the depth of the deque of in-flight futures the engine core keeps. | 05-02 |
pp_max_micro_batch_size | max_running_requests // pp_size. Why PP silently divides your per-step batch while leaving your total concurrency alone. | 05-02 |
| Ring deadlock | Every PP rank sending before receiving. Safe on CUDA, fatal elsewhere; the fix is parity-ordered send/recv. The symptom is a hang with no error and every rank at 0% SM utilisation. | 05-02 |
| Plain data parallelism | N complete engine replicas behind a router: linear throughput, no inter-replica communication, N copies of the weights — and no shared prefix cache, which is why cache-aware routing exists. | 05-03 |
ScatterMode | SGLang's per-layer tensor-layout enum — TP_ATTN_FULL, FULL, SCATTERED — deciding whether the DP-attention all-gather and reduce-scatter run at all. An all-to-all MoE backend makes sparse layers SCATTERED and the gather disappears. | 05-03 |
attn_tp_size | The attention block's tensor-parallel degree when it differs from the MoE block's. Derived, not a flag; the canonical DeepSeek tp = dp = 8 configuration gives 1, i.e. attention is pure DP. | 05-03 |
| Idle batch / dummy batch | A fabricated empty forward pass a rank with no work must still run, because the FFN is a collective and every DP rank enters every pass. DP attention is one engine in lockstep, not N independent replicas. | 05-03 |
| Balancedness | Mean over max per-rank token load in an MoE layer — the fraction of purchased MoE FLOPs actually used. 0.55 means 45% of expert compute is stall. Both engines define it identically. | 05-03 |
| EPLB | DeepSeek's Expert-Parallel Load Balancer: replicate the highest-load-per-replica experts into spare physical slots, then permute slots onto GPUs hierarchically so traffic prefers NVLink. Replication is what makes it work — permutation alone cannot fix one hot expert. | 05-03 |
| Redundant experts | Physical expert slots beyond the logical count, which EPLB hands to whichever expert has the highest load per replica. HBM traded for balancedness, 44 MB per slot for DeepSeek-V3. | 05-03 |
| Dispatch / combine | The MoE all-to-all pair: send each token to the ranks holding its experts, then send the partial outputs back and weight-sum them. Real libraries deduplicate by destination rank, so a token crossing costs once per rank rather than once per expert. | 05-03 |
DpPaddingMode | SGLang's choice between MAX_LEN (pad every rank to the longest, all-gather) and SUM_LEN (zero-fill a sum-length buffer, all-reduce), picked by whichever moves fewer rows. The gap between them is exactly the cost of DP-attention imbalance. | 05-03 |
| Elastic expert parallelism | Surviving a dead rank or growing the EP world without a restart: an active-rank mask tensor, pre-reserved slots past the launch world size, a scale state machine, and a separate process holding expert weights in DRAM. | 05-03 |
| α (collective fixed cost) | The per-invocation cost of a collective — kernel launch plus rendezvous — independent of message size. Roughly 92% of a small decode all-reduce's cost, which is why both engines' custom kernels attack α and not bandwidth. | 05-04 |
| One-shot / two-shot P2P all-reduce | Direct peer reads over NVLink instead of a ring. One-shot spends (N−1)S bytes to need one barrier; two-shot achieves ring's 2(N−1)S/N in two barriers instead of 2(N−1). | 05-04 |
| Crossover S* | The message size at which two-shot overtakes one-shot. Both engines hard-code it — vLLM as two literals, SGLang as a per-architecture, per-world-size table with separate graph and eager columns. | 05-04 |
| Push vs pull | Whether a rank writes its contribution into every peer's workspace or reads every peer's buffer. Push wins at the smallest sizes; only pull has a CUDA-graph mode. | 05-04 |
| Graph zero-copy input registration | Under capture the input's address is fixed by definition, so it can be exported by CUDA IPC and dereferenced from a device-side pointer table — deleting Θ(S) staging copies and two kernels from a captured two-shot all-reduce. | 05-04 |
| Rank-invariant predicate | The rule that every fallback condition on a collective path must evaluate identically on all ranks. A rank-variant one deadlocks with no error: one rank spins on peer flags nobody will write while the others block in NCCL. | 05-04 |
| NVLink SHARP / multicast | In-switch reduction on Blackwell and later: a rank issues one multicast store and the fabric performs the N-way sum, removing the (N−1) factor from the byte count. Cannot use the zero-copy path, because it routes through the workspace address. | 05-04 |
| Broadcast the plan / replicate the scheduler | The two ways to keep N ranks in lockstep. vLLM's serving path broadcasts one SchedulerOutput; SGLang runs a full scheduler per rank over broadcast inputs. Everything else about each engine's process architecture follows. | 05-05 |
collective_rpc | vLLM's single executor primitive — a (method, args, kwargs, output_rank) tuple enqueued once on a shared-memory ring. Its docstring restricts it to control messages; tensors travel on NCCL or a KV connector. | 05-05 |
ExecutorWithExternalLauncher | vLLM's one replicate-the-scheduler path, for offline torchrun jobs only, guarded by an assert — because the design assumes every rank's scheduler decides identically and V1 scheduling is not guaranteed to. | 05-05 |
| Headless mode | A non-zero node running a worker executor with no scheduler above it. How vLLM does multi-node TP and PP without Ray at this SHA. | 05-05 |
| Rank divergence | Two ranks reaching different scheduling decisions from the same inputs. It does not produce a wrong answer — it produces a permanent stall inside a collective, because the fabric is doing exactly what it was told. | 05-05 |
| Rank consensus checker | SGLang's opt-in divergence detector: each rank hashes the sequence of marked decisions it observed and MIN/MAX-reduces the digest. A mismatch calls os._exit(1) rather than letting a diverged server keep answering. | 05-05 |
| Soft watchdog | A timeout that py-spy-dumps every rank's stack without killing anything. Set it below the hard watchdog so you get a stack trace of the near-miss instead of only the kill. | 05-05 |
| Zombie service | An HTTP server that accepts requests and never answers — what a C++-level abort leaves behind when it kills a child without running any Python handler. The failure mode SubprocessWatchdog exists to prevent. | 05-05 |
§8
Architectures that change the inference story
| Term | One line | Chapter |
|---|---|---|
| Total vs active parameters | Total parameters determine resident model capacity; active parameters estimate per-token arithmetic. Neither alone determines capability or batch traffic: expert reuse, routing, precision and the number of distinct experts touched also matter. | 07-01 |
| Routed expert / shared expert | The E experts a token may be routed to, versus the always-on FFN every token runs. Its relative cost depends on active routed experts, widths, precision, batch size and caching; resident-byte and compute shares are different quantities. | 07-01 |
| Experts touched | The coupon-collector count of distinct experts a batch reaches. Under independent uniform routing, passes half of E near batch 22 and 95% near batch 94 for the stated DeepSeek-V3 setup. Small or correlated production batches can touch far fewer experts; skew can reduce weight traffic while increasing rank imbalance. | 07-01 |
| Grouped GEMM | A collection of expert-specific matrix multiplications grouped into a launch. Each expert has its own routed row count; this is not one ordinary dense GEMM over a shared weight matrix. | 07-01 |
Alignment block / BLOCK_SIZE_M | The tile height each expert's row segment is rounded up to so a grouped GEMM can run. That single ceiling is the source of MoE's tensor-core padding tax, and the tile must track M_e = Bk/E, not B. | 07-01 |
| Contiguous (compact) layout | Pack each expert's rows into an aligned segment. Rounding adds fewer than one tile of rows per nonempty expert; routing imbalance can still increase compute and communication latency even when padding is small. | 07-01 |
| Masked layout | Reserve m_max rows per expert, giving capacity proportional to E*m_max. The cited default can reserve a rounded whole-chunk bound; an optional path caps it using measured expert loads. Reserved capacity is not executed work: masks and kernel tile bounds matter. | 07-01 |
| Padding factor | Padded rows divided by real routed rows. The chapter's examples give 2-16 times as many allocated rows at small batches. Actual extra MMA work and memory traffic depend on the layout, masks and kernel; allocation alone does not establish a runtime penalty. | 07-01 |
e_score_correction_bias | DeepSeek's per-expert scalar trained by a control loop rather than by gradient. It shifts which experts win the top-k while the routing weight is gathered from the unbiased scores — auxiliary-loss-free balancing. | 07-01 |
| Group-limited routing | Partition the experts into n_group groups, score each group by its top-2 sum, keep topk_group of them, then take top-k inside. Bounds how many destination ranks a token can reach under a group-aligned placement. | 07-01 |
| Capacity factor | A multiplier setting a per-expert token budget, commonly discussed in MoE training. The cited inference routing paths do not drop tokens to enforce a training-style quota. Their logical loads vary, while graph-compatible buffers may be statically reserved; this is not a claim about every engine backend. | 07-01 |
| Absorbed-weight trick | Reassociate MLA's key and value up-projections with the query and attention output so attention can use cached latents. This avoids context-wide expansion but changes query-side projection and pairwise attention costs; the total FLOP comparison depends on prefix and query lengths. | 07-02 |
| Decoupled RoPE | MLA's split: the 128-wide nope half absorbs cleanly, the 64-wide rope half is projected straight from the hidden state, rotated and cached verbatim with one copy shared by every head. The bribe that keeps the other 512 dimensions position-free. | 07-02 |
kv_b_proj / w_kc / w_vc | The checkpoint tensor holding [W_UK; W_UV] concatenated per head, split at load time into the two folded weights. The absorbed decode path never calls it. | 07-02 |
reorder_batch_threshold | vLLM's per-backend query length below which a row counts as "decode" and takes the absorbed kernel — 1 for Triton MLA, 128 for FlashMLA, 512 for FlashAttn-MLA. Three different definitions of decode for one model. | 07-02 |
| DSA indexer | DeepSeek-V3.2's lightweight scorer that top-k selects cached positions for the MLA kernel. Needs its own second cache — 132 bytes per token per indexing layer — on top of the latent. | 07-02 |
| Recurrent state / SSM temporal state | The per-sequence summary a selective state-space layer absorbs each token into and then overwrites. O(1) in context but not small: 97.4 MiB per sequence for Nemotron-H-8B at fp32. The 779-token crossover against Llama-3-8B compares this component alone; including the hybrid's attention KV moves the total-cache crossover to about 891 tokens. | 07-03 |
| Short-convolution state | A rolling window of the last W−1 token activations keeping the depthwise convolution causal across step boundaries. A FIFO, not a summary, and 1.4% of the state. | 07-03 |
mamba_cache_mode | vLLM's none / align / all: one resident state per sequence, two (the previous must live while the next is written), or a checkpoint chain over the whole model length. | 07-03 |
| State checkpoint / snapshot point | A stored copy of the recurrent state at a chosen position — the only place prefix reuse can resume, because a summary contains no sub-object equal to an earlier summary. A hit rounds down to the last checkpoint and the remainder is replayed. | 07-03 |
| Page-size alignment (hybrid) | Inflating the attention block size until an attention page is at least as large as one Mamba state page — 1,040 or 1,280 tokens for Nemotron-H-8B, which quantises prefix caching two orders of magnitude coarser for the attention layers too. | 07-03 |
mamba_branching_seqlen | SGLang's record of the last chunk-aligned position that has KV but no state — a proven-shared point where a checkpoint would turn a future replay into a hit. | 07-03 |
MambaAttentionBackendEnum | vLLM's second, parallel backend enum for SSM layers. Selected by the layer's declared mamba_type rather than by capability, so nothing is validated and nothing is negotiated — an unsupported combination surfaces as a kernel failure, not a startup refusal. | 07-03 |
| Placeholder expansion | One <image> token becoming N placeholder ids that the text embedding table fills and the encoder output then clobbers. The counts must match exactly, which is what the "assign N multimodal tokens to M placeholders" error means. | 07-04 |
| Encoder cache | A third budget beside weights and KV, holding encoder output embeddings. vLLM counts embedding slots and refcounts them because its scheduler admission-controls on the number; SGLang counts bytes with a plain LRU because its lookup happens after admission. vLLM's size is max(max_num_batched_tokens, max_tokens_per_mm_item) slots — five 1,369-token images at an 8,192-slot budget, eleven once a 16,384-token per-item cap lifts the floor; SGLang fits nine in its 100 MB. | 07-04 |
Multimodal identifier (mm_hash) | The digest keying both the encoder cache and the block hash's extra keys. vLLM hashes the raw item plus the per-request processor kwargs; SGLang hashes the preprocessed feature tensor. Inconsistent max_pixels silently forks vLLM's. | 07-04 |
| Encode-prefill (EPD) disaggregation | Splitting a request one stage earlier than P/D and shipping embeddings. Ships 2Lh_kv d_h/d times fewer bytes than KV — 16× for Llama-3-8B shapes — and moves the CPU decode/resize pipeline off the box that owns the KV cache. | 07-04 |
| SGMV | Segmented Gather Matrix-Vector multiplication: one base GEMM for the whole batch, then a per-row indexed low-rank correction. The correction is [tokens, r]-shaped in the middle and never touches a d × d matrix. | 07-05 |
| Shrink / expand | The two halves of that correction — [tokens, d] → [tokens, r] then back out — accumulated into the base output in place. | 07-05 |
| Chunked SGMV | SGLang's fixed-size segments over the adapter-sorted token stream, so segment count is predictable and prefill can be CUDA-graph captured. vLLM instead gives each adapter a variable-length row range. | 07-05 |
| Active slots vs registered adapters | Registered adapters, resident GPU slots and distinct adapters active in one batch are different counts. The scheduler may defer requests or evict reloadable adapters when slots fill; low-level allocation can raise if a requested active set cannot fit. | 07-05 |
| Drainer | SGLang's fix for a starved adapter: any adapter whose requests have waited past a threshold triggers a controlled drain of one running adapter to free its slot. Disabled by default. | 07-05 |
| Per-adapter cache partition | Because the adapter name is a block-hash extra key, identical prompts under different adapters cannot share a KV block. Fifty tenants sharing one 2,048-token system prompt burn 22% of a 60 GB pool on redundant copies. | 07-05 |
| Pooling model | One forward pass, output a vector or a scalar, no decode loop. No autoregressive decode loop; compute/traffic behavior and completion time depend on input shape and backend, with no autoregressive KV growth — an encoder-only KV spec reports zero bytes. | 07-05 |
| Pooling type | LAST, CLS or MEAN specifies the reduction from token states to an output. Follow the checkpoint's trained pooling and normalization contract; an arbitrary override can degrade task quality even when output shapes remain valid. | 07-05 |
| Cross-encoder | A joint query-document sequence produces a relevance score, so candidate documents form a natural batch. In a bidirectional encoder, query states depend on the document and cannot be reused as an independent cached query prefix; a causal reranker has a different reuse contract. | 07-05 |
§9
Compilation and runtime
| Term | One line | Chapter |
|---|---|---|
| Capture / instantiate / replay | CUDA stream capture records operations and dependencies rather than executing the captured GPU work. Instantiation prepares an executable graph; replay submits it. Python recording, warmup and allocation are distinct startup costs, and one graph launch still has nonzero submission overhead. | 08-01 |
CUDAGraphMode | vLLM's NONE / PIECEWISE / FULL plus two tuple composites whose first element is the decode mode and second the mixed-batch mode. The default is FULL_AND_PIECEWISE. | 08-01 |
| Capture size / bucket ladder | A set of captured batch shapes; eligible batches can be padded to a supported bucket. Extra rows may share weight traffic, but attention, activations, metadata and kernel choices still cost work. Being below a weight-only roofline ridge does not make padding free. | 08-01 |
| Graph memory pool | The shared allocation captured graphs replay into. vLLM estimates it by capturing the two largest buckets and extrapolating, then subtracts it from the KV budget before the block manager sees a byte. | 08-01 |
| Breakable CUDA graph | SGLang's default prefill backend on CUDA — a graph that may be broken at op boundaries — as against full, tc_piecewise and disabled. Not the same axis as vLLM's piecewise, which is an FX-graph split. | 08-01 |
| Dynamo / AOTAutograd / Inductor | Dynamo traces Python into FX graphs with guards; AOTAutograd performs ahead-of-time graph transformations including functionalization; Inductor schedules code generation. Generated code can include GPU Triton, CPU C++ and external library calls. Fusion is limited by supported operations and graph boundaries. | 08-02 |
| Graph break | A point Dynamo cannot represent as a tensor op, ending the traced region. vLLM compiles fullgraph=True, so a break is a hard error rather than a silent slowdown. | 08-02 |
| Guard | A predicate validating reuse of compiled code. vLLM's optimized path can bypass guard evaluation under engine-managed invariants, while stock compile and compatible guard-evaluation configurations retain it. Disabled checks make coverage of shape-dependent behavior the engine's responsibility; they do not imply every branch is wrong. | 08-02 |
| Splitting op | An operator that ends a compiled subgraph — attention and the KV-cache update. vLLM keeps a central allowlist in config; SGLang registers with a decorator at the definition site. | 08-02 |
| Piecewise compilation | Compiling around the opaque attention op rather than through it: 33 compiled pieces and 32 opaque regions for Llama-3-8B, each compiled piece separately CUDA-graph-capturable. | 08-02 |
| Compile range | An interval of [1, max_num_batched_tokens] that gets its own Inductor artefact. Pieces times ranges is where the startup minutes go. | 08-02 |
| Compile cache key | A key derived from selected configuration, environment, traced source and compiler/build factors. A changed key can miss an existing artifact and require compilation unless a matching artifact is already available. The exact factor set is version-specific; a cache hit does not eliminate process startup. | 08-02 |
| Inductor graph partition | Deferring the split to codegen time, after FX passes have run, so a whole-graph rewrite like sequence parallelism can coexist with piecewise capture. | 08-02 |
| Program (Triton) | One kernel instance in the launch grid, indexed by tl.program_id, roughly corresponding to a thread block. Source expresses block-shaped values; launch configuration still controls execution resources, including num_warps and thus 32 times that many threads on NVIDIA GPUs. | 08-03 |
| Block value | A tensor-shaped value such as tl.arange(0, BLOCK_N) whose elements are distributed across a program's threads. The compiler selects layouts, registers and instructions; a vector expression does not promise simultaneous execution or freedom from spills. | 08-03 |
| Masking idiom | mask=cols < n_cols, other=0.0 against a power-of-two BLOCK_N. Forgetting other= poisons any reduction; dividing by BLOCK_N instead of n_cols passes every power-of-two test and is wrong everywhere else. | 08-03 |
num_stages | Software-pipeline depth over the K-loop — how many BLOCK_K slabs are in flight in shared memory at once. More hides latency and costs shared memory. | 08-03 |
| Autotune cache key | The tuple of runtime arguments whose values change which config wins. Too narrow and you tune for a shape that behaves differently; too wide and you stall on every new request shape. | 08-03 |
| Grouped tile ordering | Walking a GROUP_SIZE_M-tall column of output tiles instead of a full row, so the concurrently resident blocks share both a small set of A rows and a small set of B columns in L2. Pure scheduling, expressible because you own the grid. | 08-03 |
| Epilogue fusion | Applying output operations such as bias, activation and suitable scaling to an accumulator before storing the result, avoiding separate output passes. Group scales that vary along the inner reduction dimension must be applied within that reduction; they cannot generally be deferred to a final epilogue. Split reductions or other kernel stages can still require intermediate traffic. | 08-03 |
| Epilogue visitor tree (EVT) | CUTLASS's compile-time expression tree for that epilogue, so any composition of broadcasts, elementwise ops and casts is a type — one compiled kernel covering per-tensor, per-channel and per-token scaling. | 08-03 |
CuTe Layout | A shape-and-stride object composable, permutable and coalescible at compile time. Turns "how wide a vector may I legally load" into an if constexpr rather than a runtime branch. | 08-03 |
| Tile hierarchy | CUTLASS's nested tile levels — cluster, CTA, warpgroup, instruction — each a template parameter you instantiate rather than a runtime argument you pass. | 08-03 |
| Collective builder | CUTLASS's mainloop/epilogue constructor, which solves for pipeline depth at compile time from whatever shared memory the epilogue already claimed. Triton's num_stages is a number you guess; CUTLASS's the type system computes. | 08-03 |
| Safetensors | An 8-byte header length, a JSON header mapping tensor name to dtype, shape and byte offsets, then a raw payload. No code runs, every tensor has an O(1) offset, and the file can be mapped rather than copied. | 08-04 |
WeightsMapper | vLLM's declarative checkpoint-name rewriter. It also staples a shard_id onto the tensor object, so the tree walk that routes the weight can stay name-agnostic. | 08-04 |
weight_loader callback | The per-parameter function that narrows a checkpoint tensor to this rank's slice and copies it in. Choosing the parallel Linear class is choosing the sharding, because each class ships its own loader. | 08-04 |
AutoWeightsLoader | vLLM's prefix-grouping walk down the module tree, delegating at any module that defines its own load_weights. Replaces per-model string matching. | 08-04 |
stacked_params_mapping | SGLang's flat table mapping checkpoint suffixes onto a fused parameter plus a shard id — the same job as the mapper, spelled as an explicit loop. The legacy path warns rather than raising on an unmapped tensor. | 08-04 |
EntryClass | The module attribute SGLang scans for instead of keeping a registry table; the class name must equal the checkpoint's architecture string. An import error inside a model file removes that architecture with only a warning. | 08-04 |
| Read amplification (TP) | Physical checkpoint bytes read relative to the useful local shard. Mapping a whole tensor and narrowing a view does not necessarily touch every page; shard orientation, layout, iterator behavior, readahead, prefetching and page-cache reuse determine actual I/O. | 08-04 |
| Pre-sharded checkpoint | sharded_state / presharded: per-rank files so each rank reads only what it keeps. Bound to one exact (TP, PP, quantization) tuple — a build artefact keyed by config, not a checkpoint. | 08-04 |
| Weight cache daemon | SGLang's separate process retains loaded, TP-sharded, post-quantization GPU weights and exports CUDA IPC handles. Consumers can share the same allocation rather than duplicate the weights; retained residency and ownership lifetime remain costs, and non-weight startup work still runs. | 08-04 |
§10
The serving system around the engine
| Term | One line | Chapter |
|---|---|---|
| Tier (of API compatibility) | Which of honoured / clamped / ignored / extension a request field falls into. A property of the deployment, not the API: logit_bias is honoured on vLLM until you enable speculative decoding, after which it is a startup warning and a silent no-op. | 09-01 |
| Chat template | A Jinja file that ships with the checkpoint, not with the engine, turning messages into the string the model sees — including system messages the client never sent. Change it in the model repo and your server's behaviour changes with no code deployed. | 09-01 |
add_generation_prompt | The template variable that appends the assistant header. Set it false by accident and the model continues the user's turn — grammatical, plausible and completely wrong. | 09-01 |
| Tool parser / format detector | The per-model reverse of whatever tool-call format a fine-tune learned, selected by a server startup flag rather than by the request. Choose wrong and tool calls arrive as ordinary content with a 200. | 09-01 |
| Terminal SSE frame | The literal data: [DONE] sentinel. Once the first frame is on the wire the 200 is frozen, so an engine failure afterwards arrives as a data: error object followed dutifully by [DONE]. Streaming clients must parse frames, not statuses. | 09-01 |
stream_options.include_usage | The field that adds a final usage-only chunk. vLLM requires it before honouring continuous_usage_stats and rejects stream_options without stream; SGLang treats them independently and accepts the latter silently. | 09-01 |
| Incremental detokenisation | Decoding a window of recent tokens and emitting only the newly stable suffix, because one character can span four tokens and a token is a byte string rather than a character. | 09-02 |
| Prefix offset / read offset | The two nested window edges, p ≤ r ≤ n. Both decodes share a left edge so any decoder-side space cleanup cancels in the subtraction. Both engines trail by an admittedly arbitrary five tokens. | 09-02 |
| Byte-fallback token | A vocabulary entry holding one raw byte, emitted for characters whose bytes never earned a merge. Decoding an incomplete multibyte prefix alone can yield a replacement character; an ASCII byte is already valid, which is why per-token decoding is wrong and only on the inputs your tests lack. | 09-02 |
| Hold latch | Emitting nothing and leaving the offsets unchanged when the decoded window ends in a replacement character. A mid-string one is a genuinely invalid id and is emitted — get the distinction backwards and you either garble output or hang the stream forever. | 09-02 |
| Stop-string hold-back buffer | Withholding max_s |s| − 1 characters from every chunk, because SSE cannot retract and the longest dangerous unsent suffix is a proper prefix of a stop string. A long stop string is a latency knob you did not know you were turning. | 09-02 |
DecodeStream | The Rust tokenizers object vLLM's fast detokeniser drives one token at a time. When its internal state goes wrong it cannot be repaired in place — the stream is thrown away and rebuilt, losing the left context. | 09-02 |
| Stream interval | Emitting a chunk every k tokens rather than every token. Can amortise network emission across k tokens and coarsen observed ITL; internal per-step detokenization may remain. Both engines default it to 1 in the pinned paths. | 09-02 |
| Per-request mailbox | The one-slot object a batched engine output is demultiplexed into. vLLM merges on put, bounding container count but not necessarily merged payload size; SGLang appends and merges on get, so it can see the depth and warn about it. | 09-03 |
| Coalescing (output stream) | What that mailbox does when the consumer falls behind. Lossless for tokens — they all arrive, concatenated — and lossy for timing, which is exactly the thing ITL is computed from. Not memory coalescing, below. | 09-03 |
| Coalescing (memory) | Unrelated sense, from the GPU side: a warp's 32 lanes touching one contiguous run so the hardware services them in a single transaction. It is what a paged KV layout puts at risk, and it survives as long as a tile does not straddle two blocks. The two senses share a word and nothing else. | 03-03 |
| Frontend budget | The host CPU a frontend may spend per token, one step's period divided by the batch: 140 µs at batch 32, 35 µs at batch 128 against the 4.48 ms floor. Every serving system becomes frontend-bound at a batch size, not at a token rate. | 09-03 |
| Batch-notify quantum | How many per-request tasks the output fan-out wakes before yielding — 16 in SGLang, 128 in vLLM. A deliberate throughput trade that shows up as ITL variance rather than in the mean. | 09-03 |
| Deferred free | Returning KV only after outstanding users can no longer read or write it. Whether reclamation is immediate or delayed depends on queued/in-flight work and synchronization, not a universal exactly-one-forward delay. | 09-03 |
| Deprecation shim | A file kept at an old import path that re-exports and warns. At these SHAs vLLM's entrypoints/openai/api_server.py and both projects' root-level benchmark scripts are shims — every tutorial citing them describes code that no longer runs. | 09-03 |
| Cache dilution | Repeated placement can duplicate hot prefixes and reduce effective hit probability. Under uniform independent references Hrr=min(1,M/W), but the union of N caches can still contain up to NM distinct keys; popularity and routing affect reuse. | 09-04 |
| Approximate cache model | The router's own radix tree of what it sent where — keyed on raw characters to keep a tokenizer off the routing path, and decayed by its own LRU sweep. It never learns what a replica actually kept. | 09-04 |
| Tenant (routing) | The worker a prefix is currently assigned to in that tree. A stale or unhealthy tenancy silently disables affinity and sends every such request to the same fallback worker. | 09-04 |
cache_threshold | Matched characters over total input characters, below which the router falls back to shortest queue. Being a ratio, a long unique suffix suppresses cache-aware routing exactly when the absolute shared prefix is most valuable. | 09-04 |
| Balance guard | The two-condition test — an absolute in-flight gap and a ratio — that suspends affinity when the fleet is uneven. With the shipped defaults it rarely fires at moderate load, which is a deliberate bias toward hit rate. | 09-04 |
| Cold prefix cache | A newly started replica may have fewer reusable prefixes and initially higher prefill cost. Partial caches are useful; arrival mix, routing, and eviction determine warming, not a mandatory wait until the whole pool fills. | 09-04 |
| Time to parity | An illustrative estimate combining weight load and cache-fill work, not a readiness deadline. The 09-04 example includes 61 seconds of aggregate prefill work that can serve traffic; useful service can start before full cache parity. | 09-04 |
| MiniLB | SGLang's 462-line Python prefill/decode load balancer, explicitly a debugging tool: random pair selection, dual dispatch, PD mode only. The readable version of what the Rust gateway does. | 09-04 |
sgl-model-gateway / smg | SGLang's router as a separately versioned product — one Rust crate producing three binaries plus the sglang-router wheel, on its own release cadence and its own Docker image. | 09-04 |
| Busy vs overloaded | A bounded queue with KV parked near the watermark is the design point; a sustained positive queue derivative plus a rising preemption rate is not. Alert on the second, never page on the first. | 09-05 |
| Blast radius | Which shell a fault happened in. A fault in the HTTP process kills only itself; a fault in a GPU worker propagates outward and kills everything. Determined by process topology, not by how bad the log looks. | 09-05 |
| Worker-level fault recovery | Neither engine has it. vLLM's fault-tolerance package rebuilds a stateless data-parallel process group; a dead tensor-parallel rank marks the engine DEAD. Restart is the recovery mechanism, so the operational work is fast detection and a clean restart. | 09-05 |
§11
Benchmarking and performance engineering
| Term | One line | Chapter |
|---|---|---|
| Workload shape | Four distributions, not a product category: input length, output length, arrival process, and prefix-sharing structure. Same silicon and same weights, and the spread across ordinary workloads is 29× in output tokens per second. | 10-01 |
| Prefill:decode split | The share of engine time each phase takes. It swings from 82% prefill on summarisation to 2% on agentic traffic, and it decides which optimisation in this book is worth anything on your deployment. | 10-01 |
| Arrival Ca2 | Squared coefficient of variation of inter-arrival times, the term Kingman's formula multiplies queueing by. vLLM's --burstiness is exactly its reciprocal; SGLang samples exponential only, so it is pinned at 1. | 10-01 |
| Sharing rate | Cached tokens over prompt tokens — and, separately, the absolute p50 of matched tokens, because the router only ever sees the ratio and the GPU only ever feels the tokens. | 10-01 |
--random-range-ratio | The same flag with inverted meaning: vLLM centres the length interval on L, SGLang right-anchors it at L with ⌊Lr⌋ as the floor. At the shared default of 0.0 the mean input length differs by 2×. | 10-01 |
| Open loop / closed loop | Fixing the arrival process versus fixing the number of requests in flight. A closed loop cannot represent overload — the client is the admission control — so it converts every capacity deficit into a proportional latency increase. | 10-02 |
| Amortised ITL | Dividing an SSE chunk's gap by the tokens it carried and emitting that many samples. SGLang's native backend does this; vLLM records one sample per frame. Their medians depend on integer frame sizes and sample weighting, not only mean acceptance. A pooled ITL mean and a mean of request TPOTs can differ without speculation. | 10-02 |
| Load curve | What a benchmark actually produces: one (throughput, latency-percentile) point per offered load. A throughput number without the latency it was achieved at is a coordinate with one axis missing. | 10-03 |
| Knee | The offered load at which the SLO binds — the one interesting point per configuration, and not where the curve happens to look bent. | 10-03 |
| Relaxation time (trel) | A stationary M/M/1 spectral-relaxation scale, not a convergence guarantee for batched serving. The worked three-scale heuristic gives about 1025 arrivals; wall time depends on the service-time constant. Check queue, cache, compiler, and hardware warmup separately. | 10-03 |
| MDE (minimum detectable effect) | A planning effect size at specified significance and power under a variance/design model. Three independent runs per arm and CV = 5% give about 11.4% MDE under the normal approximation; actual significance requires the observed data, not comparison with that threshold. | 10-03 |
| Run-to-run vs batch-invariant | Same input in the same batch across two runs (usually holds already) versus the same input in a different batch (holds in neither engine by default). Conflating them is why "re-running the failing test alone passes". | 10-04 |
| ULP | Units in the last place for a specified floating-point format and value range. Reduction-order differences can change close decisions; ULP distance alone does not guarantee an argmax flip. JSON binary64 logprob distance is not original tensor precision. | 10-04 |
| Truncation align size | Snapping the chunked-prefill length to a multiple of the attention split tile, so prefill split boundaries do not move with the leftover token budget. Batch invariance is a scheduling constraint, not only a kernel property. | 10-04 |
| M-independent dispatch | A quantized-GEMM dispatcher that selects its CUTLASS config from N and ignores M, so the tile — and therefore the reduction tree — does not change with the token count. | 10-04 |
| Attribution budget | The vertical distance between a measured step and its derived floor. Every microsecond of it belongs to a named cause, and the trace says which. | 10-05 |
| Trace shape | Four common decode-step geometries — launch-bound, bandwidth-bound, compute-bound, occupancy-limited. The middle two are indistinguishable in a timeline and differ only in achieved DRAM throughput, which is what picks the tool. | 10-05 |
| Detailed annotations | Per-step span aggregates ΣN_Q, ΣN_KV, ΣN_Q², ΣN_Q·N_KV — the roofline inputs for the batch that step actually ran. Both engines independently landed on the same four. | 10-05 |
| Effective ITL | Step time divided by the fraction of wall clock a replica spends in decode rather than in someone else's prefill. Why a 10.8 ms step can read as a 59 ms ITL at 77% prefill occupancy. | 10-05 |
§12
Engine internals and comparison
| Term | One line | Chapter |
|---|---|---|
| MRV2 / Model Runner V2 | vLLM's second worker rewrite, in vllm/v1/worker/gpu/. Default for every dense model despite the "experimental" README; the 8,008-line runner is the fallback. No log line says which one you got. | 11-01 |
idx_mapping | MRV2's per-step array mapping batch position to a permanent request-state row. Because state rows never move, there is no condense() and no per-request Python mirror. | 11-04 |
StagedWriteTensor | MRV2's GPU-resident base tensor plus accumulated ragged CPU diffs, applied with one Triton launch — replacing the per-step block-table copy that forced the persistent batch to stay dense. | 11-04 |
BatchExecutionDescriptor | MRV2's explicit CUDA-graph selection object, keyed on padded token count and effective LoRA count, returning cg_mode = NONE when nothing matches. V1's equivalent decision is implicit in the forward context. | 11-04 |
libtorch_stable | The csrc subtree vLLM is migrating kernels into for PyTorch's stable ABI. Most active kernels now live there, so the old paths hold residue rather than code. | 11-01 |
| vLLM IR | vllm/ir/ — a functional dialect separating operator semantics from implementation and dispatch, successor to CustomOp, migrated piecewise. | 11-01 |
EngineCoreRequest / EngineCoreOutputs | The two hot-path msgspec structs crossing vLLM's process boundary. array_like=True makes field order wire ABI, so a new field may only be appended at the end. | 11-02 |
RequestOutputCollector | vLLM's one-slot mailbox with merge-on-overflow. A slow consumer changes chunk shape, never token content — making put drop instead of merge would be a correctness bug. | 11-02 |
| Zero-copy threshold | The byte size above which a tensor rides as its own ZMQ frame instead of being copied inline — and above which its backing buffer must be kept alive until ZMQ is done with it. | 11-02 |
ENGINE_CORE_DEAD | The single sentinel frame vLLM's engine pushes on its way out. The frontend translates it into EngineDeadError and fails every in-flight stream at once — correct, because the engine's state is gone. | 11-02 |
| Input budget | The scheduler's second counter, beside the token budget: rows the model runner's input buffers must hold. Larger, because the worker appends draft tokens the scheduler never issued. | 11-03 |
| Skipped-waiting queue | vLLM's second waiting queue, holding requests blocked on grammar compilation, a remote KV transfer, stale in-flight output, or the max_loras cap — so they do not block the head of the main queue. | 11-03 |
allocate_slots returning None | The only admission-failure signal in vLLM's scheduler: no exception, no log line, no counter. The correct breakpoint for every "why was my request not scheduled" question. | 11-03 |
| Entry-point group | The Python packaging group an engine discovers plugins through. Registry dicts are per-process, so on a TP=8 node your entry-point function runs ten times in ten interpreters. | 11-05 |
| Re-entrant plugin | Re-entrancy means an invocation remains correct when entered again before it returns. Repeated initialization is a different contract: idempotence means repeating it has no additional effect. Unconditional list append is not idempotent; use guarded registration and explicit resource ownership. | 11-05 |
is_argmax_invariant | A logits processor's claim about whether it can move the argmax. It is a dispatch decision, not documentation: declare it wrongly and greedy requests silently ignore your parameter. | 11-05 |
can_implement contract | A linear kernel answering "can I run this configuration" with a reason string. The selector concatenates every rejection into one error, which is why a bad quantization config produces a wall of text rather than one line. | 11-05 |
| Seam | A place you can insert code without forking the engine. The published guarantee in either project covers exactly one interface; everything else is a seam that exists because somebody needed it, not a contract that holds. | 11-05 |
| SRT | SGLang RunTime, python/sglang/srt/. Everything this book calls the engine lives under it; the lang/ frontend that named the project is marked deprecated in-tree. | 12-01 |
PortArgs | The dataclass allocating all five SGLang IPC endpoint names up front, with field comments naming both ends of each socket. ipc:// by default, TCP when DP attention needs cross-node reach. | 12-01 |
NS(...) namespace bag | SGLang's marker assigning a ServerArgs field to one of eleven published namespaces. Each process calls publish(server_args, role=...) once at start, and everything else reads a frozen snapshot rather than a mutable global. | 12-01 |
| DWDP | Distributed Weight Data Parallelism: keep tokens where they are and prefetch expert weights over NVLink into a virtual-memory composite address space. The inverse of the usual all-to-all trade; prefill-only and disaggregation-restricted. | 12-01 |
SenderWrapper | The no-op output sender every non-zero rank gets. Because sending is a no-op rather than a branch, no if tp_rank == 0 is scattered through the result path — which is what keeps a replicated scheduler tractable. | 12-02 |
FutureMap | What makes SGLang's one-iteration overlap lag legal: the scheduler leaves input_ids unset and the forward stream gathers the previous step's sampled token from a device-resident buffer indexed by request-pool slot. The scheduler never learns the token, only the slot. | 12-02 |
| Tree core | The UnifiedTreeCore behind UnifiedRadixCache's facade. Reading the 2,887-line facade as the tree wastes an afternoon; most of it is HiCache plumbing. | 12-03 |
Component (ComponentType) | The per-node dimension — FULL, SWA, MAMBA — that used to be a whole cache subclass. Model shape now selects a component tuple inside one class rather than a class. | 12-03 |
| Node handle | The opaque integer the scheduler holds instead of a tree node, so cache internals stay invisible above the interface. | 12-03 |
| D-leaf / device leaf | A node holding device KV with no locked component and no child holding device KV. Eviction candidates are a maintained set of these, heapified per pass — interior nodes are protected by not being offered, which is the invariant vLLM expresses as a calling convention. | 12-03 |
| Insert barrier | The suspend point in SGLang's resumable insert state machine: the tree emits an action — free this KV, back this node up — and the facade performs it, because the tree does not own the allocator. | 12-03 |
| Hook registry | SGLang's general plugin mechanism, and not a set of registries at all: a monkey-patch framework with four hook types keyed on any fully-qualified dotted name in the codebase — the module docstring's own example targets the scheduler's schedule. | 12-04 |
| Serve backend | The one seam in either project carrying an explicit compatibility version, for plugging a different runtime into sglang serve. Plugins are told to hard-code the version they implement rather than read SGLang's at runtime. | 12-04 |
| Complexity budget | Where a project spends its irregularity — vLLM in the block layer, SGLang in the index. The place it spends it is where its bugs live, and a data-structure choice three layers down sets the cost of a policy choice at the top. | 13-01 |
| Handshake compatibility hash | The digest a KV connector compares before interpreting remote memory as a KV tensor. It includes the engine version, so a disaggregated vLLM fleet has no rolling upgrade — only "every node identical" or "unchecked". | 13-03 |
| Bytes per link | The quantity that actually governs disaggregation, rather than bytes per node. It depends on the prefill/decode TP mapping, which is why the same model gives both 14% and 1.7% of prefill and why no published overhead ratio ports to another cluster. | 13-03 |
| Prefill fraction (φ) | Share of a request's GPU time spent in prefill, T_p / (T_p + T_d), both amortised to a per-request share of a batched step. It bounds how much any prefill-side optimisation can possibly buy you. | 13-02 |
| Prefix hit rate (η) | Fraction of a request's prefill tokens a perfect prefix cache would serve from an existing hit. A property of prompt structure, not of the engine — which is why it can be estimated before you pick one. | 13-02 |
Multi-item scoring (--enable-mis) | SGLang reranking mode — one request scores many candidate passages against one query. Setting it auto-disables CUDA graphs, so throughput on that path bears no relation to what a standard benchmark predicted. | 13-02 |
| Keyspace fork | A shared KV store silently partitioned by config_suffix (model, TP rank and size, PP, CP), so replicas of differing shape cannot reuse each other's pages even though they are pointed at the same store. | 13-03 |
index_topk | The checkpoint-declared sparse-attention budget that sets the dense/sparse crossover in SGLang's DSA path; used as the default for the dense-attention KV-length threshold. | 13-03 |
Rebase debt (R) | R = r · T — upstream commits touching your file set (rate r) accruing over a project of T weeks. An upper bound on conflicts, and the reason to extend an abstract base rather than the concrete class that churns 16× faster. | 13-04 |