连续批处理:推理吞吐的范式革命

在 LLM 推理系统中,如何高效处理并发请求是决定服务成本的核心问题。连续批处理(Continuous Batching)配合 PagedAttention,让 vLLM 实现了比传统推理框架高 5-20 倍的吞吐量。本文深入解析这项技术的原理与工程实现。

一、批处理的演进

1.1 静态批处理(Static Batching)

最简单的批处理方式:等待凑齐一批请求,然后一起推理,直到所有请求完成。

┌────────────────────────────────────────────────────┐
│              静态批处理                              │
├────────────────────────────────────────────────────┤
│                                                    │
│  时间→ 0    1    2    3    4    5    6    7       │
│                                                    │
│  Req1  [G][G][G][G][G][D]              ✓          │
│  Req2  [G][G][G][G][G][G][G][G]        ✓          │
│  Req3  [G][G][D]                       ✓          │
│  Req4  [G][G][G][G][G][G][G][D]        ✓          │
│                                                    │
│  G=生成, D=完成(停止生成)                           │
│                                                    │
│  问题: Req3 在 t=2 就完成了,                       │
│        但要等 Req2 到 t=7 整批才结束               │
│  GPU 利用率: ~35%                                  │
└────────────────────────────────────────────────────┘

核心问题:Padding 浪费 + 等待浪费。不同请求长度差异大,短请求要等长请求完成。

1.2 动态批处理(Dynamic Batching)

Orum 等人 2022 年提出连续批处理(也叫 iteration-level batching 或 in-flight batching):

核心思想:不是等到整批完成,而是每一步都可以加入新请求或移除完成的请求。

┌────────────────────────────────────────────────────┐
│              连续批处理                              │
├────────────────────────────────────────────────────┤
│                                                    │
│  时间→ 0    1    2    3    4    5    6    7       │
│                                                    │
│  Req1  [G][G][G][G][G][D]                         │
│  Req2  [G][G][G][G][G][G][G][G][D]               │
│  Req3  [G][G][D]                                  │
│  Req4               [G][G][G][G][G][G][D]        │
│  Req5                  [G][G][G][G][G][D]        │
│  Req6                        [G][G][G][G][G]...   │
│                                                    │
│  t=2: Req3 完成, 移除                              │
│  t=3: Req4 加入                                    │
│  t=4: Req5 加入                                    │
│  t=6: Req6 加入                                    │
│                                                    │
│  GPU 始终满载! 利用率: ~90%                        │
└────────────────────────────────────────────────────┘

1.3 优势对比

指标静态批处理连续批处理
GPU 利用率20-40%80-95%
平均延迟高(等凑批)低(即来即处理)
P99 延迟不稳定稳定
吞吐量1x5-20x
实现复杂度简单

二、连续批处理的技术挑战

2.1 变长序列的 KV Cache 管理

连续批处理中,不同请求的长度不同,且随时有请求加入和离开。这导致 KV Cache 的分配变得复杂:

传统分配 (连续分配):
┌───────────────────────────────────────────────┐
│Req1 KV Cache (4KB)│Req2 KV Cache (8KB)│  空闲  │
└───────────────────────────────────────────────┘
问题: Req1 完成后, 释放的空间可能不够新请求使用 → 碎片化

PagedAttention (分页分配):
┌────┬────┬────┬────┬────┬────┬────┬────┐
│P:1 │P:1 │P:2 │P:2 │P:2 │P:3 │P:3 │Free│
│blk0│blk1│blk2│blk3│blk4│blk5│blk6│    │
└────┴────┴────┴────┴────┴────┴────┴────┘
Req1 → [blk0, blk1]
Req2 → [blk2, blk3, blk4]
Req3 → [blk5, blk6]

Req1 完成, 释放 blk0, blk1:
┌────┬────┬────┬────┬────┬────┬────┬────┐
│Free│Free│P:2 │P:2 │P:2 │P:3 │P:3 │Free│
│blk0│blk1│blk2│blk3│blk4│blk5│blk6│    │
└────┴────┴────┴────┴────┴────┴────┴────┘
新请求可以立即使用 blk0, blk1 → 无碎片!

2.2 PagedAttention 原理

PagedAttention 借鉴操作系统的虚拟内存管理:

┌──────────────────────────────────────────────────┐
│           PagedAttention vs OS 虚拟内存           │
├──────────────────────────────────────────────────┤
│                                                  │
│  OS 虚拟内存            PagedAttention           │
│  ────────────           ──────────────           │
│  进程                   请求 (Request)            │
│  虚拟地址空间           逻辑 KV 位置              │
│  页表 (Page Table)      Block Table              │
│  物理页 (4KB)           KV Block (16 tokens)     │
│  页面置换               KV Cache 驱逐            │
│                                                  │
│  请求的 Block Table:                              │
│  ┌────────────────────────────────┐              │
│  │ Logical Block │ Physical Block │              │
│  │      0        │      7         │              │
│  │      1        │      12        │              │
│  │      2        │      3         │              │
│  │      3        │      19        │              │
│  └────────────────────────────────┘              │
│                                                  │
│  逻辑上连续, 物理上分散 — 无碎片化问题             │
└──────────────────────────────────────────────────┘

每个 KV Block 存储固定数量 Token 的 KV Cache(通常 16 个 Token)。Block Table 记录逻辑块到物理块的映射。

2.3 Block 大小的选择

Block 大小 $B$ 是关键参数:

Block 大小内存碎片管理开销注意力效率
4 tokens最低最高低(频繁寻址)
16 tokens
64 tokens最高
128 tokens最低

vLLM 默认使用 16 tokens 的 Block 大小,在碎片和效率间取得平衡。

三、vLLM 的连续批处理实现

3.1 调度器架构

┌─────────────────────────────────────────────────────┐
│                  vLLM 调度器架构                      │
├─────────────────────────────────────────────────────┤
│                                                     │
│  ┌──────────┐   请求队列   ┌──────────────┐        │
│  │ Request  │─────────────►│   Scheduler  │        │
│  │ Queue    │              │              │        │
│  └──────────┘              │  - 选择请求   │        │
│                            │  - 分配 Block │        │
│                            │  - 驱逐策略   │        │
│                            └──────┬───────┘        │
│                                   │                 │
│                    ┌──────────────┼──────────────┐ │
│                    ▼              ▼              ▼ │
│              ┌──────────┐  ┌──────────┐  ┌──────┐ │
│              │ Running  │  │  Waiting  │  │Swapped│ │
│              │  Batch   │  │  Queue    │  │ Queue │ │
│              └────┬─────┘  └──────────┘  └──────┘ │
│                   │                                 │
│                   ▼                                 │
│            ┌──────────────┐                         │
│            │ Model Forward│                         │
│            │ (PagedAttn)  │                         │
│            └──────────────┘                         │
│                                                     │
└─────────────────────────────────────────────────────┘

3.2 调度策略

vLLM 的调度器在每个时间步执行以下决策:

# 简化的调度逻辑
def schedule(running, waiting, swapped, free_blocks):
    # 1. 移除已完成的请求, 释放 KV Cache
    for req in running:
        if req.is_finished():
            free_blocks += req.blocks
            running.remove(req)
    
    # 2. 尝试从 swapped 恢复请求 (优先级最高)
    for req in swapped:
        if can_allocate(req, free_blocks):
            running.append(req)
            swapped.remove(req)
            free_blocks -= req.needed_blocks
    
    # 3. 尝试从 waiting 加入新请求
    for req in waiting:
        if can_allocate(req, free_blocks):
            running.append(req)
            waiting.remove(req)
            free_blocks -= req.needed_blocks
        else:
            break  # 显存不足, 停止加入
    
    # 4. 如果显存仍然不足, 换出部分 running 请求
    while free_blocks < min_needed and running:
        victim = select_victim(running)  # 选择换出对象
        swapped.append(victim)
        running.remove(victim)
        free_blocks += victim.blocks
    
    return running

3.3 Preemption(抢占)策略

当显存不足时,vLLM 有两种抢占策略:

1. Swap(换出):将 KV Cache 移到 CPU 内存

  • 优势:恢复时不需要重新计算
  • 劣势:PCIe 传输延迟

2. Recompute(重算):丢弃 KV Cache,恢复时重新计算

  • 优势:无内存传输
  • 劣势:需要重新前向传播
def should_swap_or_recompute(req, gpu_memory, cpu_memory):
    swap_time = 2 * req.kv_cache_size / pcie_bandwidth  # 传输时间
    recompute_time = req.seq_len * forward_latency       # 重算时间
    
    if swap_time < recompute_time and cpu_memory_available:
        return "swap"
    else:
        return "recompute"

经验值:对于序列长度 <2000 的请求,Recompute 更快;对于长序列,Swap 更快。

四、注意力计算与 PagedAttention

4.1 Paged Attention 的 CUDA 实现

标准注意力计算假设 KV Cache 在内存中连续存储,但 PagedAttention 中 KV Cache 分散在不同 Block 中:

标准 Attention:
  Q @ K[0:n].T  → K 在连续内存中, 一次 matmul

Paged Attention:
  for each block b:
    Q @ K[b].T  → K 在不同 Block 中, 多次小 matmul
  合并 softmax 结果

PagedAttention 的 CUDA Kernel 实现:

// 伪代码: PagedAttention CUDA Kernel
__global__ void paged_attention_kernel(
    float* output,       // [num_tokens, d]
    float* q,            // [num_tokens, d]
    float* key_cache,    // [num_blocks, block_size, d]
    float* value_cache,  // [num_blocks, block_size, d]
    int* block_tables,   // [num_requests, max_blocks]
    int* seq_lens,       // [num_requests]
    int d, int block_size
) {
    int token_idx = blockIdx.x;
    int head_idx = blockIdx.y;
    
    // 获取该 token 的 Q 向量
    float* q_vec = q + token_idx * d + head_idx * d_k;
    
    // 遍历该请求的所有 KV Block
    int req_idx = get_req_idx(token_idx);
    int seq_len = seq_lens[req_idx];
    int num_blocks = (seq_len + block_size - 1) / block_size;
    
    float max_score = -INF;
    float sum_exp = 0;
    float* partial_output = shared_memory;  // 使用 SRAM
    
    for (int b = 0; b < num_blocks; b++) {
        int physical_block = block_tables[req_idx * max_blocks + b];
        
        // 在 Block 内计算注意力
        for (int i = 0; i < block_size; i++) {
            int kv_pos = b * block_size + i;
            if (kv_pos >= seq_len) break;
            
            float* k_vec = key_cache[physical_block * block_size * d 
                                     + i * d + head_idx * d_k];
            float score = dot_product(q_vec, k_vec) / sqrt(d_k);
            
            // 在线 softmax
            if (score > max_score) {
                rescale(partial_output, exp(max_score - score));
                max_score = score;
            }
            float exp_score = exp(score - max_score);
            sum_exp += exp_score;
            
            float* v_vec = value_cache[physical_block * block_size * d 
                                       + i * d + head_idx * d_k];
            axpy(partial_output, v_vec, exp_score);
        }
    }
    
    // 归一化输出
    scale(partial_output, 1.0 / sum_exp);
    copy(output + token_idx * d + head_idx * d_k, partial_output, d_k);
}

4.2 Prefix Sharing

多个请求共享相同的前缀(如 system prompt),PagedAttention 可以让这些请求共享 KV Block:

请求1: [System Prompt][User Query 1]
请求2: [System Prompt][User Query 2]
请求3: [System Prompt][User Query 3]

KV Cache 布局:
┌──────────────┬──────────────┬──────────────┐
│ Shared KV    │ Req1 KV      │ Req2 KV      │
│ (System)     │ (Query 1)    │ (Query 2)    │
└──────────────┴──────────────┴──────────────┘
       │              ▲              ▲
       └──────────────┴──────────────┘
       三个请求共享 System Prompt 的 KV Block

节省: 3x System Prompt 的 KV Cache!

五、性能优化

5.1 吞吐量 vs 延迟的权衡

连续批处理有一个核心权衡:Batch Size 越大,吞吐量越高,但延迟也越高

    吞吐量
      │          ╱ ← 吞吐量
      │        ╱
      │      ╱
      │    ╱
      │  ╱
      │╱
      ├──────────────── → Batch Size
      │\                 ← 延迟
      │ \
      │   \
      │     \
      │       \
      │         \

最优区间: Batch Size = 32-128 (取决于模型和硬件)

5.2 参数调优

参数推荐值说明
max_num_seqs64-256最大并发请求数
max_model_len32768最大序列长度
block_size16KV Block 大小
swap_space4-16 GBCPU 换出空间
gpu_memory_utilization0.9GPU 显存利用率
preemption_mode“swap”抢占策略

5.3 实测性能

在 8x A100 80GB 上使用 Llama 4-70B:

配置吞吐 (tok/s)P50 延迟P99 延迟
静态批处理 (batch=8)420180ms350ms
连续批处理 (vLLM)2,80045ms120ms
连续批处理 + GQA4,10035ms95ms
连续批处理 + Speculative6,20028ms85ms

六、2026 年推理框架对比

框架连续批处理PagedAttention投机解码MoE 支持多模态
vLLM
TensorRT-LLM
SGLang
TGI部分部分
MLServer部分部分

vLLM 和 TensorRT-LLM 是 2026 年最成熟的高吞吐推理框架。SGLang 在特定场景(如结构化生成)表现优异。

七、总结

连续批处理和 PagedAttention 是大模型推理系统的基础设施:

  1. 连续批处理解决了变长请求的效率问题
  2. PagedAttention解决了 KV Cache 的内存碎片问题
  3. Prefix Sharing解决了重复前缀的内存浪费
  4. Preemption解决了显存不足时的优雅降级

这些技术组合在一起,让大模型推理的成本降低了 5-20 倍,是 LLM 从实验室走向生产的关键技术。

加入讨论

这篇文章有姊妹讨论帖在硅基AGI论坛 — 全球首个碳基硅基认知交流平台。

碳基与硅基的智慧碰撞,认知差异创造无限可能。