vLLM: An Exhaustive Technical Study for Kernel Engineers

Audience: engineers building a competing Go inference kernel (fak). This is a mechanics-first study, grounded against the fak repository where the two systems touch. Claims are graded against the verification that produced this document; where a detail could not be independently confirmed (often because the primary source was fetch-blocked during research), it is hedged as “per vLLM’s design docs” or “reportedly” rather than stated as settled fact.

1. Executive overview

vLLM is a high-throughput LLM inference and serving engine whose entire design descends from one observation: the two things an autoregressive transformer server must manage — memory and compute — fail in two different, separable ways, and each failure has a distinct fix borrowed from operating systems.

Problem one is memory. The KV cache (the per-token keys and values every attention layer must retain) is large, grows one token at a time, and has a lifetime nobody can predict in advance (you do not know how long a generation will run). Classical serving systems allocated one contiguous KV tensor per request, sized to the worst case. That produces catastrophic waste: internal fragmentation from over-reservation, external fragmentation from variable-length contiguous allocations, and no way to share identical prefix KV between requests. The official vLLM blog reports existing systems wasted 60–80% of KV memory to fragmentation and over-reservation (https://blog.vllm.ai/2023/06/20/vllm.html). vLLM’s answer is PagedAttention: treat KV like OS virtual memory — fixed-size blocks (pages), tokens (bytes), sequences (processes) — with a per-request block table mapping logical block indices to physical blocks scattered anywhere in a global GPU pool. This drops waste to under 4% and enables refcounted, copy-on-write sharing of common prefixes.

Problem two is compute. Requests arrive and finish continuously and at wildly different lengths, so any scheme that batches at request granularity (wait for a batch, run it to completion, return) leaves the GPU idle whenever fast requests finish before slow ones. vLLM’s answer is iteration-level (continuous) scheduling: the scheduler decides, at every single forward step, which requests contribute tokens — admitting new work and retiring finished work between steps rather than between batches.

The V1 re-architecture (alpha announced 2025-01-27, https://blog.vllm.ai/2025/01/27/v1-alpha-release.html) is a ground-up rewrite that took both ideas further. It isolates the EngineCore (scheduler + model executor) in its own process with heavy multiprocessing and busy-loop workers; it fuses prefill and decode into a single unified scheduler that can put thousand-token prompt chunks and single-token decodes in the same forward pass (chunked prefill, on by default); it makes automatic prefix caching default-on; it reorganizes the block manager into a KVCacheManager → KVCacheCoordinator → BlockPool stack; and it changed several V0 defaults (notably: preemption now recomputes rather than swaps to CPU). The unified mixed batch is precisely what forced the adoption of a variable-length paged attention kernel (FlashAttention 3) as, in the authors’ words, the “final piece of the puzzle.”

The rest of this document dissects each subsystem. A recurring theme for the fak reader: vLLM’s KV management is paged and LRU; fak’s is tree-structured (RadixAttention-style) and value-aware. Those are the two axes on which the systems genuinely diverge, and Section 3 maps them file-by-file.

2. Subsystem deep-dives

2.1 PagedAttention & the block manager

How it works. The KV cache is a global pool of fixed-size physical blocks; each block holds block_size tokens’ worth of K and V for one layer group. A request’s KV is not a contiguous tensor but a list of block ids recorded in a per-request block table that maps contiguous logical block indices to physical block ids scattered anywhere in the pool. This is a direct OS virtual-memory analogy — the vLLM blog states it verbatim: “one can think of blocks as pages, tokens as bytes, and sequences as processes” (https://blog.vllm.ai/2023/06/20/vllm.html; SOSP’23, Kwon et al., arXiv:2309.06180). (Confirmed.)

Allocation is on-demand. For N tokens the system allocates exactly ceil(N/block_size) blocks. This bounds internal fragmentation to at most the last, partially-filled block per sequence (≤ block_size − 1 tokens; 15 with the legacy default of 16), and eliminates external fragmentation entirely because all blocks are the same size and interchangeable. Empirically this is “near-optimal memory usage, with a mere waste of under 4%” versus the 60–80% of prior contiguous systems. (Confirmed against the vLLM blog and SOSP’23 paper. Note one real-world nuance: vLLM reserves a slot for the next generated token and shares blocks via prefix caching/CoW, so the literal per-request count can differ in those cases.)

Sharing is refcounted with copy-on-write. Physical blocks carry a reference count. Parallel-sampling and beam-search sequences from a common prompt point their logical blocks at the same physical blocks, so the shared prompt KV is stored once; a block is duplicated only when a sequence writes to a block whose refcount is still > 1. vLLM “keeps track of the reference counts of the physical blocks and implements the Copy-on-Write mechanism,” reducing the memory overhead of complex sampling by up to 55% for ~2.2× throughput on those workloads (paper-reported figures; finer-grained: ~6–10% for parallel sampling, ~38–55% for beam search on OPT-13B/Alpaca). (Confirmed. The beam-search specifics are somewhat historical — later vLLM refactored beam search out of the block-manager sharing path.)

block_size is no longer a fixed 16. Historically the default was 16 with CLI choices restricted to --block-size {8,16,32}. In current V1 there is no static default; it is set per-platform in Platform.check_and_update_config(): CUDA caps at 32, HPU defaults to 128, and on Neuron it is forced to --max-model-len. (Confirmed against docs.vllm.ai cache config and historical engine_args docs.) The CUDA cap of 32 ties to FlashAttention/PagedAttention paged-block-size kernel constraints. (The kernel-constraint rationale is sound; the specific GitHub issue commonly cited for it — #2331 — appears to be a misattribution; the substantive “why is CUDA block_size capped at 32” discussion is vllm-project/vllm#14319. Treat the number as correct, the issue reference as uncertain.)

The V1 manager stack (per vLLM’s V1 design docs — architectural details below were not independently re-verified against pinned source in this research pass; treat as design-doc-accurate rather than line-confirmed):

V1 block tables are append-only. This is a genuine V0→V1 semantic change: in V0, on detecting a duplicate block, vLLM would free it and repoint the request at the shared block; V1 forbids rewriting an existing block-table entry, so a fully-computed block is immutable-by-convention and new tokens always land in freshly-allocated blocks. Sharing in V1 is handled purely through the block-hash → block map plus ref_cnt at allocation time. (Presented in the research as a V1 design fact but not independently re-verified against pinned source; consistent with the append-only summary throughout.)

Reaching the kernel. Block tables reach the attention kernel via the worker-side BlockTable (vllm/v1/worker/block_table.py), which supplies a block_table tensor plus a slot_mapping tensor inside each backend’s attention metadata, where a write slot is device_block_number * block_size + block_offset. There is one BlockTable per KV cache group; hybrid Mamba+attention models keep a separate pool per group because the page size (block_size × num_hidden_layers × kv_hidden_size) differs across attention types (see the Hybrid KV Cache Manager design, https://docs.vllm.ai/en/stable/design/hybrid_kv_cache_manager/). (Design-doc-level; a single-pool mental model breaks for hybrid architectures.)

Preemption defaults reversed V0→V1. Under KV pressure V0 could preempt by swap (move blocks to a CPU pool sized by --swap-space GiB/GPU, then swap back) or recompute (discard GPU blocks, rerun prefill on prompt+generated tokens as one larger prompt), selected by preemption_mode. V1 reportedly sets CPU blocks to 0 and always recomputes by default to avoid PCIe traffic. (Both preemption claims are marked unverified in the research — sourced from a secondary deep-dive citing vllm/v1/engine/core.py, not the file directly. Present as “reportedly.” Note --swap-space and --cpu-offload-gb are distinct knobs; --cpu-offload-gb defaults to 0.)

Key knobs: --block-size, --swap-space, --cpu-offload-gb, --gpu-memory-utilization (governs how much VRAM the KV pool may claim), --enable-prefix-caching (default-on in V1), --max-model-len (sets max_num_blocks_per_req = cdiv(max_model_len, block_size)).

Gotchas: don’t hardcode 16 — it’s a legacy default. A partially-filled trailing block carries all the internal fragmentation and cannot be prefix-shared until it fills (its hash isn’t assigned until full). Reasoning about block_table alone omits slot_mapping, which is where the current step’s K/V actually land.

2.2 The V1 scheduler & continuous batching

The one-line mental model. In vLLM V1 the scheduler emits, each engine step, a single map num_scheduled_tokens: dict[request_id -> int] — “process exactly this many tokens for this request this pass” — and nothing else about phase. There is no “prefill batch” vs “decode batch”: prompt tokens and generated tokens are the same currency, and one forward pass can carry a chunk of one request’s prefill alongside single-token decodes of a dozen others. The official post states it plainly: V1 “removes the traditional distinction between ‘prefill’ and ‘decode’ phases … scheduling decisions are represented as a simple dictionary, e.g. {request_id: num_tokens},” general enough “to support features such as chunked prefills, prefix caching, and speculative decoding” (https://blog.vllm.ai/2025/01/27/v1-alpha-release.html; “Anatomy of vLLM,” https://blog.vllm.ai/2025/09/05/anatomy-of-vllm.html). Code lives in vllm/v1/core/sched/scheduler.py; schedule() returns a SchedulerOutput the ModelRunner consumes. (All five key claims in this section were independently verified as confirmed.)

The per-step schedule() loop. Each EngineCore step drains the input queue into the scheduler, then schedule() walks the two queues under a fixed token budget. It is running-first (decode-first): the running list is advanced first to keep in-flight generation moving, then the waiting queue is admitted only if budget remains. Each request tracks num_computed_tokens and a target num_tokens_with_spec; the scheduler hands out min(remaining_to_catch_up, remaining_budget, chunk_cap) tokens per request until waiting empties or the budget is exhausted. Crucially, KV allocation happens inside the scheduling decision, not after: schedule() calls KVCacheManager.allocate_slots(...) while deciding the step, rather than forming a batch and hoping workers can fit it. The resulting SchedulerOutput names new requests, cached/resumed requests, per-request num_scheduled_tokens, allocated KV blocks, preempted requests, and finished ids for worker-side cleanup (https://www.ubicloud.com/blog/life-of-an-inference-request-vllm-v1).

The two budgets: max_num_batched_tokens and max_num_seqs. max_num_batched_tokens is the total tokens the scheduler may pack into one forward pass (the sum over the num_scheduled_tokens dict); max_num_seqs caps the number of sequences in the batch. A related field max_num_scheduled_tokens is documented as “usually equal to max_num_batched_tokens, but can be smaller … when the model might append tokens into the batch (such as speculative decoding)” (https://docs.vllm.ai/en/latest/api/vllm/config/scheduler/). Reportedly max_num_seqs has no static default in recent V1 — if unset it is derived in EngineArgs.create_engine_config from the usage context. max_num_batched_tokens defaults are version-dependent (reportedly 512 under chunked prefill in v0.6.0, 2048 by v0.8.2), so treat any single number as a moving target (https://docs.vllm.ai/en/stable/configuration/optimization/).

Queues and policy — FCFS deque vs. priority heap. The scheduler holds a waiting queue (new requests plus preempted ones awaiting resume) and a running list. Policy is chosen by SchedulerConfig.policy / the --scheduling-policy CLI flag, swapping the queue class (vllm/v1/core/sched/request_queue.py): FCFSRequestQueue is a collections.deque subclass where popleft() enforces strict arrival order (the default "fcfs"); PriorityRequestQueue is a heapq-backed min-heap keyed on the tuple (priority, arrival_time) — lower priority value is served first, earlier arrival_time breaks ties (https://docs.vllm.ai/en/latest/api/vllm/v1/core/sched/request_queue/). Note the sign convention: a smaller priority number is more important. Priority scheduling in the V1 engine is comparatively recent (PR #19057) and has a known gap — when running is already at max_num_seqs, a high-priority waiting request reportedly cannot preempt to force its way in (issue #40004).

Mixing prefill and decode; chunked prefill. Because there is no phase flag, one batch freely mixes a compute-bound prefill chunk with many memory-bound decodes — exactly the throughput win V1 targets (co-locating compute- and memory-bound work in one pass). Chunked prefill is on by default whenever possible; the scheduler batches all pending decodes first, then fills leftover budget with prefill tokens, chunking the last prefill that doesn’t fit. The cap is long_prefill_token_threshold: per the Anatomy post, “cap the number of new tokens per step. If the requested number exceeds long_prefill_token_threshold, reset it to exactly that value.” Two companion knobs shape partial-prefill concurrency: max_num_partial_prefills (max sequences partially prefilling at once) and max_long_partial_prefills (max prompts longer than the threshold prefilling concurrently — setting it below max_num_partial_prefills lets short prompts jump ahead of long ones for latency) (https://docs.vllm.ai/en/latest/api/vllm/config/scheduler/).

Preemption under KV pressure — allocate_slots → None. When the KV cache cannot hold every running request’s next step, the scheduler preempts. The signal is mechanical: KVCacheManager.allocate_slots(...) returns None on allocation failure, and the scheduler evicts a victim, frees its blocks, and re-attempts. The victim is moved to the front of the waiting queue with status PREEMPTED. Victim selection is policy-dependent: under FCFS it is effectively LIFO — the most-recently-scheduled running request is evicted first — while under priority it evicts the request with the largest (priority, arrival_time) (lowest importance / newest) (https://deepwiki.com/vllm-project/vllm/2.5-request-scheduling). (This claim verified at medium confidence: the exact returns None wording traces to a secondary deep-dive rather than a direct source quote — confirm against scheduler.py before hard-coding.)

Recompute vs. swap — V1’s default is recompute, CPU blocks = 0. V0 offered two recovery modes for a preempted request: SWAP (copy its KV blocks to a CPU-side allocator over PCIe and back) and RECOMPUTE (discard the KV and re-prefill). V1 defaults to RECOMPUTE because it has lower overhead in the V1 architecture, so the CPU block allocator / swap_space are effectively 0 — there is no host-side KV mirror; a preempted request simply re-runs prefill when readmitted (https://docs.vllm.ai/en/stable/configuration/optimization/). The user-facing symptom is the warning that a request was “preempted by PreemptionMode.RECOMPUTE … not enough KV cache space,” with the documented remedy being to raise gpu_memory_utilization or tensor_parallel_size. (Beam-search / multi-sequence groups reportedly still fall back to swap in V0 — not the V1 path.)

The V0 → V1 architectural break. V0 modeled each iteration as either a prefill step or a decode step (a persistent-batch, prefill-XOR-decode model), which forbade mixing and complicated chunked prefill. V1 deletes that step model in favor of the unified token-budget scheduler, and relocates scheduling into a dedicated EngineCore process: a busy loop that pulls from an input queue and runs one “step” (schedule + one model forward pass) per iteration. EngineCore runs in a separate process from the API front-end / AsyncLLM, communicating over ZeroMQ, so tokenization/detokenization/HTTP overlap GPU work and dodge the GIL (https://blog.vllm.ai/2025/01/27/v1-alpha-release.html; https://developers.redhat.com/articles/2025/01/28/vllm-v1-a-major-upgrade-vllms-core-architecture). The dynamism of one-pass prefill+decode also forced a kernel change — V1 leans on FlashAttention 3 to handle heterogeneous batches.

Prefix caching, spec decode, structured output — all expressed through the same dict. These are not separate scheduler modes; they are just different ways the per-request token count is computed. Prefix caching: KVCacheManager.get_computed_blocks(...) returns cached-hit blocks, so num_computed_tokens jumps forward and the scheduler simply schedules fewer new tokens for that request (in V1 prefix caching is on by default and near-zero-overhead). Speculative decoding: the draft proposes k tokens verified in one target pass, so a request’s target becomes num_tokens_with_spec (> 1 token/step) — the same budget arithmetic, which is exactly why max_num_scheduled_tokens can drop below max_num_batched_tokens (early V1 supported ngram spec decode first, with Eagle/MTP next). Structured output: a StructuredOutputManager compiles the grammar (e.g. via xgrammar) into an FSM, produces a _grammar_bitmask tensor once per step, broadcasts it to workers (instead of recomputing per-worker as in V0), and masks disallowed logits to −∞ before sampling — deliberately off the critical path (https://blog.vllm.ai/2025/01/14/struct-decode-intro.html).

Gotchas / version notes. (1) The Anatomy post pins to commit 42172ad (≈2025-08-09); field names like num_tokens_with_spec, long_prefill_token_threshold, and the request_queue.py split are stable around v0.9–v0.11 but predate that in flux. (2) max_num_seqs has no fixed default (usage-context derived); max_num_batched_tokens defaults shifted across releases (512 → 2048 → higher) — never assume a literal. (3) Priority policy’s (priority, arrival_time) uses smaller = more important; easy to invert. (4) The running-first ordering means a flood of cheap decodes can starve new prefills under a tight max_num_batched_tokens; conversely too-large a budget hurts ITL. (5) A reported prepend-on-skip ordering bug in schedule() can shuffle waiting-queue order and hurt tail latency (issue #27441).

2.3 Automatic prefix caching & KV reuse

How it works. V1 automatic prefix caching (default-on in V1; off by default in V0) reuses KV of shared prompt prefixes at the granularity of fixed-size token blocks, indexed by a chained block hash. Each full block’s key is computed by hash_block_tokens(hash_function, parent_block_hash, curr_block_token_ids, extra_keys) — i.e., over the parent block’s hash, the block’s token ids, and optional extra keys. Chaining forces an exact token-by-token match of the entire prefix up to and including a candidate block for a hit. Only full blocks are cached/hashed; a block’s hash is assigned when it fills and reset on eviction. (Confirmed against the Automatic Prefix Caching design doc, https://docs.vllm.ai/en/stable/design/prefix_caching/.)

extra_keys carry the LoRA id, a per-request cache_salt (injected into the first block’s hash for multi-tenant isolation), and multimodal image hashes (mm_hash). (Confirmed as the set of extra keys.)

Data structures. The V1 KV cache manager pre-allocates a pool of KVCacheBlock objects, each with an immutable block_id, a block_hash (set-when-full / reset-on-evict), a ref_cnt, and prev_free_block/next_free_block pointers forming a doubly-linked FreeKVCacheBlockQueue kept in LRU order (LRU block at the head, popped on eviction). A separate map — cached_block_hash_to_block on the BlockPool — supports O(1) hash lookup. The pool is pre-allocated “to avoid Python object creation overhead.” (Confirmed; note the actual identifier is cached_block_hash_to_block, and it is a nested block_hash → {block_id → KVCacheBlock} map, not a flat hash→block.)

Lookup and eviction. Lookup walks block hashes in order; each hit “touches” the physical block and bumps ref_cnt; the first miss triggers fresh allocation for the remainder — so identical prefixes dedup to the same physical block. When a request finishes, its ref_cnt == 0 blocks are appended to the free-queue tail in reverse order (the last, longest-prefix block is least reusable, so it’s evicted first); allocation under pressure pops the LRU head; a ref_cnt > 0 block is immune.

Hit accounting is token-granular — but read the metric carefully. The manager records prefix_cache_stats.record(num_tokens=request.num_tokens, num_hits=num_new_computed_tokens) (PrefixCacheStats in vllm/v1/metrics/stats.py), surfaced as the cumulative counters vllm:prefix_cache_queries and vllm:prefix_cache_hits. Correction: the older gauge vllm:gpu_prefix_cache_hit_rate is not co-emitted with these counters in V1 — it was deprecated and replaced by the counters (deprecated 0.8, hidden 0.9, removal ~0.10). vLLM does not report a cumulative hit-rate directly; operators derive a windowed rate via PromQL, e.g. rate(vllm:prefix_cache_hits[5m]) / rate(vllm:prefix_cache_queries[5m]). A naive all-time hits_total/queries_total ratio does climb toward 1 over repeated identical requests — that dilution is precisely the anti-pattern the counter design (PR #12592) moved away from. (This is the corrected form of a partially-correct claim; see https://docs.vllm.ai/en/latest/design/metrics/ and PR #12592.)

Hash algorithm is selectable via --prefix-caching-hash-algo: builtin (Python hash(), per-process randomized by PYTHONHASHSEED, fast), sha256 (collision-resistant, pickle-serialized), and sha256_cbor / sha256_cbor_64bit (reproducible, cross-language-stable via CBOR). Reportedly sha256 is the current stable default, with builtin the earlier default (~v0.8.3). (The selectability and option set are documented; the default-flip and the ~100–200 ns/token cost figure are marked unverified — hedge.)

Cascade attention (V1, reportedly PR #11635, on by default under V1) is the compute-side complement: when every request in a batch shares one common prefix, it computes the shared-prefix attention once and merges it with each request’s suffix attention via FlashAttention/FlashInfer log-sum-exp renormalization (metadata: common_prefix_len, use_cascade). It reportedly does not yet support a forest of distinct prefixes in one batch — mixed batches fall back to normal paged attention. (Marked unverified — hedge; derives from FlashInfer’s Cascade Inference technique.)

Multimodal prefix caching (reportedly PR #11187, simplified in #11646 to drop a redundant offset) folds a per-image mm_hash into extra_keys so image tokens are keyed by pixel content, not just placeholder token ids; it requires the mm-cache preprocessor to supply image hashes. (Unverified — hedge.)

Contrast with SGLang RadixAttention. vLLM’s fixed-block chained hashing needs block-boundary alignment for a hit; SGLang’s RadixAttention (arXiv:2312.07104) indexes KV in a token-level radix tree that discovers arbitrary-length shared prefixes without alignment. Both evict by LRU. This distinction is directly load-bearing for fak, whose radixkv is the SGLang-style design (§3).

Gotchas: a 207-token prompt at block_size=16 caches only 192 tokens (12 full blocks); the trailing 15 never form a block and are always recomputed — so per-request hit rate caps below 100% and is block-size-dependent. builtin hashing is per-process-randomized and even sha256 (pickle-based) is not reproducible across Python/vLLM versions; only the sha256_cbor variants are deterministic cross-process — relevant if you want host-stable prefix keys. With weak hashing, distinct prefixes can theoretically collide and serve wrong KV, so multi-tenant deployments should use sha256 and/or a per-request cache_salt. (A token-sequence radix match, as fak uses, is collision-free by construction — a genuine correctness edge over hash-keyed lookup.)

2.4 Attention kernels & backends

The seam. In V1 the attention layer is where the unified mixed batch meets the GPU. Because a single step can contain thousand-token prefills and single-token decodes together, the kernel must be a variable-length (varlen) paged kernel handling arbitrary per-request query/KV lengths against a paged KV cache.

FlashAttention 3 was the “final piece.” The V1 alpha post states FA3 integration “was the final piece of the puzzle” and that “given the high level of dynamism in V1 (e.g., combining prefill and decode within the same batch) a flexible and high-performance attention kernel was essential.” FA3 support was added by Lucas Wilkinson (PR #12093), and it handles both prefill and decode tokens of standard decoder models in one unified flash_attn_varlen_func call. (Confirmed against https://blog.vllm.ai/2025/01/27/v1-alpha-release.html and PR #12093.)

Backend selection is automatic and overridable. vLLM tries FlashAttention first and lets you override with --attention-backend (env VLLM_ATTENTION_BACKEND), where the string "auto" normalizes to None (auto-select). Per-architecture defaults differ: FlashAttention is the default on Hopper (SM90, H100/H200); FlashInfer is the default on Blackwell (SM100, B200/B300) (opt into FlashInfer on Hopper with VLLM_ATTENTION_BACKEND=FLASHINFER). On Blackwell the auto fallback order is TRT-LLM Ragged → FlashInfer → MLA path. The FlashAttention major version is itself configurable via --attention-config.flash_attn_version, with recent builds defaulting to FA4 on SM100+, FA3 on SM90, FA2 otherwise. (Confirmed against https://docs.vllm.ai/en/latest/design/attention_backends/. These defaults are version- and hardware-sensitive — pin any “X is the default” statement to a vLLM version + SM level.)

The backend abstraction. Each backend has an Impl (the kernel call) plus a metadata builder that turns the scheduler’s per-step plan into device tensors: slot_mapping, block tables, query_start_loc (prefix-sum of scheduled tokens — e.g. [3,2,5] → [0,3,5,10]), seq_lens, num_computed_tokens, and num_prefills/num_prefill_tokens/num_decode_tokens. The paged block table is address translation (logical block → physical device block); the per-token write slot is device_block_number * block_size + block_offset. (The block-table/slot-mapping abstraction is confirmed at the PagedAttention level; the exact V1 metadata field list is design-doc-level.)

FlashInfer is a broader operator library that consumes vLLM’s paged block tables directly, with no copy into contiguous memory, via BatchPrefillWithPagedKVCacheWrapper / BatchDecodeWithPagedKVCacheWrapper (it models PagedAttention as a block-sparse layout and passes only CSR-style metadata — paged_kv_indices/paged_kv_indptr — not KV payloads). It is the natural home for block-sparse and MLA attention; on Blackwell it can route through NVIDIA TRT-LLM attention kernels (which support attention sinks — relevant to gpt-oss). (Confirmed against the attention-backends doc and PR #14061.) Correction on the disable knob: TRT-LLM attention can be disabled either via the long-standing env var VLLM_USE_TRTLLM_ATTENTION=0 (with _DECODE_/_CONTEXT_ variants) or the newer config flag --attention-config.use_trtllm_attention=0; both are valid.

The Triton “unified attention” backend is ~800 lines implemented entirely in Triton, native to vLLM, depends only on PyTorch + Triton, runs the same source across NVIDIA/AMD/Intel GPUs, and is always importable — so it is the universal fallback when FlashAttention/FlashInfer are unavailable, and the default on AMD ROCm. Its GQA optimization processes all query heads sharing one KV head together and groups multiple query tokens into a single “Q block” work item. (Confirmed against the 2026-03-04 Triton backend deep-dive, https://blog.vllm.ai/2026/03/04/vllm-triton-backend-deep-dive.html. One sub-detail — that Intel XPU uses Triton for fp32 because Flash Attention lacks fp32 there — is plausible and consistent with vLLM’s XPU platform but was not independently confirmed; treat it as unverified.)

MLA is a separate, non-unified path. For DeepSeek-V2/V3-family Multi-head Latent Attention, KV is low-rank-compressed into a single shared latent vector cached per token and attended MQA-style on that latent, while multi-head attention is “simulated” during the compute-heavy prefill. Shared logic lives in MLACommonImpl (splits the query into q_nope no-positional and q_pe positional components, updates the latent KV cache, output-projects). Crucially, MLA uses separate prefill and decode backends — the “one unified FA3 kernel” story does not extend to DeepSeek-family models. The optimized FlashMLA backend (Hopper/Blackwell) supports standard and FP8 KV caches and strictly requires block_size=64 (vs the common 16), so MLA models cannot share the standard decoder path’s block size — a real KV-cache-manager constraint. A portable TRITON_MLA backend also exists. MLA reportedly reduces KV-cache memory dramatically (DeepSeek’s own claim: up to ~93% vs MHA). (The MLA structural claims are marked unverified — sourced via search summaries; hedge. The specific latent dimension often cited — kv_lora_rank=512 + 64-dim decoupled RoPE = 576 cached scalars/token — was not confirmed; verify against the model config before relying on it.)

DeepSeek V4 sparse MLA (DSA) adds its own decode/prefill backends (FLASHMLA_SPARSE_DSV4, FLASHINFER_MLA_SPARSE_DSV4), selected by --attention-backend. Reportedly selection depends on KV dtype: FP8 KV always prefers FLASHINFER_MLA_SPARSE; with BF16 KV, FLASHINFER_MLA_SPARSE is preferred for low query-head counts (≤16) and FLASHMLA_SPARSE otherwise. (Version-sensitive, 2026-era; hedge.)

Gotchas: “FlashAttention handles unified prefill+decode” is true only for standard decoder attention — MLA splits the phases. Backend defaults are a moving target across releases and GPU generations (FA2 pre-Hopper → FA3 Hopper → FA4 Blackwell; FlashAttention-default Hopper vs FlashInfer-default Blackwell). FlashMLA’s block_size==64 is a hard cache-manager constraint, not a tuning knob.

2.5 Distributed execution & parallelism (TP/PP/DP/EP/EPLB)

vLLM (V1 era) composes four orthogonal parallelism axes over a centralized-scheduler control plane.

Tensor parallel (TP) — --tensor-parallel-size/-tp, default 1 — uses Megatron-LM’s tensor-parallel algorithm to shard attention heads and FFN matrices intra-node, synchronizing with all-reduce; convention TP = GPUs-per-node. Pipeline parallel (PP) — --pipeline-parallel-size/-pp, default 1 — splits the model by layer stages across nodes and, uniquely, supports uneven layer splits, so it is the recommended axis when the GPU count doesn’t divide the model evenly (set TP=1, PP=#GPUs); convention PP = #nodes (e.g. 16 GPUs over 2 nodes → TP=8, PP=2). Data parallel (DP) — --data-parallel-size/-dp — runs independent engine-core replicas, each with its own KV cache, coordinated for multi-node via --data-parallel-size-local, --data-parallel-address, --data-parallel-rpc-port, --data-parallel-start-rank. (The TP/PP/DP mechanics are marked unverified in the research extraClaims — sourced from distributed_serving docs; the Megatron algorithm and uneven-PP-split are well-established, so state those and hedge the exact convention framing.)

Expert Parallel (EP) — --enable-expert-parallel — layers on top for MoE. Confirmed: EP width = data_parallel_size × tensor_parallel_size; enabling EP shards the MoE experts across all EP ranks while the router/gate runs replicated on each rank (attention/non-expert weights are replicated). The DeepSeek-V3 single-node example --tensor-parallel-size 1 --data-parallel-size 8 --enable-expert-parallel yields exactly 8 EP ranks. (Confirmed against https://docs.vllm.ai/en/latest/serving/expert_parallel_deployment/ and the fak runbook. Two nuances: EP only activates when TP×DP > 1, and the clean “replicated on each rank” picture holds at TP=1 — with TP>1, attention is TP-sharded within each DP group.)

Expert Parallel Load Balancer (EPLB) — --enable-eplb — fights the highly-skewed token routing of auxiliary-loss-free MoEs (DeepSeek-V3/R1). It collects a sliding window of per-physical-expert token-count load and periodically runs rebalance_experts (an algorithm adapted from DeepSeek’s EPLB) to recompute physical↔logical expert placement, optionally with redundant (duplicated) experts for both balance and fault tolerance. Confirmed config surface: EPLB is configured via --eplb-config (a JSON string) or dotted --eplb-config.<key> args; ParallelConfig holds eplb_config: EPLBConfig = Field(default_factory=EPLBConfig). Real keys include window_size, step_interval, num_redundant_experts, log_balancedness, policy, use_async (and also log_balancedness_interval, communicator). The documented example is vllm serve Qwen/Qwen3-30B-A3B --enable-eplb --eplb-config '{"window_size":1000,"step_interval":3000,"num_redundant_experts":2,"log_balancedness":true}'. The old standalone flags (--num-redundant-experts, --eplb-window-size, --eplb-step-interval) were consolidated under --eplb-config and are deprecated (removal ~v0.12.0). Observed defaults: window_size=1000, step_interval=3000, num_redundant_experts=0, log_balancedness=False, policy='default', use_async=True (latest; use_async defaulted False in older v0.20.1). (Confirmed against the EP deployment doc, the parallel config API, and PR #18343.)

The balancedness metric — corrected. log_balancedness is off by default because logging it requires cross-rank communication overhead (confirmed: EPLBConfig default log_balancedness=False). The metric is a mean/max load ratio (1.0 = perfectly balanced), but vLLM computes it across EP ranks, not across individual experts: each rank’s load is the summed token count of the physical experts assigned to it, and balancedness = mean(per-rank load) / max(per-rank load), summed across MoE layers. Load derives from per-expert token counts, but the mean/max aggregation axis is ranks (GPUs). (This corrects a partially-correct claim; verified against vllm/distributed/eplb/eplb_state.py. The fak runbook’s own “across experts” wording is imprecise on this point.)

The load window — confirmed. EPLB records a sliding window of expert load, shape (window_size, num_moe_layers, num_physical_experts) — recording all physical experts, not just local ones (changed via PR #22167). step_interval controls rearrangement cadence; if step_interval > window_size, only the last window_size steps’ metrics are used. (Confirmed verbatim against v0.11.0 source: expert_load_window docstring and EPLBConfig.step_interval.)

Rearrangement and redundancy (reportedly). The entry point is rebalance_experts(weight, num_replicas, num_groups, num_nodes, num_gpus) → (Tensor, Tensor, Tensor), producing the physical↔logical index maps; weights are then rearranged in place. Redundant experts let each routed expert keep multiple parameter copies across ranks (initial arrangement [original routed experts, redundant experts]) for finer balance and fault tolerance, with memory overhead NUM_MOE_LAYERS × BYTES_PER_EXPERT × (NUM_TOTAL_EXPERTS + NUM_REDUNDANT_EXPERTS) / NUM_EP_RANKS (~2.4 GB for one redundant expert/rank on DeepSeek-V3, per PR #18343). (Both marked unverified — hedge.)

Executor/worker hierarchy (reportedly). EngineCore hosts the Scheduler plus a Model Executor that progresses UniProcExecutor (single worker/GPU) → MultiProcExecutor (one worker process per GPU). Worker init reportedly has three phases: Init Device (assign CUDA device, check dtype/VRAM via gpu_memory_utilization, build model_runner + InputBatch), Load Model, and Initialize KV Cache (per-layer KV spec + VRAM profiling to compute allocatable KV blocks). (Per the Anatomy blog; unverified — hedge.)

Distributed backend selection (reportedly). --distributed-executor-backend ∈ {mp, ray, uni, external_launcher}. Multiprocessing (mp) is the single-node default; Ray for multi-node. Auto-selection uses mp “when not running in a Ray placement group and if there are sufficient GPUs available on the same node for the configured tensor_parallel_size, otherwise Ray.” (Per parallelism_scaling docs; unverified — hedge.)

Gotchas: confirm EPLBConfig numeric defaults from vllm/config/parallel.py, not the docs example (1000/3000/2 are illustrative). The balancedness axis is ranks, not experts — a per-GPU balance number, not a per-expert one.

2.6 Disaggregated prefill & KV transfer

The idea. V1 disaggregates the two phases — compute-bound prefill (prompt → full KV, no tokens emitted) and memory-bandwidth-bound decode (autoregressive generation) — onto physically separate GPU pools connected by a KV-transfer layer. This eliminates prefill→decode interference (the head-of-line blocking that spikes inter-token latency), trading it for extra prefill queuing.

The V1 seam: KVConnectorBase_V1 (vllm/distributed/kv_transfer/kv_connector/v1/base.py, introduced by PR #15960, ApostaC). Each connector splits into a scheduler-side and a worker-side role selected by KVConnectorRole (SCHEDULER/WORKER). Confirmed method set:

The KV exchange is bracketed around the forward pass: start_load_kv on entry (load external KV into paged memory for decode), wait_for_save on exit (block until KV uploaded for prefill). Layer-by-layer save/wait enables pipelined transfer. (Confirmed against https://docs.vllm.ai/en/stable/api/vllm/distributed/kv_transfer/kv_connector/v1/base/ and base.py on main. The int|None return — async remote lookup — was broadened by PR #23620; an older v0.10.1 snapshot may show int.)

NixlConnector is the primary production transport. Confirmed: it uses NVIDIA NIXL (Inference Xfer Library, open-sourced at GTC 2025) moving KV over RDMA/UCX (InfiniBand, RoCE), TCP, NVMe-oF, and S3; decode PULLS KV blocks from prefill (pull model); it is “fully asynchronous send/receive”; it uses an out-of-band ZMQ handshake side channel (background thread _nixl_handshake_listener) to exchange agent handles, block counts and lengths once per P–D pair; HND is the default KV layout for non-MLA models (NHD works but disables heterogeneous-TP head splitting; MLA is exempt — KV replicated, no head split); and a cache_dtype compatibility hash is enforced at handshake (disable via kv_connector_extra_config’s enforce_handshake_compat=false, or VLLM_USE_TRTLLM_ATTENTION-style env). Heterogeneous TP: when P’s TP > D’s TP, one D worker pulls from multiple P workers (PR #22663). (Confirmed against the NixlConnector compatibility matrix + usage docs; the specificity of the exact thread name and config key is itself evidence of grounding. NVMe-oF/S3 are NIXL-library backend capabilities; the production PD hot path in practice runs over UCX/RDMA.)

Other connectors (reportedly): LMCacheConnectorV1 (+ LMCacheMPConnector standalone server; uses NIXL underneath), MooncakeConnector, MultiConnector (load from the first connector advertising tokens, save to all), OffloadingConnector/CPUOffloadingConnector (GPU→CPU), SharedStorageConnector (minimal reference impl), DecodeBenchConnector.

LMCache’s multi-tier hierarchy — confirmed: GPU HBM (~3.35 TB/s) → pinned CPU DRAM (hot cache, LRU, NUMA-aware) → local disk/NVMe (GDS) → remote store (Redis/Mooncake Store/InfiniStore), with a default chunk granularity of 256 tokens, async offload/load off the critical path, and — unlike vLLM’s ephemeral --swap-space — persistence across engine restarts plus cluster-scale cross-engine sharing. Reported 3×–10× latency reductions when reuse is high. (Confirmed against https://docs.lmcache.ai/developer_guide/architecture.html, arXiv:2510.09665, and the production-stack --swap-space-vs-LMCache tutorial. The 256-token/3–10× figures are LMCache-doc-reported, not independently benchmarked.)

The win condition is goodput, not throughput (reportedly): goodput = max request rate such that TTFT < T_ttft and ITL < T_itl (DistServe methodology, arXiv:2401.09670; vLLM used TTFT<1s, ITL<50ms/token). Disaggregation eliminates the ITL-violation cluster that collocated continuous batching produces, leaving TTFT-under-load as the residual failure mode. It is not universally better: per BentoML’s handbook it can regress ~20–30% on small/untuned workloads, short prompts, or high local prefix-hit, where collocated + chunked prefill is simpler and faster. Chunked prefill is the collocated middle ground — it reduces but cannot eliminate prefill/decode interference without a separate GPU pool. (Both goodput and the regression figures are marked unverified — hedge.)

Gotchas: V0→V1 is a hard cut — V0 used SimpleConnector/PyNcclConnector (over PyNcclPipe/MooncakePipe + a lookup buffer) with kv_role/kv_rank/kv_parallel_size, supported only 1P1D, and PyNcclConnector actively fails under V1. Don’t cite V0 flags as current. Disaggregation is still officially experimental; production PD relies on third-party connectors (LMCache, NIXL/Dynamo, Mooncake), not the in-tree SharedStorageConnector reference. Directionality is a lease/handshake lifecycle: prefill must hold its KV blocks until the consumer pulls (request_finished decides sync vs async free) — not fire-and-forget. Heterogeneous-TP transfer has correctness gates (HND required for het-TP head splitting; matching cache_dtype on both sides or the compat-hash fails).

2.7 Speculative decoding

Restructured around a “drafter” in the model runner. In V1, speculative decoding moved out of the engine core (PR #13363, “[V1][Spec decode] Move drafter to model runner,” WoosukKwon) into the GPU model runner. At init the runner instantiates a proposer chosen by speculative_config.method (dispatch on self.speculative_config.method in vllm/v1/worker/gpu_model_runner.py: ngram → NgramProposer, eagle → EagleProposer, medusa → MedusaProposer, draft_model → DraftModelProposer, …) plus a custom rejection sampler (partially Triton). Each decode step runs the target’s forward+sample, then calls propose_draft_token_ids(...) to produce k draft tokens for the next step; the following target forward verifies all k+1 positions at once (one forward yields up to k+1 accepted tokens). (Confirmed against PR #13363, the gpu_model_runner source, and the Anatomy blog. “(k)” is shorthand — the method receives sampler output/hidden-state context and delegates to the selected proposer’s propose().)

Proposal methods and checkpoint resolution — corrected/confirmed. Methods split by whether they need a separate checkpoint:

N-gram is configured with num_speculative_tokens plus prompt_lookup_min/prompt_lookup_max. (This is the corrected form; confirmed against https://docs.vllm.ai/en/latest/features/speculative_decoding/ and vllm/config/speculative.py. Caveat: on current main, generic draft_model speculation raises NotImplementedError — it is the checkpoint-based method described but not yet fully wired.)

The n-gram algorithm — confirmed: take the last prompt_lookup_max tokens of the sequence, find a prior occurrence in the same sequence, and if found propose the k tokens that followed; on no match, shrink the window and retry down to prompt_lookup_min. It requires zero draft weights and zero extra VRAM — “effectively free” — but only helps on repetitive/structured output (2×–4× on input-grounded tasks; poor on roleplay). (Confirmed against the Anatomy blog and issue #46977.)

Verification is lossless-in-expectation (reportedly): the custom rejection sampler does left-to-right verification with the accept rule accept token i with prob min(1, P_target/P_draft), else resample a corrected token from the residual, so the output distribution equals plain autoregressive sampling in expectation. Greedy verification reduces to accepting the longest matching argmax prefix plus one bonus/correction token. Exact bitwise reproducibility is NOT guaranteed — FP precision and batch-size-dependent logprobs make spec-on vs spec-off outputs differ token-by-token even at temperature 0. (Marked unverified — but the “lossless-in-expectation, not bitwise” framing is well-established; state it with the reproducibility caveat.)

Measured wins are workload- and batch-dependent (reportedly): EAGLE/EAGLE-3 give ~2–3× single-request / low-concurrency latency reductions (ITL cut ~3–4×, TTFT unchanged) with accepted-lengths ~1.7–2.8, but the gain collapses toward zero at high concurrency (batch ~32–64+), which is why ngram (free on repetitive output) and batch-size fallbacks exist. The EAGLE-3.1 blog reports 2.03× per-user throughput at concurrency 1 falling to ~1.66× at C=16. (Marked unverified — hedge; the specific numbers are third-party/vendor-reported and workload-specific.)

Continuous-batching integration (reportedly): V1 avoids V0’s batch-expansion padding by using variable-length packing over PagedAttention, with per-request scatter-gather and rollback of both KV-cache and position-ids for rejected draft tokens; chunked prefill (default-on) interleaves long prefills with decode, so spec decode targets ITL not TTFT. (Unverified — hedge. V0 used batch expansion + an MQA scorer.)

Gotchas: speculative decoding is a latency optimization, not throughput — gains evaporate once the GPU is compute-bound at high concurrency, so any speedup claim must state the concurrency/batch. “Lossless” ≠ bitwise-reproducible. Early V1 (Jan 2025 alpha) dropped V0’s separate draft_model method and shipped only ngram/EAGLE/Medusa; it was reintroduced later — a “V1 doesn’t support draft model” snippet is version-stale. EAGLE benchmark numbers were themselves once buggy (PR #25916 fixed a chat-template preamble bug that swung measured EAGLE-3-over-EAGLE-1 from ~5% to ~32% on MT-Bench) — treat old numbers with suspicion. Raising num_speculative_tokens (gamma) only helps when acceptance length is already high. --speculative-disable-by-batch-size, batch expansion, and the MQA scorer are V0 concepts — don’t assume the same flags/mechanism in V1.

2.8 Quantization & compute kernels

The selection surface is a single flag plus config auto-detection. vLLM chooses a weight-quant path from --quantization / -q (alias LLM(quantization=...)). Accepted tokens historically include aqlm, awq, awq_marlin, gptq, gptq_marlin, gptq_marlin_24, marlin, fp8, compressed-tensors (aka compressed_tensors), bitsandbytes, modelopt/nvfp4, gguf, deepspeedfp, moe_wna16, None (https://docs.vllm.ai/en/v0.5.0/models/engine_args.html). When omitted, vLLM reads the checkpoint’s config.json → quantization_config and dispatches on its quant_method; else weights are treated as unquantized and fall back to --dtype. QuantizationConfig (.../quantization/base_config.py) exposes get_from_keys(...) to pull quant_method/bits/group_size, and override_quantization_method(...) — a base-class hook (default no-op, “only overwritten by subclasses in exceptional circumstances”) that lets e.g. compressed-tensors claim a checkpoint or promotes awq→awq_marlin / gptq→gptq_marlin when the kernel is supported. An explicit --quantization does not unconditionally override auto-detect (corrected during verification): after the override/promotion pass, _verify_quantization compares the arg against the detected quant_method and raises ValueError on an incompatible mismatch (e.g. config fp8 vs arg gguf, issue #19050) — the arg takes effect only when it matches or is reconciled through the override/promotion path. Unknown methods register via register_quantization_config.

Weight-quant schemes cluster into three lineages. (1) Integer weight-only — GPTQ and AWQ, both INT4 (occasionally INT3) with group scales, group_size typically 128 (also −1/per-column, 32, 64); GPTQ carries optional desc_act/act-order, AWQ activation-aware channel scales. (2) compressed-tensors, the llm-compressor on-disk format (a safetensors extension) — a config layer, not one scheme: CompressedTensorsConfig maps config_group target/scheme tuples to concrete CompressedTensorsScheme subclasses (W8A8Int8, W8A8Fp8, WNA16). (3) Float — FP8 (E4M3, float8_e4m3fn) weight+activation, and FP4/NVFP4 (E2M1) plus MXFP4. Additional loaders: bitsandbytes (INT8/NF4, supports inflight load-time quant via load_in_4bit=True), gguf (llama.cpp Q2_K…Q8_0, reportedly experimental/slow), and vendor formats TorchAO, AMD Quark, Intel Neural Compressor, NVIDIA ModelOpt.

The kernel matters as much as the scheme — Marlin is the workhorse mixed-input GEMM. For W4A16/W8A16, the fast path is Marlin (Neural Magic), a fused-dequant GEMM for medium batch sizes; a runtime selector (reportedly choose_mp_linear_kernel) picks among Marlin, Machete, ExLlamaV2 from layer config + arch. Reported benchmarks: ~2.6× (GPTQ) and ~10.9× (AWQ) vs naive kernels — AWQ’s own kernel can be slower than FP16 without it. marlin_utils.py enforces MARLIN_SUPPORTED_GROUP_SIZES = [-1, 32, 64, 128] and gates arch via check_marlin_supported(...). Machete (PR #7174, merged 2024-08-20, vLLM 0.6.x) is the “spiritual successor to Marlin” for Hopper sm_90a: wgmma warp-group tensor-core instructions where Marlin’s mma tops out at ~75% of Hopper peak, with CUTLASS/CUTE weight-layout descriptions (https://developers.redhat.com/articles/2024/10/14/introducing-machete-mixed-input-gemm-kernel). Activation-quantized paths use CUTLASS INT8/FP8 W8A8 GEMMs; Hopper/Blackwell block-scaled FP8 uses DeepGEMM (128×128 block scales, DeepSeek-style).

AWQ/GPTQ checkpoints are repacked to Marlin layout at load, not on disk. Marlin needs a specific interleaved order, so GPTQ tensors are repacked into Marlin format during init, and AWQ weights (non-standard output-dim packing) are first normalized to a GPTQ-like layout (_convert_awq_to_standard_format) then repacked — hence awq_marlin/gptq_marlin add load cost but win at runtime. Caveat: desc_act: true (act-order) models fall off the fast Marlin path into a slower dequant path, so llm-compressor now defaults to static activation ordering (act-order accuracy, no runtime cost) (issue #5596).

KV-cache quantization is a separate axis (--kv-cache-dtype). Values: auto (default), fp8 = fp8_e4m3, fp8_e5m2 (CUDA 11.8+); ROCm only fp8(e4m3); Gaudi fp8_inc. E4M3 (±240) is higher precision but needs an FP32 scale; E5M2 wider range, lower precision (https://docs.vllm.ai/en/latest/features/quantization/quantized_kvcache/). Scales come three ways: default all-1.0; on-the-fly from a warmup batch (--calculate-kv-scales); or calibrated k_scale/v_scale (and for FA3, q_scale/prob_scale) baked in by llm-compressor — recommended. Granularity is per-tensor only for the KV cache today (per-channel WIP), unlike weight FP8.

FP8 KV directly shrinks the PagedAttention block. Blocks are fixed-size (default 16 tokens; MLA/DeepSeek larger, e.g. 128). Per-token bytes = f(hidden_size, num_kv_heads, num_hidden_layers, head_dim, dtype_size), so fp8 KV halves block bytes vs BF16 → ~2× KV capacity. The Anatomy post gives a worked figure: DeepSeek-R1 FP8 on 8×H200, ~45.7 GB×8 ≈ 365 GB KV pool, each 128-token block ≈ 8.6 MB. Latency caveat: most backends get no speedup from FP8 KV (dequant isn’t fused into attention) — the exception is FlashAttention-3 (runs attention in the FP8 domain, also quantizing Q). FlashInfer “always supports FP8 KV cache” and is recommended for FP8 serving (VLLM_ATTENTION_BACKEND=FLASHINFER); FlashMLA (auto for DeepSeek MLA) accepts FP8 KV but its path is less validated — vLLM’s FP8-KV writeup flags a systematic downward accuracy shift for uncalibrated FP8 on FlashMLA and recommends calibrating there (https://blog.vllm.ai/2026/04/22/fp8-kvcache.html; issue #12543).

Arch gating decides native vs emulated, and vLLM prefers fallback over refusal (corrected during verification — Turing removed, consumer-Blackwell nuance added). FP8 W8A8 (activation FP8) runs natively only on Hopper sm_90 and Ada sm_89; on Ampere sm_80/86/87 vLLM downgrades FP8 to weight-only W8A16 via “FP8 Marlin” (PR #5975, merged 2024-07-03) — a fused FP8→BF16/FP16 dequant kernel that actually gives higher accuracy since activations stay unquantized. Turing sm_75 is not covered — FP8, including the Marlin weight-only path, requires compute capability ≥8.0 (Ampere), the same floor as the int4 GPTQ/AWQ Marlin kernels. NVFP4 executes on native FP4 tensor cores (tcgen05) only on datacenter Blackwell sm_100/sm_103; on Hopper/Ampere (no FP4 hardware) an NVFP4 checkpoint loads as W4A16 through a Marlin FP4 dequant fallback (memory savings, no FP4 FLOPS). Consumer Blackwell sm_120/sm_121 (RTX 5090 / PRO 6000 / GB10 systems) physically has FP4 tensor cores, but vLLM’s backend oracle currently gates it out (e.g. the family(100) check; MoE experts often still fall to Marlin) — a gating limitation, not hardware absence (issue #31085). General rule: unsupported/gated arch → usually a silent fallback to a dequant path (Marlin) with a warning, not a hard refusal — though not universal (some MoE backends, e.g. FLASHINFER_CUTLASS, error on sm_120 instead, issue #33333), and the fallback has had correctness bugs (issue #34694).

Weight-only vs activation-quant, and accuracy recovery. Weight-only (W4A16 GPTQ/AWQ, W8A16 FP8-Marlin) shrinks memory but pays dequant overhead — it can lose throughput when not VRAM-bound. Activation+weight (W8A8 FP8/INT8) is the throughput sweet spot with minimal quality loss on Hopper+. Scale granularity: per-tensor cheapest; per-channel and per-group (group_size 128) recover accuracy; DeepGEMM FP8 uses 128×128 block scales. Accuracy-recovery algorithms live in llm-compressor (successor to AutoGPTQ/AutoAWQ/AutoFP8): GPTQ (Hessian error comp), AWQ (activation-aware channel scaling), SmoothQuant (migrates activation outliers into weights — typical INT8 recipe SmoothQuantModifier(smoothing_strength=0.8) + GPTQModifier(scheme="W8A8") ignoring lm_head), plus SparseGPT/RTN and 2:4-sparse+W4A16 combos.

Gotchas / version notes. (1) gptq_marlin_24 / 2:4-sparse-marlin_24 is deprecated/being removed — use compressed-tensors 2:4-sparse + W4A16/INT8. (2) desc_act=true silently drops you off fast Marlin; prefer desc_act=false or static ordering. (3) INT8 W8A8 models are BOS-sensitive — eval with add_bos_token=True. (4) V0 vs V1: quant scheme code is largely shared, but V1 owns the rewritten paged KV manager and the FP8-attention backends (FA3/FlashInfer FP8 KV is V1-era); FlashMLA-on-V1 FP8 was still shaking out ~0.9.2. (5) A --quantization conflicting with the checkpoint’s quantization_config errors (ValueError, issue #19050) rather than overriding. (6) NVFP4 emulation “works” but forfeits the FP4 FLOPS and has shipped correctness bugs — verify outputs on non-native-FP4 GPUs.

2.9 Serving architecture & the API server

Posture. vLLM exposes an OpenAI-compatible HTTP server (/v1/chat/completions, /v1/completions, stream=true SSE) fronting the isolated EngineCore. The public control/observability surface that downstream systems integrate against — and the only surface fak rides (§3) — is three things:

  1. OpenAI-compatible HTTP (chat/completions, streaming).
  2. The V1 KV-cache-events stream — BlockStored / BlockRemoved / AllBlocksCleared (native transport is ZMQ/msgpack).
  3. The Prometheus /metrics scrape — including vllm:time_to_first_token_seconds, vllm:time_per_output_token, vllm:inter_token_latency, vllm:request_queue_time, vllm:kv_cache_usage_perc, vllm:num_requests_running/waiting/swapped, and the prefix-cache counters vllm:prefix_cache_queries/vllm:prefix_cache_hits (§2.3).

Control-plane note (corrected). vLLM’s prefix-cache reset endpoint, POST /reset_prefix_cache, flushes the entire local prefix cache (whole-prefix, not span-granular) and is currently a dev-mode-gated administrative endpoint (VLLM_SERVER_DEV_MODE=1, grouped with /sleep, /wake_up, /collective_rpc), not yet a stable public API (open request: vLLM issue #32593). There is no exact-span / middle-span KV eviction endpoint on the vLLM control plane — per-tenant isolation is offered via cache_salt, and internal eviction is block-level LRU with no operator-addressable span API. This is load-bearing for the fak boundary in §3. (Corrected/confirmed against docs.vllm.ai cache api_router and issue #32593.)

Determinism. Absent batch-invariant kernels or a deterministic offline scheduler, a temperature-0 request is not token-for-token reproducible: batch-size-dependent reduction kernels yield non-identical outputs (Thinking Machines Lab’s “Defeating Nondeterminism in LLM Inference,” integrated into vLLM, demonstrated 80 unique completions of 1000 on Qwen3-235B at temp 0; enabling batch invariance made all 1000 identical). vLLM offers two reproducibility guarantees: batch invariance (VLLM_BATCH_INVARIANT=1, the only online option) and deterministic offline scheduling (VLLM_ENABLE_V1_MULTIPROCESSING=0), both only on identical hardware + vLLM version. (Confirmed against docs.vllm.ai reproducibility/batch-invariance and the Thinking Machines research.)

Compilation. vLLM V1 enables torch.compile by default, keeps a compile cache on by default (VLLM_DISABLE_COMPILE_CACHE=1 to disable), completes all compilation before serving (no request triggers compilation), and uses piecewise CUDA graphs with a warmup step by default. This is why a benchmark window that catches a cold/disabled cache, incomplete warmup, or request-time compilation is a cold-start reading, never a tuned baseline (§3). (Confirmed against the torch.compile blog and PR #10528.)

3. fak ↔ vLLM: the concrete seams

fak’s relationship to vLLM is “fak governs, vLLM serves” (issue #40). internal/engine.VLLMEngine registers the vllm EngineDriver/LifecycleEngine (admit/step/stream, issue #46) and rides a vLLM V1 worker over the three documented public surfaces above — no vLLM source is vendored, forked, or patched. That no-fork contract is executable: TestVLLMAdapterConstructsOnlyPublicVLLMEndpoints (internal/engine/vllm_public_surface_test.go) drives the real path constructors (buildOpenAIRequest, deriveMetricsURL) and fails closed if any emitted HTTP path leaves the allowlist {/v1/chat/completions, /v1/completions, /metrics} or hits /internal, /debug, .., /private, /admin. (Confirmed by direct repo verification: the test passes, no vLLM appears in go.mod or a vendored tree, and docs/serving/vllm-v1-adapter.md states the three-surface contract as “not merely a review promise.”)

The mapping, mechanism by mechanism:

KV eviction policy — the real architectural divergence. vLLM’s V1 BlockPool evicts strictly LRU off the free-block queue (recency, single dimension). fak’s eviction is value-aware: internal/compute/kvcost.go defines KVEvictionCost = recomputeCost(tokens) × reuseProbability(hits+1, Laplace) ÷ bytes — “expected recompute work per byte freed.” PickEvictionVictim picks the cheapest-to-lose evictable span, hard-excludes Pinned/Leased (fak’s analogue of vLLM ref_cnt > 0), and breaks ties by oldest LastUsed, so on uniform-cost inputs it provably reduces to LRU (a strict generalization). The doc comment explicitly names “vLLM’s priority block pool, SGLang’s radix LRU, Dynamo KVBM” as the recency-only baselines it aims to beat (#2239). ReplayKVCache is the harness measuring hit-rate of LRU vs cost-aware at a fixed token budget.

Prefix cache structure — SGLang-style, not vLLM-style. internal/radixkv/radixkv.go is a RadixAttention port (token-level compressed trie with longest-prefix walk() + mid-run edge split()), not vLLM’s block-level chained hash. It matches vLLM on refcounting (node.refs, refs>0 immune to eviction) and LRU (node.lastUsed, evictToBudget/lruLeaf), and diverges by keying on token-id sequences — no block_size, no block_hash, no cache_salt, collision-free by construction. Its explicit faithfulness note admits each node stores the full-prefix KVCache (length == path length), not vLLM’s paged per-block slabs — so a single N-token chain holds N(N+1)/2 positions (O(depth²) resident overshoot), made measurable via Stats.PrefixTokens (true resident) vs Tokens (SGLang per-segment LRU budget). EvictNode/EvictPrefix add policy (quarantine) eviction of a named subtree regardless of recency — a capability vLLM’s refcount+LRU pool structurally lacks. prewarm.go’s WarmInsert adds a lowest-priority eviction class (a prefetch that is always the first LRU victim), which vLLM’s flat free-queue has no tier for.

Tiered residency / offload economics — fak’s overlap with LMCache/Dynamo-KVBM. fak reifies the same GPU→CPU→disk lifecycle as vLLM/LMCache, but as a min-cost action per span rather than LRU + a single-tier connector: internal/compute/kvresidency.go (KVResidencySplit: hot on device/VRAM, cold spilled to host/RAM, fail-open — always keeps an equal-or-larger effective context), internal/compute/kvdemote.go (the demote-not-drop ladder Keep/SpillHost/Evict, where Evict’s restore cost is full re-prefill and SpillHost’s is a per-byte transfer-back; #2671/#2236), and internal/compute/kvprecision.go (quantize-in-place, the rung below spill and above evict; #1474 — the analogue of FP8 KV quantization as a density-delta planner). internal/xenginekv/arena.go is the cross-engine co-residence seam and a deliberate inversion of vLLM’s free-block reuse: a bump allocator never reuses a freed offset (so a later Put can’t alias a stale handle), trading vLLM’s aggressive recycling for provable handle non-aliasing and per-span quarantine.

The exact-span honesty boundary. fak’s native bit-exact middle-span Evict does not ride vLLM, because vLLM’s public control plane offers only whole-prefix POST /reset_prefix_cache (§2.9). So enginecache.SupportsExactSpan(EngineVLLM) is false (“currently never” true for any public engine), and a span-named quarantine degrades to one auditable whole-prefix flush (Degraded=true, DegradeReason="exact_span_unsupported_whole_prefix_flush") or fails closed under --engine-cache-require-exact-span. (Confirmed by direct repo read of internal/enginecache/enginecache.go and cross-checked against the vLLM control-plane fact.) This is the sharpest fak↔vLLM seam: whole-prefix flush is safe (the poisoned span is gone) but coarse (it evicts every other resident prefix too) — do not claim exact-span governance over a ridden vLLM.

Measurement-honesty gates. internal/vllmcompile/vllmcompile.go folds recorded torch.compile/CUDA-graph/warmup state into tuned/cold-start/diagnostic (Block.Classify() precedence: request-time-compilation → cold, cache-disabled → cold, warmup-incomplete → cold, unobserved → diagnostic, else tuned), and GateRows() makes one non-tuned row poison the whole A/B — a silently cold raw-vLLM baseline makes a fak “win” meaningless. Pointer fields (nil vs *false) distinguish “not observed” from “observed false.” (Confirmed by repo verification; #1731.) internal/engine/vllm_determinism.go treats reproducibility as operator-declared, fail-closed to unavailable: ParseVLLMDeterminism maps unknown/empty to DeterminismUnavailable, so fak never infers temp-0 reproducibility from sampling params and never reimplements batch-invariant kernels — it surfaces the engine’s declared batch_invariance / deterministic_offline_scheduler capability. (Confirmed; #1734.)

Request-control identity (reportedly). internal/engine/vllm_identity.go derives a per-request cache_salt ("fak-<hex>" digest folding fak_cache_tenant/authority/family, tenant always in the salt so byte-identical prefixes across tenants never share a slot; #1841) and emits a V1 priority field only when the engine advertises priority scheduling (FAK_VLLM_PRIORITY_SCHEDULING), otherwise degrading to vLLM’s FCFS default. Prometheus counters normalize into a shared fak_serving_* L2 schema; the prefix-cache-hit ratio is a *float64 left nil for vLLM (which exports query/hit counters, not a ratio) so no false 0% line is emitted. (These extraClaims are marked unverified in the research — present as the adapter’s design; the enginecache/vllmcompile/determinism/public-surface claims above *were directly verified.)*

Benchmark positioning. Raw vLLM — especially EP + EPLB for MoE serving — is the SOTA serving floor native fak work must beat or mark NOT COMPARABLE. docs/benchmarks/VLLM-EP-EPLB-MOE-BASELINE-RUNBOOK.md ships reproducible vllm serve EP/EPLB commands (ARM A eplb-off vs ARM B eplb-on), vllm bench serve percentile capture, and the fak.vllm-ep-eplb-moe-baseline.v1 artifact schema with result_claim_allowed:false — status “pending measurement,” no numbers, and it requires each arm’s vllm_compile.class == "tuned". fak’s native EP design (docs/notes/GLM52-EXPERT-PARALLEL-MULTIGPU-2026-06-29.md) deliberately contrasts: static contiguous expert bands per rank (ExpertParallelPlan), a replicated router, per-rank [H] partials combined by a single AllReduceSum — no all-to-all dispatch, no EPLB analogue (static placement). Host multi-process residency is proven bit-exact; the device-NCCL multi-GPU tok/s witness is still “not yet.”

The generation horizon. The vLLM adapter carries gen/second-next — a promotable-or-retirable architectural option (docs/serving/vllm-v1-adapter.md, “Generation horizon”): its value collapses toward gen/future if a native fak engine ships base serving (continuous batching, paged KV, prefix cache) first. The invalidating assumption is vLLM staying whole-prefix-reset-only — if a future vLLM exposes documented exact-span eviction, SupportsExactSpan(EngineVLLM) should flip true.

Where there is no overlap: fak is a single-engine kernel with no scheduler/worker KV-connector seam — it has no analogue of NIXL/Mooncake P–D transport. The genuine overlap is only the offload/residency economics (kvcost/kvresidency/kvdemote vs LMCache/Dynamo-KVBM tiering) and the on-disk content-addressed prefix-reuse tier (internal/l3kv SpanStager, digest-keyed, fail-closed re-verify at restore) vs LMCache’s cross-request cache reuse. Do not present fak as competing with NIXL/Mooncake transport.

4. Gotchas, version notes & sharp edges (consolidated)

V0-vs-V1 traps:

Moving-target defaults (pin to a version + SM level before quoting):

Correctness edges:

Research-provenance caveat: during the verification that produced this document, direct WebFetch of vLLM primary sources (blog, arXiv, DeepWiki, parts of docs.vllm.ai) was repeatedly blocked by the local tool guard, so a number of V1 internal-architecture details rest on WebSearch summaries of the primary docs rather than direct fetches — these are flagged as “reportedly”/”per the design docs” throughout. The confirmed claims (PagedAttention mechanics, prefix-cache hashing/data-structures, EP width, EPLB config surface + load window, KV-connector method set, NixlConnector specifics, LMCache tiers, spec-decode drafter move, and every fak-repo seam) were corroborated against authoritative sources or verified directly in the repo.

5. Further reading (annotated)

Foundational papers

Official vLLM blog

Official docs (docs.vllm.ai)

Source to read against pinned commits (the research recommends these to convert “reportedly” into “confirmed”): vllm/v1/core/{kv_cache_manager,kv_cache_coordinator,block_pool,kv_cache_utils}.py, vllm/v1/worker/{gpu_model_runner,block_table}.py, vllm/v1/attention/backends/{flash_attn,flashinfer,triton_attn,mla/common,mla/flashmla}.py, vllm/v1/sample/rejection_sampler.py, vllm/distributed/eplb/{eplb_state,rebalance_algo}.py, vllm/config/{speculative,parallel}.py, and vllm/distributed/kv_transfer/kv_connector/v1/{base,nixl_connector}.py.

Third-party / ecosystem

In-repo (fak) — docs/serving/vllm-v1-adapter.md (the authoritative issue-#40 driver doc and no-fork guard), docs/serving/dual-track-serving-plan.md (RIDE vs NATIVE tracks), docs/proofs/engine-seam.md (the exact-span fail-closed theorem behind SupportsExactSpan=false), and the internal/compute/kv*.go family (the value-aware eviction/residency/demotion planners that are fak’s genuine differentiator over vLLM’s LRU block pool).