SGLang HiCache KVCache Offload 方案调研

香菜烤面包

366人浏览 · 2026-06-17 17:50:57

香菜烤面包 · 2026-06-17 17:50:57 发布

官方文档：https://docs.sglang.io/docs/advanced_features/hicache_design

总结

HiCache 是 SGLang 在 RadixAttention 基础上的分层 KV Cache 方案。它把 KV cache 组织成三层：

层级	位置	作用	是否本地	数据结构
L1	GPU 显存	推理计算直接使用的 KV cache	单实例/单 rank 私有	`TokenToKVPool` / device KV pool
L2	CPU Host 内存，通常 pinned/registered	扩大本地 cache 容量，作为 L1 与 L3 的中转层	单实例/单 rank 私有	`HostKVCache` / host KV pool
L3	外部存储或分布式 KV store	跨实例共享、持久/半持久复用	集群共享，取决于 backend	`HiCacheStorage` backend

核心传输路径：

GPU L1 <----H2D/D2H----> CPU L2 <----storage backend IO----> L3

结论：

计算只能直接消费 GPU L1 中的 KV cache。命中在 L2 或 L3 的 KV，最终都要先变成 GPU device slot。
L2 -> L1 是直接 host-to-device 传输，由 HiCacheController.load() 和 HostKVCache.load_to_device_per_layer() 发起，支持 direct 或 kernel backend。
L3 -> L1 不是直接进 GPU，而是先 L3 -> L2 host pool，再 L2 -> L1 load-back。
file / HF3FS / NIXL / Mooncake 等 L3 backend 的读写对象是** page 级 KV 数据**；高性能 backend 支持把 host pool 的地址或 host indices 直接传给 backend，避免额外 CPU Tensor 拷贝。
L1 和 L2 的元数据保存在本地 HiRadixTree / UnifiedRadixTree 节点上；L3 不做全量本地元数据同步，而是在需要时按 hash key 查询 backend。
write_through 会尽早把 GPU KV 写入 L2，并继续异步写入 L3；write_back 只有 L1 eviction 时才写 L2/L3；write_through_selective 需要 hit count 达阈值后才写。
PD 分离中，Prefill 侧可启用 HiCache 做 L3 复用；Decode 侧可以启用异步 KV offload，把 decode 产生的增量 KV 写入 L3，供后续多轮对话的 Prefill/Decode prefix 恢复使用。

HiCache 整体架构

HiCache 借鉴了 CPU 多级缓存（Multi-Level Cache）的设计思想，在 RadixAttention 的 GPU prefix cache 基础上扩展出三级 KV Cache 层次：

L1: GPU 显存
- TokenToKVPool / device KV pool：计算真正消费的 KV cache，attention 只能直接读这一层
L2: CPU Host 内存
- HostKVCache：本地大容量缓存，也是 L1 和 L3 之间的中转层
L3: 分布式存储或外部 KV store
- Mooncake / HF3FS / NIXL / AIBrix / HiCacheFile / EIC / SIMM：跨实例共享 cache，按 page hash 查询和读写

其核心目标是在 LLM 推理过程中最大化 KV Cache 复用率，减少重复 Prefill 计算，从而降低推理延迟并提升系统吞吐量。

关键点：

L3 命中后常规路径不是直接进 GPU，而是先读入 Host L2，再由 load_back() 做 L2 -> L1。
计算前，所有可复用 KV 最终都必须恢复为 GPU device slot，并写入 req.prefix_indices。
HiRadixTree 只精确维护本地 L1/L2 元数据；L3 不在本地维护完整位置元数据，需要按 hash key 实时查询 backend。

核心组件

2.1 HiRadixCache

源码：python/sglang/srt/mem_cache/hiradix_cache.py

HiRadixCache 继承自 RadixCache，是非 Hybrid 模型下 HiCache 的核心管理器。Hybrid SWA/SSM 等模型通常走 UnifiedRadixCache.init_hicache()。

它扩展了 Radix Tree 节点元数据：

字段	说明
`node.value`	GPU L1 device indices
`node.host_value`	CPU L2 host indices
`node.hash_value`	Page 级 Hash，用于 L3 Key
`node.backuped`	是否已有 Host Copy
`node.evicted`	GPU Copy 是否已被驱逐

写入路径

GPU L1 -> CPU L2 -> L3 Storage

关键函数：

write_backup()：GPU → Host，内部调用 HiCacheController.write()
write_backup_storage()：Host → L3，内部调用 HiCacheController.write_storage()
_finish_write_through_ack()：GPU → Host 完成后，标记节点 Backup 完成，并触发 L3 写入

支持三种策略：

策略	行为
`write_through`	hit threshold = 1，尽早写入 Host，并继续异步写 L3
`write_through_selective`	hit threshold = 2，热数据才写入
`write_back`	正常命中不写回，GPU Eviction 时才写 Host/L3

读取路径

L3 Storage -> CPU L2 -> GPU L1

关键函数：

match_prefix()：先匹配本地 GPU/Host Prefix
prefetch_from_storage()：对本地未命中的后缀查询 L3，并异步预取到 Host
check_prefetch_progress()：按策略检查 L3 → L2 是否可以结束，并把已取回的 Page 插入 Host Tree
load_back() / init_load_back()：把 Host 命中的 KV 加载回 GPU

注意：prefetch_from_storage() 只负责 L3 → L2，不负责 L2 → L1。真正的 Host → GPU 在调度 Admission 阶段通过 init_load_back() 完成。

驱逐

evict()：GPU L1 驱逐，选择 evictable leaves
_evict_backuped()：节点已有 Host Copy，只释放 GPU Slot，保留 Host Metadata
_evict_regular()：节点没有 Host Copy，释放 GPU 并删除缓存节点
evict_host()：Host L2 驱逐，释放 host_value，但若 L3 已有数据，未来仍可从 L3 Prefetch 回来

多卡同步

HiCache 在 TP/CP/PP 场景需要同步关键决策，避免不同 Rank 命中长度不一致：

同步内容	方式
L3 Hit Length	`all-reduce min`
Prefetch Completed Tokens	`all-reduce min`
PP 场景同步	`_pp_sync()` 做流水并行间同步

2.2 HiCacheController

源码：python/sglang/srt/managers/cache_controller.py

HiCacheController 负责实际 I/O 调度：

write()：GPU -> Host
load()：Host -> GPU
start_loading()：启动 Host -> GPU 的逐层 load-back
prefetch()：把 L3 prefetch operation 放入后台队列
write_storage()：把 Host -> L3 backup operation 放入后台队列
prefetch_thread_func()：查询 L3 连续命中 page 数
prefetch_io_aux_func()：实际执行 L3 -> Host 读取
backup_thread_func()：执行 Host -> L3 写入

L1/L2 传输

HiCacheController.load()
  -> 分配 device_indices
  -> load_queue append

HiCacheController.start_loading()
  -> merge load ops
  -> move_indices()
  -> HostKVCache.load_to_device_per_layer()
  -> LayerDoneCounter 标记每层完成

HiCacheController.write()
  -> 分配 host_indices
  -> start_writing()
  -> HostKVCache.backup_from_device_all_layer()

LayerDoneCounter 是计算-传输重叠的关键。start_loading() 逐层加载 KV，并在每层完成后标记 event，使模型可以在计算第 N 层时并行加载后续层。

TransferBuffer 类仍在 cache_controller.py 中存在，但当前主要 L1/L2 路径直接通过 write_queue/load_queue、CUDA stream、event 和 HostKVCache transfer 函数完成。

2.3 HostKVCache

源码：python/sglang/srt/mem_cache/memory_pool_host.py

HostKVCache 是 L2 Host 内存池抽象，负责：

Host KV 内存分配和释放
GPU <-> Host KV 搬运
Host page flatten / set，用于 L3 generic backend
暴露 page buffer metadata，供 Mooncake/NIXL/HF3FS 等 backend zero-copy 使用

主要实现类

类	作用
`MHATokenToKVPoolHost`	MHA/GQA 类模型
`AsymmetricMHATokenToKVPoolHost`	K/V head_dim 不同的 MHA，例如 MiMo-V2
`MLATokenToKVPoolHost`	MLA 模型，例如 DeepSeek 系列
`MambaPoolHost`	Mamba/SSM state host pool
`DeepSeekV4PagedHostPool`	DeepSeek V4 paged KV/indexer side pool
`DeepSeekV4StateHostPool`	DeepSeek V4 state side pool
`DSAIndexerPoolHost`	DSA indexer host pool
`HostPoolGroup`	hybrid 多 pool 组合

Host 内存布局

常见布局：

layer_first：[2, layer, token, head, dim]，更接近逐层计算
page_first：[2, token/page, layer, head, dim]，同页数据更连续，利于 L3 I/O
page_first_direct：[2, page, layer, page_size, head, dim]，适合 direct I/O
page_head：按 head/page 组织，主要用于异构 TP/head split 场景
page_first_kv_split：Ascend/NPU MLA 相关

兼容性会被 server_args.py 自动修正：

page_first_direct + kernel 会改成 direct
page_first + direct 会改成 page_first_direct
Mooncake 不支持 layer_first 时，会按 I/O backend 改为 page_first 或 page_first_direct

Host <-> GPU backend

--hicache-io-backend：

kernel：GPU-assisted I/O kernel，indices 通常搬到 GPU，适合 kernel gather/scatter
direct：直接拷贝/索引路径，适配 layer_first / page_first_direct
kernel_ascend：Ascend/NPU 专用

从内存到显存是否直接 H2D：是，L2 -> L1 是 Host -> Device 传输，只是具体实现可能是 direct copy，也可能是 GPU-assisted kernel。

2.4 HiCacheStorage

源码：python/sglang/srt/mem_cache/hicache_storage.py

HiCacheStorage 是 L3 backend 抽象

重要接口

接口	作用
`batch_exists()`	查询一组 page key 的最长连续命中数量
`batch_get()`	generic 路径，从 L3 读出 Tensor page
`batch_set()`	generic 路径，把 Tensor page 写入 L3
`batch_get_v1()`	单 KV pool zero-copy/host_indices 路径
`batch_set_v1()`	单 KV pool zero-copy/host_indices 路径
`batch_exists_v2()`	多 pool/hybrid 场景的存在性查询
`batch_get_v2()`	多 pool/hybrid 场景读取
`batch_set_v2()`	多 pool/hybrid 场景写入

当前 Mooncake、HF3FS、NIXL 等 backend 都会用 batch_get_v1 / batch_set_v1 做高性能单池 Host zero-copy I/O。v2 主要服务多 pool/hybrid cache，例如 Mamba、SWA、DSA、DeepSeek V4 side pools、draft KV。

关键数据结构

PoolTransfer：描述一个 pool 的 host indices、device indices、keys、hit policy
PoolName：KV、MAMBA、SWA、INDEXER、DRAFT、DeepSeek V4 side pools 等
PoolHitPolicy：
- ALL_PAGES：前缀内每页都必须存在
- TRAILING_PAGES：只要求尾部若干页存在，适合 Mamba/SWA state 类数据
PoolTransferResult：返回 KV 和 side pool 的可用 page 数

2.5 L3 backend

Backend	路径	特点
Mooncake	`storage/mooncake_store/`	RDMA/DistributedStore，注册 Host buffer，支持 zero-copy 和 PD 场景
HF3FS	`storage/hf3fs/`	文件 + metadata server，支持 host page 读写
NIXL	`storage/nixl/`	统一传输 API，可对接 FILE/OBJ plugin、O_DIRECT、3FS/GDS/S3 等
AIBrix KVCache	`storage/aibrix_kvcache/`	外部生产级 KV cache backend
HiCacheFile	`hicache_storage.py`	本地文件 backend，适合测试/演示，带 LRUFileEvictor
EIC	`storage/eic/`	EIC backend
SIMM	`storage/simm/`	SIMM backend
dynamic	factory 动态加载	用户自定义 backend

L3 读写粒度是 page。page key 由 token page hash 链式计算得到，所以它表达的是“从某 prefix 继续到当前 page”的内容哈希。

2.6 CUDA/JIT Kernel

源码：

python/sglang/jit_kernel/hicache.py
python/sglang/jit_kernel/csrc/hicache.cuh

HiCache JIT kernel 用 tvm_ffi + load_jit() 编译特化 CUDA wrapper，不是普通 Python JIT。

支持：

transfer_hicache_one_layer()：MHA 单层 K/V 传输
transfer_hicache_all_layer()：MHA 全层 K/V 传输
transfer_hicache_one_layer_mla()：MLA 单层传输
transfer_hicache_all_layer_mla()：MLA 全层传输

约束：

element_size 必须是 128B 的倍数
默认 block_quota = 2
_default_unroll()：
- element_size <= 512 使用 unroll = 4
- element_size <= 1024 使用 unroll = 2
- 更大使用 unroll = 1
使用 ld.global.L1::no_allocate / st.global.L1::no_allocate，减少对 GPU L1 Cache 的污染
MHA 搬运 K 和 V
MLA 只搬运合并后的 Cache Buffer

核心逻辑

for (uint32_t i = work_id; i < length; i += kNumWorkers) {
    pos_src = indices_src[i];
    pos_dst = indices_dst[i];

    src_k = k_cache_src + pos_src * src_stride;
    dst_k = k_cache_dst + pos_dst * dst_stride;

    vec_k = load_vec<kElementSize>(src_k);
    store_vec<kElementSize>(dst_k, vec_k);

    if (!MLA) {
        // 同样搬运 V
    }
}

全层版本是在每个 token worker 内再遍历 layer。

关键工作流程

普通请求流程

请求进入 scheduler
  -> _prefetch_kvcache()
    -> req.init_next_round_input(tree_cache)
      -> tree_cache.match_prefix()
    -> tree_cache.prefetch_from_storage()

等待调度
  -> check_prefetch_progress()
  -> req.init_next_round_input(tree_cache) 重新 match
  -> PrefillAdder.add_one_req()
    -> init_load_back()
    -> req.prefix_indices = L1 hit + load_back 后的新 GPU slots

模型 forward
  -> 直接使用 GPU device KV slots

请求完成或 chunk commit
  -> cache_finished_req()/cache_unfinished_req()
  -> insert 到 radix tree
  -> 按 write policy 触发 write_backup()
  -> GPU -> Host 完成后可继续 write_backup_storage()

Prefetch 流程

prefetch_from_storage()
  -> 构造 page-aligned suffix key
  -> 检查 enable_storage / threshold / rate limit
  -> protect_host(anchor)
  -> 分配 host_indices
  -> cache_controller.prefetch()

prefetch_thread_func()
  -> _storage_hit_query()
  -> batch_exists()
  -> all_reduce(min hit length)
  -> 小于 threshold 则 revoke
  -> 否则裁剪 host_indices 并交给 IO thread

prefetch_io_aux_func()
  -> _page_transfer()
  -> batch_get_v1 或 batch_get
  -> L3 page 写入 Host pool
  -> operation.completed_tokens 增长

check_prefetch_progress()
  -> 按 best_effort / wait_complete / timeout 判断是否停止等待
  -> terminate_prefetch()
  -> all_reduce(min completed tokens)
  -> _insert_helper_host()
  -> Host tree 中出现 L2 cache nodes

best_effort 的准确语义：调度检查时可以立刻结束等待，使用已经完成的部分；不是说发起 prefetch 后马上取消所有 I/O。

Load-back 流程

init_load_back()
  -> load_back(best_match_node)
    -> 收集需要从 Host 恢复的 nodes
    -> 拼接 host_indices
    -> cache_controller.load()
      -> 分配 GPU device_indices
    -> GPU 不足则 evict() 后重试
    -> 更新 node.value = device_indices

ready_to_load_host_cache()
  -> cache_controller.start_loading()
  -> 逐层 Host -> GPU
  -> LayerDoneCounter 通知每层完成

这一步结束后，请求的 req.prefix_indices 才包含 Host/L3 命中的 prefix 对应的 GPU slots。

Write-back 流程

insert/match 命中节点
  -> _inc_hit_count()
  -> 达到 write_through_threshold
  -> write_backup()
    -> cache_controller.write()
    -> GPU -> Host
  -> writing_check()/flush_write_through_acks()
    -> _finish_write_through_ack()
    -> node.host_value 生效
    -> write_backup_storage()
      -> cache_controller.write_storage()
      -> backup_thread_func()
      -> Host -> L3

write_back 策略下，主动 hit 不写，直到 evict() 时才 backup。

4. KV Cache 在 L1/L2/L3 的状态与传输路径

HiCache 的核心逻辑可以理解成一个分层状态机：

L1 GPU 显存
  ↑  load_back / H2D
  ↓  write_backup / D2H

L2 CPU Host 内存
  ↑  prefetch from L3
  ↓  backup to L3

L3 Storage / 分布式存储

关键原则：

只有 L1/GPU 上的 KV cache能直接参与 attention 计算
L2/Host 命中不能直接计算，必须先 load-back 到 GPU
L3/磁盘/远端存储命中也不能直接计算，常规路径必须先进入 L2，再从 L2 进入 L1
同一段 KV 可以同时存在于多层，例如 GPU 有一份，Host 有一份，L3 也有一份
HiRadixTree 精确记录** L1/L2 的本地位置**；L3 不在本地保存完整位置元数据，而是按 page hash 实时查询 backend

什么时候 KV Cache 位于 L1/GPU

KV 位于 L1 的情况：

场景	说明
当前请求正在 prefill/decode	新生成的 KV 首先写入 GPU KV pool
prefix cache 在 GPU 命中	`match_prefix()` 返回 `device_indices`
Host/L3 命中后被恢复	`init_load_back()` / `load_back()` 把 Host KV 加载回 GPU
decode 阶段持续使用	decode attention 只能读 GPU 上的 KV

对应源码路径：

HiRadixCache.match_prefix()
  -> 返回 device_indices

PrefillAdder.add_one_req()
  -> req.needs_host_load_back()
  -> tree_cache.init_load_back()
  -> req.prefix_indices += new GPU indices

L1 的本质是：

TokenToKVPool / device KV pool

RadixCache/HiRadixCache 本身不是 L1，它是管理 L1/L2/L3 metadata 的树结构。

什么时候 KV Cache 位于 L2/Host

KV 位于 L2 的情况：

场景	说明
write_through	GPU KV 生成或命中后，较积极地写入 Host
write_through_selective	节点 hit count 达阈值后写入 Host
write_back	GPU cache 被 evict 时才写入 Host
L3 prefetch 命中	Storage 中的 KV page 被读入 Host pool
GPU copy 被 evict	节点可能只保留 `host_value`，成为 Host-only cache

L2 的本质是：

HostKVCache
  - MHATokenToKVPoolHost
  - MLATokenToKVPoolHost
  - MambaPoolHost
  - DeepSeekV4PagedHostPool
  - ...

节点上对应字段：

node.host_value

只要 node.host_value is not None，就说明这段 KV 在 Host 上有 copy。

什么时候 KV Cache 位于 L3/Storage

KV 位于 L3 的情况：

场景	说明
Host backup 完成后继续写 L3	`write_backup_storage()`
write_through 策略	GPU -> Host 完成后可异步写 L3
write_back 策略	L1 eviction 时先写 Host，再写 L3
PD Decode offload	Decode 侧把增量 KV 异步写入 L3
跨实例复用	其他实例按 hash key 查询到同一 KV page

L3 存储的是 page 级 KV 数据，key 通常来自 token page 的链式 hash：

tokens page -> hash -> L3 key

L3 backend 可能是：

file / Mooncake / HF3FS / NIXL / AIBrix / EIC / SIMM

L1 与 L2 之间怎么传输

GPU -> Host：D2H

写回 Host 时：

HiRadixCache.write_backup()
  -> HiCacheController.write()
  -> HiCacheController.start_writing()
  -> HostKVCache.backup_from_device_all_layer()

这是 L1 -> L2：

GPU device_indices -> Host host_indices

用于：

write-through
write-through-selective
write-back eviction
decode offload 的第一步

Host -> GPU：H2D

Host cache 要参与计算前，必须 load-back 到 GPU：

HiRadixCache.init_load_back()
  -> HiRadixCache.load_back()
  -> HiCacheController.load()
  -> HiCacheController.start_loading()
  -> HostKVCache.load_to_device_per_layer()

这是 L2 -> L1：

Host host_indices -> GPU device_indices

内存到显存是不是直接 H2D？

是，L2 到 L1 是 Host-to-Device 传输。

但实现上有两种 backend：

backend	说明
direct	更直接的 copy/indexing 路径，常配 `page_first_direct`
kernel	GPU-assisted I/O kernel，常配 `page_first` 或 `layer_first`

kernel 模式不是简单的 PyTorch copy，而是通过专门的 KV transfer kernel 做 gather/scatter；direct 更接近直接拷贝路径。

L3 与 L1 之间怎么传输

当前 HiCache 常规路径不是 L3 -> GPU 直达，而是 L3 -> Host -> GPU。

完整路径是：

L3 Storage
  -> prefetch_from_storage()
  -> batch_get_v1 / batch_get / batch_get_v2
  -> HostKVCache
  -> init_load_back()
  -> GPU KV pool

也就是：

L3 -> L2 -> L1

源码路径：

HiRadixCache.prefetch_from_storage()
  -> HiCacheController.prefetch()
  -> HiCacheController.prefetch_thread_func()
  -> HiCacheController.prefetch_io_aux_func()
  -> storage_backend.batch_get_v1() / batch_get()
  -> HostKVCache.set_from_flat_data_page() 或 zero-copy 写入 Host pool

之后：

PrefillAdder.add_one_req()
  -> tree_cache.init_load_back()
  -> Host -> GPU

需要特别说明：

Mooncake/NIXL/HF3FS 的 zero-copy 是 L3 与 Host 之间的 zero-copy，不是 L3 与 GPU 之间的 direct load。

例如：

Mooncake batch_get_into()
NIXL READ
HF3FS batch_read

这些可以直接把数据写入注册好的 Host KV buffer，但 attention 计算前仍然需要：

Host KV buffer -> GPU KV buffer

KV 的生命周期

可以用这个流程统一理解：

请求 prefill：KV 首先生成在 GPU L1
请求结束或 chunk commit：KV 插入 HiRadixTree，node.value = GPU device indices
根据 write policy：GPU L1 -> Host L2，node.host_value = Host indices
如果启用 L3 backend：Host L2 -> L3 Storage，L3 保存 page hash -> page data
GPU 内存紧张：L1 copy 被 evict，node.value = None，但 node.host_value 可能仍保留
后续请求命中同一 prefix
- 如果 L1 命中：直接复用 GPU indices
- 如果 L2 命中：Host -> GPU load_back
- 如果 L3 命中：L3 -> Host prefetch，再 Host -> GPU load_back
forward 前：req.prefix_indices 必须全部是 GPU device indices

总结

当前 KV 所在位置	是否能直接参与计算	需要做什么	传输路径
L1 GPU	可以	直接复用 device_indices	无
L2 Host	不可以	`load_back()` 到 GPU	Host -> GPU
L3 Storage	不可以	先 prefetch 到 Host，再 load-back 到 GPU	Storage -> Host -> GPU
L1 + L2	可以	优先用 L1，L2 是 backup	可选
L2 + L3	不可以	先 Host -> GPU	Host -> GPU
只有 L3	不可以	先 L3 -> Host，再 Host -> GPU	Storage -> Host -> GPU

HiCache 中 KV cache 可以同时存在于 GPU L1、Host L2 和 Storage L3，但 attention 计算只能消费 GPU L1。L2 命中时通过 load_back() 做 Host-to-Device 恢复；L3 命中时先通过 storage prefetch 读入 Host，再走同样的 Host-to-Device 恢复。因此 L3/磁盘并不直接进入 Device，Host L2 是 L3 和 GPU 之间的必要中转层。

5. 请求生命周期与策略

5.1 请求入队：本地 match + L3 prefetch 发起

普通非 PD 路径中，请求进入 scheduler 时会执行：

Scheduler._add_request_to_queue()
  -> Scheduler._prefetch_kvcache()
    -> req.init_next_round_input(tree_cache)
      -> tree_cache.match_prefix()
    -> tree_cache.prefetch_from_storage(...)

match_prefix() 做本地 L1/L2 match：

找到 GPU 上连续命中的 prefix，返回 device_indices
继续向上/向后统计 Host 命中的连续 suffix，返回 host_hit_length
返回 last_device_node、last_host_node、best_match_node

如果 L3 backend 开启，且本地 L1/L2 之后还有未命中的 suffix，scheduler 会使用 last_hash + 后续 token 计算 page hash，查询 L3。

5.2 L3 prefetch 触发条件

HiRadixCache.prefetch_from_storage() 里会先构造 page-aligned prefetch_key。触发 L3 prefetch 需要满足：

enable_storage == True，即设置了 --hicache-storage-backend
prefetch 长度达到 prefetch_threshold
未被 prefetch_rate_limited() 限流
Host pool 能分配足够或部分足够的 host slots

默认 prefetch_threshold 在 _parse_storage_backend_extra_config() 里是 256 tokens，也可以放到 --hicache-storage-backend-extra-config：

--hicache-storage-backend-extra-config '{"prefetch_threshold":512}'

5.3 L3 prefetch 停止策略

参数：

--hicache-storage-prefetch-policy {best_effort,wait_complete,timeout}

源码入口：

HiRadixCache.can_terminate_prefetch()
HiRadixCache.check_prefetch_progress()
HiCacheController.prefetch_thread_func()
HiCacheController.prefetch_io_aux_func()

策略：

策略	行为	适用
`best_effort`	调度需要执行时尽快结束，能取多少用多少	低 TTFT/SLO 优先
`wait_complete`	等 L3 命中的 pages 全部取完再继续	cache hit 最大化，适合离线/高复用
`timeout`	完成或超时即停	线上推荐，平衡 hit 和延迟

timeout 配置默认来自 PrefetchTimeoutConfig：

prefetch_timeout_base = 2.0 秒
prefetch_timeout_per_ki_token = 0.1 秒 / 1024 tokens
prefetch_timeout_max = 30.0 秒

计算逻辑：

timeout = min(
  prefetch_timeout_max,
  prefetch_timeout_base + prefetch_timeout_per_ki_token * tokens / 1024
)

配置示例：

--hicache-storage-prefetch-policy timeout \
--hicache-storage-backend-extra-config \
'{"prefetch_threshold":256,"prefetch_timeout_base":0.05,"prefetch_timeout_per_ki_token":0.02,"prefetch_timeout_max":0.5}'

5.4 调度前：L2/L3 命中必须 load-back 到 GPU

请求真正被加入 prefill batch 时，PrefillAdder.add_one_req() 会检查：

if req.needs_host_load_back():
    new_indices, req.last_node = self.tree_cache.init_load_back(...)
    req.prefix_indices = torch.cat([req.prefix_indices, new_indices])

这一步是关键：只要 prefix 命中在 Host/L3，最后都要通过 init_load_back() 得到新的 GPU device_indices，并追加到 req.prefix_indices。然后 forward 才能把这段 prefix 当作已缓存 KV 使用。

因此，L2/L3 hit 的最终效果是：

原先 GPU prefix:
  req.prefix_indices = [L1 hit slots]

load_back 后:
  req.prefix_indices = [L1 hit slots] + [newly allocated GPU slots loaded from L2]

5.5 Prefill 结束后：新 KV 插入树并按策略写回

请求 prefill/decode 产生的新 KV 初始在 GPU。完成/中途 chunk cache 时会插入 tree：

RadixCache.cache_finished_req()
RadixCache.cache_unfinished_req()
HiRadixCache._insert_helper()

HiCache 额外计算 page hash，并按 write policy 决定是否 backup。

6. 写回策略

参数：

--hicache-write-policy {write_back,write_through,write_through_selective}

6.1 write_through

含义：尽早把 GPU KV 写到 Host，并在开启 L3 backend 时继续写 L3

代码：

HiRadixCache.write_through_threshold = 1
_inc_hit_count() 中 hit count 达阈值就调用 write_backup(node)
write_backup() 调 HiCacheController.write()，执行 GPU -> Host
GPU -> Host ack 完成后 _finish_write_through_ack()
如果 enable_storage，继续 write_backup_storage()，进入 Host -> L3 backup queue

优点：

后续请求更容易命中 L2/L3
跨实例共享更积极

代价：

D2H 和 L3 写入更多，可能影响高并发/低延迟

6.2 write_through_selective

含义：不是每个节点都立即写回，而是热度达到阈值后写回

代码：

write_through_threshold = 2
_inc_hit_count() 每次命中增加 node.hit_count
达阈值才 write_backup(node)

优点：

减少冷数据写回

代价：

首次/低频复用可能没有 L2/L3 缓存

6.3 write_back

含义：只有上层被逐出时才写回下层

代码：

_inc_hit_count() 对 write_back 不主动 backup
HiRadixCache.evict() 中，如果节点要从 GPU evict 且策略是 write_back，先 write_backup(x, write_back=True)
写完 host 后再释放 GPU
L3 backend 开启时，Host -> L3 仍由 write_backup_storage() 继续处理

优点：

减少正常请求路径上的写放大

代价：

只有 eviction 后才进入 L2/L3，复用时机更滞后

7. Eviction 策略

7.1 L1 GPU eviction

GPU device pool 不够时，scheduler/allocator 会要求 tree cache evict：

HiRadixCache.evict()
  -> 选择 evictable leaves
  -> write_back 策略下先 backup 到 host
  -> _evict_backuped() 或 _evict_regular()

如果节点已经 backup 到 Host：

_evict_backuped() 释放 GPU slot
节点 value = None，但 host_value 仍在
后续请求 match 时能统计为 Host hit

如果节点没有 backup：

_evict_regular() 直接释放 GPU slot 并删除树节点
这段 cache 不再可复用

7.2 L2 Host eviction

Host pool 不够时：

HiRadixCache.evict_host(num_tokens)

规则：

只从 evictable_host_leaves 里选择
host_ref_counter > 0 的节点不能 evict，因为正在 prefetch/backup/load-back 使用
evict 后释放 host_value，节点不再有 L2 copy
如果 L3 已经写入，未来仍可能通过 L3 查询重新 prefetch

7.3 L3 eviction

L3 eviction 取决于 backend：

file backend 有 LRUFileEvictor，可以按文件目录容量做 LRU/空间控制。
HF3FS 由 metadata + file space 管理。
Mooncake/NIXL/AIBrix 由各自 backend 管理容量、淘汰或失败行为。

HiCache 主逻辑不会维护一份全局 L3 radix metadata，它只按 hash key 查询 backend 是否连续存在。

8. L1/L2 传输

L2 Host -> L1 GPU 是直接 H2D，但具体实现取决于 --hicache-io-backend 和 --hicache-mem-layout。

入口：

HiCacheController.load()
  -> 分配 device_indices
  -> load_queue append

HiCacheController.start_loading()
  -> move_indices()
  -> 对每层调用 HostKVCache.load_to_device_per_layer()

8.1 direct backend

参数：

--hicache-io-backend direct

特点：

使用更直接的拷贝/索引路径
move_indices() 会把 device_indices 放到 CPU 或按 layout 排序，适配 direct indexing
支持 layer_first 和 page_first_direct

MHA 示例代码路径：

MHATokenToKVPoolHost.load_to_device_per_layer()
transfer_kv_direct()
transfer_kv_per_layer_direct_pf_lf()

8.2 kernel backend

参数：

--hicache-io-backend kernel

特点：

使用 GPU-assisted I/O kernels 或 JIT/AOT kernel
move_indices() 会把 host indices 搬到 GPU，kernel 使用 indices 做 gather/scatter
可和 page_first / layer_first 配合
设计上更适合小 page 或复杂 stride 场景，官方文档称 GPU-assisted kernel 可显著提升传输速度

MHA 示例代码路径：

transfer_kv_per_layer()
transfer_kv_per_layer_pf_lf()
transfer_kv_all_layer()
transfer_kv_all_layer_lf_pf()
JIT：sglang.jit_kernel.hicache.transfer_hicache_*

8.3 按层 load-back 与 compute-transfer overlap

start_loading() 不是一次性全部层同步拷完，而是在 load_stream 中逐层：

for i in range(self.layer_num):
    self.mem_pool_host.load_to_device_per_layer(..., layer_id=i, ...)
    producer_event.complete(i)

LayerDoneCounter 会记录每层完成事件；tp_worker 在 batch 上设置 hicache_consumer_index。这样模型计算可等待对应层 KV load 完成，支持 L2 -> L1 传输和 prefill 计算重叠。

8.4 GPU -> Host D2H

写回 L2 时走：

HiCacheController.write()
  -> start_writing()
  -> HostKVCache.backup_from_device_all_layer()

它通常按 all-layer backup，把某个节点的所有层 KV 从 GPU 写到 Host。

9. L3/磁盘到显存

常规路径不是 L3 -> GPU 直达，而是：

L3 backend -> Host L2 pool -> GPU L1 pool

9.1 generic L3 get

通用路径：

HiCacheController._generic_page_get()
  -> storage_backend.batch_get(hash_values, dummy_page_dst)
  -> mem_pool_host.set_from_flat_data_page(host_index, data_page)

这类路径会存在中间 CPU Tensor / flatten page。

9.2 zero-copy / v1 L3 get

高性能路径：

HiCacheController._page_get_zero_copy()
  -> storage_backend.batch_get_v1(hash_values, host_indices, extra_info)

backend 可以根据 host_indices 找到 Host KV pool 中的真实地址，直接把数据读入注册好的 host memory。

例如：

Mooncake：batch_get_v1() 调 batch_get_into() / batch_get_into_multi_buffers()
NIXL：预注册 host kv_buffer 或 bounce buffer，batch_get_v1() 发 NIXL READ
HF3FS：batch_get_v1() 将 host pages 作为读目标，读文件 offsets

即使是 zero-copy，也是 L3 -> Host KV pool zero-copy，不是 L3 -> GPU device pool

一些 backend 支持 RDMA / registered memory / direct I/O：

Mooncake 可注册 host buffer 并远程读写
NIXL 可预注册 host region，并在 page-first/page-first-direct 下直接对 host kv_buffer 做 READ/WRITE
HF3FS 支持 page buffer 作为读写目标

这些优化减少的是 L3 与 Host 之间的额外拷贝和注册开销；最终计算前仍要 L2 -> L1。

10. Host Memory Layout

参数：

--hicache-mem-layout {layer_first,page_first,page_first_direct,page_first_kv_split,page_head}

10.1 layer_first

[K/V, layer, token, head, dim]

优点：

和 GPU 计算的 layer-by-layer 模式直观兼容

缺点：

L3 page 级 I/O 时，同一 page 的数据跨 layer 分散，不利于 zero-copy 大块传输

10.2 page_first

[K/V, token/page, layer, head, dim]

优点：

同一 KV page 的数据更连续，适合 L3 I/O
和 kernel backend 搭配

限制：

server_args._resolve_layout_io_compatibility() 中，如果 page_first + direct，会自动改成 page_first_direct

10.3 page_first_direct

[K/V, page, layer, page_size, head, dim]

优点：

适合 direct I/O
仍保持 page 级连续性
文档和测试里常作为 Mooncake/HF3FS/NIXL 推荐配置之一

限制：

如果 page_first_direct + kernel，会自动把 I/O backend 改成 direct

10.4 page_head

用于 MHA/GQA 异构 TP 场景，便于按 head split，支持 tp_lcm_size 让不同 TP size 的集群共享同一 L3 namespace。

示例：

--hicache-storage-backend-extra-config '{"tp_lcm_size": 8}'

11. L3 Backend 接口与实现差异

统一抽象在 HiCacheStorage：

batch_exists(keys) -> 连续命中 page 数
batch_get(keys, target_locations) -> 读 page
batch_set(keys, values) -> 写 page
batch_get_v1(keys, host_indices) -> 直接读入 host pool
batch_set_v1(keys, host_indices) -> 直接从 host pool 写出
batch_get_v2 / batch_set_v2 -> 多 pool/hybrid transfer

11.1 file backend

源码：hicache_storage.py::HiCacheFile

特点：

每个 page 存为一个 .bin 文件
key 会加模型名、TP rank/size、PP rank/size、CP rank/size suffix，避免不同 rank 冲突
set() 已存在则 skip rewrite 并 refresh recency
LRUFileEvictor 负责空间控制

适合功能验证、单机调试，不适合高性能线上 L3。

11.2 Mooncake

源码：storage/mooncake_store/mooncake_store.py

特点：

使用 Mooncake DistributedStore
注册 host buffer，使用 batch_put_from / batch_get_into
支持 multi-buffer、RDMA、metadata server、可选 SSD offload
standalone_storage=True 时要求 MooncakeHostTensorAllocator
对 MHA/MLA 生成不同 key suffix

适合跨节点低延迟 KV cache sharing。

11.3 HF3FS

源码：storage/hf3fs/storage_hf3fs.py

特点：

将 page 存在 HF3FS 文件中，用 metadata server 管理 key -> page index
batch_get_v1() / batch_set_v1() 支持按 host indices 获取 host page 并读写文件 offset
MLA 模型只 rank0 backup，避免重复

11.4 NIXL

源码：storage/nixl/hicache_nixl.py

特点：

支持 FILE/OBJ plugin
zero-copy 条件：host layout 是 page_first 或 page_first_direct，并且 O_DIRECT 时 base/stride page aligned
不满足 zero-copy 时使用预注册 bounce buffer
对 file backend 可以使用 O_DIRECT

11.5 AIBrix / EIC / SIMM / dynamic

通过 StorageBackendFactory 注册。dynamic 允许不改 SGLang 源码，传入：

--hicache-storage-backend dynamic \
--hicache-storage-backend-extra-config \
'{"backend_name":"xxx","module_path":"your.module","class_name":"YourBackend"}'

12. PD 分离场景

PD 分离中需要区分三类机制：

Prefill 节点的 HiCache：普通 HiCache，用于 prefill prefix 复用
Decode 节点的 decode radix cache：decode 侧也可本地 match prefix
Decode 节点的 KV offload：decode 产生的增量 KV 异步写入 L3，供多轮对话复用

12.1 Prefill-only HiCache

Prefill 节点启用：

--disaggregation-mode prefill \
--enable-hierarchical-cache \
--hicache-storage-backend hf3fs \
--hicache-storage-prefetch-policy wait_complete

Decode 节点不启用 decode offload。

对重复 system prompt / shared prefix，Prefill 节点可以从 L1/L2/L3 命中
Decode 输出不会写回 L3，因此多轮对话中“上一轮生成内容”不一定可被 Prefill 复用

12.2 Decode async KV offload

Decode 节点启用：

--disaggregation-mode decode \
--disaggregation-decode-enable-offload-kvcache \
--hicache-storage-backend hf3fs

源码入口：

Scheduler.__init__() 创建 DecodeKVCacheOffloadManager
DecodeKVCacheOffloadManager.offload_kv_cache(req)
DecodeKVCacheOffloadManager.check_offload_progress()

行为：

Decode 过程中 KV 初始在 GPU
每累计 offload_stride 个 page-aligned tokens，触发异步 GPU -> Host：

cache_controller.write(device_indices)

D2H 完成后触发 Host -> L3：

cache_controller.write_storage(host_indices, incremental_tokens, hash_value)

L3 backup 完成后释放 decode host pool 中对应 host indices
请求 finish 时释放 GPU slots

offload_stride：

默认等于 page_size
可通过环境变量 SGLANG_HICACHE_DECODE_OFFLOAD_STRIDE 设置
会向下对齐到 page size 的倍数，并至少为 page size

12.3 PD 下当前 Cache 在内存或磁盘时如何传输

Decode 侧如果启用 decode radix cache 和 HiCache restore，流程在 decode_hicache_mixin.py：

DecodePreallocQueue:
  match_prefix_for_req()
  -> 得到 L1 prefix_len、L2 host hit、L3 storage hit
  -> 如果有 L3 hit，先发起 prefetch_from_storage()

DecodeTransferQueue:
  _process_hicache_local_restores()
    -> 等 L3 -> L2 prefetch 完成
    -> 重新 match
    -> init_load_back() 执行 L2 -> L1
    -> ready_to_load_host_cache() 启动 H2D
  HiCacheRestoreGatedKVReceiver:
    -> PD KV receiver 成功也要等 local restore READY

因此，如果 PD 下 prefix 的 Cache 集中在 L3 或 Host：

在 Host/L2：需要 init_load_back() 分配 GPU slot，然后 H2D
在 L3：先 prefetch 到 Host/L2，再 init_load_back() 到 GPU/L1
PD network transfer 的 KV receiver 成功不代表请求可以马上 decode；HiCache restore 没 READY 时 HiCacheRestoreGatedKVReceiver.poll() 会把状态保持为 Transferring

换句话说：PD 场景下仍然需要先把本地可复用 prefix 恢复到 Decode GPU device pool，不能直接在 Host/L3 上参与 attention。

13. 启用 HiCache 示例

13.1 只启用 L2 Host Cache

python3 -m sglang.launch_server \
  --model-path /path/to/model \
  --tp 1 \
  --host 0.0.0.0 \
  --port 30000 \
  --page-size 64 \
  --enable-hierarchical-cache \
  --hicache-ratio 2 \
  --hicache-write-policy write_through \
  --hicache-mem-layout page_first \
  --hicache-io-backend kernel

13.2 启用 file L3，本地功能验证

export SGLANG_HICACHE_FILE_BACKEND_STORAGE_DIR=/tmp/sglang_hicache_file

python3 -m sglang.launch_server \
  --model-path /path/to/model \
  --tp 1 \
  --host 0.0.0.0 \
  --port 30000 \
  --page-size 64 \
  --enable-hierarchical-cache \
  --hicache-ratio 2 \
  --hicache-write-policy write_through \
  --hicache-storage-backend file \
  --hicache-storage-prefetch-policy wait_complete \
  --enable-cache-report

13.3 线上方向：page_first_direct + direct + L3 backend

python3 -m sglang.launch_server \
  --model-path /path/to/model \
  --tp 8 \
  --host 0.0.0.0 \
  --port 30000 \
  --page-size 64 \
  --mem-fraction-static 0.85 \
  --enable-hierarchical-cache \
  --hicache-ratio 2 \
  --hicache-size 0 \
  --hicache-mem-layout page_first_direct \
  --hicache-io-backend direct \
  --hicache-write-policy write_through \
  --hicache-storage-backend mooncake \
  --hicache-storage-prefetch-policy timeout \
  --hicache-storage-backend-extra-config \
  '{"prefetch_threshold":256,"prefetch_timeout_base":0.05,"prefetch_timeout_per_ki_token":0.02,"prefetch_timeout_max":0.5}'

实际 backend 配置需要按 Mooncake/HF3FS/NIXL 文档补齐。

13.4 PD Prefill 示例

python3 -m sglang.launch_server \
  --model-path /path/to/model \
  --tp 8 \
  --host 0.0.0.0 \
  --port 10000 \
  --page-size 64 \
  --enable-hierarchical-cache \
  --hicache-ratio 2 \
  --hicache-mem-layout page_first_direct \
  --hicache-io-backend direct \
  --hicache-write-policy write_through \
  --hicache-storage-backend hf3fs \
  --hicache-storage-prefetch-policy wait_complete \
  --disaggregation-mode prefill \
  --disaggregation-transfer-backend mooncake

13.5 PD Decode async offload 示例

python3 -m sglang.launch_server \
  --model-path /path/to/model \
  --tp 8 \
  --host 0.0.0.0 \
  --port 11000 \
  --page-size 64 \
  --hicache-ratio 2 \
  --hicache-mem-layout page_first_direct \
  --hicache-io-backend direct \
  --hicache-write-policy write_through \
  --hicache-storage-backend hf3fs \
  --hicache-storage-prefetch-policy wait_complete \
  --disaggregation-mode decode \
  --disaggregation-decode-enable-offload-kvcache \
  --disaggregation-transfer-backend mooncake

如果希望 decode 侧也用本地 radix/HiCache restore，需要同时关注：

--disaggregation-decode-enable-radix-cache
--enable-hierarchical-cache

具体兼容性由 server_args.py 和 kv_cache_builder.py 检查；hybrid SWA/SSM decode radix cache 当前有限制。

14. 运行时 attach/detach L3 backend

SGLang 支持服务启动后动态 attach/detach L3 backend。

要求：

服务必须 idle
没有 running requests
没有 waiting/queued requests
HiCache 已启用，即启动时有 --enable-hierarchical-cache

查询：

curl -s http://127.0.0.1:30000/hicache/storage-backend

Attach：

curl -s -X PUT http://127.0.0.1:30000/hicache/storage-backend \
  -H 'Content-Type: application/json' \
  -d '{
    "hicache_storage_backend": "mooncake",
    "hicache_storage_backend_extra_config_json": "{\"master_server_address\":\"127.0.0.1:50051\",\"protocol\":\"tcp\",\"global_segment_size\":\"4gb\",\"prefetch_threshold\":256}",
    "hicache_storage_prefetch_policy": "timeout",
    "hicache_write_policy": "write_through"
  }'

Detach：

curl -s -X DELETE http://127.0.0.1:30000/hicache/storage-backend

实现路径：

HTTP server
  -> TokenizerManager
  -> Scheduler.attach_hicache_storage_wrapped()
  -> HiRadixCache.attach_storage_backend()
  -> HiCacheController.attach_storage_backend()
  -> StorageBackendFactory.create_backend()
  -> start prefetch/backup threads

参数清单

参数	默认	说明
`--enable-hierarchical-cache`	False	开启 HiCache
`--hicache-ratio`	2.0	Host KV pool / Device KV pool token 数比例
`--hicache-size`	0	Host KV pool GB，设置后覆盖 ratio
`--hicache-write-policy`	`write_through`	写回策略
`--hicache-io-backend`	`kernel`	L1/L2 传输 backend
`--hicache-mem-layout`	`layer_first`	Host KV pool 布局
`--hicache-storage-backend`	None	L3 backend
`--hicache-storage-prefetch-policy`	源码默认 `timeout`，文档部分示例用 `best_effort`	L3 prefetch 停止策略
`--hicache-storage-backend-extra-config`	None	backend 和 prefetch JSON/TOML/YAML 配置
`--page-size`	依 server 默认	KV page 粒度，强烈影响 hit 粒度和 I/O 效率
`--enable-cache-report`	False	输出 cache 命中统计
`--disaggregation-decode-enable-offload-kvcache`	False	PD decode 侧异步 offload KV 到 L3

注意：本地源码 ServerArgs.hicache_storage_prefetch_policy 默认是 timeout；部分文档表格/旧文档可能写 best_effort。以当前源码为准。

注意事项

--enable-hierarchical-cache 与 --disable-radix-cache 互斥
Decode offload 只能在 --disaggregation-mode decode 下使用，并且必须配置 L3 backend
Hybrid SWA/SSM 的 decode radix cache 有兼容限
Host pool 分配要求内存充足；不足会启动失败。
runtime attach/detach 要求服务 idle
page_first_direct + kernel 会自动改 I/O backend 为 direct
page_first + direct 会自动改 layout 为 page_first_direct
Mooncake 不支持 layer_first 时，源码会根据 I/O backend 自动改 layo
L3 backend 的跨实例共享依赖一致的 model name、TP/PP/CP suffix、backend namespace 和 rank 组织。
MLA 下为了避免重复写，部分 backend 只让 rank0 backup。

SGLang HiCache 的核心不是让 GPU 直接从磁盘读 KV，而是用本地 HiRadixTree 管理 L1 GPU 和 L2 Host 的精确位置，用 L3 backend 做 page 级跨实例共享。请求到来时先本地 match，再按策略从 L3 prefetch 到 Host，调度前把 Host 命中 load-back 到 GPU。写回则由 write policy 决定何时从 GPU 降级到 Host，并进一步异步写入 L3。PD 分离下也是同一原则：无论 Cache 当前在 Host 还是 L3，最终都要先恢复到 Decode/Prefill 当前 worker 的 GPU device pool，才能参与 attention 计算。

加入AMD AI开发者计划！

免费领 200 小时云算力，进群参与显卡、AI PC 幸运抽奖

更多推荐

大模型推理显存不够用，试试 AMD MI300X 上的 PagedAttention 优化

本文详解如何在 AMD MI300X 上利用 PagedAttention 优化解决大模型推理显存不足难题。通过 ROCm 7.x 环境搭建、vLLM 参数调优及 FP8 量化实战，成功在单卡部署 Llama3-70B，显著提升显存利用率与并发性能，为低成本大模型推理提供高效方案。

AMD开发者中国社区

GitHub Desktop中文界面终极配置指南：3分钟快速上手

还在为GitHub Desktop的英文界面而烦恼吗？想要快速实现界面本地化却不知从何下手？这篇完整的GitHub Desktop汉化教程将为您提供最简单实用的解决方案。通过界面语言切换和本地化配置，即使是零基础开发者也能轻松完成中文版本配置，享受母语操作体验。## 为什么你需要中文界面对于习惯中文环境的开发者来说，英文界面不仅增加了学习成本，还降低了操作效率。GitHub Desktop

AMD开发者中国社区

GitHub Desktop中文界面终极配置指南：5步完成专业级汉化

还在为GitHub Desktop的英文界面而烦恼吗？想要快速实现界面本地化却不知从何下手？GitHubDesktop2Chinese项目为您提供了一套完整的解决方案，让您能够在几分钟内将GitHub Desktop界面完全转换为中文，享受母语操作体验。## 快速入门：从零开始的中文化之旅### 第一步：获取汉化工具首先，您需要获取GitHubDesktop2Chinese汉化工具。最

AMD开发者中国社区

所有评论(0)

查看更多评论

香菜烤面包

@daydayup858

已为社区贡献3条内容

SGLang HiCache KVCache Offload 方案调研

香菜烤面包

总结

HiCache 整体架构

相关源码索引

核心组件

2.1 HiRadixCache

写入路径

读取路径

驱逐

多卡同步

2.2 HiCacheController

L1/L2 传输

2.3 HostKVCache

主要实现类

Host 内存布局

Host <-> GPU backend

2.4 HiCacheStorage

重要接口

关键数据结构

2.5 L3 backend

2.6 CUDA/JIT Kernel

核心逻辑

关键工作流程

普通请求流程

Prefetch 流程

Load-back 流程

Write-back 流程

4. KV Cache 在 L1/L2/L3 的状态与传输路径

什么时候 KV Cache 位于 L1/GPU

什么时候 KV Cache 位于 L2/Host

什么时候 KV Cache 位于 L3/Storage

L1 与 L2 之间怎么传输

GPU -> Host：D2H

Host -> GPU：H2D

L3 与 L1 之间怎么传输

KV 的生命周期

总结

5. 请求生命周期与策略

5.1 请求入队：本地 match + L3 prefetch 发起

5.2 L3 prefetch 触发条件

5.3 L3 prefetch 停止策略

5.4 调度前：L2/L3 命中必须 load-back 到 GPU

5.5 Prefill 结束后：新 KV 插入树并按策略写回

6. 写回策略

6.1 write_through

6.2 write_through_selective

6.3 write_back

7. Eviction 策略

7.1 L1 GPU eviction

7.2 L2 Host eviction

7.3 L3 eviction

8. L1/L2 传输

8.1 direct backend

8.2 kernel backend

8.3 按层 load-back 与 compute-transfer overlap

8.4 GPU -> Host D2H

9. L3/磁盘到显存

9.1 generic L3 get

9.2 zero-copy / v1 L3 get

10. Host Memory Layout

10.1 layer_first

10.2 page_first

10.3 page_first_direct

10.4 page_head

11. L3 Backend 接口与实现差异

11.1 file backend

11.2 Mooncake

11.3 HF3FS

11.4 NIXL

11.5 AIBrix / EIC / SIMM / dynamic

12. PD 分离场景

12.1 Prefill-only HiCache

12.2 Decode async KV offload

12.3 PD 下当前 Cache 在内存或磁盘时如何传输

13. 启用 HiCache 示例

13.1 只启用 L2 Host Cache

13.2 启用 file L3，本地功能验证

13.3 线上方向：page_first_direct + direct + L3 backend

13.4 PD Prefill 示例