BlockSpaceManagerV1 概览 — KV Cache 块管理的核心职责
在 vLLM 的推理流水线中,KV Cache(Key-Value Cache)是 GPU 显存的主要消费者。
每个 token 在每一层 Transformer 都会产生一对 Key/Value 张量,这些张量需要被持久化以避免重复计算。
BlockSpaceManagerV1 正是负责管理这块"KV Cache 显存池"的核心组件。
它采用分页内存管理(Paged Memory Management)思想,将连续的 KV Cache 显存切割成固定大小的"物理块"(Physical Block), 再通过逻辑块到物理块的映射表(Block Table)为每个序列虚拟出连续的地址空间。 这一设计直接来源于操作系统的分页虚拟内存机制。
- 物理块分配/释放:通过 GPU/CPU 两个
BlockAllocator管理显存和内存中的物理块 - Block Table 维护:每个
Sequence拥有独立的逻辑→物理块映射表 - Copy-on-Write(CoW):Beam Search / 并行采样时,子序列共享父序列物理块,写入时触发复制
- GPU ↔ CPU 换入换出:调度器决定 swap 时,搬移整个序列的 KV Cache
- Prefix Caching:通过
CachedBlockAllocator对相同前缀的块进行哈希匹配复用
BlockSpaceManagerV1 继承自抽象基类 BlockSpaceManager(定义于 core/interfaces.py),
实现了调度器所需的完整显存管理接口:
class AllocStatus(enum.Enum):
"""Result for BlockSpaceManager.can_allocate
1. Ok: seq_group can be allocated now.
2. Later: seq_group cannot be allocated.
The capacity of allocator is larger than seq_group required.
3. Never: seq_group can never be allocated.
The seq_group is too large to allocated in GPU.
"""
OK = enum.auto()
LATER = enum.auto()
NEVER = enum.auto()
class BlockSpaceManager(ABC):
@abstractmethod
def can_allocate(self, seq_group: SequenceGroup) -> AllocStatus: ...
@abstractmethod
def allocate(self, seq_group: SequenceGroup) -> None: ...
@abstractmethod
def can_append_slots(self, seq_group: SequenceGroup,
num_lookahead_slots: int) -> bool: ...
@abstractmethod
def append_slots(self, seq: Sequence,
num_lookahead_slots: int) -> Dict[int, List[int]]: ...
@abstractmethod
def fork(self, parent_seq: Sequence, child_seq: Sequence) -> None: ...
@abstractmethod
def swap_in(self, seq_group: SequenceGroup,
num_lookahead_slots: int) -> Dict[int, int]: ...
@abstractmethod
def swap_out(self, seq_group: SequenceGroup) -> Dict[int, int]: ...
@abstractmethod
def free(self, seq: Sequence) -> None: ...
AllocStatus 的三态返回值设计非常精巧:OK 表示立即可以分配,LATER 表示当前显存不足但等待后可能成功(调度器会让请求继续等待),NEVER 表示请求所需 KV Cache 超过总显存容量,永远无法满足(直接拒绝)。
PhysicalTokenBlock / LogicalTokenBlock — 物理块与逻辑块数据结构
块管理的基础是两种数据结构:逻辑块(LogicalTokenBlock)代表序列视角的连续 token 存储,
物理块(PhysicalTokenBlock)代表 GPU/CPU 内存中实际分配的 KV Cache 存储单元。
两者均定义在 vllm/block.py 中。
LogicalTokenBlock — 逻辑视角
class LogicalTokenBlock:
"""A block that stores a contiguous chunk of tokens from left to right.
Logical blocks are used to represent the states of the corresponding
physical blocks in the KV cache.
"""
def __init__(
self,
block_number: int,
block_size: int,
) -> None:
self.block_number = block_number
self.block_size = block_size
self.token_ids = [_BLANK_TOKEN_ID] * block_size
self.num_tokens = 0
def is_empty(self) -> bool:
return self.num_tokens == 0
def get_num_empty_slots(self) -> int:
return self.block_size - self.num_tokens
def is_full(self) -> bool:
return self.num_tokens == self.block_size
def append_tokens(self, token_ids: List[int]) -> None:
assert len(token_ids) <= self.get_num_empty_slots()
curr_idx = self.num_tokens
self.token_ids[curr_idx:curr_idx + len(token_ids)] = token_ids
self.num_tokens += len(token_ids)
LogicalTokenBlock 存储实际的 token ID 列表,维护已填充 token 数量(num_tokens)。
它本身不涉及任何 GPU 内存,只是序列逻辑状态的抽象。
token_ids 初始化为 _BLANK_TOKEN_ID = -1,通过 append_tokens() 逐步填充。
PhysicalTokenBlock — 物理视角
class PhysicalTokenBlock:
"""Represents the state of a block in the KV cache."""
def __init__(
self,
device: Device,
block_number: int,
block_size: int,
block_hash: int,
num_hashed_tokens: int,
) -> None:
self.device = device # GPU 或 CPU
self.block_number = block_number # 在显存池中的槽位编号
self.block_size = block_size # 每块容纳的 token 数
self.block_hash = block_hash # 内容哈希(Prefix Caching 用)
self.num_hashed_tokens = num_hashed_tokens # 参与哈希的前缀长度
self.ref_count = 0 # 引用计数(CoW 核心)
self.last_accessed = DEFAULT_LAST_ACCESSED_TIME # LRU 时间戳
self.computed = False # 是否已计算过 KV(Prefix Cache 命中标记)
# Mapping: logical block number -> physical block.
BlockTable = List[PhysicalTokenBlock]
- ref_count:引用计数。当多个序列(如 Beam Search 的子序列)共享同一物理块时,
ref_count > 1。只有ref_count == 0的块才能被释放或淘汰。 - block_hash:对该块中所有 token ID 计算的哈希值。
CachedBlockAllocator用此实现 Prefix KV Cache 复用。 - num_hashed_tokens:参与哈希的 token 总数(包含前缀),用于 LRU 淘汰时优先保留更长的前缀块(哈希 token 越多,复用价值越高)。
- computed:标记该块是否已经做过 forward pass 计算。Prefix Caching 命中时直接跳过,无需重算。
- last_accessed:上次访问时间戳,驱动 LRU 淘汰算法。
BlockTable = List[PhysicalTokenBlock] 是每个序列的块映射表——索引即逻辑块号,值为对应的物理块对象。
调度器通过 get_block_table(seq) 获取物理块号列表,传给 CUDA kernel 进行 attention 计算。
BlockAllocator — GPU/CPU 块分配器实现
BlockSpaceManagerV1 持有两个分配器实例——gpu_allocator 和 cpu_allocator,
分别管理 GPU 显存和 CPU 内存中的物理块池。
vLLM V1 提供两种分配器实现,通过 enable_caching 参数选择:
UncachedBlockAllocator — 简单空闲列表
class UncachedBlockAllocator(BlockAllocatorBase):
def __init__(
self,
device: Device,
block_size: int,
num_blocks: int,
) -> None:
self.device = device
self.block_size = block_size
self.num_blocks = num_blocks
# Initialize the free blocks.
self.free_blocks: BlockTable = []
for i in range(num_blocks):
block = PhysicalTokenBlock(device=device,
block_number=i,
block_size=block_size,
block_hash=-1,
num_hashed_tokens=0)
self.free_blocks.append(block)
def allocate(self,
block_hash: Optional[int] = None,
num_hashed_tokens: int = 0) -> PhysicalTokenBlock:
if not self.free_blocks:
raise ValueError("Out of memory! No free blocks are available.")
block = self.free_blocks.pop()
block.ref_count = 1
return block
def free(self, block: PhysicalTokenBlock) -> None:
if block.ref_count == 0:
raise ValueError(f"Double free! {block} is already freed.")
block.ref_count -= 1
if block.ref_count == 0:
self.free_blocks.append(block)
最简实现:初始化时预分配所有物理块对象并放入 free_blocks 列表,
allocate() 从列表尾部弹出,free() 递减引用计数后归还。
不支持 Prefix Caching,block_hash 参数被忽略。
CachedBlockAllocator — 支持 Prefix KV Cache
class CachedBlockAllocator(BlockAllocatorBase):
def __init__(self, device, block_size, num_blocks,
eviction_policy=EvictionPolicy.LRU) -> None:
self.device = device
self.block_size = block_size
self.num_blocks = num_blocks
self.current_num_blocks = 0
self.cached_blocks: Dict[int, PhysicalTokenBlock] = {}
self.evictor: Evictor = make_evictor(eviction_policy)
self.default_hash_ctr = count()
def allocate_block(self, block_hash, num_hashed_tokens):
if self.current_num_blocks == self.num_blocks:
# 显存已满:从 evictor 淘汰一块
block = self.evictor.evict()
block.block_hash = block_hash
block.num_hashed_tokens = num_hashed_tokens
return block
# 显存有余:分配新的物理块槽位
block = PhysicalTokenBlock(device=self.device,
block_number=self.current_num_blocks,
block_size=self.block_size,
block_hash=block_hash,
num_hashed_tokens=num_hashed_tokens)
self.current_num_blocks += 1
return block
def allocate(self, block_hash=None, num_hashed_tokens=0):
if block_hash is None:
block_hash = next(self.default_hash_ctr)
# 命中 evictor(已释放但未驱逐的缓存块)
if block_hash in self.evictor:
block = self.evictor.remove(block_hash)
assert block.ref_count == 0
self.cached_blocks[block_hash] = block
block.ref_count += 1
return block
# 命中 cached_blocks(正在使用的块)→ 共享引用
if block_hash not in self.cached_blocks:
self.cached_blocks[block_hash] = self.allocate_block(
block_hash, num_hashed_tokens)
block = self.cached_blocks[block_hash]
block.ref_count += 1
return block
def free(self, block: PhysicalTokenBlock) -> None:
if block.ref_count == 0:
raise ValueError(f"Double free! {block} is already freed.")
block.ref_count -= 1
if block.ref_count == 0:
# 引用归零:不立即释放,交给 evictor 管理(延迟淘汰)
self.evictor.add(block)
del self.cached_blocks[block.block_hash]
allocate(block_hash) 按优先级顺序检查三个位置:
- evictor 中:该哈希对应的块已被释放(ref_count=0)但还未被驱逐——直接"复活",避免重算 KV。
- cached_blocks 中:该哈希对应的块正被其他序列使用——增加 ref_count 共享,这正是 Prefix Caching 的核心。
- 全新分配:调用
allocate_block(),若显存满则先淘汰一块再使用。
注意 free() 在引用归零时不会立即销毁物理块,而是调用 evictor.add(block) 将其转入"缓存待淘汰"状态。
这样下次遇到相同前缀的请求时,可以从 evictor 中直接复活该块(命中第一级),实现零拷贝的 KV Cache 复用。
allocate() — 为新序列分配块的完整流程
当调度器决定将一个新请求(SequenceGroup)从 Waiting 队列调度执行时,
会先调用 can_allocate() 检查可用物理块是否充足,再调用 allocate() 完成实际分配。
can_allocate() — 预检
def can_allocate(self, seq_group: SequenceGroup) -> AllocStatus:
# 取 waiting 列表中第一个(代表性)序列
seq = seq_group.get_seqs(status=SequenceStatus.WAITING)[0]
# 计算该序列需要多少个物理块
num_required_blocks = math.ceil(
seq.get_len() / self.block_size
)
# 加上 lookahead 块(投机解码预留)
num_required_blocks += self.block_sliding_window or 0
if self.block_sliding_window is not None:
num_required_blocks = min(num_required_blocks,
self.block_sliding_window)
num_free_gpu_blocks = self.gpu_allocator.get_num_free_blocks()
# 若可用块数低于水位线,拒绝调度
if (self.num_total_gpu_blocks - num_required_blocks
< self.watermark_blocks):
return AllocStatus.NEVER
if num_free_gpu_blocks - num_required_blocks >= self.watermark_blocks:
return AllocStatus.OK
else:
return AllocStatus.LATER
watermark_blocks 是一个安全余量(默认 watermark=0.01,即 1% 的 GPU 块数)。
保留这部分块作为缓冲,防止 allocate 与 append 之间出现竞态导致 OOM。
只有可用块数减去所需块数仍高于水位线,才返回 OK;否则返回 LATER(等待)。
若总容量本身不够,则返回 NEVER(直接拒绝该请求)。
allocate() — 执行分配
def allocate(self, seq_group: SequenceGroup) -> None:
# 取代表性序列(所有序列共享相同 prompt,块布局相同)
seq = seq_group.get_seqs(status=SequenceStatus.WAITING)[0]
# 为该序列的每个逻辑块分配一个物理块
block_table: BlockTable = []
num_prompt_blocks = math.ceil(seq.get_len() / self.block_size)
for logical_idx in range(num_prompt_blocks):
# 计算该块的哈希(用于 Prefix Caching)
block_hash = ... # 基于 token ids 内容
block = self.gpu_allocator.allocate(
block_hash, num_hashed_tokens=...)
block_table.append(block)
# 将 block_table 注册给序列组中的每个序列
for seq in seq_group.get_seqs(status=SequenceStatus.WAITING):
self.block_tables[seq.seq_id] = block_table.copy()
# 对于并行采样(num_seqs > 1),子序列共享同一批物理块
# 通过增加 ref_count 实现共享(CoW 的准备)
for block in block_table:
block.ref_count += len(seq_group.get_seqs()) - 1
当一个请求使用 n > 1 的并行采样(如 top-p 采样出多个候选序列)时,
所有子序列的 prompt 部分完全相同,因此共享同一份物理块——ref_count 被增加到 n。
Decode 阶段各子序列生成不同的 token,需要各自独立的物理块,此时通过 CoW 机制触发写时复制。
can_append_slots() / append_slots() — 追加 token 的块管理
在 Decode 阶段,每次迭代为序列生成一个新 token,需要追加到 KV Cache 中。
这由 can_append_slots() 预检 + append_slots() 执行来完成。
can_append_slots() — 检查空间
def can_append_slots(self,
seq_group: SequenceGroup,
num_lookahead_slots: int = 0) -> bool:
"""Simple heuristic: a sequence can append if there are enough
free blocks for all sequences in the group that need a new block.
"""
num_free_gpu_blocks = self.gpu_allocator.get_num_free_blocks()
# 统计当前 running 中需要新块的序列数
num_seqs = seq_group.num_seqs(status=SequenceStatus.RUNNING)
# 每个序列最多需要 1 个新块(当前最后一块已满时)
# 加上 lookahead 预留(投机解码)
num_required_blocks = num_seqs * (1 + num_lookahead_slots)
return num_free_gpu_blocks >= num_required_blocks
append_slots() — 追加执行
def append_slots(
self,
seq: Sequence,
num_lookahead_slots: int = 0,
) -> Dict[int, List[int]]:
"""Allocate a physical slot for the new token(s).
Returns a dict of {src_block: [dst_block]} for Copy-on-Write.
"""
block_table = self.block_tables[seq.seq_id]
# 可能需要升级(promote)最后一块:
# 若最后一块已满且是只读的前缀缓存块,需要 CoW
block_table, maybe_promote_block = \
self._maybe_promote_last_block(seq, block_table)
# 检查最后一块是否已满
if self._is_last_block_full(seq):
# 分配新的物理块
new_block = self._allocate_last_physical_block(seq)
block_table.append(new_block)
# lookahead 预留:额外分配 num_lookahead_slots 个块
for _ in range(num_lookahead_slots):
new_block = self.gpu_allocator.allocate()
block_table.append(new_block)
# 返回 CoW 映射表(供 CUDA kernel 执行内存复制)
return maybe_promote_block
_is_last_block_full() 与 _allocate_last_physical_block()
def _is_last_block_full(self, seq: Sequence) -> bool:
token_ids_len = seq.data.get_len()
return token_ids_len > 0 and token_ids_len % self.block_size == 0
def _allocate_last_physical_block(self, seq: Sequence) -> PhysicalTokenBlock:
# 计算新块对应的 token 哈希(Prefix Caching)
block_table = self.block_tables[seq.seq_id]
num_hashed_tokens = self.block_size * len(block_table)
# 计算内容哈希(含前缀块哈希,形成链式哈希)
new_block = self.gpu_allocator.allocate(
block_hash=hash of token slice,
num_hashed_tokens=num_hashed_tokens
)
return new_block
- 最后一块未满 & ref_count=1:直接在该块的空槽写入新 token KV,无需分配新块。
- 最后一块未满 & ref_count>1(共享块):触发 CoW,复制一个新块,在新块上写入。
- 最后一块已满:调用
_allocate_last_physical_block()分配新物理块追加到 block_table。
Copy-on-Write 机制 — 共享块的写时复制
Copy-on-Write(CoW)是 vLLM 支持 Beam Search 和并行采样的关键机制。
当多个序列共享同一物理块(ref_count > 1)且需要写入新内容时,
必须先复制一份私有副本再写入,否则会污染其他序列的 KV Cache。
触发时机:_promote_last_block()
def _promote_last_block(
self,
seq: Sequence,
last_block: PhysicalTokenBlock,
) -> Tuple[PhysicalTokenBlock, Dict[int, List[int]]]:
"""If the last block is shared, copy it to a new block.
Returns:
- new_block: 专属的新物理块
- cow_map: {旧块号: [新块号]},传给 CUDA kernel 执行内存复制
"""
if last_block.ref_count == 1:
# 独占块,无需 CoW
return last_block, {}
# 触发 CoW:分配新块
new_block = self.gpu_allocator.allocate()
# 释放对旧块的引用(ref_count-1)
self.gpu_allocator.free(last_block)
# 返回 CoW 映射:调度器将此映射传给 Worker 执行 GPU 内存复制
return new_block, {last_block.block_number: [new_block.block_number]}
def _maybe_promote_last_block(
self,
seq: Sequence,
block_table: BlockTable,
) -> Tuple[BlockTable, Dict[int, List[int]]]:
if self._is_last_block_full(seq):
# 最后一块已满,即将分配新块,不需要 promote
return block_table, {}
last_block = block_table[-1]
new_last_block, cow_map = self._promote_last_block(seq, last_block)
if cow_map:
block_table[-1] = new_last_block
return block_table, cow_map
fork() — Beam Search 子序列创建
def fork(self, parent_seq: Sequence, child_seq: Sequence) -> None:
"""Fork a sequence by sharing all its physical blocks.
Used for Beam Search: child sequence shares parent's blocks.
"""
# 子序列的 block_table 是父序列的浅拷贝(共享物理块对象)
src_block_table = self.block_tables[parent_seq.seq_id]
self.block_tables[child_seq.seq_id] = src_block_table.copy()
# 每个物理块的引用计数加 1
for block in src_block_table:
block.ref_count += 1
以 Beam Search 为例,完整的 CoW 流程如下:
- Prefill 完成,prompt 的物理块由父序列独占(
ref_count=1) - 调用
fork(parent, child),子序列与父序列共享所有物理块,每块ref_count=2 - Decode 步骤,两个序列各自调用
append_slots() - 若当前最后一块未满且
ref_count=2,_promote_last_block()触发 CoW:- 分配新物理块
new_block - 旧块
ref_count -= 1(变为 1,另一序列仍持有) - 返回
{old_block_number: [new_block_number]}CoW 映射
- 分配新物理块
- 调度器将 CoW 映射表传给 GPU Worker,Worker 执行
cudaMemcpy复制旧块内容到新块 - 此后两个序列各自持有独立的最后一块,可以独立写入
swap_in() / swap_out() — GPU↔CPU 块交换
当 GPU 显存不足时,调度器会选择将部分序列的 KV Cache 从 GPU 换出(swap out)到 CPU 内存, 腾出空间给更高优先级的序列。之后在合适的时机再换入(swap in)继续执行。
swap_out() — GPU → CPU
def swap_out(self, seq_group: SequenceGroup) -> Dict[int, int]:
"""Swap out all sequences in the group from GPU to CPU.
Returns:
mapping: {gpu_block_number: cpu_block_number}
"""
mapping: Dict[PhysicalTokenBlock, PhysicalTokenBlock] = {}
for seq in seq_group.get_seqs(status=SequenceStatus.RUNNING):
new_block_table: BlockTable = []
block_table = self.block_tables[seq.seq_id]
for gpu_block in block_table:
if gpu_block in mapping:
# 该 GPU 块已被映射(共享块,只换出一次)
cpu_block = mapping[gpu_block]
cpu_block.ref_count += 1
else:
# 在 CPU 上分配对应的块
cpu_block = self.cpu_allocator.allocate(
gpu_block.block_hash, gpu_block.num_hashed_tokens)
mapping[gpu_block] = cpu_block
new_block_table.append(cpu_block)
# 释放 GPU 块
self.gpu_allocator.free(gpu_block)
self.block_tables[seq.seq_id] = new_block_table
# 返回 {gpu_block_number: cpu_block_number} 供 Worker 执行 cudaMemcpy
return {gpu_block.block_number: cpu_block.block_number
for gpu_block, cpu_block in mapping.items()}
swap_in() — CPU → GPU
def swap_in(self, seq_group: SequenceGroup,
num_lookahead_slots: int = 0) -> Dict[int, int]:
"""Swap in all sequences in the group from CPU to GPU.
Returns:
mapping: {cpu_block_number: gpu_block_number}
"""
mapping: Dict[PhysicalTokenBlock, PhysicalTokenBlock] = {}
for seq in seq_group.get_seqs(status=SequenceStatus.SWAPPED):
new_block_table: BlockTable = []
block_table = self.block_tables[seq.seq_id]
for cpu_block in block_table:
if cpu_block in mapping:
gpu_block = mapping[cpu_block]
gpu_block.ref_count += 1
else:
gpu_block = self.gpu_allocator.allocate(
cpu_block.block_hash, cpu_block.num_hashed_tokens)
mapping[cpu_block] = gpu_block
new_block_table.append(gpu_block)
self.cpu_allocator.free(cpu_block)
self.block_tables[seq.seq_id] = new_block_table
return {cpu_block.block_number: gpu_block.block_number
for cpu_block, gpu_block in mapping.items()}
两个方法结构完全对称:遍历序列组中的每个序列 → 遍历每个物理块 → 在目标设备分配对应块 → 释放源设备块 → 更新 block_table。
返回值是一个块编号映射字典,传给 GPU Worker 后执行批量 cudaMemcpy(CPU↔GPU 之间的物理内存复制)。
注意共享块的处理:多个序列共享的块只会被换出/换入一次(通过 mapping 字典去重),
其他序列直接共享同一目标块并递增 ref_count,避免重复拷贝。
can_swap_in() / can_swap_out() — 预检
def can_swap_in(self, seq_group: SequenceGroup,
num_lookahead_slots: int = 0) -> AllocStatus:
# 计算换入所需的 GPU 块数
blocks = self._get_physical_blocks(seq_group)
num_swapped_seqs = seq_group.num_seqs(status=SequenceStatus.SWAPPED)
num_free_blocks = self.gpu_allocator.get_num_free_blocks()
# 加上 lookahead 预留
num_required = len(blocks) + num_swapped_seqs * num_lookahead_slots
if self.num_total_gpu_blocks - num_required < self.watermark_blocks:
return AllocStatus.NEVER
if num_free_blocks - num_required >= self.watermark_blocks:
return AllocStatus.OK
return AllocStatus.LATER
def can_swap_out(self, seq_group: SequenceGroup) -> bool:
# 检查 CPU 是否有足够空闲块
blocks = self._get_physical_blocks(seq_group)
return len(blocks) <= self.cpu_allocator.get_num_free_blocks()
LRU Evictor — 块淘汰策略
当 CachedBlockAllocator 的显存已满(current_num_blocks == num_blocks),
需要分配新块时,必须从"已释放但未驱逐"的块中淘汰一个以腾出空间。
LRUEvictor(定义于 core/evictor_v1.py)实现了这一淘汰策略。
LRUEvictor 数据结构
class LRUEvictor(Evictor):
"""Evicts in a least-recently-used order using the last_accessed timestamp
that's recorded in the PhysicalTokenBlock. If there are multiple blocks with
the same last_accessed time, then the one with the largest num_hashed_tokens
will be evicted.
"""
def __init__(self):
# OrderedDict 保持插入顺序,块按 free 时间排列
self.free_table: OrderedDict[int, PhysicalTokenBlock] = OrderedDict()
def evict(self) -> PhysicalTokenBlock:
if len(self.free_table) == 0:
raise ValueError("No usable cache memory left")
evicted_block = next(iter(self.free_table.values()))
# 遍历找到最旧(last_accessed 最小)且前缀最长的块
for _, block in self.free_table.items():
if evicted_block.last_accessed < block.last_accessed:
break # 时间戳更大,跳过
if evicted_block.num_hashed_tokens < block.num_hashed_tokens:
evicted_block = block # 同时间戳,优先淘汰更长前缀块
self.free_table.pop(evicted_block.block_hash)
evicted_block.computed = False
return evicted_block
def add(self, block: PhysicalTokenBlock):
# 块被 free() 时调用,加入待淘汰队列
self.free_table[block.block_hash] = block
def remove(self, block_hash: int) -> PhysicalTokenBlock:
# 块被重新激活时调用(从 free 状态复活)
if block_hash not in self.free_table:
raise ValueError("Attempting to remove block not in evictor")
block = self.free_table[block_hash]
self.free_table.pop(block_hash)
return block
evict() 并非简单的 LRU——它结合了两个维度的优先级:
-
首要维度:last_accessed(时间戳)——优先淘汰最久未使用的块。
last_accessed由access_all_blocks_in_seq()在每次序列执行时更新。 - 次要维度:num_hashed_tokens(前缀长度)——时间戳相同时,优先淘汰包含更多 hashed token 的块(即更长前缀的块)。 直觉上,包含更长前缀的块在被复用时节省更多计算,但在时间戳相同时说明它们同样"过时",淘汰更长的块可以一次性清理更多缓存空间。
淘汰后,evicted_block.computed = False——因为该物理块将被用于新内容,
旧的"已计算"标记必须清除,防止 Prefix Caching 错误命中。
Evictor 与 CachedBlockAllocator 的交互
(等待被淘汰或复活) Note over A,E: 复活路径(同 hash 命中) A->>E: evictor.remove(block_hash) E-->>A: block(ref_count=0) A->>A: ref_count += 1 Note over A,E: 淘汰路径(显存满) A->>E: evictor.evict() E-->>A: LRU block(最旧/最长) A->>A: 重置 block 内容供新用
块管理流程图
下图展示了 BlockSpaceManagerV1 在整个推理生命周期中的核心操作流:
从新请求入队、Prefill 分配、Decode 追加,到 CoW 触发、Swap 换出/换入,以及最终的块释放。