Latent Sparse Attention: Towards Subquadratic Ultra-Long Context Attention
Authors: Ales Kubicek, Andrea Caciolai, Liangzu Peng, Sukrit Sumant, Benjamin Hoover, Fraser Greenlee, Ruairidh M. Battleday, Radim Urban, Torsten Hoefler, James C. R. Whittington
Organizations: Prefrontal AI, London, United Kingdom · ETH Zurich, Zurich, Switzerland
Agentic workflows are increasingly shifting towards long-horizon tasks, such as open-ended scientific work, where models benefit from long context windows. However, at ultra-long context (above 1M tokens), full attention is prohibitively expensive. A promising approach is sparse attention, in which a dedicated indexer selects a subset of keys for the attention. Yet current methods of indexing either still have quadratic cost in context size or lose accuracy at ultra-long context scale. We address this with Latent Sparse Attention (LSA), an efficient attention method that maintains accuracy at ultra-long context. In LSA, keys are hierarchically clustered in a learned latent space, and queries score the hierarchy of clusters instead of scoring every key, making prefill, decode and training faster at long context. We propose two variants, Distilled-LSA (D-LSA), which can be applied to pre-trained models via distillation, and Native-LSA (N-LSA), which is trained natively on the next-token loss alone. On GLM-5.3-FLASH (QWEN-3.8-27B) at 10M tokens, D-LSA speeds up one attention layer by 22.7x (72.4x) in prefill and 4.6x (42.5x) in decode over its original sparse attention (full attention), while keeping similar accuracy across a variety of tasks. N-LSA matches the accuracy of full attention with a training speedup of 1.6x per layer at 128k tokens. LSA offers practical and efficient attention for models to benefit from ultra-long context.
Figures & tables
Figure 1: D-LSA and N-LSA overview. Top: one attention layer. The keys are pooled and clustered into buckets. A query ( ) scores the bucket centroids ( ) and attends only to the selected long-range keys (plus sink and window). Bottom: accuracy and performance (P: prefill, D: decode, T: training) for both D-LSA and N-LSA. D-LSA is shown with flat search. Its subquadratic variant, D-LSA (Tree), follows a similar trend (Figure 5 ).
Figure 2: Traditional sparse attention at ultra-long context. One global attention layer of GLM-5.3-Flash (512k to 10M tokens) on one H200, prefilling the last chunk of 32,768 queries and decoding one token. Speedup of sparse attention over full attention ( blue ), share of the indexer ( % ), and potential speedup over sparse attention if the indexer were ‘free’ ( green ).
Table 1: Sparse attention methods compared. Rows are grouped by the unit the indexer scores: single keys, blocks of consecutive keys, or clusters of similar keys. Design column: compression of the sequence the indexer scans by compressing consecutive keys (C. Seq.), or clustering keys by content (C. Sem.), and its search: flat ( ), a fixed number of levels ( ), or levels efficiently growing as logT ( ). Evaluation : the largest context at which accuracy (Acc.) and performance (Perf.) are reported. supported, partial support, not supported.
Figure 3: D-LSA design. Left: D-LSA (Flat). Indexer keys: 1 projection and L2 norm, 2 pooling (4:1), 3 bottom-up agglomerative clustering (Ward’s method), 4 bucket cache (centroids and index table). Indexer queries: A sum over query heads, projection and L2 norm, B scoring the cached centroids and selecting the top-p buckets per query, C expanding the buckets into indices (sink + top-k long-range + window), D attention over those indices (one index list per query shared by all query heads). Right: D-LSA (Tree) groups the buckets into a hierarchy of clusters, a query descends using beam search (close in the latent space but scattered through the KV cache).
Figure 4: N-LSA design. The attention layer of N-LSA, which can be trained from scratch on the next-token loss alone.
Table 2: Parameters. Block Sparse Attention (Flat), Block Sparse Attention (Tree), D-LSA (Flat) and D-LSA (Tree) per model (symbols in Table 5 ). The search is flat ( ) or tree ( ), over blocks of npool consecutive keys or over buckets. Block Sparse Attention (Tree) always keeps the first, previous and current blocks among its p blocks. On GLM-5.3-Flash , D-LSA (Flat) uses the first k and p below 1M tokens and the second beyond.
Figure 5: D-LSA accuracy is on par with full attention and sparse attention baselines on downstream tasks. Shown on knowledge, reasoning and long-context tasks up to 10M tokens. Baseline scores are given with their 95% confidence interval.
Figure 6: D-LSA speedup of one global attention layer. Prefill and decode speedup over the baseline attention (log scale), Full Attention for Qwen-3.8-27B and Block Sparse Attention (Flat) for GLM-5.3-Flash , required memory at 10M with one index key per block of four keys for Block Sparse Attention (Flat), and prefill scaling up to 100M (log-log), extrapolated beyond 10M (shaded).
Table 3: D-LSA speedup in production setting. Time to first token (TTFT) and time per output token (TPOT) of one request under SGLang. Full Attention, Block Sparse Attention (Flat), D-LSA (Flat) and D-LSA (Tree). Each block’s first row gives the baseline’s times and the others their speedup over it, with D-LSA at 10M in bold. Qwen-3.8-27B is served with tensor parallelism (TP8) and decode context parallelism (DCP8) on H200 and with TP8 on B300, GLM-5.3-Flash with either context parallelism (CP8, with DCP8 and expert parallelism EP8) or TP8. Without DCP, Block Sparse Attention’s prefill indexer runs whole on every TP rank. All models run on a single node, except TP8 for GLM-5.3-Flash at 10M on H200 (2 nodes with PP2); a B300 node holds it on one.
Figure 7: N-LSA training speedup. Models trained natively on the NTP loss with Full Attention, SWA, NSA, PISA and N-LSA. Left: one attention layer, forward and backward, on one H200. Right: validation loss over the training time of the whole model in the decay phase (all on the same data and number of tokens).
Table 4: Native training accuracy. Accuracy (%) of models trained from scratch with Full Attention, SWA, NSA, PISA and N-LSA. Every sparse method attends to about 2,560 keys per query.
Appendix figures & tables25 assets
Supplementary material from the paper’s appendix.
Appendix
Symbol
Description
T
context size
i , j , h
query position, key position, head
WQ,K,V(attn)
attention projections
qi,h(attn) , kj(attn)
attention query (head h ) and key
nh , dh
query heads, head dimension
WQ,K(idx) , d(idx)
index projections, index dimension
Appendix
Table 5: Notation. D-LSA symbols and their description.
Algorithm 1 Ward clustering of one cluster block
Algorithm 2 Beam search of one query in D-LSA (Tree)
Figure 9: LongDocs. Tokens of the distillation mix by domain and source (left) and by sequence length (right), with the in-house agent traces hatched.
Figure 10: PaperTrail. 1 A sample, its fabricated, host and background papers. 2 One trail, introduced, revised and renamed, and still quoted in its old form after the revision. 3 Questions by tier, and one whose answer is not in the papers. 4 Context sizes, each holding every paper of the size below. 5 Controls.
Figure 11: D-LSA accuracy on RULER individual tasks. The 13 tasks of RULER from 4k to 10M tokens on Qwen-3.8-27B and GLM-5.3-Flash , by task group, with their average.
Questions
Repeats
Reasoning
Max output
Metric
MMLU-Pro
2,000
1
on
163,840
Exact match
GPQA-Diamond
198
1
on
163,840
Exact match
AIME25
30
4
on
163,840
Math-Verify
MATH500
100
1
on
163,840
Math-Verify
LongBench v2
60
1
on
163,840
Accuracy
RULER
100 ( ≤ 1M) / 20 ( > 1M)
1
off
4,096
String match
Appendix
Table 6: Evaluation setup. Questions, repeats, generation and metric of every benchmark, with RULER questions per task.
Table 7: Accuracy compared to the baseline. Baseline scores are in bold. Other rows show the difference to the baseline in points. Differences smaller than the 95% confidence interval (CI) of the baseline are not significant.
Figure 12: Reasoning tokens against baselines. Reasoning tokens of each method relative to the baseline of its model, on the same questions. Grey bands mark the expected variation, and markers inside them show no measurable change in reasoning length.
Figure 13: D-LSA accuracy on PaperTrail per tier. Tier 1, tier 2 and tier 3 questions, each averaged over the three domains and three repeats.
Table 8: PaperTrail not-in-corpus questions and controls. Score (%) on the questions whose answer is not in the papers at 512k tokens, for Full Attention, Block Sparse Attention (Flat) and (Tree), and D-LSA (Flat), (Tree) and (Training-Free). The controls ask the scored questions over parts of a sample, with GLM-5.3-Flash and D-LSA (Training-Free).
Figure 14: Block size of Block Sparse Attention (Tree). RULER and PaperTrail up to 1M tokens with blocks of 8, 16 and 64 keys ( npool ), each selecting the same 4,608 keys, against the baseline of each model and D-LSA (Flat).
H200
B300
Architecture
Hopper (sm_90)
Blackwell Ultra (sm_103)
SMs
132
148
Shared memory per SM
228 KiB
228 KiB
Matrix instructions
WGMMA
tcgen05
Memory (HBM3e)
141 GB
288 GB
HBM bandwidth
4.8 TB/s
8 TB/s
Appendix
Table 9: Hardware specs. Specifications of the GPUs we use, as given by NVIDIA ( NVIDIA, 2026c ; NVIDIA, 2026d ; NVIDIA, 2026a ; NVIDIA, 2026b ) .
Figure 15: Implementation of the steps of D-LSA. Badges as in Figure 3 .
Figure 16: Threshold estimate. For one query i . Top: the sampled scores, with θi at or just below the nθ -th best. Bottom: the row, of which about γp buckets score at or above θi and are listed, the top- p among them.
Figure 17: B Bucket scan on H200. Left: the thread blocks (colours) over tiles of tq index queries and spans of centroid tiles, and below, candidate scoring and C expansion for one query. Right: one thread block (Algorithm 3 , lines 3 - 9 ).
Algorithm 3 Flat prefill selection of one chunk
Algorithm 4 Flat decode selection of one request
Table 10: Attention kernels. Kernel through which each method attends, by model, phase and GPU, at TP8 or on one GPU. Full Attention, Block Sparse Attention (Flat), D-LSA (Flat and Tree). Kernels in italics are non-default SGLang options.
Algorithm 5 Sparse attention backward of N-LSA, for one query head
Step
CuTe DSL (ms)
Triton (ms)
Speedup
Regime
Flat prefill
Threshold estimate
2.02
8.85
4.38 ×
compute (28%)
Bucket scan
11.39
24.82
2.18 ×
compute (49%)
Candidate scoring
1.63
1.66
1.02 ×
memory (85%)
Expansion
2.37
4.14
1.75 ×
selection
Total
17.41
39.48
2.27 ×
Appendix
Table 11: Indexer kernels. Time of each step of the indexer of D-LSA on one H200, with its CuTe DSL and its Triton kernels, and the regime of its main CuTe DSL kernel.
128k
1M
10M
Requests
Flat
Tree
Flat
Tree
Flat
Tree
1
30
38
35
49
86
68
4
32
40
56
54
211
81
16
42
47
109
82
700
125
Appendix
Table 12: Tree search in decode. Time of the selection of one layer for one decode step, from the scoring of the centroids to the index lists, in µs, with the flat and the tree search, on one H200, replayed from a CUDA graph. The faster search is in bold.
Figure 18: D-LSA with tensor parallelism (TP8). Top: MLA ( GLM-5.3-Flash ). Bottom: GQA ( Qwen-3.8-27B ). Indexer keys: 1 projection (from the replicated latents with MLA, a partial per KV head with GQA), 2 bucket cache (the same on every rank). Indexer queries: A sum over the rank’s query heads and projection (a partial, summed across ranks by one AllReduce , before the projection with MLA), B selecting the index list (the same on every rank), C attention over the rank’s own query heads.
Phase
Collective
Size
Frequency
TP, MLA
both
AllReduce of query head sums
2 KiB / token
every chunk and step
TP, GQA
both
AllReduce of partial index queries and keys
512 B / token
every chunk and step
PCP, MLA
prefill
-
-
-
DCP
decode
AllReduce of index keys
64 KiB / block
every ncluster steps
Appendix
Table 13: Communication added by the indexer. Collectives that the indexer of D-LSA adds to one layer for TP alone, PCP and DCP, with MLA ( GLM-5.3-Flash ) and GQA ( Qwen-3.8-27B ), beyond those of the attention. Sizes are per rank, and decode frequencies per request.
Figure 19: D-LSA with context parallelism (CP8). Top: prefill of GLM-5.3-Flash (PCP8 + DCP8), each rank holding a contiguous span of the query rows. Indexer keys: 1 projection of the whole chunk (on every rank), 2 bucket cache (the same on every rank). Indexer queries: A projection of the rank’s own rows, B selecting their index lists (other ranks’ rows lighter), C marking the selected prefix rows (a 1 B mask, merged across ranks), D gathering those rows from the ranks that hold them (padded), E attention of the rank’s own rows. Bottom: decode of GLM-5.3-Flash (PCP8 + DCP8) and Qwen-3.8-27B (TP8 + DCP8). Indexer keys: 1 projection of a complete cluster block (summed across ranks, once per block), 2 bucket cache. Indexer queries: A projection, B selecting the index list (the same on every rank), C attention over the entries the rank holds, D merging the partial outputs of the ranks (by log-sum-exp).
Scaling modern large language models (LLMs) to long contexts is limited by the quadratic computation cost, and poor length extrapolation of dense attention. Chunk-wise sparse attention offers a promising alternative, but all existing methods fall short of full attention because of their inaccurate chunk selection. We propose Hierarchical Landmark Sparse (HiLS) Attention, a chunk-wise sparse attention mechanism that learns chunk selection end-to-end under the language-modeling (LM) loss. HiLS factorizes attention hierarchically: each query performs attention independently with each retrieved chunk to extract chunk-specific information, and the resulting outputs are fused according to chunk retrieval scores. By incorporating retrieval scores into the forward attention computation, HiLS optimizes them directly with the LM loss, enabling end-to-end retrieval learning and native sparse training. Experimental results show that HiLS-Attention achieves performance comparable to, and in some cases better than, full attention at in-domain context lengths. Meanwhile, HiLS-Attention extrapolates more than 64× the training context length with 90% retrieval accuracy, far beyond full attention. Moreover, existing full-attention models can be converted to HiLS-Attention with lightweight continued pretraining, preserving in-domain performance while acquiring ultra-long-context extrapolation. Together with its sparse KV access and computation, HiLS-Attention breaks the usual efficiency-performance trade-off, enabling long-context LLMs that are both more efficient and more effective on general long-context tasks than their full-attention counterparts.
Xiang Hu, Xinyu Wei, Hao Gu +10
Tencent HY Team · ShanghaiTech University · The Hong Kong University of Science and Technology +1
Long contexts have become standard in pretrained LLMs, yet they remain expensive to run: prefill compute grows quadratically with sequence length, and every decode step re-reads a key-value cache that grows linearly with it. Sparse attention cuts these costs by attending only to a relevant subset of past tokens, but selecting that subset is itself expensive. We present SpotAttention, a lightweight selector that attaches to a frozen pretrained transformer and learns by KL distillation to estimate its attention distribution. The selector picks the top-K keys each query attends to, and because its estimate is a calibrated distribution, a dual top-p rule reads the per-query, per-layer budget directly from it. Across Qwen3 (dense, 4B-32B) and Qwen3.5 (hybrid linear/full attention, 4B-9B), SpotAttention matches dense accuracy at contexts up to 128K tokens, eight times the training length. Decode at L=128K runs 3.9x faster than FlashAttention and 1.8x faster than Twilight, the strongest training-free baseline. Quantizing the selector's K-cache to INT4 or FP4 microscale shrinks it 3.5x at no accuracy cost.
DeepSeek Sparse Attention (DSA) enables efficient long-context modeling through its Lightning Indexer. However, practical deployment remains constrained by the indexer's expensive O(L2) scoring overhead and the hardware-inefficient, discontinuous memory-access patterns induced by its outputs. To address these system-level bottlenecks, we introduce LongCat Sparse Attention (LSA), a hardware-algorithm co-designed framework comprising three complementary and orthogonal strategies: (1) Streaming-Aware Indexing, which selectively converts scattered KV entries into hardware-aligned contiguous layouts to enable coalesced HBM access; (2) Cross-Layer Indexing, which amortizes indexing overhead by reusing the results produced by a single layer across consecutive layers, supported by cross-layer distillation; and (3) Hierarchical Indexing, which adopts a coarse-to-fine scoring scheme to progressively narrow the candidate set for each query, thereby substantially reducing indexing computation. Extensive scaling experiments, ranging from 69B-A3B to 560B-A27B models, demonstrate that LSA consistently achieves performance on par with full attention across both general-purpose and long-context benchmarks. Moreover, LSA supports native training with context lengths of up to one million tokens and underpins the development of LongCat-2.0 (1.6T-A48B). To facilitate further research, we also introduce and open-source LongCat-Flash-Lite-Sparse (69B-A3B), which integrates LSA into LongCat-Flash-Lite and incorporates an updated long-context training corpus.