Skip to content

PagedAttention:vLLM 系统论文

本页速览 vLLM 论文(SOSP 2023)是本领域必读:把操作系统虚拟内存分页的思想引入 KV Cache 管理,消除显存碎片,配合连续批处理让吞吐提升 2-4 倍。本文逐节精读。

PagedAttention:vLLM 系统论文

论文:Efficient Memory Management for Large Language Model Serving with PagedAttention(Woosuk Kwon, Zhuohan Li et al., SOSP 2023,UC Berkeley 等)

vLLM 是模型部署领域必读的系统论文:它把操作系统虚拟内存的分页思想搬进 GPU 显存管理,解决了 LLM 推理最痛的「KV Cache 显存浪费」,配合连续批处理让吞吐提升 2-4 倍。今天几乎每个 LLM 推理框架(TGI、TensorRT-LLM、SGLang、乃至自研引擎)都在吸收它的设计。

读本文前建议先掌握 大模型推理优化(KV Cache、prefill/decode 两阶段)与 推理服务化(批处理概念);本文的精读是「为什么 vLLM 长这样」的答案,与之配套的落地细节见 vLLM 实战案例

背景与动机:KV Cache 是 LLM 推理的隐形瓶颈

LLM 自回归生成时,每个新 token 都要看之前所有 token 的 attention 信息。为了不重复计算,系统把每个已生成 token 的 Key/Value 向量缓存在显存里——这就是 KV Cache。问题在于:

  • 随序列动态增长:生成几个 token 就增加几个 KV 向量,长度事前不知道;
  • 单请求占比巨大:13B 模型在 A100(40GB)上,权重约占 65%KV Cache 约占 30%,激活只占一小部分(论文图 1);
  • 多请求竞争:KV Cache 是「批里每个请求各一份」,batch 越大、序列越长,显存越吃紧。

而传统系统(FasterTransformer、Orca)的做法是:按请求的 max_length 预分配一整块连续显存。论文的实测触目惊心:这种静态分配下,只有 20.4%-38.2% 的 KV Cache 显存真正存了 token 状态,其余全浪费了。

text
现有系统:按请求最大长度预分配连续块(三种浪费)
┌───────────────────────────────────────────────────────┐
│ 请求 A(预分配 2048 token,实际只生成 512)              │
│ ████████░░░░░░░░░░░░░░░░░░░░░░░░░░░░░░░░  ← 内部碎片   │
│ 请求 B(预分配 2048,还没开始生成)                      │
│ ░░░░░░░░░░░░░░░░░░░░░░░░░░░░░░░░░░░░░░░░  ← reserved  │
│ 请求 C(长度与 A/B 不同)                               │
│ ██░░░░░░░░░░░░░░░░░░░░░░░░░░░░░░░░░░░░░░  ← 外部碎片   │
└───────────────────────────────────────────────────────┘
显存利用率实测只有 20.4%-38.2%

这三种浪费(reserved / 内部碎片 / 外部碎片)直接限制了 batch 大小——而 batch 大小就是 LLM 服务的吞吐上限。把 KV Cache 用明白,比多买几块卡便宜得多。

核心思想:PagedAttention——OS 虚拟内存的 GPT 复刻

论文的洞察朴素而深刻:「动态增长、长度未知、碎片化、需要共享」的 KV Cache,和操作系统的进程内存是一模一样的问题——而 OS 早在 60 年前就靠「分页」解决了。于是有了 PagedAttention:把 KV Cache 切成固定大小的块(block),按需分配,用**块表(block table)**记录每个请求的块映射。

text
PagedAttention:KV Cache 按固定大小分块,非连续存放

 GPU 显存块池
 ┌────┬────┬────┬────┬────┬────┬────┬────┐
 │B0  │B1  │B2  │B3  │B4  │B5  │B6  │B7  │  每个块 = 16 个 token 的 KV
 ├────┼────┼────┼────┼────┼────┼────┼────┤
 │A-1 │B-1 │A-2 │    │B-2 │    │A-3 │B-3 │
 └────┴────┴────┴────┴────┴────┴────┴────┘
         ↑按需领取,用完归还,可被任意请求复用

 请求 A 的块表(逻辑 → 物理):[0, 2, 6]
 请求 B 的块表(逻辑 → 物理):[1, 4, 7]

注意 A 的块是 0 → 2 → 6,物理上不连续——注意力计算通过块表间接寻址(PagedAttention 的 kernel 支持非连续 KV 读取)。这和 MMU 让用户程序以为内存连续、实际散落在物理帧里,是同一个原理。

与 OS 分页的类比

操作系统虚拟内存PagedAttention / vLLM
页(page)块(block,固定 16 个 token 的 KV)
物理内存帧(frame)GPU 显存块
页表(page table)块表(block table)
按需分页(demand paging)按需分配块,随生成增长
进程地址空间请求的 KV Cache 序列
碎片整理 / 紧凑(compaction)不需要——块大小统一,无外部碎片
页共享(fork、写时复制)块共享(并行采样、beam search、前缀缓存)
换出(swap)KV 块换到 CPU 内存(或丢弃重算)

这一设计带来三个直接收益:

  1. 近零浪费:块按需分配,只有最后一块可能不满,浪费被压到块粒度;
  2. 消除外部碎片:所有块大小一致,无「空洞」;
  3. 块级共享:同一个 prompt 前缀、或一次请求的多个采样序列,可以共享同一批 KV 块——内存复用进一步压低峰值占用。

连续批处理:请求级调度 vs 迭代级调度

光有显存还不够,吞吐还取决于「batch 什么时候换」。LLM 自回归生成一个 token 需要跑一次完整前向,一次请求要跑几十上百次前向。传统服务按「请求」调度:整批请求全部生成完才换下一批——慢的拖慢快的,新请求只能干等Orca 的论文 精确描述了这个问题)。

连续批处理(continuous batching / iteration-level scheduling)把调度粒度降到「一次迭代」:

  • 迭代级调度:每生成一个 token 后,已完成的请求立刻出队返回,新请求立刻加入下一个迭代的 batch;
  • Batch 始终满:GPU 不再空转等最慢请求,吞吐因此大幅提升。
text
请求级调度(Orca 之前):           迭代级调度(连续批处理):
batch = {A,B,C} 固定                batch 每迭代更新
A 完成 → 整个 batch 等 B、C         A 完成 → A 出队,D 入队
新请求 D 排队等 batch 结束          batch = {B,C,D}
→ GPU 空转、延迟高                  → GPU 满载、吞吐高

这一机制由 Orca(OSDI 2022)首次提出,vLLM 直接继承为标配。

系统实现要点

  • 块管理器(Block Manager):显存池预分配、块分配/回收、引用计数(refcount)实现跨序列共享;显存不足时触发抢占(preemption)——把低优先级请求的块 swap 到 CPU 或整体丢弃、稍后重算(类似 OS 的 swap)。
  • 预分配显存池:启动时一次性分配 KV cache 显存池,推理中不动态 cudaMalloc(动态分配既慢又加剧碎片)。
  • 前缀共享:并行采样(parallel sampling)的多个序列共享 prompt 的 KV 块,节省 prompt 的重复计算。
  • 为投机解码预留:预留可容纳 draft token 的块,避免投机解码额外分配开销。

关键结果

指标数字
吞吐(vs FasterTransformer / Orca)2-4x,且延迟不劣化(论文摘要原话)
吞吐(vs Orca,ShareGPT 长序列 / OPT-13B 配置)最高 8.5x(论文图 10 等具体实验配置)
KV Cache 浪费从 60%-80% 降到 近零(块粒度)
提升规律序列越长、模型越大、解码算法越复杂(并行采样/beam search),收益越明显
分布式支持超过单卡容量的模型(张量并行下各卡独立管理 KV Cache)

8.5x 的正确读法

2-4x 是「同延迟下相对 SOTA 系统的综合吞吐提升」;8.5x 是特定配置(长序列工作负载如 ShareGPT、OPT-13B)下的峰值对比。工程报告里如果只说「8.5x」而不提负载特征,容易误导容量规划。生产评估务必用你自己的流量分布重测,参考性能优化

影响与后续:KV Cache 成为新战场

vLLM 之后,KV Cache 从「没人管的缓存」变成推理系统的核心基础设施,衍生出三大方向:

  1. 前缀缓存(prefix caching):把共享 prompt 前缀(系统提示词、few-shot 模板)的 KV 缓存起来跨请求复用,命中时预填近乎免费——vLLM 原生支持,命中即省大块延迟。
  2. Chunked prefill(分块预填):长 prompt 的 prefill 计算量大、会「堵住」decode,把 prefill 拆成小块与 decode 交错执行,降低首个 token 延迟(TTFT)、提高 GPU 利用率。
  3. 更极致的 KV 优化:SGLang 的 RadixAttention(把 KV 缓存建成可共享的前缀树)、KV Cache 量化、GQA/MQA 减少 KV 规模,以及 prefill/decode 分离部署(PD 分离)等。

vLLM 在站内的位置

局限

  • KV Cache 再省,权重和长上下文仍可能吃爆显存:块管理优化的是「浪费」,不解决「总量」;超长上下文仍要配合量化(量化经典)与并行(并行与分布式推理)。
  • 抢占的开销:显存不足时 swap/重算请求会带来延迟尖刺,调度策略需谨慎调参。
  • 非连续块的 kernel 复杂度:PagedAttention kernel 的实现与调优成本高,也是各家引擎差异化的主战场之一。

延伸阅读

参考资料