KV 块的三张映射表：逻辑块号、物理块号与 KV 槽号如何串成一条 64 位索引

追加式的行写入

每次调度给出新增的块，执行器并不会重建整张表，而是调用 append_row：

1
def append_row(self, block_ids, row_idx):
2
    if not block_ids:
3
        return
4
    num_blocks = len(block_ids)
5
    start = self.num_blocks_per_row[row_idx]
6
    self.num_blocks_per_row[row_idx] += num_blocks
7
    self.block_table.np[row_idx, start : start + num_blocks] = block_ids

为什么是追加而不是全量写？因为一条请求的块序列是前缀单调增长的。tokens 只能往后加，块只能往后借，前面已分配好的块号不会改（除非发生抢占，见下一段）。而前缀缓存命中时，命中的那一段前缀块号也已经确定，同样只需要把新算那一段的块号接在后面。于是 CPU 侧的追加工作量是 $O(\text{新块数})$，而不是 $O(\text{表宽})$。注意这不等于拷贝量：提交仍然按整行宽拷贝，这笔账留到字节账一节再算。

append_row 的代价是：它假定原表这一行的前 start 个位置已经被正确填写。这个假定在第一块新块被追加时成立，但在下面两种情况下会被破坏。

move_row 为什么要把源行清零

批次里请求的行号并不是恒定的。一条请求结束后，执行器会把后面的请求往前挪一格以保持紧凑——这正是 GPUModelRunner 如何把动态请求拼成一次 GPU 前向 里讲过的「压紧」动作。对应的块表操作是 move_row(src, tgt)：

1
def move_row(self, src, tgt):
2
    num_blocks = self.num_blocks_per_row[src]
3
    block_table_np = self.block_table.np
4
    block_table_np[tgt, :num_blocks] = block_table_np[src, :num_blocks]
5
    self.num_blocks_per_row[tgt] = num_blocks
6
    # Clear the vacated source row: dummy-run batches dereference stale
7
    # rows as mamba state slots and write state in place there, possibly
8
    # after the blocks have been freed and reallocated.
9
    block_table_np[src, :num_blocks] = 0
10
    self.num_blocks_per_row[src] = 0

清零被移出的源行，源码注释给出的理由是：dummy 批次会把这些已经搬空的旧行当成 mamba 状态槽来解引用，并且原地写状态，而这些块可能已经被释放并分配给别的请求。这一句非常值得单独讨论，因为它揭示了一个在别的框架里也广泛存在的事实：表里的无效项不是「无人问」的，而是「有人在读」的。dummy 行（为图捕获或对齐而构造的假请求）总要有一条块表，而它的行号可能是任意一个刚被搬空、当前无主的旧行。如果不清零，dummy 的去引用就会落到已被回收的块上，可能把垃圾写进别的请求的 KV。

这里有一个表内坐标的漏洞：move_row 把整行搬过去了，但它并没有把表搬过去之后继续维持「源行的内容为 0」这个不变量在所有其它操作上都成立。例如 swap_row 就是一次成对的数据交换，它交换两行的有效长度与内容；clear_row 只清有效长度范围内的项，不碰之后的项。所以「未初始化项等于 0」在 row 语义层面并不成立：块表行的末尾永远是上一次写入留下的旧值。这引出后文那块表末尾残留的坑——它之所以无害，靠的是「注意力内核按 seq_lens 遍历、槽映射对无效 token 写 -1」这两道闸门，而不是块表本身干净。

双缓冲：CPU 上的 numpy 才是权威

_make_buffer 返回的是 CpuGpuBuffer，它把同一份数据在 CPU（pinned）和 GPU 上各存一份，并提供 np 视图直接改 CPU 那份、copy_to_gpu(n) 把前 $n$ 行拷过去：

1
class CpuGpuBuffer:
2
    def __init__(self, *size, dtype, device, pin_memory=PIN_MEMORY, with_numpy=True):
3
        with torch.inference_mode(False):
4
            self.cpu = torch.zeros(*size, dtype=dtype, device="cpu", pin_memory=pin_memory)
5
            self.gpu = torch.zeros_like(self.cpu, device=device)
6
        if with_numpy:
7
            self.np = self.cpu.numpy()
8

9
    def copy_to_gpu(self, n=None):
10
        cpu, gpu = self.cpu, self.gpu
11
        if n is not None:
12
            cpu, gpu = cpu[:n], gpu[:n]
13
        return gpu.copy_(cpu.pin_memory() if PIN_MEMORY else cpu, non_blocking=True)

这个结构回答了块表映射的一个关键设计问题：权威数据放在 CPU 还是 GPU？ 放在 CPU，因为追加、搬移、清零这些修改都发生在控制面，Python 里改 numpy 数组是最便宜的方式。GPU 那份是只读镜像，每次前向之前整体刷新一次。pin_memory=True 让这次刷新走的是不需要额外暂存的分页锁定内存（pinned memory），non_blocking=True 则让它挂到当前流上异步执行。

异步就带来一个顺序问题：CPU 那边改完之后，GPU 需要多久才能看到新值？vLLM 的答案是把这件事提前。执行器的 _prepare_inputs 里有一句注释说得很直白：

1
# OPTIMIZATION: Start copying the block table first.
2
# This way, we can overlap the copy with the following CPU operations.
3
self.input_batch.block_table.commit_block_table(num_reqs)

commit_block_table(n) 在CPU 侧所有后续准备工作开始之前先发起拷贝，这样后续计算 positions、query_start_loc、slot_mapping 等 CPU 计算的几微秒里，拷贝已经在 DMA 引擎上跑着。这是一种把「拷表」隐进关键路径缝隙的调度技巧，用一句话概括：把表刷新做到最前面，让它后面的每一条 CPU 指令都变成它的掩护。

不在 CPU 上算的理由

这个映射可以完全在 CPU 上算出：

对 $j$ 落在第 $r$ 条请求内、在该请求内位置为 $i$ 的 token，CPU 只要查一次块表乘以 $B$ 再加上偏移。为什么 vLLM 反而把它做成一个 Triton 内核？表面上的答案是「减少 CPU 开销」，但真正的原因更深：

块表在 GPU 上已经有了权威副本。 一旦 commit_block_table 把表拷到 GPU，在 GPU 上直接算 slot 就不需要再读 CPU 侧的表，也不需要为每个 token 维护一份 CPU 侧的 num_blocks_per_row 之外的中间结果。CPU 只需要把 positions 张量准备好，剩下的「乘 $B$ 加偏移再查表」全部交给 GPU 的 ComputeSlotMappingKernel：

1
@triton.jit
2
def kernel(num_tokens, max_num_tokens, query_start_loc_ptr, positions_ptr,
3
           block_table_ptr, block_table_stride, slot_mapping_ptr,
4
           KV_CACHE_BLOCK_SIZE: tl.constexpr, TOTAL_CP_WORLD_SIZE: tl.constexpr,
5
           TOTAL_CP_RANK: tl.constexpr, CP_KV_CACHE_INTERLEAVE_SIZE: tl.constexpr,
6
           PAD_ID: tl.constexpr, BLOCK_SIZE: tl.constexpr):
7
    req_idx = tl.program_id(0)
8

9
    if req_idx == tl.num_programs(0) - 1:
10
        # Pad remaining slots for CUDA graph compatibility.
11
        for i in range(num_tokens, max_num_tokens, BLOCK_SIZE):
12
            offsets = i + tl.arange(0, BLOCK_SIZE)
13
            tl.store(slot_mapping_ptr + offsets, PAD_ID, mask=offsets < max_num_tokens)
14
        return
15

16
    start_idx = tl.load(query_start_loc_ptr + req_idx).to(tl.int64)
17
    end_idx = tl.load(query_start_loc_ptr + req_idx + 1).to(tl.int64)
18

19
    virtual_block_size = KV_CACHE_BLOCK_SIZE * TOTAL_CP_WORLD_SIZE
20
    row_offset = req_idx * block_table_stride
21
    for i in range(start_idx, end_idx, BLOCK_SIZE):
22
        offsets = i + tl.arange(0, BLOCK_SIZE)
23
        mask = offsets < end_idx
24
        pos = tl.load(positions_ptr + offsets, mask=mask, other=0)
25
        virtual_block_indices = pos // virtual_block_size
26
        virtual_block_offsets = pos - virtual_block_indices * virtual_block_size
27
        is_local = (virtual_block_offsets // CP_KV_CACHE_INTERLEAVE_SIZE) % TOTAL_CP_WORLD_SIZE == TOTAL_CP_RANK
28
        local_block_offsets = (
29
            (virtual_block_offsets // (TOTAL_CP_WORLD_SIZE * CP_KV_CACHE_INTERLEAVE_SIZE))
30
            * CP_KV_CACHE_INTERLEAVE_SIZE
31
            + (virtual_block_offsets % CP_KV_CACHE_INTERLEAVE_SIZE)
32
        )
33
        block_numbers = tl.load(block_table_ptr + row_offset + virtual_block_indices,
34
                                mask=mask & is_local, other=0).to(tl.int64)
35
        slot_ids = block_numbers * KV_CACHE_BLOCK_SIZE + local_block_offsets
36
        slot_ids = tl.where(is_local, slot_ids, PAD_ID)
37
        tl.store(slot_mapping_ptr + offsets, slot_ids, mask=mask)

逐个部分读：

• 每个 program 处理一条请求（req_idx = tl.program_id(0)），再用 query_start_loc 把这条请求在这一步产生的 token 区间取出来。为什么按请求而不是按 token？因为块表的行偏移是 req_idx * block_table_stride，按请求切分意味着每一行只算一次行偏移，块表读取的索引计算在寄存器里完成一次即可。
• positions 是全局位置，不是请求内位置。 virtual_block_indices = pos // virtual_block_size 直接用全局位置做除法，这意味着这个 kernel 假设块表是按请求从 0 号位置连续编号的，因此 pos // B 恰好等于该请求内的逻辑块号。如果某条请求因为抢占而从头重算过，positions 已经归零重新开始，这个假定依旧成立。
• PAD_ID = -1 是无效槽号。它被写进两处：数组尾部多出来的部分（最后那个 program 负责 pad，这部分填充由槽映射独立承担，第三张表不必知道它）；以及不属于本 rank 的 token（CP 相关的那几个 is_local 判断）。对绝大多数非 CP 部署，TOTAL_CP_WORLD_SIZE = 1、is_local 恒为真，这个 kernel 退化成一个非常朴素的「查表乘 $B$ 加偏移」。
• do_not_specialize=["num_tokens", "max_num_tokens"] 让 Triton 不把这两个长度当成编译期常量，否则每换一个 batch 形状就要重新编译一次。

最大的坑：提交宽度与有效项的错位

槽映射的计算依赖两个不变量：

1. positions 是这条请求从上次计算点继续往后数的全局位置；
2. 块表行在每步 commit 前已经被正确追加。

任何一个不成立，算出来的槽号就会是错的。其中最容易踩的是第二个：commit_block_table(n) 只拷 n 行，但 get_device_tensor(num_reqs) 返回的是 block_table.gpu[:num_reqs]——切片的尾部没有拷贝。如果某次提交前，GPU 这一行的前 num_blocks_per_row 项里残留的是上一次的旧值（例如一个已经释放的块号），那么本步 kernel 会把一个 token 写进不在本请求页表内的块，而不会报任何错。注意这里说的是「有效范围内」的残留：内核按 positions 推出块号下标后按 mask 只查有效项，因此真正被读出并使用的，永远是提交时应该被刷新到的那一段。这正是为什么 move_row 必须清零源行，为什么 clear_row 要维护 num_blocks_per_row，也是为什么每个分配与提交点都要严格对齐。

FlashAttention：直接吃二维块表

vLLM 的 FlashAttention 后端把公共元数据里的 block_table_tensor 原样接收，再在自己的 forward 里把它传给底层 kernel：

1
block_table_tensor = common_attn_metadata.block_table_tensor
2
...
3
attn_metadata = FlashAttentionMetadata(
4
    ...
5
    block_table=block_table_tensor,
6
    ...
7
)

传进去的正是 block_table.gpu[:num_reqs] 的切片，宽度是 max_num_blocks_per_req。这意味着内核看到的是一个规整的二维张量：每一行对应一条请求，行宽是最大块数，行内有效元素之后的内容可能是旧值或 0。为什么这里右侧未初始化不致命？ 因为内核并不按「行宽」去遍历，它按的是 seq_lens（或 block_table 中有效块数）来遍历；行宽只是地址步长（stride）。只要行宽对齐、strides 符合期待，后面那一段读到什么都无所谓——在地址层面它永远不会被访问。这是「块表映射」与「页索引映射」最本质的一点：同一份信息，二维表靠 stride 表达，CSR 表靠指针表达，两侧的「无效区域」有不同的容错方式。

FlashInfer：先做一次 CSR 重打包

FlashInfer 的 paged attention 接口要的是 CSR 格式的三张表：paged_kv_indptr（每条请求的页起始偏移）、paged_kv_indices（拼接后的页号）、paged_kv_last_page_len（最后一页的有效长度）。vLLM 用 _compute_flashinfer_kv_metadata 从二维块表推导出这三张：

1
np.cumsum(num_blocks_np, dtype=np.int32, out=self.paged_kv_indptr.np[1 : num_reqs + 1])
2
paged_kv_indptr = self.paged_kv_indptr.copy_to_gpu(num_reqs + 1)
3

4
num_actual_pages = self.paged_kv_indptr.np[num_reqs]
5
paged_kv_indices = self.paged_kv_indices[:num_actual_pages]
6
_copy_page_indices_kernel[(num_reqs,)](
7
    paged_kv_indices, block_table_tensor, block_table_tensor.stride(0),
8
    paged_kv_indptr, BLOCK_SIZE=1024,
9
)
10

11
paged_kv_last_page_len_np = seq_lens_np % page_size
12
self.paged_kv_last_page_len.np[:num_reqs] = np.where(
13
    (paged_kv_last_page_len_np == 0) & (seq_lens_np != 0),
14
    page_size,
15
    paged_kv_last_page_len_np,
16
)
17
self.paged_kv_last_page_len.copy_to_gpu(num_reqs)

这里值得讲的不是这段代码本身，而是它的三步里每一步都在 CPU 与 GPU 之间搬一次账：

• cumsum 在 CPU 上做，因为 num_blocks_np 是「每条请求需要多少页」，它恰好是块表 num_blocks_per_row 在 CPU 上的另一份推导，CPU 侧已经有这个信息。
• 页号拷贝在内核里做，但不跨请求。 _copy_page_indices_kernel 的每个 program 只负责一条请求：从该请求行里按 [indptr[r], indptr[r+1]) 的区间把块号拷进平铺的 page_indices 中对应位置。它不会做跨程序的数据依赖，因为每个 program 的写入区间互不重叠——这正是 CSR 重打包能做成并行的原因：

1
@triton.jit
2
def _copy_page_indices_kernel(page_indices, block_table, block_table_stride,
3
                              cu_num_blocks, BLOCK_SIZE: tl.constexpr):
4
    req_idx = tl.program_id(0)
5
    row_ptr = block_table + req_idx * block_table_stride
6
    start_idx = tl.load(cu_num_blocks + req_idx)
7
    end_idx = tl.load(cu_num_blocks + req_idx + 1)
8
    num_blocks = end_idx - start_idx
9
    offset = tl.arange(0, BLOCK_SIZE)
10
    for i in tl.range(0, num_blocks, BLOCK_SIZE):
11
        block_ids = tl.load(row_ptr + i + offset, mask=i + offset < num_blocks)
12
        tl.store(page_indices + start_idx + i + offset, block_ids, mask=i + offset < num_blocks)

• paged_kv_last_page_len 的 np.where 是一个不显眼但常见的坑：当序列长度恰好是页大小的整数倍时，seq_lens % page_size == 0，按语义最后一页应该是满的，不是 0。FlashInfer 用 0 表示「空页」，所以必须显式改成 page_size。这条规则和块表里的「一行末尾是旧值」不是同一类问题：后者是无效区域，前者是有效信息。

CSR 打包的代价是额外一次 GPU 内核与几次 CPU 侧同步（paged_kv_indptr.np[num_reqs] 需要读到 CPU 值以确定 num_actual_pages），换来的是内核只需线性读取 page_indices，不需要再处理二维 stride。同一个信息在两条路线上被物化成了两种形状，本质上是「谁是查表的受害者」的不同选择：FlashAttention 让内核承担 stride，FlashInfer 让引擎承担打包。

从槽号到字节地址：布局决定最后一步

槽号到这里还只是一个「第几格」的编号，真正落到显存地址还要乘上 head 与 head_dim 两个维度。设块大小 $B$、KV 头数 $H$、每头维度 $D$，一个 token 在单层的 K 或 V 占 $H \cdot D$ 个元素。两种常见布局给出的地址公式不同。

NHD 布局（[num_blocks, B, H, D]，vLLM 的默认 KV 布局）把一个块内的 token 维放在 head 维之前：

注意 $p \cdot B + o$ 正是 slot，所以这个式子可以写成：

HND 布局（[num_blocks, H, B, D]）把 head 维放到 token 维之前：

此时 $p = \lfloor\text{slot}/B\rfloor$、$o = \text{slot}\bmod B$，所以：

两个式子的差别不在「槽号是多少」，而在「谁乘谁」：NHD 让同一个 token 的所有 head 连续，HND 让同一个 head 的所有 token 连续。槽映射本身与布局无关——它只回答「第几格」；把格号变成地址是内核的事，内核必须知道自己用的是哪种布局。这也解释了为什么同一份 slot_mapping 可以在不同 attention backend 之间复用，而不同 backend 的块表/页索引视图却不能直接互换：前者是逻辑编号，后者已经带有布局与 stride 的约定。

这也是一处常见的性能取舍：NHD 布局里，一次写入覆盖「一个 token 的全部 head」，天然与 reshape_and_cache 的按 token 处理对齐，因此 vLLM 把它作为默认；HND 布局里，一次寻址覆盖「一个 head 的全部历史 token」，更贴近部分注意力内核与 MLA 类模型对 latent 连续遍历的需求。vLLM 的 KV cache 张量形状与后端的 stride 必须严格对齐，任何一处对不齐，症状都是「写进去了但读出来错位」，而不是报错。

页索引的第三种形态：DeepInfer 的压缩块

DeepInfer 在「压缩 KV」之上又加了一层块级筛选：它维护一份压缩 KV 的候选块 topk 列表（见 候选块 topk 管线），并且对 SWA 使用 ring 槽位与块粒度装载计划。它的页索引并不等同于块表的平铺，而是块级的描述符列表：每个描述符是一段「起始页、长度」的区间。它的写与读两个方向在 InputMetadata 的构造 中有完整拆解。它的存在说明了第三张表的另一种可能性：页索引不一定是逐页的，也可以是一组区间——前提是内核能理解区间语义。

混合模型下的多张块表

当模型的不同层使用不同的 KV 结构时，一张块表不够用。vLLM 用 MultiGroupBlockTable 管理「每个 KV cache group 一张表」，每张表的块大小可以不同：

1
self.block_tables = [
2
    BlockTable(block_size, max_num_reqs, max_num_blocks_per_req, ...)
3
    for block_size, max_num_blocks_per_req, slot_mapping_mode
4
    in zip(block_sizes, max_num_blocks, slot_mapping_modes)
5
]

两种典型情况：

• 滑窗层与全注意力层共存（如 Gemma 2/3 一类交替注意力模型）。滑窗层的 KV 只需要保留窗口内的内容，块粒度与全注意力层不同，各 group 的块表宽度与块大小都不同。
• 线性注意力 / Mamba 状态层。这类层没有「KV 块」的概念，块表被当作状态槽的下标使用，因此在构造时把槽映射模式设为 SlotMappingMode.NONE，compute_slot_mapping 直接返回：

1
if self.slot_mapping_mode == SlotMappingMode.NONE:
2
    # Mamba/GDN groups consume the block table as recurrent state
3
    # indices and do not use per-token slot mappings.
4
    return

也就是说，「块表」这个名字本身在混合模型里已经不够准确：对 Mamba 组，它是「状态槽表」；对滑窗组，它是「窗口块表」；只有对全注意力组，它才是真正的分页 KV 块表。三者在同一个 BlockTable 类里共用一套追加/搬移/清零逻辑，但它们的语义由 slot_mapping_mode 区分。这也是为什么 move_row 偏偏要在注释里提到 mamba 状态槽——它是这套共用逻辑里唯一会「原地写状态」的组，也因此对源行残留最敏感。

对推理服务来说，混合模型的多组块表意味着每步的提交要对每个 group 各走一遍：每个 group 有自己的 BlockTable（块大小可能不同），共享同一组 num_reqs、query_start_loc 与 positions，元数据构建循环按 group 逐个进行：

1
for kv_cache_gid, kv_cache_group in enumerate(kv_cache_groups):
2
    cm = copy(cm_base)  # shallow copy
3
    if kv_cache_gid > 0:
4
        cm.block_table_tensor = _get_block_table(kv_cache_gid)
5
        cm.slot_mapping = slot_mappings[kv_cache_gid]
6
    ...

注意这里只换「块表与槽映射」两项，其余公共元数据（seq_lens、query_start_loc、positions）全部复用。这就是「一张块表」和「多组块表」在实现上的最小差异：公共的 batch 几何只有一份，分页的记法按 KV 结构分家。

无效槽号：-1

PAD_SLOT_ID = -1，它进入槽映射有两个来源：一是 ComputeSlotMappingKernel 里 pad 尾部与 cross-rank 的 token，二是执行器 _get_slot_mappings 对尾部 token 的显式填充：

1
slot_mapping[num_tokens_unpadded:num_tokens_padded].fill_(-1)

-1 被接受的前提是写 KV 的内核认识它，并会跳过它。 vLLM 的注释写得很清楚：Needed for reshape_and_cache in full cuda graph mode。在无图的 eager 模式下，slot_mapping 的长度恰好等于真实 token 数，尾部压根不存在，-1 也不会出现。只有在图捕获（形状必须固定）时，num_tokens_padded > num_tokens_unpadded，才会有一小段槽号落进「假 token」的范围：

1
reshape_and_cache_flash(key, value, key_cache, value_cache, slot_mapping,
2
                        self.kv_cache_dtype, layer._k_scale, layer._v_scale)

这个 op 用的是 slot_mapping 的形状来决定写多少个 token，所以它必须能对 -1 项做「不写」的处理。这同时是一个顺序约定：KV 写入发生在注意力计算之前，而注意力只读 seq_lens 范围内的 KV，写进去的假 token 不会被读到。于是即使假 token 真的被写了，影响也只是覆盖了几个以后会被覆盖的 KV 落点。

0 号块：被保留的 null block

物理块号 $0$ 在 vLLM 中不是普通块，而是被保留的 null block。执行器在给 pad 行填块表时，用的就是 NULL_BLOCK_ID = 0：

1
def _get_block_table(kv_cache_gid: int):
2
    ...
3
    blk_table = self.input_batch.block_table[kv_cache_gid]
4
    blk_table_tensor = blk_table.get_device_tensor(num_reqs_padded)
5
    # Fill unused block table entries with NULL_BLOCK_ID (null block)
6
    # for CUDAGraph padding. Block 0 is reserved for padding.
7
    blk_table_tensor[num_reqs:num_reqs_padded].fill_(NULL_BLOCK_ID)
8
    return blk_table_tensor

这里的设计逻辑比 -1 更微妙。当 num_reqs < num_reqs_padded 时，多出来的行是一条假的请求：块表全 0，seq_lens 是一个约定值（由后端决定，见 CUDA Graph 的桶与垫），它的注意力输出是垃圾，采样结果会被丢弃。但由于它是一行真实存在于批次中的请求，注意力内核会去读它的块表，去计算它的注意力。

为什么不能把它填成「别的块」，而必须是 0？有三个理由：

1. 0 号块是整个池里唯一一个必须「谁都别写」的块。 物理分配器从 0 号块之后开始发块，因此 0 号块永远不会被任何请求的真实 KV 写入。假行的 attention 读到的是一块恒为零的 KV，输出自然也是零或极小的值，不会污染真实的 batch 内别的任何东西。
2. 0 号块是常数，不随步态变化。 CUDA Graph 捕获的时候不能绑定一个会变的块号（否则图内地址计算全乱），0 是唯一保证永远不变的选择。
3. 它把「无效」表达成一个有效的数据值，而不是一个越界值。 如果填的是 max_num_blocks 这种越界值，内核会真的越界访存；填 0 则是一次完全合法的读取。

0 号块是由 BlockPool 在初始化时从空闲队列里直接摘下来「没收」的，它带一个 is_null 标记，并且引用计数不维护，因此永不会被回收：

1
# To represent a placeholder block with block_id=0.
2
# The ref_cnt of null_block is not maintained, needs special care to
3
# avoid freeing it.
4
self.null_block = self.free_block_queue.popleft()
5
self.null_block.is_null = True

这带来了 null block 的第二类用途：被跳过的块也用 0 号块占位。调度器在稀疏注意力或 remove_skipped_blocks 时，会把「不再需要、但表里位置仍然要有内容」的块名替换成 null block——KVCacheManager.remove_skipped_blocks 的文档里写得很直接：

Remove the blocks that are no longer needed from blocks and replace the removed blocks with null_block.

也就是说，块表中出现 0 号的两种情形在语义上是统一的：「这里读什么都无所谓」。无论它是 pad 行、是被淘汰的稀疏页、还是被前缀命中替换掉的占位，注意力内核都可以安全地读它。这种「用一个保留块表示空」的做法与软件工程里用 NULL 指针表示空引用是同构的；唯一要注意的是，它并没有把空引用变成「不引用」——内核仍然会走一遍访存路径，只是访存目标被钉在一个可写的常量块里。

这里也有一个与上文「块表右侧未初始化」呼应的点：fill_(NULL_BLOCK_ID) 只填了新增 row 的整行，却没有清除已有 row 右侧的残留。前者是必须清，后者可以不清，区别在它是否会在某个路径里被实际解引用。这正是 move_row 注释里说到的 dummy 行问题：dummy 行有时会落到一条被搬空的行上，而那条行右侧的旧值可能就是它要解引用的位置。

dummy 行与 slot 的「不写」语义

把 -1 与 0 拼起来，才能看清 pad 行在图里的完整语义：

这三者必须一起对齐。比如 DeepInfer 的 dummy 用 seqlen_q=1, seqlen_k=1 并显式跳过槽位 0，TensorRT-LLM 的 dummy 行 block_table=[0]、mtp_info_write_slot 落在 pad 槽位 0。它们是在各自的引擎里，靠三张表各自选了一个「最不容易被撞上的数」来保护真实算子的。

一个容易忽略的点：跨块的边界位置

考虑一条请求在两次分块预填充之间切换：

• 第一轮算到 token 1000，块表有 63 块（63 × 16 = 1008），第 63 块（逻辑块 62）已用 8 个槽。
• 第二轮从 token 1000 开始算 500 个 token。第二轮新增的块从逻辑块 63 起，但它们被追加到块表行的末尾，与逻辑块 62 相邻。
• 槽映射对 token 10001007 仍落在块表第 62 项指向的块里（偏移 815），token 1008 起才落进第 63 项。

这段推导里唯一的风险是块表行在第二轮开始时是否完整。第一轮结束后块表的 63 项已经提交给 GPU；第二轮的追加从第 64 项开始，不影响前 63 项。所以跨块边界正确的前提是「第一轮提交之后没有发生块表搬行」——一旦中间发生了一次请求被抢占或压紧，第二轮的 positions 会从 0 重来，块表也会被重写，这时「前缀部分」与前一轮的物理块可能已经不一致。这正是为什么抢占必须把整条请求从头重算：只重算残块、不重算前面已提交的块，就等于假定物理块映射没有变，而抢占恰恰收回了这些块。

query_start_loc 在映射中的双重身份

query_start_loc 是「每条请求在这步 token 流中的起始位置」，但它在两个不同阶段扮演了两个角色：

• 在 compute_slot_mapping 里，它把 program 映射到「哪条请求、哪一段 token」，因此它必须是在 GPU 上已经写好的张量；
• 在元数据里，它是注意力内核的 cu_seqlens_q，因此它必须描述有效 token 的真实边界，不能包含 pad token。

两者之所以能共用一个张量，是因为 pad token 永远排在 token 流的最末尾（num_tokens_unpadded 之后）。这也解释了为什么执行器的填充是这一段：

1
slot_mapping[num_tokens_unpadded:num_tokens_padded].fill_(-1)

fill 从「未 pad 的末尾」开始，而不是从「每条请求的末尾」开始。batch 的几何是「真实 token 在前，pad 在后」的单一前缀形结构，三张表都建立在这个结构上。如果某天为了负载均衡把 pad 插进中间，这三张表都要跟着改。

一个可手算的例子：为什么槽映射必须每步重算

设上面那条请求在第 1001 步之后被抢占，恢复时重新分配到的物理块与之前不同——逻辑块 62 指向的物理块换成了另一个。如果槽映射是持久状态，它里面保存的 $62\cdot 16 + o$ 就全部失效了。

这就是槽映射与块表生命周期不同的根本原因：

块表是持久的，槽映射是瞬时的。 把两个东西合成一张表意味着每步都要拷一次整张上下文长度的映射（SGLang 的 page size 1 在需要持久表的路径上就是这个代价的具体形式），而把两个表分开、用 Triton 每步重算，意味着每步都重新做一次「乘 $B$ 加偏移」的算术。哪个更划算取决于 $B$：$B$ 越大，重算的那次乘法省得越多；$B=1$ 时乘法消失，表退化成逐 token 的映射（SGLang 的 req_to_token），「重算」与「持久保存」的边界也随之改变。这个分界与 KV Cache 的页粒度 里对页大小的分析是同一笔账。

另一个不对称：seq_lens_cpu 的代价

FlashInfer 走原生分页路径时，需要先在 CPU 上算出每条请求的页数 num_blocks_np：

1
seq_lens_np = seq_lens_cpu.numpy()
2
num_blocks_np = (seq_lens_np + (page_size - 1)) // page_size

而在投机解码开启时，seq_lens_cpu_upper_bound 不再是精确值，于是它必须走一次真正的 GPU→CPU 同步：

1
with gpu_sync_allowed():
2
    seq_lens_cpu = common_attn_metadata.seq_lens.cpu()

一次 tensor.cpu() 就是一次流同步，它会把整条流水线的重叠打穿。这正是为什么 FlashInfer 把「把块表重打包成 CSR」这件事限定在原生分页路径：needs_native_paged_prefill、needs_native_paged_decode 判断下，XQA 与 trtllm-gen 路径直接吃 GPU 上的 block_tables，完全不进这条分支。同一条映射信息，不同内核接口的代价差别可能大到一个同步点。判断第三张表是否值得物化，首先要看它会不会把一个 GPU 张量拉回 CPU。