KVCache offloading 系统的 CPU 三大瓶颈(细粒度缓存管理、PCIe 带宽浪费、CPU-centric 同步开销)通过算法-系统协同解决:head 粒度近似缓存 + 零拷贝传输引擎 + GPU 中心同步,解码吞吐提升 9.3%–66.6%,精度几乎无损。
长序列 LLM 推理中 KVCache 被 offload 到 CPU 内存后,现有系统(RetroInfer、InfiniGen、PQCache)引入三个 CPU 端瓶颈,导致实际解码延迟远超理想值(全 KVCache 驻留 GPU HBM):
cudaStreamSynchronize 阻塞 CPU kernel launch 线程,产生数十至数百微秒空闲。
Paper Figure 2:各系统单层解码延迟分解(BSZ=1, SeqLen=128K, top-k=10%, Llama3-8B-1048K)。RetroInfer 延迟为理想值的 137%,InfiniGen 337%,PQCache 354%。
Figure 2 量化了三大瓶颈的相对权重:RetroInfer 的性能损失主要来自缓存管理(橙色),InfiniGen 主要来自未隐藏的数据传输(蓝色),PQCache 则两者叠加。这一分解直接驱动了 CLO 的设计优先级。
CLO 的核心策略是 牺牲缓存命中率换取管理零开销,再用系统级优化补偿 miss 代价:
核心技术壁垒:Head 粒度缓存 + query 余弦相似度判定的组合看似简单,但其可行性依赖于一个非显然的实证发现——相邻步 query 在绝大多数 head 上的余弦相似度高到足以支撑 ~79% 的 head 级缓存命中率。这个 observation 需要跨模型、跨序列长度的系统性验证,且阈值设计(Eq. 7–9)需要与 GQA 聚合(Eq. 10)和 reuse difficulty 判定(Eq. 11–13)配合才能在精度-效率间取得平衡。单独复制其中任一环节均无法复现整体效果。

Paper Figure 8:CLO 系统总览。三大组件——Cache Manager、Transfer Engine、Data Synchronization Controller——围绕 head 粒度缓存协同工作。
CLO 的三个核心模块清晰分离职责:Cache Manager 管理 KVCache 的分区(persistent vs. offloaded)、similarity cache 的 lookup/update、以及 top-$k$ retrieval(集成 HATA);Transfer Engine 通过 GDRCopy 零拷贝通道执行 CPU→GPU 传输;Data Synchronization Controller 通过 GPU polling 共享内存变量协调数据就绪与 kernel 执行。

Paper Figure 9:解码阶段推理流程。展示 layer $l$ 的 KV 数据如何在 layer $l-1$ 执行期间被预取和准备。
流程的关键是 一层提前量:在 layer $l-1$ 开始时用 $q_l^{\text{approx}} = x_{l-1} W_Q^l$ 近似下层 query → cache lookup → 对 miss head 发起 prefetch → layer $l-1$ 的 attention + FFN 计算期间传输完成 → layer $l$ 执行时从 cache buffer + persistent cache 双源取 KV。

Paper Figure 7:Query similarity-based approximate cache 的 lookup / update 流程。
与传统 LRU/LFU 缓存的关键区别:CLO 的 cache lookup 是一次 GPU tensor core 上的余弦相似度计算($q_i^t$ vs. $q_i^{\text{label}}$),update 只需替换一个 query vector + 写入新 KV 数据到最终位置(无需 temp buffer → merge 的二次拷贝)。

Paper Figure 1:RetroInfer 的 block 粒度 LRU 缓存流程。对比 CLO 的 head 粒度方案,可见三阶段(lookup → access → update)中每一步都需要 CPU 介入。
Table 1 总结了核心差异:RetroInfer/PQCache 的 cache metadata 是 CPU 上的链表(高开销、需额外 buffer merge),CLO 的 metadata 是 GPU 上的 query vectors(开销可忽略、无额外拷贝)。
| 符号 | 定义 | 来源 |
|---|---|---|
| $q_i^t$ | Layer $l$ head $i$ 在步 $t$ 的 query vector | §3.1 |
| $q_i^{\text{label}}$ | Head $i$ 缓存的 query label | §3.1 |
| $\tau_i$ | Head $i$ 的余弦相似度阈值 | Eq. 9 |
| $s_i$ | Head $i$ 的 importance score(DuoAttention) | §3.2 |
| $\lambda_i$ | $s_i^n$,importance-weighted coefficient | Eq. 8 |
| $\eta$ | 全局相似度阈值上界(实验中 = 0.8) | Eq. 7 |
| $n$ | 重要性指数(实验中 = 3) | §3.2 |
| $D_i$ | KV head $i$ 的复用困难度 | Eq. 11 |
| $\hat{s}_i$ | Offline profiled 平均余弦相似度 | Eq. 11 |
| $\epsilon$ | 容差余量(Llama3: 0.1, Qwen2.5: 0.05) | §6.1 |
| $N_p$ | 每层可 prefetch 的 head 数 | Eq. 12 |
| $N_{\text{persist}}^l$ | Layer $l$ 需常驻 HBM 的 head 数 | Eq. 13 |
| $T_{\text{comp}}$ | 单层 GPU 计算延迟 | Eq. 12 |
| $B_{\text{PCIe}}$ | PCIe 峰值带宽 | Eq. 12 |
| $\text{mem}_{\text{head}}$ | 单 head top-$k$ KV 数据量 | Eq. 12 |
| # | 检查项 | 结果 | ||||||
|---|---|---|---|---|---|---|---|---|
| 1 | Eq. 6 的 $\ | q\ | $ 不变性声明 | ✓ 标准线性代数性质。qk-score 对所有 k 的排序确实与 $\ | q\ | $ 无关,因 $\ | q\ | $ 是公共因子。 |
| 2 | Eq. 7–9 在极端 $s_i$ 处的行为 | ✓ $s_i = 1 \Rightarrow \lambda_i = 1 \Rightarrow \theta_i = \theta^* \Rightarrow \tau_i = \eta = 0.8$(最严格)。$s_i = 0 \Rightarrow \lambda_i = 0 \Rightarrow \theta_i = \pi \Rightarrow \tau_i = -1$(永远命中)。边界行为符合设计意图。 | ||||||
| 3 | Eq. 10 的两条规则 | ✓ 规则 1:importance 越大 → 权重越大 → 该 head 的 sim 贡献越多。规则 2:所有 sim 相同时,加权平均退化为该值本身(分子 = sim·Σs_i,分母 = Σs_i)。 | ||||||
| 4 | Eq. 12 的量纲一致性 | ✓ $T_{\text{comp}}$(秒)/ ($B_{\text{PCIe}}$(字节/秒)· $\text{mem}_{\text{head}}$(字节))= 无量纲整数。 | ||||||
| 5 | Eq. 13 非负性 | ✓ $\max(\cdot, 0)$ 保证当 prefetch 能力足够时 persistent head 数为 0。 | ||||||
| 6 | 缓存命中率与管理开销的 tradeoff 可行性 | ✓ CLO 命中率 ~79%(vs RetroInfer ~91%),但管理开销从 36.2% 降至可忽略(<1%)。净效果:单层延迟降低 2.1–5.2×。命中率损失被管理开销消除 + prefetch + persistent cache 三重补偿。 |

Paper Figure 10:端到端 prefill + decoding 吞吐。CLO 在所有配置下解码性能领先,短序列下 FullAttn 仍占优(attention 非瓶颈时 top-k retrieval 反成开销)。
关键观察:(1) CLO 的 prefill 吞吐几乎匹配 FullAttn(HATA 检索元数据构建开销极低);(2) 解码吞吐随 batch size 增长,CLO 始终保持优势;(3) InfiniGen 在所有配置下最慢,因为大量 PCIe 传输无法被完全隐藏。
| 方法 | LongBench | RULER 16K | RULER 32K | RULER 64K | RULER 128K |
|---|---|---|---|---|---|
| Qwen2.5-14B | |||||
| FullAttn | 53.34 | 94.35 | 94.48 | 92.29 | 88.85 |
| RetroInfer | 54.71 | 94.73 | 94.41 | 92.37 | 89.49 |
| InfiniGen | 52.74 | 92.78 | 92.19 | 89.74 | 85.26 |
| CLO + HATA | 53.15 | 94.36 | 94.22 | 92.35 | 88.94 |
| Llama3-8B | |||||
| FullAttn | 41.06 | 86.07 | 80.80 | 76.26 | 72.96 |
| RetroInfer | 41.11 | 86.36 | 80.64 | 76.26 | 72.73 |
| InfiniGen | 40.22 | 79.70 | 76.76 | 72.96 | 69.37 |
| CLO + HATA | 41.15 | 85.65 | 80.70 | 75.94 | 73.44 |
CLO 的精度损失在所有配置下 ≤ 0.42(相对 FullAttn),与 RetroInfer 基本持平,显著优于 InfiniGen。这验证了 head 粒度近似缓存在 importance-aware 阈值控制下不引入实质精度退化。

Paper Figure 11:单层解码延迟分解(BSZ=1, SeqLen=128K)。CLO 在 Llama3 上实现 2.1–5.2× 加速,Qwen2.5 上 1.6–4.2×。
分解揭示 CLO 的优势来源:(1) 缓存管理开销几乎为零(vs RetroInfer 26.9–36.2%);(2) 数据传输成本低(speculative prefetch 隐藏 miss 传输);(3) kernel launch 开销可忽略(GPU-centric sync vs RetroInfer 114.5 µs, InfiniGen 273.2 µs)。
| 方法 | Llama3 32K | 64K | 128K | 256K | 512K | Qwen2.5 16K | 32K | 64K | 128K | 256K |
|---|---|---|---|---|---|---|---|---|---|---|
| RetroInfer | 1.20 | 2.40 | 4.80 | 9.56 | 19.12 | 0.90 | 1.79 | 3.59 | 7.20 | 14.34 |
| CLO | 0.53 | 1.09 | 2.62 | 8.19 | 14.09 | 0.37 | 0.74 | 1.49 | 3.44 | 8.03 |
(单位:GB)CLO 平均仅使用 RetroInfer 47.6% 的 GPU 缓存内存,因为 head 粒度缓存无需为每个 block 分配独立 metadata/buffer 空间。
| 方法 | Llama3 8K (BSZ 4/8/16/32/64) | Qwen2.5 8K (BSZ 2/4/8/16/32) |
|---|---|---|
| RetroInfer | 90.41/90.15/90.08/90.07/90.05 | 92.65/91.33/91.42/91.34/91.35 |
| CLO | 79.54/74.69/76.69/84.99/80.00 | 80.62/80.65/80.65/81.53/80.75 |
CLO 命中率平均 79.22%(vs RetroInfer ~91%),但 RetroInfer 的高命中率被其 CPU 管理开销抵消。CLO 通过 speculative prefetch 隐藏 miss 传输,净效果更优。

Paper Figure 3:InfiniGen 和 PQCache 在不同传输量下的 PCIe 实际带宽,始终低于峰值。原因:CPU 侧 gather 操作(先聚合分散 KV 再传输)引入额外开销。
CLO 的零拷贝引擎消除了 CPU gather 步骤。4 个 CPU 线程使用 AVX SIMD + cache prefetch 即可逼近 GDRCopy 峰值 21.21 GB/s,且跨序列长度(4K–128K)保持稳定。对比之下,InfiniGen/PQCache 即使使用 64 CPU 线程仍大幅低于峰值。
消融实验(+T = adaptive threshold, +G = GQA aggregation, +P = persistent caching):
| 配置 | Llama3 延迟降幅 | Qwen2.5 延迟降幅 |
|---|---|---|
| Baseline → +T | 35.2% | 30.7% |
| +T → +T+G | 48.0% | 43.2% |
| +T+G → +T+G+P | 56.9% | 47.5% |
三个组件贡献递增且互补。Adaptive threshold 通过差异化阈值提升总体命中率;GQA aggregation 避免 group 内保守取 min 导致的过频繁 miss;persistent caching 消除 hard-to-reuse head 的传输等待。
| 步骤 | 论点 | 依据 | 作用 |
|---|---|---|---|
| 1 | 长序列 LLM 推理中 KVCache 远超 GPU HBM 容量,必须 offload 到 CPU 内存 | Qwen2.5-14B 512K 序列 → 93.75 GB KVCache,3.3× 模型权重(§1) | 确立 offloading 必要性 |
| 2 | 现有 offloading 系统集成 top-$k$ attention + GPU caching / prefetching,但三个 CPU 瓶颈导致延迟远超理想值(137%–354%) | Figure 2 单层延迟分解;RetroInfer 缓存管理占 36.2%,InfiniGen 传输占 62.2%(§2.3) | 诊断:CPU 是被忽视的瓶颈 |
| 3 | 相邻解码步的 query 向量余弦相似度高 → top-$k$ 集合高度重叠 → head 粒度缓存复用可行 | Eq. 6(qk-score 排序与 ‖q‖ 无关)+ Figure 4/5/6(实证分布)(§3.1) | 算法基础:缓存复用的理论与实证支撑 |
| 4 | Head-wise 近似缓存将管理开销从链表遍历降至单次余弦计算,消除 CPU 瓶颈 1 | Table 1 对比(CLO: GPU 上 query vectors vs RetroInfer: CPU 上 LRU linked list)(§3.1) | 解决瓶颈 1 |
| 5 | Importance-aware 阈值 + GQA 聚合保持精度,同时最大化缓存复用率 | Eq. 7–10 + ablation(+T 降 35%, +T+G 再降 48%)(§3.2, §6.5) | 精度-效率平衡 |
| 6 | Hard-to-reuse heads 通过 selective residency(persistent cache)+ speculative prefetch 补偿 | Eq. 11–13 量化 reuse difficulty + prefetch 能力;+T+G+P 再降 56.9%(§3.3, §6.5) | 解决 cache miss 的性能代价 |
| 7 | 零拷贝传输引擎消除 CPU gather,4 线程饱和 PCIe 4.0 | GDRCopy + AVX SIMD;Figure 13 带宽稳定逼近 21.21 GB/s(§4.3, §6.5) | 解决瓶颈 2 |
| 8 | GPU-centric 同步消除 kernel launch overhead | UVA shared memory polling 替代 cudaStreamSynchronize(§4.4) | 解决瓶颈 3 |
| 9 | 端到端验证:9.3%–66.6% 吞吐提升,精度 ≤ 0.42 损失 | Figure 10(吞吐)+ Table 5(精度)跨两个模型、多个序列长度/batch size(§6) | 全系统验证 |
CLO 基于 FlashInfer + HuggingFace Transformers 实现:3417 行 CUDA/C++ + 1555 行 Python(§5)。
[实现未公开] — 论文未公开代码仓库。实现基于 FlashInfer(github.com/flashinfer-ai/flashinfer)和 Transformers(github.com/huggingface/transformers),但 CLO 自身的 CUDA/C++ 和 Python 代码未开源。
| 维度 | CLO 的定位 |
|---|---|
| 服务阶段 | Prefill + Decode 均覆盖,decode 为核心优化目标 |
| 并发模式 | 单请求(BSZ=1)到中等并发(BSZ≤64)均有优化 |
| 硬件亲和性 | PCIe 4.0 GPU(非 NVLink 互连);CPU 负载极轻(仅 4 线程);受益于 GDRCopy(需 BAR 映射支持) |
| 生态集成 | 基于 FlashInfer + Transformers,非 vLLM/SGLang 插件,需独立部署 |
| 运行模式 | 单机单 GPU(scale-up),与分布式 KVCache(scale-out)正交 |
| 工作负载 | CLO | FullAttn | RetroInfer | 原因 |
|---|---|---|---|---|
| 短序列 (≤8K), 低并发 | 次优 | 最佳 | 中等 | Attention 非瓶颈,top-k retrieval 反成开销 |
| 长序列 (≥64K), BSZ=1 | 最佳 | OOM | 中等 | KVCache 超出 HBM,CLO 的 CPU-light 策略优势最大 |
| 中序列 (8K–64K), 高 BSZ | 最佳 | 受限 | 中等 | KVCache × BSZ 增长快,CLO 缓存内存效率 (47.6%) 释放更多 BSZ 空间 |
对于 128K 序列 BSZ=1(FP16):