01

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), 实现了调度器所需的完整显存管理接口:

core/interfaces.py — BlockSpaceManager 抽象接口 L23-100
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 超过总显存容量,永远无法满足(直接拒绝)。

02

PhysicalTokenBlock / LogicalTokenBlock — 物理块与逻辑块数据结构

块管理的基础是两种数据结构:逻辑块LogicalTokenBlock)代表序列视角的连续 token 存储, 物理块PhysicalTokenBlock)代表 GPU/CPU 内存中实际分配的 KV Cache 存储单元。 两者均定义在 vllm/block.py 中。

LogicalTokenBlock — 逻辑视角

vllm/block.py — LogicalTokenBlock L13-53
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 — 物理视角

vllm/block.py — PhysicalTokenBlock L56-90
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 计算。

03

BlockAllocator — GPU/CPU 块分配器实现

BlockSpaceManagerV1 持有两个分配器实例——gpu_allocatorcpu_allocator, 分别管理 GPU 显存和 CPU 内存中的物理块池。 vLLM V1 提供两种分配器实现,通过 enable_caching 参数选择:

UncachedBlockAllocator — 简单空闲列表

block_manager_v1.py — UncachedBlockAllocator.__init__() / allocate() L156-210
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

block_manager_v1.py — CachedBlockAllocator.allocate() L105-140
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]
CachedBlockAllocator 的三段式分配逻辑

allocate(block_hash) 按优先级顺序检查三个位置:

  1. evictor 中:该哈希对应的块已被释放(ref_count=0)但还未被驱逐——直接"复活",避免重算 KV。
  2. cached_blocks 中:该哈希对应的块正被其他序列使用——增加 ref_count 共享,这正是 Prefix Caching 的核心。
  3. 全新分配:调用 allocate_block(),若显存满则先淘汰一块再使用。

注意 free() 在引用归零时不会立即销毁物理块,而是调用 evictor.add(block) 将其转入"缓存待淘汰"状态。 这样下次遇到相同前缀的请求时,可以从 evictor 中直接复活该块(命中第一级),实现零拷贝的 KV Cache 复用。

04

allocate() — 为新序列分配块的完整流程

当调度器决定将一个新请求(SequenceGroup)从 Waiting 队列调度执行时, 会先调用 can_allocate() 检查可用物理块是否充足,再调用 allocate() 完成实际分配。

can_allocate() — 预检

block_manager_v1.py — BlockSpaceManagerV1.can_allocate() L264-292
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 水位线机制

watermark_blocks 是一个安全余量(默认 watermark=0.01,即 1% 的 GPU 块数)。 保留这部分块作为缓冲,防止 allocate 与 append 之间出现竞态导致 OOM。 只有可用块数减去所需块数仍高于水位线,才返回 OK;否则返回 LATER(等待)。 若总容量本身不够,则返回 NEVER(直接拒绝该请求)。

allocate() — 执行分配

block_manager_v1.py — BlockSpaceManagerV1.allocate() L292-320
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 机制触发写时复制。

05

can_append_slots() / append_slots() — 追加 token 的块管理

在 Decode 阶段,每次迭代为序列生成一个新 token,需要追加到 KV Cache 中。 这由 can_append_slots() 预检 + append_slots() 执行来完成。

can_append_slots() — 检查空间

block_manager_v1.py — BlockSpaceManagerV1.can_append_slots() L321-355
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() — 追加执行

block_manager_v1.py — BlockSpaceManagerV1.append_slots() L398-444
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()

block_manager_v1.py — 辅助方法 L353-397
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
追加流程的三种情况
  1. 最后一块未满 & ref_count=1:直接在该块的空槽写入新 token KV,无需分配新块。
  2. 最后一块未满 & ref_count>1(共享块):触发 CoW,复制一个新块,在新块上写入。
  3. 最后一块已满:调用 _allocate_last_physical_block() 分配新物理块追加到 block_table。
06

Copy-on-Write 机制 — 共享块的写时复制

Copy-on-Write(CoW)是 vLLM 支持 Beam Search 和并行采样的关键机制。 当多个序列共享同一物理块(ref_count > 1)且需要写入新内容时, 必须先复制一份私有副本再写入,否则会污染其他序列的 KV Cache。

触发时机:_promote_last_block()

block_manager_v1.py — CoW 触发点 L333-368
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 子序列创建

block_manager_v1.py — BlockSpaceManagerV1.fork() L445-456
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
CoW 完整流程

以 Beam Search 为例,完整的 CoW 流程如下:

  1. Prefill 完成,prompt 的物理块由父序列独占(ref_count=1
  2. 调用 fork(parent, child),子序列与父序列共享所有物理块,每块 ref_count=2
  3. Decode 步骤,两个序列各自调用 append_slots()
  4. 若当前最后一块未满且 ref_count=2_promote_last_block() 触发 CoW:
    • 分配新物理块 new_block
    • 旧块 ref_count -= 1(变为 1,另一序列仍持有)
    • 返回 {old_block_number: [new_block_number]} CoW 映射
  5. 调度器将 CoW 映射表传给 GPU Worker,Worker 执行 cudaMemcpy 复制旧块内容到新块
  6. 此后两个序列各自持有独立的最后一块,可以独立写入
07

swap_in() / swap_out() — GPU↔CPU 块交换

当 GPU 显存不足时,调度器会选择将部分序列的 KV Cache 从 GPU 换出(swap out)到 CPU 内存, 腾出空间给更高优先级的序列。之后在合适的时机再换入(swap in)继续执行。

swap_out() — GPU → CPU

block_manager_v1.py — BlockSpaceManagerV1.swap_out() L527-552
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

block_manager_v1.py — BlockSpaceManagerV1.swap_in() L490-522
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()}
swap_out / swap_in 的对称设计

两个方法结构完全对称:遍历序列组中的每个序列 → 遍历每个物理块 → 在目标设备分配对应块 → 释放源设备块 → 更新 block_table。 返回值是一个块编号映射字典,传给 GPU Worker 后执行批量 cudaMemcpy(CPU↔GPU 之间的物理内存复制)。

注意共享块的处理:多个序列共享的块只会被换出/换入一次(通过 mapping 字典去重), 其他序列直接共享同一目标块并递增 ref_count,避免重复拷贝。

can_swap_in() / can_swap_out() — 预检

block_manager_v1.py — 换入/换出可行性检查 L470-530
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()
08

LRU Evictor — 块淘汰策略

CachedBlockAllocator 的显存已满(current_num_blocks == num_blocks), 需要分配新块时,必须从"已释放但未驱逐"的块中淘汰一个以腾出空间。 LRUEvictor(定义于 core/evictor_v1.py)实现了这一淘汰策略。

LRUEvictor 数据结构

core/evictor_v1.py — LRUEvictor L55-102
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
LRU 淘汰的双重优先级

evict() 并非简单的 LRU——它结合了两个维度的优先级:

  1. 首要维度:last_accessed(时间戳)——优先淘汰最久未使用的块。 last_accessedaccess_all_blocks_in_seq() 在每次序列执行时更新。
  2. 次要维度:num_hashed_tokens(前缀长度)——时间戳相同时,优先淘汰包含更多 hashed token 的块(即更长前缀的块)。 直觉上,包含更长前缀的块在被复用时节省更多计算,但在时间戳相同时说明它们同样"过时",淘汰更长的块可以一次性清理更多缓存空间。

淘汰后,evicted_block.computed = False——因为该物理块将被用于新内容, 旧的"已计算"标记必须清除,防止 Prefix Caching 错误命中。

Evictor 与 CachedBlockAllocator 的交互

sequenceDiagram participant A as CachedBlockAllocator participant E as LRUEvictor Note over A,E: 释放路径(free) A->>E: evictor.add(block) Note right of E: block 进入 free_table
(等待被淘汰或复活) 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 内容供新用
09

块管理流程图

下图展示了 BlockSpaceManagerV1 在整个推理生命周期中的核心操作流: 从新请求入队、Prefill 分配、Decode 追加,到 CoW 触发、Swap 换出/换入,以及最终的块释放。

flowchart TD A([新请求进入 waiting 队列]) --> B{can_allocate?} B -- NEVER --> Z1([拒绝请求]) B -- LATER --> A B -- OK --> C[allocate:\n为 prompt 分配物理块\n建立 block_table] C --> D{并行采样 n>1?} D -- 是 --> E[fork: 共享物理块\nref_count += 1] D -- 否 --> F E --> F([序列进入 running 队列]) F --> G{can_append_slots?} G -- 否 --> H{can_swap_out?} H -- 是 --> I[swap_out:\n物理块 GPU→CPU\n更新 block_table] I --> J([序列进入 swapped 队列]) H -- 否 --> K([抢占 / 等待]) G -- 是 --> L[append_slots] L --> M{最后一块满?} M -- 是 --> N[_allocate_last_physical_block\n分配新物理块] N --> O M -- 否 --> P{ref_count > 1?} P -- 是 --> Q[CoW: _promote_last_block\n复制旧块 → 新块\n返回 cow_map] Q --> O P -- 否 --> O([写入 token KV Cache]) J --> R{can_swap_in?} R -- OK --> S[swap_in:\n物理块 CPU→GPU\n更新 block_table] S --> F O --> T{生成结束?} T -- 否 --> G T -- 是 --> U[free:\n释放所有物理块\n清除 block_table] U --> V([请求完成])

物理块状态机

stateDiagram-v2 [*] --> Free : 初始化 Free --> Active : allocate()\nref_count=1 Active --> Active : fork()\nref_count++ Active --> Evictable : free()\nref_count==0\n→ evictor.add() Active --> Active : append_slots()\n追加写入 Active --> Active : CoW\n_promote_last_block() Evictable --> Active : 同 hash 命中\nevictor.remove() Evictable --> Free : evictor.evict()\n被淘汰复用 Active --> CpuActive : swap_out()\ngpu_allocator.free() CpuActive --> Active : swap_in()\ncpu_allocator.free() Active --> Free : free()\nref_count==0

关键数据结构关系图

classDiagram class BlockSpaceManagerV1 { +gpu_allocator: BlockAllocatorBase +cpu_allocator: BlockAllocatorBase +block_tables: Dict[int, BlockTable] +block_size: int +watermark_blocks: int +can_allocate() +allocate() +append_slots() +swap_in() +swap_out() +fork() +free() } class CachedBlockAllocator { +cached_blocks: Dict[int, PhysicalTokenBlock] +evictor: Evictor +current_num_blocks: int +allocate() +free() } class UncachedBlockAllocator { +free_blocks: BlockTable +allocate() +free() } class LRUEvictor { +free_table: OrderedDict +evict() +add() +remove() } class PhysicalTokenBlock { +device: Device +block_number: int +block_size: int +block_hash: int +ref_count: int +last_accessed: float +computed: bool } class LogicalTokenBlock { +block_number: int +block_size: int +token_ids: List[int] +num_tokens: int +append_tokens() } BlockSpaceManagerV1 --> CachedBlockAllocator : gpu_allocator BlockSpaceManagerV1 --> UncachedBlockAllocator : cpu_allocator CachedBlockAllocator --> LRUEvictor : evictor CachedBlockAllocator --> PhysicalTokenBlock : manages UncachedBlockAllocator --> PhysicalTokenBlock : manages BlockSpaceManagerV1 --> PhysicalTokenBlock : block_tables