概述

 大语言模型推理时，KV Cache 的大小往往超过模型权重本身。在长上下文场景下，KV Cache 可能占用数十甚至上百 GB 显存，成为限制并发数和吞吐量的主要瓶颈。
 针对sparse Attention 场景，HiSparse 针对 GLM‑5/DSA 稀疏 MLA 的热缓冲区 decode 方案。完整 KV 存放在CPU pinned host 内存；GPU 只保留小块 hot‑buffer；decode 时 top‑k token 优先命中 GPU 热缓冲，未命中就从主机内存 PCIe 拷贝到 GPU 热缓冲，使用 LRU 淘汰旧缓存。
 针对deepseek的csa场景，博客给出HiSparse的工作原理：

 Hisparse需要实现两个接口：

• SwapIn：从内存中拷贝kvcache到HBM。
• BackUp：推理引擎每执行一个step，把新生成的kvcache拷贝到主机内存。

vllm PageAttention的KV Cache管理

 pic source: vLLM框架解析一：vLLM Engine 分析开篇
 如上图，在一个layer，block_size=16，Block中的每个slot可以存储一个token对应的key或value，其数据长度为row_bytes。每个slot有唯一的编号slot id。根据slot id和 cache基地址， 可以定位内存存储空间。

char dst = dst_cache + slot_id * row_bytes

dst_cache是某个层整个k cache的起始地址。row_bytes是一个token对应k cache的占用字节数。

Hisparse host cache的初始化

 Host cache位于主机内存，按照block组织。
 _allocate_kv_cache_tensors 分配一块连续的内存作为host cache。

FlashMLASparseBackend.get_kv_cache_shape 定义的形状：

class FlashMLASparseBackend(AttentionBackend):
    @staticmethod
    def get_kv_cache_shape(
        num_blocks: int,
        block_size: int,
        num_kv_heads: int,  # assumed to be 1 for MLA
        head_size: int,
        cache_dtype_str: str = "auto",
    ) -> tuple[int, ...]:
        if cache_dtype_str == "fp8_ds_mla":
            # V3.2 main MLA: 656-byte custom storage format. See module docstring.
            return (num_blocks, block_size, 656)
        else:
            return (num_blocks, block_size, head_size)

Hisparse hot cache初始化

 hot cache 申请，HiSparseCoordinator.init

class HiSparseCoordinator:
    """Per-layer decode-time hot buffer for sparse MLA KV rows.

    The pinned host-resident KV pool is the only full-size store; misses are
    always served from it. Hot-buffer hits are keyed by global KV slot id, so
    correctness relies on one invariant: a recycled slot's stale state is
    dropped before reuse — the model runner invalidates all blocks
    (re)assigned to incoming requests (covering connector RDMA loads of any
    kind) before any step can select them.
    """

    def __init__(
        self,
        config: HiSparseConfig,
        max_num_reqs: int,
        row_width: int,
        kv_dtype: torch.dtype,
        device: torch.device | str,
    ) -> None:
        self.region_stride = round_up(config.device_buffer_size + 1, HOT_REGION_ALIGN)

        row_bytes = row_width * kv_dtype.itemsize
        if row_bytes % 16 != 0:
            raise ValueError(
                f"HiSparse requires 16-byte aligned KV rows, got {row_bytes}B."
            )

        hot_rows = max_num_reqs * self.region_stride
        # Allocated eagerly so vLLM's memory profiling accounts for it when
        # sizing the main KV cache.
        self.hot_cache = torch.zeros(
            (hot_rows, row_width), dtype=kv_dtype, device=self.device
        )

 hot cache位于HBM。初始化配置可以处理的最大请求个数为max_num_reqs，每个请求申请的k cache的个数为 2 * top_k + 1。最后一个槽位用于存储last token对应的k cache。
 一个k cache占用字节数：row_bytes = row_width * sizeof(kv_dtype)。
 根据华为的测试，Layerwise and Sparse KV cache offloading to support longer sequence length：Based on DeepSeek-V3.2, we achieve 80%-90% cache hit rate by using a device buffer size of 2 * topk.
 row_width的取值，hisparse_row_width。

class FlashMLASparseImpl(SparseMLAAttentionImpl[FlashMLASparseMetadata]):
    def __init__(
        self,
        num_heads: int,
        head_size: int,
        scale: float,
        num_kv_heads: int,
        alibi_slopes: list[float] | None,
        sliding_window: int | None,
        kv_cache_dtype: str,
        logits_soft_cap: float | None,
        attn_type: str,
        kv_sharing_target_layer_name: str | None,
        # MLA Specific Arguments
        topk_indices_buffer: torch.Tensor | None = None,
        indexer: "Indexer | None" = None,
        **mla_args,
    ) -> None:
        if attention_config.hisparse_config is not None:
            if kv_cache_dtype == "fp8_ds_mla":
                hisparse_row_width = FP8_DS_MLA_ROW_BYTES
                hisparse_kv_dtype = torch.uint8
            else:
                hisparse_row_width = head_size
                hisparse_kv_dtype = kv_cache_dtype_str_to_dtype(
                    kv_cache_dtype, get_current_vllm_config().model_config
                )
            model_top_k = (
                indexer.topk_tokens
                if indexer is not None
                else get_current_vllm_config().model_config.hf_config.index_topk
            )
            self.hisparse_coordinator = create_hisparse_coordinator(
                get_current_vllm_config(),
                model_top_k,
                row_width=hisparse_row_width,
                kv_dtype=hisparse_kv_dtype,
            )

FP8_DS_MLA_ROW_BYTES

# fp8_ds_mla KV row: 512 B quantized NoPE + 16 B scales + 128 B RoPE.
FP8_DS_MLA_ROW_BYTES = 656

swap_in 方法流程详解

swap_in 方法的输入包括：

• KV 位置信息：topk_indices（由 Indexer 为每个 token 选出的 Top-K 位置）、req_id_per_token 和 block_table。
• 用于转换的辅助信息：block_size。
• slot_mapping：用于将最新 token（newest token）直接映射到其预留的热槽位。

HiSparseCoordinator.swap_in

class HiSparseCoordinator:
    def swap_in(
        self,
        *,
        kv_cache: torch.Tensor,
        req_id_per_token: torch.Tensor,
        block_table: torch.Tensor,
        topk_indices: torch.Tensor,
        block_size: int,
        slot_mapping: torch.Tensor | None,
        return_valid_counts: bool = False,
        produce_plan: bool = False,
    ) -> (
        tuple[torch.Tensor, torch.Tensor]
        | tuple[torch.Tensor, torch.Tensor, torch.Tensor]
    ):

        converted = triton_convert_req_index_to_global_index(
            req_id_per_token[:num_tokens],
            block_table,
            topk_indices,
            BLOCK_SIZE=block_size,
            NUM_TOPK_TOKENS=top_k,
            BLOCK_N=128 if top_k % 128 == 0 else top_k,
            return_valid_counts=return_valid_counts,
        )
        if return_valid_counts:
            global_indices, valid_counts = converted
        else:
            global_indices = converted
            valid_counts = None


        # Padded rows are skipped by the kernel (num_real_reqs) and must
        # come out as -1 so the attention kernel masks them.
        torch.ops._C_cache_ops.hisparse_swap_in(
            self._host_cache,
            self.hot_cache,
            global_indices,
            newest_global,
            hot_indices,
            self.device_global_indices,
            self.lru_slots,
            self.num_real_reqs,
            self.region_stride,
            miss_mask,
            self._swap_stats,
        )

    return self.hot_cache_paged(block_size), hot_indices, valid_counts

将 Token 位置转换为全局 KV 槽位 ID

triton_convert_req_index_to_global_index

def triton_convert_req_index_to_global_index(
    req_id: torch.Tensor,  # int32 [num_tokens]
    block_table: torch.Tensor,  # int32 [num_requests, max_num_blocks_per_req]
    token_indices: torch.Tensor,  # int32 [num_tokens, NUM_TOPK_TOKENS]
    BLOCK_SIZE: int = 64,
    NUM_TOPK_TOKENS: int = 2048,
    BLOCK_N: int = 128,  # tile width along columns
    HAS_PREFILL_WORKSPACE: bool = False,
    prefill_workspace_request_ids: torch.Tensor | None = None,
    prefill_workspace_starts: torch.Tensor | None = None,
    return_valid_counts: bool = False,
) -> torch.Tensor | tuple[torch.Tensor, torch.Tensor]:

 triton_convert_req_index_to_global_index 内核将 topk_indices（这是模型输出的、在请求内的 token 索引）转换为全局唯一的 KV 缓存槽位 ID（global_indices）。

 转换逻辑：利用 block_table 将请求内的位置索引映射到物理的 block_id，再结合 block_size 计算出最终的全局槽位 ID。
_convert_req_index_to_global_index_kernel

@triton.jit
def _convert_req_index_to_global_index_kernel(
    req_id_ptr,  # int32 [num_tokens]
    block_table_ptr,  # int32 [num_requests, max_num_blocks_per_req]
    token_indices_ptr,  # int32 [num_tokens, NUM_TOPK_TOKENS]
    out_ptr,  # int32 [num_tokens, NUM_TOPK_TOKENS]
    valid_count_ptr,  # int32 [num_tokens] - output valid count per row
    prefill_request_id_ptr,  # int32 [num_tokens], -1 for decode, >=0 for prefill
    workspace_starts_ptr,  # int32 [num_prefill_reqs+1] or nullptr
    # shapes (compile-time where possible)
    max_num_blocks_per_req: tl.constexpr,
    BLOCK_SIZE: tl.constexpr,
    BLOCK_N: tl.constexpr,  # tile width along columns
    HAS_PREFILL: tl.constexpr,
    COUNT_VALID: tl.constexpr,  # whether to count valid indices
    # strides (in elements)
    bt_stride0,
    bt_stride1,
    ti_stride0,
    ti_stride1,
    out_stride0,
    out_stride1,
):
    # Load token indices for this tile
    ti_ptr = token_indices_ptr + token_id * ti_stride0 + indice_id * ti_stride1
    tok = tl.load(ti_ptr)  # int32

    # Only token == -1 should propagate as -1
    is_invalid_tok = tok < 0
    is_prefill = False
    if HAS_PREFILL:
        prefill_req_id = tl.load(prefill_request_id_ptr + token_id)
        is_prefill = prefill_req_id >= 0
    # Compute block id and in-block offset
    block_id = tok // BLOCK_SIZE
    inblock_off = tok % BLOCK_SIZE

    # Guard block_table access
    valid_block = (block_id < max_num_blocks_per_req) & (block_id >= 0)
    bt_ptr = block_table_ptr + req * bt_stride0 + block_id * bt_stride1
    is_invalid_tok |= ~valid_block
    base = tl.load(bt_ptr, mask=valid_block & ~is_prefill, other=0)
    out_val = base * BLOCK_SIZE + inblock_off

调用核心 Swap-In 内核

 torch.ops._C_cache_ops.hisparse_swap_in
输入参数：

• self._host_cache: 源数据，即主机内存中的完整 KV 池。
• self.hot_cache: 目标数据，即设备端的“热缓冲区”。
• global_indices: 阶段 1 转换出的全局槽位 ID 列表。
• newest_global: 当前步骤最新 token 的全局槽位 ID（来自 slot_mapping）。
• hot_indices: 输出 Tensor，用于存储每个 Top-K 条目在 hot_cache 中的实际行索引。
• self.device_global_indices: 设备端 LRU 表，记录每个请求的 hot_cache 中存储的全局槽位 ID。这是判断命中的关键。
• self.lru_slots: 设备端 LRU 顺序表，存储每个请求中 hot_size 个槽位的 LRU 顺序（int16 索引）。
• self.num_real_reqs: 当前批次中真实请求的数量，用于跳过 CUDA Graph 中的填充（padding）行。
• self.region_stride: 每个请求在 hot_cache 中占用的行数（hot_size + 1，并对齐到 128）。
• miss_mask: 输出 Tensor，标记哪些 Top-K 条目在 hot_cache 中未命中，需要从主机加载。
• self._swap_stats: 用于记录命中/未命中次数的统计 Tensor。

 lru_slots中存储的槽位顺序：[可驱逐的槽位] + [新加载的条目] + [命中的条目]。

 hisparse_swap_in 调用 hisparse_swap_in_kernel 。每个 CUDA block 处理一个请求（batch row）。
 hisparse_swap_in_kernel 是一个单内核解决方案，在一个 GPU 内核中完成了：

1. 缓存查找（Cache Lookup）
2. LRU 状态管理
3. 缓存驱逐（Eviction）
4. 数据加载（从主机内存到设备内存）

 hisparse_swap_in 中有 5 个阶段。假设当前批次有 1 个请求 (num_rows = 1)：

• top_k = 8：索引器选择 8 个最相关的 token
• hot_size = 16：热缓冲区大小为 16（2 * top_k）
• hash_size = 2 * top_k = 16

 初始状态：

# 从最久未使用到最新使用
lru_slots = [3, 1, 5, 0, 7, 2, 4, 6, 8, 9, 10, 11, 12, 13, 14, 15]
# 对应的设备全局索引（槽位 0-15 存储的全局 ID）
device_global_indices = [100, 200, 300, 400, 500, 600, 750, 800, 
                         900, 1000, 1050, 1200, 1300, 1400, 1500, 1600]
# 索引器输出的 8 个全局 ID
global_indices = [150, 200, 700, -1, 300, 800, 1100, 900]
# 当前步新 token 的全局 ID
newest_global = 150
# hot_base = row * region_stride = 0

Phase 1: 初始化与哈希

 预处理 top‑k，过滤无效、最新 token，构建哈希表

const int32_t* row_topk = global_indices + static_cast<int64_t>(row) * top_k;
// Phase 1: 每个线程处理一个 top_k 条目
for (int i = tid; i < top_k; i += blockDim.x) {
    const int32_t g = row_topk[i];  // 读取全局索引
    if (row_miss != nullptr) row_miss[i] = 0;
    
    if (g < 0) {
        // 无效条目：直接标记为完成
        row_out[i] = -1;
        s_topk[i] = kTokenDone;
        atomicAdd(&s_counters[1], 1);  // 已解决计数+1
    } else if (g == newest_id) {
        // 新 token：使用保留槽位 (hot_size)
        row_out[i] = static_cast<int32_t>(hot_base) + hot_size;  // = 0 + 16
        s_topk[i] = kTokenDone;
        atomicAdd(&s_counters[1], 1);
    } else {
        // 常规条目：插入哈希表
        int slot = hash_slot(g, hash_size);
        while (true) {
            const int32_t old = atomicCAS(&s_hash_keys[slot], kHashEmpty, g);
            if (old == kHashEmpty || old == g) {
                s_hash_vals[slot] = static_cast<int16_t>(i);
                break;
            }
            slot = (slot + 1) % hash_size;
        }
        s_topk[i] = g;  // 保存供后续处理
    }
}
__syncthreads();

 如果 g = -1（无效），直接标记为完成 (kTokenDone)。
 如果 g = newest_id，则将其在 hot_cache 中的位置设置为预留的最后一个槽位（hot_size），并标记为完成。
 构建哈希表：对于其他有效的 global_indices，使用一个开放寻址哈希表 (s_hash_keys/s_hash_vals) 将其存储起来，value 是它在 row_topk 列表中的索引 (i)。
按照上面的实例，phase 1执行后的，哈希表的结果：

// 假设 hash(200)=0, hash(700)=1, hash(300)=2, hash(800)=3, hash(1100)=4, hash(900)=5
s_hash_keys = [200, 700, 300, 800, 1100, 900, -1, -1, -1, -1, -1, -1, -1, -1, -1, -1]
s_hash_vals = [1,   2,   4,   5,   6,    7,   -1, -1, -1, -1, -1, -1, -1, -1, -1, -1]

Phase 2: 缓存查找

 LRU 遍历：按照 lru_slots 记录的 LRU 顺序（从最久未使用到最新使用），遍历该请求的 hot_size 个热槽位。
 查找命中：对于每个槽位，从device_global_indices读取其中存储的 global_indices (cached_g)，并用它在 Phase 1 构建的哈希表中查找。

• 命中 (Hit)：如果在哈希表中找到了 cached_g，则将该 Top-K 条目的输出 row_out 设置为 hot_base + slot（当前热槽位），并将该 Top-K 条目标记为完成。
• 未命中 (Miss)：如果没找到，或者槽位无效，则此槽位为“可驱逐的”(is_evictable)。

 为了高效地进行后续的驱逐和 LRU 更新，该阶段使用并行前缀和 (warp_inclusive_scan) 将命中的槽位和可驱逐的槽位分别放置到共享内存数组 s_lru_out 的两端

int16_t* row_lru = lru_slots + static_cast<int64_t>(row) * hot_size;
// Phase 2: 扫描 LRU 链表，分类命中/可驱逐
const int iters_buffer = (num_buffer_chunks + NUM_WARPS - 1) / NUM_WARPS;
int total_hit_count = 0;
int total_evict_count = 0;

for (int iter = 0; iter < iters_buffer; iter++) {
    const int chunk_idx = warp_id + iter * NUM_WARPS;
    const bool has_valid_chunk = chunk_idx < num_buffer_chunks;
    
    const int pos = chunk_idx * kWarpSize + lane_id;
    const bool has_valid_pos = has_valid_chunk && (pos < hot_size);
    const int16_t slot = has_valid_pos ? row_lru[pos] : int16_t(-1);
    // 从 LRU 中读取槽位号，然后读取该槽位存储的全局 ID
    const int32_t cached_g = (slot >= 0 && slot < hot_size) ? row_dgi[slot] : -1;
    
    // 在哈希表中查找 cached_g
    int found_topk_idx = -1;
    if (cached_g >= 0) {
        int h = hash_slot(cached_g, hash_size);
        while (true) {
            const int32_t k = s_hash_keys[h];
            if (k == cached_g) {
                found_topk_idx = static_cast<int32_t>(s_hash_vals[h]);
                break;
            }
            if (k == kHashEmpty) break;
            h = (h + 1) % hash_size;
        }
    }
    
    const bool is_hit = found_topk_idx >= 0;
    const bool is_evictable = has_valid_pos && !is_hit;
    
    // 命中：标记对应的 top_k 条目为已完成
    if (is_hit) {
        s_topk[found_topk_idx] = kTokenDone;
        row_out[found_topk_idx] = static_cast<int32_t>(hot_base) + slot;
    }
    
    // 使用 warp 级同步进行压缩（详细代码见原 kernel）
    // ... 计算偏移并写入 s_lru_out ...
}
__syncthreads();

lru_slots 遍历，当前 LRU 状态（16 个槽位）：

lru_slots = [3, 1, 5, 0, 7, 2, 4, 6, 8, 9, 10, 11, 12, 13, 14, 15]
             0  1  2  3  4  5  6  7  8  9  10  11  12  13  14  15
             ↑ 从最久未使用到最新使用 ↑

索引pos	slot = row_lru[pos]	cached_g = row_dgi[slot]	在哈希表中？
0	3	400	no
1	1	200	yes
2	5	600	no
3	0	100	no
4	7	800	yes
5	2	300	yes
…	…	…	…
8	8	900	yes

命中槽位：1 (200), 7 (800), 2 (300), 8 (900)。
total_hits = 4
s_counters[0] = 4
更新后的 s_topk：

// T0(150新token)已解决, T1(200)已命中, T2(700)未解决, T3(-1)已解决,
// T4(300)已命中, T5(800)已命中, T6(1100)未解决, T7(900)已命中
s_topk = [kTokenDone, kTokenDone, 700, kTokenDone, kTokenDone, kTokenDone, 1100, kTokenDone]
         索引0      索引1      索引2  索引3      索引4      索引5      索引6  索引7

处理后 s_lru_out 数组

命中项（从 s_lru_out 前端开始放）：
  s_lru_out[0] = 1  (slot 1)
  s_lru_out[1] = 7  (slot 7)
  s_lru_out[2] = 2  (slot 2)
  s_lru_out[3] = 8  (slot 8)
  命中项总数 = 4

可驱逐项（从 s_lru_out 后端开始放，hot_size=16）：
  s_lru_out[15] = 3  (slot 3, 第1个可驱逐)
  s_lru_out[14] = 5  (slot 5, 第2个可驱逐)
  s_lru_out[13] = 0  (slot 0, 第3个可驱逐)
  s_lru_out[12] = 4  (slot 4, 第4个可驱逐)
  s_lru_out[11] = 6  (slot 6, 第5个可驱逐)
  s_lru_out[10] = 9  (slot 9, 第6个可驱逐)
  s_lru_out[9]  = 10 (slot 10, 第7个可驱逐)
  s_lru_out[8]  = 11 (slot 11, 第8个可驱逐)
  s_lru_out[7]  = 12 (slot 12, 第9个可驱逐)
  s_lru_out[6]  = 13 (slot 13, 第10个可驱逐)
  s_lru_out[5]  = 14 (slot 14, 第11个可驱逐)
  s_lru_out[4]  = 15 (slot 15, 第12个可驱逐)

最终 s_lru_out 完整布局：
s_lru_out = [1, 7, 2, 8, 15, 14, 13, 12, 11, 10, 9, 6, 4, 0, 5, 3]
             0  1  2  3  4   5   6   7   8   9   10  11 12 13 14 15
             ↑ 命中项 (4个) ↑  ↑     可驱逐项 (12个)          ↑

Phase 3: 处理未命中与分配新槽位 (Miss Handling & Slot Assignment)

 Phase 1 和 Phase 2完成后，s_topk 中剩下的是未命中条目（既不是无效/最新，也没有在热缓存中找到）。
 分配驱逐槽位：再次使用并行前缀和，将未命中的条目“压缩”到 s_topk 的前面。同时，从 s_lru_out 的尾部（即最久未使用的槽位）按顺序为这些未命中的条目分配要驱逐的槽位 (evict_slot)。
 更新 LRU 表：更新 row_dgi[evict_slot]，将其记录为新的全局 ID (g)，并设置 row_out[i] 为该槽位的 hot_base + evict_slot。

// Phase 3: 处理缺失项，分配可驱逐槽位
// 首先重置前缀和缓冲区
for (int i = tid; i < num_token_chunks + 1; i += blockDim.x) {
    s_chunk_off[i] = 0;
}
__syncthreads();

const int iters_token = (num_token_chunks + NUM_WARPS - 1) / NUM_WARPS;
int miss_running_total = 0;

for (int iter = 0; iter < iters_token; iter++) {
    const int chunk_idx = warp_id + iter * NUM_WARPS;
    const bool has_valid_chunk = chunk_idx < num_token_chunks;
    
    const int i = chunk_idx * kWarpSize + lane_id;
    const bool has_valid_token = has_valid_chunk && (i < top_k);
    
    int32_t g = 0;
    bool is_miss = false;
    if (has_valid_token) {
        // 检查是否未解决 (既不是新token也不是命中)
        is_miss = s_topk[i] != kTokenDone;
        if (is_miss) {
            g = s_topk[i];
        }
    }
    
    // 计算本地偏移（使用 warp 级同步）
    int local_miss_off = 0;
    if (has_valid_chunk) {
        const unsigned int miss_mask_val = __ballot_sync(0xFFFFFFFF, is_miss);
        local_miss_off = __popc(miss_mask_val & lanes_before);
        if (lane_id == 0) {
            s_chunk_off[chunk_idx + 1] = __popc(miss_mask_val);
        }
    }
    __syncthreads();
    
    // 计算全局偏移（前缀和）
    if (warp_id == 0) {
        miss_running_total = warp_inclusive_scan(s_chunk_off, lane_id, chunk_idx + 1,
                                                 num_token_chunks + 1, miss_running_total);
    }
    __syncthreads();
    
    // 分配槽位
    if (is_miss) {
        const int m = s_chunk_off[chunk_idx] + local_miss_off;
        // 从可驱逐列表的后端取槽位
        const int16_t evict_slot = s_lru_out[hot_size - 1 - m];
        if (evict_slot < 0 || evict_slot >= hot_size) {
            row_out[i] = -1;
        } else {
            // 紧凑存储全局 ID，供 Phase 5 使用
            s_topk[m] = g;
            row_out[i] = static_cast<int32_t>(hot_base) + evict_slot;
            if (row_miss != nullptr) row_miss[i] = 1;
            row_dgi[evict_slot] = g;
        }
    }
}
__syncthreads();

针对s_topk中的值：
i=2: 700 → 是缺失（第1个缺失项）
i=6: 1100 → 是缺失（第2个缺失项）

槽位分配
第1个缺失项 (m=0，来自 i=2，g=700)：

• 从可驱逐列表取最久未使用的槽位： s_lru_out[hot_size - 1 - 0] = s_lru_out[15] = 3
• row_out[2] = 0 + 3 = 3
• miss_mask[2] = 1
• row_dgi[3] = 700（更新设备全局索引）
• s_topk[0] = 700

第2个缺失项 (m=1，来自 i=6，g=1100)：

• 从可驱逐列表取次久未使用的槽位：s_lru_out[hot_size - 1 - 1] = s_lru_out[14] = 5
• row_out[6] = 0 + 5 = 5
• miss_mask[6] = 1
• row_dgi[5] = 1100（更新设备全局索引）
• s_topk[1] = 1100

Phase 4: LRU 列表更新

Phase 4构建新的 LRU 顺序。row_lru 存储的槽位顺序：[可驱逐的槽位] + [新加载的条目] + [命中的条目]。

// Phase 4: 写入新的 LRU 顺序
const int total_evictable = hot_size - total_hits;  // 16 - 4 = 12
for (int i = tid; i < hot_size; i += blockDim.x) {
    if (i < total_misses) {  // i < 2
        // 新加载的数据放在可驱逐项之后，命中项之前
        row_lru[total_evictable - total_misses + i] = s_lru_out[hot_size - 1 - i];
    } else if (i < total_evictable) {  // 2 <= i < 12
        // 被驱逐的项（最久未使用）
        row_lru[i - total_misses] = s_lru_out[hot_size - 1 - i];
    } else {  // i >= 12
        // 命中项（最近使用）
        row_lru[i] = s_lru_out[i - total_evictable];
    }
}

 s_lru_out 回顾：

s_lru_out = [1, 7, 2, 8, 15, 14, 13, 12, 11, 10, 9, 6, 4, 0, 5, 3]
             0  1  2  3  4   5   6   7   8   9   10  11 12 13 14 15
             ↑ 命中项 (4个) ↑  ↑     可驱逐项 (12个)          ↑

 计算新 row_lru （hot_size=16, total_hits=4, total_misses=2, total_evictable=12）。s_lru_out 的最后两个slot值分配为新的加载项。

i	条件	计算	含义
0	i < total_misses (0<2)	row_lru[12 - 2 + 0] = row_lru[10] = s_lru_out[15] = 3	被驱逐的第1项
1	i < total_misses (1<2)	row_lru[12 - 2 + 1] = row_lru[11] = s_lru_out[14] = 5	被驱逐的第2项
2	i < total_evictable (2<12)	row_lru[2 - 2] = row_lru[0] = s_lru_out[13] = 0	可驱逐项（最久）
3	i < total_evictable (3<12)	row_lru[3 - 2] = row_lru[1] = s_lru_out[12] = 4	可驱逐项
4	i < total_evictable (4<12)	row_lru[4 - 2] = row_lru[2] = s_lru_out[11] = 6	可驱逐项
5	i < total_evictable (5<12)	row_lru[5 - 2] = row_lru[3] = s_lru_out[10] = 9	可驱逐项
6	i < total_evictable (6<12)	row_lru[6 - 2] = row_lru[4] = s_lru_out[9] = 10	可驱逐项
7	i < total_evictable (7<12)	row_lru[7 - 2] = row_lru[5] = s_lru_out[8] = 11	可驱逐项
8	i < total_evictable (8<12)	row_lru[8 - 2] = row_lru[6] = s_lru_out[7] = 12	可驱逐项
9	i < total_evictable (9<12)	row_lru[9 - 2] = row_lru[7] = s_lru_out[6] = 13	可驱逐项
10	i < total_evictable (10<12)	row_lru[10 - 2] = row_lru[8] = s_lru_out[5] = 14	可驱逐项
11	i < total_evictable (11<12)	row_lru[11 - 2] = row_lru[9] = s_lru_out[4] = 15	可驱逐项
12	i >= total_evictable (12>=12)	row_lru[12] = s_lru_out[12 - 12] = s_lru_out[0] = 1	命中项（最近）
13	i >= total_evictable (13>=12)	row_lru[13] = s_lru_out[13 - 12] = s_lru_out[1] = 7	命中项（最近）
14	i >= total_evictable (14>=12)	row_lru[14] = s_lru_out[14 - 12] = s_lru_out[2] = 2	命中项（最近）
15	i >= total_evictable (15>=12)	row_lru[15] = s_lru_out[15 - 12] = s_lru_out[3] = 8	命中项（最近）

 row_lru 最终结果：

row_lru = [0, 4, 6, 9, 10, 11, 12, 13, 14, 15, 3, 5, 1, 7, 2, 8]

Phase 5: 从host cache拷贝缺失数据到hot cache

// Phase 5: 从主机内存拷贝缺失数据到热缓冲区
for (int m = warp_id; m < total_misses; m += NUM_WARPS) {
    const int32_t g = s_topk[m];  // 紧凑存储的全局 ID
    const int16_t evict_slot = s_lru_out[hot_size - 1 - m];
    if (evict_slot < 0 || evict_slot >= hot_size) {
        continue;  // 跳过无效槽位
    }
    char* dst = hot_cache + (hot_base + evict_slot) * row_bytes;
    
    if (g >= 0 && g < host_rows) {
        // 从主机内存拷贝数据
        copy_row_warp(lane_id, host_cache + static_cast<int64_t>(g) * row_bytes,
                      dst, row_bytes);
    } else {
        // 无源数据，清零热槽位
        zero_row_warp(lane_id, dst, row_bytes);
        __syncwarp();
        if (lane_id == 0) {
            row_dgi[evict_slot] = -1;  // 撤销所有权
        }
    }
}

 copy任务（每个 warp 处理一个缺失项）：

m (缺失项索引)	g (全局 ID)	evict_slot	目标地址	操作
0	700	s_lru_out[15] = 3	hot_cache[3]	host_cache[700] → hot_cache[3]
1	1100	s_lru_out[14] = 5	hot_cache[5]	host_cache[1100] → hot_cache[5]

 以拷贝全局 ID 700 为例，假设 KV 行大小为 656 字节（41 个 16 字节向量），对应 FP8_DS_MLA_ROW_BYTES = 656。
 copy_row_warp 函数分析:

__device__ __forceinline__ void copy_row_warp(int lane_id, const char* src,
                                              char* dst, int64_t row_bytes) {
    const int64_t num_vec = row_bytes / 16;  // 656 / 16 = 41
    const uint64_t* src8 = reinterpret_cast<const uint64_t*>(src);
    uint64_t* dst8 = reinterpret_cast<uint64_t*>(dst);
    
    for (int64_t j = lane_id; j < num_vec; j += kWarpSize) {
        // 每个 lane 负责拷贝一个 16 字节向量
        uint64_t lo, hi;
        const uint64_t* s = src8 + j * 2;
        
        // 使用 PTX 指令：从全局内存加载 16 字节（两个 64 位）
        asm volatile("ld.global.nc.v2.b64 {%0,%1},[%2];"
                     : "=l"(lo), "=l"(hi)
                     : "l"(s)
                     : "memory");
        
        // 使用 PTX 指令：存储 16 字节到全局内存（L2 缓存，非相干）
        uint64_t* d = dst8 + j * 2;
        asm volatile("st.global.cg.v2.b64 [%0],{%1,%2};"
                     ::"l"(d), "l"(lo), "l"(hi)
                     : "memory");
    }
}

 warp 0（32 个线程，lane_id 0-31）处理缺失项 m=0（g=700）：

row_bytes = 656 字节
num_vec = 41 个 16 字节向量
kWarpSize = 32 个线程

每个 lane 负责拷贝的向量索引 j：
lane 0:  j = 0, 32          (2 个向量 = 32 字节)
lane 1:  j = 1, 33          (2 个向量 = 32 字节)
lane 2:  j = 2, 34          (2 个向量 = 32 字节)
...
lane 8:  j = 8, 40          (2 个向量 = 32 字节)
lane 9:  j = 9              (1 个向量 = 16 字节)
lane 10: j = 10             (1 个向量 = 16 字节)
...
lane 31: j = 31             (1 个向量 = 16 字节)

reference

Tair 联手 SGLang 共建 DeepSeekV4 分层缓存架构
HiSparse: host-resident sparse-MLA decode hot-buffering-pr46326
HiSparse: host-resident sparse-MLA decode hot-buffering-pr51323
HiSparse: Turbocharging Sparse Attention with Hierarchical Memory
HiSparse 源码解析：sglang 是怎么把 KV Cache 搬下 GPU 的
vLLM框架解析一：vLLM Engine 分析开篇