3、PagedAttention原理：从内存碎片到页级映射
上一章我们聊了KV缓存的基本概念，说白了就是存下每个token的Key和Value向量，避免重复计算。但实际部署时，你会发现一个很头疼的问题——显存明明够，却装不下几个请求。为什么？内存碎片在作怪。
我最早做长序列推理时，就踩过这个坑。一个8卡A100的节点，理论显存80GB，结果并发一高，OOM了。查了半天，发现显存利用率不到60%。嗯，问题就出在传统KV缓存的内存管理方式上。
3.1 传统KV缓存的内存碎片问题
传统做法是预分配。每个请求进来，系统就给它分配一块连续显存，大小按最大序列长度算。比如模型支持2048个token，那就一次性分配2048个slot。你想想看，如果实际只生成了100个token，剩下的1948个slot就白白浪费了。
这还不是最要命的。更严重的是，多个请求的显存块大小不一，频繁申请释放后，显存会变得支离破碎。我见过一个极端案例：显存还剩30GB，但最大的连续块只有4GB，新请求根本塞不进去。
传统预分配的核心问题：
- 内部碎片：预分配了最大长度，实际用不到，浪费显存
- 外部碎片：不同请求的显存块大小不一，释放后无法合并
- 无法共享：即使两个请求的prompt前缀相同，也得各自存一份KV缓存
为什么会这样？因为传统方案把KV缓存当成一个连续的大数组来管理。每个请求独占一块连续地址，地址空间是固定的，没法动态伸缩。
| 对比维度 | 传统预分配 | PagedAttention |
|---|---|---|
| 分配方式 | 一次性分配最大长度 | 按需分配，动态增长 |
| 内存碎片 | 严重（内部+外部） | 几乎无碎片 |
| 共享能力 | 不支持 | 支持页级共享 |
| 显存利用率 | 通常40%-60% | 可达90%以上 |
3.2 PagedAttention的核心设计
PagedAttention的思路，说白了就是把虚拟内存的分页机制搬到了GPU显存里。操作系统用页表管理物理内存，PagedAttention就用类似的方式管理KV缓存。
核心思想很简单：
- 把KV缓存切成固定大小的块（Block），每个块包含若干个token的KV数据
- 每个请求的逻辑地址空间是连续的，但物理存储可以分散在不同的块里
- 通过块表（Block Table）维护逻辑块到物理块的映射
我习惯把每个块叫做一个页（Page），大小通常是16或32个token。为什么选这个数？太小了映射开销大，太大了又浪费。我在项目中试过8、16、32、64，最后发现16是个不错的平衡点。
个人经验：块大小不是固定的。如果你的模型层数深、head数多，每个token的KV数据量大，建议用16；如果模型小、序列短，32可能更合适。我一般先跑个profiling，看看平均序列长度再定。
3.3 逻辑页与物理页的映射
这里有个关键概念：逻辑页和物理页。
- 逻辑页：请求视角的连续页号。比如第0个token到第15个token属于逻辑页0，第16到31属于逻辑页1，以此类推。
- 物理页：显存中实际存储的页。物理页的编号是全局的，不连续也没关系。
- 块表：一个映射数组，逻辑页号 -> 物理页号。
举个例子。一个请求生成了48个token，块大小16，那么它需要3个逻辑页（页0、1、2）。物理上，这些页可能分散在显存的不同位置：逻辑页0映射到物理页5，逻辑页1映射到物理页12，逻辑页2映射到物理页3。
Attention计算时，PagedAttention会通过块表找到每个逻辑页对应的物理页地址，然后按页读取数据。这个过程对上层模型是透明的——模型看到的还是连续的KV序列。
// 伪代码：PagedAttention的块表查找
// block_table: 逻辑页号 -> 物理页号
// logical_page: 当前要访问的逻辑页
// page_size: 每个页包含的token数
int physical_page = block_table[logical_page];
int offset_in_page = token_index % page_size;
// 从物理页中读取KV数据
float* k_ptr = physical_memory[physical_page] + offset_in_page * kv_dim;
float* v_ptr = physical_memory[physical_page] + offset_in_page * kv_dim + kv_dim;
注意：块表本身也占显存。如果块大小太小（比如4），块表会变得很大，反而浪费。我建议块大小不要小于8，也不要大于64。16和32是最常用的。
3.4 vLLM中的实现机制
vLLM是目前最流行的PagedAttention实现。它的核心组件有三个：
- BlockManager：管理所有物理页的分配和释放
- BlockTable：每个请求维护一个块表
- PagedAttention Kernel：支持分页读取的Attention计算核
BlockManager的工作流程是这样的：
- 初始化时，把所有显存划分成固定大小的物理页，放入空闲列表
- 新请求到来时，按需分配物理页，更新块表
- 请求结束后，回收物理页，放回空闲列表
- 支持写时复制（Copy-on-Write）：多个请求共享相同前缀时，物理页只存一份
写时复制这块我得多说两句。假设两个请求的prompt都是"今天天气真不错，我们一起去"，后面的内容不同。传统方案下，每个请求都得存一遍这10个token的KV缓存。PagedAttention呢？让两个请求的块表指向同一组物理页，只有当某个请求要修改某个页时，才复制一份出来。
我在实际项目中测试过，共享前缀的场景下，显存占用能降低40%-60%。尤其是对话场景，用户消息前面都带着system prompt和历史上下文，共享效果非常明显。
vLLM的PagedAttention优势总结：
- 显存利用率从50%提升到95%以上
- 支持动态序列长度，不浪费
- 页级共享，前缀复用零成本
- 调度灵活，可以抢占、暂停、恢复请求
下面这张图展示了PagedAttention的核心架构：
vLLM的PagedAttention Kernel也做了特殊优化。传统Attention需要连续的内存访问，PagedAttention则要处理非连续的物理页。vLLM的做法是：在kernel启动前，先把块表加载到共享内存中，然后每个线程块处理一个物理页的数据。这样既保证了访存效率，又支持了分页。
我曾经对比过vLLM和传统HuggingFace实现的推理性能。在相同显存下，vLLM的吞吐量是传统方案的2-3倍。原因很简单：显存利用率高了，能同时跑的请求就多了。
避坑指南：我曾经遇到过一个问题——块表更新不及时，导致Attention读到了旧数据。排查后发现是异步kernel执行顺序的问题。解决方案是在每次块表更新后，加一个cudaDeviceSynchronize()。虽然会损失一点性能，但保证了正确性。如果你用vLLM的官方实现，它已经处理好了这个细节，但自己写PagedAttention kernel时一定要注意。
最后总结一下。PagedAttention的核心价值就一句话：用操作系统的分页思想，解决KV缓存的内存碎片问题。它让显存利用率从50%提升到95%以上，还顺带实现了前缀共享和动态调度。vLLM把这个思想落地成了工业级的实现，现在已经成为大模型推理加速的事实标准。