系列:Inference Systems Infrastructure Notes

Inference Systems Infrastructure Notes (6): HiSparse — Treating HBM as a Cache to Break the Long-Context Capacity Wall

0. The One-Line Thread

The earlier posts (sys3 / sys5) asked “how do we make KV smaller and read less of it”; this one tackles the half that is still too big even after shrinking — demoting GPU HBM from “where KV lives” to “a cache over KV.”

HiSparse (Scaling Sparse-Attention Decoding with Hierarchical KV Cache Management, Stanford MAST Lab, arXiv:2608.07009, Aug 2026) makes a single argument: if top-k sparse attention reads only k KV entries per step, why should HBM hold L_ctx of them? Move the full KV history to host DRAM, keep a small fixed-size hot cache on the GPU, and add one fused CUDA kernel that does hit detection, LRU replacement and host-to-device fetch in a single launch — and you raise long-context concurrency while keeping the output bit-for-bit identical, token by token.

It is merged into upstream SGLang and evaluated across three sparse-attention families (DSA, NSA, Quest) on H200, B200 and GH200, with up to 4.7x higher peak generation throughput on long-context workloads.


1. Capacity Wall: Sparse Attention Saved FLOPs, Not Memory

Start with the problem. Top-k sparse attention (NSA, DSA, Quest, and friends) sells cheap compute: at layer ℓ and step t it selects only k historical positions S_t^(ℓ) — typically a few thousand, not the full context length L_ctx.

But serving stacks carry a hidden assumption that drags everything back:

Any past token may be selected by the indexer at some future step.

To avoid missing one, systems simply keep the entire KV cache resident in GPU HBM. The memory bill still grows linearly with L_ctx — so memory runs out long before compute does. That is the capacity wall the paper names.

The numbers are blunt:

ScenarioKV memory footprintConsequence
GLM-5.1, single 128K-context request≈ 13.09 GBAn 80GB card serves only a handful concurrently
Single 1M-context requestExceeds total HBMCannot be served at all

Note the mismatch: the compute side already reads only k entries, while the memory side still pays for all of L_ctx. That is exactly what HiSparse removes — HBM footprint should scale with k, not with L_ctx.

The hard part is that sparse selection is dynamic: S_t^(ℓ) changes every step and every layer, so you cannot statically “compress it once and freeze it” the way MLA does.


2. Two-Level Hierarchy: Authoritative Copy on Host, Hot Cache on GPU

HiSparse never touches model logic — it only changes where KV records live:

(1) Host KV pool (Host DRAM, pinned) Authoritative full KV history per request, written at prefill Capacity tracks host memory, on the order of L_ctx GLM-5.1 128K single request is about 13.09 GB: trivial for host RAM (2) GPU hot cache (HBM, fixed B slots) One per request-layer pair, B >= k, holds recently selected KV HBM usage = N_l x B x W_KV x s, decoupled from L_ctx Up to 30x HBM savings for GLM-5.1 at 128K (3) RESOLVE fused CUDA kernel (inside decode graph) Stage, Mark, Scan, Fetch, Publish in one launch GPU threads pull KV from pinned host memory via ld.global.nc.v2.b64 miss fetch write back / page table (4) Sparse attention kernel Receives physical device offsets and computes as usual, reading exactly what full residency would

Key invariant Only physical placement changes: no model or math change, so output is unchanged (exact) The one and only price: host-to-device IO bandwidth

Figure 1: The two-level hierarchy. Host DRAM holds the authoritative copy; GPU HBM degrades into a hot cache managed by RESOLVE.

Three components, three jobs:

  1. Host KV pool — KV records produced during prefill go straight into pinned host memory. This is the only authoritative copy.
  2. GPU hot cache — B slots per request-layer pair holding recently selected KV records. To guarantee the current attention can always proceed, B ≥ k.
  3. Metadata — compact page tables on the GPU map a logical token position to a physical slot or mark it host-only, plus recency bits for LRU.

Per-request HBM consumption becomes:

HBM usage = N_l x B x W_KV x s

where N_l is layer count, W_KV is KV elements per token, and s is bytes per element. L_ctx is gone — that is what “decoupling decoding throughput from GPU memory capacity” actually means.

2.1 How Slots Are Actually Allocated (DRAM to HBM Mapping)

Slot allocation: DRAM stores everything per token; HBM grants B slots per request-layer pair

(1) Host DRAM (authoritative full copy) request 0 request 1 request 2 Every (request, layer, pos) K/V pair is retained Written directly at prefill; never evicted, never LRU'd Capacity follows host RAM: a 13 GB, 128K request is fine Addressed contiguously by logical position; the only authority (2) Page table logical pos to slot pos to slot pos to slot pos to slot pos to slot hit means in slot miss means host only (3) GPU HBM (hot cache, fixed B slots) One set per request-layer pair, with 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 here is illustrative; real B is a few multiples of k (thousands) green = hot slot hit this step, red = LRU victim (next miss evicts it) position 17 can occupy a slot in every layer: granularity is request x layer Each slot = 1 KV record + logical position + LRU recency bits lookup HBM usage = N_l x B x W_KV x s, independent of L_ctx Allocation granularity is request x layer, not request: different layers pick different entries, so each needs its own slots. The constraint B >= k guarantees the k entries selected this step always fit, so attention can always proceed. Example: 128K context, k = 2048, take B = 2k = 4096. Per layer that is 4096 resident entries instead of 131072, about 3.1%, roughly a 32x saving (matching the 30x the paper reports). Push the context to 1M and B stays 4096: HBM usage does not grow. That is what decoupling means.

Figure 3: How DRAM and HBM allocation correspond. The DRAM side stores everything contiguously by (request, layer, pos); the HBM side grants each request-layer pair a fixed set of B slots, with a page table mapping logical positions to physical slots and LRU deciding who gets evicted.

Three points worth memorizing on their own:

  1. Allocation granularity is request x layer, not request. Different layers select different entries, so each layer needs its own set of slots — which is why N_l appears in the formula.
  2. DRAM and HBM are decoupled by the page table. The DRAM side is addressed contiguously by logical position; HBM slots have no fixed relation to logical position and simply hold “the recently selected ones,” with all mapping in the page table.
  3. B is a constant, independent of context length. This is exactly where the capacity wall comes down: as context grows from 128K to 1M, the DRAM side grows linearly (fine, host memory is large) while the HBM side does not move at all.

3. RESOLVE: Five Stages in One Launch

Hierarchical caching is an old idea. The hard part is making a miss cheap enough to tolerate, because sparse selection access patterns are scattered, and a conventional CPU-side paging path exposes the full latency.

HiSparse writes a single fused CUDA kernel named RESOLVE, launched once per sparse layer, doing five things in parallel across GPU threads:

StageWhat it doesWhy it matters
StageLoads the indexer-selected logical positions into a shared-memory hash tableLater comparisons stay on-chip
MarkChecks existing GPU cache slots against the hash table to identify hits and evictable slotsHit detection is pure metadata work
ScanParallel scan over slots, updates LRU metadata, picks victims for this step’s missesPrefix-sum-style slot allocation
FetchMiss threads use vectorized non-coherent loads (ld.global.nc.v2.b64) to pull KV from pinned host memory into their assigned slotsGPU-assisted IO: GPU threads fetch for themselves
PublishUpdates page tables and hands physical device offsets to the sparse attention kernelDownstream kernel is unaware

GPU-assisted IO is the real technical contribution here. Letting GPU threads read host memory directly saturates PCIe or NVLink bandwidth even for scattered access, bypassing the long “CPU notices a fault, notifies, then copies” path. And the whole flow runs inside the decode CUDA Graph, so dynamic control flow never breaks graph capture.


4. Locality: Why a Small B Is Enough

Every cache lives on locality, and HiSparse shows sparse selection has two exploitable properties:

Measured miss rates on GLM-5.1 (k = 2048) make the point:

30.0% B = k 13.4% B = 2k 6.7% B = 4k

GLM-5.1, k = 2048: top-k miss rate falls fast as the cache grows Doubling from k to 2k halves the miss rate; 4k halves it again. B only needs a few multiples of k.

Figure 2: The locality dividend under LRU. A cache a few times larger than k already pushes misses into the single digits.

Exact Layer-Wise Prefetching

There is a sharper trick. Some models (for example GLM-5.2) share indexer output across groups of layers — once an anchor layer fixes its selection, the system immediately knows the selections of the layers that share it.

HiSparse turns this into exact layer-wise prefetching:

  1. The anchor layer computes its selection, which yields the shared layers’ miss plans deterministically;
  2. A background copy-only kernel replays those future layers’ miss plans;
  3. Host-to-device transfer then overlaps with the current layer’s compute.

Result: roughly half of the remaining IO latency is hidden. Note the word exact — this is not a heuristic guess but deterministic knowledge, so it introduces no correctness risk.


5. Why It Is Exact and Indexer-Agnostic

These two properties are why it can go straight into production:

In short, HiSparse is a general-purpose memory backend for sparse attention: if your model already has a sparse indexer, it just plugs in.


6. Measured Results

SetupResult
DeepSeek-V4-Flash (NSA), 2x B200, 64 concurrent requests2.1x generation throughput
Qwen3 with Quest, GH200, 200K input lengthup to 4.7x generation throughput
High load (prefill and decode share the GPU)TTFT drops significantly (new requests are no longer blocked by HBM exhaustion)
Per-token latencyComparable to baseline (not faster, but not slower)
No-IO oracle experimentThe resolution mechanism adds no measurable per-token cost

That last row matters: host-device IO is the only, and the entire, price of this design. Which means link bandwidth directly sets the payoff:

Link sensitivity: KV fetch time on GH200 (NVLink-C2C) is nearly 4x lower than on H200 (PCIe Gen5), allowing a smaller GPU cache and higher concurrency.


7. When Not to Use It (Limits and Preconditions)

This section matters more than the speedups:

  1. It assumes host DRAM is much larger than GPU HBM. On Grace-based GB200 / GB300, CPU and GPU share unified memory and the host side has no capacity advantage — the whole offload-for-capacity argument does not hold. The paper lists this as a fundamental limitation.
  2. It relies on GPU-assisted IO reaching near-link bandwidth for scattered fetches. That is borrowed from the authors’ Strata work and is not independently benchmarked in this paper. Misses cost noticeably more on PCIe Gen5 than on NVLink-C2C.
  3. It buys concurrency and capacity, not single-request speed. Per-token latency is merely comparable; it will not make one request finish faster.
  4. It does nothing for dense-attention models. You need a sparse indexer first — without top-k selection there is no footing for “only k entries must be resident.”
  5. Payoff varies with load. At low concurrency or short context the baseline never hits the capacity wall, and the hierarchy only adds an IO hop.

8. The One-Line Thread, Completed

Stack sys3, sys5 and this post together and the KV story is whole:

MLA makes each token store smaller; NSA / DSA / Quest read only the important k entries; CSA compresses tokens first and then does sparse reads; HCA compresses very long history into a short summary and reads all of it — while HiCache and HiSparse answer a different question: once KV exists, which tier of the memory hierarchy should hold it?

One more distinction:

Cutting compute (sparse attention) and cutting capacity (compression plus hierarchy) are orthogonal axes. HiSparse sits at the far end of the second one, and is currently the closest to production (it is in upstream SGLang).



Appendix: Paper

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
Code: merged into upstream SGLang

This post is a close reading and engineering interpretation of the paper. It is not technical or investment advice. Figures are taken from the paper and its alphaxiv reading page (retrieved 2026-09-07).

觉得有用?欢迎点赞、收藏,或请作者喝咖啡 ☕️

支付宝收款码

支付宝

微信收款码

微信

💬 留言

评论由 Giscus 驱动(基于 GitHub Discussions)。 当前仓库 NaphJohn/LLM-blog 尚未启用 Discussions:请在 GitHub 仓库 Settings → General → Features 勾选 Discussions 后刷新本页,评论区即自动显示。