外观
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 内存(或丢弃重算) |
这一设计带来三个直接收益:
- 近零浪费:块按需分配,只有最后一块可能不满,浪费被压到块粒度;
- 消除外部碎片:所有块大小一致,无「空洞」;
- 块级共享:同一个 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 从「没人管的缓存」变成推理系统的核心基础设施,衍生出三大方向:
- 前缀缓存(prefix caching):把共享 prompt 前缀(系统提示词、few-shot 模板)的 KV 缓存起来跨请求复用,命中时预填近乎免费——vLLM 原生支持,命中即省大块延迟。
- Chunked prefill(分块预填):长 prompt 的 prefill 计算量大、会「堵住」decode,把 prefill 拆成小块与 decode 交错执行,降低首个 token 延迟(TTFT)、提高 GPU 利用率。
- 更极致的 KV 优化:SGLang 的 RadixAttention(把 KV 缓存建成可共享的前缀树)、KV Cache 量化、GQA/MQA 减少 KV 规模,以及 prefill/decode 分离部署(PD 分离)等。
vLLM 在站内的位置
局限
- KV Cache 再省,权重和长上下文仍可能吃爆显存:块管理优化的是「浪费」,不解决「总量」;超长上下文仍要配合量化(量化经典)与并行(并行与分布式推理)。
- 抢占的开销:显存不足时 swap/重算请求会带来延迟尖刺,调度策略需谨慎调参。
- 非连续块的 kernel 复杂度:PagedAttention kernel 的实现与调优成本高,也是各家引擎差异化的主战场之一。
延伸阅读
- 大模型推理优化 —— KV Cache、连续批处理、投机解码的原理底座
- vLLM 实战案例 —— 论文设计在真实部署里的参数与调优
- 推理服务系统:Clipper / Orca / Nexus 等 —— 连续批处理的来源(Orca)与更早的服务系统
- 并行与分布式推理 —— 张量并行/流水并行如何与 vLLM 叠加
- 量化经典:LLM.int8 / GPTQ / AWQ —— 权重与 KV 量化减少显存总量
- 性能优化与容量规划 —— 吞吐/延迟的测量与压测方法