Skip to content

vLLM 与 PagedAttention

本页速览 UC Berkeley 2023 年开源的 LLM 推理引擎,用 PagedAttention 把 KV Cache 像虚拟内存一样分页管理,配合 continuous batching 让 Llama-65B 的吞吐直接涨一个量级。本文拆解 PagedAttention、prefix caching、量化支持、与 SGLang/TGI 的对比。

vLLM 与 PagedAttention

一、概念定义:把 LLM 推理从"重武器"变成"标配组件"

vLLM 是 UC Berkeley Sky Computing Lab 于 2023 年 6 月开源的高性能 LLM 推理引擎,论文《Efficient Memory Management for Large Language Model Serving with PagedAttention》发表于 SOSP 2023。它解决的问题可以用一个数字概括:HuggingFace Transformers 原生推理只能跑到 GPU 理论吞吐的 10–25%,vLLM 把这个数字拉到 60–90%——典型场景下吞吐提升 10–24×

理解 vLLM 的意义要先理解 LLM 推理的瓶颈:

  • Decode 阶段是 memory-bound(详见 显存带宽Roofline 模型):每生成一个 token,都要把整个模型权重 + 历史 KV cache 读一遍,但只算很少 FLOPs。
  • KV Cache 巨大且管理粗放:Llama-2-7B 单请求 2048 序列长 KV cache ~5 GB,传统框架预分配"最大序列长"的连续显存——大量请求不满,浪费 60–80%。
  • 批处理难做:请求长度不一、生成步数不一、长短请求混跑时短请求等长请求,GPU 利用率上不去。

vLLM 用一个核心创新(PagedAttention)+ 两个工程组合(continuous batching + prefix caching)一次性把这三个问题解决了。它现在是开源 LLM 推理的事实标准,与 TensorRT-LLM、SGLang、TGI 并列主流。

二、PagedAttention:OS 虚拟内存思想搬到 LLM

PagedAttention 是 vLLM 的核心创新,灵感直接来自操作系统的分页式虚拟内存

传统 KV Cache 的问题

传统框架给每个请求预分配一段连续显存,按 max_seq_len 算:

请求 A: max 2048 → 实际用 800   [████████░░░░░░░░░░░░░░] 浪费 1248
请求 B: max 2048 → 实际用 1500  [█████████████████░░░░░] 浪费 548
请求 C: max 2048 → 实际用 600   [██████░░░░░░░░░░░░░░░░] 浪废 1448
                                                   合计 60% 显存浪费
  • 内部碎片:预分配大、实际用小,浪费严重。
  • 外部碎片:请求来去之间产生空洞,难以重新利用。
  • 无法共享:相同 prompt 前缀(system prompt)的 KV 不能复用,每个请求重算一遍。

PagedAttention 的解法

借鉴 OS 虚拟内存:把 KV cache 按 block(块大小通常 16)分页,逻辑地址 → 物理块用 block table 映射。

逻辑 KV:  [B0][B1][B2][B3]      ← 单请求的"虚拟地址"
              │  │  │  │
              ▼  ▼  ▼  ▼
block table: │ 1| 3| 7| 9 |     ← 物理块编号(物理显存池散布)
              ▼     ▼
物理显存池:  [B0][B1][B2][B3][B4][B5][B6][B7][B8][B9]...
                    ▲            ▲           ▲
                    └── 请求 A   └── 请求 B  └── 请求 C
  • 请求内部按需分配:生成到第 100 个 token 才申请第 7 个 block,没有预分配浪费。
  • 物理显存紧凑:所有请求的 block 散布在统一池子,无外部碎片。
  • 可共享前缀:相同前缀的请求指向同一物理 block,写时复制(CoW)。

注意力计算改造

传统 attention 是 "Q × K^T" 一次性算,K/V 在连续显存里。PagedAttention 改成按块加载 K/V——每个 block 单独计算 attention,再累加。这增加了少量 kernel 调度开销,但换来显存利用率的飞跃。

CUDA kernel 关键结构(伪代码):

cpp
// PagedAttention CUDA kernel 简化版
__global__ void paged_attention(
    float* output,            // [seq_len, d_model]
    const float* q,           // [num_heads, d_head]
    const float* key_cache,   // [num_blocks, block_size, num_heads, d_head]
    const float* value_cache, // [num_blocks, block_size, num_heads, d_head]
    const int* block_table,   // [max_num_blocks_per_seq]
    int context_len,          // 当前请求的历史长度
    int block_size            // 通常 16
) {
    int block_idx = block_table[block_id];   // 逻辑 → 物理
    // 在物理块内计算 attention
    ...
}

详见 核心论文 里的 PagedAttention 论文精读。

三、Continuous Batching:每步重组批次

传统批处理("static batching"):凑齐一个 batch 一起进、一起出。问题:短请求被长请求卡住,GPU 在等最后一个请求完成时空转。

Continuous batching(也叫 "iteration-level batching" 或 "in-flight batching")的做法:每个生成 step 都重新组 batch——

Step 0:  [A B C D]      4 个请求一起 prefill
Step 1:  [A B C D]      4 个一起 decode
Step 2:  [A B C D]      A 完成退出
Step 3:  [_ B C D E]    E 进来填 A 的位
Step 4:  [_ B C D E]    D 完成退出
Step 5:  [F B C _ E]    F 进来,C 还在继续
...

每个 step 都能在 prefill 队列里捞请求进来,已完成请求立刻退出。GPU 永远是满的,吞吐最大化。详见 批处理与调度

vLLM 的调度器还做了 preemption(抢占):显存压力大时把某些请求的 KV 换出到 CPU 内存(recompute 或 swap),等显存空出再换回。这是 LLM serving 的关键技术。

四、Prefix Caching:复用 system prompt 的 KV

LLM 应用普遍结构是:长 system prompt + 短 user query。比如客服 bot:

system: "你是 XX 公司的客服助手,负责回答以下问题类型..."  ← 1500 token
user:   "我的订单什么时候发货?"                          ← 8 token

每个请求都 prefill 那 1500 token system prompt 是巨大浪费。vLLM 的 Automatic Prefix Caching

  1. 第一个请求 prefill 时把 KV 存到 LRU 缓存;
  2. 后续请求 system prompt 一致 → 直接复用,只 prefill user query 部分;
  3. 节省 80–95% 的 prefill 计算。

效果(典型客服场景):

配置吞吐(tokens/s)首 token 延迟
不开 prefix caching1200220 ms
开 prefix caching850025 ms

vLLM 0.5+ 起默认开启 prefix caching。

五、量化支持:AWQ / GPTQ / FP8 / INT8 / INT4

vLLM 支持主流权重量化方案,详见 量化权重-激活混合精度

量化方法位宽vLLM 支持备注
GPTQINT4✅ 原生EXL2 / Marlin kernel 加速
AWQINT4✅ 原生通常比 GPTQ 精度好 0.3–0.5%
Marlin (GPTQ 优化)INT4极致 GPU kernel,速度快 1.5–2×
bitsandbytesINT8 / NF4兼容性好、性能一般
FP8 (E4M3)FP8✅ (H100+)Hopper 卡杀手锏,详见 TensorRT-LLM
INT8 (W8A16)INT8老卡兜底
INT4 (W4A16)INT4Llama-2-7B 单卡 8GB 显存可跑
bash
# 启动 AWQ 量化模型
python -m vllm.entrypoints.openai.api_server \
    --model TheBloke/Llama-2-7B-Chat-AWQ \
    --quantization awq_marlin \
    --dtype float16 \
    --max-model-len 4096 \
    --gpu-memory-utilization 0.9

六、版本演进:从 0.2 到 0.7+

vLLM 是个高速演进的项目,每个版本都加东西:

版本时间关键特性
v0.22023.08PagedAttention + continuous batching 首发
v0.32023.11AWQ / GPTQ 量化支持
v0.42024.02Prefix caching、tensor parallel 改进
v0.52024.06Speculative decoding(EAGLE/lookahead)、chunked prefill、默认 prefix caching
v0.62024.10CUDA graphs、FP8 量化、改进调度器
v0.7+2025.01+MTP / EAGLE-3 集成、增强多模态、reasoning 模型支持

关键的工程优化

Chunked prefill:把长 prompt 的 prefill 切成 chunk,与 decode 混合调度,避免长 prompt 把 decode 卡死。CUDA graphs:把生成 step 的整个 kernel 序列录制成 graph,launch 开销从十几次降到一次。这两项让 vLLM 在 0.6+ 后延迟进一步下降 30–50%。

七、性能数据:基线参考

下面给一组 H100 80GB 上的基线(vLLM 0.6.x,详见 基准测试 方法论):

模型精度并发吞吐(tokens/s)首 token 延迟
Llama-3-8BBF1632~500030 ms
Llama-3-8BFP832~850030 ms
Llama-3-8BAWQ-INT432~750032 ms
Llama-3-70BBF16 (2×H100 TP)32~150080 ms
Llama-3-70BFP8 (1×H100)32~170075 ms
Qwen2-72BAWQ-INT432~200075 ms
DeepSeek-R1-Distill-7BBF1632~550030 ms

对比 HuggingFace Transformers(FP16,batch=8):

Llama-3-8B HF Transformers:  ~280 tokens/s
Llama-3-8B vLLM (BF16):      ~5000 tokens/s    (18× 加速)
Llama-3-8B vLLM (FP8):       ~8500 tokens/s    (30× 加速)

八、与 SGLang / TGI / TensorRT-LLM 对比

引擎主要特色与 vLLM 关系
vLLMPagedAttention、社区生态最广、HuggingFace 模型即开即用基准
SGLang2024 年 UC Berkeley 同源作品,RadixAttention 前缀树 + 结构化输出 + 程序化编排同源竞争,SGLang 在复杂应用(结构化输出、Agent、tree search)更强;vLLM 在基础 serving 更广
TGI (HuggingFace)HF 官方、与 HF 生态深度集成性能稍弱(约 vLLM 70–85%),但 HF 接入成本低
TensorRT-LLMNVIDIA 官方、H100 上极致优化TRT-LLM 在 H100 上略快(1.1–1.3×)、build 流程重;vLLM 易用性胜出
DeepSpeed-MII微软出品、ZeRO-Inference 支持超显存模型偏研究、生态较小
MLC-LLM通用编译器路径,跨平台强项在 WebGPU / 移动端(见 移动端部署

引擎选择

  • 起步、PoC、中小团队 → vLLM(生态、易用、文档最全)
  • H100 集群极致生产 → TensorRT-LLM 或 vLLM 0.7+ CUDA graphs,benchmark 实测决胜负
  • 结构化输出 / Agent / tree search → SGLang
  • HF 生态深度集成 → TGI
  • 跨平台 / WebGPU / 移动 → MLC-LLM

九、部署示例:OpenAI 兼容 API 一行起

bash
# 最简:单卡起 vLLM OpenAI 兼容服务
python -m vllm.entrypoints.openai.api_server \
    --model meta-llama/Meta-Llama-3-8B-Instruct \
    --port 8000 \
    --max-model-len 8192 \
    --gpu-memory-utilization 0.9 \
    --enforce-eager                     # 不开 CUDA graphs(调试用)

# 多卡张量并行(TP=2)
python -m vllm.entrypoints.openai.api_server \
    --model meta-llama/Meta-Llama-3-70B-Instruct \
    --tensor-parallel-size 2 \
    --max-model-len 4096

# 调用(OpenAI 协议兼容)
curl http://localhost:8000/v1/chat/completions \
  -H "Content-Type: application/json" \
  -d '{
    "model": "meta-llama/Meta-Llama-3-8B-Instruct",
    "messages": [{"role": "user", "content": "什么是 PagedAttention?"}]
  }'
python
# Python API
from vllm import LLM, SamplingParams

llm = LLM(model="meta-llama/Meta-Llama-3-8B-Instruct",
          tensor_parallel_size=1,
          max_model_len=8192,
          quantization=None)            # None / "awq" / "gptq" / "fp8"

prompts = ["用一句话解释 PagedAttention", "什么是 continuous batching?"]
sampling = SamplingParams(temperature=0.7, top_p=0.9, max_tokens=128)

outputs = llm.generate(prompts, sampling)
for o in outputs:
    print(o.outputs[0].text)

十、局限与边界

  1. Batch 极大时性能不再增长:decode 阶段是 memory-bound,并发超过某阈值后 GPU 显存带宽打满,吞吐见顶。这时投机解码(投机解码)等"省带宽"技术更有效。
  2. 定制模型支持滞后:新架构(如 Mamba、新 MoE 变体)在 vLLM 上落地比 HF Transformers 慢几周到几月。
  3. 结构化输出弱于 SGLang:vLLM 的 guided decoding(outlines / xgrammar)在性能和灵活度上不及 SGLang 的 RadixAttention + 复杂程序化编排。
  4. NPU / CPU 性能差:vLLM 是 GPU-first,CPU / NPU 后端较弱,移动端请走 llama.cppONNX Runtime
  5. 多模态支持在演进:vLLM 0.5+ 支持图像输入(LLaVA、Qwen-VL),但生态与优化程度不如纯文本。
  6. TP 上限 8:单机 8 卡内 TP,跨机请配合 PP / DP,参见 分布式推理

十一、可继续追踪

参考资料