PaperlayerView exact revision on arXiv

Source-linked AI summary

GAAC: Lossless KV Compression for Bandwidth-Efficient LLM Inference on Personal Devices

Yipin Guo, Arun M George, Siddharth Joshi

arXiv:2610.38210v1cs.AI

TL;DR

Long-context inference on personal devices is constrained by the growing KV cache and its memory traffic. GAAC losslessly compresses BF16 KV exponents and reconstructs cache tiles on demand during attention. Across evaluated models, it reduces KV storage and improves long-context attention and serving throughput.

  • Problem

    On personal devices, the growing KV cache adds memory and bandwidth demands to long-context inference, while lossy KV compression may introduce additional approximation.

    Source: §1 · p. 1

  • Method

    GAAC uses calibrated per-channel base-plus-offset codes with adaptive group widths to compress BF16 exponents losslessly and reconstruct KV values on demand in attention.

    Source: §1 · p. 2, §3.1 · p. 3

  • Results

    GAAC reduces KV storage by approximately 26%, accelerates long-context attention by up to 1.27×, and increases equal-memory serving throughput by up to 29%.

    Source: §5 · p. 8

  • Takeaways & Limitations

    The evaluation shows that lossless KV compression can reduce storage while translating into lower decoding latency and higher serving throughput on the tested systems.

    Source: §5 · p. 8

    Evaluation covers selected LLMs on DGX Spark and Apple M4 and does not include power or energy measurements.

    Source: Limitations · p. 9

Abstract

from arXiv · show

Weight-only, lossy quantization substantially reduces the storage and memory requirements for serving large language models (LLMs), but the growing key-value (KV) cache remains a bottleneck for long-context inference on personal devices. We present Group-Adaptive Additive Codec (GAAC), a lossless KV compression and serving system that exploits exponent redundancy to reduce storage and accelerate attention without introducing KV approximation. GAAC encodes BF16 exponents using calibrated per-channel base-plus-offset codes over contiguous ranges. GAAC adapts code widths across small groups while preserving KV values bit-for-bit. A fused attention kernel reconstructs KV tiles on the fly, avoiding full-cache decompression. Block-wise compression and fixed-capacity slots support incremental cache updates and paged serving. Across five tested models, GAAC achieves 1.405-1.407$\times$ lossless BF16 KV compression. We also study compression performance on low-precision tensors FP8 E4M3 caches ( 1.09-1.14$\times$) and E5M2 caches (1.25-1.30$\times$). On NVIDIA DGX Spark, GAAC accelerates decode attention at a 64K-token context by 1.19-1.27$\times$ and end-to-end W4A16 decode steps by 1.04-1.15$\times$. In serving, it reduces resident KV memory by approximately 26% and increases aggregate throughput by 13-29% under comparable KV-memory budgets. On Apple M4, GAAC achieves 1.02-1.25$\times$ attention speedups, demonstrating portability beyond NVIDIA GPUs.

1 Introduction

Long-context inference on personal devices faces a growing KV-cache memory and bandwidth burden, even after weight quantization. GAAC compresses KV values losslessly and reconstructs them within attention while supporting incremental cache management.

1 Introduction

  • Long-context KV caches grow with sequence length and are repeatedly read during generation, making storage and memory traffic costly.Weight-only quantization reduces weight storage and traffic but leaves the KV cache unchanged.Source: §1 · p. 1
  • Lossy KV pruning and quantization can reduce costs but may introduce approximation errors requiring validation across tasks and context lengths.The paper motivates reducing KV storage and bandwidth with minimal quality impact instead.Source: §1 · p. 1
  • GAAC encodes BF16 exponents with group-adaptive codes while preserving every sign, exponent, and mantissa bit, then reconstructs KV values on demand within attention.The kernel avoids materializing the full decompressed cache, trading lightweight computation for reduced memory traffic.Source: §1 · p. 2
  • Block-wise compression and slot-based memory management support variable-length compressed blocks and integration with existing LLM serving engines.The design aims to minimize online encoding and decoding overhead.Source: §1 · p. 2

2 Related Work

Prior work compresses model weights, KV states, or communication data using different representations and serving assumptions. GAAC's context is online KV compression with reconstruction inside attention.

2.1 Lossless LLM Weight Compression

  • DFloat11 (Zhang et al., 2025), ZipNN (Hershcovitch et al., 2025), Falcon (Li et al., 2025), and ZipServ (Fan et al., 2026) exploit numerical redundancy for lossless model-weight compression.ZipServ combines Tensor-Core-Aware Triple Bitmap Encoding with a fused decompression-GEMM kernel.Source: §2.1 · pp. 2–3
  • Unlike static weights compressed offline, KV states are generated incrementally and must be compressed online while remaining efficiently reconstructable in attention.The scheme must keep encoding overhead low while reducing cache storage.Source: §2.1 · p. 3

2.2 KV Cache and State Compression

  • KV pruning and quantization reduce memory by evicting tokens or approximating values, potentially adding errors on top of weight quantization.H2O and SnapKV (Li et al., 2024) select tokens; KVQuant (Hooper et al., 2024) uses low-bit storage.Source: §2.2 · p. 3
  • ZipCCL (Lin et al., 2026) losslessly compresses states exchanged during distributed training, targeting inter-device transfers rather than KV-cache storage and memory traffic.Source: §2.2 · p. 3
  • SplitZip (Guo and Joshi, 2026) uses fixed 4-bit frequent-exponent codes and an escape stream, but requires full-cache decompression after transfer rather than piecewise decoding inside attention.Source: §2.2 · p. 3

3 Group-Adaptive Additive Codec

GAAC keeps historical KV states compressed and reconstructs them only when attention needs them. Its codec exploits contiguous BF16 exponent ranges, while fused attention and block-wise management limit data movement and online overhead.

3.1 Lossless KV Compression

  • Approximately 95% of BF16 exponent values can be covered by about six consecutive exponents on average, and contiguous windows approach the coverage of unconstrained frequent-value selection.This motivates replacing arbitrary exponent codebooks with calibrated base-plus-offset codes.Source: §3.1 · p. 3
  • Calibration selects per-head-channel windows of lengths 8 and 16, and groups of eight use 3-bit, 4-bit, or raw exponent modes according to whether offsets fit.Out-of-window values widen the entire group rather than using individually indexed exceptions; sign and mantissa bits remain unchanged, preserving BF16 bit patterns.Source: §3.1 · pp. 3–4

3.2 Attention with On-the-Fly Decompression

  • The fused attention kernel reads compressed KV tiles, reconstructs only the values needed for current attention, and avoids materializing a full decompressed cache.Compressed blocks and any raw tail contribute to the same softmax normalization.Source: §3.2 · p. 4
  • Reduced KV reads can outweigh reconstruction overhead most in long-context decoding, when historical KV memory access can dominate attention latency.Source: §3.2 · p. 4

3.3 Block-Wise Compression Scheduling

  • Fixed-capacity slots store blocks whose encoded size fits; oversized blocks remain BF16, so paged serving manages compressed slots, raw blocks, and an incomplete tail.The block table records storage mode and location, allowing attention to access both representations.Source: Figure 3 · p. 4, §3.3 · p. 5
  • New tokens accumulate in an uncompressed BF16 tail; when a block fills, it is encoded once, and finalized blocks are not recompressed as the sequence grows.During prefill, complete blocks can be encoded while the incomplete suffix stays uncompressed.Source: §3.3 · pp. 4–5

4 Experiments

Experiments use W4A16 models on DGX Spark, with Apple M4 attention tests, and compare GAAC against codec and BF16-KV baselines. Across the tested settings, GAAC compresses KV, accelerates attention and decode, and improves throughput under comparable KV-memory budgets.

4.2 Compression Effectiveness and Generality

  • 1.405–1.407× BF16 KV compression across five models accompanies approximately 26% lower resident KV memory; equal-memory aggregate throughput rises 13–29%.FP8 compression is 1.09–1.14× for E4M3 and 1.25–1.30× for E5M2.Source: §4.2 · pp. 5–6, §4.5 · p. 7

4.3 Block-Level Compression Latency

  • Approximately 1.405× compression with substantially lower encoding and decoding latency than competing codecs is reported for a single KV block and 64K-token cache.GAAC provides higher compression and faster processing than SplitZip and nvCOMP; DietGPU compresses slightly more but has significantly greater latency.Source: Figure 4 · p. 6, §4.3 · p. 6

4.4 Long-Context Attention Acceleration

  • 1.19–1.27× decode-attention speedups over FlashAttention-2 hold across four models at batch size 1 and 64K context.On Qwen3-8B, speedups are 1.31×, 1.27×, and 1.25× at 32K, 64K, and 128K, respectively.Source: Figure 5 · p. 6, §4.4 · p. 6

4.5 End-to-End Serving

  • 1.04–1.15× end-to-end W4A16 decode-step speedups occur across models and context lengths, while resident KV memory falls about 26%.At equal KV-memory budgets, GAAC fits one more sequence per model and raises aggregate throughput by 13–29%.Source: §4.5 · pp. 6–7, Figure 6 · p. 7, Table 2 · p. 7

4.6 Apple Silicon

  • 1.02–1.25× decode-attention speedups over a BF16-KV kernel occur across ten Apple M4 configurations.Speedup grows with context length and batch size; short-context gains are smaller than on DGX Spark.Source: Figure 7 · p. 7, §4.6 · p. 7

4.7.1 Calibration Dataset Generalization

  • Cross-task WikiText-2 calibration reduces the 3-bit-group share by at most 1.7 percentage points and the compression ratio by at most 0.004× versus oracle calibration.Raw groups remain at or below 0.35%; the HumanEval ratio changes from 1.406× to 1.402×.Source: §4.7.1 · pp. 7–8, Table 3 · p. 8

4.7.2 Nested Precision Modes

  • Adding a 2-bit mode improves compression ratio by 2.3% but raises decoding latency to 2.8 times the 3+4+raw design.Removing the 3-bit mode lowers decoding latency but reduces compression ratio by 7.3%.Source: Table 4 · p. 8, §4.7.2 · p. 8

4.7.3 Attention Time Breakdown

  • 294.3 µs less KV-loading time outweighs 49.2 µs of added computation, reducing attention latency from 1.36 ms to 1.11 ms for a 1.22× speedup.The breakdown measures decode attention over a 64K-token context.Source: Table 5 · p. 8, §4.7.3 · p. 8

5 Conclusion

GAAC preserves KV bit patterns while reducing memory use and data movement for long-context inference. Across evaluated models, it reduces KV storage by about 26%, speeds long-context attention by up to 1.27×, and raises equal-memory serving throughput by up to 29%.

5 Conclusion

  • Approximately 26% lower KV storage, up to 1.27× faster long-context attention, and up to 29% higher equal-memory serving throughput are reported across evaluated models.GAAC also achieves 1.02–1.25× decode-attention speedups on Apple M4 and 1.19–1.27× on NVIDIA DGX Spark.Source: §5 · p. 8
  • GAAC exploits contiguous-range BF16 exponent concentration and reconstructs KV values on demand within attention while preserving their original bit patterns.Fill-then-finalize compression and paged management of compressed and raw blocks connect the representation to lower decoding latency and higher throughput.Source: §5 · p. 8
  • Decode-attention speedups are reported on both Apple M4 and NVIDIA DGX Spark, demonstrating portability across architectures.The reported ranges are 1.02–1.25× on Apple M4 and 1.19–1.27× on NVIDIA DGX Spark.Source: §5 · p. 8

Limitations

GAAC's compression depends on exponent redundancy, and its efficiency varies with cache format and calibration context. Benefits also depend on workload and evaluation scope: prefill receives no comparable direct acceleration, and energy was not measured.

Limitations

  • GAAC's effectiveness depends on exponent-field redundancy and it provides less additional lossless compression for FP8, particularly E4M3, than for BF16.Calibrated exponent windows transfer well across evaluated tasks, but RoPE-rotated keys can exhibit positional drift when calibration contexts are shorter than deployment contexts.Source: Limitations · p. 9
  • Under shorter calibration contexts, RoPE positional drift remains lossless but can require wider encoding modes and raw fallbacks, reducing compression efficiency.Source: Limitations · p. 9
  • Benefits are smaller for shorter contexts and models with fewer full-attention layers; prefill uses standard attention and receives no comparable direct acceleration.The evaluation covers selected LLMs on DGX Spark and Apple M4 and does not include power or energy measurements.Source: Limitations · p. 9

Ethical Considerations

GAAC changes inference efficiency without changing model weights, training data, or reconstructed BF16 KV values. It does not address underlying model and application risks, and lower per-request resource use is not evidence of lower aggregate environmental impact.

Ethical Considerations

  • GAAC leaves model weights, training data, and reconstructed BF16 KV values unchanged.Source: Ethical Considerations · p. 9
  • GAAC does not address safety, privacy, bias, or reliability risks of the underlying language model and its application.Source: Ethical Considerations · p. 9
  • Lower per-request resource use does not demonstrate lower aggregate environmental impact, since deployment volume may also increase.Source: Ethical Considerations · p. 9

A Model Quality under Activation Quantization

GAAC preserves the benchmark results of the W4A16 and W3A16 baselines while losslessly reducing KV storage. The activation-quantization comparisons report additional quality degradation, especially on reasoning and coding tasks.

A Model Quality under Activation Quantization

  • The evaluation uses Qwen3.8-27B on mathematical-reasoning and code-generation benchmarks across W4A16, W4AFP8, W4A4, W3A16, and W3AFP8 variants.All variants use the same prompt template, decoding configuration, and evaluation harness.Source: A Model Quality under Activation… · p. 11
  • W4AFP8 scores 76.7% on AIME 2024 and 95.1% on HumanEval, versus W4A16's 80.0% and 97.0%.W4A4 reduces HumanEval accuracy to 93.9% and LiveCodeBench pass@1 from 37.0% to 34.9%.Source: A Model Quality under Activation… · p. 11
  • W3AFP8 reduces AIME 2024 accuracy to 50.8%, compared with 56.7% for W3A16.Source: A Model Quality under Activation… · p. 11
  • W4-GAAC and W3-GAAC preserve their respective A16-baseline benchmark results while losslessly reducing KV storage to approximately 11.8 bits per element.The results support combining aggressive weight quantization with lossless KV compression without additional numerical approximation.Source: A Model Quality under Activation… · p. 11

B Relation to FP8 KV

GAAC further compresses FP8 KV caches without adding numerical error beyond the initial FP8 conversion. For positional drift, calibration on deployment-length prompts keeps key overflow stable across the prompt.

B Relation to FP8 KV

  • FP8 KV quantization reduces per-element storage, and GAAC further compresses the quantized cache losslessly.The evaluation uses model-generated KV states converted to FP8.Source: B Relation to FP8 KV · p. 11
  • GAAC adds 1.09–1.14× compression on E4M3 and 1.25–1.30× on E5M2, without approximation beyond the initial FP8 conversion.These correspond to total compression ratios of 2.18–2.28× and 2.50–2.60× relative to BF16, respectively.Source: B Relation to FP8 KV · p. 11, Table 7 · p. 11

C RoPE Positional Drift

  • With windows calibrated on prompts of at most 8K tokens, K overflow rises from 37–38 bytes per tile up to 8K to 190 bytes at 28–32K, while value overflow stays at 38–41 bytes.RoPE’s position-dependent rotations shift key exponents outside short-prompt-calibrated windows; calibrating on 32K prompts keeps K overflow at 37–40 bytes in every segment.Source: C RoPE Positional Drift · pp. 11–12, Table 8 · p. 11

D Storage

A compressed 16-token tile averages 3,016 bytes, or 11.783 bits per element, including coding and serving overhead.

D Storage

  • A compressed 16-token tile averages 3,016 bytes, corresponding to 11.783 bits per element, including its storage overhead.The 256-dimensional head is represented as two 128-dimensional half-heads; storage includes 3-bit exponent codes, 2-bit mode tags, escape bytes, fixed-capacity escape slots, block-table entries, and occasional raw fallbacks.Source: D Storage · p. 12, Table 9 · p. 12

E Attention Baseline Selection

On DGX Spark at a 64K-token context, GAAC is at least 1.20× faster than FA2 and FlashInfer for decode attention. FA2 is used as the FlashAttention baseline because FA3 and FA4 are unsupported on this device.

E Attention Baseline Selection

  • FA3 and FA4 are unsupported on DGX Spark, so FA2 serves as the FlashAttention baseline on this device.The GB10 GPU lacks the Hopper warpgroup MMA and TMA paths used by FA3, and the tcgen05 tensor-core instructions and tensor memory required by FA4 on data-center Blackwell.Source: E Attention Baseline Selection · p. 12
  • At a 64K-token context, GAAC is at least 1.20× faster than FA2 or FlashInfer for decode attention.The comparison measures one Qwen3.8-27B layer; all methods run in the same process, with median per-call latency summarized across five runs.Source: Table 10 · p. 12, E Attention Baseline Selection · p. 12

Related papers

Loading 2610.38210v1…