What vLLM Actually Does

vLLM is a high-throughput LLM inference engine. When you send a prompt to a model like Llama or Mistral, vLLM orchestrates the entire process: scheduling requests, managing GPU memory, and running optimized kernels to generate tokens as fast as possible.

Think of vLLM like an air traffic control system. Multiple planes (requests) arrive simultaneously, each needing runway time (GPU compute) and gate space (memory). The controller must land planes efficiently without collisions.

Request Lifecycle

When you send a prompt to vLLM, here is the path your request takes through the system. Each step is a handoff between components, designed for maximum throughput.

The V1 Architecture Directory

The V1 engine lives in vllm/v1/. It was a ground-up rewrite of the original vLLM engine for better performance and simpler scheduling logic.

The Scheduling Algorithm

The V1 scheduler has no concept of separate "prefill" or "decode" phases. Instead, it treats every request uniformly: each has num_computed_tokens and a target num_tokens_with_spec. The scheduler assigns compute budget to help each request catch up.

# vllm/v1/core/sched/scheduler.py
def schedule(self) -> SchedulerOutput:
    # Each request just has num_computed_tokens and
    # num_tokens_with_spec. At each step, the scheduler
    # tries to assign tokens so that each request's
    # num_computed_tokens can catch up.
    scheduled_new_reqs: list[Request] = []
    scheduled_running_reqs: list[Request] = []
    preempted_reqs: list[Request] = []
    token_budget = self.max_num_scheduled_tokens

Scheduling Policies

The scheduler supports multiple scheduling policies to control which requests get compute time first.

# vllm/v1/core/sched/scheduler.py
# Scheduling constraints.
self.max_num_running_reqs = self.scheduler_config.max_num_seqs
self.max_num_scheduled_tokens = (
    self.scheduler_config.max_num_scheduled_tokens
    if self.scheduler_config.max_num_scheduled_tokens
    else self.scheduler_config.max_num_batched_tokens
)
self.policy = SchedulingPolicy(self.scheduler_config.policy)

Check Your Understanding

The Memory Bottleneck

LLM inference is memory-bound, not compute-bound. The GPU has enormous arithmetic power but limited memory bandwidth. The KV cache - storing key and value tensors from previous tokens - dominates GPU memory usage during inference.

Think of it like a library with limited shelf space. Each book (token) needs a reserved spot on the shelves (memory blocks). Without clever management, you run out of shelves long before you run out of reading capacity.

Block-Based Memory Allocation

vLLM pioneered paged attention - the key insight that made high-throughput LLM serving possible. Instead of reserving maximum-length contiguous memory for each request, vLLM allocates memory in small blocks that can be scattered across GPU memory.

KV Cache Data Structure

The KVCacheBlocks dataclass is the interface between the Scheduler and the KV Cache Manager. It hides internal memory details from scheduling logic.

# vllm/v1/core/kv_cache_manager.py
@dataclass
class KVCacheBlocks:
    blocks: tuple[Sequence[KVCacheBlock], ...]
    # blocks[i][j] refers to the i-th kv_cache_group
    # and the j-th block of tokens.

    def __add__(self, other) -> "KVCacheBlocks":
        return KVCacheBlocks(
            tuple(list(itertools.chain(blk1, blk2))
                  for blk1, blk2 in zip(self.blocks, other.blocks))
        )

Prefix Caching with Block Hashing

When two requests share the same prefix (e.g., the same system prompt), vLLM can reuse the computed KV cache blocks instead of recomputing them. This is done via content-addressable hashing of block contents.

# vllm/v1/core/block_pool.py
class BlockHashToBlockMap:
    """Cache of blocks used for prefix caching.
    Maps hash directly to a block or multiple blocks
    (i.e. {block_hash: KVCacheBlocks})
    A cached block is a full block with a block hash
    that can be used for prefix caching.
    The cached block may be used by running requests
    or in the free_block_queue that could
    potentially be evicted."""

Memory Access in Paged Attention

The attention kernel must handle non-contiguous blocks. For each query token, it looks up which physical blocks hold the relevant keys and values using a page table (Req_to_tokens mapping).

Check Your Understanding

Why Custom Kernels?

Standard PyTorch operations execute one operation at a time: matrix multiply, then softmax, then another multiply. Each operation reads from and writes to GPU HBM. The data movement between operations is the bottleneck, not the math.

Triton lets you fuse multiple operations into one kernel. Data stays in fast on-chip SRAM between operations instead of making round-trips to slow HBM.

The Triton Programming Model

Triton operates at the block level. You write code that processes a tile of data, and the GPU launches thousands of these program instances in parallel. Think of it like a factory assembly line where each worker handles one box of parts, not individual screws.

GPU Memory Hierarchy

Understanding where data lives is the key to writing fast Triton kernels. The hierarchy has three levels, each with dramatically different speed and capacity.

vLLM Prefill Attention Kernel

The prefill kernel processes all prompt tokens at once. Each program instance handles one tile of query positions for one attention head. It iterates over all key/value positions in tiles, computing attention scores with online softmax.

# vllm/v1/attention/ops/triton_prefill_attention.py
@triton.jit
def _fwd_kernel(Q, K, V, sm_scale, B_Start_Loc, B_Seqlen, Out,
    stride_qbs, stride_qh, stride_kbs, stride_kh, ...):
    cur_batch = tl.program_id(0)
    cur_head = tl.program_id(1)
    start_m = tl.program_id(2)
    cur_kv_head = cur_head // kv_group_num
    cur_batch_seq_len = tl.load(B_Seqlen + cur_batch)
    # initialize pointer to m and l
    m_i = tl.zeros([BLOCK_M], dtype=tl.float32) - float("inf")
    l_i = tl.zeros([BLOCK_M], dtype=tl.float32)
    acc = tl.zeros([BLOCK_M, BLOCK_DMODEL], dtype=tl.float32)

The Core Attention Loop

This is where the magic happens. The kernel iterates over tiles of keys, computing attention scores and accumulating the weighted values - all without writing intermediate results to HBM.

# The tiled attention loop (prefill kernel)
for start_n in range(start_n_limit, end_n_limit, BLOCK_N):
    k = tl.load(k_ptrs + start_n * stride_kbs,
                mask=(pos_k < cur_batch_seq_len))
    qk = tl.dot(q, k)
    qk = tl.where(mask, qk * sm_scale, -1.0e8)
    m_ij = tl.maximum(m_i, tl.max(qk, 1))
    qk -= m_ij[:, None]
    p = tl.math.exp2(qk)
    l_ij = tl.sum(p, 1)
    alpha = tl.math.exp2(m_i - m_ij)
    l_i = l_i * alpha + l_ij
    acc = acc * alpha[:, None]
    v = tl.load(v_ptrs + start_n * stride_vbs, ...)
    acc += tl.dot(p, v)
    m_i = m_ij

Check Your Understanding

The Roofline: Memory vs. Compute Bound

Every kernel operation sits somewhere on the roofline. Prefill attention is often compute-bound (large matrix multiplies). Decode attention is almost always memory-bound (loading KV cache for a single query token). Different optimizations apply to each.

Split-KV: Parallelizing Decode Attention

vLLM's decode kernel uses a two-stage approach. Stage 1 splits the KV sequence across NUM_KV_SPLITS parallel instances. Stage 2 merges their results. Think of it like parallel reading groups: each group reads a chapter of the book, then they combine their notes.

Decode Stage 1: The Split-KV Kernel

Each split instance processes a slice of the KV sequence. The key insight: it uses paged memory access to handle non-contiguous blocks while iterating over its assigned KV range.

# vllm/v1/attention/ops/triton_decode_attention.py
@triton.jit
def _fwd_kernel_stage1(Q, K_Buffer, V_Buffer, sm_scale,
    Req_to_tokens, B_Seqlen, Att_Out, ...):
    cur_batch = tl.program_id(0)
    cur_head = tl.program_id(1)
    split_kv_id = tl.program_id(2)

    kv_len_per_split = tl.cdiv(cur_batch_seq_len, NUM_KV_SPLITS)
    split_kv_start = kv_len_per_split * split_kv_id
    split_kv_end = tl.minimum(split_kv_start + kv_len_per_split,
                              cur_batch_seq_len)

Paged KV Access in the Kernel

The decode kernel handles paged memory by computing physical addresses from logical positions. This is the critical path where page table lookups happen at GPU speed.

# Inside the decode attention loop
for start_n in range(split_kv_start, split_kv_end, BLOCK_N):
    offs_n = start_n + tl.arange(0, BLOCK_N)
    kv_page_number = tl.load(
        Req_to_tokens + stride_req_to_tokens_b * cur_batch_req_idx
        + offs_n // PAGE_SIZE,
        mask=offs_n < split_kv_end)
    kv_loc = kv_page_number * PAGE_SIZE + offs_n % PAGE_SIZE
    k = tl.load(K_Buffer + kv_loc[:, None] * stride_buf_kbs
        + cur_kv_head * stride_buf_kh + offs_d[None, :],
        mask=(offs_n[:, None] < split_kv_end) & mask_d)
    qk = tl.sum(q[None, :] * k, 1)
    qk *= sm_scale

Optimization: Block Size Tuning

The BLOCK_N and BLOCK_DMODEL parameters control tile sizes. Choosing the right values is a trade-off between occupancy and per-instance efficiency.

FP8 KV Cache: Halving Memory Traffic

vLLM's decode kernel supports FP8 quantized KV cache. Since decode is memory-bound, halving the bytes read from HBM nearly doubles effective performance.

# FP8 dequantization in decode kernel
ks = tl.load(k_scale)
vs = tl.load(v_scale)
for start_n in range(split_kv_start, split_kv_end, BLOCK_N):
    k = tl.load(K_Buffer + offs_buf_k, mask=...)
    if k.dtype.is_fp8():
        k = (k.to(tl.float32) * ks).to(q.dtype)
    qk = tl.sum(q[None, :] * k, 1)

Check Your Understanding

Why Quantize?

Quantization solves two problems at once: it shrinks model size so larger models fit on fewer GPUs, and it reduces memory traffic so inference runs faster. A 70B parameter model in FP16 needs ~140GB of memory. In INT4, it fits in ~35GB - a single GPU.

Think of it like compressing a high-resolution photo to JPEG. You lose some subtle details, but the image is recognizable and loads much faster. The art is choosing compression levels that preserve quality where it matters.

Quantization Granularity

The scale factor determines accuracy. Finer granularity preserves more information but adds overhead. vLLM supports all three strategies.

vLLM Quantization Methods

vLLM supports a comprehensive set of quantization methods, each with different trade-offs between accuracy, speed, and memory savings.

FP8 Quantization in Practice

FP8 is the sweet spot for inference: minimal accuracy degradation, hardware-accelerated on H100/MI300X, and transparent to the model architecture. vLLM implements FP8 for both linear layers and KV cache.

# vllm/model_executor/layers/quantization/fp8.py
from vllm.model_executor.layers.quantization.base_config import (
    QuantizationConfig, QuantizeMethodBase)
from vllm.model_executor.layers.quantization.kv_cache import (
    BaseKVCacheMethod)
from vllm.model_executor.kernels.linear import (
    init_fp8_linear_kernel)
from vllm.model_executor.kernels.linear.scaled_mm import (
    CutlassFP8ScaledMMLinearKernel,
    MarlinFP8ScaledMMLinearKernel)

How Quantized Inference Works

The quantized forward pass adds a dequantization step before (or fused into) each linear operation. The key optimization is fusing dequantization with the GEMM so there is no extra memory traffic.

Check Your Understanding

Scaling Beyond One GPU

Large models (70B+ parameters) exceed single-GPU memory even with quantization. vLLM supports splitting models across multiple GPUs using tensor parallelism and pipeline parallelism.

Think of it like a relay race vs. parallel swimmers. Tensor parallelism is parallel swimmers: everyone works on the same lap simultaneously and syncs at the wall. Pipeline parallelism is a relay: each swimmer handles one leg sequentially.

Communication Primitives

Tensor parallelism requires GPUs to exchange partial results after each layer. vLLM uses all-reduce for this synchronization.

# vllm/distributed/communication_op.py
def tensor_model_parallel_all_reduce(
    input_: torch.Tensor) -> torch.Tensor:
    """All-reduce the input tensor across
       model parallel group."""
    return get_tp_group().all_reduce(input_)

def tensor_model_parallel_all_gather(
    input_: torch.Tensor, dim: int = -1
) -> torch.Tensor:
    """All-gather the input tensor across
       model parallel group."""
    return get_tp_group().all_gather(input_, dim)

Tensor Parallel Attention

In tensor-parallel attention, each GPU handles a subset of attention heads. After computing attention for its heads, results are combined via all-reduce before passing to the next layer.

Continuous Batching

Continuous batching is vLLM's scheduling superpower. Instead of waiting for all requests in a batch to finish before starting new ones, the scheduler inserts and removes requests at every step.

Speculative Decoding

Speculative decoding is vLLM's approach to reducing latency. A small draft model guesses the next several tokens, then the large model verifies them in parallel. When guesses are correct, you get multiple tokens from one expensive forward pass.

Profiling & Debugging Strategies

Optimizing kernels requires measurement. Here are the key tools and what they reveal about vLLM performance.

Check Your Understanding