LoopKV: Loop-Aware KV Cache Compression for Serving Looped Transformers

L OOP KV: L OOP -AWARE KV C ACHE C OMPRESSION FOR S ERVING L OOPED T RANSFORMERS Tianle Zhou * 1 Yang Zhang * 2 Tianyu Wu * 1 Zhuohan Wang * 1 William Nixon 3 Juncheng Yang 1

arXiv:2611.04137v1 [cs.LG] 3 Nov 2026

A BSTRACT Looped transformers apply a shared block several times per token to improve accuracy without adding parameters. For models such as Ouro that keep a separate key-value (KV) cache per loop, cache memory limits the batch size that amortizes repeated weight reads. Uniform-bit quantization and fixed rules per loop do not allocate precision according to each layer’s sensitivity in each loop. Measurements on Ouro show that quantization sensitivity varies across (loop, layer) cells, and the sum of individual cell damages approximates the overall output damage. Sharing damage also varies by cell, with later loops generally tolerating sharing better. Skipping low-attention values to save reads produces errors that compound across loops. LoopKV uses a one-time offline profile to choose 4-bit storage, 2-bit storage or reuse of the final loop’s KV for each cell. The planner minimizes reads under budgets for predicted damage and storage. During decoding, each loop reads the high 2 bits of each 4-bit value and fetches the low bits only where its own attention weights require 4-bit precision. A single fused kernel in vLLM executes Current Figure 1Serving in the paper the plan. Ouro-1.4B on one RTX PRO 6000, the memory-optimized LoopKV-M plan stores 35% less KV than uniform 4-bit quantization (KIVI-4) on the same fused kernel and CUDA graph. It sustains 1.7–2.2× The system overview now in the paper (figures/v4/fig_system.tex, GitHub main ecd5d97): profile and knapsack, plan, KIVI-4’s under a flow, latency objective with a 16K-token maximum length. None of its averages per-cell kernel launches, the fusedgoodput kernel’s operator and service-level launches per step. on six public benchmarks falls significantly below BF16’s.

ONLINE

OFFLINE 1

Contribution 1: Per cell optimization

3

Contribution 3: Dynamic reading

We store V in two planes with higher and lower bits of INT4. A runtime decider allocates precisions dynamically with attention scores.

Do quantization in each cellnm (layer n, loop m) according to the damage loop 1

loop 4 reads with its query q4

loop 3 reads with its query q3

K, 4 bits

loop 2

q ⋅k

Threshold

loop 3

V, high 2 bits loop 4

V, low 2 bits INT4

2

INT2

Block sharing

Contribution 2: Cell-level KV sharing

KV sharing through a direct pointer redirection

4

Combined Synergy: The contributions can be wrapped as a fused kernel for every cell type 1 BF16 at Sink/Tail

without sharing loop 4

LoopKV plan (served) loop 4

loop 4’s pool loop 3

loop 3 loop 3’s pool

BF16 sink history

loop 4’s pool freed

host launches per decode step

1 fused kernel = 1 launch per cell, the same code for every cell type 2 Get seed m 0

BF16 tail

Exact BF16 attention over the sink and the recent window.

m0 sink scores

The largest sink score, taken from step 1 with no extra read.

near

V hi V lo

504

4 merge

3 History pages hot

Dynamic V precision based on attention score higher than m 0 at compute time.

far

LSE

332

output

The last block merges all partials by log-sum-exp, exactly.

17 unfused

fused

fused + CUDA graph

Figure 1. Overview of LoopKV. The offline phase of LoopKV, run once per model, chooses a storage option for every (loop, layer) pair, which we call a cell, from the KL damage measured by profiling each cell independently (C1). A cell may reuse the final loop’s KV through a pointer when its KL damage is low, which we show costs little accuracy (C2). In the online phase, V is stored as a high and a read-optimized + V precision BF16 768 KiB KiB KiB INT4 history KV read per token, 4 loops: low 2-bit plane, so that the attention scores decide at decode timeuniform which tokens240also read the low plane215 (C3). We then combine186allKiBthree into one fused kernel that runs every cell type in a single launch. With vLLM’s whole-step CUDA graph, this cuts LoopKV-M’s host launches per decode step from 504 to 17. * Equal contribution. 1 Harvard University 2 University of California San Diego 3 University of Chicago. Correspondence to: Tianle Zhou <tianle_zhou@fas.harvard.edu>.

LoopKV: Loop-Aware KV Cache Compression for Serving Looped Transformers

I NTRODUCTION

Looped transformers, which reuse one block of layers for several iterations per token, are increasingly used as an architectural component of large language models (Dehghani et al., 2019; Saunshi et al., 2025). Because every iteration reuses the same weights, a looped model reaches the effective depth of a much deeper model, i.e., the number of layers times the number of loops, while storing far fewer layers. It therefore needs less memory to hold its parameters and achieves higher accuracy per parameter (Saunshi et al., 2025). For example, Ouro matches standard models with two to three times its parameter count by looping its layers four times (Zhu et al., 2025). Huginn, with 3.5B parameters, keeps improving its reasoning as it iterates its recurrent block more, up to a compute load equivalent to a 50B-parameter model (Geiping et al., 2025). Mixtureof-Recursions chooses for each token how many times the shared block runs and reaches lower perplexity than standard transformers at equal training compute with fewer unique parameters (Bae et al., 2025b;a). Serving these models, however, is costly, because for some of them the KV cache is large and each decode step reads many more KV blocks from memory than a standard model of the same size does. Ouro, for example, keeps a separate KV cache for every loop, so Ouro-1.4B stores 768 KiB of KV per token in BF16, four times what a dense model with the same layers would store, and every decode step re-reads its 2.5 GB of weights four times. A larger batch amortizes these weight re-reads, but then KV memory becomes the bottleneck, because the cache stored per request sets how many requests fit in a batch and the bytes read per step set the decode time (§2.2). At BF16’s peak-throughput batch for 8K-token contexts on Ouro-1.4B, the KV cache accounts for 84% of the bytes each decode step reads. In this setting, compressing the cache raises both the batch size and the decoding speed, which is the key to higher throughput and goodput. General KV quantization methods, including FP8, KIVI (Liu et al., 2024b), KVQuant (Hooper et al., 2024) and rotation-based methods (Ashkboos et al., 2024; Zandieh et al., 2025), treat each loop of a looped transformer as one more layer with the same bits. Only a few works target the features specific to looped transformers, and we find two concurrent works that exploit the loop structure. FlashLoop updates only a subset of tokens in each loop, guides each loop’s sparse attention with the previous loop and quantizes chained INT4 residuals (Yang & Liu, 2026). ResidualQuant keeps an INT4 copy of the final loop and stores only 2-bit residuals for every earlier loop (Kim et al., 2026). Both assign precision on a per-loop basis. A per-loop setting, however, treats every layer of a loop alike, which is harmful when the layers within a loop behave differently.

difference (%)

1

loop 1

0.3

loop 2

loop 3

0.2 0.1 0.0

1

6

12

18

24

layer

Figure 2. P1: damage depends on the cell. Each point shows the damage when one cell (one layer in one loop) is stored at INT2 while every other cell stays exact. The damage on the y-axis is the share of next-token predictions that change (fp32, 16 documents). The top 19 of the 72 cells hold half of the total difference, and KL ranks the cells alike (Spearman 0.95).

For example, our experiments show that under aggressive quantization a few middle layers cause most of the damage, while some later layers are much less sensitive (§2.4). Some layers can therefore use fewer bits and leave more of the budget to the layers that matter more. What remains missing is a decision at the granularity of each (loop, layer) pair, e.g., which precision each pair needs and whether a pair can reuse another loop’s KV instead of storing its own when its damage is low enough. We propose LoopKV, which compresses the KV cache of looped transformers at the granularity of a (loop, layer) pair, which we call a cell. LoopKV first profiles the damage of each cell offline and then jointly chooses every cell’s precision and whether it shares the final loop’s KV, so that the bytes saved on robust cells can be spent on the more sensitive ones. This planner minimizes the bytes read per decode step under budgets for KL damage and storage. Looping also raises two questions that a transformer with one KV cache per layer does not, i.e., whether the KV of the late loops converges enough to be shared, and whether an error introduced in one loop compounds as later loops process it again. Our measurements of Ouro’s KV cache answer both questions and show why the decision must be made per cell rather than per loop (§2.4, Figures 2–4). (P1) Sensitivity depends on the cell: a few cells cause most of the quantization damage, and which cells they are depends on the model. Cell damages also approximately add up to the model’s final KL divergence, so a single profile per model predicts a plan’s damage before it is served (Figure 2). (P2) Sharing does less damage in later loops: for some cells, reading the final loop’s KV costs less than storing their own INT2 copy, e.g., in 21 of the 24 loop-3 layers and 6 loop-2 layers, but in no loop-1 layer (Figure 3). (P3) The loops compound read errors: each loop injects nearly the same error, so skipping the pages of unimportant tokens costs more on looped models than on plain ones. In Ouro1.4B, page skipping in all four loops costs 1.8× the sum of its one-loop damages in KL, over 50% more than the at most 1.17× of plain models such as Qwen2.5 and Qwen3.

LoopKV: Loop-Aware KV Cache Compression for Serving Looped Transformers loop 1 loop 2

difference (%)

10

loop 3 sharable: less damage than INT2

1 0.1

it at different precisions without extra bytes. In vLLM, one fused kernel executes every cell type in a single launch, and the whole decode step replays as one CUDA graph, which cuts LoopKV-M’s host launches per decode step from 504 to 17 (Figure 1). Our contributions are:

0.01 1

6

12

18

layer

• Characterization. We measure three properties of the per-loop KV cache on Ouro-1.4B (P1–P3) and show that its sensitivity profile is specific to each model, since it differs across Ouro-2.6B, Dory-4B and Huginn3.5B. On each of these models, a plan computed from the model’s own per-cell profile beats hand-made plans of equal bytes (§2.4, §5.7). • Design. LoopKV turns each property into one mechanism, and one fused kernel executes all three in a single launch. They include per-cell mixed-precision allocation (C1), cross-loop KV sharing as a per-cell option, with the planner deciding which cells share (C2), and read-time, attention-aware V precision (C3) (§3, §4). Restricting the plan to one option per layer costs 1.4–1.9× the KL at equal bytes. • Evaluation. On Ouro-1.4B, Ouro-2.6B and Dory-4B, LoopKV-M reaches 1.13–4.44× the decode throughput of vLLM’s BF16 cache at 16K context on three kinds of GPUs, and none of its suite averages falls more than 0.5 percentage points below BF16’s (up to six suites per model) (Table 1, Table 7, §5.7). On Ouro1.4B, against uniform INT4 (KIVI-4) on the same kernel and CUDA graph, LoopKV-M stores 35% less KV and sustains 1.7–2.2× KIVI-4’s request rate under a latency SLO (Figure 5).

24

Figure 3. P2: some cells can share with little loss. Each point shows the damage (the y-axis, as in Figure 2) when one cell reads the final loop’s KV instead of storing its own, while every other cell stays exact. Some cells (green) do less damage when sharing than when stored at INT2

Ouro, Quest, by loop Ouro, Quest, by layer slice Qwen2.5, Quest Qwen3, Quest

damage (×10−6)

relative damage†

(a) Quest's errors compound across loops, not across layer slices

3.0 2.5 2.0 1.5 1.0 0.5

additive

1

2

3

loops (layer groups)

4

(b) damage under the same bytes read 10.0 4.7× 7.5 5.0 2.5 0.0 LoopKV Quest style

Figure 4. P3: the loops compound read errors. Quest drops the low-ranked far pages that are not in the sink and window, while LoopKV-style reading keeps them and reads the values of lowattention far tokens at 2 bits. Both read 85% of the far bytes. (a) Relative damage† of Quest as k grows, applied to the first k loops of Ouro-1.4B, to k of four layer slices of Ouro-1.4B that each span all loops, or to k of four layer groups of Qwen2.5 and Qwen3. (b) Damage of LoopKV-style reading and of Quest in all four loops of Ouro-1.4B at the same 85% of bytes read. Damage is the fp32 KL to the exact model. † Damage when the k groups are treated together, divided by the sum of their damages when page skipping is applied to each group alone.

At equal bytes, LoopKV-style reads, which read the values of unimportant tokens at 2 bits instead of skipping them, cause 4.7× less damage. Loops that share a block, moreover, attend to different tokens (Jaccard 0.83 for loops 3 and 4), so which values need full precision depends on the loop that reads them and must be decided at run time (Figure 4). In the offline phase, LoopKV profiles every cell once and measures its damage under each storage option, i.e., 4-bit codes, 2-bit codes (P1) or a pointer to the final loop’s KV (P2). It then solves a knapsack over these profiles that minimizes the KV bytes each decode step reads under a damage budget. Adding a tight storage budget gives the memory-optimized plan (LoopKV-M), and leaving storage unconstrained gives the read-optimized plan (LoopKV-R). In the online phase, LoopKV stores each 4-bit value as two 2-bit planes and reads them at a precision set by the attention each token receives, i.e., a loop reads the low plane only for the tokens that its own attention marks as important (P3). Loops that share a block can therefore read

2

BACKGROUND AND M OTIVATION

2.1

Looped and recurrent-depth transformers

A looped transformer applies a block fθ of L layers T times, h(t) = fθ (h(t−1) ), and decodes from h(T ) (Dehghani et al., 2019; Saunshi et al., 2025). In loop t, layer ℓ writes keys and (t) (t) values Kℓ , Vℓ and attends over the same loop and layer’s KV of all earlier tokens. Ouro-1.4B has T =4, L=24 and 16 KV heads of dimension 128 without grouped-query attention, so it stores 768 KiB of BF16 KV per token (Zhu et al., 2025). Huginn-3.5B iterates a 4-layer core r=32 times between a prelude and a coda (Geiping et al., 2025); Dory-4B (NVIDIA, 2026) runs T =4 loops over 20 recurrent attention layers (5 full-attention, 15 sliding-window), plus 16 non-recurrent attention layers. LoopKV keeps the model’s weights, architecture and number of loops and compresses only its KV cache, exploiting how a looped model uses its KV blocks across loops and layers. We therefore compare LoopKV with general KV cache compression methods, e.g., KIVI (Liu et al., 2024b) and TurboQuant (Zandieh et al., 2025), and with methods designed for looped transformers,

LoopKV: Loop-Aware KV Cache Compression for Serving Looped Transformers

goodput (× BF16)

8× 8× 6× 6× 5× 4× 4× 3.2× 2.6× 2× 1× 0 BF16 FP8 FP8 KIVI KIVI-4 Loop- Loop-

(b) LoopKV: no significant loss vs BF16; naive baselines lose

100 accuracy (%)

(a) LoopKV-M reaches 8× BF16's goodput, 2× KIVI-4's

bars: accuracy; error bars: 95% paired CI of the change from BF16

naive: INT2, loops 1-3 naive: last-loop KV reuse

75 50 25

+S

-4

+V KV-R KV-M

BF16

FP8

FP8 + loop-3 sharing

0

RULER-16K RULER-32K LongBench MATH500 KIVI-4

KIVI-4 + V

LoopKV-R (ours)

GSM8K

AIME + MATH-L5

LoopKV-M (ours)

Figure 5. LoopKV serves more requests with no significant accuracy loss on six public suites (tested on Ouro-1.4B). (a) Goodput relative to BF16 on one RTX PRO 6000 for the memory-optimized plan (LoopKV-M), which minimizes reads under a tight storage budget, the read-optimized plan (LoopKV-R), which minimizes reads with no storage budget, and comparable methods. Goodput is the highest request rate at which 95% of requests meet the SLO2 . In the figure, LoopKV-M sustains 8.0× BF16’s goodput and 2.0× that of KIVI-4 (§5.4). (b) Over all the public suites we tested (Table 1), neither LoopKV-M nor LoopKV-R is significantly less accurate than BF16, whereas naive INT2 on loops 1–3 in the KIVI style and naive last-loop KV reuse lose accuracy significantly on several suites. Other comparable methods: KIVI-4 + V adds our V precision to KIVI-style INT4 on every loop, on the same fused kernel and graph. BF16, FP8 and FP8 + loop-3 sharing use vLLM’s own kernels, and KIVI-4 is uniform INT4 without V precision.

e.g., FlashLoop (Yang & Liu, 2026) and ResidualQuant (Kim et al., 2026). 2.2

KV cache in decoding

A decode step processes one token for each of B requests in
t(B) = tfix + WT + B · Dread /BW, (1) the memory-bound view of decoding (Williams et al., 2009; Pope et al., 2023). In the formula, WT counts the bytes of the T -fold weight read per step, Dread the KV bytes one request reads per step, BW is the effective bandwidth of the whole step, and tfix is the step time that does not scale with bytes. In serving, the batch is bounded by B ≤ (KV pool)/Dstored , where Dstored is the KV bytes one request stores (Kwon et al., 2023). Unlike a standard model, a looped transformer reads its weights T times per step. Sharing, residuals and V precision also make the bytes it reads differ from the bytes it stores, so optimizing Dread and optimizing Dstored are separate goals. In our experiments, Eq. (1) predicts the peak throughput of the quantized formats we checked within 3%3 . For a history of N tokens, Dread =N dread + Dexact , and likewise for storage, where d is the bytes per history token over all layers and loops, which we report in KiB, and Dexact is the cost of the sink and window tokens kept in full BF16 precision (Appendix C). Our experiments show that once the KV cache is quantized, throughput is bounded by the bytes read, not the bytes 3 On our RTX PRO 6000, WT =10.07 GB, and a fit gives BW = 1.29 TB/s and tfix =6.84 ms. The formats were served by our plugin before kernel fusion and peak at 98–103% of the prediction.

stored. At the same 180 KiB stored per token, KIVI-4 with loop 3 sharing loop 4’s KV reads 240 KiB per token and decodes 1.36× as fast at 8K as a residual-coding baseline with the same KIVI-style quantization, which reads 360 KiB. In contrast, uniform KIVI-4 reads the same 240 KiB but stores 33% more, and it is only 2% slower (all before kernel fusion, Table 8 and Table 11 in Appendix C). 2.3

Measuring KV damage

We measure the damage of a KV policy with the KL divergence between the exact model and the model under that policy (Dutta et al., 2024). As in serving, the sink and the recent window stay BF16 (a 32-token sink and a 256-token window), and each 32-token page is processed by the policy when it leaves the window. We report damage in units of KLshare , the KL of the cheapest whole-loop sharing4 . All these measurements run in fp32 activations5 . Under this measurement, the memory- and read-optimized plans of LoopKV keep the exact model’s top token at 99.63% and 99.72% of positions. 2 A request meets the SLO if its first token arrives within 10 s and its mean time per output token stays below 3τ0 =55.8 ms, for AIME prompts whose contexts grow to 16K tokens (§5.1). 4 Naively letting loop 3 reuse loop 4’s BF16 KV beyond the window, which gives a KL of 1.20×10−4 . 5 In BF16, rounding noise alone produces about half the KL of KLshare and would hide the differences between cells, whereas in fp32 every effect we report lies well above the noise floor (Appendix E). Every KL in this paper is fp32 unless marked “served”. KLshare is a unit, not an acceptance threshold, and the task scores of §5 test whether a KL difference matters.

LoopKV: Loop-Aware KV Cache Compression for Serving Looped Transformers

2.4

Three loop-specific properties

We observe three loop-specific properties in our experiments, and each of them leads to a corresponding optimization in LoopKV (§3). P1: Sensitivity depends on the cell, not just the layer. When we quantize one cell of the history at a time to KIVIstyle INT2 (Figure 2), a few layers within a loop turn out to be much more sensitive than the others. Loop 1 and four middle layers (7, 8, 10 and 11) cause most of the KL damage. These four layers are on average 4.9× as sensitive as the other layers and hold half of the total damage, while the last layers add almost no damage. At INT4, three quarters of all cells cost less than 1% of KLshare , so most cells can be quantized more aggressively. Sensitivity also varies as much with the loop as with the layer, and loop 1’s cells are together 3 times as sensitive as loop 3’s. Which loop dominates depends on the model, i.e., the first on Ouro but the last on Dory and Huginn in fp32 (§5.7), so no rule fixed by loop index transfers across models without a model-specific calibration. We therefore profile each cell once before planning. Fortunately, profiling each cell once is enough, because a plan’s KL is approximately predictable from the sum of its cells’ damages (Proposition 2 in Appendix D). Over the 14 plans and residual designs we test, the summed per-cell damage and the measured KL have a correlation of 0.98. P2: Some cells can be shared with less damage than INT2 quantization. A cell can read the final loop’s KV instead of storing its own, which removes its storage entirely. For some cells, this sharing does less damage than quantizing the cell to INT2 in the KIVI style. On Ouro-1.4B, sharing changes fewer predictions than own INT2 in 21 of the 24 layers of loop 3 and in 6 layers of loop 2 (Figure 3). Letting loop 3 read loop 4’s KV stays non-inferior to the exact model under a 3-point margin on four of the five suites in Table 2, all but the MQ641 stress test. After profiling, LoopKV can therefore share such cells with the final loop without a large impact on accuracy. Sharing in earlier loops costs far more, changing about 4 times as many predictions per loop-1 cell as per loop-3 cell.6 This loss cannot be recovered by aligning the loops, because their keys have nearly equal norms but different directions, so a MiniCache-style norm correction (Liu et al., 2024a) recovers only 2.7% of the sharing KL (Appendix L). Some models are trained so that their loops share KV, e.g., Dory’s last-loop mode, whereas LoopKV serves the released model and treats sharing as one more per-cell option, deciding from each cell’s measured 6 The Ouro authors report that letting every loop read the last loop’s KV costs at most 2 points with few-shot prompts on the base model (Zhu et al., 2025). When the Thinking model generates up to 8K tokens in our engine, the same reuse loses 44.4 points on MATH500, whereas loop 3 reading loop 4 loses at most 1.0 point (Appendix L).

KL where sharing pays off. P3: The loops compound read errors. Every loop rereads the whole history at each step, and these reads bound the decoding speed. At the 8K throughput peak, the history reader takes 75% of GPU time (Table 16). Yet attention is sparse. In every loop, about one history token in six receives an attention weight of at least 2−14 of the largest weight in its row (Table 14), and the same holds for Dory and Huginn (Appendix F). Traditional methods skip or drop the low-attention values to save reads. Looped transformers, however, are more sensitive to such dropping. Each loop repeats the read with the same weights on converging states, so it injects nearly the same error, and these errors compound across loops into much larger end-to-end damage (Appendix D). Applying Quest’s page dropping in all four loops costs 1.82× the sum of its one-loop damages, whereas for plain models dropping in several layer groups together costs only 0.95–1.17× the sum. At equal bytes, page dropping also does 4.7× the damage of reading the same values at 2 bits (Figure 4). Therefore, low-attention values should be kept and read at low precision rather than skipped. Because loops attend to different tokens, each reading loop should also choose for itself, at compute time, which values to read at full precision (C3).

3

S YSTEM D ESIGN

LoopKV works in two phases (Figure 1). Offline, once per model, the planner takes the model’s profile (§2.3) and budgets for damage and storage, and generates a plan that assigns each cell one storage option, i.e., 4-bit codes, 2-bit codes or the final loop’s KV (§3.1, §3.2). Our vLLM plugin then turns the plan table into one descriptor per cell at load time. In the online phase, storage is fixed and one decision remains. LoopKV dynamically sets the precision at which each loop reads the values, based on the attention scores computed at run time (§3.3). One fused kernel then executes every descriptor (§3.4). 3.1

C1: Sensitivity-aware mixed-precision allocation

Profiling. For every cell c=(t, ℓ), profiling measures the single-cell damage dc,o = KL(pexact ∥ pc←o ), where we apply each of the three storage options o, to that cell alone, keep every other cell exact, and measure the KL on 16 of the model’s own long generations in fp32 (§2.3). This profiling is cheap and predictive. In our experiments, summing the profiled damages of a plan’s cells predicts the plan’s measured KL with a correlation of 0.98 (§2.4)7 . 7 On Ouro-1.4B, profiling covers 72 cells with three options each plus the 24 source cells in about 3 GPU-hours. The profile is also stable, i.e., two halves of the 16 documents rank the cells almost identically (rank correlation 0.985).

LoopKV: Loop-Aware KV Cache Compression for Serving Looped Transformers

Allocation. Each cell takes one option o, which has stored bytes so , read bytes ro and damage dc,o . With KIVI-style codes (Liu et al., 2024b), which quantize keys per channel and values per token in groups of 32 with an fp16 scale and zero point, own INT4 codes store and read 2.5 KiB per token and layer and own INT2 codes 1.5 KiB, metadata included P (Appendix C). With xc,o ∈ {0, 1} and o xc,o = 1 for every cell, the allocation is min

P

s.t.

P

x

c,o xc,o ro c,o xc,o dc,o ≤ ε,

P

c,o xc,o so ≤ M,

(2)

uniform INT4 + sharing

180 / 240 KiB · KL 1.51 · 505 tok/s

read-optimized (knapsack)

215 / 215 KiB · KL 0.87 · 529 tok/s

memory-optimized (knapsack)

155 / 220 KiB · KL 1.80 · 535 tok/s

loop 1 loop 2 loop 3 loop 4

loop 1 loop 2 loop 3 loop 4

loop 1 loop 2 loop 3 loop 4 1

layer own INT4

a multiple-choice knapsack with two resource constraints (Kellerer et al., 2004). Because every option’s byte cost is a multiple of 0.5 KiB, a dynamic program over stored and read bytes solves this knapsack exactly over the 72 cells of loops 1–3. We keep loop 4, which every shared cell reads, at INT4, because coding it at INT2 alone would already cost 5.8× KLshare . Its cost is therefore a constant in every plan, and ε and M bound the damage and the memory of the whole cache. Two plans for the two ends of the frontier. The two budgets let us reach the two ends of the frontier with two plans. The read-optimized plan (LoopKV-R) minimizes reads under a damage budget of 1.19× KLshare , with no storage budget. It keeps INT4 in the sensitive middle layers and uses INT2 in 25 of the 72 cells, all in the first or last six layers (Figure 6). It shares no cell, because sharing saves storage but not reads. The memory-optimized plan (LoopKV-M) minimizes reads under a damage budget of 2.0× KLshare and a storage budget of 156 KiB per token, and it shares 26 cells, 20 of them in loop 3 as P2 predicts. Both budgets equal the damage of a residual-coding baseline8 . Both plans meet their targets, i.e., their measured KL stays below the prediction (Table 12), and they reach 6× and 8× BF16’s goodput with no significant accuracy loss (Figure 5). 3.2

C2: Cross-loop KV sharing as a per-cell option

Sharing is a per-cell option because its damage, measured as KL, differs from cell to cell and from model to model (P2, Figure 3). A shared cell (t, ℓ) attends with its own query over the final loop’s history at layer ℓ. Its descriptor names the final loop’s pool, so the cell does not store a copy of its own, while its sink and window stay its own and exact. When a 32-token page leaves the window, the cell releases its own copy, and its descriptor reads the final loop’s page for those positions instead. It stores nothing but reads the 8 ε and M are parameters that a user can set as needed. Here we set them to the damage of residual-coding baselines (Table 13) so that each plan can be compared with related work at equal damage. LoopKV itself stores no residuals.

own INT2

12 shared

24 INT4 anchor

Figure 6. The read-optimized plan beats the best per-loop allocation plan of LoopKV, i.e., INT4 in every loop with loop 3 sharing loop 4’s KV, which has the lowest damage of all per-loop plans that store less than uniform INT4, on KL and throughput (42% lower KL and 5% higher throughput). The memory-optimized plan stores 14% less than this plan and decodes 6% faster. At equal bytes, the read- and memory-optimized plans have 0.36× and 0.43× the measured KL of the best per-loop plans, i.e., INT4 in loops 1–2 with INT2 in loop 3, and INT4 in loop 1, INT2 in loop 2 and sharing in loop 3.

source’s INT4 pages like an own INT4 cell, so the knapsack shares only when the storage budget binds. 3.3

C3: Attention-aware V precision

P3 shows that only a few values receive high attention, about one history token in six, and that loops sharing a KV block attend to different tokens. V precision therefore lets every loop decide at read time which values to read at 4 bits and which at 2. Bit-plane V layout. To support this design, we change the V layout. Keys, which are always read in full, stay one plane, while an INT4 cell stores its value codes c ∈ {0, . . . , 15} as two contiguous 2-bit planes, hi=⌊c/4⌋ and lo=c mod 4, followed by the scales and zero points. The cell therefore stores as many bytes as plain INT4, and skipping a token’s lo plane saves 0.5 KiB of reads per token and layer. Read-time V precision at no extra read cost. For each history page, the kernel reads all key codes first. When it reaches the values, it therefore already knows the score sj of every token for this query in this loop, and thus which values receive low attention before it reads them. The key read that attention performs anyway supplies these scores, with no predictor, extra pass or extra bytes. Methods that also skip keys must estimate importance before reading, e.g., from page bounds (Tang et al., 2024), approximate scores (Ribar et al., 2024; Lee et al., 2024) or accumulated history (Zhang et al., 2023). LoopKV reads every key, so it

LoopKV: Loop-Aware KV Cache Compression for Serving Looped Transformers

decides from the scores attention computes and saves only V’s lo plane. Ideally, a token j is cold if its weight is below 2−τ times the largest in its row, i.e., sj < smax − τ ln 2. Because smax is costly to know before the row is complete, the kernel streams the row and tests sj < m − τ ln 2 against a running maximum m ≤ smax that starts at a seed m0 9 . A cold token reads only the hi plane and decodes to the midpoint of its bucket, ĉ = 4 hi + 1.5, while a hot token reads both planes and recovers its full INT4 code. We use τ =14, which changes the fp32 KL of KIVI-4 on Ouro-1.4B by a factor of only 1.016, whereas lowering τ raises the cost quickly, e.g., τ =10 already costs a factor of 1.35. Table 14 sweeps τ on the best per-loop allocation plan of LoopKV. Proposition 1 (Error bound) Let aj be the normalized attention weights of one row and head, C the set of tokens the seeded rule marks cold, and σj the V scale of token j’s channel group. Then, relative to the INT4 reader that reads both V planes of every token, the attention output changes by at most X ∥∆o∥∞ ≤ 1.5 aj σj ≤ 1.5 σmax |C| 2−τ amax . j∈C

Appendix E gives the proof. A design opportunity specific to looped transformers. This rule applies to any INT4 cache, KIVI-4 included (§5.3). In a looped transformer, however, each loop decides from its own attention scores when to read a value’s lo plane. We therefore store every value at full INT4 once and decide at read time how much of it each loop reads. An alternative that fixes the precision in advance would have to keep every token that any reading loop marks hot at full precision. On Ouro-1.4B, for example, loops 3 and 4 mark different hot tokens, and only 83% of the tokens are overlapped (Table 15). Across all four loops, such a fixed precision reads 1.6× the lo-plane bytes of deciding at read time. In the two plans, V precision keeps 17–22% of the history V reads per loop at full precision and cuts the reads per token by 13–14%, from 215 to 186 KiB for LoopKV-R and from 220 to 189 KiB for LoopKV-M. 3.4

The unified kernel

Layout. A page holds 32 tokens of one (loop, layer) cell. An INT4 page stores the keys as interleaved 4-bit codes, and it stores the values as the two planes of different positions of the original V matrix §3.3. An INT2 page stores its own 2-bit codes. A shared cell stores no pages, and its descriptor 9

The seed m0 is the largest score over the row’s 32 exact sink keys, which the kernel’s near blocks compute anyway (§3.4). As a lower bound on the row maximum, it only lowers the threshold, so the streaming rule never marks cold a token that the ideal rule keeps hot.

points to the corresponding pages of the final loop instead. Sinks and windows stay in vLLM’s BF16 blocks. There is one descriptor per cell, built when the plan is loaded and kept in device memory, and it carries all three of our designs, i.e., which pool to read (C2), how many bits are stored (C1) and whether to read the lo plane (C3). One launch per cell. In our implementation, the CUDA launch grid has up to 8 near thread blocks and up to 64 far thread blocks per decode row (Algorithm 1 in Appendix E). The near blocks compute exact BF16 attention over the sink and window and publish their partial results and the sink maximum. Each far block reads its share of the history pages through the descriptor and uses the published maximum as the seed of V precision. To avoid a long wait, a far block uses the published maximum only if it arrives within 4 µs and sets the seed to −∞ otherwise, which would mark every token as a hot token, so the precision will not lose by a fake positive. The last block of a row merges all partial results by their log-sum-exp, which would give the exact softmax, so one launch replaces up to four per layer call (§4). Because the descriptors and every buffer the kernel reads have fixed addresses, the whole decode-only step, including the 96 attention launches of its four loops, also replays as one CUDA graph. Window-exit quantization, whose page lists depend on the data, runs just before the replay.

4

I MPLEMENTATION

vLLM integration. LoopKV is a plugin to vLLM 0.26 (Kwon et al., 2023) and leaves the engine unmodified. It gives each loop its own KV-cache group, carves the history pools out of vLLM’s KV budget, and turns the plan table into the 96 descriptors at load time. When a page leaves the window, a Triton KIVI codec (Tillet et al., 2019) quantizes it into its cell’s pool and releases the BF16 block, so that more requests can be served at the same time. History reader. LoopKV’s CUDA history reader streams KV codes at 1.59–1.63 TB/s, as fast as FlashAttention-2 streams BF16 on the same GPU10 . With V precision, V is stored as two planes, and the reader fetches a value’s lo plane only for the tokens whose attention score is high enough. On real attention scores, this makes the reader 1.12–1.17× faster than the plain INT4 reader11 . Fused kernel and whole-step graph. Without fusion, each cell runs up to four kernels, i.e., FlashAttention-2 (Dao, 2024) over the sink and window, a seed kernel on INT4 cells, the history reader and a log-sum-exp merge, so 10

An RTX PRO 6000. The reader’s techniques are described in Appendix E. 11 In a microbenchmark on an RTX PRO 4500.

LoopKV: Loop-Aware KV Cache Compression for Serving Looped Transformers Table 1. On every model and suite, no LoopKV average accuracy falls significantly below BF16 (best in bold). The other methods in the table include KIVI-4 + V, i.e., uniform 4-bit quantization plus LoopKV’s V precision, naive uniform INT2 quantization of loops 1–3, and naive last-loop KV reuse. R16K and R32K denote RULER-16K and RULER-32K (Hsieh et al., 2024), LB denotes LongBench v1 (Bai et al., 2024), and L5 denotes AIME + MATH level 5 (Hendrycks et al., 2021). R16K R32K

LB

MATH500 GSM8K L5

Ouro-1.4B BF16 71.44 54.04 25.29 LoopKV-M 72.12 53.63 25.19 LoopKV-R 72.49 53.86 25.32 KIVI-4 + V 71.37 54.57 25.25 INT2, loops 1–3 65.50 — 24.43 Last-loop KV reuse 4.14 — 7.29

84.5 86.0 84.5 84.5 86.0 47.0

76.6 77.1 76.6 76.9 77.0 76.8

70.6 71.1 70.3 72.0 69.3 —

Ouro-2.6B BF16 LoopKV-M LoopKV-R KIVI-4 + V

80.73 80.98 81.00 80.61

— — — —

33.70 34.74 34.17 34.07

89.5 89.5 90.5 90.5

79.7 79.5 79.2 79.5

— — — —

Dory-4B BF16 (last loop) LoopKV-M LoopKV-R

93.54 94.23 94.09

— — —

32.43 32.37 32.62

91.5 91.7 91.5

88.6 88.2 87.9

— — —

a decode step issues 504 host launches in all. LoopKV’s fused kernel combines these four kernels into a single launch, which takes their GPU time within 2% and cuts the launches to 337 (LoopKV-R) or 332 (LoopKV-M). It also makes it possible to replay each decode-only step as one CUDA graph, which cuts the launches further to 26 and 17. In our experiments, fusion has little effect on KL, i.e., served KL is 0.99× that of the unfused path, and graph replay keeps it within 4% of eager execution.

5

E VALUATION

5.1

Setup

Models and hardware. We evaluate LoopKV on three looped models, Ouro-1.4B and Ouro-2.6B (Zhu et al., 2025) and Dory-4B (NVIDIA, 2026), and add Huginn-3.5B (Geiping et al., 2025) for the planning results in fp32 (§5.7). For Ouro, we use the base checkpoint for retrieval and GSM8K and the Thinking checkpoint, on which the plans are profiled, for everything else. We serve with vLLM 0.26.0 (Dory with its own vLLM fork) on NVIDIA RTX PRO 6000, RTX PRO 4500 and A100-80 GPUs, and each table and figure names the GPU it uses (configuration in Appendix B). Throughput is measured on the full system, with our fused kernel and the whole-step graph. We use RULER-16K, LongBench, MATH500 and GSM8K to measure serving quality, and we also use RULER-32K and long reasoning (AIME + MATH level 5) to evaluate long-context and reasoning abilities.

Quality. Served tasks are paired against the served BF16 model with 95% bootstrap CIs. The public suites are official RULER (Hsieh et al., 2024) (13 tasks of 100 items, at 16K and 32K), the 16 English and code subsets of LongBench v1 (Bai et al., 2024), MATH500 (Hendrycks et al., 2021; Lightman et al., 2024) (first 200 items, greedy, 8K output cap), GSM8K (Cobbe et al., 2021) (1,319 items, 3-shot) and sampled long reasoning on AIME and MATH level-5 problems. Appendix A adds two stress tests at 16K and details long reasoning. No plan was selected on any task, since the plans come from the fp32 KL knapsack (Appendix A checks its documents’ overlap with MATH500). We use fp32 KL (§2.3) as the sensitivity measure because KL and answer flips expose changes that accuracy averages hide (Dutta et al., 2024), and non-inferiority means that a suite average’s lower CI bound lies above −3 pp. Throughput and SLO. We measure peak decode throughput by sending B requests at context length N ∈ {8K, 16K, 32K} to the vLLM server, which decode 128 tokens together, and report the best throughput over the values of B at which no request is queued. We evaluate the SLO workload by sending 60 prompts from AIME 2024 (the American Invitational Mathematics Examination) (Mathematical Association of America, 2024) to the Thinking checkpoint at Poisson arrivals, the standard model of independent user requests in LLM serving (Kwon et al., 2023; Zhong et al., 2024). A request meets the SLO if its time to first token (TTFT) is at most 10 s and its mean time per output token (TPOT) is at most 3τ0 .12 The goodput (Zhong et al., 2024) is therefore the highest tested arrival rate at which 95% of requests meet the SLO. Baselines and ablations. We compare with BF16, vLLM’s FP8 cache (FlashInfer (Ye et al., 2025)), KIVI4 and KIVI-2 (uniform INT4 and INT2 on all loops, groups of 32), FlashLoop (official code and our reproduction) and ResidualQuant (reimplemented). To the best of our knowledge, FlashLoop and ResidualQuant have no vLLM implementation, so we compare them on quality and bytes stored for one token. Cross-loop KV sharing itself is adopted by some of the models: Dory is trained to read the last loop’s KV in every loop (NVIDIA, 2026), and Mixture-ofRecursions shares KV across recursions (Bae et al., 2025b). We therefore also compare with BF16 + loop-3 sharing, FP8 + loop-3 sharing and the hand-made KIVI-4 + loop-3 sharing plan, in which loop 3 reads loop 4’s KV, the placement that P2 finds cheap. Configurations that use our allocation (C1) or V precision (C3) are ablations. KIVI-4 runs on the 12

τ0 is the decode step time of stock vLLM on the RTX PRO 6000 predicted by a linear fit of step time to context length (17.31 ms plus 0.495 ms per 1K tokens). The fit gives τ0 = 18.60 ms or 20.69 ms under two rules for the context length, i.e. 3τ0 = 55.8 ms or 62.1 ms, and we report both.

KIVI-4

LoopKV-R (ours)

BF16

FP8 + sharing

FP8

KIVI-4 + our V (ablation)

LoopKV-M (ours)

KIVI-4 + V

LoopKV-R

300

14

14

300

200 13

100

100

6

12

16

24

32

48

64

80

time per output token (ms, log scale; batch grows to the right)

Figure 7. LoopKV keeps more requests resident: in this sweep it matches or beats the other configurations at each time per token and keeps scaling after they run out of KV memory. Ouro-1.4B on one RTX PRO 6000: decode tok/s against time per output token as the batch grows, up to each configuration’s largest resident batch (numbers); INT4 configurations on the fused kernel and wholestep graph, BF16 and FP8 as vLLM serves them; time per token on a log scale; the dotted line marks the SLO’s time-per-token bound (3τ0 , §5.1). KIVI-4 + our V is an ablation.

same fused kernel and whole-step graph as LoopKV and is our main baseline: it is the strongest prior method we serve, and the gains we report are over it. KIVI-4 + V adds our V precision (C3) to it and is the ablation step before allocation (C1) and sharing (C2). On our bit-plane INT4 pages (§3.3), KIVI-4 reads both V planes of every token, as KIVI does. 5.2

200 0

0

Quality

No LoopKV plan loses accuracy significantly on any public suite average (Table 1). On Ouro-1.4B, LoopKV-M is non-inferior to BF16 on all six public suites under a 3point margin and LoopKV-R on five, its MATH500 interval reaching −4.0 points on 200 items (Appendix A lists every interval). LoopKV-R even gains on official RULER-16K by about one point, a change carried by a single task. LoopKV-M stores 35% less than KIVI-4 at about twice the fp32 KL of LoopKV-R, yet it sustained the accuracy over the final scores of the 6 public suites. LoopKV-R, which shares no cell, keeps variable tracking, with no RULER-32K task changing significantly, and still sustains 1.3–1.6× KIVI-4’s goodput (§5.4); it is the plan to serve when long-context tracking matters. KIVI-4 + V and FP8 + loop-3 sharing behave the same way, staying within noise on the RULER, LongBench, MATH500 and GSM8K averages, although each loses variable tracking on RULER-16K. Allocation alone leaves answers as stable as uniform quantization does: LoopKV-R changes about

181 242 268 289 303 89

300 200 100 0

(c) Ouro-2.6B, 8K

24

16

16K context

(b) Ouro-1.4B, 16K

(d) Ouro-2.6B, 16K 85 110 118 124 137

12 TPOT bound

uniform INT4 LoopKV-M

150 100 50

31

200

600 400 200 0

171 358 479 528 556 577

25

0 decode tok/s

34

168 208 222 240 264

400

(a) Ouro-1.4B, 8K

48

36 34

81

8K context

tok/s

600

BF16

tok/s

decode tok/s

LoopKV: Loop-Aware KV Cache Compression for Serving Looped Transformers

0

Figure 8. LoopKV’s storage saving becomes throughput where memory binds. Peak decode tok/s at the resident peak for Ouro1.4B (a, b) and Ouro-2.6B (c, d) on an RTX PRO 6000, each GPU’s arms in one session, fused kernel and graph for INT4 arms; Ouro-1.4B within 1.1% of the separate runs quoted in the text (repeat means: §5.3); 32K and a 32 GB GPU: Tables 8, 7 and 21; Dory-4B: Table 20.

as many RULER-16K answers as KIVI-4 + V (Table 4). FlashLoop’s official code, in contrast, loses 12.75 points against its dense path on the same RULER-16K items in a run on an H100, mostly on variable tracking and multi-key needles. Its residual codec alone loses nothing in our earlier reproduction (Appendix A), which points to its sparse, lazily updated attention as the cause. Allocation buys quality where bytes are scarce. Uniform INT2 in loops 1–3 stores 8% more than LoopKV-M (168 vs 155 KiB) yet loses 5.9 points on RULER-16K, where LoopKV-M stays within noise (Table 1); a hand-made plan of LoopKV-M’s storage has 2.3× its fp32 KL (4.21 vs 1.80; Table 12). At KIVI-4’s own bytes, the knapsack’s minimum-damage plan only matches KIVI-4’s KL (0.47 vs 0.46; Table 12). 5.3

Throughput

LoopKV-M decodes 575.9 and 301.9 tok/s at 8K and 16K, 3.37× and 3.39× BF16, while storing 46% less KV than FP8 with sharing. LoopKV-R reaches 550.2 tok/s at 8K, 3.22× BF16 (all peaks: Figure 8 and Table 8 in Appendix B). LoopKV-M decodes 1.20× as fast as KIVI-4 at 8K: V precision gives about half of the gain, and allocation and sharing the rest. On the fused path on one RTX PRO 6000, V precision, which needs no plan, lifts uniform INT4 by 10–12% at 1.06× its served KL (Figure 8a, b). At the resident peak, allocation alone adds 5.6–7.6% to KIVI-4 + V (LoopKV-R) and allocation with sharing 8.4% at 8K, 12.5% at 16K and 15.6% at 32K (LoopKV-M; means of three repeats at 8K and 16K, one run at 32K, Table 8). The gains keep their shape on a second GPU type: on an A100-80, LoopKV-M decodes 2.2×, 2.2× and 4.7× faster than BF16 at 8K, 16K and 32K and 1.21–1.30× faster than

LoopKV: Loop-Aware KV Cache Compression for Serving Looped Transformers

uniform INT4 (Table 7). The mechanism is residency (Figure 7), since at KIVI-4 + V’s peak batch the plans gain 4.3–6.1% from fewer reads and the rest from the requests that the saved storage keeps resident; at 16K LoopKV-M peaks at batch 24 against 14, and at a matched time per token both plans remain faster than KIVI-4 + V (Appendix B). Where memory binds, the gain grows further. On a 32 GB RTX PRO 4500, LoopKV-M keeps 4 requests resident at 32K against uniform INT4’s 2 and decodes 29% faster than KIVI-4 + V, against 3–4% at equal batch (Table 21 in Appendix H).

On an H100. A port of the fused kernel to an H100 SXM (80 GB; tensor-core key path; stored codes byte-identical and served KL within 2% of ours) keeps the ordering: LoopKV-M decodes 2.69× as fast as BF16 at 8K and at 16K, against 2.27× and 1.96× for KIVI-4 and 2.08× and 2.05× for FP8, and served MATH500 and GSM8K stay within noise of BF16 for every INT4 arm (Table 10). On this GPU, V precision adds no speed: the H100 reader is bound by latency and instructions rather than bandwidth, and a variant that fetches the low plane only for attended tokens reads 17% fewer bytes per page but runs no faster. 5.4

SLO capacity (goodput)

LoopKV-M sustains 1.7–2.2× the SLO capacity of KIVI4 on the same GPU. On the second RTX PRO 6000, which ran every SLO arm, LoopKV-M sustains 2.00×, 1.72× and 2.20× KIVI-4’s rate in the three runs at τ0 =18.60 ms and 1.80–2.00× at 20.69 ms (Figure 11, Table 9); over the ablation KIVI-4 + V, which adds our V precision, the ratios are 1.60×, 1.72× and 1.83×. Again the mechanism is residency. Wherever both arms ran the same trace at the same rate, only KIVI-4 + V failed; at 0.0714 req/s it preempts 100 times and misses the SLO for 11.6% of requests on time to first token, while LoopKV-M preempts nothing. The ratio to KIVI-4 + V holds under resampling and grows as the time-per-token bound loosens: at τ0 =18.60 ms it stays at 1.5× or more in at least 86% of 10,000 bootstrap draws in every run, and with time per token relaxed to 5τ0 it rises to at least 1.80× in every run (Appendix B). LoopKV-R sustains 1.3–1.6× KIVI-4’s rate at 18.60 ms, and on an A100-80 the plans sustain 1.20–1.56× uniform INT4’s (Appendix B). FP8 and BF16, as vLLM serves them with piecewise graphs, are limited by time per token, and LoopKV-M sustains 2.5× and 8.0× their rates (2.8× and 7.5× at 20.69 ms); FP8 + loop-3 sharing sustains less than plain FP8 at 18.60 ms (3.1× below LoopKV-M), because its pointer path costs time per token. V precision alone moves uniform INT4’s capacity by one rate step or none.

5.5

The quality–throughput frontier

KIVI-4, KIVI-4 + V, LoopKV-R and LoopKV-M form the measured frontier from 0.46× to 1.8× KLshare , and the per-layer optimized baseline KIVI-4 + loop-3 sharing plan and every residual design lie behind it (Figure 9; Figure 10 shows one point per comparable method and omits KIVI-4 + V). 5.6

Ablation

Each contribution buys a different resource at a stated KL (Table 6): allocation (C1) cuts reads by 10% while KL rises from 0.46× to 0.87× KLshare , sharing (C2), with a larger damage budget, cuts storage by 28% and raises the resident batch from 36 to 48, and V precision (C3) cuts reads by 14% at unchanged KL. Fusion and the graph buy latency more than peak throughput, cutting host launches from 504 to 17 for LoopKV-M and speeding up batch-1 decoding of the two plans by 1.07–1.19× at 8K and 16K, against 0.5–2.5% at the peak. Sharing alone buys only part of the capacity: with V precision on the same fused path, the hand-made KIVI-4 + loop-3 sharing plan (180 KiB stored per token, against LoopKV-M’s 155) sustains 0.75× LoopKV-M’s 16K SLO capacity, the same as LoopKV-R. Comparison with codec baselines. No other codec reading less than 216 KiB, TurboQuant-3.5 (Zandieh et al., 2025) and KVQuant-3b (Hooper et al., 2024) included, reaches the read-optimized plan’s KL in our implementations (Figure 13 in Appendix J). The loop axis matters as well: restricting the knapsack to one option per layer, as layer-wise allocators do (Li et al., 2025a), raises the KL at equal bytes by 1.52× at the read-optimized plan’s bytes and by 1.43× at the memory-optimized plan’s bytes with loop 3 shared (Appendix D). 5.7

Generality: Ouro-2.6B, Dory and Huginn

On Ouro-2.6B, loop 1 again dominates, damages add up, and the memory-optimized plan computed from its own profile has 0.48× the KL of the hand-made plan of equal bytes (Appendix F). Served by us on the RTX PRO 6000, Ouro-2.6B’s LoopKV-M decodes 3.25×, 4.44× and 11.0× faster than BF16 at 8K, 16K and 32K, because BF16 holds only 6, 3 and 1 requests (Table 7 adds an A100-80). No Ouro-2.6B arm loses accuracy significantly on any suite average (Table 1). On Dory-4B and Huginn-3.5B, where the last loop carries 79–93% of the INT2 damage (P1), the knapsack at 10% fewer bytes than uniform INT4 has 1.33× uniform INT4’s KL on Dory, against 2.00× for INT2 in the first loops and 29.3× in the last; Huginn shows the same order against the first loops (Table 18). Dory’s trained lastloop mode, in which all loops read loop 4’s KV, confirms the cross-loop redundancy that LoopKV exploits post hoc.

LoopKV: Loop-Aware KV Cache Compression for Serving Looped Transformers

LoopKV’s descriptor path and fused reader serve this cache with per-plane plans from the same knapsack, which pass the in-serving byte checks and show no significant loss on any suite average (Table 1); on a 32 GB GPU, LoopKV-M decodes 3% and 13% faster than Dory’s BF16 cache at 8K and 16K, and the plans hold their KL budget, whereas INT4 on all nine full-attention planes buys further speed with a 35–52% higher served KL (Appendix G).

6

R ELATED W ORK

Looped models and cross-loop compression. In looped models, later loops refine earlier ones (Blayney et al., 2026). Several designs train the model so that loops share their KV: MELT keeps one gated cache per layer (Conchello Vendrell et al., 2026), the Parallel Loop Transformer reuses the first loop’s KV in every loop (Wu et al., 2025), and pretraining in which only the first recursion writes KV cuts context memory by 76–79% (Monea et al., 2026); these change the model, whereas LoopKV serves the released model. Looped latent attention keeps the model frozen and learns a lowrank cross-loop codec by distillation (O’Neill & Reid, 2026); LoopKV needs no training, only a profile. Without retraining, Huginn reads and writes a fixed number of cache slots across recurrence steps (Geiping et al., 2025), and a parallel sampler for it matches GSM8K accuracy with a single slot (Geiping et al., 2026); Huginn was trained so that keys and values from different depths are compatible, whereas on Ouro such reuse loses accuracy over long outputs (P2), and WaveFront Decoding, a self-speculative decoder for looped models, keeps Ouro’s per-recurrence cache for the same reason (Ha & Kim, 2026). FlashLoop introduced closed-loop residual coding of cross-loop KV with lazily updated sparse attention in a batch-1 research harness (Yang & Liu, 2026), and ResidualQuant stores 2-bit residuals against an INT4 final loop, a fixed loop-wise precision rule, and reports up to 4.15× higher peak throughput from larger batches on an RTX 5090 with its own kernels (Kim et al., 2026); we reimplement its codec in our harness. Building on their observation that adjacent loops are similar, LoopKV measures where this pays (P1, P2) and what a looped step reads (P3), serving the released model unchanged. KV quantization, sharing and eviction. Uniform-bit codecs (Liu et al., 2024b; Hooper et al., 2024; Ashkboos et al., 2024; Zandieh et al., 2025; Zhou et al., 2026; Lin et al., 2025) and rate–distortion allocation (Zhang et al., 2026) follow sensitivity-based mixed precision (Dong et al., 2019), and KVTuner (Li et al., 2025a) and AutoQuant-KV, which picks one K/V format per layer by KL sensitivity under an average-bit budget (NVIDIA Model Optimizer Team, 2026), allocate over layers; LoopKV adds the loop axis and a sharing option, so its knapsack separates stored from read bytes. MiKV (Yang et al., 2024a) and

ZipCache (He et al., 2024) set token precision at store time by importance heuristics, and DiffKV (Zhang et al., 2025) treats keys, values, tokens and heads differently. SpAtten fetches low-order bits and recomputes only when attention from the high-order bits is flat (Wang et al., 2021), and AnyPrecision LLM and Matryoshka Quantization nest lower precisions in a weight code’s high-order bits (Park et al., 2024; Nair et al., 2025); LoopKV stores INT4 values as two such planes, so a cold token reads only the high plane, and decides per loop and token from the exact scores its fused kernel already computes. AlignedKV also chooses how many bits of each value to read from the scores of the same attention computation, to cut memory access to an FP16 cache (Tan et al., 2024), and QuantSpec stores INT8 KV as two INT4 parts so that a draft model reads the upper part and the target reads both (Tiwari et al., 2025); in LoopKV the readers of one INT4 value are the loops of one model. Concurrent ReadKV also reads key and value prefixes of a progressive code according to the query and its attention (Bhandari et al., 2026); LoopKV applies the value side per loop, inside the fused kernel that serves the knapsack’s plan. Cross-layer sharing, trained (Brandon et al., 2024; Sun et al., 2024; Wu & Tu, 2024) or post hoc (Liu et al., 2024a; Chang et al., 2025; Shutova et al., 2025; Yang et al., 2024b; Qiao et al., 2024), aliases layers; LoopKV aliases loops, cell by cell. Eviction and page selection (Xiao et al., 2024; Zhang et al., 2023; Li et al., 2024; Ge et al., 2024; Tang et al., 2024; Li et al., 2025b) drop tokens, while LoopKV keeps StreamingLLM’s sink and window, evicts nothing and is evaluated in the engine (Gao et al., 2025).

7

C ONCLUSION

LoopKV turns three properties of Ouro’s looped KV into allocation, sharing and V precision, executed by one fused kernel. On Ouro-1.4B, on the same kernel and graph as uniform INT4 (KIVI-4), LoopKV-M stores 35% less KV and sustains 1.7–2.2× KIVI-4’s goodput under a 16K latency SLO across three runs, and LoopKV-R sustains 1.3–1.6×. No public suite average of either plan falls significantly below BF16’s. The lesson is that a looped model’s KV is a grid, not a stack, in which a cheap profile and an exact knapsack decide once per model where each cell spends its bits and whose KV it reads. We will release the plugin, the kernels, the plans and the profiles. Limitations. LoopKV targets looped models that keep a KV cache per loop: one such family, Ouro, is served in full; as transfer evidence, Dory-4B is served in its trained last-loop mode and Huginn-3.5B is evaluated in fp32. We measure speed on RTX PRO 6000, RTX PRO 4500 and A100-80 GPUs, and SLO capacity on one decode-heavy

LoopKV: Loop-Aware KV Cache Compression for Serving Looped Transformers

workload (60 AIME prompts at 16K, up to three runs). On RULER-32K, LoopKV-M loses variable tracking (−21.2 points, 95% CI [−29.6, −12.8]) while its average stays within noise; LoopKV-R keeps it.

R EFERENCES Ashkboos, S., Mohtashami, A., Croci, M. L., Li, B., Cameron, P., Jaggi, M., Alistarh, D., Hoefler, T., and Hensman, J. QuaRot: Outlier-free 4-bit inference in rotated LLMs. In Advances in Neural Information Processing Systems (NeurIPS), 2024. Bae, S., Fisch, A., Harutyunyan, H., Ji, Z., Kim, S., and Schuster, T. Relaxed recursive transformers: Effective parameter sharing with layer-wise LoRA. In International Conference on Learning Representations (ICLR), 2025a. Bae, S., Kim, Y., Bayat, R., Kim, S., Ha, J., Schuster, T., Fisch, A., Harutyunyan, H., Ji, Z., Courville, A., and Yun, S.-Y. Mixture-of-recursions: Learning dynamic recursive depths for adaptive token-level computation. arXiv preprint arXiv:2507.10524, 2025b. Bai, Y. et al. LongBench: A bilingual, multitask benchmark for long context understanding. In Annual Meeting of the Association for Computational Linguistics (ACL), 2024. Bhandari, S., Gretta, L., Balasubramanian, K., et al. Read what matters: Query-adaptive quantization for KV caches. arXiv preprint arXiv:2610.11245, 2026. Blayney, H., Arroyo, Á., Obando-Ceron, J., Castro, P. S., Courville, A., Bronstein, M. M., and Dong, X. A mechanistic analysis of looped reasoning language models. arXiv preprint arXiv:2604.11791, 2026. Brandon, W., Mishra, M., Nrusimha, A., Panda, R., and Ragan-Kelley, J. Reducing transformer key-value cache size with cross-layer attention. In Advances in Neural Information Processing Systems (NeurIPS), 2024. Chang, C.-C. et al. xKV: Cross-layer SVD for KV-cache compression. arXiv preprint arXiv:2503.18893, 2025. Cobbe, K. et al. Training verifiers to solve math word problems. arXiv preprint arXiv:2110.14168, 2021. Conchello Vendrell, V., Padres Masdemont, A., Grillo, N., et al. Memory-efficient looped transformer: Decoupling compute from memory in looped language models. arXiv preprint arXiv:2605.07721, 2026. Dao, T. FlashAttention-2: Faster attention with better parallelism and work partitioning. In International Conference on Learning Representations (ICLR), 2024.

Dehghani, M., Gouws, S., Vinyals, O., Uszkoreit, J., and Kaiser, Ł. Universal transformers. In International Conference on Learning Representations (ICLR), 2019. Dong, Z., Yao, Z., Gholami, A., Mahoney, M. W., and Keutzer, K. HAWQ: Hessian AWare quantization of neural networks with mixed-precision. In IEEE/CVF International Conference on Computer Vision (ICCV), 2019. Dutta, A., Krishnan, S., Kwatra, N., and Ramjee, R. Accuracy is not all you need. In Advances in Neural Information Processing Systems, 2024. Gao, W., Zhou, X., Sun, P., Zhang, T., and Wen, Y. Rethinking key-value cache compression techniques for large language model serving. In Proceedings of Machine Learning and Systems (MLSys), 2025. Ge, S., Zhang, Y., Liu, L., Zhang, M., Han, J., and Gao, J. Model tells you what to discard: Adaptive KV cache compression for LLMs. In International Conference on Learning Representations (ICLR), 2024. Geiping, J., McLeish, S., Jain, N., Kirchenbauer, J., Singh, S., Bartoldson, B. R., Kailkhura, B., Bhatele, A., and Goldstein, T. Scaling up test-time compute with latent reasoning: A recurrent depth approach. In Advances in Neural Information Processing Systems (NeurIPS), 2025. Geiping, J., Yang, X., and Su, G. Efficient parallel samplers for recurrent-depth models and their connection to diffusion language models. In Proceedings of the 43rd International Conference on Machine Learning (ICML), volume 306 of Proceedings of Machine Learning Research, pp. 34442–34462, 2026. Ha, H. and Kim, J.-J. WaveFront decoding: Parallelized self-speculative decoding for looped language models. arXiv preprint arXiv:2609.23033, 2026. He, Y., Zhang, L., Wu, W., Liu, J., Zhou, H., and Zhuang, B. ZipCache: Accurate and efficient KV cache quantization with salient token identification. In Advances in Neural Information Processing Systems (NeurIPS), 2024. Hendrycks, D., Burns, C., Kadavath, S., Arora, A., Basart, S., Tang, E., Song, D., and Steinhardt, J. Measuring mathematical problem solving with the MATH dataset. In NeurIPS Datasets and Benchmarks Track, 2021. Hooper, C., Kim, S., Mohammadzadeh, H., Mahoney, M. W., Shao, Y. S., Keutzer, K., and Gholami, A. KVQuant: Towards 10 million context length LLM inference with KV cache quantization. In Advances in Neural Information Processing Systems (NeurIPS), 2024.

LoopKV: Loop-Aware KV Cache Compression for Serving Looped Transformers

Hsieh, C.-P., Sun, S., Kriman, S., Acharya, S., Rekesh, D., Jia, F., Zhang, Y., and Ginsburg, B. RULER: What’s the real context size of your long-context language models? In Conference on Language Modeling (COLM), 2024. Kellerer, H., Pferschy, U., and Pisinger, D. Knapsack Problems. Springer, 2004. Kim, H., Lee, J., Cho, S., Han, D., Han, I., and Kim, S. ResidualQuant: KV cache quantization for looped transformers with 2-bit residuals. arXiv preprint arXiv:2610.10381, 2026. Kwon, W., Li, Z., Zhuang, S., Sheng, Y., Zheng, L., Yu, C. H., Gonzalez, J. E., Zhang, H., and Stoica, I. Efficient memory management for large language model serving with PagedAttention. In Proceedings of the 29th Symposium on Operating Systems Principles (SOSP), 2023. Lee, W., Lee, J., Seo, J., and Sim, J. InfiniGen: Efficient generative inference of large language models with dynamic KV cache management. In USENIX Symposium on Operating Systems Design and Implementation (OSDI), 2024. Li, X. et al. KVTuner: Sensitivity-aware layer-wise mixed precision KV cache quantization for efficient and nearly lossless LLM inference. In International Conference on Machine Learning (ICML), 2025a.

Mathematical Association of America. American invitational mathematics examination (AIME) 2024. https://maa.org/math-competitions/ american-invitational-mathematics-examination-aim 2024. Monea, G., Ramji, K., El-Kurdi, Y., Lastras, L. A., Artzi, Y., Godey, N., and Fernandez Astudillo, R. The surprising effectiveness of shared memory in looped transformers. arXiv preprint arXiv:2610.02383, 2026. Nair, P. A., Datta, P., Dean, J., Jain, P., and Kusupati, A. Matryoshka quantization. In International Conference on Machine Learning (ICML), 2025. NVIDIA. NVIDIA-Nemotron-Labs-Dory1.04B. https://huggingface.co/nvidia/ NVIDIA-Nemotron-Labs-Dory1.0-4B-BF16, 2026. Model card. NVIDIA Model Optimizer Team. AutoQuant KV: Automatic mixed-precision KV cache quantization. https://nvidia.github. io/Model-Optimizer/announcements/ 20261008-autoquant-kv.html, 2026. Announcement, October 8, 2026. O’Neill, J. and Reid, F. Looped latent attention: Cross-loop KV compression for looped transformers. arXiv preprint arXiv:2607.15456, 2026.

Li, Y., Jiang, H., Wu, Q., Luo, X., Ahn, S., Zhang, C., Abdi, A. H., Li, D., Gao, J., Yang, Y., and Qiu, L. SCBench: A KV cache-centric analysis of long-context methods. In International Conference on Learning Representations (ICLR), 2025b.

Park, Y., Hyun, J., Cho, S., Sim, B., and Lee, J. W. Any-precision LLM: Low-cost deployment of multiple, different-sized LLMs. In International Conference on Machine Learning (ICML), 2024.

Li, Y. et al. SnapKV: LLM knows what you are looking for before generation. In NeurIPS, 2024.

Pope, R., Douglas, S., Chowdhery, A., Devlin, J., Bradbury, J., Heek, J., Xiao, K., Agrawal, S., and Dean, J. Efficiently scaling transformer inference. In Proceedings of Machine Learning and Systems (MLSys), 2023.

Lightman, H. et al. Let’s verify step by step. In International Conference on Learning Representations (ICLR), 2024. Lin, Y., Tang, H., Yang, S., Zhang, Z., Xiao, G., Gan, C., and Han, S. QServe: W4A8KV4 quantization and system co-design for efficient LLM serving. In Proceedings of Machine Learning and Systems (MLSys), 2025. Liu, A., Liu, J., Pan, Z., He, Y., Haffari, G., and Zhuang, B. MiniCache: KV cache compression in depth dimension for large language models. In Advances in Neural Information Processing Systems (NeurIPS), 2024a. Liu, Z., Yuan, J., Jin, H., Zhong, S., Xu, Z., Braverman, V., Chen, B., and Hu, X. KIVI: A tuning-free asymmetric 2bit quantization for KV cache. In International Conference on Machine Learning (ICML), 2024b.

Qiao, A. et al. SwiftKV: Fast prefill-optimized inference with knowledge-preserving model transformation. arXiv:2410.03960, 2024. Ribar, L., Chelombiev, I., Hudlass-Galley, L., Blake, C., Luschi, C., and Orr, D. SparQ attention: Bandwidthefficient LLM inference. In International Conference on Machine Learning (ICML), 2024. Rouhani, B. D. et al. Microscaling data formats for deep learning. arXiv preprint arXiv:2310.10537, 2023. Saunshi, N., Dikkala, N., Li, Z., Kumar, S., and Reddi, S. J. Reasoning with latent thoughts: On the power of looped transformers. In International Conference on Learning Representations (ICLR), 2025.

LoopKV: Loop-Aware KV Cache Compression for Serving Looped Transformers

Shutova, A. et al. Cache me if you must: Adaptive key-value quantization for large language models. In International Conference on Machine Learning (ICML), 2025.

Yang, W. and Liu, S. FlashLoop: Fast and memory-efficient looped transformers via lazy updates. arXiv preprint arXiv:2609.29812, 2026.

Sun, Y. et al. You only cache once: Decoder-decoder architectures for language models. In NeurIPS, 2024.

Yang, Y. et al. KVSharer: Efficient inference via layer-wise dissimilar KV cache sharing. arXiv:2410.18517, 2024b.

Tan, Y., Wang, H., Yan, C., and Deng, Y. AlignedKV: Reducing memory access of KV-cache with precisionaligned quantization. arXiv preprint arXiv:2409.16546, 2024.

Ye, Z. et al. FlashInfer: Efficient and customizable attention engine for LLM inference serving. In Proceedings of Machine Learning and Systems (MLSys), 2025.

Tang, J., Zhao, Y., Zhu, K., Xiao, G., Kasikci, B., and Han, S. Quest: Query-aware sparsity for efficient long-context LLM inference. In International Conference on Machine Learning (ICML), 2024. Tillet, P., Kung, H. T., and Cox, D. Triton: An intermediate language and compiler for tiled neural network computations. In Proceedings of the 3rd ACM SIGPLAN International Workshop on Machine Learning and Programming Languages (MAPL), 2019. Tiwari, R., Xi, H., Tomar, A., Hooper, C. R. C., Kim, S., Horton, M., Najibi, M., Mahoney, M. W., Keutzer, K., and Gholami, A. QuantSpec: Self-speculative decoding with hierarchical quantized KV cache. In Proceedings of the 42nd International Conference on Machine Learning (ICML), volume 267 of Proceedings of Machine Learning Research, pp. 59668–59686, 2025. Wang, H., Zhang, Z., and Han, S. SpAtten: Efficient sparse attention architecture with cascade token and head pruning. In IEEE International Symposium on HighPerformance Computer Architecture (HPCA), 2021. Williams, S., Waterman, A., and Patterson, D. Roofline: An insightful visual performance model for multicore architectures. Communications of the ACM, 52(4):65–76, 2009. Wu, B., Chen, M., Luo, X., et al. Parallel loop transformer for efficient test-time computation scaling. arXiv preprint arXiv:2510.24824, 2025. Wu, H. and Tu, K. Layer-condensed KV cache for efficient inference of large language models. In ACL, 2024. Xiao, G. et al. Efficient streaming language models with attention sinks. In ICLR, 2024. Yang, J. Y., Kim, B., Bae, J., Kwon, B., Park, G., Yang, E., Kwon, S. J., and Lee, D. No token left behind: Reliable KV cache compression via importance-aware mixed precision quantization. arXiv preprint arXiv:2402.18096, 2024a.

Zandieh, A. et al. TurboQuant: Online vector quantization with near-optimal distortion rate. arXiv:2504.19874, 2025. Zhang, J., Guo, H., Benini, L., and Li, Y. RDKV: Ratedistortion bit allocation for joint eviction and quantization of the KV cache. arXiv preprint arXiv:2605.08317, 2026. Zhang, Y., Hu, Y., Zhao, R., Lui, J. C. S., and Chen, H. DiffKV: Differentiated memory management for large language models with parallel KV compaction. In SOSP, 2025. Zhang, Z. et al. H2O: Heavy-hitter oracle for efficient generative inference of large language models. In NeurIPS, 2023. Zhong, Y., Liu, S., Chen, J., Hu, J., Zhu, Y., Liu, X., Jin, X., and Zhang, H. DistServe: Disaggregating prefill and decoding for goodput-optimized large language model serving. In USENIX Symposium on Operating Systems Design and Implementation (OSDI), 2024. Zhou, Z., Zhuang, D., Li, J., Chen, Z., Song, S. L., Athiwaratkun, B., and Wu, X. OSCAR: Offline spectral covariance-aware rotation for 2-bit KV cache quantization. arXiv preprint arXiv:2605.17757, 2026. Zhu, R.-J., Wang, Z., Hua, K., Zhang, T., Li, Z., Que, H., Wei, B., Wen, Z., Yin, F., Xing, H., et al. Scaling latent reasoning via looped language models. arXiv preprint arXiv:2510.25741, 2025.

LoopKV: Loop-Aware KV Cache Compression for Serving Looped Transformers

A

F ULL QUALITY RESULTS

Table 3 gives the ∆s of the suites run on the fused path, Table 2 those of the runs before fusion, including the in-house stress tests MQ641 and RULER-style (§5.1), and Table 5 those of Ouro-2.6B served; the remaining cells of Table 1 follow below. On the 300-item RULER-style set, the intervals are too wide for any quantized configuration to establish the 3-point margin. MQ641 has 641 multi-key items with seven-digit values at 16K in public PG19 books; its pre-registration was written before any run, its generator, seeds and hashes are released, its CIs bootstrap over about 131 book clusters, and an independent run reproduced it on another GPU (BF16 85.96%, KIVI-4 87.05, the two plans 86.27–86.43, the hand-made sharing plans 81.75–84.09). No plan was selected on it. In the independent runs, the memory-optimized plan scores +4.7 points above the hand-made INT4/INT2 + sharing plan of equal storage. Protocols. Long reasoning samples 194 AIME and MATH level-5 problems at temperature 1.0 and top-p 0.7, 6 times per AIME and 8 times per MATH problem, with a 16,384-token cap; its accuracy is the mean of the items’ sample accuracies, with CIs paired by item. In LongBench, narrativeqa, musique and passage_count are at the BF16 floor. The 16 documents of the fp32 KL profile are generations for 16 of the 200 MATH500 prompts, and removing them moves no pre-fusion MATH500 delta by more than 0.4 points. On the fused path, MATH500 caps each INT4 arm’s concurrency at its pool’s capacity, because without the cap the pages that overflow the pool stay in BF16 (a pool audit counted 15,008 for LoopKV-R); uncapped, each of these arms scored +2.5 , and the MATH500 runs of Table 2 predate the cap. Remaining cells of Table 1. On the fused path, long reasoning (AIME + MATH-L5) changes by −0.26 for LoopKV-R and +0.54 for LoopKV-M; before fusion it changed by +1.14 and −0.43 (Table 2), and an independent run of that version on H100 and H200 GPUs with the same protocol gives +1.25 and −0.62. On the A100-80, against BF16 on the same GPU, long reasoning changes by +1.10 for KIVI-4 + V and by −0.13 for FP8 + loop-3 sharing. On Ouro-2.6B, RULER-16K changes by −0.12 for KIVI-4 + V, +0.27 for LoopKV-R and +0.25 for LoopKV-M, and LongBench by +0.37, +0.47 and +1.05. On Dory-4B, against its last-loop BF16 cache, LoopKV-M and LoopKV-R change RULER-16K by +0.69 and +0.55, LongBench by −0.06 and +0.19, and GSM8K by −0.5 and −0.7; their MATH500 results are in Table 20. Replication on another RTX PRO 6000. An independent run on an RTX PRO 6000 Server Edition repeats the final-path quality suites against its own BF16. It passes the same gate first: decode KL within 2% of ours and no differing byte. Every comparison stays within noise. On MATH500-200 and GSM8K, the differences for KIVI-4 + V, LoopKV-R, LoopKV-M and FP8 + loop-3 sharing lie between −1.5 and +2.5 points, and LongBench for LoopKV-R is +0.00. On RULER-32K, LoopKV-M scores −0.78 and again loses variable tracking (−18.0) while gaining frequent-word extraction (+15.7), whereas KIVI-4 + V scores +0.62 and keeps variable tracking (−0.4). The variable-tracking loss therefore follows loop-3 sharing rather than V precision or the codec. FP8 + loop-3 sharing loses it too on an A100-80 (Table 3), whereas KIVI-4 + V, which shares nothing, keeps variable tracking at 32K, and its frequent-word gain (+5.7) is about a third as large. In our earlier reproduction, FlashLoop’s residual codec alone costs 0.0 points on MQ641 and its sparse, lazily updated attention alone −7.8 (full method −6.4). Answer flips. Following Dutta et al. (2024), Table 4 counts items whose score changed against BF16 on the same items. On RULER-16K LoopKV-R changes about as many items as KIVI-4 + V (70 vs 67 of 1,300), LoopKV-M 125 and FP8 + loop-3 sharing 146. Before fusion, LoopKV-M also lost two RULER-16K needle tasks (niah_multikey_2 −7.0, niah_multivalue −2.75); on the fused path neither loss is significant. Ouro-2.6B, served. Ouro-2.6B (48 layers) is served on the fused path (fused kernel, whole-step graph, τ =14, production seed) with the two plans that the same allocator computes from its profile (Table 19). GSM8K (all 1,319 items, base checkpoint, 3-shot, strict match) ran on A40 GPUs, with every arm of an item slice on the same GPU; we ran MATH500 (the same 200 items as Ouro-1.4B, Thinking checkpoint, greedy, 8K output cap) on one RTX PRO 6000, with each arm’s concurrency capped at its pool’s capacity for 8K requests, and the engine counted no full pool (Table 5). Before the A40 runs, served decode KL matched the H200 graph-path values of the independent run within 5% (1.00–1.03×; 1.008× / 1.000× for LoopKV-R, which ran later on other A40s against a BF16 re-run there), and their in-serving byte checks of 12 pages differ in 1, 3, 4 and 2 bytes for KIVI-4, KIVI-4 + V, LoopKV-M and LoopKV-R, all on the shared source pages, the rounding class also seen on the RTX PRO 6000 (Appendix I).

LoopKV: Loop-Aware KV Cache Compression for Serving Looped Transformers Table 2. Quality before kernel fusion, including the in-house stress tests MQ641 and RULER-style: neither knapsack plan loses significantly against BF16; read-optimized + V meets the 3-point margin on MQ641, MATH500, GSM8K and long reasoning but not on RULER-style. fp32 KL / KLshare (KLshare = 1.20×10−4 ; 16 documents); tasks: ∆ (pp) vs served BF16, + / − where the paired 95% CI lies above / below zero (single runs; greedy except long reasoning). “—” = not run. † Earlier RTX PRO 4500 harness, vs its own BF16. ‡ Codec only, no sparse path. FP8 / FP8 + loop-3 sharing: fp32 KL 0.19× / 1.20×, no runs here.

LoopKV

plan

fp32 KL

read-optimized + V precision memory-optimized + V precision

0.87 0.87 1.80 1.80

KIVI-4 KIVI-2 baselines residual, INT4 final + INT4 residual, INT4 final + INT2 FlashLoop (our reproduction) ResidualQuant (reimpl.) + sharing

KIVI-4 + loop-3 sharing (hand-made) + our V precision (ablation) BF16 + loop-3 sharing (KLshare )

MQ641 (stress)

RULER-style 16K

MATH500 GSM8K long reasoning (200) (1,319) (194)

−1.6 −1.1 −0.2 —

— +3.0 +2.0 +1.5

— −0.5 +0.8 —

— +1.14 — −0.43

0.46 +0.9 −1.7 21.5 −13 to −18† — 1.42 −0.8 −2.0 4.38 −3.0− — ‡ 0.30 −6.4† −8.2−† 2.01 −0.8† niah_mq −9.0−†

— — — — — —

— — — — — —

— — — — — —

+2.5 — +0.41†

+0.2 — +0.3†

— — −0.45

1.51 1.49 1.00

+0.8 +0.0 −0.8 +0.3

−1.6 −1.2 −1.6

−0.2 — +0.8

Table 3. Quality on the fused path (fused kernel, whole-step graph, V-precision seed m0 as served): ∆ (pp) vs served BF16, + / − where the paired 95% CI lies above / below zero; tasks listed if significant (unadjusted); ∗ BF16-floor task. FP8 + loop-3 sharing runs as vLLM serves it (no LoopKV kernel); its RULER-32K run is on an A100-80 against BF16 on the same GPU. RULER tasks: vt variable tracking, fwe frequent-words extraction, cwe common-words extraction, qa_1 single-hop QA; LongBench subsets by name. arm

official RULER-16K official RULER-32K

KIVI-4 + V

−0.07; vt −3.0

LoopKV-R LoopKV-M FP8 + loop-3 sharing (as served)

B

LongBench v1 (16)

−0.04; qasper −0.87, gov_report −0.64 +1.05+ ; qa_1 +7.0 −0.18 +0.03; multifieldqa_en +0.93 +0.68; fwe +5.0, −0.41; vt −21.2, fwe −0.10 qa_1 +8.0 +18.3, cwe +0.6∗ −0.22; vt −6.8, +0.29 (A100-80); vt +0.07; narrativeqa −0.44, cwe +1.5 −14.2, fwe +19.7 musique +1.64∗ +0.53

MATH500 GSM8K (200) (1,319)

MQ641 (stress)

+0.0

+0.3

+0.5

+0.0

+0.0

+0.6

+1.5

+0.5

+0.5

+2.0

−1.1

—

T HROUGHPUT AND SLO DETAILS

Part of V precision’s lift recovers the cost of the two-plane V layout: uniform INT4 decodes 479.4 tok/s with two planes on the fused path, against 494.4 with interleaved V on the reader before fusion. On the RTX PRO 6000, vLLM 0.26.0 runs with PyTorch 2.11 (CUDA 13), memory utilization 0.90. Table 8 lists the peak decode throughput of every configuration. The SLO prompts have a mean output of 10,999 tokens, and attainment is scored over the requests that arrive in the middle 80% of the arrival window, 43, 49 and 47 of 60 for seeds 101–103. The SLO’s τ0 comes from a stock-vLLM step-time fit of 17.31 ms plus 0.495 ms per 1K tokens of context, evaluated at a shorter workload’s mean context for the 18.60 ms rule and at this workload’s own for the 20.69 ms rule (at a maximum model length of 17,408), which therefore allow 3τ0 =55.8 ms and 62.1 ms per output token. Capacity uses the 95% rule, fixed before the ladders ran. Table 7 repeats the own-peak protocol for Ouro-2.6B and on a second GPU type, an A100-80. On the A100-80 and for Ouro-2.6B, every arm of a context runs on one physical GPU, so each ratio is a same-card ratio. Ouro-2.6B’s plugin arms need a smaller code pool than the plugin’s default share to stay resident up to their formula batch (0.80, and 0.75 for the memory-optimized plan at 8K). At 32K, Ouro-1.4B’s plugin arms use a share of 0.88, as on the RTX PRO 6000. Under the SLO protocol of §5.1 on one A100-80, with this GPU’s τ0 rules of 16.60 and 26.34 ms and every arm on this GPU, the

LoopKV: Loop-Aware KV Cache Compression for Serving Looped Transformers Table 4. Answer flips against BF16: items whose score went down / up (graded suites: items whose official score changed, down / up). † Before fusion. Long reasoning: items whose sample accuracy went down / up. suite (items) RULER-16K (1,300) RULER-32K (1,300) LongBench v1 (3,750) MATH500 (200) GSM8K (1,319) long reasoning (194) MQ641 (641)

KIVI-4 + V LoopKV-R LoopKV-M 30 / 37 37 / 47 499 / 480 8/8 21 / 25 — 13 / 16

26 / 44 44 / 45† 513 / 540 9/9 22 / 22 34 / 39† 12 / 16

KIVI-4† KIVI-4 + loop-3 sharing† FP8 + loop-3 sharing

54 / 71 29 / 31 84 / 89 40 / 47 556 / 596 483 / 495 7 / 10 — 27 / 34 — 33 / 38† — 18 / 21 9 / 15

45 / 87 — 579 / 570 9 / 14 27 / 30 — 29 / 19

66 / 80 — 681 / 719 5/9 59 / 45 — —

Table 5. Ouro-2.6B served (fused path): no quantized arm differs significantly from BF16; LoopKV-M misses the 3-point margin on MATH500 by half a point. MATH500 on one RTX PRO 6000, every arm at our 8,192-token output cap; GSM8K on A40 GPUs, same-GPU pairs. Accuracy (%) and ∆ vs BF16 on the same GPU (pp, from unrounded scores); history KV stored per token. § Second set of A40s. arm KiB / token MATH500 (200) GSM8K (1,319) BF16 KIVI-4 + V LoopKV-M LoopKV-R

1,536 480 (31%) 310 (20%) 430 (28%)

89.5 90.5, +1.0 89.5, +0.0 90.5, +1.0

BF16§ LoopKV-R§

1,536 — 430 (28%) —

79.7 79.5, −0.2 79.5, −0.2 — 79.8 79.2, −0.5

read-optimized plan sustains 0.0244 / 0.0381 req/s, the memory-optimized plan 0.0244 / 0.0476, uniform INT4 0.0204 / 0.0305 and BF16 0.0083 / 0.0133, so the two plans sustain 1.20–1.56× uniform INT4’s request rate and 2.9–3.6× BF16’s. Throughput at a matched time per token. Every decode step gives each of the B resident requests one token, so the time per output token at a peak is B divided by its tok/s. In the fused block of Table 8 it is 64.4, 52.3 and 56.0 ms for KIVI-4 + V at 8K, 16K and 32K, 61.1, 55.4 and 65.2 ms for LoopKV-R and 83.2, 79.2 and 84.8 ms for LoopKV-M; BF16 and FP8 + sharing take 46.8 and 92.1 ms at their 8K peaks. To compare the arms at the SLO’s bound of 3τ0 =55.8 ms, we interpolate the step time, which is linear in B as Eq. 1 predicts, between the two measured batches that bracket the bound (batches 1, 8, 16, 34 and each arm’s peak at 8K; 1, 8, 14 and the peak at 16K; 1, 4, 6, 8 and the peak at 32K), and keep the peak of an arm whose peak already meets the bound. At 55.8 ms, KIVI-4 + V decodes 508, 268 and 142 tok/s at 8K, 16K and 32K, LoopKV-R 543, 289 and 152 (1.069×, 1.078× and 1.069×) and LoopKV-M 539, 286 and 152 (1.061×, 1.067× and 1.067×). Eq. 1, with the constants of §2.2 and the read bytes of Table 8 plus the BF16 sink and window, gives 1.05–1.09× for the same comparisons, and LoopKV-M measured at 16K and batch 16 in an earlier block on the same GPU (56.05 ms, 285.5 tok/s) agrees with the interpolation. At a matched time per token both plans are therefore 6–8% faster than KIVI-4 + V at every context, and LoopKV-M’s larger resident batches lift its peak gains to 8.4–15.6%. Under the same bound BF16 and FP8 + sharing decode 171 and 331 tok/s at 8K and 77 and 167 tok/s at 16K, so LoopKV-M’s multiples are 3.2× and 1.6× at 8K and 3.7× and 1.7× at 16K; at 32K, BF16 takes 81.5 ms per token already at batch 1. At the 62.1 ms bound of the 20.69 ms rule the plans gain 6–9%. SLO ladders. Table 9 and Figure 12 give every ladder of Figure 11, rescored from the stored per-request records with the scorer’s own code, which reproduces every reported attainment exactly. The INT4 ladders step by 1/84≈0.0119 req/s. Resampling the scored requests of every cell (10,000 draws) keeps the LoopKV-M / KIVI-4 + V ratio at 1.5× or more in 86%, 97% and 96% of the draws for seeds 101–103 at 18.60 ms; in a draw in which KIVI-4 + V passes its highest tested rung, its capacity is censored at that rung. At their first failing rungs, KIVI-4 + V misses mostly on time to first token after preemptions (11 of 13 misses at 18.60 ms, 20 of 21 at 20.69 ms), and so does LoopKV-R (22 of 22 at 18.60 ms), whereas LoopKV-M at 18.60 ms misses mostly on time per token (8 of 12), with its 95th-percentile TPOT within 0.5 ms of the 55.8 ms bound; BF16 and the FP8 arms miss on time per token only. With TPOT relaxed to 5τ0 , KIVI-4 + V keeps its capacity on two seeds and LoopKV-M rises, to 1.80×, at least 1.88× and 2.00×; moving the TTFT bound to 5 or 20 s keeps the ratio within 1.50–1.83×. V precision alone adds one rung to uniform INT4 on seeds 101 and 103 at 18.60 ms and on

LoopKV: Loop-Aware KV Cache Compression for Serving Looped Transformers Table 6. Ablation of LoopKV’s components on one RTX PRO 6000: each step improves throughput, KV footprint, or batch-1 latency at a stated KL. Columns report stored / read history KV (KiB/token), fp32 KL / KLshare , peak decode throughput at 8K and 16K (tok/s), relative speedup over the preceding row measured back-to-back in the same job (“same block ×”, at 8K peak / B=1), and host launches per step. The read- and memory-optimized plans are separate knapsack solutions, and the indented rows incrementally add C3, kernel fusion, and whole-step CUDA graphs to the memory-optimized plan. configuration

fp32 KL / KLshare (lower is better)

KIVI-4 read-optimized plan (C1) memory-optimized plan (C1 + C2) + V precision (C3) + fused kernel + whole-step CUDA graph (LoopKV-M)

frontier with V precision

frontier without

bytes

KL

8K

240/240 215/215 155/220 155/189 155/189 155/189

0.46 0.87 1.80 1.80 1.80 1.80

494.4 528.9 535.1 561.6 571.2 575.9

same block ×

host launches

249.1 — 271.5 — 280.7 — 295.5 1.044 / 0.98 300.0 1.019 / 1.033 301.9 1.006 / 1.067

— 433 — 504 332 17

16K

knapsack plans

+V, fused kernel and graph

hand-made + sharing

INT2 + sharing

101 BF16 + sharing

0

U4 + sharing

10

FP8 + sharing

R

BF16 (KL 0)

10

150

200

uniform INT2

INT4/INT2 + sharing M+V R+V U4+V

uniform INT4 −1

M

INT4 + residuals

FP8 250

300

350

400

450

500

550

600

650

700

750

peak decode tok/s at 8K context, vLLM (higher is better)

Figure 9. All served configurations at 8K context, including LoopKV’s own variants (Figure 10 shows one point per comparable method): knapsack plans without V precision, KIVI-4 + V, hand-made placements with sharing, and INT4 designs with residuals. KIVI-4 + V adds our V precision (C3) to KIVI-4 and is an ablation; the sharing points combine loop-3 sharing with BF16, FP8, KIVI-4 or hand-made INT4 / INT2 placements. The frontier with V precision runs through uniform INT4, KIVI-4 + V and LoopKV-R and LoopKV-M; every hand-made and residual design with at most LoopKV-M’s KL lies behind it. Points without +V use our unfused reader, except BF16 and FP8 (vLLM) and uniform INT2 (a separate run).

seed 102 at 20.69 ms, and none in the other three cases. Plain FP8 without sharing, run on seed 101 on the same GPU after the other ladders, sustains 0.0381 req/s under both rules and fails 0.0476 on time per token (0.535 and 0.791), one rung more than FP8 + sharing at 18.60 ms, as the decode sweeps suggest (at 55.8 ms it matches FP8 + sharing at 8K, 329 against 331 tok/s, and leads at 16K, 171 against 167); against plain FP8, LoopKV-M’s multiple is 2.5× at 18.60 ms and 2.8× at 20.69 ms.

C

B YTE MODEL

Ouro-1.4B stores, per token and layer, 2 × 16 × 128 = 4096 K and V elements per loop. KIVI-style codes with group size 32 carry one fp16 scale and one fp16 zero point per 32 elements, 1 bit per element of metadata: INT4 is 5 bits per element, 2.5 KiB per token and layer, and INT2 3 bits, 1.5 KiB. Over 24 layers and 4 loops, INT4 on every loop is 240 KiB per token and BF16 768 KiB. A shared cell stores nothing and reads the source’s 2.5 KiB; a residual cell stores its residual and reads the residual plus the source (INT4 residual: 2.5 / 5.0 KiB; INT2: 1.5 / 4.0). V precision skips the 2-bit lo plane of V, 0.5 KiB per INT4 cell, for the cold fraction of tokens. The sink and window (288 tokens, BF16, 4 loops) add 216 MiB per request and step to Dread in Eq. 1, 10% of the 8K read of the KIVI-4 + loop-3 sharing plan and 5% at 16K.

D

A DDITIVITY OF CELL DAMAGES AND FURTHER PLANNER PLANS

The knapsack scores a plan by the sum of its cells’ single-cell damages. Proposition 2 gives the reason this sum predicts a plan’s KL, and the measurements below test it. Proposition 2 (Approximate additivity) Let xc be the exact model’s history KV of cell c. A plan stores xc + δc , with

fp32 KL / KLshare (lower is better)

LoopKV: Loop-Aware KV Cache Compression for Serving Looped Transformers measured frontier

vLLM BF16 / FP8

KIVI-4 / KIVI-2

KIVI-4 + loop sharing

ResidualQuant-style

LoopKV (ours)

loop sharing (BF16 / FP8)

KIVI-2 101 ResidualQuant-style KIVI-4 + loop sharing

loop sharing (BF16; = 1)

0

10

LoopKV-M (ours)

LoopKV-R (ours)

loop sharing + FP8

KIVI-4 BF16 (KL 0)

−1

10

150

200

250

300

FP8 (vLLM) 350 400

450

500

550

600

650

700

750

peak decode tok/s at 8K context, vLLM (higher is better)

Figure 10. At 8K context, LoopKV-R and LoopKV-M (fused kernel and graph) lie on the measured frontier of served configurations between KIVI-4 (0.46× KLshare ) and KIVI-2 (21.5× KLshare ), the only faster method. Comparable methods, one point each: BF16 and FP8 as vLLM serves them, KIVI-4 and KIVI-2, and INT2 residuals for loops 1–2 against an INT4 loop 4, with loop 3 shared, in the style of ResidualQuant (Kim et al., 2026) (4.38× KLshare ; our reimplementation of ResidualQuant itself has 2.01×, Table 2); the loop-sharing points combine loop-3 sharing with BF16 (KLshare ), FP8 and KIVI-4. The KIVI, sharing and residual points use our unfused reader. Figure 9 adds KIVI-4 + V, also on the frontier, and our other variants. Single runs. LoopKV-M (ours) LoopKV-R (ours)

8.0× 6.0× 5.0×

KIVI-4 + our V KIVI-4 FP8 (vLLM) FP8 + loop-3 sharing BF16 (vLLM)

4.0× 3.2× 2.6× 1.0× 0

2×

4×

6×

8×

16K SLO goodput, multiple of BF16 (seed 101)

Figure 11. LoopKV-M sustains 8.0× BF16’s 16K SLO goodput on one RTX PRO 6000, and 2.0× that of KIVI-4 on the same kernel. Highest request rate with 95% of requests in SLO (τ0 =18.60 ms, §5.1), arrival seed 101, which every arm ran, relative to BF16’s 0.0119 req/s; on seeds 102 and 103, LoopKV-M sustains 1.72× and 2.20× KIVI-4’s rate and LoopKV-R 1.29× and 1.60×.

δc = Q(xc ) − xc for the cell’s own codes and δc = Q(xs ) − xc for a cell that reads its source s, coded as when the cell was profiled; each δc is supported on cell c and does not depend on the other cells’ options. Let the KL D be twice differentiable in the stored KV, D(2) (δ) = 21 δ ⊤ Hδ its second-order part, so that D(δ) = D(2) (δ) + O(∥δ∥3 ), and Dc = D(2) (δc ), which P
P P ⊤ ′ the profile measures up to O(∥δc ∥3 ). Then (a) D(2) c δc = c Dc + c<c′ δc H δc , with H ⪰ 0; (b) if the errors of a set I of cells are random with zero mean and uncorrelated with each other and with every error, every
P Pother cell’s interaction term that involves a cell of I has zero mean, so E D(2) (δ) = c∈I E Dc + E D(2) δ . c c∈I / Proof. The KL is zero and minimal at δ = 0, so D(0) = 0, ∇D(0) = 0, and its Hessian there is H = J ⊤ F J ⪰ 0: the Fisher information F of the output distribution pulled back through the Jacobian J of the logits with respect to the stored KV, averaged over positions and documents. The term with second derivatives of the P logits vanishes, because the gradient of the KL with respect to the logits is zero at δ = 0. Expanding the quadratic form of c δc gives (a). Under (b), a term δc⊤ Hδc′ with c ∈ I has expectation tr(H E[δc′ δc⊤ ]) = 0, since the two errors are uncorrelated and at least one has zero mean. □ Here rounding errors are modeled as random, the usual additive-noise model of quantization, and E is the expectation under this model. Each interaction term sums products of errors over many tokens and channels, so it concentrates near its zero mean, and the sum of single-cell damages becomes a stable predictor of the plan’s KL. That separately coded cells (own codes or residuals) satisfy (b) is an assumption, which the measurements test. A shared cell does not satisfy it: its error (xs − xc ) + (Q(xs ) − xs ) contains the systematic difference between the two loops, which has nonzero mean, and the rounding error of the source, which the source and every cell that reads it share, so shared cells add interactions that the sum omits. Neither do higher-order terms enter the sum. The measurements show that the sum predicts the KL nonetheless. Over the ten plans of Table 12 and the four residual designs of Table 13, the measured KL is linear in the sum on a log scale (slope 1.05, correlation 0.98, rank correlation 0.98) and lies within 0.67–1.20× of it. The deviation is systematic rather than random: the five plans without shared cells measure

LoopKV: Loop-Aware KV Cache Compression for Serving Looped Transformers Table 7. Peak decode tok/s at the resident peak (batch in parentheses) and the ratio to BF16 on the same GPU, fused kernel and whole-step graph, one run per cell. Uniform INT4 is KIVI-4 without V precision; LoopKV-R and -M are the read- and memory-optimized plans with V precision; Ouro-2.6B uses its own allocator plans. model, GPU

context

BF16

Ouro-1.4B, A100-80

8K 16K 32K

188.4 (10) 336.6 (27) 1.79× 377.6 (30) 2.00× 408.5 (42) 2.17× 95.0 (5) 161.3 (11) 1.70× 189.9 (13) 2.00× 209.2 (20) 2.20× 22.7 (2) 84.6 (6) 3.73× 105.4 (8) 4.65× 107.0 (11) 4.72×

Ouro-2.6B, RTX PRO 6000 8K 16K 32K

81.3 (6) 208.3 (16) 2.56× 240.4 (18) 2.96× 263.9 (23) 3.25× 30.8 (3) 109.5 (8) 3.56× 124.1 (8) 4.03× 136.8 (11) 4.44× 6.2 (1) 52.7 (4) 8.51× 58.6 (4) 9.46× 68.1 (5) 11.0×

Ouro-2.6B, A100-80

78.5 (5) 133.2 (13) 1.70× 148.5 (14) 1.89× 174.1 (19) 2.22× 20.2 (2) 68.6 (6) 3.40× 76.6 (7) 3.80× 87.3 (10) 4.33× 5.7 (1) 36.2 (3) 6.35× 39.4 (3) 6.91× 46.4 (5) 8.13×

8K 16K 32K

uniform INT4

LoopKV-R

LoopKV-M

Table 8. With the fused kernel and whole-step graph, LoopKV-R / -M decode 5.4 / 9.3% (8K), 7.8 / 13.2% (16K) and 7.4 / 15.6% (32K) more than KIVI-4 + V given the same kernel and graph (single runs; repeat means at 8K and 16K: §5.3). Peak decode tok/s (batch); KiB per token, stored / read. Top two blocks: before fusion, uniform INT4 with interleaved V; bottom: fused + graph, two-plane V, one block at 8K and 16K (read-optimized + V at batch 36: 552.7) and a second block on the same GPU at 32K. BF16 peaks at batch 8 at 8K: at batch 12 stock FlashAttention-2 falls to one split per head (154.9 tok/s). At batch 1 (8K, bottom block) the plans match KIVI-4 + V: 79.1 and 78.9 vs 80.3 tok/s. × BF16 / plan bytes 8K 16K 32K FP8+S, 8K BF16 FP8 FP8 + sharing (FP8+S) uniform INT4

768/768 170.8 (8) 89.0 (6) 384/384 328.8 (16) 176.3 (13) 288/384 358.4 (33) 180.7 (16) 240/240 494.4 (34) 249.1 (14)

— — — —

1.00 / 0.48 1.93 / 0.92 2.10 / 1.00 2.89 / 1.38

uniform-INT4 + sharing + V precision read-optimized + V precision memory-optimized + V precision

180/240 180/200 215/215 215/186 155/220 155/189

505.4 (40) 521.1 (40) 528.9 (36) 544.5 (36) 535.1 (48) 561.6 (48)

264.2 (20) 277.4 (20) 271.5 (16) 282.9 (16) 280.7 (24) 295.5 (24)

— — — — — —

2.96 / 1.41 3.05 / 1.45 3.10 / 1.48 3.19 / 1.52 3.13 / 1.49 3.29 / 1.57

uniform INT4, fused + graph + V precision read-opt. + V, fused + graph memory-opt. + V, fused + graph

240/240 240/200 215/186 155/189

479.4 (34) 527.7 (34) 556.3 (34) 577.0 (48)

241.7 (14) 127.2 (8) 267.8 (14) 142.9 (8) 288.6 (16) 153.5 (10) 303.2 (24) 165.2 (14)

2.81 / 1.34 3.09 / 1.47 3.26 / 1.55 3.38 / 1.61

0.67–0.79× the sum, so the sum is conservative for them, and the nine with 24–28 shared cells measure 0.73–1.20×. Forty random plans, drawn and with their predictions fixed before any was measured (20 with own codes and residuals only, 20 with shared cells in loops 2–3; predicted 0.73–9.47× KLshare ), show the same on plans the planner did not choose: their measured KL is linear in the sum (log–log slope 1.11, correlation 0.98, rank correlation 0.94) and lies within 0.65–1.37× of it. All twenty without shared cells measure 0.65–0.94× the sum, and the two largest ratios (1.37× and 1.03×) belong to plans with 27 and 26 shared cells. Two plans re-measured in the same run on another GPU reproduce their earlier KL within 0.4%. A two-family correction fitted on the 14 plans of Table 12 and frozen before the random plans were scored, 0.73 S for plans without shared cells and 0.97 S for plans with them (S the sum of single-cell damages), predicts the 40 random plans with a mean absolute error of 15%, against 28% for S itself; its largest miss is a factor of 1.41. Table 12 compares the measured fp32 KL of every plan the planner was asked to build, and of the hand-made plans, with the sum of single-cell damages; every planner plan meets its target, and for the residual designs the sums rank the designs in the same order as their served KL. Plans a, b, c and d use a menu that also contains residual options; c and d are alternative budgets to the memory- and read-optimized plans. Layer-only allocation. To test whether the loop axis matters, we solve the same knapsack with one option per layer, shared by loops 1–3, as layer-wise allocators do, on the same profile and codes (an exact dynamic program over the 24

LoopKV: Loop-Aware KV Cache Compression for Serving Looped Transformers Table 9. SLO ladders on one RTX PRO 6000 (Ouro-1.4B, 16K maximum length, fused path for the INT4 arms; 60 requests per cell, of which 43, 49 and 47 are scored for seeds 101–103). Rungs: lowest and highest rate tested; INT4 rungs step by 0.0119 req/s. Capacity: highest tested rate before the first rung below 95% attainment, with TTFT ≤ 10 s and TPOT ≤ kτ0 ; <: the lowest rung fails; ≥: no tested rung fails. First failing rung (the rung above the 3τ0 capacity at τ0 =18.60 ms): attainment, misses by component (TTFT only / TPOT only / both) and preemptions. Ratio rows: capacity ratios on the same trace, and the share of 10,000 request-bootstrap draws in which the 3τ0 ratio stays at 1.5× or more. capacity (req/s), τ0 =18.60 ms or as noted first failing rung, 3τ0 , 18.60 ms arm

seed

rungs (req/s)

2τ0

BF16 FP8 FP8 + sharing uniform INT4

101 101 101 101 102 103 101 102 103

0.0119–0.0170 0.0305–0.0476 0.0305–0.0476 0.0476–0.0714 0.0476–0.0952 0.0476–0.0833 0.0381–0.0714 0.0595–0.1071 0.0595–0.0833

<0.0119 <0.0305 <0.0305 <0.0476 0.0476 <0.0476 0.0381 <0.0595 <0.0595

KIVI-4 + V

uniform INT4 101 + sharing + V LoopKV-R 101 102 103 LoopKV-M 101 102 103 LoopKV-M / KIVI-4 + V LoopKV-R / KIVI-4 + V

3τ0

0.0119 ≥0.0170 0.0381 ≥0.0476 0.0305 ≥0.0476 0.0476 0.0595 0.0833 0.0833 0.0595 0.0714 0.0595 0.0595 0.0833 0.0952 0.0714 0.0714

0.0714–0.0952 <0.0714 0.0714 0.0714–0.0833 0.0714–0.1190 0.0714–0.1071 0.0595–0.1190 0.0952–0.1786 0.0952–0.1548

<0.0714 <0.0714 <0.0714 <0.0595 <0.0952 <0.0952

5τ0 3τ0 , 20.69 attainment T / P / B preempt. 0.0143 0.0381 0.0381 0.0595 0.0833 0.0714 0.0595 0.0952 0.0714

0.605 0.535 0.930 0.930 0.857 0.936 0.884 0.939 0.894

0 / 17 / 0 0 / 20 / 0 0/3/0 1/2/0 5/0/2 2/1/0 3/0/2 2/1/0 4/1/0

3 6 0 92 93 112 100 44 80

0.0833

0.0714

0.907

1/2/1

189

0.0714 0.0714 0.1071 0.1071 0.0952 0.0952 0.0952 0.1071 0.1429 ≥0.1786 0.1310 0.1429

0.0714 0.1071 0.0952 0.1071 0.1667 0.1429

0.791 0.857 0.872 0.930 0.918 0.894

8/0/1 7/0/0 6/0/0 1/1/1 1/3/0 2/3/0

209 220 209 128 117 182

101

<1.56

1.60

1.80

1.80 ≥1.5× in 86% of draws

102 103

– –

1.72 1.83

≥1.88 2.00

1.75 ≥1.5× in 97% of draws 2.00 ≥1.5× in 96% of draws

101

–

1.20

1.20

1.20

102 103

– –

1.29 1.33

1.13 1.33

1.13 1.33

layers), and measure the plans in the fp32 harness in one run with the two shipped plans (16 documents, paired bootstrap). At the read-optimized plan’s bytes, the layer-only plan (INT4 or INT2 per layer, 213 KiB) has 1.30× KLshare against 0.86× for the read-optimized plan in the same run (0.87× in Table 12), a paired ratio of 1.52×. At the memory-optimized plan’s bytes, it has 2.58× KLshare with loop 3 shared in every layer (154 / 214 KiB) and 3.35× with sharing as a per-layer option (151.5 / 219 KiB), against 1.80×: ratios of 1.43× and 1.86×. Residual designs. Table 13 lists the four closed-loop residual designs of this analysis (INT4 final loop, residuals for loops 1–2, loop 3 shared), the design family of FlashLoop and ResidualQuant, served with our reader. Why the loops compound read errors (P3). Proposition 2(a) also explains P3. Split into k groups of sites, P a perturbation P with damages Di alone and logit errors ei . To second order the joint damage is i Di + i<j ⟨ei , ej ⟩F , so the relative damage of Figure 4(a) is p P P
Di Dj D i̸=j ρij i δi P P = 1+ , i Di i Di where ρij is the cosine of ei and ej in the Fisher metric: 1 for orthogonal errors and k for k identical ones. We record every arm’s log-probability change on 64 documents and form the Fisher Gram of the single-group errors. The Gram predicts each measured ratio without a joint run: for Quest on Ouro-1.4B grouped by loop, 1.81 predicted against 1.82 measured (95% CI 1.62–2.02), with a mean pairwise cosine of 0.31 (CI 0.25–0.37); for the same approximated layers grouped into four layer slices that each span all four loops, 1.21 against 1.19 (CI 1.08–1.27), with a mean cosine of 0.07 (CI 0.03–0.09); and for four layer groups of Qwen2.5-1.5B and Qwen3-1.7B, 0.97 and 1.17 against 0.95 and 1.17, with mean cosines of −0.00 and 0.06. The alignment is already present where the errors are injected. The residual-stream error that dropping pages injects at

attainment (3τ0 , 18.60 ms)

LoopKV: Loop-Aware KV Cache Compression for Serving Looped Transformers seed 101

seed 102

seed 103

1.00 0.95 0.90

1.00 0.95 0.90

1.00 0.95 0.90

0.80

0.80

0.80

0.70

0.70 0.04

0.08

0.12

0.16

0.70 0.04

0.08

0.12

0.16

0.04

arrival rate (req/s)

arrival rate (req/s) uniform INT4

LoopKV-R

KIVI-4 + V

0.08

0.12

0.16

arrival rate (req/s) LoopKV-M

Figure 12. SLO attainment against arrival rate for the INT4 arms on the fused path (one RTX PRO 6000, 16K maximum length; TTFT ≤ 10 s and TPOT ≤ 3τ0 at τ0 =18.60 ms; red line: the 95% rule). Each point is one cell of 60 requests, of which 43–49 are scored; Table 9 gives the scored counts and the other SLO bounds. Table 10. On an H100 SXM (80 GB), LoopKV-M decodes 2.69× as fast as BF16 at 8K and 16K, against 2.27× and 1.96× for KIVI-4, and no INT4 arm differs significantly from BF16 on MATH500 or GSM8K. Own-peak decode throughput relative to BF16 (each arm on its own H100; a rerun of every arm on one GPU agrees within 2%), step time at B=1 and 8K, and served accuracy on MATH500 (500 items) and GSM8K (1,319 items). The H100 kernel reads both V planes. arm BF16 FP8 KIVI-4 KIVI-4 + V LoopKV-R LoopKV-M

8K peak 16K peak B=1, 8K (ms) MATH500 (%) GSM8K (%) 1.00× 2.08× 2.27× 2.26× 2.54× 2.69×

1.00× 2.05× 1.96× 1.96× 2.41× 2.69×

10.68 — 10.63 10.55 10.35 10.33

83.2 — 83.2 85.4 84.6 84.4

76.6 — 77.0 76.5 76.8 77.0

one layer has a cosine of 0.45, 0.71 and 0.82 between loops 1 and 2, 2 and 3, and 3 and 4, but 0.06 between two adjacent layers of one loop (0.04–0.05 on the plain models): the loops apply the same weights to states that converge, so each loop injects nearly the same error, and aligned errors add in amplitude rather than in energy. Coarse reads compound as well: by loop, LoopKV-style 2-bit reads cost 1.51× the sum of their one-loop damages (CI 1.13–2.10; 1.93× without one outlier document) and 1.39× by layer slice, because decoding each code to its bucket midpoint shrinks the extreme channels of a group by the same amount in every loop. What keeps their damage low is the size of each read’s error, not its direction: at equal bytes they inject 4.7× less error energy than dropping on Ouro-1.4B. The advantage of coarse reads over dropping is therefore not specific to looped models, and the plain models show an even larger one; once the outlier document is excluded, it does not grow with the number of loops. What is specific to looped models is the compounding itself.

E

A DDITIONAL RESULTS

Full sensitivity profile. Summed over layers, own INT2 costs 5.6×, 3.0× and 1.9× KLshare in loops 1–3 and own INT4 0.36×, 0.11× and 0.05×, and at own INT4, 54 of the 72 history cells cost less than 0.01× KLshare ; sharing is prohibitive in loop 1 (1074× KLshare at layer 1) and cheap in loop 3 (0.69×, summed over its cells). Proof of Proposition 1. Scores use the full INT4 keys, so the weights equal those of the INT4 reader. A decoded value is z + σc, and |c − (4⌊c/4⌋ + 1.5)| ≤ 1.5 for c ∈ {0, . . . , 15}. Every cold token has sj < m − τ ln 2 ≤ smax − τ ln 2 because m never exceeds the row’s true maximum, hence aj < 2−τ amax . □ V precision threshold. Table 14 sweeps τ on loops 1–2 of the KIVI-4 + loop-3 sharing plan and on all loops. From τ =10 upward on loops 1–2, and at τ =14 on all loops, the KL change is not significant; smaller τ costs 0.1–0.76× KLshare . Overlap of the loops’ hot sets. Table 15 compares the tokens each loop attends to above the threshold, on Ouro-1.4B in fp32 over the last 256 query positions of 16 documents (a different sample from Table 14). Per query row and head, each loop’s set is a quarter of the far history at these positions (Table 14’s sample gives about one in six) and the union over four loops 37%; the hot sets of loop 1 and each later loop share 39–44% of their union, those of loops 3 and 4 share 76%. A

LoopKV: Loop-Aware KV Cache Compression for Serving Looped Transformers Table 11. Once quantized, throughput follows bytes read: at equal storage, own INT4 codes decode 1.36× faster than residuals at 8K. KiB per token of history (4 loops); measured peaks in vLLM on one RTX PRO 6000, except the 16K peak of INT4 residuals, run on the second GPU (§5.1). The cells of the two quantized rows that ran on this GPU, served by our plugin before fusion, lie at 98–103% of Eq. 1; vLLM’s own BF16 and FP8 rows run faster than Eq. 1 predicts (BF16 at 8K: 170.8 against about 146 tok/s). KV format

stored read tok/s 8K / 16K

BF16 BF16 + loop-3 sharing (KLshare ) FP8 FP8 + loop-3 sharing INT4 final + INT4 residuals KIVI-4 + loop-3 sharing

768 576 384 288 180 180

768 768 384 384 360 240

170.8 / 89.0 182.1 / 95.8 328.8 / 176.3 358.4 / 180.7 370.3 / 192.0 505.4 / 264.2

Table 12. Stored / read: KiB per token (4 loops). Predicted (sum of single-cell damages plus source) and measured fp32 KL / KLshare ; served KL relative to the served BF16 sharing reference; peak tok/s in vLLM; no V precision, before fusion. plan

rule

min reads, damage ≤ pred. of INT4 residuals memory-optimized min reads, stored ≤ 156, damage ≤ pred. of layer-split residuals d min reads, damage ≤ 1.42 c min damage, stored ≤ 156 a min damage, read ≤ 240 b min damage, read ≤ 312 KIVI-4 — KIVI-4 + loop-3 sharing — INT4/INT2 + sharing — INT2 + sharing —

read-optimized

stored / read predicted measured served KL tok/s 8K / 16K 215 / 215

1.15

0.87

1.28

528.9 / 271.5

155 / 220

1.97

1.80

1.77

535.1 / 280.7

210 / 210 156 / 321 237 / 240 239 / 312 240 / 240 180 / 240 156 / 216 132 / 192

1.41 1.09 0.67 0.49 0.68 1.32 4.21 9.43

1.11 1.08 0.47 0.37 0.46 1.51 4.21 8.48

1.42 1.41 1.01 0.98 1.02 1.54 3.02 5.11

519.4 / 276.2 401.9 / 209.1 479.8 / 213.0 398.0 / 179.3 494.4 / 249.1 505.4 / 264.2 543.2 / 285.4 604.2 / 319.9

precision fixed per token at write time must serve every loop that reads the page: on a page that loops 3 and 4 share it keeps both loops’ sets at 4 bits, and a single write-time precision for all four loops would keep their union and read 1.6× the V lo-plane bytes of the per-loop rule at τ =14. History reader. LoopKV’s CUDA history reader streams KV codes at 1.59–1.63 TB/s, against 1.57–1.62 TB/s for FlashAttention-2 on BF16, both on an RTX PRO 6000 at 16K (batch 8 and 32). It uses three techniques. First, one block of 512 threads covers all 16 heads of a request row and split, so it loads whole rows with coalesced 16-byte vector loads, and it double-buffers the values in registers. Second, KIVI’s per-channel key scale and zero point are constant over a page, so they are multiplied into the query once per page, and the per-token value scale and zero point fold into the attention weight pj and a running sum. Third, each code decodes in two instructions, one LOP3 bit operation and one fused multiply-add. With V precision on real attention scores, the reader’s 1.12–1.17× speedup over the plain INT4 reader (§4) compares with 1.19–1.23× predicted from its bytes. The fused kernel (Algorithm 1) keeps this thread mapping with 128 registers and no spills and runs INT2 cells within 1–3% of the dedicated reader; in a microbenchmark on the RTX PRO 4500, it is 1.19–1.30× faster than the four launches it replaces at batch 1 and 1.00–1.04× at batch 8–48. Where the decode step goes. Table 16 breaks one decode step of the KIVI-4 + loop-3 sharing plan at 8K and batch 40 into GPU time per component, with and without V precision. The history reader is 75% of the step; V precision makes it 1.06× faster, and the sink seed costs 0.51 ms. BF16 measurement floor. In BF16 activations, coding a single cell as an INT2 or INT8 residual against the INT4 source, or perturbing it by ±1 ulp, gives the same KL, 0.9–1.6×10−4 ; in fp32, INT2 rises 27–561× above the ±1-ulp control while INT8 stays at or below it (Table 17). Over all cells at once, a random ±1-ulp change gives 1.37×10−4 in BF16 activations, 0.56× the 2.44×10−4 that the BF16 sharing configuration has with BF16 activations, but 2.2×10−6 in fp32 activations, 0.02× KLshare (1.20×10−4 ). This is why all damage in this paper is measured with fp32 activations.

LoopKV: Loop-Aware KV Cache Compression for Serving Looped Transformers Table 13. Residual designs served in vLLM before fusion (stored / read: KiB per token; served KL relative to the served BF16 sharing reference; fp32 KL / KLshare ; MQ641: ∆ in pp vs BF16). The 16K peak of INT4 residuals ran on the second GPU (§5.1). design

stored / read tok/s 8K / 16K served KL fp32 KL MQ641

INT4 residuals loop 1 INT4, loop 2 INT2 layer split: layers 1–12 INT4, 13–24 INT2 INT2 residuals

180 / 360 156 / 336 156 / 336 132 / 312

370.3 / 192.0 386.3 / 199.7 385.6 / 200.1 416.8 / 215.9

1.53 1.88 1.79 2.97

1.42 2.12 1.83 4.38

−0.8 −0.8 −2.3 −3.0

Table 14. V precision threshold sweep (base: KIVI-4 + loop-3 sharing, 1.51× KLshare , 1.50× with explicit attention; fp32, 16 documents; ∆ against the same base with explicit attention). loops τ 1–2 1–2 1–2 1–2 1–2 1–2 1–4 1–4

F

KL / KLshare

3 6 8 10 12 16 10 14

2.261 1.768 1.606 1.502 1.492 1.501 1.580 1.488

∆ full-precision V fraction read KiB +0.760 +0.267 +0.105 +0.001 −0.009 −0.000 +0.080 −0.012

.001 / .000 .008 / .002 .020 / .007 .042 / .024 .090 / .067 .321 / .303 .042 / .024 / .024 / .025 .181 / .158 / .167 / .173

216.0 216.1 216.3 216.8 217.9 223.5 193.4 200.1

G ENERALITY: FULL RESULTS

Table 18 lists the Dory and Huginn plans. Additivity holds on random cell sets (measured / summed 0.57–1.09); Dory’s plans measure 0.49–0.58× their predicted KL, and Huginn’s, predicted from interpolated steps, 1.5–2.3×. V precision at τ =10, which keeps 17% (Dory) and 16% (Huginn) of far tokens hot, cuts far reads by 16–17% and changes KL over INT4 by 0.999× on Dory and 0.978× on Huginn; τ is set per model from an fp32 sweep over 10, 14 and 16. Table 19 applies the same profile and knapsack to Ouro-2.6B-Thinking (T =4, 48 layers, 144 history cells; 8 documents, KLshare of this model 5.75×10−5 ). Loop 1 carries the most damage again (own-INT2 loop sums 3.7 : 2.5 : 1; 1.4B: 3.0 : 1.6 : 1), and two random 8-cell sets measure 0.87× and 1.00× their summed single-cell damage. The knapsack’s plans measure 0.85× and 0.98× their prediction; with 8 documents the per-plan intervals are wide and overlap.

G

D ORY-4B SERVED

Dory-4B (NVIDIA, 2026) ships a trained last-loop mode in which all four loops read loop 4’s KV, which cuts its KV per token by 2.67× by design. LoopKV serves on top of this mode, as a plugin to Dory’s vLLM fork: when a page of the shared loop-4 KV of a coded full-attention layer leaves the window, it is coded as KIVI-4 (groups of 32), and all four loops read it as shared cells through the same descriptor model as on Ouro; sliding-window layers are unchanged. The plan step is the same as Ouro’s: we measure in fp32 the single-layer damage of each of the nine full-attention planes at INT4 and INT2, and an exact knapsack chooses BF16, INT4 or INT2 per plane, with ε set to the KL of all nine planes at INT4; the memory-optimized plan minimizes stored bytes under ε, and the read-optimized plan minimizes read bytes under ε/2. INT2 costs 4–20× INT4’s damage on every plane, so neither plan selects it; the memory-optimized plan codes six planes at INT4 and keeps three at BF16 (0.54× BF16’s bytes), and the read-optimized plan codes two (0.85×). Two fixed configurations, which use no profile or knapsack, code the five recurrent full-attention layers (rec5) or all nine full-attention layers (full9). All four pass the in-serving byte check with no differing byte, and with the fused reader on the RTX PRO 4500 their served KL keeps the fp32 order: read-optimized 4.4×10−4 , memory-optimized 4.9×10−4 , rec5 6.2×10−4 , full9 6.6×10−4 (4 documents, 256-token exits). Table 20 gives MATH500 accuracy and the 32 GB GPU’s own peaks under vLLM’s graphs, with the fused kernel templated to Dory’s four KV heads of dimension 256. LoopKV-M’s largest resident batch is 1.33× and 2.0× last-loop BF16’s (64 vs 48 requests at 8K, 48 vs 24 at 16K), and it decodes 3% and 13% faster at 8K and 16K; LoopKV-R, which keeps seven of the nine planes in BF16, decodes as fast as the BF16 cache. On MATH500, no plan loses significantly against last-loop BF16, and with 484 items both planner plans meet a 3-point margin (lower bounds −2.1 and −1.9).

LoopKV: Loop-Aware KV Cache Compression for Serving Looped Transformers Table 15. Loops attend to overlapping but different history tokens, so a per-loop read-time V precision reads fewer lo-plane bytes than a write-time per-token precision that all loops share and that keeps every token any loop marks hot. Ouro-1.4B, fp32, 16 documents, last 256 query positions; hot = attention weight above the τ threshold; 3s = loop 3’s query on loop 4’s keys (a shared cell). hot fraction of the far history, τ =14

Jaccard of hot sets, τ =14

loop 1 / 2 / 3 / 4 .27 / .24 / .25 / .25 union of 4 loops .37 shared cell 3s .25

loop 1 vs 2 / 3 / 4 loops 2–3 / 2–4 / 3–4 3s vs 4

.44 / .40 / .39 .63 / .57 / .76 .83

V lo-plane bytes, write-time per-token precision / per-loop read-time rule: 2.2× (τ =10), 1.6× (τ =14), 1.3× (τ =18)

Algorithm 1 Fused cell kernel for loop t, layer ℓ (one launch). d ← desc[t][ℓ]; row ← blockIdx.x if blockIdx.y < n_near then (o0 , m, l) ← exact BF16 attention over sink + window publish (o0 , m, l); mrow ← m; set ready flag else m0 ← mrow if ready within 4 µs, else −∞ for page p of this block’s split, in d.pool do fold K scales of p into q; read K codes → scores sj update running max m, online softmax hotj ← d.v_planes_read=2 ∧ sj ≥ max(m, m0 )−τ ln 2 read V hi plane; read V lo plane only if hotj accumulate pj σj (4 hi + (hot ? lo : 1.5)) end for publish partial (oy , m, l) end if if last block of the row then o ← LSE-merge of all partials

H

S ERVING ON A 32 GB GPU

Serving Ouro-1.4B on an RTX PRO 4500 (32 GB; 24.5 GiB of KV memory at utilization 0.90), the same final path (fused kernel, whole-step graph, τ =14 on loops 1–4) and protocol, with a 1024-token prefill chunk, give the peaks of Table 21. In-serving byte checks (12 pages per arm) find no differing byte on own pages; on the shared INT4 source pages, 0 bytes differ for LoopKV-M, 1 for LoopKV-R and 3–5 for KIVI-4 and KIVI-4 + V.

I

C ORRECTNESS CHECKS

Served KL agrees with the fake-quantized reference within ±15%; on the RTX PRO 6000, in-serving byte checks of 12 pages per run never differ on a LoopKV cell’s own pages, unfused, fused or graph-replayed; on the shared loop-4 source pages, 0–5 bytes differ per run, on the same pages and in the same counts under every seed mode and plugin copy (one code step at a round-half boundary of the codec), and KIVI-4 differs in at most 2 bytes, own or source; unit tests cover the codec, the reader, mixed prefill and decode batches, and V precision with τ =∞, which is bit-identical to the plain reader. Plain BF16 vLLM already differs between two identical processes from the first generated token, so served comparisons are paired against one BF16 run. All fused-kernel numbers use the production seed of V precision, which waits at most 4 µs for the sink maximum (§3.4); a deterministic seed that keeps the sink maximum makes graph replay match eager execution bit for bit and costs 0.1–0.7% at the served resident peaks, while one that starts at −∞ costs 5–9%.

J

C ODEC BASELINES

Figure 13 places the plans among published KV codecs in the fp32 harness, and Table 22 compares them with uniform KIVI-4 designs on the same codes (32-token groups, as served). The table predicts each design’s 16K resident batch from its stored bytes and its peak decode throughput from Eq. 1 with its reads. These predictions reproduce the measured 16K throughput ratios over KIVI-4 within 4%, before and after fusion (2.2% on the fused path; up to 6.4% at 8K), although at 16K the served designs keep 2–11% more requests than the byte ratio predicts, because near the peak a 10% larger batch raises Eq. 1’s throughput by about 2%.

LoopKV: Loop-Aware KV Cache Compression for Serving Looped Transformers Table 16. GPU time per decode step (ms), KIVI-4 + loop-3 sharing, 8K, batch 40 (torch.profiler), before fusion: seed, history reader and sink/window attention run as separate kernels; the two V-precision columns take m0 from the window or from the 32 sink keys. component

plain + V precision (window seed) + V precision (sink seed)

history reader 53.69 seed — sink + window FlashAttention-2 7.17 GEMMs 8.23 GPU busy 71.38

50.38 2.02 7.21 8.29 70.42

50.71 0.51 7.21 8.28 69.07

Table 17. KL of single-cell perturbations in BF16 and fp32 activations (16 documents). INT2 and INT8 code the cell as a residual against loop 4’s INT4 source, unlike the own codes of Figure 2; ±1 ulp changes the cell’s KV by one unit in the last place. cell loop 1, layer 1 loop 1, layer 12 loop 2, layer 12 loop 3, layer 12 loop 3, layer 24

K

BF16 INT2 BF16 INT8 BF16 ±1 ulp fp32 INT2 fp32 INT8 fp32 ±1 ulp fp32 INT2 / ulp 1.63e-4 1.44e-4 1.18e-4 0.98e-4 0.92e-4

1.37e-4 1.23e-4 1.10e-4 0.98e-4 0.90e-4

1.37e-4 1.25e-4 1.08e-4 0.98e-4 0.90e-4

2.89e-5 1.78e-5 7.53e-6 8.30e-7 9.26e-8

7.5e-8 3.0e-8 4.2e-9 1.0e-9 1.2e-10

8.3e-8 3.2e-8 1.7e-8 7.2e-9 3.4e-9

348 561 433 115 27

W HY INTEGER CODES

LoopKV uses KIVI-style integer codes. Decode attention is memory-bound, so FP4 tensor-core throughput does not shorten the step, and FP4 block formats, although they carry less metadata (NVFP4 0.5 bit per element, MXFP4 (Rouhani et al., 2023) 0.25 bit, KIVI 1 bit), are less accurate on this KV (Table 23): NVFP4 has 2.08× KLshare at 216 KiB and MXFP4 6.40× at 204 KiB, against 0.46× for KIVI-4 at 240 KiB and 0.87× for the read-optimized plan, whose KIVI-style INT4 and INT2 codes read 215 KiB. Rotation helps FP4 most but leaves it behind INT4 (NVFP4 0.97×, MXFP4 2.46×, KIVI-4 0.31×). V precision needs a 2-bit read of a 4-bit code; at four hot cells, reading V’s top two bits does 5–10× the KL damage of the full INT4 cell with integer codes, and with NVFP4 codes 2.7–3.6× the damage of INT4’s top-two-bit read.

L

A LTERNATIVES THAT DID NOT PAY

Previous-loop guidance (for C3). Reading only the previous loop’s top quarter of pages at INT4 and the rest at 2 bits costs +0.69× KLshare over its KIVI-4 + loop-3 sharing base (paired, 16 documents), because the previous loop’s top quarter holds 62–67% of the next loop’s attention per head (49–50% when one page set serves all heads). Aligning unconverged loops (for C2). Keys of loops 2 and 3 differ mainly by a rotation (norm ratio 0.996, cosine 0.976); a MiniCache-style norm correction recovers 2.7% of the sharing KL, and fitted rotations raise it from 4.17× to 4.75× KLshare (BF16 harness). Decode-time last-step reuse. Ouro’s authors let every loop read the last loop’s KV for generated tokens and report 78.85 vs 78.92 on GSM8K and 80.40 vs 82.40 on MATH-500 (Zhu et al., 2025). Under one length-controlled protocol in our engine (rotated INT4 reference cache, greedy decoding, the same 484 MATH500 items at every output cap), this reuse loses 44.4 points at an 8K cap, and at 16K it loses 56.0 points on MATH level 5 and 31.7 on AIME. Letting only loop 3 read loop 4’s KV, outside an exact recent region, loses 1.0 at 8K and stays within noise at 16K (+0.7 and +1.7). Long outputs therefore expose what short benchmarks hide, and the early loops cannot share. Fewer loops or early exit. The exit head after loops 1, 2 and 3 agrees with loop 4 on only 90%, 96% and 99% of tokens, so stopping after fewer loops changes the model’s output.

LoopKV: Loop-Aware KV Cache Compression for Serving Looped Transformers Table 18. The same knapsack has lower mean KL than uniform plans of equal bytes on Dory and Huginn; against INT2 in the first loops, the CIs separate only for Dory at −10% (Huginn: 6 short documents). fp32 KL relative to uniform INT4 on all loops; Dory 16 documents, Huginn 6. At Dory −10% the knapsack has 1.33 [1.25, 1.40] against 2.00 [1.59, 2.35] for INT2 in the first loops (95% CIs); the other rows’ CIs overlap and are not shown. Uniform plans of the same bytes put INT2 in the first or the last loops. model

bytes

knapsack uniform, first uniform, last

Dory-4B Dory-4B Huginn-3.5B Huginn-3.5B

−10% −25% −10% −25%

1.33 2.44 1.04 2.78

2.00 2.92 1.20 3.62

29.3 33.4 85.0 98.3

Table 19. Ouro-2.6B: at equal bytes the memory-optimized plan has 0.48× the KL of the hand-made INT4/INT2 + sharing plan (95% document-bootstrap CIs overlap; not shown), and the read-optimized plan 1.39× against KIVI-4’s 0.82× at 10% fewer reads. fp32 KL / KLshare of Ouro-2.6B, 8 documents; KiB per token (4 loops); V precision at τ =14 on all loops, ∆ paired. plan

stored / read

KIVI-4 KIVI-4 + loop-3 sharing INT4/INT2 + sharing INT2 + sharing read-optimized (knapsack) memory-optimized (knapsack)

KL + V precision (read) ∆ V precision

480 / 480 0.82 360 / 480 1.86 312 / 432 6.30 264 / 384 12.70 430 / 430 1.39 310 / 430 3.04

— — — — 1.41 (374.1) 3.04 (373.9)

— — — — +0.021 +0.001

Table 20. Dory-4B with allocator and fixed plans on one RTX PRO 4500 (32 GB), graph mode, fused GQA reader. MATH500: 484 items, greedy, 8,192-token cap, paired against last-loop BF16. Bytes: KV stored per token relative to last-loop BF16. Peaks: tok/s at the throughput peak (batch at the peak, which can lie below the largest resident batch). arm

bytes MATH500 (∆)

last-loop BF16 LoopKV-R LoopKV-M rec5 (fixed) full9 (fixed)

1.00 0.85 0.54 0.62 0.31

91.5 91.5 (+0.0) 91.7 (+0.2) 91.7 (+0.2) —

8K peak

16K peak

352.0 (40) 352.9 (40) 362.5 (56) 367.7 (56) 370.3 (48)

219.7 (24) 221.6 (24) 247.9 (48) 254.8 (40) 259.6 (56)

Table 21. Ouro-1.4B on a 32 GB RTX PRO 4500, 1024-token prefill chunk: peak decode tok/s (resident batch). The plans’ gain over KIVI-4 + V grows with context because they keep more requests resident. 8K: one repeat; 16K and 32K: mean of three fresh processes at the peaks. context 8K 16K 32K

KIVI-4 KIVI-4 + V LoopKV-R LoopKV-M 190.4 (10) 100.7 (5) 47.2 (2)

202.9 (10) 109.1 (5) 50.6 (2)

218.9 (11) 122.3 (6) 60.2 (3)

242.2 (15) 132.6 (8) 65.5 (4)

gain, own peak gain, same batch (LoopKV-R / -M) (LoopKV-R / -M) 1.078 / 1.194 1.121 / 1.215 1.190 / 1.293

1.042 / 1.036 1.037 / 1.031 1.042 / 1.033

Table 22. At the 32-token groups we serve, the read-optimized plan has lower KL and fewer reads than KIVI-4 + loop-3 sharing, the memory-optimized plan stores less at a higher KL (1.80× against 1.51× KLshare ), and stored bytes set the resident batch. Bytes: KiB per token of history (4 loops), stored / read without V precision; KL: fp32 KL / KLshare of the stored codes, which V precision changes by at most 1.6% (§3.3). Resident batch and peak decode tok/s at 16K, relative to KIVI-4 + V (14 requests, 267.8 tok/s on the fused path), are predicted for every row with V precision: the batch scales with stored bytes from KIVI-4 + V’s 14, and the step time follows Eq. 1 with the measured reads. Parentheses: measured, from the fused-path block of Table 8 and, for the hand-made plan’s batch, from its run before fusion. resident tok/s design bytes KL batch × × 32-token groups with V precision: served KIVI-4 + V 240/240 + loop-3 sharing (hand-made) 180/240 LoopKV-R (knapsack) 215/215 LoopKV-M (knapsack) 155/220

0.46 1 1 1.51 1.33 (1.43) 1.07 0.87 1.12 (1.14) 1.08 (1.08) 1.80 1.55 (1.71) 1.16 (1.13)

LoopKV: Loop-Aware KV Cache Compression for Serving Looped Transformers

fp32 KL / KLshare

2-bit codecs (KIVI-2, TurboQuant-2.5) 101 3-bit codecs (TurboQuant-3.5, KVQuant-3b, rotated 3-bit) ResidualQuant (reimpl.) FP8 + sharing

M rot-4b R+V R M+V

100

KIVI-4 g32

FlashLoop chain FP8

100

150

200

250 300 350 400 450 history KV read per token, KiB (4 loops)

500

550

Figure 13. No codec reading less than 216 KiB matches the read-optimized plan’s KL (0.87× KLshare at 215 KiB without and 186 KiB with V precision). Ouro-1.4B, 16 documents, fp32 KL / KLshare against history bytes read per token; published recipes in our implementation: TurboQuant without QJL, KVQuant without Fisher weighting. Grey: codecs on all four loops; orange: cross-loop residual codecs (FlashLoop’s chain reads all earlier planes); purple and blue: LoopKV plans without and with V precision.

Table 23. 4-bit formats on all four loops (fp32 KL / KLshare , 16 documents). Bits per element include scales and zero points; KL with Hadamard: the same format after a Hadamard rotation. format bits / element KiB per token KL KL with Hadamard KIVI-4, group 32 (asymmetric INT4) INT4 symmetric, group 16 NVFP4 (blocks of 16, E4M3 scale) MXFP4 (blocks of 32, E8M0 scale) FP8 E4M3 (reference)

5.0 5.0 4.5 4.25 8.0

240 240 216 204 384

0.46 1.25 2.08 6.40 0.19

0.31 0.67 0.97 2.46 —

Related documents

• PulseInfer: I/O-Centric Sparse KV Cache Offloading for Efficient Long-Context LLM Decoding#1128820
• Beyond Compression Ratio: High-Throughput KV Compression for Disaggregated LLM Serving#1527901
• UltraQuant: 4-bit KV Caching for Context-Heavy Agents#999816
• DeepSeek-V4.1-Flash: Pushing the Limits of KV Cache Compression#1033751
• SVRF: Efficient Register Storage for Long-Vector Architectures#1216453