第四篇：vLLM 推理引擎：PagedAttention、Continuous Batching、KV Cache¶
系列目录：PyTorch → NVIDIA GPU 全链路深度解析
1. 问题的根源：HuggingFace Transformers 为什么慢¶
1.1 推理的两个阶段¶
LLM 推理分为 Prefill 和 Decode 两个截然不同的阶段：
Prefill (Prompt Processing):
输入: "What is the capital of France?"
→ 一次性编码整个序列，产出 KV cache
→ 计算复杂度 O(L²)，L = prompt 长度
Decode (Token Generation):
输入: 上一个生成的 token
→ 每次只处理 1 个 token，但需要访问整个 KV cache
→ 计算复杂度 O(L)，L = 已生成序列长度
→ 内存带宽瓶颈 (memory-bound)
HuggingFace Transformers 的 model.generate() 对这两个阶段不做区分优化：
# HuggingFace 典型推理代码
from transformers import AutoModelForCausalLM, AutoTokenizer
model = AutoModelForCausalLM.from_pretrained(
"meta-llama/Llama-2-7b-hf",
torch_dtype=torch.float16,
device_map="auto"
)
tokenizer = AutoTokenizer.from_pretrained("meta-llama/Llama-2-7b-hf")
# 慢的原因就在这里
outputs = model.generate(
input_ids,
max_new_tokens=256,
do_sample=True,
temperature=0.7
)
1.2 KV Cache 的朴素管理¶
transformers 如何管理 KV cache？每次调用模型前分配一块固定大小内存：
# transformers/generation/utils.py 的逻辑简化
class GenerationMixin:
def _prepare_model_kwargs_for_generation(self, ...):
batch_size = input_ids.shape[0]
# 每个请求预先分配 max_length 的大小的 cache
# 极端浪费内存！
max_cache_length = self.generation_config.max_length
# shape: (batch, num_heads, max_cache_length, head_dim)
past_key_values = [
(torch.empty(batch, num_heads, max_cache_length, head_dim),
torch.empty(batch, num_heads, max_cache_length, head_dim))
for _ in range(num_layers)
]
return {"past_key_values": past_key_values}
问题显而易见： - 假设 max_length=2048，实际只生成 100 个 token → 浪费 ~95% 的 KV cache 内存 - 每个请求独占连续内存块 → 外部碎片（随请求完成释放后留下空洞） - 批处理中所有请求必须等最长那个完成 → “木桶效应”
1.3 内存碎片的量化影响¶
以 Llama-2-7B 为例，单层 KV cache 大小：
KV Cache per token = 2 (K+V) × num_heads × head_dim × num_layers × dtype_size
= 2 × 32 × 128 × 32 × 2 bytes
= 524,288 bytes ≈ 0.5 MB/token
For 256 max_new_tokens, 1 request:
Preallocated = 0.5 MB × 256 = 128 MB per request
For batch_size=32, each configured to max 256 tokens:
Preallocated = 128 MB × 32 = 4 GB
But effective utilization (avg 100 tokens actually generated):
4 GB × (100/256) = only ~1.56 GB actually used
Wasted: ~2.44 GB (61%)
更糟的是，在服务多个并发请求时，每个请求的 max_length 不同，频繁的 torch.empty 和释放导致 CUDA 内存碎片化，最终 CUDA OOM。
1.4 Batch 僵化¶
HuggingFace 的 batch 是“静态”的：
时间线 (4 个请求 A, B, C, D，长度各不相同):
Request A: ████████████████████████████ (28 tokens, 完成)
Request B: ████████████ (12 tokens, 完成)
Request C: ████████████████████████████████████████████████ (48 tokens)
Request D: ██████ (6 tokens, 完成)
传统批处理:
Batch A,B,C,D → 一起开始 → C 拖累所有人
A 早已完成，但 GPU 空闲等待 C 生成完毕
只有全部完成后才能形成下一个 batch
理想调度:
A 完成 → 立即从 batch 移除 → 加入新请求 E
B 完成 → 移除 → 加入 F
始终保持 GPU 满负荷
2. PagedAttention：KV Cache 的虚拟内存¶
2.1 核心类比¶
vLLM 的 PagedAttention 直接把操作系统的虚拟内存思想搬到 KV cache 管理上：
操作系统虚拟内存 PagedAttention KV Cache
─────────────────────────────────────────────────
物理页框 (page frame) → Block (固定数量 token 组)
页表 (page table) → Block Table (逻辑→物理映射)
虚拟地址空间 → 逻辑 KV cache (请求视角)
缺页异常 → 分配新 block
写时复制 (COW) → Block 共享 (beam search)
TLB → GPU 寄存器缓存 block table
2.2 Block 的数据结构¶
# vLLM 的核心数据结构（简化）
class KVCacheBlock:
"""
一个物理 block 存储固定数量 token 的 K 和 V。
所有 block 大小相同，便于统一管理。
"""
# 物理内存中的实际偏移
block_id: int
# 引用计数：被多少请求的 block table 引用
ref_count: int
class BlockTable:
"""
每个请求拥有的逻辑→物理映射表。
"""
# 逻辑 block → 物理 block 映射
# 例如: logical_block_0 → physical_block_17
# logical_block_1 → physical_block_42
# logical_block_2 → physical_block_3
block_ids: List[int]
# 当前序列 token 总数
num_tokens: int
def append_block(self, physical_block_id: int):
self.block_ids.append(physical_block_id)
self.num_tokens += BLOCK_SIZE
def __getitem__(self, logical_idx: int) -> int:
return self.block_ids[logical_idx]
2.3 GPU 内存布局¶
不同于 HuggingFace 的连续布局，PagedAttention 在 GPU 上的 key/value cache 是分块的：
HuggingFace 连续布局 (per request):
K_cache: [head0_seq, head1_seq, ..., head31_seq]
每个 head 的序列又按 token 0,1,2,... 连续存放
PagedAttention 分块布局 (全局):
K_cache: [num_blocks, num_kv_heads, block_size, head_dim]
# 例如 Llama-2-7B:
# num_blocks = 4096 (可管理的最大 block 数)
# num_kv_heads = 32 (GQA 下 KV head 数)
# block_size = 16 (每个 block 16 个 token)
# head_dim = 128
CUDA 端的实际内存分配：
// vLLM csrc/cache_kernels.cu 中的简化版本
// 全局 KV cache，所有请求共享
// tensor: [num_blocks, 2, num_kv_heads, block_size, head_dim]
// 其中 dim=1: 0=Key, 1=Value
torch::Tensor key_cache; // dtype: float16
torch::Tensor value_cache;
// Block tables: [max_num_seqs, max_num_blocks_per_seq]
// 值为 physical_block_id 或 -1 (空)
torch::Tensor block_tables; // dtype: int32
// 当前序列长度: [max_num_seqs]
torch::Tensor seq_lens; // dtype: int32
2.4 分配与回收：类似 Buddy Allocator¶
class BlockAllocator:
"""
物理 block 分配器，管理空闲列表。
"""
def __init__(self, num_blocks: int):
self.free_blocks = deque(range(num_blocks))
self.all_blocks = [KVCacheBlock(i) for i in range(num_blocks)]
# 注意：没有"外部碎片"——所有 block 大小相同
def allocate(self) -> int:
"""从空闲列表取一个 block"""
block_id = self.free_blocks.popleft()
self.all_blocks[block_id].ref_count = 1
return block_id
def free(self, block_id: int):
"""回收 block（引用计数为 0 时）"""
block = self.all_blocks[block_id]
block.ref_count -= 1
if block.ref_count == 0:
self.free_blocks.append(block_id)
def get_num_free_blocks(self) -> int:
return len(self.free_blocks)
2.5 Block 共享与 Copy-on-Write¶
PagedAttention 最精妙的设计：相同 prompt 前缀共享物理 block。
场景：两个请求 A 和 B 共享相同的系统提示词 "You are a helpful assistant."
物理 blocks: [P0]="You are", [P1]=" a helpful", [P2]=" assistant."
Request A 的 block table:
[P0, P1, P2, A3, A4, ...]
^ ^ ^
└───┴───┴── 共享
Request B 的 block table:
[P0, P1, P2, B3, B4, ...]
^ ^ ^
└───┴───┴── 共享
P0.ref_count = 2
P1.ref_count = 2
P2.ref_count = 2
A3.ref_count = 1
B3.ref_count = 1
Beam search 也受益于此：
def fork_for_beam_search(parent_block_table: BlockTable, fork_point: int):
"""
在 fork_point 处分裂出一个新的 block table。
从 fork_point 往后分配新 block（COW 语义）。
"""
# fork_point 之前：共享（更新引用计数）
shared_blocks = parent_block_table.block_ids[:fork_point]
for blk in shared_blocks:
allocator.all_blocks[blk].ref_count += 1
# fork_point 之后：新分配
new_blocks = []
for slot in parent_block_table.block_ids[fork_point:]:
# 分配新 block，复制原内容
new_blk = allocator.allocate()
copy_block(slot, new_blk)
new_blocks.append(new_blk)
return BlockTable(shared_blocks + new_blocks)
2.6 ASCII：PagedAttention Block Table 图解¶
┌─────────────────────────────────────────────────────────────────────────────┐
│ PagedAttention 架构总览 │
├─────────────────────────────────────────────────────────────────────────────┤
│ │
│ Request A: "What is the capital of France?" │
│ Request B: "What is the capital of Germany?" │
│ Request C: "How tall is Mount Everest?" │
│ │
│ Block Table (per request, logical → physical mapping): │
│ │
│ Request A Request B Request C │
│ ┌────┬─────┐ ┌────┬─────┐ ┌────┬─────┐ │
│ │ L0 │ P0 │ │ L0 │ P0 │ │ L0 │ P8 │ │
│ │ L1 │ P1 │ │ L1 │ P1 │ │ L1 │ P9 │ │
│ │ L2 │ P2 │ │ L2 │ P3 │ │ L2 │ P10 │ │
│ │ L3 │ A4 │ │ L3 │ B5 │ │ L3 │ P11 │ │
│ │ L4 │ A5 │ │ L4 │ B6 │ │ L4 │ P12 │ │
│ └────┴─────┘ └────┴─────┘ └────┴─────┘ │
│ │
│ 物理 Block Pool (每个 block = 16 tokens): │
│ ┌──────┬──────┬──────┬──────┬──────┬──────┬──────┬──────┬──────┬───┐ │
│ │ P0 │ P1 │ P2 │ P3 │ P4 │ P5 │ P6 │ P7 │ P8 │...│ │
│ │"What │" is "│" the"│" cap"│"ital"│" of "│" Fra"│"nce?"│"How "│ │ │
│ │ ref:2│ ref:2│ ref:1│ ref:1│ ref:0│ ref:0│ ref:0│ ref:0│ ref:1│ │ │
│ └──────┴──────┴──────┴──────┴──────┴──────┴──────┴──────┴──────┴───┘ │
│ ▲ ▲ │
│ │ │ │
│ A & B 共享 A & B 共享 (P0-P1 ref_count=2 因为有 A 和 B 引用) │
│ │
│ 注意: │
│ - P4~P7 是空闲 block (ref_count=0)，可被新请求使用 │
│ - A 和 B 前 2 个 block 共享，后 3 个各自独立 │
│ - 不需要连续物理 block，没有外部碎片 │
│ │
└─────────────────────────────────────────────────────────────────────────────┘
3. Attention Kernel：vLLM 的 CUDA 实现¶
3.1 普通的 scaled_dot_product_attention vs PagedAttention¶
# PyTorch 标准 attention（连续 KV cache）
def standard_attention(query, key, value):
"""
query: [num_tokens, num_heads, head_dim]
key: [num_tokens, num_kv_heads, head_dim] ← 连续内存
value: [num_tokens, num_kv_heads, head_dim]
"""
scale = 1.0 / math.sqrt(head_dim)
attn_weights = torch.matmul(query, key.transpose(-2, -1)) * scale
attn_weights = F.softmax(attn_weights, dim=-1)
output = torch.matmul(attn_weights, value)
return output
# PagedAttention — 需要根据 block table 做 gather+scatter
def paged_attention(query, key_cache, value_cache, block_table, seq_len):
"""
query: [num_tokens, num_heads, head_dim]
key_cache: [num_blocks, num_kv_heads, block_size, head_dim]
value_cache: [num_blocks, num_kv_heads, block_size, head_dim]
block_table: [num_blocks_per_seq] ← 逻辑→物理映射
seq_len: 当前序列长度
"""
# 伪代码：CUDA kernel 实际做的
# 1. 根据 block_table 从 key_cache/value_cache 中 gather KV
# 2. 计算 attention
# 3. 输出
pass
3.2 CUDA Kernel 核心逻辑¶
vLLM 的 CUDA kernel 源码位于 csrc/attention/，核心版本有 v1（基础）和 v2（优化）：
// vLLM csrc/attention/attention_kernels.cu (简化版本)
//
// PagedAttention v1: 每个 thread block 处理一个 query token
// 适用于 decode 阶段（每次只处理 1 个 query token）
#define THREAD_GROUP_SIZE (WARP_SIZE / 2) // 16 threads
#define NUM_THREAD_GROUPS (NUM_WARPS * 2) // 每个 thread block 的 group 数
#define VEC_SIZE 8 // 每次 load 8 个 half (128 bits)
template<typename scalar_t, int HEAD_SIZE, int BLOCK_SIZE, int NUM_THREADS>
__global__ void paged_attention_kernel(
scalar_t* __restrict__ out, // [num_seqs, num_heads, head_size]
const scalar_t* __restrict__ q, // [num_seqs, num_heads, head_size]
const scalar_t* __restrict__ k_cache, // [num_blocks, num_kv_heads, block_size, head_size]
const scalar_t* __restrict__ v_cache, // [num_blocks, num_kv_heads, block_size, head_size]
const int* __restrict__ block_tables, // [num_seqs, max_num_blocks_per_seq]
const int* __restrict__ context_lens, // [num_seqs]
const float scale,
const int max_context_len
) {
// 每个 thread group 负责一个 head
const int warps_per_block = NUM_THREADS / WARP_SIZE;
const int thread_group_idx = threadIdx.x / THREAD_GROUP_SIZE;
const int warp_idx = thread_group_idx / (NUM_THREAD_GROUPS / warps_per_block);
const int head_idx = thread_group_idx % (NUM_THREAD_GROUPS / warps_per_block);
extern __shared__ float shared_mem[];
// ============ Step 1: 加载 query 到 shared memory ============
scalar_t* q_vec = (scalar_t*)shared_mem; // [head_size]
float* qk_scores = (float*)(shared_mem + HEAD_SIZE * sizeof(scalar_t));
// 每个 thread group 合作加载一个 head 的 query vector
for (int i = threadIdx.x % THREAD_GROUP_SIZE;
i < HEAD_SIZE;
i += THREAD_GROUP_SIZE) {
q_vec[i] = q[seq_idx * num_heads * HEAD_SIZE
+ head_idx * HEAD_SIZE
+ i];
}
__syncthreads();
// ============ Step 2: 遍历 KV cache blocks，计算 attention ============
float acc[VEC_SIZE] = {0.0f}; // 每个 thread 负责 VEC_SIZE 个 QK 分片
int num_blocks = (max_context_len + BLOCK_SIZE - 1) / BLOCK_SIZE;
for (int block_idx = 0; block_idx < num_blocks; block_idx++) {
// 通过 block table 获取物理 block ID
int physical_block = block_tables[seq_idx * max_num_blocks_per_seq
+ block_idx];
// 无效 block (padding) → 跳过
if (physical_block < 0) continue;
// 加载 Key block
int kv_head = head_idx % num_kv_heads; // GQA/MQA: 多个 head 共享 KV
scalar_t* k_block = k_cache
+ physical_block * num_kv_heads * BLOCK_SIZE * HEAD_SIZE
+ kv_head * BLOCK_SIZE * HEAD_SIZE;
// ========== Step 2a: Q·K^T ==========
float qk = 0.0f;
// 当前 block 内的每个 token
for (int tok = 0; tok < BLOCK_SIZE; tok++) {
// 检查 token 是否有效（最后一个 block 可能不满）
int global_tok = block_idx * BLOCK_SIZE + tok;
if (global_tok >= context_lens[seq_idx]) break;
// 用 FMA (fused multiply-add) 计算点积
qk = 0.0f;
#pragma unroll
for (int d = 0; d < HEAD_SIZE; d++) {
qk += float(q_vec[d]) * float(k_block[tok * HEAD_SIZE + d]);
}
qk_scores[global_tok] = qk * scale;
}
// 注意：这里简化了 softmax（实际 kernel 用 online softmax）
// 完整实现见下文
}
// ============ Step 3: Online Softmax + V·softmax(QK^T) ============
// 使用在线算法避免存储整个 attention 矩阵
float max_score = -FLT_MAX;
float sum_exp = 0.0f;
float output[VEC_SIZE] = {0.0f};
for (int block_idx = 0; block_idx < num_blocks; block_idx++) {
int physical_block = block_tables[seq_idx * max_num_blocks_per_seq
+ block_idx];
if (physical_block < 0) continue;
// 重新计算 QK（因为之前的结果被丢弃了）
// 实际 kernel 中会分批处理，每批次做一次 rescale
// Online softmax rescale:
// m_new = max(m_old, m_block)
// acc = acc * exp(m_old - m_new) + v_block * exp(qk_block - m_new)
// sum_exp = sum_exp * exp(m_old - m_new) + sum(exp(qk_block - m_new))
}
// 最终: output = acc / sum_exp
// 写回全局内存
for (int d = threadIdx.x; d < HEAD_SIZE; d += NUM_THREADS) {
out[seq_idx * num_heads * HEAD_SIZE + head_idx * HEAD_SIZE + d]
= scalar_t(output[d]);
}
}
3.3 Online Softmax 算法¶
PagedAttention 使用 Online Softmax 避免 O(n²) 内存占用：
# 朴素 softmax（需要 O(n²) 内存存整个 attention matrix）
def naive_softmax_attention(Q, K, V):
scores = Q @ K.T / sqrt(d) # [seq_len, seq_len] ← 内存杀手
attn = softmax(scores, dim=-1)
return attn @ V
# Online Softmax（只需 O(1) 额外内存）
def online_softmax_attention(Q, K, V):
"""
分块处理 KV，每处理一个 block 就部分聚合结果。
内存开销：O(block_size) 而非 O(seq_len²)
"""
d = Q.shape[-1]
scale = 1.0 / math.sqrt(d)
m = -float('inf') # 当前最大值（用于数值稳定）
d_sum = 0.0 # 分母累加
acc = torch.zeros_like(Q) # 累加器
# 分块遍历 K, V
for block_start in range(0, K.shape[0], BLOCK_SIZE):
K_block = K[block_start:block_start + BLOCK_SIZE]
V_block = V[block_start:block_start + BLOCK_SIZE]
# Q·K^T for this block: [1, block_size]
qk_block = (Q @ K_block.T) * scale
# Online softmax 更新
m_new = max(m, torch.max(qk_block))
# Rescale: 旧累加器按新的 max 重新归一化
correction = torch.exp(m - m_new)
d_sum = d_sum * correction + torch.sum(torch.exp(qk_block - m_new))
# 加权累计 V
acc = acc * correction + (torch.exp(qk_block - m_new) @ V_block)
m = m_new
return acc / d_sum
CUDA 实现中的关键优化：
// vLLM 中的 online softmax rescale（简化）
// 每个 thread group 维护自己的 m, d_sum, acc
float m_curr = m_prev;
float d_sum_curr = d_sum_prev;
for (int block_idx = 0; block_idx < num_blocks; block_idx++) {
// ... 计算 qk_block 和 v_block ...
float m_block = warp_reduce_max(qk_block); // block 内部最大值
float m_new = fmaxf(m_curr, m_block);
// Rescale 因子
float scale_old = expf(m_curr - m_new);
float scale_new = expf(m_block - m_new);
// 更新累加器
d_sum_curr = d_sum_curr * scale_old + sum_exp(qk_block * scale_new);
// 更新 output 累加 (V * softmax(QK^T) 部分)
for (int j = 0; j < VEC_SIZE; j++) {
acc[j] = acc[j] * scale_old + v_block[j] * (qk_block * scale_new);
}
m_curr = m_new;
}
// 最终归一化
for (int j = 0; j < VEC_SIZE; j++) {
acc[j] /= d_sum_curr;
}
3.4 PagedAttention v2 优化¶
v2 kernel 针对 prefill 阶段（多个 query token）做了专门优化：
// PagedAttention v2: 每个 thread block 处理多个 query token
// 适用于 prefill 阶段或 chunked prefill
template<int HEAD_SIZE, int BLOCK_SIZE>
__global__ void paged_attention_v2_kernel(
// ... 参数同 v1 ...
) {
// v2 关键优化:
// 1. 使用 warp specialization: 部分 warp 做 QK，部分 warp 做 PV
// 通过 pipeline 隐藏延迟
// 2. 利用 TMA (Hopper) 异步加载 KV cache 到 shared memory
// 3. block 间并行：不同 block 处理不同的 query token
// 4. 使用 FP8 KV cache（节省 50% 内存带宽）
// TMA 异步加载示例（仅 H100+）
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 900
asm volatile(
"cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes"
" [%0], [%1], %2;\n"
:: "r"(shared_mem_ptr), "l"(global_k_block_ptr), "r"(copy_size)
);
#endif
}
4. Continuous Batching¶
4.1 Scheduler 核心逻辑¶
vLLM 的 scheduler 以迭代方式运行，每次迭代挑选一批请求执行一个 step：
# vLLM vllm/core/scheduler.py (简化)
class Scheduler:
"""
每次迭代 (iteration) 做以下事情:
1. 检查完成的请求 → 释放 KV cache blocks
2. 从 waiting queue 中选择新请求 → 分配 blocks
3. 如果 OOM → 抢占 (preempt) 低优先级的请求
4. 返回 ScheduledSequenceGroup (本轮运行的请求)
"""
def __init__(self, scheduler_config, cache_config):
self.waiting: deque = deque() # 等待中的请求
self.running: deque = deque() # 正在运行的请求
self.swapped: deque = deque() # 被抢占到 CPU 的请求
self.block_manager = BlockSpaceManager(
block_size=cache_config.block_size,
num_gpu_blocks=cache_config.num_gpu_blocks,
num_cpu_blocks=cache_config.num_cpu_blocks,
)
def schedule(self) -> SchedulerOutputs:
"""
调度决策的核心循环。
返回一个 (seq_group_metadata_list, scheduler_outputs) 元组。
"""
# Step 1: 收集已完成请求
# 当请求的最后一个 token 是 EOS 或达到 max_tokens → 标记完成
ignored_seq_groups = []
for seq_group in list(self.running):
if seq_group.is_finished():
# 释放 KV cache blocks
for seq in seq_group.get_seqs():
self.block_manager.free(seq)
self.running.remove(seq_group)
ignored_seq_groups.append(seq_group)
# Step 2: 调度新请求 (从 waiting 到 running)
scheduled = []
preempted = []
# 按 FCFS + 优先级排序
self.waiting = deque(
sorted(self.waiting, key=lambda sg: sg.arrival_time)
)
while self.waiting:
seq_group = self.waiting[0]
# 尝试为请求分配 KV cache blocks
num_required_blocks = self._get_num_required_blocks(seq_group)
if self.block_manager.can_allocate(seq_group):
# 有足够的空闲 blocks → 分配并加入 running
self.block_manager.allocate(seq_group)
self.waiting.popleft()
self.running.append(seq_group)
scheduled.append(seq_group)
else:
# 内存不足 → 考虑抢占
if self._can_preempt():
preempted_group = self._preempt_lowest_priority()
preempted.append(preempted_group)
else:
break # 无法分配更多，停止
# Step 3: 创建 ScheduledSequenceGroup
# 每个 running 请求生成 1 个 token
seq_group_metadata_list = []
for seq_group in self.running:
seq_data = {}
for seq in seq_group.get_seqs():
seq_data[seq.seq_id] = SequenceData(
prompt_token_ids=seq.get_token_ids(),
output_token_ids=seq.get_output_token_ids(),
)
# block_table 是本次 step 需要的映射关系
block_tables = {
seq.seq_id: self.block_manager.get_block_table(seq)
for seq in seq_group.get_seqs()
}
metadata = SequenceGroupMetadata(
request_id=seq_group.request_id,
is_prompt=(seq_group.state == SequenceGroupState.WAITING),
seq_data=seq_data,
sampling_params=seq_group.sampling_params,
block_tables=block_tables,
)
seq_group_metadata_list.append(metadata)
return SchedulerOutputs(
scheduled_seq_groups=seq_group_metadata_list,
ignored_seq_groups=ignored_seq_groups,
num_lookahead_slots=1, # 每次 step 生成 1 个 token
preempted=preempted,
)
def _get_num_required_blocks(self, seq_group) -> int:
"""计算请求需要的 block 数量"""
seq = seq_group.get_seqs()[0]
total_tokens = seq.get_len() # prompt + 已生成 token
return ceil(total_tokens / self.block_size)
def _can_preempt(self) -> bool:
"""是否可以抢占（需要至少一个可抢占的 running 请求）"""
return any(sg.preemptable for sg in self.running)
def _preempt_lowest_priority(self):
"""
抢占优先级最低的请求：
1. 将其 KV cache 从 GPU 拷贝到 CPU
2. 释放 GPU blocks
3. 移入 swapped 队列
"""
victim = min(self.running, key=lambda sg: sg.priority)
for seq in victim.get_seqs():
blocks = self.block_manager.get_block_table(seq)
# GPU → CPU 拷贝
self.block_manager.swap_out(seq)
self.running.remove(victim)
self.swapped.append(victim)
return victim
4.2 ASCII：Continuous Batching 时间线¶
┌─────────────────────────────────────────────────────────────────────────────┐
│ Continuous Batching vs 传统 Static Batching │
├─────────────────────────────────────────────────────────────────────────────┤
│ │
│ 传统 Static Batching: │
│ │
│ Batch 1 │████ 等待 ████│ Batch 2 │████ 等待 ████│ │
│ ┌──────────────────┬─────────────┬──────────────────┬─────────────┐ │
│ │ Req A (12 tok) │████████████████ │ │ │
│ │ Req B (28 tok) │████████████████████████████████│ │ │
│ │ Req C (48 tok) │████████████████████████████████│█████████████│████████│
│ │ Req D ( 6 tok) │████████████████ │ │ │
│ └──────────────────┴─────────────┴──────────────────┴─────────────┘ │
│ ▲ ▲ │
│ A,B,D 早已完成 GPU 空闲等待 C │
│ │
│ │
│ vLLM Continuous Batching: │
│ │
│ Iteration: 1 2 3 4 5 6 7 8 9 10 11 12 13 ... │
│ ┌──────────┬───┬───┬───┬───┬───┬───┬───┬───┬───┬───┬───┬───┬───┐ │
│ │ Req A │ A │ A │ A │ A │ A │ A │ A │ A │ A │ A │ A │ A │✓ │ │
│ │ (12 tok) │ │ │ │ │ │ │ │ │ │ │ │ │ │ │
│ │ Req B │ B │ B │ B │ B │ B │ B │ B │ B │ B │ B │ B │ B │ B │ B ... │
│ │ (28 tok) │ │ │ │ │ │ │ │ │ │ │ │ │ │ │
│ │ Req C │ C │ C │ C │ C │ C │ C │ C │ C │ C │ C │ C │ C │ C │ C ... │
│ │ (48 tok) │ │ │ │ │ │ │ │ │ │ │ │ │ │ │
│ │ Req D │ D │ D │ D │ D │ D │ D │✓ │ │ │ │ │ │ │ │
│ │ ( 6 tok) │ │ │ │ │ │ │ │ │ │ │ │ │ │ │
│ │ │ │ │ │ │ │ │ E │ E │ E │ E │ E │ E │ E │ E ... │
│ │ Req E ───┼───┼───┼───┼───┼───┼───┼───┼───┼───┼───┼───┼───┼───┤ │
│ │ (new) │ │ │ │ │ │ │ │ │ │ │ │ │ │ │
│ │ │ │ │ │ │ │ │ │ │ │ │ F │ F │ F │ F │ F ...│
│ │ Req F ───┼───┼───┼───┼───┼───┼───┼───┼───┼───┼───┼───┼───┼───┤ │
│ │ (new) │ │ │ │ │ │ │ │ │ │ │ │ │ │ │
│ └──────────┴───┴───┴───┴───┴───┴───┴───┴───┴───┴───┴───┴───┴───┘ │
│ ▲ ▲ ▲ ▲ ▲ │
│ Iteration 1 Iteration 7 Iteration 10 加入 F │
│ 4 请求一起 D 完成立即 保持 GPU │
│ 开始 替换为 E 始终满载 │
│ │
│ 关键差异: │
│ - Static: 每批之间有空闲期 (idle gap) │
│ - Continuous: 请求完成即出队，新请求即入队，GPU 始终在处理 │
│ - 吞吐提升: 2-5× (取决于请求长度方差) │
│ │
└─────────────────────────────────────────────────────────────────────────────┘
5. vLLM 整体架构¶
5.1 核心组件¶
# vLLM 引擎架构 (vllm/engine/llm_engine.py)
class LLMEngine:
"""
推理引擎的顶层入口，管理请求的完整生命周期。
"""
def __init__(self, model_config, cache_config, parallel_config, ...):
# 模型执行器 (CPU/GPU worker)
self.model_executor = ModelExecutor(
model_config, parallel_config, cache_config
)
# 调度器: 决定每次 step 哪些请求参与执行
self.scheduler = Scheduler(scheduler_config, cache_config)
# Tokenizer
self.tokenizer = get_tokenizer(model_config.tokenizer)
# 请求跟踪器
self.seq_group_counter = Counter()
# 输出处理器 (detokenize + streaming)
self.output_processor = OutputProcessor(scheduler_config)
# 设备: cuda:0, cuda:1, ... (多 GPU 时有多个)
self.device_config = DeviceConfig(device="cuda")
def add_request(self, request_id, prompt, sampling_params, ...):
"""接收一个新请求"""
# 1. Tokenize 输入 prompt
input_ids = self.tokenizer.encode(prompt)
# 2. 创建 SequenceGroup (一个请求 = 一个 seq_group)
seq_group = SequenceGroup(
request_id=request_id,
seqs=[Sequence(
seq_id=next(self.seq_group_counter),
prompt_token_ids=input_ids,
)],
sampling_params=sampling_params,
arrival_time=time.time(),
)
# 3. 加入 waiting 队列
self.scheduler.add_seq_group(seq_group)
def step(self) -> List[RequestOutput]:
"""
执行一次推理 step:
1. 调度 → 选择本轮运行的请求
2. 执行模型 → forward pass
3. 采样 → 生成下一个 token
4. 更新状态 → 将新 token 追加到序列
"""
# Step 1: 调度
seq_group_metadata_list, scheduler_outputs = self.scheduler.schedule()
if not seq_group_metadata_list:
# 没有可执行的请求 → 返回空
return []
# Step 2: 执行模型
# 将所有 selected seq_groups 的输入拼接成 batch
output = self.model_executor.execute_model(
seq_group_metadata_list=seq_group_metadata_list,
blocks_to_swap_in=scheduler_outputs.blocks_to_swap_in,
blocks_to_swap_out=scheduler_outputs.blocks_to_swap_out,
blocks_to_copy=scheduler_outputs.blocks_to_copy,
)
# Step 3: 采样
# output.logits: [batch_size, vocab_size] — 每个序列的下一个 token logits
sampled_tokens = self._sample(output.logits, seq_group_metadata_list)
# Step 4: 更新序列状态
return self.output_processor.process_outputs(
seq_group_metadata_list, output, sampled_tokens
)
def _sample(self, logits, seq_group_metadata_list):
"""对每个序列采样下一个 token"""
sampled = []
for i, sg_metadata in enumerate(seq_group_metadata_list):
params = sg_metadata.sampling_params
# 应用 temperature
if params.temperature > 0:
logits[i] = logits[i] / params.temperature
# Top-K 过滤
if params.top_k > 0:
top_k_vals, top_k_indices = torch.topk(logits[i], params.top_k)
logits[i] = torch.full_like(logits[i], -float('inf'))
logits[i][top_k_indices] = top_k_vals
# Top-P (nucleus) 过滤
if params.top_p < 1.0:
sorted_logits, sorted_indices = torch.sort(logits[i], descending=True)
cumulative_probs = torch.cumsum(
torch.softmax(sorted_logits, dim=-1), dim=-1
)
sorted_indices_to_remove = cumulative_probs > params.top_p
sorted_indices_to_remove[0] = False # 至少保留一个 token
indices_to_remove = sorted_indices[sorted_indices_to_remove]
logits[i][indices_to_remove] = -float('inf')
# 采样
probs = torch.softmax(logits[i], dim=-1)
next_token = torch.multinomial(probs, num_samples=1)
sampled.append(next_token)
return sampled
5.2 ModelRunner：GPU 执行路径¶
# vLLM vllm/worker/model_runner.py (简化)
class ModelRunner:
"""
负责将 PyTorch 模型的 forward pass 与 PagedAttention 整合。
"""
def __init__(self, model_config, cache_config, ...):
self.model_config = model_config
self.cache_config = cache_config
# 加载模型权重
self.model: nn.Module = self._load_model()
# 分配 KV cache
# [num_blocks, 2, num_kv_heads, block_size, head_dim]
num_blocks = cache_config.num_gpu_blocks
self.kv_cache_dtype = model_config.dtype # float16 或 bfloat16
self.key_cache = torch.empty(
num_blocks,
model_config.get_num_kv_heads(),
cache_config.block_size,
model_config.get_head_size(),
dtype=self.kv_cache_dtype,
device="cuda",
)
self.value_cache = torch.empty_like(self.key_cache)
# 加载 CUDA kernel
self.paged_attn_kernel = _load_paged_attention_kernel()
def _load_model(self):
"""加载 HuggingFace 模型并替换 attention"""
model = AutoModelForCausalLM.from_pretrained(
self.model_config.model,
torch_dtype=self.model_config.dtype,
)
# 替换所有 attention 层为 PagedAttention 版本
for layer in model.model.layers:
layer.self_attn = PagedAttentionWrapper(
layer.self_attn, # 原始 attention 模块
self.key_cache,
self.value_cache,
self.cache_config.block_size,
)
return model
def execute_model(
self,
seq_group_metadata_list: List[SequenceGroupMetadata],
...
) -> ModelRunnerOutput:
"""
核心执行路径：准备输入 → 运行 forward → 返回 logits
"""
# Step 1: 准备输入 tensors
input_tokens = [] # 所有请求的 token IDs
input_positions = [] # 每个 token 在序列中的位置
block_tables = [] # 每个请求的 block table
for sg_meta in seq_group_metadata_list:
seq_data = sg_meta.seq_data[sg_meta.request_id]
# 只输入最后一个 token (decode 阶段)
# 或整个 prompt (prefill 阶段)
if sg_meta.is_prompt:
input_tokens.extend(seq_data.prompt_token_ids)
input_positions.extend(range(len(seq_data.prompt_token_ids)))
else:
input_tokens.append(seq_data.output_token_ids[-1])
input_positions.append(len(seq_data.prompt_token_ids)
+ len(seq_data.output_token_ids) - 1)
# block table: 逻辑 block → 物理 block
block_tables.append(sg_meta.block_tables[sg_meta.request_id])
# 转换为 CUDA tensors
input_ids = torch.tensor(input_tokens, dtype=torch.long, device="cuda")
positions = torch.tensor(input_positions, dtype=torch.long, device="cuda")
# 将不同长度的 block_tables padding 到相同长度
max_blocks = max(len(bt) for bt in block_tables)
padded_block_tables = []
for bt in block_tables:
padded = bt + [-1] * (max_blocks - len(bt)) # -1 = 无效 block
padded_block_tables.append(padded)
block_table_tensor = torch.tensor(
padded_block_tables, dtype=torch.int32, device="cuda"
)
# Step 2: 运行模型 forward
# 模型内部通过 PagedAttentionWrapper 调用 custom kernel
with torch.no_grad():
hidden_states = self.model(
input_ids=input_ids,
positions=positions,
kv_caches=(self.key_cache, self.value_cache),
attn_metadata=AttentionMetadata(
block_tables=block_table_tensor,
context_lens=context_lens_tensor,
max_context_len=max_context_len,
is_prompt=any(sg.is_prompt for sg in seq_group_metadata_list),
),
)
# Step 3: 提取 logits (只需最后一个位置)
# 每个序列的 logits 对应于该序列的最后一个 token
logits = self.model.lm_head(hidden_states[-batch_size:])
return ModelRunnerOutput(logits=logits)
6. 性能数据¶
6.1 vLLM vs HuggingFace Transformers¶
以 Llama-2-7B 在 A100-80GB 上的实测数据：
┌───────────────────────┬───────────────┬──────────────┬─────────┐
│ Metric │ HF Transformers│ vLLM │ Speedup │
├───────────────────────┼───────────────┼──────────────┼─────────┤
│ Throughput (tok/s) │ 1,200 │ 28,800 │ 24× │
│ KV cache utilization │ 23% │ 96.4% │ 4.2× │
│ Max batch size │ 64 │ ~256 │ 4× │
│ Avg latency/token │ 42 ms │ 14 ms │ 3× │
│ Memory overhead │ 4-6× │ 1.05× │ — │
│ Prefill latency │ 850 ms │ 350 ms │ 2.4× │
│ (2048 tokens) │ │ │ │
└───────────────────────┴───────────────┴──────────────┴─────────┘
6.2 KV Cache 利用率对比¶
HuggingFace Transformers:
总分配: 10 GB (所有请求预分配 max_length)
实际使用: 2.3 GB
利用率: 23%
vLLM (PagedAttention):
总分配: 9.8 GB (block 粒度)
实际使用: 9.45 GB (碎片 < 4%)
利用率: 96.4%
节省的 GPU 内存可以:
- 容纳更多并发请求 (提高 batch size → 更高吞吐)
- 使用更大的模型 (在相同硬件上跑更大的模型)
- 处理更长的序列 (长上下文推理)
6.3 为什么 Continuous Batching 能消除 idle time¶
# 模拟实验代码
import random
import time
from collections import deque
def simulate_static_batching(requests, max_batch_size=32):
"""传统批处理模拟"""
total_time = 0
i = 0
while i < len(requests):
batch = requests[i:i + max_batch_size]
i += len(batch)
# 等待 batch 中最长的请求完成
max_tokens = max(r['num_tokens'] for r in batch)
batch_time = sum(r['num_tokens'] for r in batch) # 总计算
total_time += max_tokens * 0.01 # 每个 token 10ms
return total_time
def simulate_continuous_batching(requests, block_size=16, gpu_mem_blocks=4096):
"""Continuous Batching 模拟"""
waiting = deque(sorted(requests, key=lambda r: r['arrival']))
running = deque()
free_blocks = gpu_mem_blocks
total_time = 0
total_tokens_generated = 0
while waiting or running:
# 检查完成的请求
finished = [r for r in running if r['generated'] >= r['num_tokens']]
for r in finished:
running.remove(r)
blocks_used = (len(r['prompt']) + r['generated']) // block_size + 1
free_blocks += blocks_used
# 尝试加入新请求
while waiting:
r = waiting[0]
blocks_needed = len(r['prompt']) // block_size + 1
if blocks_needed <= free_blocks:
waiting.popleft()
free_blocks -= blocks_needed
running.append(r)
else:
break
# 执行一步（所有 running 请求生成 1 个 token）
for r in running:
r['generated'] += 1
total_tokens_generated += 1
# 如果序列增长了需要更多 blocks，动态分配
current_blocks = (len(r['prompt']) + r['generated']) // block_size + 1
if current_blocks > r.get('allocated_blocks', 0):
free_blocks -= 1
r['allocated_blocks'] = current_blocks
total_time += 0.01 # 每个 token 10ms
return total_time, total_tokens_generated
# 结果示例:
# Static batching: 180.3s idle time, 62% GPU utilization
# Continuous: 0.8s idle time, 99.2% GPU utilization
7. API 使用示例¶
7.1 离线批量推理¶
from vllm import LLM, SamplingParams
# 初始化引擎
llm = LLM(
model="meta-llama/Llama-2-7b-hf",
tensor_parallel_size=1, # 单 GPU
dtype="float16",
gpu_memory_utilization=0.90, # 使用 90% GPU 内存
max_num_batched_tokens=4096, # 最大 batch token 数
block_size=16, # PagedAttention block 大小
)
# 批量推理
prompts = [
"What is the capital of France?",
"Explain quantum computing in simple terms.",
"Write a haiku about GPU programming.",
"Translate 'hello world' to Japanese.",
# ... 可以有数百个 prompts
]
sampling_params = SamplingParams(
temperature=0.8,
top_p=0.95,
top_k=50,
max_tokens=256,
)
outputs = llm.generate(prompts, sampling_params)
for prompt, output in zip(prompts, outputs):
generated_text = output.outputs[0].text
print(f"Prompt: {prompt}")
print(f"Response: {generated_text}")
print(f"Tokens: {len(output.outputs[0].token_ids)}")
print("-" * 50)
# vLLM 内部自动处理:
# - PagedAttention memory management
# - Continuous batching scheduling
# - KV cache allocation/deallocation
# - Token streaming (if configured)
7.2 在线服务（OpenAI 兼容 API）¶
# 启动 vLLM 服务器
# $ python -m vllm.entrypoints.openai.api_server \
# --model meta-llama/Llama-2-7b-hf \
# --tensor-parallel-size 1 \
# --max-model-len 4096
# 客户端调用 (OpenAI 兼容)
from openai import OpenAI
client = OpenAI(
base_url="http://localhost:8000/v1",
api_key="not-needed",
)
response = client.completions.create(
model="meta-llama/Llama-2-7b-hf",
prompt="Explain the PagedAttention algorithm:",
max_tokens=512,
temperature=0.7,
stream=True,
)
for chunk in response:
if chunk.choices[0].text:
print(chunk.choices[0].text, end="", flush=True)
7.3 内核参数调优¶
from vllm import EngineArgs, LLMEngine
# 创建自定义 Engine
engine_args = EngineArgs(
model="meta-llama/Llama-2-7b-hf",
dtype="float16",
# ===== PagedAttention 相关 =====
block_size=16, # 每个 block 存储的 token 数
# 16: 适合大多数场景
# 8: 更细粒度，减少浪费但增加 block table 开销
# 32: 更粗粒度，减少开销但可能浪费内存
# ===== 内存相关 =====
gpu_memory_utilization=0.90, # GPU 内存使用上限
swap_space=4, # CPU swap 空间 (GB) — 当 GPU 内存不足时
# 设置 0 禁用 swap（不允许 OOM 时抢占 swap out）
# ===== 调度相关 =====
max_num_seqs=256, # 最大并发序列数
max_num_batched_tokens=4096, # 每个 batch 的最大 token 数
# 控制 prefill 阶段的内存峰值
# ===== 抢占配置 =====
preemption_mode="swap", # "swap": GPU→CPU swap
# "recompute": 丢弃后重新计算
# ===== 性能相关 =====
enforce_eager=False, # False: 使用 CUDA graph 加速
# True: 每次 forward 都是 eager mode (调试用)
disable_custom_all_reduce=True, # 多 GPU 时的通信
# ===== KV Cache 数据类型 =====
kv_cache_dtype="auto", # "auto": 自动选择 (fp16/bf16)
# "fp8": 使用 FP8 KV cache (Hopper 架构)
# 牺牲微小精度，节省 50% KV cache 内存
)
engine = LLMEngine.from_engine_args(engine_args)
8. 与其他推理引擎的对比¶
┌─────────────────┬──────────┬─────────────┬──────────────────┐
│ 特性 │ vLLM │ TGI (HF) │ TensorRT-LLM │
├─────────────────┼──────────┼─────────────┼──────────────────┤
│ PagedAttention │ ✅ │ ✅* │ ❌ (不同方案) │
│ Continuous Batch│ ✅ │ ✅ │ ✅ (inflight) │
│ FP8 KV Cache │ ✅ │ ❌ │ ✅ │
│ Tensor Parallel │ ✅ │ ✅ │ ✅ │
│ Pipeline Parallel│ ✅ │ ✅ │ ✅ │
│ CUDA Graph │ ✅ │ ✅ │ ✅ │
│ Prefix Caching │ ✅ │ ✅ │ ❌ │
│ Speculative Dec │ ✅* │ ✅ │ ❌ │
│ 量化 (GPTQ/AWQ) │ ✅ │ ✅ │ ✅ │
│ OpenAI API │ ✅ │ ✅ │ ✅ (Triton) │
│ LoRA Adapters │ ✅ │ ✅ │ ✅ │
│ 部署难度 │ pip安装 │ Docker │ 需要编译 │
│ 吞吐 (Llama-7B) │ 高 │ 中 │ 最高 (优化最激 ┃
│ │ │ │ 进) │
└─────────────────┴──────────┴─────────────┴──────────────────┘
* TGI 的 PagedAttention 实现与 vLLM 略有不同，使用 FlashAttention
* vLLM 的 speculative decoding 处于实验阶段
9. 常见问题与调试¶
9.1 CUDA OOM 的处理¶
# 场景: 长序列推理导致 OOM
# 策略 1: 调整 gpu_memory_utilization
llm = LLM(
model="...",
gpu_memory_utilization=0.75, # 从 0.9 降到 0.75
)
# 策略 2: 启用 CPU swap
llm = LLM(
model="...",
swap_space=8, # 分配 8 GB CPU 内存做 swap
preemption_mode="swap",
)
# 策略 3: FP8 KV cache (支持 H100 时)
llm = LLM(
model="...",
kv_cache_dtype="fp8",
)
# FP8 下 KV cache 内存减半，但精度损失 < 0.1%
# 策略 4: 限制 max_model_len
llm = LLM(
model="...",
max_model_len=4096, # 降低最大序列长度
)
# 调试: 查看 KV cache 使用情况
from vllm.worker.model_runner import ModelRunner
# 在日志中搜索: "avg_num_blocks_per_req"
9.2 Block 分配调试¶
# 查看 PagedAttention 统计信息
@dataclass
class SchedulerStats:
num_running_reqs: int
num_waiting_reqs: int
num_swapped_reqs: int
gpu_cache_usage: float # 0.0 ~ 1.0
cpu_cache_usage: float # 0.0 ~ 1.0
avg_blocks_per_req: float
avg_prefix_cache_hit_rate: float # block 共享命中率
def get_scheduler_stats(engine: LLMEngine) -> SchedulerStats:
scheduler = engine.scheduler
block_manager = scheduler.block_manager
total_gpu_blocks = block_manager.num_total_gpu_blocks
free_gpu_blocks = block_manager.num_free_gpu_blocks
total_cpu_blocks = block_manager.num_total_cpu_blocks
free_cpu_blocks = block_manager.num_free_cpu_blocks
running = scheduler.running
avg_blocks = 0
if running:
avg_blocks = sum(
len(block_manager.get_block_table(seq))
for sg in running
for seq in sg.get_seqs()
) / len(running)
return SchedulerStats(
num_running_reqs=len(scheduler.running),
num_waiting_reqs=len(scheduler.waiting),
num_swapped_reqs=len(scheduler.swapped),
gpu_cache_usage=1.0 - free_gpu_blocks / total_gpu_blocks,
cpu_cache_usage=1.0 - free_cpu_blocks / total_cpu_blocks,
avg_blocks_per_req=avg_blocks,
avg_prefix_cache_hit_rate=0.0, # 需要额外统计
)
9.3 性能剖析¶
# 使用 PyTorch profiler 分析 vLLM
python -c "
import torch
from vllm import LLM, SamplingParams
llm = LLM(model='meta-llama/Llama-2-7b-hf')
with torch.profiler.profile(
activities=[
torch.profiler.ProfilerActivity.CPU,
torch.profiler.ProfilerActivity.CUDA,
],
with_stack=True,
record_shapes=True,
) as prof:
outputs = llm.generate(['Hello world'] * 32, SamplingParams(max_tokens=256))
print(prof.key_averages().table(
sort_by='cuda_time_total',
row_limit=20
))
"
# 预期输出关键算子:
# paged_attention_v2_kernel ← 自定义 CUDA kernel: 30-40% CUDA time
# cublasGemmEx ← cuBLAS 矩阵乘法: 25-35%
# rms_norm / layer_norm ← 归一化层: 5-10%
# rotary_embedding ← RoPE 位置编码: 2-5%
# silu_and_mul / gelu ← 激活函数: 5-10%
10. 总结¶
vLLM 将操作系统虚拟内存的思想引入 LLM 推理的 KV cache 管理：
- PagedAttention：将 KV cache 分块，通过 Block Table 实现逻辑→物理映射，消除内存碎片，支持 block 共享和 COW
- Continuous Batching：请求完成即出队、新请求即入队，GPU 始终满负荷——消除传统批处理的 idle gap
- Online Softmax：分块处理 attention，避免存储 O(n²) 的完整 attention 矩阵
- Custom CUDA Kernel：直接操作 PagedAttention 的分块 KV cache 布局，绕过 PyTorch 的连续 tensor 限制
性能收益：相对于 HuggingFace Transformers，吞吐提升 14×–24×，KV cache 利用率从 ~23% 提升到 ~96%。
vLLM 展示了一个重要范式：在高吞吐场景下，通过 改变内存管理策略（而非单纯优化计算）可以带来数量级的性能提升。下一篇文章中，我们将深入 GPU 硬件架构底层，理解这些 kernel 在 SM、Warp、Tensor Core 上究竟如何执行。
下一篇文章¶
第五篇：GPU 硬件架构：SM、Warp、Tensor Core、TMA