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:
| Scenario | KV memory footprint | Consequence |
|---|---|---|
| GLM-5.1, single 128K-context request | ≈ 13.09 GB | An 80GB card serves only a handful concurrently |
| Single 1M-context request | Exceeds total HBM | Cannot 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:
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:
- Host KV pool — KV records produced during prefill go straight into pinned host memory. This is the only authoritative copy.
- GPU hot cache — B slots per request-layer pair holding recently selected KV records. To guarantee the current attention can always proceed, B ≥ k.
- 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)
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:
- 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.
- 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.
- 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:
| Stage | What it does | Why it matters |
|---|---|---|
| Stage | Loads the indexer-selected logical positions into a shared-memory hash table | Later comparisons stay on-chip |
| Mark | Checks existing GPU cache slots against the hash table to identify hits and evictable slots | Hit detection is pure metadata work |
| Scan | Parallel scan over slots, updates LRU metadata, picks victims for this step’s misses | Prefix-sum-style slot allocation |
| Fetch | Miss threads use vectorized non-coherent loads (ld.global.nc.v2.b64) to pull KV from pinned host memory into their assigned slots | GPU-assisted IO: GPU threads fetch for themselves |
| Publish | Updates page tables and hands physical device offsets to the sparse attention kernel | Downstream 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:
- Temporal locality — consecutive decoding steps re-select many of the same tokens;
- Cross-layer correlation — different layers tend to attend to nearby regions of the context.
Measured miss rates on GLM-5.1 (k = 2048) make the point:
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:
- The anchor layer computes its selection, which yields the shared layers’ miss plans deterministically;
- A background copy-only kernel replays those future layers’ miss plans;
- 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:
- Exact — it changes only the physical placement of KV records: no model surgery, no change to attention math, no approximate recall. Model output is unchanged token by token, which separates it fundamentally from “trade memory for approximate retrieval” schemes.
- Indexer-agnostic — it does not care how you picked those k positions. The paper evaluates DSA, NSA and Quest and all three benefit.
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
| Setup | Result |
|---|---|
| DeepSeek-V4-Flash (NSA), 2x B200, 64 concurrent requests | 2.1x generation throughput |
| Qwen3 with Quest, GH200, 200K input length | up 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 latency | Comparable to baseline (not faster, but not slower) |
| No-IO oracle experiment | The 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:
- 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.
- 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.
- It buys concurrency and capacity, not single-request speed. Per-token latency is merely comparable; it will not make one request finish faster.
- 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.”
- 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:
- HiCache asks “should KV sit on GPU, CPU, or even lower tiers” — a storage-tier question;
- HiSparse asks “under sparse decoding, how much must actually be resident on the GPU” — making the capacity bill scale with k rather than L_ctx.
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).
9. Links to Earlier Posts
- Full KV compression landscape (MHA to HiCache) — see
sys5-attention-evolution-kvcache - MSA / CSA / HCA: three attention redesign routes — see
sys3-attention-evolution - HiCache and the NUMA x PCIe x NIC data path — see
sys1-pcie-numa-nic-hicache - SGLang internals (prefix reuse and throughput optimization) — see
fw3-sglang-internals - HiCache / Mooncake KV reuse and compositional correctness — see
tr2-vllm-sglang-20260807
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).
💬 留言
NaphJohn/LLM-blog尚未启用 Discussions:请在 GitHub 仓库 Settings → General → Features 勾选 Discussions 后刷新本页,评论区即自动显示。