推理系统基础设施手记（六）：HiSparse —— 把 HBM 当缓存用，长上下文服务的容量墙怎么拆

0. 一句话主线

前三篇（sys3/sys5）讲的是「怎么让 KV 变小、让每次读得更少」；这一篇讲的是更小之后仍然装不下的那一半问题——把 GPU HBM 从「KV 的居住地」降级成「KV 的缓存」。

HiSparse（Scaling Sparse-Attention Decoding with Hierarchical KV Cache Management，Stanford MAST Lab，arXiv:2608.07009，2026-08）的核心主张只有一句：既然 top-k 稀疏注意力每步只读 k 个 KV 条目，那 HBM 里凭什么要常驻 L_ctx 个？ 把全量 KV 挪到主机 DRAM，GPU 只留一个固定大小的热缓存，再配一个融合 CUDA 核把「命中检测 + LRU 替换 + 主机抓取」一次做完——就能在输出逐 token 完全不变的前提下，把长上下文的并发度拉上去。

它已经合入 SGLang 上游，在 H200 / B200 / GH200 上跑通 DSA、NSA、Quest 三类稀疏注意力，长上下文峰值生成吞吐最高 4.7×。

1. 容量墙：稀疏注意力省了算力，却没省显存

先把问题摆清楚。Top-k 稀疏注意力（NSA、DSA、Quest 等）的卖点是计算便宜：第 ℓ 层、第 t 步只挑出 k 个历史位置 S_t^(ℓ) 来做注意力，k 通常是几千，而不是整个上下文长度 L_ctx。

但服务系统有个隐含假设拖了后腿：

任意一个过去的 token，都可能在未来某一步被索引器选中。

为了不漏，系统干脆把整个 KV cache 常驻 GPU HBM。于是显存账单依旧随 L_ctx 线性增长——算力还没用完，显存先没了。这就是论文说的 capacity wall。

数字很直白：

场景	KV 显存占用	后果
GLM-5.1，单请求 128K 上下文	≈ 13.09 GB	80GB 卡只能并发个位数请求
单请求 1M 上下文	超过整卡 HBM	根本无法服务

注意这里的尴尬：计算侧已经优化到只读 k 条了，内存侧却还按 L_ctx 全量付费。HiSparse 要拆的就是这个错配——让 HBM 占用与 k 成比例，而不是与 L_ctx 成比例。

难点在于稀疏选择是动态的：S_t^(ℓ) 每一步、每一层都在变，你没法像 MLA 那样静态地”压缩完就固定”。

2. 两级内存层次：主机存权威副本，GPU 只留热缓存

HiSparse 的解法是不碰模型逻辑，只改 KV 记录放在哪：

关键不变式只改 KV 的物理存放位置，不改模型结构、不改计算 → 输出逐 token 不变（exact）代价只有一个：host ↔ device 的 IO 带宽

图 1：HiSparse 的两级层次。主机 DRAM 是「权威副本」，GPU HBM 退化成一个由 RESOLVE 核管理的热缓存。

三个部件各司其职：

主机 KV 池：prefill 阶段生成的 KV 记录直接写进固定的主机内存，这是唯一权威副本。
GPU 热缓存：为每个「请求 × 层」保留 B 个槽位，装最近被选中的 KV 记录。为保证当前注意力一定算得动，约束 B ≥ k。
元数据：GPU 上维护紧凑页表，把逻辑 token 位置映射到物理槽位，或标记为「仅主机」；同时维护 LRU 所需的最近性信息。

于是单请求 HBM 消耗变成：

HBM 占用 = N_ℓ × B × W_KV × s

N_ℓ 是层数，W_KV 是每 token 的 KV 元素数，s 是每元素字节数。L_ctx 消失了——这就是”把解码吞吐与显存容量解耦”的实质。

2.1 槽位到底怎么分（DRAM 与 HBM 的对应关系）

槽位分配：DRAM 按 token 全量存，HBM 按「请求 × 层」各给 B 槽

① 主机 DRAM（权威全量副本） request 0 request 1 request 2 每个 (request, layer, pos) 的 K/V 全量保留 prefill 直接写入，永不淘汰、永不做 LRU 容量随主机内存走，128K 请求约 13 GB 也毫无压力按逻辑位置连续寻址，是唯一权威副本 ② 页表逻辑位置 → 物理槽 pos→slot pos→slot pos→slot pos→slot 命中 = 槽里有未命中 = 仅主机 ③ GPU HBM（热缓存 · 固定 B 槽）每个「请求 × 层」一组，B ≥ k L0 17 9 44 3 21 8 L1 31 17 6 52 12 40 L2 5 28 17 63 1 35 L3 22 47 14 17 59 30 图中 B = 6 仅示意；实际 B 为 k 的数倍（几千量级）绿 = 本步命中的热槽，红 = LRU 选中的 victim（下一个 miss 顶掉它）同一逻辑位置 17 可在多层各占一槽：分配粒度是「请求 × 层」每槽 = 1 条 KV 记录 + 逻辑位置 + LRU 最近性位查 HBM 占用 = N_ℓ × B × W_KV × s —— 与 L_ctx 无关分配粒度是「请求 × 层」而不是「请求」：不同层的索引器选择不同，各自要有一组槽。约束 B ≥ k：保证当前步选出的 k 条一定放得下，注意力永远算得动。算例：128K 上下文、k = 2048、取 B = 2k = 4096。单层常驻 4096 条 vs 全量 131072 条 ≈ 3.1%，约 32× 节省（与论文实测"最高 30×"吻合）。上下文再涨到 1M，B 仍是 4096，HBM 占用不增——这才叫解耦。

图 3：DRAM 与 HBM 的分配对应关系。DRAM 侧按 (request, layer, pos) 连续全量存放；HBM 侧给每个「请求 × 层」分配一组固定 B 槽，靠页表做逻辑位置到物理槽的映射，LRU 决定谁被顶掉。

三个要点值得单独记住：

分配粒度是「请求 × 层」，不是「请求」。不同层的索引器选择不同，所以每层都要有一组自己的槽——这也是为什么公式里要乘 N_ℓ。
DRAM 与 HBM 解耦靠页表。DRAM 侧按逻辑位置连续寻址；HBM 侧的槽位与逻辑位置无关，装的是”最近被选中的那几条”，映射关系全在页表里。
B 是常数，不随上下文长。这正是容量墙被拆掉的地方：上下文从 128K 涨到 1M，DRAM 侧线性增长（无所谓，主机内存大），HBM 侧一动不动。

3. RESOLVE 融合核：五个阶段一次 launch 做完

分层缓存的想法不新，难的是怎么让 miss 的代价低到可以接受。稀疏选择的访存是零散的（scattered），传统 CPU 侧分页搬运会直接把延迟暴露出来。

HiSparse 的做法是写一个叫 RESOLVE 的单一融合 CUDA 核，每个稀疏层启动一次，在 GPU 线程里并行做完五件事：

阶段	做什么	关键点
暂存 Stage	把索引器选出的逻辑位置载入共享内存哈希表	后续比对全在片上，不碰全局内存
标记 Mark	对照哈希表检查 GPU 缓存槽，识别”命中”和”可驱逐”槽	命中判定是纯元数据操作
扫描 Scan	并行扫描槽位、更新 LRU 元数据，为当前步的 miss 挑 victim	并行前缀和式的槽位分配
抓取 Fetch	miss 线程用矢量化非一致性加载（`ld.global.nc.v2.b64`）把 KV 从 pinned 主机内存拉进分配到的 GPU 槽	GPU-assisted IO：GPU 线程自己去拉
发布 Publish	更新页表，把物理设备偏移交给稀疏注意力核	下游核无感知

“GPU 辅助 IO”是这篇论文真正的技术贡献：让 GPU 线程直接从主机内存取数，即使是稀疏选择这种零散访存，也能吃满 PCIe / NVLink 带宽——绕开了”CPU 发现缺页 → 通知 → 再搬运”这条长路径。而且整个流程在 decode CUDA Graph 内部完成，不会因为动态控制流破坏图捕获。

4. 局部性：为什么一个小小的 B 就够用

任何缓存都靠局部性吃饭。HiSparse 的实测发现，稀疏注意力的选择序列有两个可利用的性质：

时间局部性：连续解码步骤会重复选中大量相同的 token；
跨层相关性：不同层往往关注上下文里相近的区域。

用 GLM-5.1（k = 2048）跑出来的 miss rate 很能说明问题：

GLM-5.1 · k = 2048：top-k miss rate 随缓存放大快速下降缓存从 k 放大到 2k，miss 直接腰斩；放大到 4k 再腰斩一次。B 只需几倍 k，不必随上下文增长。

图 2：LRU 管理下的局部性红利。缓存只要几倍于 k，就能把 miss 压到个位数百分比。

精确逐层预取

还有一招更狠的。有些模型（如 GLM-5.2）在多组层之间共享索引器输出——一旦某个”锚点层”确定了它的选择，系统立刻就能知道后面那些共享层的选择是什么。

HiSparse 据此做 exact layer-wise prefetching：

锚点层算出选择 → 同步推出后续共享层的 miss 计划；
一个后台纯复制核重放这些未来层的 miss 计划；
主机到设备的传输与当前层的计算重叠。

结果：隐藏掉约一半的剩余 IO 延迟。注意这里用的是”精确”（exact）——不是启发式猜，是确定性地知道，所以不会引入正确性风险。

5. 为什么它是「精确」且「与索引器无关」

这两点是 HiSparse 能直接进生产的关键：

精确（exact）：只改变 KV 记录的物理放置，不碰模型结构、不碰注意力计算、不做近似召回。所以模型输出逐 token 完全不变——这是和”用近似检索换显存”那类方案的根本分野。
与索引器无关（indexer-agnostic）：它不关心你是怎么挑出那 k 个位置的。论文在 DSA、NSA、Quest 三类稀疏注意力上都做了评估，都能吃。

换句话说，HiSparse 是稀疏注意力的通用内存后端——只要你的模型已经有稀疏索引器，它就能接上。

6. 实测效果

配置	结果
DeepSeek-V4-Flash（NSA），双 B200，64 并发	生成吞吐 2.1×
Qwen3 + Quest，GH200，20 万输入长度	生成吞吐最高 4.7×
高负载场景（prefill / decode 共享 GPU）	TTFT 显著下降（不再因 HBM 耗尽阻塞新请求）
单 token 延迟	与基线相当（不是更快，但没变慢）
no-IO oracle 实验	解析机制本身无可测量的 per-token 开销

最后一行很重要：它说明 host-device IO 是这套方案唯一的、也是全部的代价。所以这条链路的带宽直接决定收益：

链路敏感性：GH200（NVLink-C2C）上的 KV 抓取耗时比 H200（PCIe Gen5）少近 4 倍，因而可以配更小的 GPU 缓存、跑更高的并发。

7. 什么时候不该用它（局限与前提）

这部分比效果更值得记：

前提是主机 DRAM 远大于 GPU HBM。在 Grace 系的 GB200 / GB300 上，CPU 与 GPU 共享统一内存、主机侧并无容量优势——这套”卸载换容量”的逻辑不成立。论文自己把这点列为根本性限制。
依赖 GPU-assisted IO 对零散访存也能接近链路带宽。这一点借自作者的 Strata 工作，本文未独立 benchmark。PCIe Gen5 上 miss 代价明显高于 NVLink-C2C。
换来的是并发度与容量，不是单请求速度。per-token latency 只是”可比”，想让单个请求更快它帮不上忙。
对稠密注意力模型无效。必须先有稀疏索引器；没有 top-k 选择，就没有”只需常驻 k 条”的立足点。
收益随负载而变：低并发、短上下文场景下，基线根本没撞到容量墙，分层只会平白添一层 IO。

8. 一句话串起来

把 sys3 / sys5 和本篇叠起来看，KV 这条主线就完整了：

MLA 是「每个 token 存得更小」，NSA / DSA / Quest 是「每次只读重要的 k 条」，CSA 是「先把 token 压小再做稀疏读取」，HCA 是「把很长的历史压成很短的摘要后全读」——而 HiCache 与 HiSparse 是在 KV 已经产生之后，解决「放在哪一层存储」的问题。

再往前一步区分二者：

HiCache 解决”KV 放 GPU、CPU 还是更低层存储”，是存储层级的问题；
HiSparse 解决”稀疏解码下 GPU 里到底需要常驻多少”，是容量账单与 k 而非 L_ctx 挂钩的问题。

减计算（稀疏注意力）与减容量（压缩 + 分层）是两条正交的线，HiSparse 站在后者的下游，也是目前最接近生产落地的那一环（已在 SGLang 上游）。

9. 与已有文章的衔接

KV Cache 压缩方向全景（MHA → HiCache） → 见 sys5-attention-evolution-kvcache
MSA / CSA / HCA 三种注意力改造路线 → 见 sys3-attention-evolution
HiCache 与 NUMA × PCIe × NIC 数据链路 → 见 sys1-pcie-numa-nic-hicache
SGLang 原理（前缀复用与吞吐优化） → 见 fw3-sglang-internals
HiCache / Mooncake 的 KV 复用与组合正确性 → 见 tr2-vllm-sglang-20260807

附：论文信息

HiSparse: Scaling Sparse-Attention Decoding with Hierarchical KV Cache Management
Zhiqiang Xie, Zhangheng Huang, Tingwei Huang, Ziyi Xu, Ruiyang Ma, Christos Kozyrakis
Stanford MAST Lab · 2026-08-07 · arXiv:2608.07009
代码状态：已合并至上游 SGLang

本篇为论文精读与工程解读，不构成任何技术选型或投资建议。数据引自原论文与 alphaxiv 中文解读页（2026-09-07 检索）。