编程 vLLM 推理引擎源码级拆解：PagedAttention 显存虚拟化、Continuous Batching 调度与量化落地全指南

2026-07-27 03:44:38 +0800 CST views 6

vLLM 推理引擎源码级拆解：PagedAttention 显存虚拟化、Continuous Batching 调度与量化落地全指南

一句话先把话说透：LLM 推理慢、贵、显存爆，90% 的锅不在模型本身，而在 KV Cache 的显存管理和请求调度这两件事上。vLLM 之所以能在同样一张 H100 上把吞吐量拉高好几倍，靠的不是什么魔法算子，而是把操作系统里搞了几十年的虚拟内存分页思想搬进了 GPU 显存。这篇文章我们不谈 PPT 概念，直接从显存字节、调度队列、CUDA kernel 一路啃到生产部署。

一、背景：为什么裸跑 HuggingFace `model.generate()` 一定会爆显存

先摆一个所有做过 LLM 服务化的人都踩过的坑。

你拿一个 7B 模型，FP16 权重大约 14GB，塞进一张 24GB 的 4090，看起来绰绰有余。结果一上并发，稍微来几个长 prompt，CUDA out of memory 立刻糊你一脸。你会很困惑：权重才 14GB，剩下 10GB 去哪了？

答案是 KV Cache。

自回归解码（autoregressive decoding）的本质是：生成第 N 个 token 时，需要用到前面 N-1 个 token 在每一层 attention 里算出来的 Key 和 Value。如果每次都重算，复杂度是 O(N²)，慢到无法接受。所以工程上必然要把历史的 K、V 缓存下来，这就是 KV Cache。

我们算一笔账。一个 token 的 KV Cache 占多少显存？

单 token KV Cache 字节数
= 2 (K 和 V)
× num_layers (层数)
× num_kv_heads (KV 头数)
× head_dim (每个头的维度)
× dtype_size (FP16 = 2 字节)

以 Llama-2-13B 为例：num_layers=40，num_kv_heads=40，head_dim=128，FP16：

2 × 40 × 40 × 128 × 2 = 1,638,400 字节 ≈ 1.6 MB / token

一个请求如果上下文 2048 token，光 KV Cache 就要 3.2GB。你要是想同时服务 20 个这样的请求，那就是 64GB —— 一张 H100 80GB 也就勉强扛住 20 来个并发，而且这还没算权重和激活值。

问题的关键不在"总量大"，而在于传统实现对这块显存的管理极其浪费。这就是 vLLM 要解决的核心矛盾。

1.1 传统 KV Cache 管理的三宗罪

在 vLLM 出现之前（也就是各种朴素的 HF generate + 静态 batch 实现），KV Cache 的显存分配方式是连续预分配：为每个请求预留一整块连续显存，大小按 max_seq_len 来算。

这带来三种浪费：

第一宗罪：内部碎片（internal fragmentation）。 你按 max_seq_len=2048 预留，但一个请求实际只生成了 100 个 token 就遇到 EOS 停了，剩下 1948 个 token 的空间白占着，谁也用不了。研究显示，传统方案里真正被有效利用的 KV Cache 显存往往只有 20%~40%，剩下全在打水漂。

第二宗罪：外部碎片（external fragmentation）。 不同请求要的连续块大小不一，请求来来去去，显存被切成一堆大小不等的碎片。明明总空闲量够，就是找不出一块足够大的连续区域给新请求。

第三宗罪：无法共享。 并行采样（一个 prompt 生成多个候选 n>1）、beam search、以及多个请求共享同一段 system prompt 时，这些完全相同的 KV Cache 被重复存了 N 份，纯纯的浪费。

这三宗罪加起来，就是"显存明明还有，但就是塞不下更多请求"的根本原因。

二、核心概念：PagedAttention —— 把虚拟内存搬进显存

vLLM 的第一发子弹，也是最出名的一发，叫 PagedAttention。它的思想一句话能说清：别再要求 KV Cache 连续存放，把它切成固定大小的块（block），像操作系统管理内存页那样管理它。

如果你写过操作系统或者了解过虚拟内存，那这套东西会让你会心一笑——这就是 分页（paging） 的原样照搬。

2.1 从连续到分块

传统方式：一个请求的 KV Cache 是逻辑上连续、物理上也连续的一大块。

PagedAttention：

把 KV Cache 切成固定大小的 KV block，每个 block 存固定数量（比如 16 个）token 的 K、V。
每个请求维护一张 block table（块表），记录"逻辑块 → 物理块"的映射，等价于操作系统里的页表。
逻辑上连续的 token 序列，物理上可以散落在显存的任意位置。

这样一来：

内部碎片消失了：最多浪费"最后一个 block 里没填满的部分"，也就是不到 1 个 block 的量（16 token × 1.6MB/token 也就几十 MB 甚至更小），碎片率从 60% 直接降到 4% 以下。
外部碎片消失了：所有 block 大小一致，显存变成一个整齐的 block 池，随便分配随便回收，永远不会"有空间但凑不齐连续块"。
共享成为可能：多个请求/多个采样序列可以把 block table 的某几项指向同一个物理 block，配合引用计数（refcount）和写时复制（copy-on-write），实现 KV Cache 共享。

2.2 一张图说清 block table

假设 block size = 4（为了画图方便，真实默认是 16），一个请求已经处理了 "程序员茄子是一个" 这 7 个 token（假设分成 7 个 token）：

逻辑视图（请求看到的连续序列）:
[程序][员][茄][子]  [是][一][个][很]

逻辑块 0: [程序][员][茄][子]   -> 物理块 #7
逻辑块 1: [是][一][个][很]     -> 物理块 #2

Block Table (页表):
  logical_block 0 -> physical_block 7
  logical_block 1 -> physical_block 2

物理显存 (block 池，乱序):
  #0 [空闲]
  #1 [别的请求的数据]
  #2 [是][一][个][很]   <- 本请求逻辑块1
  #3 [空闲]
  ...
  #7 [程序][员][茄][子] <- 本请求逻辑块0

生成下一个 token 时，如果当前最后一个 block 还没满，就直接往里填；满了就从空闲 block 池里申请一个新 block，追加到 block table 末尾。按需分配，用多少给多少。

2.3 PagedAttention 的 CUDA kernel 做了什么

普通 attention kernel 假设 K、V 在显存里是连续的，可以用简单的指针步进访问。PagedAttention 的 kernel 不能这么干，它必须先查 block table，再去物理块里取数据。

伪代码大致长这样（简化到能看懂原理的程度）：

// PagedAttention kernel 的核心逻辑（伪代码）
// query: 当前 token 的 Q 向量
// block_table: 该序列的逻辑块->物理块映射
// k_cache/v_cache: 全局物理 block 池
__global__ void paged_attention_kernel(
    const float* query,          // [num_heads, head_dim]
    const int*   block_table,    // [max_num_blocks_per_seq]
    const float* k_cache,        // [num_blocks, block_size, num_heads, head_dim]
    const float* v_cache,
    int seq_len, int block_size)
{
    float acc[HEAD_DIM] = {0};
    float running_max = -INFINITY, running_sum = 0;

    // 遍历该序列所有 token（分块进行）
    for (int tok = 0; tok < seq_len; tok++) {
        int logical_block = tok / block_size;      // 属于第几个逻辑块
        int offset        = tok % block_size;      // 块内偏移
        int physical_block = block_table[logical_block];  // 查页表拿到物理块号

        // 用物理块号定位到全局 k_cache 里的真实地址
        const float* k = k_cache + physical_block * block_size * H * D
                                 + offset * H * D;

        // 计算 attention score，在线 softmax（FlashAttention 那套）
        float score = dot(query, k) * scale;
        // ... online softmax 更新 running_max / running_sum / acc ...
    }
    // 写回输出
}

关键就在 block_table[logical_block] 这次间接寻址——它把"逻辑上连续"翻译成"物理上离散"，代价是每个 block 多一次查表。实践证明这点开销完全可以被显存利用率的大幅提升覆盖，净收益巨大。

现代 vLLM 里，这套 kernel 已经进化到集成 FlashAttention / FlashInfer，把在线 softmax、分块计算、paged KV 访问揉在一起，并支持 GQA（Grouped Query Attention）、MQA 等变体。

三、架构分析：vLLM 的四大件如何协同

光有 PagedAttention 只解决了"显存怎么存"，还没解决"请求怎么调度"。vLLM 的完整架构由四个核心组件构成，我们一个个拆。

                ┌─────────────────────────────────┐
   HTTP 请求 →  │           API Server            │  (OpenAI 兼容接口)
                └──────────────┬──────────────────┘
                               │
                ┌──────────────▼──────────────────┐
                │           LLMEngine              │  中央协调器
                │  ┌────────────┐  ┌────────────┐  │
                │  │  Scheduler │  │BlockManager│  │
                │  │  (调度器)  │  │ (块管理器) │  │
                │  └────────────┘  └────────────┘  │
                └──────────────┬──────────────────┘
                               │
                ┌──────────────▼──────────────────┐
                │           Worker(s)              │  每 GPU 一个
                │   ModelRunner + CacheEngine      │
                │   (跑前向 + 管理物理 KV block)   │
                └─────────────────────────────────┘

3.1 Scheduler：连续批处理的大脑

调度器是 vLLM 吞吐量的真正引擎，它实现了 Continuous Batching（连续批处理），这是相对传统 static batching 的又一次范式升级。

Static batching 的问题： 假设一个 batch 有 8 个请求，其中 7 个生成 20 个 token 就结束了，1 个要生成 500 个。传统做法必须等整个 batch 全部跑完才能返回、才能接新请求。于是那 7 个"提前毕业"的槽位就一直空转，GPU 利用率被最长的那个请求拖死。这叫 head-of-line blocking。

Continuous batching 的解法： 调度以 迭代（iteration） 为单位，而不是以请求为单位。每生成一步（一个 token），调度器就重新审视一遍：

哪个请求结束了（遇到 EOS 或达到 max_tokens）→ 立刻踢出 batch、释放它的 KV block、返回结果。
有空位了 → 立刻从 waiting 队列拉新请求进来填补。

于是 batch 的组成是动态流动的，GPU 几乎永远满载。这一招通常就能带来 2~4 倍的吞吐量提升。

调度器内部维护三个队列：

# 简化版调度器状态机
class Scheduler:
    def __init__(self):
        self.waiting  = deque()   # 等待首次调度（还没算 prefill）
        self.running  = deque()   # 正在解码中
        self.swapped  = deque()   # 被抢占换出的请求

    def schedule(self):
        scheduled = []
        # 1. 优先把 running 里的请求继续往下解码
        for seq in self.running:
            if self.block_manager.can_append(seq):
                self.block_manager.append_slot(seq)  # 可能要新增 block
                scheduled.append(seq)
            else:
                # 显存不够了 -> 触发抢占
                self.preempt(seq)

        # 2. 显存有余则从 waiting 拉新请求做 prefill
        while self.waiting and self.block_manager.can_allocate(self.waiting[0]):
            seq = self.waiting.popleft()
            self.block_manager.allocate(seq)
            self.running.append(seq)
            scheduled.append(seq)

        return scheduled

3.2 抢占与恢复：显存不够了怎么办

当正在跑的请求越来越长、KV block 池被耗尽时，调度器必须抢占（preempt） 一些请求给别人让路。vLLM 有两种策略：

Recomputation（重算）：直接把被抢占请求的 KV block 全部释放，等以后有空间了，把它当作新请求重新 prefill 一遍。适合 prompt 不太长的场景，重算成本可控。
Swapping（换出）：把 KV block 从 GPU 显存拷贝到 CPU 内存（host memory），腾出 GPU 空间；恢复时再拷回来。适合 prompt 很长、重算代价高的场景，但受 PCIe 带宽限制。

这又是操作系统的老朋友——页面换入换出（swap）。vLLM 把 OS 内存管理的整套心智模型完整复刻了一遍，这也是为什么理解它对做过系统编程的人格外亲切。

3.3 BlockManager 与 CacheEngine

BlockManager 管理逻辑层面的账本：谁占了哪些物理 block、每个 block 的引用计数是多少、空闲 block 池还剩多少。它不碰真实显存数据，只管"记账"。

CacheEngine 管实际的物理显存：在启动时一次性把 GPU 上能用于 KV Cache 的显存全部预分配成一个巨大的 block 池（由 gpu_memory_utilization 参数控制比例，默认 0.9），之后所有分配/回收都在这个池内进行，避免运行时反复向 CUDA 申请释放显存带来的开销和碎片。

启动时 vLLM 会做一次 profiling run：跑一个 dummy 的极限 batch，测出峰值激活显存，然后用 总显存 × utilization − 权重 − 峰值激活 反推出能开多少个 KV block。这个数字（num_gpu_blocks）直接决定了你的服务能扛多大并发。

四、Prefix Caching：把重复的前缀只算一次

第三发子弹叫 Automatic Prefix Caching（APC，自动前缀缓存），它把 block 共享从"同一请求内"扩展到了"跨请求"。

4.1 场景

想想这些真实场景：

你的服务有一段几百 token 的固定 system prompt，每个请求都带着它。
Few-shot 提示，前面几个示例对所有查询都一样。
多轮对话，第二轮请求的前缀就是第一轮的全部内容。

这些公共前缀在传统实现里，每来一个请求就要重新 prefill 一遍，纯纯的重复计算。

4.2 原理：block 级的哈希去重

APC 的实现优雅得令人拍案：既然 KV Cache 已经按 block 切分，那就给每个 block 算一个哈希，哈希由"该 block 的 token 内容 + 它前面所有 block 的哈希（即前缀）"共同决定。

# 前缀哈希：block 的哈希依赖它之前所有内容
def compute_block_hash(prev_block_hash, token_ids_in_block):
    return hash((prev_block_hash, tuple(token_ids_in_block)))

当新请求进来，逐 block 算哈希，如果发现某个 block 的哈希在缓存里已存在，说明"这段前缀之前有人算过"，直接把 block table 指向那个已有的物理 block，跳过这段的 prefill 计算，refcount +1 即可。

于是共享的 system prompt 只在第一次被完整计算，之后所有请求都是"零成本复用"。在 system prompt 很长、QPS 很高的生产场景，APC 能把 TTFT（首 token 延迟）砍掉一大截。

开启方式极简：

from vllm import LLM
llm = LLM(model="meta-llama/Llama-3-8B", enable_prefix_caching=True)

被缓存的 block 采用 LRU 淘汰，refcount 归零且长期不用的 block 会被回收还给空闲池。

五、代码实战：从起步到榨干性能

理论讲完，上手。

5.1 离线批量推理

from vllm import LLM, SamplingParams

# 一次性把模型加载进显存，预分配 KV block 池
llm = LLM(
    model="Qwen/Qwen2.5-7B-Instruct",
    dtype="bfloat16",
    gpu_memory_utilization=0.90,   # 用 90% 显存做权重+KV Cache
    max_model_len=8192,            # 上下文上限，直接影响单请求最大 block 数
)

sampling = SamplingParams(
    temperature=0.7,
    top_p=0.9,
    max_tokens=512,
)

prompts = [
    "用一句话解释什么是 PagedAttention。",
    "写一个 Python 快速排序。",
    "解释 TCP 三次握手为什么不能是两次。",
]

# vLLM 会自动把这些请求组成连续批处理，一起跑
outputs = llm.generate(prompts, sampling)
for out in outputs:
    print("Q:", out.prompt)
    print("A:", out.outputs[0].text)
    print("-" * 40)

注意：你不需要手动 padding、不需要手动组 batch、不需要管显存——这些全被引擎接管了。你只管把 prompt 丢进去。

5.2 起一个 OpenAI 兼容的在线服务

vLLM 最实用的能力是一行命令起一个和 OpenAI API 完全兼容的服务：

python -m vllm.entrypoints.openai.api_server \
  --model Qwen/Qwen2.5-7B-Instruct \
  --dtype bfloat16 \
  --gpu-memory-utilization 0.9 \
  --max-model-len 8192 \
  --enable-prefix-caching \
  --tensor-parallel-size 1 \
  --port 8000

然后你现有的 OpenAI SDK 代码改个 base_url 就能直接用：

from openai import OpenAI

client = OpenAI(
    base_url="http://localhost:8000/v1",
    api_key="EMPTY",  # 本地服务随便填
)

resp = client.chat.completions.create(
    model="Qwen/Qwen2.5-7B-Instruct",
    messages=[
        {"role": "system", "content": "你是一个资深后端工程师。"},
        {"role": "user", "content": "Redis 和 Valkey 该怎么选？"},
    ],
    stream=True,
)
for chunk in resp:
    delta = chunk.choices[0].delta.content or ""
    print(delta, end="", flush=True)

那段固定的 system prompt，配合 --enable-prefix-caching，从第二个请求开始就是复用的，白赚一波延迟优化。

5.3 张量并行：单卡装不下就切开

70B 级别的模型单卡装不下，用张量并行（Tensor Parallelism, TP）把权重切到多卡：

python -m vllm.entrypoints.openai.api_server \
  --model meta-llama/Llama-3.1-70B-Instruct \
  --tensor-parallel-size 4 \
  --gpu-memory-utilization 0.92

tensor-parallel-size=4 表示把每一层的权重矩阵按列/行切成 4 份分到 4 张卡，前向时通过 NCCL 做 all-reduce 聚合。注意 TP 对卡间互联带宽很敏感，有 NVLink 的机器收益远好于走 PCIe 的。超大模型还可以叠加 --pipeline-parallel-size 做流水线并行跨节点。

六、性能优化：生产环境的调参清单

把 vLLM 跑起来只是及格线，调好才是真本事。下面是我从实战里总结的关键旋钮。

6.1 `gpu_memory_utilization`

这是性价比最高的一个参数。它决定给 KV block 池留多少显存。调高（如 0.95）→ 更多 block → 更高并发；但留太满容易被激活值峰值撞爆 OOM。建议从 0.9 起步，压测观察峰值再往上试探，独占卡可以摸到 0.95。

6.2 `max_num_seqs` 与 `max_num_batched_tokens`

max_num_seqs：一个 iteration 里最多同时跑多少个序列。太小浪费吞吐，太大会因 KV block 不够频繁触发抢占。
max_num_batched_tokens：一个 iteration 里 prefill + decode 加起来最多处理多少 token。这个值和 chunked prefill 强相关。

6.3 Chunked Prefill：平衡延迟与吞吐

prefill（处理输入 prompt）是计算密集型，一个长 prompt 的 prefill 会长时间霸占 GPU，把正在 decode 的其他请求全卡住，导致它们的 token 间延迟（ITL）飙升。

Chunked prefill 的解法：把长 prompt 的 prefill 切成小块，穿插在 decode 步骤之间执行。这样既不让 decode 请求饿死，又能填满 GPU 算力。新版本 vLLM 默认开启：

--enable-chunked-prefill --max-num-batched-tokens 2048

max_num_batched_tokens 调小 → decode 优先、延迟更平滑；调大 → prefill 吞吐更高。这是一个延迟 vs 吞吐的经典权衡旋钮，按你的 SLA 来定。

6.4 量化：用精度换显存和速度

量化把权重从 FP16 压到 INT8/INT4/FP8，直接砍显存、提吞吐。vLLM 支持一大堆量化格式：

# 加载一个 AWQ INT4 量化模型
python -m vllm.entrypoints.openai.api_server \
  --model Qwen/Qwen2.5-7B-Instruct-AWQ \
  --quantization awq

# Hopper/Blackwell 架构上用 FP8，几乎无损且吞吐高
python -m vllm.entrypoints.openai.api_server \
  --model meta-llama/Llama-3.1-8B-Instruct \
  --quantization fp8

选型经验：

FP8（需 H100/H200/Blackwell）：精度损失极小，速度提升明显，生产首选。
AWQ / GPTQ (INT4)：显存省得最狠（权重砍到 1/4），适合显存吃紧的场景，精度损失可接受但对某些任务敏感，务必用你自己的评测集验一遍。
注意：量化的是权重，KV Cache 默认还是 FP16/BF16。想进一步省显存可以开 KV Cache 量化 --kv-cache-dtype fp8，但要评估对长上下文精度的影响。

6.5 投机解码：用小模型给大模型"打草稿"

Speculative decoding（投机解码） 是压低延迟的杀手锏：用一个小的 draft 模型（或 n-gram、EAGLE 等方法）一次性猜出未来若干个 token，再让大模型一次前向并行验证这些草稿。猜对的部分直接采纳，等于一步解码出多个 token。

python -m vllm.entrypoints.openai.api_server \
  --model meta-llama/Llama-3.1-70B-Instruct \
  --speculative-model meta-llama/Llama-3.2-1B-Instruct \
  --num-speculative-tokens 5

在 draft 模型和目标模型分布接近、接受率高的场景，能把延迟压下去一大截；但如果接受率低，验证的开销反而拖后腿，所以一定要在你的真实流量上测接受率。

6.6 观测：别当黑盒用

vLLM 暴露了 Prometheus 指标，生产必接。重点盯这几个：

vllm:gpu_cache_usage_perc：KV block 池使用率。长期贴近 100% → 该扩容或降并发了。
vllm:num_requests_running / waiting：waiting 队列持续堆积 → 容量不足。
vllm:time_to_first_token_seconds（TTFT）：首 token 延迟，用户体感最强。
vllm:time_per_output_token_seconds（TPOT/ITL）：token 间延迟，决定"打字流畅度"。
抢占次数：频繁 preempt 说明显存吃紧，要么调低 max_num_seqs，要么加卡。

curl http://localhost:8000/metrics | grep vllm:

七、一个容易被忽视的细节：block size 到底怎么选

--block-size 默认 16（部分 kernel 支持 8/32）。这个值不是越大越好也不是越小越好：

block 越小：内部碎片越小（浪费更少），但 block table 越长、间接寻址开销越多、kernel 调度粒度越碎。
block 越大：寻址开销小、kernel 更高效，但最后一个不满的 block 浪费更多，且 prefix caching 的共享粒度变粗（要整块完全一致才能共享）。

经验上 16 是绝大多数场景的甜点值，除非你有非常规整的短序列负载再去动它。

八、常见踩坑清单（血泪总结）

启动就 OOM：gpu_memory_utilization 太高，或 max_model_len 设得离谱大导致 profiling 阶段峰值激活撑爆。先把 utilization 降到 0.85 试。
吞吐上不去：多半是 max_num_seqs 太小，或者 KV block 不够（看 gpu_cache_usage_perc）。别忘了这是 KV block 池的容量在限制并发，不是算力。
长请求把短请求拖慢：没开 chunked prefill，或 max_num_batched_tokens 太大。开启 chunked prefill 立竿见影。
量化后精度崩了：INT4 对某些任务（尤其代码、数学）敏感，务必用业务评测集验证，别只看 loss。
多卡 TP 反而更慢：卡间走 PCIe 没有 NVLink，all-reduce 通信成了瓶颈。TP 只在高带宽互联下划算。
prefix caching 没生效：前缀必须 token 级完全一致才能命中，system prompt 里带了时间戳/随机 ID 这种动态内容会让缓存全部失效。把动态部分挪到 prompt 后面去。

九、总结与展望

回过头看，vLLM 真正的天才之处，不是发明了什么全新算法，而是精准识别出 LLM 推理的瓶颈在显存管理，然后把计算机系统领域被验证了几十年的成熟思想（分页、页表、swap、写时复制、哈希去重）系统性地移植到了 GPU KV Cache 上。这是一种典型的"跨领域降维打击"。

把这篇文章的核心浓缩成几条，方便你随时回顾：

PagedAttention = KV Cache 分块 + 块表映射，把碎片率从 60% 干到 4% 以下，这是显存利用率飞跃的根基。
Continuous Batching = 以迭代为粒度动态进出请求，消灭 head-of-line blocking，吞吐 2~4 倍起。
Prefix Caching = block 级哈希去重，公共前缀只算一次，狠砍 TTFT。
抢占/换出 = 显存版的 OS swap，让系统在过载时优雅降级而非直接崩。
生产调参核心四件套：gpu_memory_utilization、chunked prefill、量化（FP8 优先）、投机解码。

展望往后，推理引擎的战争还在升级：PD 分离（prefill 和 decode 拆到不同实例/硬件，各自优化）、KV Cache 跨请求跨节点共享（把 KV 当成一种可调度的分布式资源）、更激进的量化与稀疏、以及 vLLM V1 架构对调度器和执行路径的彻底重写，都在把"每一分 GPU 显存和算力都不浪费"这件事推向极致。

对我们工程师来说，结论很实在：LLM 服务化早已不是"把模型 load 进来 generate 一下"的玩具阶段。真正决定你的服务能不能扛住流量、成本能不能压下来的，是你对显存、调度、缓存这些系统级细节的理解深度。而 vLLM，恰好是这套系统思维最好的一本活教材——它值得你 clone 下来，对着这篇文章的脉络，去读一遍它的 scheduler.py 和 block_manager 源码。读完你会发现，所谓 AI 基础设施，底子还是那套扎实的计算机系统功夫。