摘要:PagedAttention 是 vLLM 推理系统的核心技术,灵感来自操作系统的虚拟内存分页机制。它将连续的 KV 缓存分割成固定大小的非连续块,实现了近零浪费的内存利用率和 2-4 倍的吞吐量提升。本文深入解析 PagedAttention 的设计原理、实现细节及实际应用效果。
一、背景与动机
1.1 LLM 推理的内存挑战
大语言模型(LLM)的高吞吐量服务需要对足够多的请求进行批处理(batching)。然而,现有系统面临一个核心问题:
KV Cache 内存巨大:每个请求的键值缓存(KV cache)占用大量 GPU 内存
对于 LLaMA-2 70B 模型:
- 每个 token 的 KV 缓存:约 1MB(FP16)
- 4096 序列长度:约 4GB/请求
- 80GB 显存的 A100:最多服务 20 个并发请求!
动态增长与收缩:KV cache 随序列生成动态变化,长度不可预测
内存浪费严重:
- 碎片化:传统预分配方式导致大量内存碎片
- 冗余复制:相同前缀的 KV 状态被重复存储
1.2 vLLM 的突破
vLLM 是一个基于 PagedAttention 的 LLM 推理服务系统,实现了:
- 近零浪费的 KV cache 内存利用率
- 灵活的 KV cache 共享机制(同一请求内及跨请求)
- 2-4 倍吞吐量提升(相比 FasterTransformer、Orca 等 SOTA 系统)
二、PagedAttention 核心原理
2.1 灵感来源:操作系统虚拟内存
PagedAttention 的设计灵感来自操作系统的虚拟内存和分页(paging):
| 操作系统概念 | PagedAttention 对应 |
|---|---|
| 虚拟地址空间 | 逻辑 KV 缓存空间 |
| 物理内存页 | GPU 内存 Block |
| 页表 | Block 表(Block Table) |
| 按需分页 | 动态 Block 分配 |
2.2 核心思想
将连续的 KV cache 分割成固定大小的非连续块(Block):
传统方式:
[请求 1: 连续分配 100 个 token 空间] ████████████████████ (浪费 50%)
[请求 2: 连续分配 100 个 token 空间] ████████ (浪费 20%)
PagedAttention 方式:
[请求 1: Block 3 → Block 7 → Block 1] ███ ███ ███ (按需分配)
[请求 2: Block 5 → Block 2] ███ ███ (按需分配)
2.3 关键数据结构
2.3.1 Block(块)
- 定义:KV cache 的基本存储单元
- 大小:固定容纳
BLOCK_SIZE个 token(通常 16 或 32) - 内容:每个 Block 存储特定 head 上固定数量 token 的 K 和 V 数据
2.3.2 Block Table(块表)
- 作用:维护每个序列的逻辑 token 到物理 Block 的映射
- 结构:每个序列维护一个 Block 索引列表
# 序列 A 的 Block Table
seq_a_blocks = [3, 7, 1, 9] # 逻辑顺序:Block 3 → 7 → 1 → 9
# 序列 B 的 Block Table
seq_b_blocks = [5, 2, 8] # 逻辑顺序:Block 5 → 2 → 8
2.3.3 内存布局
# KV 缓存的物理存储
k_cache: [num_blocks, num_kv_heads, head_size/x, block_size, x]
v_cache: [num_blocks, num_kv_heads, head_size, block_size]
# 典型配置(A100 80GB)
num_blocks = 32768 # 可动态调整
block_size = 16 # 每个 block 存储 16 个 token
三、实现细节
3.1 注意力计算流程
def paged_attention(
q: torch.Tensor, # [batch, num_heads, head_size]
k_cache: torch.Tensor, # [num_blocks, num_kv_heads, head_size/x, block_size, x]
v_cache: torch.Tensor, # [num_blocks, num_kv_heads, head_size, block_size]
block_tables: torch.Tensor, # [batch, max_num_blocks_per_seq]
seq_lens: torch.Tensor, # [batch]
scale: float,
) -> torch.Tensor:
"""
PagedAttention 前向传播
"""
batch_size, num_heads, head_size = q.shape
output = torch.zeros_like(q)
for i in range(batch_size):
# 获取当前序列的 block table
seq_block_table = block_tables[i]
seq_len = seq_lens[i]
# 计算需要访问的 block 数量
num_blocks = (seq_len + BLOCK_SIZE - 1) // BLOCK_SIZE
# 遍历所有 block
for j in range(num_blocks):
block_id = seq_block_table[j]
block_start = j * BLOCK_SIZE
block_end = min(block_start + BLOCK_SIZE, seq_len)
actual_block_len = block_end - block_start
# 从物理 block 中加载 K, V
k_block = k_cache[block_id, :, :, :actual_block_len, :]
v_block = v_cache[block_id, :, :, :actual_block_len]
# 计算注意力
qk = torch.matmul(q[i], k_block.transpose(-1, -2)) * scale
attn_weights = torch.softmax(qk, dim=-1)
output[i] += torch.matmul(attn_weights, v_block)
return output
3.2 CUDA 内核优化
vLLM 实现了高度优化的 CUDA 内核:
__global__ void paged_attention_kernel(
float* out, // [batch, num_heads, head_size]
const float* q, // [batch, num_heads, head_size]
const float* k_cache, // [num_blocks, num_kv_heads, head_size/x, block_size, x]
const float* v_cache, // [num_blocks, num_kv_heads, head_size, block_size]
const int* block_tables, // [batch, max_num_blocks_per_seq]
const int* seq_lens, // [batch]
int batch_size,
int num_heads,
int head_size,
int max_num_blocks_per_seq,
float scale
) {
// 每个线程块处理一个序列的一个 head
int batch_idx = blockIdx.x;
int head_idx = blockIdx.y;
// 获取序列信息
int seq_len = seq_lens[batch_idx];
int* block_table = block_tables + batch_idx * max_num_blocks_per_seq;
// 共享内存存储 Q
__shared__ float q_shared[HEAD_SIZE];
// 加载 Q 到共享内存
for (int i = threadIdx.x; i < head_size; i += blockDim.x) {
q_shared[i] = q[batch_idx * num_heads * head_size + head_idx * head_size + i];
}
__syncthreads();
// 累加器
float acc[HEAD_SIZE] = {0};
float max_val = -INFINITY;
float sum_exp = 0;
// 遍历所有 block
for (int block_idx = 0; block_idx < (seq_len + BLOCK_SIZE - 1) / BLOCK_SIZE; block_idx++) {
int physical_block_id = block_table[block_idx];
// 加载 K, V 块并计算注意力
// ... (详细实现省略)
}
// 写回结果
// ...
}
3.3 内存分配器
vLLM 实现了专门的 GPU 内存分配器:
class BlockAllocator:
def __init__(self, num_blocks, block_size, dtype):
self.num_blocks = num_blocks
self.block_size = block_size
self.free_blocks = list(range(num_blocks))
self.allocated_blocks = set()
# 预分配 GPU 内存
self.k_cache = torch.empty(
(num_blocks, num_kv_heads, head_size, block_size),
dtype=dtype,
device='cuda'
)
self.v_cache = torch.empty(
(num_blocks, num_kv_heads, head_size, block_size),
dtype=dtype,
device='cuda'
)
def allocate(self, num_blocks_needed):
"""分配指定数量的 block"""
if len(self.free_blocks) < num_blocks_needed:
raise RuntimeError("Out of memory")
allocated = self.free_blocks[:num_blocks_needed]
self.free_blocks = self.free_blocks[num_blocks_needed:]
self.allocated_blocks.update(allocated)
return allocated
def free(self, block_ids):
"""释放 block"""
for block_id in block_ids:
self.allocated_blocks.remove(block_id)
self.free_blocks.append(block_id)
四、KV Cache 共享
4.1 序列内共享
在 beam search 或 parallel sampling 中,多个序列共享相同的前缀:
序列 A: [token1, token2, token3, token4_A]
序列 B: [token1, token2, token3, token4_B]
↑ 共享前缀
传统方式:重复存储 token1-3 的 KV 状态 PagedAttention:多个序列的 block table 指向相同的物理 block
4.2 跨序列共享
对于相同的 prompt,不同请求可以共享 KV 缓存:
# 缓存池
prompt_cache = {
"prompt_hash_1": [block_1, block_2, block_3],
"prompt_hash_2": [block_5, block_6],
}
# 新请求到达时
if prompt_hash in prompt_cache:
# 直接复用已计算的 KV 缓存
block_table = prompt_cache[prompt_hash].copy()
else:
# 预填充并缓存
block_table = prefill_and_cache(prompt)
这使得首个 token 延迟(TTFT)大幅降低!
五、性能对比
5.1 内存利用率
| 系统 | 内存利用率 | 浪费原因 |
|---|---|---|
| 传统预分配 | ~50% | 碎片化 + 过度预留 |
| vLLM PagedAttention | ~96% | 仅最后一个 block 可能未填满 |
5.2 吞吐量对比
在 A100 80GB 上的实测结果(LLaMA-2 70B, seq_len=4096):
| 系统 | 吞吐量 (tokens/s) | 相对提升 |
|---|---|---|
| FasterTransformer | 450 | 1.0× |
| Orca | 520 | 1.2× |
| vLLM | 1850 | 4.1× |
5.3 并发请求数
相同显存下可服务的并发请求数:
| 系统 | 并发请求数 |
|---|---|
| 传统方式 | 20 |
| vLLM | 80+ |
vLLM 可以服务 4 倍更多的并发请求!
六、使用指南
6.1 安装 vLLM
pip install vllm
6.2 基本使用
from vllm import LLM, SamplingParams
# 创建 LLM 实例
llm = LLM(model="meta-llama/Llama-2-70b-chat-hf")
# 设置采样参数
sampling_params = SamplingParams(
temperature=0.7,
max_tokens=512,
top_p=0.9
)
# 生成
prompts = [
"Hello, my name is",
"The president of the United States is",
"The capital of France is",
]
outputs = llm.generate(prompts, sampling_params)
for output in outputs:
print(output.outputs[0].text)
6.3 API 服务
# 启动 API 服务器
python -m vllm.entrypoints.api_server \
--model meta-llama/Llama-2-70b-chat-hf \
--host 0.0.0.0 \
--port 8000
# 调用 API
curl http://localhost:8000/generate \
-d '{"prompt": "Hello, my name is", "max_tokens": 512}'
七、总结
PagedAttention 通过借鉴操作系统虚拟内存思想,实现了 LLM 推理系统的内存革命:
✅ 核心创新:
- 分页管理 KV 缓存,消除内存碎片
- Block Table 映射逻辑 - 物理地址
- 支持序列内和跨序列 KV 共享
✅ 实际效果:
- 内存利用率从 50% 提升到 96%
- 吞吐量提升 2-4 倍
- 并发请求数提升 4 倍
✅ 生态影响:
- 被广泛集成到推理框架(vLLM、TGI、SGLang 等)
- 成为大模型服务的标准技术
- 显著降低部署成本
参考文献
- Kwon et al. “Efficient Memory Management for Large Language Model Serving with PagedAttention.” SOSP 2023.
- vLLM GitHub: https://github.com/vllm-project/vllm
- vLLM Documentation: https://docs.vllm.ai
本文基于技术文档整理,如有错误欢迎指正。