WorksComputer Science
GAAC: Lossless KV Compression for Bandwidth-Efficient LLM Inference on Personal Devices
Abstract
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 lossless BF16 KV compression. We also study compression performance on low-precision tensors FP8 E4M3 caches ( 1.09-1.14) and E5M2 caches (1.25-1.30). On NVIDIA DGX Spark, GAAC accelerates decode attention at a 64K-token context by 1.19-1.27 and end-to-end W4A16 decode steps by 1.04-1.15. 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 attention speedups, demonstrating portability beyond NVIDIA GPUs.
Introduction
Running large language models (LLMs) on personal devices requires fitting both model weights and the growing key-value (KV) cache within a limited memory budget. On personal platforms such as NVIDIA DGX Spark and Apple Silicon, decoding is typically constrained by memory bandwidth rather than available arithmetic throughput [3]. Weight-only quantization [7, 10, 28, 8] has been the common, practical solution for this setting. By reducing weight precision while retaining higher-precision activations, model storage and memory traffic from weight reads can be substantially reduced, with minimal quality.
However, shrinking the weights leaves the KV cache unchanged, increasing its relative contribution to memory consumption and decoding traffic. KV storage and memory traffic become particularly costly in long-context applications: the cache grows with sequence length, and full attention repeatedly reads the cached keys and values during generation [13]. Consequently, fitting a longer context into memory does not necessarily make processing that context fast.
KV pruning and quantization [30, 20] can reduce these costs, but introduce additional approximation and possibly performance loss [29]. For deployments already using aggressively quantized weights, this creates another source of error that requires validation across tasks and context lengths. Instead, alternative methods to reduce KV storage and bandwidth requirements with minimal impact on quality would be critical to efficient serving of LLMs.
Lossless compression offers one such method, but storage savings alone do not imply faster attention. Prior lossless compression methods exploit floating-point redundancy to reduce model weight storage [29] or accelerate KV transfer between server GPUs [9]. However, accelerating long-context attention has different needs, compressed KV must remain resident in device memory and be reconstructed on demand, avoiding full-cache decompression and its associated memory traffic. On bandwidth-constrained devices, the codec and attention kernel must therefore be co-designed to be memory-layout aware.
We introduce Group-Adaptive Additive Codec (GAAC), a lossless KV compression system for bandwidth-efficient inference on personal devices. GAAC encodes BF16 exponents using short codes adapted to each channel, increasing the code width for groups that require more bits while preserving every sign, exponent, and mantissa bit. The attention kernel reconstructs KV values on demand from compressed storage, trading lightweight computation for reduced memory traffic without materializing the full decompressed cache. For serving, we compress the KV cache in blocks aligned with Tensor Core processing granularity. This supports efficient compression and memory management as new KV entries are added during decoding. Our contributions are:
A bandwidth-oriented lossless KV codec. We design a group-adaptive lossless BF16 KV compression method for devices with a high compute-to-memory-bandwidth ratio, balancing compact storage with efficient, bit-exact reconstruction.
Attention acceleration through reduced data movement. We fuse KV reconstruction with attention, reducing memory traffic and reducing attention latency, especially during long-context decoding.
Efficient compressed KV cache management. We design a block-wise compression strategy to minimize online encoding and decoding overhead. A slot-based memory management scheme accommodates variable-length compressed blocks within two fixed-size storage classes, enabling integration into existing LLM serving engines.
Related Work
Lossless LLM Weight Compression
General-purpose GPU libraries such as nvCOMP ([24]) provide lossless compression primitives. LLM-oriented methods further exploit the structure of floating-point model parameters. DFloat11 ([29]) uses variable-length exponent coding and GPU decompression to reduce BF16 weight storage without changing the original values. ZipNN ([12]) develops neural-network-specific lossless compression for reducing model storage and distribution costs. Falcon ([17]) accelerates floating-point compression through GPU-oriented, precision-preserving conversion and adaptive bit-plane encoding. ZipServ ([6]) targets inference by combining Tensor-Core-Aware Triple Bitmap Encoding with a fused decompression-GEMM kernel that processes compressed weights without materializing an intermediate decompressed tensor.
These studies highlight the benefits of exploiting numerical redundancy and co-designing compression together with the kernels that use the compressed data. Unlike static model weights, which can be compressed offline, KV states are generated incrementally during inference and must be compressed online as the cache grows. The compression scheme must therefore keep encoding overhead low, reduce cache storage and support efficient reconstruction within attention.
KV Cache and State Compression
Lossy KV compression primarily reduces either the number of cached tokens or the numerical precision of the stored KV. H2O [30] retains recent tokens and attention-heavy hitters while evicting less important KV entries; SnapKV [16] uses attention patterns within an observation window to identify important positions for each attention head. For quantization, KVQuant [13] combines per-channel, pre-RoPE key quantization with nonuniform representations and outlier handling to support low-bit KV storage. These methods reduce KV memory by discarding tokens or approximating stored values, potentially introducing additional errors on top of weight quantization.
Lossless compression also reduces the volume of data transferred when communicating other model states. ZipCCL [19] compresses activations, gradients, and parameters exchanged by communication collectives during distributed training, targeting transfers between devices rather than the storage and memory traffic of the KV cache.
SplitZip [9] is a high-throughput codec targeting lossless KV compression. It encodes frequent exponents with fixed 4-bit codes and handles rare values through an explicit escape stream, prioritizing encoding and decoding throughput for prefill–decode disaggregation. However, it is designed to decompress the full cache once after transfer, and its format cannot be decoded piece by piece inside attention.
Group-Adaptive Additive Codec
GAAC couples a bandwidth-oriented lossless codec with compressed-cache attention and block-wise cache management. It keeps historical KV states compressed in memory and reconstructs
them only when needed by attention.
Lossless KV Compression
Contiguous exponent concentration. A BF16 element contains one sign bit, an 8-bit exponent, and a 7-bit mantissa. Consistent with prior observations on lossless compression [9, 29, 6], BF16 exponents in LLMs typically have substantially lower entropy than their allocated bit width, whereas the mantissa fields contain little exploitable redundancy. We therefore store each element’s sign and mantissa unchanged in a single byte and compress only its exponent.
Our KV analysis reveals that frequently occurring exponents are highly concentrated within short contiguous intervals, as illustrated in Figure 2. Across different LLM architectures, approximately 95% of exponent values can be covered by only 6 consecutive exponents on average. Moreover, contiguous intervals achieve nearly the same coverage as selecting the most frequent exponents without contiguity constraints, motivating our base-plus-offset encoding scheme.

Figure 2. Concentration and contiguity of BF16 KV exponents. Most channels require fewer than eight consecutive exponents to cover 95% of their values.
This structure allows us to replace an arbitrary exponent codebook with a base and a small integer offset. For each head–channel pair, calibration selects the contiguous windows of lengths 8 and 16 that cover the most exponent values, separately for K and V in each layer. We denote their bases for element by and . These bases are shared across tokens in each KV cache block rather than fitted during every compression operation.
Group-adaptive additive coding. As shown in Figure 1(a), we partition KV elements into groups of eight in attention-access order. Each group shares a 2-bit mode tag selecting 3-bit offsets, 4-bit offsets, or raw 8-bit exponents. Writing , the exponent width of group is

Figure 1. Overview of GAAC. (a) GAAC exploits the concentration of BF16 exponent values in contiguous high-frequency regions. Each group of eight exponents shares a 2-bit tag indicating 3-bit, 4-bit, or raw mode. The sign and mantissa bits of each element are stored unchanged in a single byte. Each code is an offset from a calibrated per-channel base, . In the illustrated example, one exponent outside the 3-bit window causes the entire group to use 4-bit mode. Bits per element include the sign-mantissa byte and the mode tag. (b) Compared with standard attention, GAAC reduces KV loading time by reading fewer bytes from the compressed cache, with only a small computational overhead from on-the-fly decompression. denotes the attention output. Example values and timelines are illustrative.
GAAC handles out-of-window exponents by increasing the code width for the entire group, avoiding individually indexed exceptions. With , encoding and reconstruction are
Decoding therefore employs integer addition rather than exponent-codebook lookups, making it better suited for personal devices with a high computation-to-bandwidth ratio.
All modes retain the low three bits of each code in a common packed stream. A 4-bit group stores its eight additional high bits in one extra byte; a raw group stores the remaining five bits per exponent in five extra bytes. The mode tag determines how many additional bytes to read and whether to add a calibrated base. Grouping limits metadata to bits per element. Therefore the common 3-bit mode occupies bits per element. Because every exponent is reconstructed exactly and all sign and mantissa bits are retained, the original BF16 bit patterns are preserved. Calibration mismatch affects compression efficiency, not correctness: values outside both windows use the raw mode.
Attention with On-the-Fly Decompression
Compressing a cache but fully expanding it before attention would reintroduce full-size memory writes and reads. GAAC instead fuses reconstruction with decode attention, as shown in Figure 1(b). The attention kernel reads compressed KV tiles and their metadata, reconstructs the required BF16 values, and immediately uses them to compute attention scores and weighted sums of values. Only the current working tiles are reconstructed; no full-size decompressed cache is materialized in device memory. Compressed blocks and any raw tail contribute to the same softmax normalization, preserving the original attention formulation.
This design trades lightweight reconstruction for reduced KV memory traffic. The benefit is most pronounced in long-context decoding, where reading historical KV states can dominate attention latency and the saved memory-access time can outweigh reconstruction overhead.
Block-Wise Compression Scheduling

Figure 3. System design of GAAC. (a) New tokens are appended to an uncompressed BF16 tail until the block is full. (b) The full block is encoded once. If its compressed size does not exceed the slot capacity , it is stored in a compressed slot; otherwise, it remains in BF16 format. (c) The paged KV cache holds compressed and raw finalized blocks in block-table order, followed by the incomplete BF16 tail.
Inspired by paged KV management ([15]), we align compression units with token blocks, with K and V stored independently for each head. Each block follows a fill-then-finalize lifecycle. New tokens are appended to an uncompressed BF16 tail while the block is incomplete. When its last token arrives, the completed block is encoded once, and its compressed segments become the persistent cache representation. After successful compression, the memory holding the uncompressed block is released for reuse, and finalized blocks are not recompressed as the sequence grows. During prefill, already complete blocks can be encoded directly, leaving only the incomplete suffix uncompressed.
The two levels of fallback in our codec produce variable-length compressed blocks, complicating memory allocation and management. We use fixedcapacity slots for compressed blocks and calibrate their capacity to maximize the compression ratio, balancing unused slot space against the frequency of raw-block fallbacks. If the encoded size of a block exceeds its slot capacity, the block remains in its original BF16 representation. Consequently, the serving system only manages two physical storage classes: compressed slots and raw BF16 blocks. The block table records the storage mode and physical location of each block, allowing attention to access both representations transparently. This design preserves losslessness while supporting efficient memory management without arbitrary variable-size allocations. The storage analysis in Appendix § 12 shows that fixed-capacity slots leave very little unused space.
Deferring compression until block completion amortizes encoding over multiple token appends and avoids repeatedly repacking partially filled blocks. Historical blocks remain compressed across subsequent decoding steps, reducing resident KV storage and repeated attention traffic. The fused path targets decoding; prefill retains standard attention, temporarily expanding cached context when required.
Experiments
Setup
Hardware. We primarily evaluate GAAC on an NVIDIA DGX Spark equipped with a GB10 GPU, 48 SMs, and 128 GB of unified LPDDR5x memory. The measured memory bandwidth is approximately 240 GB/s. We additionally evaluate the attention implementation on an Apple M4 with a 10-core GPU and 16 GB of unified memory to study the portability of the proposed design across unified memory architectures.
Models and workloads. We evaluate dense models with 4B to 32B parameters from the Qwen3 and Llama-3.1 families ([26, 21] (Qwen Team 2024)). We also evaluate newer Qwen models with hybrid attention or mixture-of-experts (MoE) layers: Qwen3.8-27B and Qwen3.6-35B-A3B. Experiments use AWQ-quantized checkpoints ([18]) with 4-bit weights and 16-bit activations (W4A16). We calibrate exponent windows on a subset of the WikiText ([23]) training data. We test calibration transfer to code, math, knowledge, and language-modeling tasks with HumanEval ([1]), GSM8K ([2]), MMLU ([11]), and PTB ([22]).
Baselines and metrics. We compare GAAC with five GPU lossless codecs at the block level: nvCOMP ([24]), DietGPU ([14]), DFloatt11 ([29]), ZipServ ([6]) and SplitZip ([9]). For decode attention, we compare GAAC with FlashAttention ([5, 4]) using a BF16 KV cache. For serving, we use the same vLLM ((vLLM Team, 2026)) 0.17.1 engine with a BF16 paged KV cache as the baseline. Decode-step latency comparisons use the same model, batch size and context length. The compression ratio is the BF16 size divided by the complete compressed size, including mode tags and raw fallbacks. Speedup is the baseline latency divided by the GAAC latency. End-to-end serving accounts for online compression overhead and all other system overheads.
Compression Effectiveness and Generality
The exponent structure exploited by GAAC remains consistent across model families, sizes, and attention architectures (Table 1). The empirical entropy of the BF16 exponent stream remains low across all five models, ranging from 2.69 to 2.78 bits. As a result, 88.3–89.9% of the 8-element groups use the 3-bit mode, 9.9–11.5% use the 4-bit mode, and at most 0.3% fall back to raw exponents. GAAC reaches a compression ratio of 1.405–1.407× on every model. Accounting for the storage overhead of fixed-capacity slots, the effective reduction in resident KV memory is approximately 26%, as detailed in Appendix § 12. This approaches the 1.422× compression ratio of a cache encoded entirely in the 3-bit mode (16/11.25), showing that wider modes add little storage overhead. GAAC maintains consistent compression ratios across dense, hybrid, and MoE models with different numbers of KV heads.
| Model | KV exp. entropy (bit) | Ratio | 3-bit | 4-bit | Raw | |
| Qwen3-4B | 8 | 2.78 | 1.405× | 89.0% | 10.8% | 0.2% |
| Qwen3-8B | 8 | 2.78 | 1.405× | 88.9% | 10.9% | 0.2% |
| Qwen3.8-27B | 4 | 2.69 | 1.405× | 88.3% | 11.5% | 0.2% |
| Qwen3.6-35B-A3B | 2 | 2.71 | 1.405× | 88.4% | 11.4% | 0.2% |
| Llama-3.1-8B | 8 | 2.70 | 1.407× | 89.9% | 9.9% | 0.2% |
Table 1. Lossless KV compression across models. Entropy denotes the empirical entropy of the BF16 exponent stream. Ratio is the BF16 size divided by encoded size. Mode percentages are measured over 8-element groups.
We also evaluate how much additional lossless compression GAAC provides on FP8 KV caches. GAAC achieves 1.09–1.14× compression on E4M3 and 1.25–1.30× on E5M2. Detailed results are provided in Appendix § 10.
Block-Level Compression Latency
We compare compression latency for a single KV block and decompression latency for a 64K-token cache, using KV states from one layer of Qwen3.8-27B. We evaluate DFloat11 and ZipServ only for decompression, as both are designed for fast decoding.
As shown in Figure 4, GAAC achieves a compression ratio of approximately 1.405× with substantially lower encoding and decoding latency than competing codecs. Compared with SplitZip and nvCOMP, it provides both higher compression and faster processing. Although DietGPU achieves a slightly higher compression ratio, it incurs significantly greater latency. These results show that additive coding combines compact storage with low decoding latency, enabling efficient on-the-fly KV reconstruction within attention.

Figure 4. Compression ratio and codec latency on KV states from one layer of Qwen3.8-27B. Compression latency is measured for a single KV block; decompression latency is measured for a 64K-token cache.
Long-Context Attention Acceleration
We next measure decode attention separately to determine whether the time saved by reducing KV traffic outweighs the overhead of on-the-fly reconstruction. Figure 5 compares GAAC on compressed KV with FlashAttention-2 on BF16 KV at a batch size of 1 and a 64K context. GAAC is 1.19–1.27× faster on all four models. The dense Qwen3-8B and Llama-3.1-8B gain the most, at 1.27× and 1.25×. GAAC achieves 1.19–1.22× speedups on the hybrid and MoE models despite their smaller per-layer KV data volume.

Figure 5. Decode-attention speedup of GAAC using compressed KV over FlashAttention-2 using BF16 KV on DGX Spark, with a batch size of 1 and a 64K-token context.
The speedup also holds across context lengths and batch sizes. On Qwen3-8B, GAAC is 1.31×, 1.27×, and 1.25× faster at 32K, 64K, and 128K, and 1.29× faster with two 64K sequences. GAAC translates the 1.41× reduction in KV data volume into substantial attention speedups, with memory-access savings outweighing the overhead of on-the-fly reconstruction.
End-to-End Serving
Equal-batch decode. We measure complete W4A16 decode-step latency at a batch size of 1, including model computation and KV-cache access (Figure 6). GAAC speeds up every model at every context length by 1.04–1.15×. The speedup increases with context length because KV reads account for a larger share of decode-step latency. Qwen3-8B improves from 1.10× at 32K to 1.15× at 128K, and Qwen3-32B from 1.05× to 1.12×. Llama-3.1-8B reaches 1.15× at 128K. The hybrid Qwen3.8-27B and MoE Qwen3.6-35B-A3B achieve 1.04–1.06× speedups, as their KV caches are small relative to their model weights and only

Figure 6. End-to-end W4A16 decode-step speedup of GAAC over the BF16-KV baseline on DGX Spark at a batch size of 1.
25% of their layers use full attention. In all cases GAAC also cuts resident KV memory by about 26%, for example from 19.3 GB to 14.2 GB for Qwen3-8B at 128K.
Equal-memory throughput. By reducing KV storage, GAAC also allows more sequences to be served concurrently within the same memory budget. We fix the KV-memory budget and increase the number of concurrent long-context sequences to the maximum each system can accommodate (Table 2). Under the same budget, GAAC fits one more sequence for every model, and aggregate throughput rises by 13–29%. The largest gains appear on the hybrid Qwen3.8-27B (+27%) and the MoE Qwen3.6-35B-A3B (+29%), even though their equal-batch speedups are the smallest. For these models, the main serving benefit comes from using the saved KV memory to support more concurrent sequences.
| Model | Context | Budget | BF16 KV Batch | BF16 KV tok/s | GAAC Batch | GAAC tok/s | Gain |
| Qwen3-8B | 16K | 10 GB | 4 | 60.3 | 5 | 68.1 | +13% |
| Qwen3-8B | 64K | 20 GB | 2 | 19.0 | 3 | 22.8 | +20% |
| Qwen3-32B | 64K | 35 GB | 2 | 8.47 | 3 | 10.51 | +24% |
| Llama-3.1-8B | 64K | 18 GB | 2 | 20.0 | 3 | 24.9 | +25% |
| Qwen3.8-27B | 64K | 14 GB | 3 | 23.3 | 4 | 29.6 | +27% |
| Qwen3.6-35B-A3B | 64K | 5 GB | 3 | 62.9 | 4 | 81.3 | +29% |
Table 2. Serving throughput under approximately equal KV-memory budgets.
Apple Silicon
To test portability beyond NVIDIA GPUs, we implement the same compressed-KV attention path in Metal. On an Apple M4, we compare it with a BF16-KV kernel of the same structure, using KV
states generated by Qwen3-8B. Figure 7 shows that the GAAC attention kernel is 1.02–1.25× faster than the BF16-KV baseline across all ten configurations. The speedup grows with both context length and batch size. With 8 sequences it rises from 1.09× at 4K to 1.25× at 64K, and with 4 sequences from 1.09× to 1.19×. For a single sequence, GAAC achieves speedups of 1.02× at 8K and 1.10× at 64K. At short contexts, speedups are smaller than on DGX Spark because reconstruction accounts for a larger share of execution time on the M4 GPU. Longer contexts and larger batches require each kernel to read more KV data, increasing the benefit of reduced memory traffic.

Figure 7. Decode-attention speedup of GAAC’s Metal implementation over a BF16-KV kernel on Apple M4.
Ablation Study
Calibration Dataset Generalization
We test whether exponent windows calibrated on one dataset generalize to other tasks. Let denote the WikiText-2 training data used for calibration. For each target task , we compare cross-task calibration with oracle calibration , reporting the proportions of groups in 3-bit, 4-bit, and raw modes and the resulting compression ratio.
Table 3 shows that windows calibrated once on WikiText-2 transfer well to code, math, knowledge, and language-modeling tasks. Cross-task calibration reduces the proportion of 3-bit groups by at most 1.7 percentage points, while raw groups remain at or below 0.35%. The compression ratio drops by at most 0.004 relative to oracle calibration, from 1.406 to 1.402 on HumanEval. A single offline calibration pass is sufficient for most workloads.
| Evaluation task | 3-bit groups (%) | 3-bit groups (%) | 4-bit groups (%) | 4-bit groups (%) | Raw groups (%) | Raw groups (%) | Ratio () | Ratio () |
| WikiText-2 / language modeling | 88.3 | 88.3 | 11.5 | 11.5 | 0.20 | 0.20 | 1.405 | 1.405 |
| HumanEval / code | 87.1 | 88.8 | 12.6 | 11.0 | 0.35 | 0.16 | 1.402 | 1.406 |
| GSM8K / math | 87.5 | 89.0 | 12.3 | 10.9 | 0.24 | 0.14 | 1.404 | 1.406 |
| MMLU / knowledge | 87.5 | 87.8 | 12.2 | 11.9 | 0.28 | 0.28 | 1.403 | 1.404 |
| PTB / language modeling | 88.3 | 88.9 | 11.5 | 11.0 | 0.23 | 0.14 | 1.405 | 1.406 |
Table 3. Cross-task calibration generalization. Paired columns compare calibration on WikiText-2 () with oracle calibration on the target task (). The lengths of the calibration prompts are consistent.
Nested Precision Modes
We compare different precision-mode configurations to understand the tradeoff between compression efficiency and decoding overhead. We measure compression latency for a 16-token block and decoding latency for a 64K-token cache, excluding transmission latency.
GAAC uses a 2-bit mode tag to select among 3-bit, 4-bit, and raw encoding. The unused tag value could support a fourth mode with 2-bit codes. However, adding this mode improves the compression ratio by only 2.3% while increasing decoding latency to 2.8 times that of the 3+4+raw design (Table 4). Removing the 3-bit mode reduces decoding latency but lowers the compression ratio by 7.3%. The 3+4+raw design retains most of the compression benefit of the four-mode scheme at substantially lower decoding latency.
| Scheme | Ratio | Compress (s) | Decode (s) |
| 2+3+4+raw | 1.439 | 42.6 | 526.7 |
| 3+4+raw (ours) | 1.406 | 42.1 | 187.1 |
| 4+raw | 1.304 | 36.7 | 158.9 |
Table 4. Precision-mode ablation on Qwen3.8-27B.
Attention Time Breakdown
To identify the source of GAAC’s attention speedup, we separately measure the time spent on KV loading and computation during decode attention over a 64K-token context.
Table 5 shows that GAAC reduces KV loading time by 294.3 s while adding 49.2 s of computation. Total attention latency therefore decreases from 1.36 ms to 1.11 ms, yielding a 1.22 speedup.
| Component | FA2 | Decomp.-first | GAAC |
| KV loading | 1117.8 | 1117.8 | 823.5 |
| Exponent decoding | – | 930.8 | – |
| Computation | 241.5 | 241.5 | 290.7 |
| Measured total | 1359.3 | 2290.1 | 1114.2 |
Table 5. Decode-attention time breakdown over a 64K-token context. Times are in s. Decomp.-first decompresses the cache before running attention.
The time saved on KV loading outweighs the overhead of on-the-fly reconstruction. Decompressing the cache before attention adds the cost of reading and decompressing the compressed KV data while retaining the full KV loading and computation costs of standard attention. Fusing reconstruction with attention avoids this separate decompression step and reduces KV loading time.
Conclusion
We presented Group-Adaptive Additive Codec, a lossless KV compression system that reduces both the memory footprint and the memory traffic of long-context LLM inference on bandwidth-constrained devices. By exploiting the concentration of BF16 exponents in contiguous ranges, GAAC preserves the original KV bit patterns while enabling lightweight on-the-fly reconstruction within attention. Combined with fill-then-finalize compression and paged management of compressed and raw blocks, this representation translates storage savings into lower decoding latency and higher serving throughput. Across the evaluated models, 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%. GAAC accelerates decode attention on both Apple M4 (1.02–1.25) and NVIDIA DGX Spark (1.19–1.27), demonstrating portability across architectures.
Limitations
The effectiveness of GAAC depends on redundancy in the exponent fields of the stored KV format. GAAC provides less additional lossless compression on FP8 KV caches, particularly E4M3, than on BF16 KV caches. Exponent windows are calibrated offline and transfer well across the evaluated tasks. However, RoPE-rotated keys can exhibit positional drift when calibration contexts are shorter than deployment contexts. The codec remains lossless under such shifts, but wider encoding modes and raw fallbacks reduce compression efficiency.
The serving design also trades some space efficiency for predictable allocation. Performance gains depend on the workload: on-the-fly reconstruction is most beneficial when KV memory reads dominate attention latency, whereas shorter contexts and models with fewer full-attention layers benefit less. Our fused kernel accelerates decoding; prefill uses standard attention and does not receive the same direct acceleration. Our evaluation covers selected LLMs on DGX Spark and Apple M4 and does not include power or energy measurements.
Ethical Considerations
GAAC is an inference-efficiency method and does not alter model weights, training data, or the reconstructed BF16 KV values. It therefore introduces no additional approximation at the KV representation level, but it does not address the safety, privacy, bias, or reliability risks of the underlying language model and its application. Efficiency improvements may also increase overall deployment volume, so lower per-request resource use should not be interpreted as a demonstrated reduction in aggregate environmental impact.
References
Model Quality under Activation Quantization
Our target deployment uses aggressive weight quantization to reduce memory traffic from weight reads while retaining higher-precision activations whenever possible. To measure the additional quality loss from activation quantization, we evaluate Qwen3.8-27B on mathematical reasoning and code-generation benchmarks using W4A16, W4AFP8, W4A4, W3A16, and W3AFP8. W4-GAAC and W3-GAAC denote the use of activations compressed via GAAC. All variants use the same prompt template, decoding configuration, and evaluation harness.
Table 6 shows that reducing activation precision introduces additional quality degradation, particularly on challenging reasoning and coding tasks. Under 4-bit weights, W4AFP8 achieves 76.7% on AIME 2024 and 95.1% on HumanEval, compared with 80.0% and 97.0% for W4A16. More aggressive W4A4 quantization further reduces HumanEval accuracy to 93.9% and LiveCodeBench pass@1 from 37.0% to 34.9%. The degradation becomes more pronounced with 3-bit weights: W3AFP8 drops AIME 2024 accuracy from 56.7% to 50.8% relative to W3A16. In contrast, W4-GAAC and W3-GAAC preserve the benchmark results of their respective A16 baselines while losslessly reducing KV storage to approximately 11.8 bits per element. These results motivate combining aggressive weight quantization with lossless KV compression to reduce memory traffic without introducing additional numerical approximation.
| Benchmark | Metric | W4A16 | W4AFP8 | W4A4 | W4-GAAC | W3A16 | W3AFP8 | W3-GAAC |
| GSM8K | Acc. (%) | 96.1 | 96.0 | 95.4 | 96.1 | 95.8 | 94.5 | 95.8 |
| MATH-500 | Acc. (%) | 84.0 | 83.4 | 83.2 | 84.0 | 78.8 | 77.8 | 78.8 |
| AIME 2024 | Acc. (%) | 80.0 | 76.7 | 76.7 | 80.0 | 56.7 | 50.8 | 56.7 |
| HumanEval | pass@1 (%) | 97.0 | 95.1 | 93.9 | 97.0 | 90.8 | 88.6 | 90.8 |
| MBPP | pass@1 (%) | 78.8 | 78.6 | 77.4 | 78.8 | 74.8 | 73.6 | 74.8 |
| LiveCodeBench | pass@1 (%) | 37.0 | 36.4 | 34.9 | 37.0 | 37.8 | 35.8 | 37.8 |
Table 6. Qwen3.8-27B performance under different weight and activation precisions, with GAAC results for comparison.
Relation to FP8 KV
FP8 KV quantization reduces the storage cost of each element, while GAAC can further compress the quantized cache without introducing additional numerical error. We evaluate lossless exponent coding on model-generated KV states converted to FP8 (Table 7). GAAC achieves an additional 1.09–
| KV Format | Additional Ratio | Total vs. BF16 |
| BF16 | 1.405–1.406× | 1.405–1.406× |
| FP8 E4M3 | 1.09–1.14× | 2.18–2.28× |
| FP8 E5M2 | 1.25–1.30× | 2.50–2.60× |
Table 7. Lossless compressibility after FP8 KV quantization.
1.14× compression on E4M3 and 1.25–1.30× on E5M2, corresponding to total compression ratios of 2.18–2.28× and 2.50–2.60× relative to BF16, respectively. These results demonstrate that lossless KV compression complements FP8 quantization, providing further memory savings without any approximation beyond the initial FP8 conversion.
RoPE Positional Drift
Table 8: Overflow storage in layer 5 of Qwen3-8B over a 32K-token prompt. For each successive 4K-token segment, we report the maximum across heads of the mean overflow bytes per block. Parentheses indicate calibration prompt lengths.
| Segment | K (≤8K) | K (32K) | V (≤8K) |
| 0–4K | 38 | 40 | 39 |
| 4–8K | 37 | 37 | 38 |
| 8–12K | 52 | 37 | 39 |
| 12–16K | 117 | 37 | 41 |
| 16–20K | 177 | 38 | 38 |
| 20–24K | 183 | 37 | 40 |
| 24–28K | 185 | 37 | 40 |
| 28–32K | 190 | 39 | 40 |
Table 8.
We isolate positional distribution shift by applying windows calibrated on short prompts to a 32K prompt. Table 8 shows that K overflow grows with position when the windows are calibrated on prompts of at most 8K tokens. Mean overflow per tile stays at 37–38 bytes up to 8K, then rises to 117 bytes at 12–16K and 190 bytes at 28–32K, about five times the short-context level. RoPE applies position-dependent rotations to keys, causing their exponents to shift outside windows calibrated on short prompts. Values are not rotated, and their overflow remains at 38–41 bytes throughout the prompt. Calibrating on 32K prompts removes the drift: K overflow then stays at 37–40 bytes in every segment. We therefore calibrate on prompts as long as the deployment context.
Storage
Table 9 details the storage layout of a compressed 16-token tile, with a 256-dimensional head represented as two 128-dimensional half-heads. Beyond the unchanged sign and mantissa bits, each tile stores 3-bit exponent codes, 2-bit mode tags, and escape bytes for 4-bit and raw groups. Additional overhead comes from fixed-capacity escape slots, block-table entries, and occasional raw fallbacks. Including these costs, the average storage requirement is 3,016 bytes per tile, corresponding to 11.783 bits per element.
| Field | Bytes | Bits / elem. |
| Sign and mantissa | 2048 | 8.000 |
| 3-bit exponent codes | 768 | 3.000 |
| Mode tags (2 bits / group) | 64 | 0.250 |
| Escape bytes (mean) | 32 | 0.125 |
| Block-table entry | 4 | 0.016 |
| Unused escape slot | 96 | 0.375 |
| Fallback raw tiles | 4.4 | 0.017 |
| Exponent bases (mean) | 0.06 | 0.000 |
| Total | 3016 | 11.783 |
Table 9. Storage breakdown for a compressed 16-token tile covering 128 head dimensions, measured on Qwen3.8-27B.
Attention Baseline Selection
We compare GAAC with FA2 and FlashInfer in the 64K decode-attention setting of § 4.4 (Table 10). All methods are measured in the same process. We report the mean and standard deviation of the median per-call latency across five runs. FlashInfer achieves similar latency to FA2 on DGX Spark, and GAAC is at least faster than either. FA3 and FA4 are not supported on DGX Spark. Its GB10 GPU (compute capability 12.1) lacks the Hopper warpgroup MMA and TMA paths used by FA3, as well as the tcgen05 tensor-core instructions and tensor memory required by FA4 on data-center Blackwell. We therefore use FA2 as the FlashAttention baseline on this device.
| Method | Latency () | Speedup |
| FA2 (block 256) | 1359.3 ±0.8 | |
| FlashInfer (page 16) | 1343.4±0.5 | |
| FlashInfer (page 256) | 1339.8±0.8 | |
| FA3, FA4 | not supported on GB10 | |
| GAAC | 1114.2±2.1 | – |
Table 10. Decode-attention baseline comparison over a 64K-token context on one Qwen3.8-27B layer. Speedup is the baseline latency divided by GAAC latency.
References
- [1]Mark Chen, Jerry Tworek, Heewoo Jun, Qiming Yuan, Henrique Pondé de Oliveira Pinto, Jared Kaplan, Harri Edwards, Yuri Burda, Nicholas Joseph, Greg Brockman, Alex Ray, Raul Puri, Gretchen Krueger, Michael Petrov, Heidy Khlaaf, Girish Sastry, Pamela Mishkin, Brooke Chan, Scott Gray, and 39 others. 2021. Evaluating large language models trained on code. CoRR, abs/2107.03374.arxiv.org/abs/2107.03374
- [2]Karl Cobbe, Vineet Kosaraju, Mohammad Bavarian, Mark Chen, Heewoo Jun, Lukasz Kaiser, Matthias Plappert, Jerry Tworek, Jacob Hilton, Reïchiro Nakano, Christopher Hesse, and John Schulman. 2021. Training verifiers to solve math word problems. CoRR, abs/2110.14168.arxiv.org/abs/2110.14168
- [3]Guohao Dai, Ke Hong, Qiuli Mao, Xiuhong Li, Jiaming Xu, Haofeng Huang, Hongyu Xia, Xuefei Ning, Shengen Yan, Yun Liang, and Yu Wang. 2025. FlashDecoding++Next: High Throughput LLM Inference With Latency and Memory Optimization. IEEE Transactions on Computers, 74(10):3263–3276.
- [4]Tri Dao. 2024. FlashAttention-2: Faster attention with better parallelism and work partitioning. In The Twelfth International Conference on Learning Representations, ICLR 2024, Vienna, Austria, May 7-11, 2024. OpenReview.net.arxiv.org/abs/2307.08691
- [5]Tri Dao, Daniel Y. Fu, Stefano Ermon, At r i Rudra, and Christopher Ré. 2022. FlashAttention: fast and memory-efficient exact attention with IO-awareness. In Advances in Neural Information Processing Systems 35: Annual Conference on Neural Information Processing Systems 2022, NeurIPS 2022, New Orleans, LA, USA, November 28 - December 9, 2022.arxiv.org/abs/2205.14135
- [6]Ruibo Fan, Xiangrui Yu, Xinglin Pan, Zeyu Li, Weile Luo, Qiang Wang, Wei Wang, and Xiaowen Chu. 2026. ZipServ: fast and memory-efficient LLM inference with hardware-aware lossless compression. In Proceedings of the 31st ACM International Conference on Architectural Support for Programming Languages and Operating Systems, Volume 2, ASPLOS 2026, Pittsburgh, PA, USA, March 22-26, 2026, pages 2264–2280. ACM.
- [7]Elias Frantar, Saleh Ashkboos, Torsten Hoefler, and Dan Alistarh. 2022. GPTQ: accurate post-training quantization for generative pre-trained transformers. CoRR, abs/2210.17323.arxiv.org/abs/2210.17323
- [8]Yipin Guo, Arun M. George, Jie Fu, Tareq Mahmoud, Sixue Xing, and Siddharth Joshi. 2026. QTEA: ternary LLMs with sparse residual salient weight and by-column optimization. CoRR, abs/2609.00224.arxiv.org/abs/2609.00224
- [9]Yipin Guo and Siddharth Joshi. 2026. SplitZip: ultra fast lossless KV compression for disaggregated LLM serving. CoRR, abs/2605.01708.arxiv.org/abs/2605.01708
- [10]Yipin Guo, Yilin Lang, and Qinyuan Ren. 2024. GP-TQT: quantize large language models twice to push the efficiency. In IEEE International Conference on Cybernetics and Intelligent Systems, CIS 2024, and IEEE International Conference on Robotics, Automation and Mechatronics, RAM 2024, Hangzhou, China, August 8-11, 2024, pages 368–373. IEEE.
- [11]Dan Hendrycks, Collin Burns, Steven Basart, Andy Zou, Mantas Mazeika, Dawn Song, and Jacob Steinhardt. 2021. Measuring massive multitask language understanding. In 9th International Conference on Learning Representations, ICLR 2021, Virtual Event, Austria, May 3-7, 2021. OpenReview.net.
- [12]Moshik Hershcovitch, Andrew Wood, Leshem Choshen, Guy Girmonsky, Roy Leibovitz, Or Ozeri, Ilias Enn-mouri, Michal Malka, Sang (Peter) Chin, Swami-nathan Sundararaman, and Danny Harnik. 2025. ZipNN: lossless compression for AI models. In 18th IEEE International Conference on Cloud Computing, CLOUD 2025, Helsinki, Finland, July 7-12, 2025, pages 186–198. IEEE.
- [13]Coleman Hooper, Sehoon Kim, Hiva Mohammadzeh, Michael W. Mahoney, Yakun Sophia Shao, Kurt Keutzer, and Amir Gholami. 2024. KVQuant: towards 10 million context length LLM inference with KV cache quantization. In Advances in Neural Information Processing Systems 37: Annual Conference on Neural Information Processing Systems 2024, NeurIPS 2024, Vancouver, BC, Canada, December 10 - 15, 2024.arxiv.org/abs/2401.18079
- [14]Jeff Johnson. 2022. DietGPU: GPU-based lossless compression for numerical data. GitHub repository.
- [15]Woosuk Kwon, Zhuohan Li, Siyuan Zhuang, Ying Sheng, Lianmin Zheng, Cody Hao Yu, Joseph Gonzalez, Hao Zhang, and Ion Stoica. 2023. Efficient memory management for large language model serving with PagedAttention. In Proceedings of the 29th Symposium on Operating Systems Principles, SOSP 2023, Koblenz, Germany, October 23-26, 2023, pages 611–626. ACM.arxiv.org/abs/2309.06180
- [16]Yuhong Li, Yingbing Huang, Bowen Yang, Bharat Venkit esh, Acyr Locatelli, Hanchen Ye, Tianle Cai, Patrick Lewis, and Deming Chen. 2024. SnapKV: LLM knows what you are looking for before generation. In Advances in Neural Information Processing Systems 37: Annual Conference on Neural Information Processing Systems 2024, NeurIPS 2024, Vancouver, BC, Canada, December 10 - 15, 2024.arxiv.org/abs/2404.14469
- [17]Zheng Li, Weiyan Wang, Ruiyuan Li, Chao Chen, Xianlei Long, Linjiang Zheng, Quanquing Xu, and Chuanhui Yang. 2025. A high-throughput GPU framework for adaptive lossless compression of floating-point data. CoRR, abs/2511.04140.arxiv.org/abs/2511.04140
- [18]Ji Lin, Jiaming Tang, Haotian Tang, Shang Yang, Wei-Ming Chen, Wei-Chen Wang, Guangxuan Xiao, Xingyu Dang, Chuang Gan, and Song Han. 2024. AWQ: activation-aware weight quantization for on-device LLM compression and acceleration. In Proceedings of the Seventh Annual Conference on Machine Learning and Systems, MLSys 2024, Santa Clara, CA, USA, May 13-16, 2024. mlsys.org.DOI
- [19]Wenxiang Lin, Xinglin Pan, Ruibo Fan, Shaohuai Shi, and Xiaowen Chu. 2026. ZipCCL: efficient lossless data compression of communication collectives for accelerating LLM training. In Proceedings of the ACM SIGCOMM 2026 Conference, SIGCOMM 2026, Colorado Convention Center, Denver, CO, USA, August 17-21, 2026, pages 207–221. ACM.
- [20]Zirui Liu, Jiayi Yuan, Hongye Jin, Shaochen (Henry) Zhong, Zhaozhuo Xu, Vladimir Braverman, Beidi Chen, and Xia Hu. 2024. KIVI: A tuning-free asymmetric 2bit quantization for KV cache. In Forty-first International Conference on Machine Learning, ICML 2024, Vienna, Austria, July 21-27, 2024, volume 235 of Proceedings of Machine Learning Research, pages 32332–32344. PMLR / OpenReview.net.DOI
- [21]Llama Team. 2024. The Llama 3 herd of models. CoRR, abs/2407.21783.
- [22]Mitchell P. Marcus, Beatrice Santorini, and Mary Ann Marcinkiewicz. 1993. Building a large annotated corpus of English: The Penn Treebank. Comput. Linguistics, 19(2):313–330.
- [23]Stephen Merity, Caiming Xiong, James Bradbury, and Richard Socher. 2017. Pointer sentinel mixture models. In 5th International Conference on Learning Representations, ICLR 2017, Toulon, France, April 24-26, 2017, Conference Track Proceedings. OpenReview.net.arxiv.org/abs/1609.07843
- [24]NVIDIA. 2026. nvCOMP. GPU compression library. Accessed October 4, 2026.
- [26]Qwen Team. 2025. Qwen3 technical report. CoRR, abs/2505.09388.
- [28]Haoran You, Yipin Guo, Yichao Fu, Wei Zhou, Huihong Shi, Xiaofan Zhang, Souvik Kundu, Amir Yazdanbakhsh, and Yingyan (Celine) Lin. 2024. ShiftAddLLM: accelerating pretrained LLMs via post-training multiplication-less reparameterization. In Advances in Neural Information Processing Systems 37: Annual Conference on Neural Information Processing Systems 2024, NeurIPS 2024, Vancouver, BC, Canada, December 10 - 15, 2024.arxiv.org/abs/2406.05981
- [29]Tianyi Zhang, Mohsen Hariri, Shaochen (Henry) Zhong, Vipin Chaudhary, Yang Sui, Xia Hu, and Anshumali Shrivastava. 2025. 70% size, 100% accuracy: Lossless LLM compression for efficient GPU inference via dynamic-length float (DFloat11). In Advances in Neural Information Processing Systems 38: Annual Conference on Neural Information Processing Systems 2025, NeurIPS 2025, San Diego, CA, USA, December 2-7, 2025 / Mexico City, Mexico, November 30 - December 5, 2025.arxiv.org/abs/2504.11651
- [30]Zhenyu Zhang, Ying Sheng, Tianyi Zhou, Tianlong Chen, Lianmin Zheng, Ruisi Cai, Zhao Song, Yuandong Tian, Christopher Ré, Clark W. Barrett, Zhangyang Wang, and Beidi Chen. 2023. H2O: heavy-hitter oracle for efficient generative inference of large language models. In Advances in Neural Information Processing Systems 36: Annual Conference on Neural Information Processing Systems 2023, NeurIPS 2023, New Orleans, LA, USA, December 10 - 16, 2023.arxiv.org/abs/2306.14048