Say it like a staff engineer.

Prefill pushes every prompt token through the model in one pass, so the weight matmuls are large and arithmetic intensity is high. Decode emits one token per sequence per step, so each step streams all the weights and that sequence's KV cache from HBM to do very little math.

Batching fixes half of this. Weight reads are shared by every sequence in the batch, so intensity grows roughly with batch size. KV reads are private to each sequence, so decode attention stays bandwidth-bound at any batch size.

KV bytes per token = 2 (K and V) × layers × KV heads × head dim × bytes per element. Llama 3.1 70B has 80 layers, 8 KV heads and head dim 128, so BF16 costs about 320 KiB per token. One 128K-token context therefore needs about 40 GiB.

Maximum concurrency ≈ free KV memory ÷ (average context × bytes per token). vLLM does this at startup: it profiles a forward pass, claims --gpu-memory-utilization of the GPU (default 0.9), gives what is left after weights and activations to KV blocks, and logs the resulting maximum concurrency.

Earlier servers reserved one contiguous KV region per request, sized for the maximum length. The vLLM paper found 60–80% of KV memory wasted that way. PagedAttention stores KV in fixed-size blocks (16 tokens by default) reached through a per-request block table, cutting waste to under 4%.

The block table then became the base for most later features:

• Sharing: parallel samples and common prefixes point at the same physical blocks, with copy-on-write.
• Prefix caching: full blocks are hashed, so a new request can reuse any cached prefix.
• Movement: CPU offloading, P/D transfer and remote KV stores all move KV block by block.
• Hybrid models: one allocator can hold attention KV, sliding-window KV and Mamba state if their pages share a size.

Continuous (iteration-level) batching reschedules after every forward step: finished sequences leave and waiting ones join at once. Static batching holds the whole batch until its longest sequence ends, idling the GPU for everyone else.

The new problem is interference. A long prompt admitted into a running batch slows that step for every decoding request, which shows up as inter-token latency spikes. vLLM answers with chunked prefill under a per-step token budget, a cap on prefill tokens per request per step, and, at scale, P/D disaggregation.

• TTFT: queueing plus prefill; how long the user waits before anything appears.
• ITL / TPOT: time between output tokens; ITL is the per-token distribution, TPOT the per-request mean.
• End-to-end latency: TTFT + TPOT × (output tokens − 1).
• Throughput: output and total tokens per second per GPU, and tokens per dollar.
• Goodput: requests per second that meet every SLO at once, the target DistServe optimizes.

Weights are a fixed cost, and the KV cache gets whatever memory is left. In the January 2025 session, Llama 70B went from TP=1 to TP=2 and gained 13.9× more KV cache blocks and 3.9× more token throughput. More KV allows bigger batches, which raises decode arithmetic intensity.

Because CPU overhead, not the GPU, had become the limit. With Llama-8B on an H100 a forward pass takes about 5 ms, so HTTP handling, scheduling, input prep, detokenization and streaming ate a large share of each step. V1 runs an isolated EngineCore busy loop that only schedules and executes, while tokenization, multimodal preprocessing, detokenization and streaming overlap in the frontend process over ZeroMQ.

Tensor parallelism changed too. V0 kept the scheduler and worker 0 in one process, an asymmetric design. V1 workers cache request state and receive only per-step diffs, so the scheduler lives apart and every worker runs the same code on one GPU or many.

It hands out a token budget. The decision is a map of {request_id: num_tokens}, and prompt tokens and generated tokens are treated alike, so there is no separate prefill or decode phase. A decoding request asks for one token; a prefilling request asks for its remaining prompt, clipped to the budget. Chunked prefill falls out of this for free, and prefix caching simply lowers the tokens a request still needs.

Only full blocks are cached. Each block's hash covers its parent block's hash, its own token IDs and extra keys (LoRA ID, multimodal input hashes, cache salt), so one hash identifies the whole prefix up to that block. A new request walks its block hashes and reuses the longest cached run.

Freed blocks join a doubly linked free queue in LRU order. Eviction pops the head, and a cache hit can pull a block from the middle in O(1). V0's version cost enough CPU to lower throughput at low hit rates, so it shipped off. V1 made eviction constant-time and cut Python object churn: under 1% throughput loss at a 0% hit rate and several-fold gains at high hit rates.

CUDA graphs record a sequence of kernel launches once and replay it, removing per-kernel CPU launch cost, which matters when a decode step lasts a few milliseconds. Graphs need static shapes, and attention's per-step metadata varies. Piecewise graphs capture everything except attention and run attention eagerly between the captured pieces.

Full graphs also capture attention, which works for uniform decode-only batches. FULL_AND_PIECEWISE uses full graphs for decode-only steps and piecewise graphs for mixed steps. For hybrid models whose Mamba kernels are written in Triton, which has high launch overhead, it recovered V0 performance and gave up to 91% more throughput on granite-4.0-h-tiny.

Without it, the CPU schedules step N+1 only after step N's tokens come back, so the GPU idles during scheduling and input prep. Async scheduling plans step N+1 while step N still runs. The catch is that the scheduler does not yet know step N's results: which requests hit a stop token, how many draft tokens were accepted, what grammar state follows.

It must schedule optimistically and correct a step later, and every feature has to tolerate that lag. In V1, async was retrofitted, so many features needed awkward special cases. Model Runner V2 makes async the core assumption: GPU-side input prep consumes rejection-sampling results directly, with no CPU–GPU sync.

MRV2 (March 2026) is a ground-up rewrite of the component that turns scheduler output into GPU inputs, runs the model and samples. V1's runner had grown past 6,700 lines in one file. It coupled persistent request state to per-step inputs, had async scheduling bolted on, and did input prep and sampling as many small CPU operations.

• Stable state table: each live request owns a fixed row, and a gather builds each step's ordered inputs.
• GPU-native input prep: Triton kernels build input_ids, positions, query_start_loc and seq_lens on the device.
• Async-first: designed for zero CPU–GPU sync, including with speculative decoding.
• Triton sampler: Gumbel-max sampling without materializing softmax; top-k logprobs computed only for the selected candidates.
• ModelState: model-specific logic (multimodal embeddings, attention metadata, CUDA graph capture) is isolated; the largest file is under 1,300 lines.

Qwen3-0.6B on one GB200 went from 16K to 25K output tokens/s (+56%). GLM-4.7-FP8 with MTP on 4× GB200 saw 6.3% lower TPOT.

Layers of the same type (full attention, sliding window, Mamba or linear attention) form a KV cache group, and all groups draw from one shared block pool with a single page size. Memory then flows between layer types on demand instead of being split up front. V0 instead preallocated Mamba state per sequence, sized by max_num_seqs, so a wrong guess meant either OOM or low concurrency.

A uniform page is hard for Mamba. On Nemotron-Nano-12B-v2 one sequence's Mamba state is about 2.57 MiB, while a 16-token attention block is about 64 KiB. vLLM grows the attention block size until the pages match (672 tokens, for example) and pads Mamba pages slightly. It later decoupled the kernel's block size from the manager's, so TRT-LLM attention kernels still run on Blackwell.

The connector API lets vLLM ask an external store for KV before scheduling a request and hand it newly computed KV afterwards. Since v0.9.0 it is asynchronous, so loads and stores overlap model compute. The offloading connector (v0.11.0) adds a pluggable backend whose core is one transfer function; the CPU backend uses cudaMemcpyAsync, which runs on the GPU's DMA engines.

A copy kernel wins on tiny blocks but competes with the model for SMs. Once blocks were 0.5–2 MB, DMA reached 83.4 GB/s bidirectional versus 68.5 GB/s for the kernel. End to end, DMA gave 5.5–15% more throughput on the worst-case model and up to 32% on Llama-3.1-8B; at a 0% hit rate the kernel was 6% slower than no offloading at all.

Compare transfer time with prefill time. Llama-3.1-8B stores 128 KiB of KV per token, so 10K tokens is about 1.3 GB, roughly 26 ms at the 50 GB/s one-way rate measured on H100. Prefilling those 10K tokens costs about 1.6×10¹⁴ FLOPs (2 × 8B × 10K), roughly 0.3 s at a realistic 500 TFLOPS. Reload wins by about 10×, and the gap widens with context because attention compute grows quadratically while transfer grows linearly.

They make KV reusable across engines and extend capacity past one machine. LMCache is a KV layer outside the engine for cross-query reuse and P/D transfer. It moves KV in 256-token chunks instead of 16–64 KB pages, reaching 400 Gbps from CPU where small transfers manage about 88 Gbps. It adds zero-copy writes, layer-by-layer pipelining with a one-layer GPU buffer, and lookup, move, pin and compress APIs for routers.

Mooncake Store, integrated in 2026, is a distributed KV pool. In standalone-store mode an external Mooncake client owns each node's CPU memory and disk, even on CPU-only nodes, and vLLM workers just request blocks. Routers such as Dynamo and llm-d can then send a request to any instance and still get a hit.

A Mamba or linear-attention layer keeps one recurrent state that is valid only at one position. Reusing a prefix needs a snapshot of that state exactly at the prefix boundary, and snapshotting every token is too costly. vLLM combines two retention policies:

• Interval-based: keep the prompt-end state at every turn, so the next turn or a forked subagent can extend it.
• Marconi-style selective: when a prefix is seen a second time with no checkpoint, recompute the state and save one at that boundary for later requests.

Megatron-style TP splits the QKV and MLP up/gate projections by columns, so each GPU owns a slice of heads or features. Element-wise ops such as SiLU run on those slices. The attention output and MLP down projections split by rows, each GPU produces a partial sum, and an all-reduce combines them: two all-reduces per layer.

Each GPU reads only 1/N of the weights per step, which effectively multiplies memory bandwidth and cuts decode latency. The price is an all-reduce on every layer of every step, which needs NVLink-class links; decode messages are small, so latency rather than bandwidth dominates.

Work through it in order:

1. Fit: the smallest TP inside one node that holds the weights plus enough KV for your target concurrency and context.
2. Too big for a node: PP across nodes when inter-node links are slow; stretch TP across nodes only on a fast fabric such as an NVL72 domain.
3. Scale out: data-parallel replicas, which need no cross-replica traffic, behind a cache-aware router.
4. MoE: expert parallelism for the experts, usually with data-parallel attention (DEP), especially for MLA models.
5. Long context: decode context parallelism (DCP) for long-context decode, prefill context parallelism (PCP) for long prompts.

MLA compresses K and V into one shared latent, in effect a single KV head, so TP cannot split the cache by heads and every rank stores the full latent. Data-parallel attention fixes this: each rank runs attention only for its own requests with its own KV, so KV capacity grows with the number of ranks.

The MoE layers then use expert parallelism. Each token is sent all-to-all to the ranks holding its routed experts and combined afterwards; DeepSeek-R1 activates only 37B of its 671B parameters per token. Wide EP combines EP with DP across many GPUs (--enable-expert-parallel). vLLM supports DeepEP's high-throughput and low-latency all-to-all kernels, Perplexity's kernels and an NCCL all-gather/reduce-scatter path.

DBO splits a step's batch into two micro-batches and overlaps one micro-batch's MoE all-to-all with the other's compute, following DeepSeek's micro-batching design. With --enable-dbo, ranks first all-reduce to agree that splitting is worth it (threshold --dbo-decode-token-threshold). Worker threads then run each micro-batch, each yielding while its dispatch or combine is in flight.

It helps when communication dominates, as with high EP degrees, where profiles showed dispatch and combine outweighing the small decode compute. It hurts tiny batches, because splitting halves GEMMs that are already small, hence the threshold.

Real traffic routes tokens unevenly across experts, so in wide EP some ranks idle while others overload, and lockstep means the busiest rank sets the pace. vLLM implements DeepSeek's expert-parallel load balancer (--enable-eplb). Each MoE forward pass records per-expert load, and a sliding window aggregates it across ranks. At each rebalance interval, EPLB computes a new logical-to-physical expert map, optionally with redundant copies of hot experts, and shuffles weights while serving.

Both shard along the sequence. DCP gives each rank 1/N of every request's KV, so long-context decode attention, which is memory-bound and grows with context, is split across ranks without replicating MLA latents. Each decode layer then needs a query gather before attention and a merge of partial outputs afterwards. PCP instead shards a long prompt across ranks: for DeepSeek V4, PCP8 prefilled a 32K prompt 2.65× faster than TP8.

For Kimi K3, DCP8 beat TP8 on decode latency and concurrency. vLLM fused DCP's communication into the attention kernels through symmetric-memory buffers, bypassing NCCL and cutting about 13% per layer. PCP still replicates decode-side state, so it suits dedicated prefill workers.

Every MoE layer's all-to-all needs all ranks at once, because a token on one rank may need an expert on another. The group moves at the pace of its slowest rank, and ranks with no work run dummy passes just to join the collectives.

• Prefill interference: one long, compute-bound prefill on any rank delays decode on every rank, which is why P/D disaggregation pays off more for EP deployments.
• Prefill cadence: prefills landing on different steps stall the group repeatedly, so --prefill-schedule-interval admits prefills only every Nth step, on a counter aligned across ranks.
• Imbalance: hot experts or uneven requests per DP rank turn directly into idle time, so EPLB and DP-aware routing matter.

The vLLM docs give two reasons: tune TTFT and ITL separately, and control tail ITL, because prefills can no longer interrupt decodes. They are also blunt that disaggregated prefill "DOES NOT improve throughput" on its own. The throughput gains seen at scale come from specializing each pool, for example PCP prefill workers and wide-EP decode workers with low-latency all-to-all kernels.

It is rarely worth it for short prompts, small fleets where a dedicated pool sits idle, or networks without RDMA-class bandwidth for KV transfer. It pays most for long-prompt workloads and MoE deployments, where one prefill stalls a whole lockstep EP group.

One interface serves P/D transfer, CPU offloading and remote caches, and it has two halves:

• Scheduler side: get_num_new_matched_tokens says how many tokens can be loaded instead of computed, update_state_after_alloc reacts once blocks exist, build_connector_meta tells workers what to load and save this step, and request_finished manages transfers still in flight.
• Worker side: start_load_kv begins async loads, wait_for_layer_load blocks per layer so transfer can pipeline with compute, save_kv_layer and wait_for_save handle stores, and get_finished reports completed transfers.

In-tree connectors include NIXL, LMCache, Mooncake, MoRI-IO (ROCm), FlexKV, the offloading connector and MultiConnector, which composes several, such as NIXL for P/D plus CPU offloading. NIXL, NVIDIA's inference transfer library, hides the transport (RDMA, NVLink, storage) behind one asynchronous read/write API.

Rate-match the pools: prefill capacity in requests per second must equal decode capacity at your SLOs, or one pool queues while the other idles. The vLLM team uses two phases:

1. Saturation profiling: benchmark prefill-only and decode-only deployments separately, sweeping parallelism (TP or wide EP) and size (8, 16, 32 GPUs) until throughput saturates.
2. P/D sweep: derive the ratio from those saturation points, then sweep concurrency on the combined deployment.

The llm-d design lays out four reasons:

• Expensive, uneven requests: RAG has long inputs and short outputs, reasoning the reverse. Imbalance compounds: an overloaded replica's ITL rises, requests stay longer, and load grows further.
• Cached state: sending a multi-turn or agentic request to the replica holding its prefix skips most of its prefill.
• Different phases: compute-bound prefill and bandwidth-bound decode favour specialized replicas.
• Different QoS: code completion needs milliseconds, while batch summarization tolerates hours.

llm-d plugs into Kubernetes' Gateway API Inference Extension. For each request the gateway, such as Envoy, calls an Endpoint Picker (EPP). The EPP runs filters (healthy pods, prefill or decode role) and then weighted scorers (prefix-cache affinity, KV cache utilization, queue depth) and returns the best vLLM pod. Prefix awareness can be approximate, from the router's own history, or precise, from KV cache events that vLLM emits as blocks are stored and evicted.

On two 8×H100 nodes with a long-input, short-output multi-turn benchmark, llm-d cut mean TTFT about 3× at 4 QPS (Llama 4 Scout, 2 replicas). It served about 50% more QPS within a P95 TTFT of 2 s (Scout, 4 replicas) and 2× the QPS within that SLO (Llama 3.1 70B, 4 replicas).

All four run vLLM underneath; they differ in the platform they assume and what they optimize.

Decode is memory-bound, so scoring k+1 positions in one target forward pass costs about the same as generating one token. A cheap drafter proposes k tokens and the target checks them all in one pass. Rejection sampling accepts a prefix of the drafts and samples a correction at the first rejection, and the acceptance rule makes the output distribution identical to sampling from the target alone.

It spends compute to save bandwidth, so it loses once batches are large enough to be compute-bound. In 2024 tests on Llama3-70B with 4×H100, a draft model was 1.5× faster at QPS 1 on ShareGPT but 1.4× slower at high QPS. n-gram drafting went from 2.8× faster to 1.8× slower on CNN/DailyMail. Content matters too: an EAGLE-3 head trained on chat data reached 2.1× on math and RAG but did poorly on German-to-English translation.

• Measure: draft acceptance rate, per-position acceptance and mean acceptance length, all exported since v0.9.1.
• Tune by load: longer drafts at low request rates, shorter at high ones; translation-like traffic may want k of 0 or 1, RAG 5 or more.
• Scope it: enable it on latency-sensitive routes and the decode pool, not on batch pools.
• Test at real load: compare TPOT and throughput per GPU at your target concurrency, never only at QPS 1.

Train one against your target on traffic that looks like yours. Speculators, a Red Hat library in the vLLM project, standardizes the loop: extract target hidden states offline with vLLM, train an EAGLE-3, P-EAGLE, DFlash or DSpark drafter, and save it in a Hugging Face-compatible format. vllm serve <speculator> then reads the config and turns speculation on, or you pass --speculative-config with the method and num_speculative_tokens.

• Scheduler: a request's k drafts are extra tokens in the step's budget, and rejected positions are rolled back afterwards.
• KV cache: blocks must have room for k+1 new tokens; KV written for rejected drafts is simply overwritten later.
• CUDA graphs: captured sizes must cover (k+1) × batch tokens. EAGLE gained CUDA graph support in v0.9.1, and Model Runner V2 can capture several draft passes in one graph.
• Async scheduling: the CPU does not know how many drafts were accepted until the step ends, so MRV2 lets GPU-side input prep consume the rejection-sampling results directly.

Match the scheme to the bottleneck.

• Weight-only (W4A16, W8A16; GPTQ, AWQ): 4-bit weights cut weight memory about 4×, which is what memory-bound decode and small batches need. A mixed-input GEMM dequantizes weights inside the kernel and does the maths in 16-bit.
• Weights and activations (W8A8 in FP8 or INT8): runs on low-precision tensor cores, doubling peak FLOPs over BF16, which helps compute-bound prefill and large batches. FP8 needs Hopper or Ada; Blackwell adds FP4 formats such as NVFP4 and MXFP4.

Hopper needed a new mixed-input kernel. Marlin, built for Ampere, used older mma instructions and lost about 37% of Hopper's peak compared with wgmma. Machete (CUTLASS 3.5.1) pre-shuffles weights offline for 128-bit shared-memory loads, keeps weights in registers through a transposed product, and uses TMA with warp specialization. It served Llama 3.1 70B on one H100 at 5 requests/s with TTFT under 250 ms and TPOT under 100 ms.

Pass two gates on your own workload: accuracy and performance. Red Hat's validated-model process is a good template: capacity-planning runs with GuideLLM, accuracy runs with the Language Model Evaluation Harness, and serving with vLLM across accelerators.

• Accuracy: compare against the BF16 baseline on suites that match real use (reasoning, maths, code, long context, languages), plus product checks such as tool-call validity and JSON-schema adherence.
• Performance: benchmark at your real prompt and output lengths and concurrency, because a scheme that wins at batch 1 can lose at batch 256.
• Calibration: produce checkpoints with versioned LLM Compressor recipes (GPTQ, SmoothQuant, AWQ and others), calibrated on data that resembles production traffic.

It halves KV bytes per token, roughly doubling the concurrency or context that fits, and halves the bytes each decode step reads, which speeds memory-bound attention. The risks are accuracy and support: quantization error accumulates over long contexts, scales must be calibrated or computed on the fly, and backends and GPUs differ in what they support.

For portability and maintainability. Hand-tuning kernels for every GPU generation and vendor does not scale, while one Triton kernel of about 800 lines, against roughly 70,000 for FlashAttention 3, runs unchanged on NVIDIA, AMD and Intel. On H100 it reached 100.7% of FA3's performance on long decode requests, and on MI300 it was about 5.8× faster than earlier implementations. It is the default on ROCm and the fallback everywhere else.

• Q-blocks: group all query heads that share a KV head, plus several query tokens, into one work item so tl.dot gets large enough tiles.
• Parallel tiled softmax ("3D kernel"): split the KV traversal across instances for decode, then reduce in a second kernel; heuristics decide when the extra launch pays.
• Persistent kernels: launch a fixed grid sized to the GPU and pull work from metadata, so CUDA graph replays do not repeat wasted waves.

Through hardware plugins. A vendor implements vLLM's platform interface (device, attention backend, communication, custom ops) in its own package, such as vllm-ascend for Huawei Ascend NPUs, the Tenstorrent plugin or vllm-metal for Apple Silicon, while the core stays shared. AMD ROCm lives in-tree with its own kernels and the Triton attention backend. By 2024 vLLM already ran on NVIDIA, AMD, Google TPU, AWS Inferentia and Trainium, Intel Gaudi and GPUs, and x86, ARM and PowerPC CPUs.

Memory scales with total parameters while compute scales with active ones: DeepSeek-R1 activates 37B of its 671B parameters per token. Every expert must be resident somewhere, yet each token touches only a few. So per-expert batches are small: tokens per expert ≈ batch × top-k ÷ number of experts. DeepSeek-V3 sends each token to 8 of 256 routed experts, so a 256-token decode batch gives each expert about 8 tokens, a bandwidth-bound GEMM over full expert weights.

The responses follow: large global batches through wide EP with DP attention, grouped-GEMM expert kernels, block-scaled FP8 or MXFP4 expert weights, and EPLB for skewed routing.

DeepSeek-V3.2 adds a lightning indexer in front of MLA. For each query token the indexer scores earlier tokens with its own small key cache and keeps the top 2,048; sparse MLA then attends only to those, or to everything when the context is shorter. Beyond 2K tokens, attention cost per token stops growing, apart from the indexer's cheap scan.

• A second cache: indexer keys live in their own buffer, separate from MLA's, and this model requires 64-token blocks.
• A new FP8 layout: each token's MLA entry is 656 bytes: 512 FP8 latent values, 16 bytes of scales and 128 bytes of unquantized BF16 RoPE values.
• Batching: the indexer has separate prefill and decode paths, and batched prefill tracks each request's start and end so tokens never score another request's context; shorter results are padded with -1.
• Kernels: DeepGEMM's indexer kernels, FlashMLA's sparse attention and a fused top-k.

Hybrids such as Qwen3-Next, Nemotron Nano 2, MiniMax-Text-01 and Granite 4.0 keep a few full-attention layers and make the rest linear, because attention's KV grows linearly with context and its prefill quadratically. The engine needs three things:

1. State management: a fixed-size recurrent state per sequence, held in the same block pool as attention KV (Q13).
2. Launch overhead control: many Mamba and linear-attention kernels are Triton, with high CPU launch cost. V1 with piecewise graphs was slower than V0 at low concurrency; FULL_AND_PIECEWISE fixed it, with 2–18% more throughput on Nemotron-Nano-12B-v2 and up to 91% on granite-4.0-h-tiny.
3. Feature adaptations: prefix caching needs state checkpoints (Q17), P/D must transfer the state, and speculative decoding must roll the state back, because rejected tokens have already updated it.

It treats preprocessing and the encoder as separate, cacheable work:

• Off-loop preprocessing: image decoding, resizing and cropping run in a separate process, with a cache for repeated inputs, so the GPU loop never waits on them.
• Encoder cache: vision embeddings are kept after the encoder runs, so a long text prefill can be chunked across steps without re-running the encoder. V0 had to process the image and text in one step.
• Multimodal prefix caching: image hashes join token IDs in block hashes, so multi-turn chats about the same image hit the cache.

V1's speedup over V0 was larger on Qwen2-VL than on text models.

Omni models are pipelines of different models. Qwen3-Omni has a Thinker (about 30B, MoE, with audio and SigLIP2 vision encoders) for reasoning and text, a Talker (about 3B, MoE) that turns the Thinker's hidden states into audio codec codes, and Code2Wav, a ConvNet vocoder. vLLM-Omni describes the model as a graph of stages, each with its own devices, memory share and parallelism, behind one OpenAI-compatible API (vllm serve <model> --omni).

Autoregressive stages reuse PagedAttention, continuous batching, CUDA graphs and the scheduler. A connector splits control from data: on one node, payloads move through shared memory and only a handle crosses the control plane; across nodes, RDMA or TCP. Chunks stream so later stages start before earlier ones finish. On Qwen3-TTS-1.7B it reached a real-time factor of 0.17 against 2.64 for Hugging Face Transformers.

The engine compiles the constraint (choice list, JSON schema, regex, grammar or structural tags) into a grammar state machine. Every step it turns the allowed next tokens into a bitmask, applies it to the logits before sampling, and advances the grammar state with the chosen token. In V0, compiling a grammar stalled the whole engine; in V1 compilation is non-blocking, handled by the scheduler while other requests keep decoding.

• XGrammar: caches compiled grammars; lowest TPOT for repeated schemas and long generations.
• Guidance (llguidance): computes constraints token by token; faster TTFT on complex or changing schemas, suiting multi-tenant traffic.
• auto (default): picks a backend per request.

With reused schemas, TPOT was only marginally higher than unconstrained generation.

The chat template renders the tool definitions into the prompt, the model emits a call in its own format, and a model-specific parser (--tool-call-parser; more than 20 ship in tree, including Hermes, Mistral, Llama 3 JSON, Qwen and DeepSeek) converts it into OpenAI-style tool_calls. --enable-auto-tool-choice switches this on. The guarantee depends on tool_choice:

• required or a named function: vLLM uses structured outputs, so calls always match the tool's JSON schema: valid, not necessarily sensible.
• auto: the model decides, and arguments can occasionally be malformed unless strict constraints apply, which needs VLLM_ENFORCE_STRICT_TOOL_CALLING (on by default) and at least one tool marked strict: true.
• none: tool calling is off.

It is a routing layer for mixture-of-models systems: it picks which model serves each request. Requests flow client → Envoy → an ExtProc filter running the Go router → the chosen backend. Version 0.1 (Iris, January 2026) extracts six signal types (domain, keyword, embedding similarity, factual, user feedback and preference), and a decision engine combines them with AND/OR rules and priorities. Plugins add semantic caching, jailbreak and PII detection, hallucination detection, system-prompt injection and header rewriting.

Use one when you run models of very different cost and quality, want safety checks at one chokepoint, or need to cut spend. In Red Hat's walkthrough, 86% of requests stayed on a local quantized Qwen model and 14% went to Gemini, with a P50 routing latency of 40 ms against 800–11,000 ms of inference.

SemiAnalysis's AgentX benchmark, built from real agentic coding traces, shows the shape:

• Long sessions: a median of 43 turns per session.
• Long inputs, short outputs: a median of 142K input tokens and 444 output tokens per request.
• Heavy reuse: a prefix-cache hit rate above 96%.
• Subagents: 44% of sessions spawn at least one, with a median of four among those.

Each turn appends the latest tool result and resends the whole context, so inputs keep growing while each turn adds only a short new prefill. The vLLM team names three challenges: keeping many sessions' KV warm between turns, executing long-context steps fast enough, and finding a P:D ratio when hit rates and context lengths vary by session.

That is head-of-line blocking. With first-in, first-out chunked prefill, one long prefill can claim the entire token budget step after step, so short cached turns on the same rank never get scheduled. Cap the prefill tokens one request may take per step with --long-prefill-token-threshold: at 512, short turns join every step. On DeepSeek V4 Pro with B300s, this raised total tokens per GPU-second by up to 93% and P90 interactivity by about 2.3×.

The trade-off is a higher TTFT for the long request itself, so TTFT-sensitive deployments should use a larger threshold. In DEP deployments, also set --prefill-schedule-interval so prefills land on the same steps across ranks and the steps between are decode-only.

In aggregated DEP deployments the team saw large KV-usage imbalance across ranks and tried balancing by queue depth, running tokens or KV utilization. On AgentX traces, all three lost to simple session-aware sticky routing. Inter-turn gaps are short, so the next turn often arrives while its prefix is still resident on the previous GPU. Moving the session forces a KV fetch from the shared pool, and although the transfer overlaps compute, the prefetched blocks occupy GPU capacity, so the destination admits fewer sequences.

The queue looked balanced while the system served fewer concurrent requests. Route on the state already resident on each worker, not only on queued work.

The vLLM team published three "bitter lessons":

• Pipeline parallelism, including chunked PP: it scales well on long, fresh prompts, but warm agentic turns add only a few hundred to a few thousand new tokens. There is too little work to fill the pipeline, and bubbles eat the gain.
• Decode context parallelism on DeepSeek V4: it works for pure and hybrid MLA models such as Kimi K3, but V4's attention has an indexer, a compressor and the main attention to partition. After heavy optimization it only matched DEP.
• Load balancing: it lost to session-sticky routing (Q51).

Make the agent's structure explicit to the engine:

• Agent hints: harnesses pass session structure, branch points, cache positions, expected tool-call latency and session lifecycle; the Q3 2026 roadmap starts with Session-ID and Correlation-ID.
• Programmable KV cache: workloads choose how KV is prefetched, evicted or soft-pinned.
• Session-based KV management: use the gap between turns to move retained KV towards the worker likely to serve the next turn.
• Split routing: send first turns and later turns to differently configured pools.

The same roadmap targets production-ready distributed, multi-tier KV offloading and speculative decoding with acceptance lengths above 5 tokens.

• Realistic workload: traces or datasets that match your prompt and output length distributions and your prefix sharing; uniform random prompts hide cache and scheduling effects.
• Load sweep: raise request rate or concurrency to saturation, and at each point report P50, P90 and P99 TTFT and ITL plus throughput per GPU; plot throughput against interactivity.
• Controlled state: say whether the prefix cache is warm, keep versions, flags and hardware fixed, and record them.
• Tools: vllm bench serve, GuideLLM for capacity planning, and public references such as SemiAnalysis's AgentX harness.
• Accuracy beside speed: quantization, speculation and new kernels can all change outputs.

Scale on LLM-native signals: waiting-queue depth (vllm:num_requests_waiting), KV cache usage (vllm:kv_cache_usage_perc), SLO attainment for TTFT and ITL, and preemption rate. GPU utilization misleads, because a GPU can be busy yet stalled on memory, and it says nothing about KV headroom. With disaggregation, scale each pool on its own signal: queue and TTFT for prefill, KV usage and ITL for decode.

Cold start is the hard part. A replica must load tens to hundreds of GB of weights, compile and capture CUDA graphs before it can serve, which takes minutes. Keep warm pools, reuse compile caches, stream weights from peers or fast storage, and scale ahead of predictable peaks; reducing cold start is on vLLM's Q3 2026 roadmap.

vLLM moves fast: from v0.16 in February 2026 to v0.22 in June, a release every two to three weeks, and defaults change along the way.

1. Pin everything: vLLM, PyTorch, CUDA and kernel libraries, in immutable images.
2. Diff behaviour: read release notes for changed defaults and deprecations, and compare the effective engine config.
3. Accuracy gate: per-model evals, including tool calls and structured outputs, against the current version.
4. Performance gate: fixed benchmark scenarios per model at your SLO points, with regression thresholds.
5. Canary, then progressive rollout by pool, with rollback images kept warm.
6. Upgrade coupled parts together: prefill and decode pools (the NIXL handshake checks versions and config), drafters with their targets, LoRA adapters with their bases.

Split TTFT into its parts: queueing, KV load or transfer, prefill compute, and frontend work such as tokenization, image preprocessing and grammar compilation. Then check, roughly in order:

1. Queueing: is the waiting queue longer? Fewer KV blocks after a memory change, or more preemptions, both cut capacity.
2. Cache: did the prefix hit rate drop? Look for a template change, such as a timestamp at the top of the prompt, a tokenizer change, or routing that lost session stickiness.
3. Scheduling: did token-budget or long-prefill defaults change, or did speculation turn on at high load?
4. Frontend: is the API server CPU-bound on long prompts, images or new schemas?
5. Disaggregation: is KV transfer slower, or has the P:D ratio drifted with traffic?

Use multi-LoRA on shared base weights. vLLM batches requests for different adapters in the same step with batched LoRA kernels and loads adapters on demand, capped by --max-loras on the GPU, --max-cpu-loras in host memory and --max-lora-rank. The LoRA ID is part of every block hash, so adapters never share KV by mistake. vLLM has since added fused MoE LoRA kernels for MoE bases.

Clarify first: traffic per model (requests/s, prompt and output length distributions, prefix sharing), SLO tiers (interactive TTFT and ITL versus batch deadlines), the GPU inventory and vendors, isolation and compliance needs, and how often models change.

Design:

• Edge: an Envoy gateway with auth, per-tenant quotas and priority classes behind one OpenAI-compatible API, with stable model aliases so backends can change.
• Placement by traffic: the top few models get dedicated vLLM pools with cache-aware routing; long-tail fine-tunes share bases through multi-LoRA; rarely used models share GPUs through sleep mode or scale to zero.
• Engine defaults: the smallest TP with KV headroom, prefix caching on, FP8 where evals pass, speculation only on interactive pools.
• Routing: llm-d's EPP picks the pod; add a semantic router only if teams want automatic model choice.
• Scaling and release: per-pool autoscaling on queue depth and KV usage, warm replicas for top models, and release gates for every engine or model change.

Trade-offs to name: dedicated pools isolate tenants but strand capacity, while shared pools raise utilization but need quotas and priorities. Multi-LoRA is cheap until adapter churn thrashes batches. P/D is worth it only for the one or two models with long prompts at scale.

What you would measure: SLO attainment per tier, GPU-hours per million tokens (cached and uncached separately), prefix hit rate, preemptions and cold-start time.

Likely follow-up: One tenant floods long prompts. What protects everyone else? Per-tenant token-rate quotas at the gateway, the long-prefill cap in the engine, a lower priority class for that tenant, and a separate pool if it persists.

Clarify first: the interactivity target (tokens/s per user), prompt and context lengths, daily token volume, and whether the hardware is NVL72-class racks or 8-GPU nodes on InfiniBand.

Design:

• Disaggregate: separate prefill and decode pools linked by NIXL over RDMA, rate-matched with saturation sweeps.
• Decode pool: wide EP with data-parallel attention so MLA latents are not replicated, DeepEP low-latency all-to-all, dual-batch overlap, EPLB with a few redundant hot experts, and MTP speculation for latency.
• Prefill pool: DeepEP high-throughput kernels, or PCP for very long prompts.
• Precision: block-scaled FP8 weights and an FP8 latent cache, validated on long-context evals.
• Routing: llm-d or Dynamo, cache-aware across DP ranks.

Trade-offs to name: a wider EP group gives more KV capacity and throughput, but more synchronization and a bigger failure domain. On NVL72-class systems, wide DEP scaled better than context parallelism once each rank held a few requests. The reference point is 2.2k tokens/s per H200 in a production-like llm-d deployment.

What you would measure: per-rank load imbalance, share of step time spent in all-to-all, P90 TPOT, and queue balance between the two pools.

Likely follow-up: One GPU in a 32-rank EP group fails. What happens? The whole group stalls, because every MoE layer needs every rank. You need fast health detection, routing away at the EPP, spare capacity to restart the group, and eventually elastic EP, which is on vLLM's roadmap.

Clarify first: session length, turns and subagent fan-out, the interactivity floor, the cost target per completed task, the model's attention architecture, and data-residency rules.

Design:

• Routing: session-sticky until a replica saturates, then spill to the best prefix overlap; the harness passes session IDs.
• KV tiers: GPU, CPU offloading, a shared Mooncake or LMCache pool, then disk, with interval and selective retention for hybrid models.
• Scheduling: cap long prefills, for example at 512 tokens per request per step, so short turns are never blocked; send first turns and later turns to differently configured pools.
• Parallelism: follow the architecture, such as DCP for MLA-style models on 8-GPU nodes and DEP on NVL72.
• P/D: size the pools with saturation sweeps, expecting decode-heavy ratios because most input is cached.

Trade-offs to name: stickiness against balance, retention storage against recompute, and the long request's TTFT against every short turn's latency.

What you would measure: hit rate by tier, tokens per GPU-second at a P90 interactivity floor, time to complete a task, and KV-pool bandwidth.

Likely follow-up: The CPU tier is full. What do you evict? Sessions least likely to return, judged by lifecycle hints and TTLs. Keep prompt-end checkpoints for active sessions, and consider FP8 KV to fit more before evicting.

A chat service peaks at 50 requests/s with 2,000-token prompts, half of them a shared, cached system prompt, and 500-token outputs. The target is 40 ms TPOT on Llama 3.1 70B with FP8 weights and FP8 KV on H100 80 GB. These are practice assumptions, not measurements.

1. Concurrency (Little's law): each request decodes for about 500 × 40 ms = 20 s, so about 50 × 20 = 1,000 sequences are in flight at peak.
2. KV capacity: FP8 KV is 160 KiB per token. The shared prompt is stored once per replica, so each sequence holds about 1,250 unique tokens (about 195 MiB), or about 190 GiB for 1,000 sequences.
3. Decode replicas: a TP=2 replica has 160 GB × 0.9 = 144 GB; minus 70 GB of weights and about 10 GB of activations, that leaves about 64 GB, roughly 300 sequences. Four replicas (8 GPUs) hold the working set.
4. Step-time check: each GPU reads about 35 GB of weights and 47 GB of KV per step, because every sequence still reads the shared prefix unless cascade attention applies. That is about 25 ms at 3.35 TB/s, or 33 ms at 75% of peak: inside 40 ms, with little slack.
5. Prefill compute: 50 × 1,000 uncached tokens × 2 × 70B ≈ 7 PFLOP/s. At about 900 effective FP8 TFLOPS per H100 (45% of peak), that is about 8 GPUs of compute.
6. Combine: folding that prefill into the decode steps would push them well past 40 ms, so plan about 16 GPUs, either aggregated or split into about 8 prefill and 8 decode GPUs. Burst and N+1 headroom takes it to about 20–24 GPUs, three 8-GPU nodes.

What makes it a staff answer: stated assumptions; decode sized by KV and step time, prefill by FLOPs; noticing that a shared prefix saves memory but not bandwidth; and closing with "then I would confirm it with a load sweep at these shapes".

On evidence at your operating point, not on reputation.

• Model coverage: how fast new open models run well; vLLM shipped day-0 support for DeepSeek-V3.2's sparse attention.
• Hardware breadth: vLLM covers NVIDIA and AMD GPUs, Google TPUs, AWS Trainium and Intel, plus vendor plugins; TensorRT-LLM is NVIDIA-only.
• Performance at your SLO: run your own traces. vLLM's Q3 2026 roadmap claims TensorRT-level performance on top models, and public benchmarks such as SemiAnalysis's include it.
• Ecosystem: llm-d, Dynamo, Ray Serve LLM and KServe build on vLLM; Dynamo can also run SGLang or TensorRT-LLM underneath.
• Governance and backing: vLLM began at UC Berkeley's Sky Computing Lab in 2023, went to the Linux Foundation in July 2024 and became PyTorch Foundation-hosted in 2025. Its creators founded Inferact in January 2026 with $150M in seed funding.

Upstream by default; never fork.

• Carry no patches you can avoid: each one taxes every upgrade, and vLLM ships a release every few weeks.
• Engage through the SIGs: Core, Large-Scale Serving, RL, Model Performance and Quantization each have a Slack channel and tracking issues; bring an RFC before a big change.
• Invest in CI: contribute tests and hardware coverage for the models and accelerators you depend on.
• Track direction: the quarterly roadmap issues and the Office Hours.

Frame it as cost per million tokens at a fixed SLO, with a staged plan:

1. Baseline: measure today's throughput-per-GPU versus interactivity curve and the cost per million tokens at your SLO.
2. Target: model the gain from published references. DeepSeek decode on H200 rose from about 1.5k to 2.2k tokens/s per GPU through runtime and kernel work, so the same load needs about 32% fewer GPUs.
3. Cost: engineering time, two pools to operate, RDMA networking for KV transfer, and new failure modes.
4. Plan: pilot one high-volume model with success and kill criteria agreed up front, then roll out pool by pool.

Treat hardware as a validation and routing problem:

• Validation matrix: model × hardware × feature, with accuracy and performance gates per cell.
• Portable layers: vLLM's Triton attention backend (one kernel source across NVIDIA, AMD and Intel, default on ROCm), torch.compile and hardware plugins.
• Placement: route each model or traffic class to the hardware where it is validated and cheapest per token at SLO; llm-d has written up serving across three GPU vendors in one cluster.
• Trade-off: procurement leverage and supply resilience against the cost of running several driver, kernel and debugging stacks.

RL generates rollouts from a policy that changes every step, so the engine becomes part of the training loop:

• Fast weight sync from trainer to inference workers without restarts; vLLM's Q1 2026 roadmap lists modular weight-sync APIs for its RL SIG.
• Memory handoff: sleep mode frees or offloads weights and KV so training and rollouts can share GPUs.
• Cache invalidation: after a weight update, cached KV was computed with old weights, so the prefix cache must be reset.
• Determinism and batch invariance: the same prompt should get the same logprobs whatever its batchmates; it is on vLLM's roadmap, and the Triton backend already supports it.
• Scheduling: rollouts are long, batch-like generations; llm-d's co-operative time-slicing reports 40% more post-training experiments on the same GPUs.