多级缓存、分布式缓存等话题在传统计算机领域里讨论已久,对于 LLM Serving 有诸多可借鉴的地方。总体来说,多级缓存的核心就是判断:GPU 上的 prefix cache 被淘汰后,再次遇到相同前缀时,该选择重新 Prefill,还是从较慢的存储层把 KV Cache 搬回 GPU。
引入 host memory 后,系统可以把 GPU 上较冷的 KV Cache 保存在 DRAM,需要时再搬回 GPU。这样延长了 KV Cache 的驻留时间,也提高了整体容量,但同时引入了 H2D/D2H 数据传输。一次命中能否真正带来加速,还得评估额外开销的耗时。这笔账最终归结为两个时间:从 host memory 恢复 KV 的实际延迟,以及重新计算这段 prefix 所需的时间。 前者更短,恢复 host 缓存才有收益。
这仍然是经典的 memory hierarchy 问题。本文把 GPU HBM、host memory 和外部存储分别记为 L1、L2、L3。越往下容量越大,但带宽通常越低,访问路径也更长。判断 L2/L3 是否值得使用,先看新增容量能换来多少复用,再比较一次命中的 load 与重算,最后检查并发 I/O 是否抵消单请求收益。Cache 布局、传输路径和拷贝方式共同决定 load 的成本与优化空间。
store 将 KV 写向下层,load 将命中的 KV 搬回 L1,evict 则释放空间。KV 只要还存在于任意一层,就能继续复用,但下层命中的数据必须先回到 GPU,才能参与 Attention。SGLang HiCache 直接采用 L1/L2/L3 这套命名1。vLLM 将 CPU 称为 primary tier,将 filesystem、P2P、object store 等称为 secondary tiers,后者通常也要经过 CPU 才能回到 GPU2。
写入策略
host memory 的容量只决定最多能保留多少数据,写入策略决定这些空间留给哪些 prefix,以及何时产生 D2H。常见策略有三种:
| 策略 | 提交时机 | 收益 | 代价 |
|---|---|---|---|
| write-through | KV 在上层生成后立即向下一层提交 | 保留所有潜在复用 | 写放大最大,低复用数据挤占下一层 |
| write-through-selective | 观察到足够复用信号后提交 | 过滤一次性前缀 | 达到阈值前仍要重算或访问更慢的缓存层 |
| write-back | 上层淘汰时提交 | 写流量最小 | 写入集中在上层容量紧张时,与加载竞争 |
从层间关系看,在 L2 容量足够且写入已经完成时,write-through 可以视为 L1 ⊆ L2:L1 里的完整 blocks 都能在 L2 找到副本;write-through-selective 只写入满足复用条件的 blocks;write-back 则等 L1 淘汰时再写。
大量请求共享同一个 system prompt 时,一次 D2H 可以换来多次复用。对于长尾对话历史,如果大多数前缀只复用一两次,过低的写入阈值会让 L2 塞满低复用的 KV,并持续消耗 D2H 带宽。
异步写的实现
把 KV 写入下一层会占用存储空间和传输带宽。L1→L2 产生 D2H,L2→L3 则产生 I/O。异步写避免了同步阻塞当前请求。典型实现会异步提交写入,并让每段更新自己的缓存状态。这样,不同批次的推理、D2H 和 I/O 可以并行。异步队列还能并发处理多个读写任务,提高 I/O 吞吐。
| 阶段 | 异步提交 | 完成判断 | 完成前必须保护的状态 |
|---|---|---|---|
| L1→L2 | cudaEventRecord 标记 KV 生成完成,cudaStreamWaitEvent 让 write stream 等待,再异步执行 D2H | 用 cudaEventQuery 轮询完成 event,或由 cudaLaunchHostFunc 触发 callback | GPU 源 slot 和目标 host page |
| L2→L3 | 把任务放进队列,由线程池调用 pwrite 等同步 I/O 接口,也可以直接使用 io_uring、RDMA 等异步接口 | 线程池以同步调用返回为准,io_uring、RDMA 则读取完成队列 | host page、存储 key 和相关元数据 |
异步写要区分任务已经提交和数据已经可用。任务提交后先进入 pending 状态,确认完成后再转为 ready。cache 查找只返回 ready 数据,避免请求读到尚未写完或写入失败的 KV。
后台线程可以直接调用 pwrite 等同步 I/O 接口,实现简单、兼容性好,但每个阻塞 I/O 都会占用线程。io_uring、RDMA 等异步接口通过完成队列处理更多并发任务,代价是取消、重试和状态管理更复杂。
不过,异步执行不会减少 D2H 和 I/O。系统应让 load 优先于低价值写入,避免写入持续占用 PCIe、host memory 带宽和存储带宽,反而拖慢命中请求。
数据传输
一次 load 可能依次经过存储、网络、host memory、PCIe 和 GPU memory。优化前先确认实际搬运多少数据、经过哪些链路,再选择具体的拷贝方式。
Cache 布局决定实际传输量
KV Cache 可能按 head 或序列分片,也可能在多个 ranks 上保存完整副本。这一差别决定 host memory 的容量、传输量和热点位置。
设模型有 个 decoder layers、 个 KV heads,每个 head 宽度为 ,KV 元素占 bytes。定义 是单个 rank 实际保存的 KV heads 数,恢复一段包含 个 token 的前缀时,单个 rank 需要搬运:
KV 完全按 TP 切分时,各 rank 保存的 KV heads 合计为 ,整个节点只有一份完整 KV。GQA 的 超过 后,框架通常会让每个 rank 至少保存一个 KV head。这样,各 rank 保存的 KV heads 合计会超过 ,同一个 KV head 也会出现在多个 ranks 上。
以 Qwen3-235B-A22B3 为例,、、、KV Cache 使用 BF16。 时,常见实现会让每个 rank 至少持有 1 个 KV head,于是每个 KV head 在 8 个 ranks 中各有两份。对于 32K prefix,单个 rank 需要搬运 1.47 GiB,整个节点合计 11.75 GiB,其中不重复的 Cache 只有 5.88 GiB。
CP 沿序列维同时切分 Attention 计算与 Cache。Prefill 可以通过 Ring Attention 交换 KV blocks,Decode 可以在各 ranks 计算局部 Attention output 和 LSE 后再归并。两种方案都让每个 rank 只保存和恢复自己的序列分片,整个 CP group 合计保留一份完整 Cache。
因此,D2H/H2D 要按各 ranks 实际保存的副本计算。 个完整副本会让节点搬运 份相同数据,序列分片只需搬运一份。
最慢链路决定带宽上限
从 L2 恢复 KV,数据通常要经过 host DRAM → PCIe → GPU HBM。从 L3 恢复时,前面还会增加存储或网络传输。单个 rank 受本地链路限制,多个 ranks 并发时还会共同争用 host memory、PCIe switch、CPU 上游链路或 NIC。因此,带宽估算既要使用对应方向和拓扑下的有效值,也要区分单卡上限与节点聚合上限。
下表给出各链路的带宽量级:
| 位置或链路 | 带宽口径 | 在 Cache 路径中的作用 |
|---|---|---|
| H200 HBM3e | 4.8 TB/s/GPU | L1 中 Attention 读取 Cache 的带宽上限 |
| GPU 间 NVLink | H200 900 GB/s/GPU,B300 1.8 TB/s/GPU | broadcast、AllGather 等 GPU 间数据传输 |
| PCIe Gen5 x16 | 单向理论约 64 GB/s,有效 H2D 可先按 50 GB/s 估算 | L2 与单卡之间的数据传输 |
| InfiniBand NDR | 单向理论约 50 GB/s/端口 | InfiniBand RDMA 访问远端 Cache |
| 400 GbE 上的 RoCE | 单向理论约 50 GB/s/端口 | Ethernet RDMA 访问远端 Cache |
RDMA 是访问方式,InfiniBand 和 RoCE 是承载网络。GPUDirect RDMA 允许 NIC 直接访问 GPU memory,省去 host memory 中转,但有效带宽仍受 NIC、PCIe 和拓扑限制4。
因此,拓扑必须和带宽一起看。host memory、GPU、NIC 或 NVMe 位于不同 CPU socket 或 PCIe 分支时,数据可能多经过一段 CPU 互连。nvidia-smi topo -m、topo -cpu 和 lspci -t 可以检查这些连接关系54。数据来自 L3 时还要处理延迟波动和失败。系统通常会设置超时,超时后放弃尚未到达的 KV,转而重算对应的 prefix。
执行方式决定固定开销与资源竞争
传输路径确定后,拷贝方式决定实际能接近多少带宽。描述一次 KV load,至少要同时记录总字节数、分段数量、单段大小、是否连续、对齐方式、传输方向和并发度。总字节数相同,几段大块连续数据与几千段离散 pages 的最佳实现可能完全不同。
host pool 通常使用 pinned memory,也叫 page-locked memory。它让异步 H2D/D2H 成为可能,GPU kernel 也可以直接访问映射后的 host memory6。不过,后者每次访问仍要经过 PCIe 或 NVLink C2C,延迟和带宽都不如 HBM,适合只读一次且访问连续的数据,不适合替代完整的 L1 Cache。
| 方式 | 主要执行资源 | 更适合的形状 | 主要代价 |
|---|---|---|---|
cudaMemcpyAsync | 通常使用 copy engine,但不保证 | 少量大块、连续的数据 | 多次小拷贝会累积 API 和调度开销 |
cudaMemcpyBatchAsync | 由 CUDA runtime 选择 copy engine 或 SM | 多段相互独立的拷贝 | 需要准备地址与大小数组,收益取决于分段 |
| CUDA copy kernel | SM | 离散 pages、gather/scatter、layout 转换 | 占用 SM 和显存带宽,可能干扰计算 |
NVIDIA 的 CUDA 文档明确说明,cudaMemcpy* 可能使用 copy engine,但不保证一定使用。当 copy engine 数量或带宽成为限制时,由 SM 执行的显式 copy kernel 可能更快。反过来,大块连续数据通常更适合 copy engine,因为它不占用 SM,也更容易与计算重叠7。cudaMemcpyBatchAsync 则把多段拷贝合并为一次 API 调用,用来摊薄准备开销,并允许 runtime 根据属性选择实现8。
SGLang HiCache 因此同时提供 direct 和 kernel 两种 I/O backend。page-first direct backend 会在运行环境支持时调用 cudaMemcpyBatchAsync,否则退回逐段拷贝(transfer.cu)。kernel backend 则启动 CUDA kernel,由 GPU threads 根据索引搬运 pages(transfer.cu)。page-first write-back 当前以 128 KiB 为阈值,较大的 page 才会尝试 batch memcpy(staged_write_back.cuh)。这个选择也说明传输 backend 要与数据 shape 匹配。
执行方式也决定传输能否与计算重叠。除了调用异步 API,还要使用不同的 non-default streams,建立正确的 event 依赖,并确保两项工作使用的硬件资源能够并发9。copy engine 是数量有限的 DMA 资源,多条 streams 只是提供并发调度机会,不会增加引擎数量。copy engine 更容易避开 SM 计算,copy kernel 则会与 Attention 和 GEMM 争用 SM 与显存带宽。
异步执行不会减少物理传输量,也无法隐藏尚未完成就必须使用的数据。把前面的数据量、路径和执行方式合在一起,请求实际感知的 load 延迟可以写成:
包含批量描述符准备、API 调用和 kernel launch 等固定开销。分段太小时,这部分占比会上升,应该合并或批量提交。单次传输过大又会推迟第一批 KV 到达 GPU,减少逐层加载与计算重叠的机会。单个 copy kernel 的带宽用于定位局部瓶颈,最终比较仍以端到端的 为准。
Load 与重算
已经包含 Cache 布局产生的物理数据量、各段传输、固定开销和排队。将它与重算比较时,可以直接沿用 Roofline model:计算强度等于 FLOPs 与通信字节数之比,硬件临界计算强度等于峰值算力与通信带宽之比。这里把通信对象换成需要恢复的 cache,把 HBM 或卡间带宽换成对应缓存路径的带宽即可10。
当 时,load 的理论耗时低于重算。简化估算可以使用 GPU 峰值算力和缓存路径的带宽估值,实测时则换成对应的实际值。L3 的 应取整条 L3→L2→L1 路径的端到端带宽。
以 Qwen3-235B-A22B、、8×H200 为例。记激活参数量为 ,Q heads 数为 ,重算 32K prefix 的 Prefill 计算量约为:
平均到每个 rank, PFLOPs, GiB,因此 FLOPs/byte。按单张 H200 的 BF16 dense 峰值 989 TFLOPS、pinned-memory H2D 有效带宽 50 GB/s 计算, FLOPs/byte。前者高于硬件临界值,所以 load 的理论耗时更低,约为 31 ms,而重算约为 390 ms。
换用其他 Attention 结构时,判断方法不变。MLA、DSA 等架构会压缩 ,也可能稀疏降低 。前者下降得更多,load 更有优势;后者下降得更多,重算更有优势。Roofline model 给出理论边界,生产环境还要比较包含排队、固定 I/O 延迟和传输重叠后的 与 。
上述临界条件针对单请求。并发服务中,后台 I/O 仍会占用 PCIe、host memory 带宽和 GPU 拷贝引擎。即使命中的单请求变快,争用也可能降低吞吐并推高 P99 TTFT。
综合分析:SGLang CP 的 Cache 去重
当前 SGLang Prefill CP 先对各 rank 的 token 分片执行 projection,随后 AllGather latent KV、RoPE 状态和 sparse Attention 的 indexer K。每个 rank 最终都保存完整序列,序列分片没有延续到 Cache 存储。在 、attn_cp_size=8、attn_tp_size=1 的配置下,L1 和 HiCache L2 都有 8 份完整 Cache。
临时方案:用 NVLink broadcast 替代重复 H2D
针对这部分重复,现有 Cache 布局可以采用三种恢复方案,也可以直接改为序列分片。表中的 D2H、H2D 表示每个逻辑 Cache page 搬运的份数。
| 方案 | 定位 | host 数据 | D2H | H2D | GPU 间数据传输 | 主要问题 |
|---|---|---|---|---|---|---|
| 每 rank 独立缓存 | 当前实现 | 无 | L2 容量与 PCIe 流量都重复 | |||
| 共享 host 副本,各 rank 完整读取 | L2 去重 | 无 | H2D 重复,多个 ranks 读取同一段 host DRAM | |||
| 固定 source rank + broadcast | 临时方案 | 一次 broadcast | source rank 与 broadcast 压力集中 | |||
| 按序列分片恢复 | 理想方案 | 分布式 Attention 通信 | 需要 Cache 布局与 Attention 后端配合 |
PR #26691 采用固定 source rank + broadcast 方案:host memory 只保存一份完整 Cache,source rank 执行一次 H2D,再通过 NVLink 将数据复制到其他 ranks。对照方案 shared-memory direct load 同样只保存一份 host Cache,但 8 个 CP ranks 会分别执行完整 H2D,节点合计经过 PCIe 搬运 8 份数据。
两种方案的 cache hit rate 都约为 74%。broadcast 同时把 H2D 从 8 份降到 1 份,并消除了 8 个 ranks 对同一 host buffer 的并发读取。平均 TTFT 从 10.2 秒降到 6.7 秒,请求总耗时从 499 秒降到 291 秒。端到端结果表明,一次 H2D 加 NVLink broadcast 比八次完整 H2D 更快。
这套临时方案也适用于多个 Attention replicas 加载同一个共享 prefix 的场景。
临时方案的单点压力与同步
固定 source rank 会把 host 传输和 broadcast 都集中到一张 GPU。H2D 占用它的 PCIe 带宽和 copy engine,broadcast 还要从它的 HBM 读取数据,再通过 NVLink/NVSwitch 发送给其他 GPU。所有 ranks 都要等待 broadcast 完成,source rank 的排队会直接拖慢整体。
可以为每个 block 生成在所有 ranks 上一致的随机 ID,再对 rank 数 取模:source_rank = block_id % g。block ID 均匀分布时,各 ranks 承担的 block 数也会接近。这个方法分散了 H2D 与 broadcast 任务,但 PCIe 和 GPU 互连的总通信量不变。它缓解的是单卡热点,共享互连的带宽压力仍然存在。
broadcast 还要求所有 ranks 以一致顺序进入 collective。该 PR 曾让 H2D 在 load stream 上执行,broadcast 却落到 default stream。不同 ranks 因而可能以不同顺序进入 cache broadcast 和 forward collective,两个 NCCL 通信组形成环形等待。修复后,H2D 与 broadcast 都在 load stream 上执行,forward 则等待该 stream 完成。
定位这类问题时,先确认各 rank 最后进入的 collective,再对齐它们使用的 NCCL 通信组和 CUDA stream。NCCL_DEBUG=TRACE 配合 NCCL_DEBUG_SUBSYS=CALL 可以记录 NCCL 调用顺序,Nsight Systems 可以还原 H2D、broadcast 和 forward 的 stream 时间线。将各 ranks 的日志按时间对齐,第一个 collective 顺序分叉的位置通常就是问题起点:
| rank 0 | rank 1 |
|---|---|
2026-09-03 10:18:42.481 [Rank 0] ncclBroadcast(comm=cache, stream=default) ENQUEUE2026-09-03 10:18:42.484 [Rank 0] ncclAllReduce(comm=tp, stream=forward) ENQUEUE | 2026-09-03 10:18:42.482 [Rank 1] ncclAllReduce(comm=tp, stream=forward) ENQUEUE2026-09-03 10:18:42.485 [Rank 1] ncclBroadcast(comm=cache, stream=default) ENQUEUE |
rank 0 先进入 cache broadcast,rank 1 却先进入 forward AllReduce。两边分别等待对方进入自己的 collective,请求因而无法继续。
理想方案:让序列分片贯穿 Cache
broadcast 减少了 L2 副本和重复 H2D,但每个 rank 仍然保存完整的 L1 Cache,每次恢复也要向所有 ranks 复制完整数据。它适合在不修改 Attention 后端的前提下快速消除主要冗余。
CP 的理想方案是让序列分片贯穿 L1、L2 和恢复过程。每个 rank 只保存并加载自己负责的序列分片,Prefill 通过 Ring Attention 交换 KV blocks,Decode 则归并各 rank 的局部 Attention output 和 LSE。这样可以同时消除 L1、L2 和 H2D 冗余,并将 host 传输分散到各 ranks。
如果现有 Attention kernel 必须读取完整 Cache,可以先让各 ranks 分片执行 H2D,再在 GPU 内 AllGather。这种过渡方案保留了完整 L1 Cache,但已经消除了重复 PCIe 传输。
总结
host memory 增加了 KV 的驻留时间,也在重算之外增加了一条 load 路径。是否使用 L2/L3,最终仍要回答三个问题:新增容量带来了多少复用,load 是否比重算更快,并发 I/O 是否降低了整体服务能力。Cache 布局、链路带宽和拷贝方式解释 load 的成本从哪里来,也为后续优化提供方向。
上线后可以从三个层级观察结果:
| 层级 | 核心问题 | 观测指标 |
|---|---|---|
| 复用 | 增加的容量是否换来了更多 KV 复用 | L1、L2、L3 命中的 token 数 |
| 单请求 | load 是否真的比重算更快 | 、、实际传输量、端到端有效带宽、传输与计算的重叠时间 |
| 服务整体 | 后台 I/O 是否引入共享瓶颈或单卡热点 | 吞吐、平均与 P99 TTFT、I/O 排队时间、各 rank 的传输字节数和耗时、host memory、PCIe 和 GPU 互连的带宽利用率 |
是否启用 L2/L3,最终应通过受控实验比较端到端 TTFT 与吞吐,再用这些指标定位收益和瓶颈。