ZipServ: GPU-Optimized Lossless Compression

• ZipServ is a lossless compression framework that replaces variable-length entropy coding with a fixed-length TCA-TBE format, enabling efficient GPU-based LLM inference.
• The system fuses decompression with GEMM operations via its ZipGEMM kernel, achieving up to 2.21× speedup over NVIDIA cuBLAS while preserving BF16 numerical precision.
• Leveraging GPU Tensor Core architecture, ZipServ reduces weight storage by about 30% and alleviates memory bandwidth bottlenecks, particularly in decode-heavy regimes.

Searching arXiv for ZipServ and closely related systems to ground the article in current papers. ZipServ is a lossless compression framework for LLM inference that is explicitly co-designed for GPU execution. It targets the memory-bandwidth bottlenecks of bit-exact serving by replacing variable-length entropy-coded weight representations with a fixed-length format, Tensor-Core-Aware Triple Bitmap Encoding (TCA-TBE), and by coupling that format to a fused decompression-GEMM kernel, ZipGEMM. In the reported evaluation, ZipServ reduces model size by up to 30%, achieves up to 2.21× kernel-level speedup over NVIDIA cuBLAS, and improves end-to-end inference by an average of 1.22× over vLLM, while preserving exact BF16 weights and therefore exact model numerics (Fan et al., 18 Mar 2026).

1. Problem setting and design objective

ZipServ is motivated by the observation that Transformer inference is dominated by dense GEMMs,

$Y_{M\times N} = W_{M\times K} X_{K\times N},$

with distinct operating regimes in prefill and decode. Prefill uses large $N=\text{batch}\times\text{prompt length}$ and is comparatively compute-intensive, whereas decode generates one token per step per batch element and is therefore memory-bandwidth bound because the same weights are streamed repeatedly from HBM with little compute per byte moved (Fan et al., 18 Mar 2026).

The framework is situated within the broader problem of bit-exact model serving. Lossy schemes such as 8-bit or 4-bit quantization can reduce memory substantially, but they alter weights. ZipServ instead addresses the deployment regime in which every BF16 bit must be recovered exactly. The system-level claim is that lossless compression need not be merely a storage optimization: if the encoding and the execution kernel are aligned with GPU architecture, it can also accelerate inference (Fan et al., 18 Mar 2026).

The paper identifies two reasons existing lossless schemes slow inference. At the kernel level, variable-length bitstreams from Huffman, ANS, or related entropy coders induce data-dependent decoding, thread divergence, serialized pointer advancement, and irregular memory access under SIMT execution. At the system level, many prior pipelines decompress into a full BF16 buffer in HBM and then invoke cuBLAS, which introduces redundant global-memory traffic. ZipServ addresses both mismatches by using a fixed-length code and a fused compute path (Fan et al., 18 Mar 2026).

2. Tensor-Core-Aware Triple Bitmap Encoding

The statistical basis of TCA-TBE is the exponent distribution of BF16 LLM weights. Across LLaMA-3, Mistral, Qwen2.5, and Gemma3, the exponent field is reported to be highly skewed and near-Gaussian: the top-3 exponents cover more than 67% of weights, the top-7 exponents cover more than 95%, and the exponent entropy is 2.57–2.74 bits rather than the 8 bits stored in BF16. This implies a theoretical lower bound of about 10.6 bits per value when the 8-bit sign-plus-mantissa payload is combined with exponent entropy, corresponding to a maximum lossless compression ratio of about $16/10.6 \approx 1.51\times$ (Fan et al., 18 Mar 2026).

A further empirical observation is structural rather than merely entropic: for 99.6% of matrices, the top-7 exponents are numerically contiguous. ZipServ therefore defines a layer-wise base exponent,

$e_{\text{base}}=\min(E_{\text{top}})-1,$

where $E_{\text{top}}$ is the set of the 7 most frequent exponents. Any exponent in that top set is encoded by a 3-bit codeword $c\in[1,7]$, while all other exponents use a fallback code $000$. High-frequency values store only sign plus mantissa explicitly, and the exponent is reconstructed arithmetically as

$e=e_{\text{base}}+c.$

Rare values store the full 16-bit BF16 word (Fan et al., 18 Mar 2026).

The codeword length is selected by minimizing the expected storage

$$\text{AverageBits}(n)=r_n\cdot(n+8)+(1-r_n)\cdot(n+16),$$

where $r_n$ is the fraction covered by the top $N=\text{batch}\times\text{prompt length}$0 exponents. The reported empirical result is that $N=\text{batch}\times\text{prompt length}$1, giving about 11.3 bits per value for $N=\text{batch}\times\text{prompt length}$2, whereas $N=\text{batch}\times\text{prompt length}$3 or $N=\text{batch}\times\text{prompt length}$4 yields about 12.1–12.4 bits. ZipServ therefore adopts 3-bit codes as the balance point between compression ratio and decoding simplicity (Fan et al., 18 Mar 2026).

The “triple bitmap” structure replaces packed variable-length streams with three fixed 64-bit bit planes per $N=\text{batch}\times\text{prompt length}$5 FragTile. For each of the 64 positions, the three planes store the three bits of the codeword. The tile also contains a PackedSignMantissa buffer of 8-bit values for high-frequency exponents and a FullValue buffer of 16-bit BF16 values for fallback cases. Tiles are organized hierarchically as $N=\text{batch}\times\text{prompt length}$6 FragTiles, $N=\text{batch}\times\text{prompt length}$7 TensorCoreTiles, and larger BlockTiles, with layout chosen to match the operand fragment organization expected by Tensor Core instructions (Fan et al., 18 Mar 2026).

3. Execution model and ZipGEMM

ZipServ is stage-aware. Offline, a CPU-side compressor profiles exponent distributions per layer and emits TCA-TBE-compressed weights plus metadata such as BaseExp. Online, inference uses two different execution paths: a decoupled decompression-plus-cuBLAS strategy for prefill, where decompression overhead is reported to remain below 4%, and a fused path for decode, where the memory bottleneck is most acute (Fan et al., 18 Mar 2026).

The fused decode kernel, ZipGEMM, implements a “load-compressed, compute-decompressed” dataflow. Compressed weights are read from HBM into shared memory; a warp then decompresses directly into registers that are immediately consumed by Tensor Cores; the decompressed weights are never materialized as a full BF16 matrix in global memory. This removes the intermediate decompressed buffer that is characteristic of decoupled pipelines (Fan et al., 18 Mar 2026).

The kernel follows a conventional tiled GEMM structure only at the outer level. For each $N=\text{batch}\times\text{prompt length}$8-slice, thread blocks load compressed weight tiles and BF16 activation tiles into shared memory using vectorized instructions and cp.async. At warp scope, each warp reconstructs the relevant $N=\text{batch}\times\text{prompt length}$9 FragTiles from the three bitmaps and the two value buffers, packs the result into the fragment layout expected by mma.sync.m16n8k16, loads the activation fragments with LDSM.M88, and issues Tensor Core instructions (Fan et al., 18 Mar 2026).

The decompression procedure is built around a spatial indicator mask

$16/10.6 \approx 1.51\times$0

For a position $16/10.6 \approx 1.51\times$1, $16/10.6 \approx 1.51\times$2 means a top-7 exponent and $16/10.6 \approx 1.51\times$3 means fallback. Each thread computes the index into the PackedSignMantissa or FullValue buffer using prefix counts implemented with bit masks and __popc. The 3-bit code itself is reconstructed from the three planes, and the exponent is recovered by the constant-time arithmetic rule $16/10.6 \approx 1.51\times$4. The critical point is that the per-element work is fixed and thread-local; no thread performs data-dependent parsing of a variable-length stream (Fan et al., 18 Mar 2026).

ZipGEMM further uses a two-level software pipeline. At coarse granularity, double-buffered shared-memory tiles overlap prefetch and compute. At finer granularity, ALU-side decompression of slice $16/10.6 \approx 1.51\times$5 is interleaved with Tensor Core execution on slice $16/10.6 \approx 1.51\times$6. The intended effect is to hide integer and bit-manipulation overhead behind Tensor Core activity (Fan et al., 18 Mar 2026).

4. Hardware rationale and roofline interpretation

A central argument of ZipServ is that fixed-length coding is more important than maximal entropy efficiency in the GPU decode path. Huffman-, ANS-, and rANS-based codecs can approach the Shannon limit more closely, but they produce variable-length symbol boundaries and data-dependent control flow. ZipServ accepts a modest gap between 11.3 bits per value and the theoretical 10.6-bit bound in order to obtain warp-synchronous, branch-free decompression (Fan et al., 18 Mar 2026).

The roofline analysis formalizes the system-level difference. For an uncompressed BF16 GEMM, the compute intensity is

$16/10.6 \approx 1.51\times$7

For a decoupled compressed pipeline with compression ratio $16/10.6 \approx 1.51\times$8, ZipServ gives

$16/10.6 \approx 1.51\times$9

The extra writeback and reread of decompressed weights reduce compute intensity substantially (Fan et al., 18 Mar 2026).

For the fused design, the paper gives

$e_{\text{base}}=\min(E_{\text{top}})-1,$0

This is reported as about 50% higher than even the uncompressed GEMM case, because the decode path loads fewer weight bytes from HBM and avoids the decompressed intermediate. A plausible implication is that ZipServ’s speedups are strongest precisely in the decode-heavy regimes where conventional BF16 serving is most bandwidth-limited (Fan et al., 18 Mar 2026).

The implementation is correspondingly hardware-specific. The paper targets NVIDIA Tensor Core GPUs, aligns TensorCoreTiles with mma.sync fragment structure, uses shared-memory layouts chosen to avoid bank conflicts, and relies on coalesced loads of fixed-size bitmaps and compact value arrays. The reported micro-profiling shows DRAM reads dropping by 29.3% relative to cuBLAS on a representative kernel, while Tensor Core utilization remains 71.6% of cuBLAS and shared-memory bank conflicts are nearly eliminated relative to entropy-code baselines (Fan et al., 18 Mar 2026).

5. Empirical performance

The reported memory savings are consistent with the 11.3-bit effective storage target. For several BF16 models, compressed weight footprints fall to about 71–73% of the original. Examples include LLaMA-3.1-8B from 14.96 GB to 10.83–11.18 GB, Mistral-24B from 43.92 GB to 31.30 GB, and LLaMA-3.1-70B from 131.56 GB to 93.52 GB (Fan et al., 18 Mar 2026).

At kernel level, the principal comparison is against NVIDIA cuBLAS Tensor Core GEMM and several prior lossless codecs. On RTX4090, ZipGEMM achieves 1.31× average speedup over cuBLAS and up to 1.71×. On L40S, the average speedup is 1.36× and the maximum is 2.21×. By contrast, DietGPU, nvCOMP rANS, and DFloat11 are reported as slower than cuBLAS in the same setting, with normalized performance ranges of 0.17×–0.34× depending on GPU and baseline (Fan et al., 18 Mar 2026).

Setting	Reported result	Citation
Model size reduction	Up to 30%	(Fan et al., 18 Mar 2026)
Kernel-level speedup over cuBLAS	Up to 2.21×	(Fan et al., 18 Mar 2026)
End-to-end throughput over vLLM	1.22× average	(Fan et al., 18 Mar 2026)

End-to-end serving measurements were integrated into vLLM and compared against vLLM, HuggingFace Transformers, and a DFloat11-based pipeline. The paper reports 17.6% average latency reduction versus vLLM, 60.79% versus Transformers, and 82.13% versus DFloat11, together with throughput gains of 1.22×, 3.18×, and 8.52× respectively. For LLaMA-3.1-8B at 2048 generated tokens and batch size 32, ZipServ reaches about 1105 tokens/s, or 1.66× vLLM (Fan et al., 18 Mar 2026).

The freed memory is also used operationally. Weight compression releases about 28–30% of GPU memory, and the paper gives an example in which KV-cache capacity grows from 5.07 GB to 8.60 GB, a 1.70× increase. This suggests that the throughput gain is not solely a kernel effect; it also arises from larger feasible batches and longer retained contexts (Fan et al., 18 Mar 2026).

6. Relation to adjacent compression systems and limitations

ZipServ occupies a specific point in the compression-for-ML design space. ZipNN targets lossless storage and distribution of AI model weights by exploiting exponent-byte skew and using Huffman-only coding, often saving 33% and sometimes more than 50% of model size, with better compression and decompression speed than vanilla compressors on stored models (Hershcovitch et al., 2024). DeltaZip addresses a different serving problem: concurrent serving of many full-parameter fine-tuned LLM variants by compressing model deltas by up to 10× and improving throughput by 2× to 12× (Yao et al., 2023). SplitZip focuses not on weights but on BF16 KV-cache transfer in prefill-decode disaggregation, achieving 613.3 GB/s compression throughput, 2181.8 GB/s decompression throughput, and up to 1.30× TTFT speedup (Guo et al., 3 May 2026). A plausible systems interpretation is that these methods are complementary rather than interchangeable: ZipServ targets the in-GPU decode path for bit-exact weight access, whereas the others target model distribution, multi-variant serving, or inter-node KV movement.

The framework also has explicit limitations. It is designed around NVIDIA Tensor Cores and BF16 weight distributions; extension to FP16 or FP8 is described as non-trivial and would require re-profiling exponent structure. It is shape-sensitive: some small $e_{\text{base}}=\min(E_{\text{top}})-1,$1 layers can be slower than cuBLAS because the kernel configuration is not tuned for tiny matrices. Its benefits are strongest on bandwidth-constrained consumer and inference-oriented GPUs such as RTX4090, L40S, and RTX5090; on bandwidth-rich accelerators such as A100 or H800, some workloads do not show the same advantage (Fan et al., 18 Mar 2026).

The conceptual significance of ZipServ is therefore not merely the introduction of another lossless codec. Its substantive claim is that lossless compression can be turned into a first-class inference optimization if the format is fixed-length, tile-aligned, and directly consumable by Tensor Core kernels. In that sense, ZipServ reframes “lossless model compression” from an archival or download problem into a GPU kernel design problem (Fan et al., 18 Mar 2026).