ZipServ: Fast and Memory-Efficient LLM Inference with Hardware-Aware Lossless Compression

Abstract: Lossless model compression holds tremendous promise for alleviating the memory and bandwidth bottlenecks in bit-exact LLM serving. However, existing approaches often result in substantial inference slowdowns due to fundamental design mismatches with GPU architectures: at the kernel level, variable-length bitstreams produced by traditional entropy codecs break SIMT parallelism; at the system level, decoupled pipelines lead to redundant memory traffic. We present ZipServ, a lossless compression framework co-designed for efficient LLM inference. ZipServ introduces Tensor-Core-Aware Triple Bitmap Encoding (TCA-TBE), a novel fixed-length format that enables constant-time, parallel decoding, together with a fused decompression-GEMM (ZipGEMM) kernel that decompresses weights on-the-fly directly into Tensor Core registers. This "load-compressed, compute-decompressed" design eliminates intermediate buffers and maximizes compute intensity. Experiments show that ZipServ reduces the model size by up to 30%, achieves up to 2.21x kernel-level speedup over NVIDIA's cuBLAS, and expedites end-to-end inference by an average of 1.22x over vLLM. ZipServ is the first lossless compression system that provides both storage savings and substantial acceleration for LLM inference on GPUs.

• The paper introduces ZipServ, a hardware-aware lossless compression system that uses fixed-length 3-bit bitmap encoding and fused Tensor Core decoding to avoid SIMT divergence and redundant memory traffic.
• ZipGEMM outperforms cuBLAS by an average of 1.31× on RTX 4090 and 1.36× on L40S, while reducing DRAM reads by 29.3% and preserving bit-exact model weights.
• Integrated with vLLM, ZipServ achieves an average 1.22× speedup and up to 1.66× higher throughput, shrinks weights to roughly 71–72% of their original size, and expands KV-cache capacity by 1.70×.

ZipServ is a system presented at ASPLOS '26 that makes lossless compression of LLM weights a source of inference acceleration rather than overhead, by co-designing a compression format with GPU architecture and Tensor Core execution (2603.17435). The central claim, supported empirically across multiple GPUs and models, is that prior lossless approaches fail at inference not because compression itself is incompatible with serving, but because of two specific algorithm–hardware mismatches: variable-length entropy-coded bitstreams conflict with SIMT parallelism, and decoupled decompress-then-compute pipelines introduce redundant global-memory traffic.

Motivation: why existing lossless pipelines are slow

Lossy techniques such as GPTQ, AWQ, and SparseGPT dominate LLM acceleration but risk accuracy degradation; the paper cites evidence that MXFP4 quantization drops LiveCodeBench accuracy from 56.0% to 36.2%, and even GPTQ-int8 can lose up to 11.1% on long-context reasoning tasks. This motivates bit-exact inference. Existing lossless systems (LMC, ZipNN for storage; NeuZip, DietGPU for training; DFloat11 for inference) exploit the well-known redundancy in BF16 exponent fields, but when integrated into serving they incur decompression costs of 1.56–3.44× the core GEMM time.

The paper identifies two structural causes. Kernel level: Huffman- or ANS-coded bitstreams require serialized, data-dependent decoding — bitstream partitioning with boundary metadata, hierarchical LUT lookups whose varying symbol lengths cause warp divergence, and serial pointer advancement. Measured on an L40S, optimized ANS (DietGPU) and Huffman (DFloat11) decompressors reach only 43.7% and 76.5% of peak memory bandwidth, respectively. System level: materializing fully decompressed weights in global memory before consumption reduces compute intensity. Under a Roofline analysis with the observed compression ratio of ~1.51×, a decoupled pipeline degrades compute intensity by roughly 62% relative to standard BF16 GEMM for a 4096×4096 weight matrix at batch sizes 8–64 — meaning the compressed format actually makes the memory-bound decode phase slower per byte saved unless fusion is used.

Compressibility structure and TCA-TBE

The empirical foundation is a statistical property the authors call exponent contiguity. Across 3,875 weight matrices from Gemma-3, Mistral, Qwen2.5, and LLaMA-3.1 families, the top-7 most frequent exponents form a numerically contiguous window ${e^\star,\dots,e^\star+6}$ in 99.6% of matrices, covering 97.1% of weights on average; measured exponent-field entropy is only 2.57–2.74 bits against its 8-bit allocation. An appendix proves this follows from unimodality of the exponent probability distribution under a zero-mean Gaussian weight assumption — though the proof relies on the Gaussian approximation and ignores mantissa rounding effects, an assumption stated plainly in the paper.

This contiguity enables Tensor-Core-Aware Triple Bitmap Encoding (TCA-TBE): each weight receives a fixed-length 3-bit codeword. Codewords 001–111 encode membership in the top-7 contiguous exponents, decoded implicitly as exponent = base_exp + codeword with a single integer add; codeword 000 marks outliers stored at full BF16 precision. The expected cost is 11.3 bits per element versus a theoretical lower bound of 10.6 bits, beating both 2-bit (12.4) and 4-bit (12.1) alternatives. Rather than packing codes into a dense bitstream, each $8\times8$ tile's 3-bit codes are split into three independent 64-bit bitmaps (one bit-plane each), guaranteeing coalesced, aligned accesses and branch-free, constant-time decoding — directly addressing the SIMT divergence problem. A three-level tiling hierarchy (FragTile $8\times8$, TensorCoreTile $16\times16$ matching mma.m16n8k16 operands, BlockTile $64\times64$) aligns the layout with register fragment ordering so no runtime coordinate transformation is needed.

ZipGEMM: fused decompression into Tensor Core registers

The system-level remedy is ZipGEMM, a fused kernel following a "load-compressed, compute-decompressed" model. Each thread block iterates over split-K chunks through four stages: asynchronous vectorized loading of compressed tiles to shared memory, warp-level thread-local decoding, activation transfer via LDSM.M88, and mma.sync execution. Decoding uses a spatial indicator mask (bitwise OR of the three bitmaps), dynamic addressing via warp-local prefix sums computed with __popc()/__shfl_sync(), and arithmetic exponent reassembly — all register-resident and branch-free. A two-level software pipeline (tile-wise double buffering plus slice-wise interleaving along K) overlaps decompression ALU work with Tensor Core execution.

The stage-aware strategy applies ZipGEMM only in the memory-bound decode phase; during compute-bound prefill it falls back to a standalone decompression kernel followed by cuBLAS GEMM, where amortized overhead is only ~4% at $N=8192$ and ~2% at $N=16384$. This design choice implies the end-to-end benefit is concentrated exactly where bandwidth matters most, while avoiding regression on large-batch prefill.

Evaluation results

Kernel benchmarks span linear layers from LLaMA3.1 (8B–405B), Qwen2.5 (7B–72B), Gemma3, and Mistral (24B–123B), against cuBLAS_TC, DietGPU, nvCOMP (rANS), and DFloat11:

Method	RTX4090 avg	L40S avg
ZipGEMM	1.31×	1.36× (peak 2.21×)
DietGPU	0.17×	0.20×
nvCOMP	0.19×	0.23×
DFloat11	0.28×	0.34×

All baselines fall below 1× because their decoupled decompression exceeds baseline GEMM time; ZipGEMM is the only variant that beats cuBLAS. Nsight Compute profiling shows the mechanism explicitly: DRAM reads drop 29.3%, ALU utilization rises to 66.0%, yet Tensor Core throughput retains 71.6% of the cuBLAS baseline because the pipeline hides decoding latency; shared-memory bank conflicts are nearly eliminated (~4.7K vs. millions for DietGPU). The standalone ZipServ-Decomp kernel also outperforms DietGPU, nvCOMP, and DFloat11 by 2.14×, 1.83×, and 1.10× respectively.

End-to-end, integrated as a vLLM backend (~3.5K LOC), ZipServ yields an average 1.22× speedup over vLLM (up to 1.66× at batch 32 with 2048-token outputs on LLaMA3.1-8B, reaching 1105 tokens/s), 3.18× over Transformers, and 8.52× over DFloat11, with latency reductions averaging 82.13% versus DFloat11. Weights shrink to ~71–72% of original size (e.g., LLaMA3.1-70B: 131.56 GB → 93.52 GB); freed memory expands the KV cache 1.70× via PagedAttention, converting static savings into throughput.

A notable result concerns hardware tiers: ported unchanged to Blackwell RTX5090, ZipGEMM achieves 1.34× (LLaMA3.1-8B) and 1.87× (Mistral-24B) over cuBLAS, and on RTX4090 it matches or slightly exceeds cuBLAS on A100. On RTX5090, ZipGEMM reduces the gap to H800 from 53.3%/125.7% down to 14.1%/20.8%. Against lossy Marlin W8A16-FP8, ZipGEMM trails by only 1.36× — approximately matching the effective bit-width ratio (~11 vs. 8 bits) — indicating decompression cost is essentially hidden within memory access latency.

Limitations and open questions

The paper concedes several boundaries candidly. Small-matrix layers can regress (O_proj of LLaMA3.1-8B drops to 0.79× on L40S) because fine-grained split-K/tiling tuning was out of scope, though these layers contribute little total FLOP. On HBM-rich training GPUs (A100/H800), ZipGEMM does not always match cuBLAS — attributed to relaxed memory bottlenecks and lower core clocks making the added ALU workload harder to hide — so the approach is targeted at consumer/inference-optimized GPUs rather than universally superior. The compressibility analysis depends on Gaussian weight assumptions and excludes mantissa-rounding effects. Open questions left explicit include adapting TCA-TBE to KV-cache compression, porting to Intel AMX and AMD Matrix Cores, and combining lossless coding atop quantized weights.

Conclusion

ZipServ demonstrates that lossless compression can be transformed from a storage utility into a genuine inference accelerator when the encoding format is fixed-length, bitmap-structured, and aligned to Tensor Core fragment layouts, and when decompression is fused directly into the GEMM kernel. Its strongest evidence is the consistent superiority over cuBLAS in memory-bound regimes on consumer hardware, together with near-lossless parity against FP8 kernels despite preserving bit-exactness — establishing a defensible case for lossless methods wherever numerical reproducibility is required.