---
title: 'ZipGEMM: Dual Hardware-Aware GEMM Approaches'
url: https://www.emergentmind.com/topics/zipgemm
type: topic
---

# ZipGEMM: Dual Hardware-Aware GEMM Approaches

Searching arXiv for the cited papers to ground the article in current literature.
“ZipGEMM” is used in the arXiv literature for two distinct matrix-computation mechanisms. In "ZipServ: Fast and Memory-Efficient LLM Inference with Hardware-Aware Lossless Compression" [2603.17435], ZipGEMM denotes a fused decompression-GEMM kernel for BF16 Large Language Model serving on GPUs, built around Tensor-Core-Aware Triple Bitmap Encoding (TCA-TBE) and a “load-compressed, compute-decompressed” execution pattern. In "SparseZipper: Enhancing Matrix Extensions to Accelerate SpGEMM on CPUs" [2502.11353], the same term is used in the detailed exposition to describe a zipper-based approach to sparse-sparse GEMM on CPUs, where sparse index/value streams are sorted and merged through minimally extended matrix instructions and a 16×16 systolic array. The shared naming reflects a common emphasis on aligning compressed or sparse representations with the execution substrate, but the two systems target different workloads, data formats, and hardware.

## 1. Scope of the term

In [2603.17435], ZipGEMM is a kernel-level mechanism inside ZipServ for bit-exact LLM inference. Its purpose is to avoid the inference slowdowns associated with variable-length entropy codecs by using fixed-length coding and by decompressing weights on-the-fly directly into Tensor Core registers. The relevant operands are dense BF16 weights and activations, and the core system claim is that the fused design removes intermediate buffers while maximizing compute intensity.

In [2502.11353], ZipGEMM refers to the SparseZipper approach to accelerated SpGEMM on CPUs. There, each sparse vector is treated as a stream of `(key,value)` tuples, split into parallel streams, and processed in fixed-size chunks of $R$ tuples using matrix registers and new zipper instructions. The objective is not lossless compression of dense weights, but efficient processing of highly sparse matrices with unstructured sparsity structures.

This suggests that “ZipGEMM” is not a single standardized algorithmic object across the literature. Rather, the term names two hardware-conscious GEMM schemes that each “zip” representation and computation together, but in materially different ways.

## 2. ZipGEMM in ZipServ: TCA-TBE and fused dense GEMM on GPUs

The ZipServ version of ZipGEMM is inseparable from Tensor-Core-Aware Triple Bitmap Encoding. The motivating observation is that in BF16 LLM weights “the 8-bit exponent field is extremely skewed,” and that “the top-7 most frequent exponent values form a contiguous range covering ≈96 % of all weights” [2603.17435]. Instead of a variable-length entropy code, TCA-TBE uses a fixed-length 3-bit codeword per weight for exponents inside that top-7 window, plus the special “000” code for fallback cases. Those 3-bit codewords are stored as three separate 64-bit bitmaps per 8×8 tile, called the “FragTile.”

For a BF16 weight $w$ with sign $s\in\{0,1\}$, exponent $e\in[0,255]$, and mantissa $m\in[0,127]$, the offline encoding procedure first computes the exponent histogram for the entire weight matrix and chooses a length-7 contiguous window
$$
E_{\mathrm{top}}=\{e^\star,e^\star+1,\dots,e^\star+6\}
$$
that maximizes coverage. It then sets a global $\mathrm{BaseExp}=e^\star-1$. If $e_i\in E_{\mathrm{top}}$, the encoder stores
$$
c_i=e_i-\mathrm{BaseExp}\in[1..7]
$$
together with the 1-bit sign and 7-bit mantissa in the dense 8-bit buffer $H$; otherwise it uses fallback code $c_i=0$ and stores the full 16-bit BF16 value in buffer $L$.

The triple-bitmap layout for an 8×8 tile is
$$
B_0=\sum_{i=0}^{63}(c_i\bmod 2)\ll i,\quad
B_1=\sum_{i=0}^{63}(\lfloor c_i/2\rfloor \bmod 2)\ll i,\quad
B_2=\sum_{i=0}^{63}(\lfloor c_i/4\rfloor \bmod 2)\ll i.
$$
Tile-specific offsets into $H$ and $L$ are recorded in a small Offset table. No other padding or variable-length streams are used.

The decode path is intentionally SIMT-friendly. The format is fixed-length; bitmap bit-planes avoid unaligned, data-dependent bit-shifts; and all threads follow the same control flow. Exponent reconstruction is
$$
e_i=\mathrm{BaseExp}+\bigl((B_2\gg i\ \&\ 1)\cdot 4+(B_1\gg i\ \&\ 1)\cdot 2+(B_0\gg i\ \&\ 1)\bigr).
$$
The paper characterizes this as constant-time, parallel decoding aligned to GPU SIMT warps [2603.17435].

## 3. GPU execution model: “load-compressed, compute-decompressed”

ZipGEMM in ZipServ implements the pattern “load-compressed, compute-decompressed.” The kernel splits GEMM into 64×64 block-tiles and 16×16 TensorCore tiles, and “a single CUDA thread block processes one 64×64 BT, sweeping over K in chunks of 16” [2603.17435]. Compressed weights and activations are asynchronously loaded into shared memory; warps then decompress weight fragments into registers and issue Tensor Core MMA instructions.

The thread and warp organization is explicit. Each 8×8 FragTile has 64 elements, and one warp with 32 threads decodes exactly one FragTile by letting each thread handle two elements at positions $p_0=2\cdot \mathrm{laneId}$ and $p_1=2\cdot \mathrm{laneId}+1$. A 16×16 TensorCore tile comprises four FragTiles arranged as a 2×2 grid; four warps decode one FragTile each to produce the full $A_{\mathrm{frag}}$ in registers. Two additional warps load the 16×8 activation fragment from shared memory, after which the group issues two `mma.sync` instructions, one per 16×8 slice. The design is described as preserving SIMT parallelism because all warps move in lock-step.

The memory path is equally central. Double-buffered `cp.async` loads bring compressed bitmaps and value buffers into shared memory, aligned to 128 bits to saturate DRAM bandwidth and bypass L1. No global-memory scratch space for decompressed weights is ever written. Instead, weights are decompressed directly into the register files backing $A_{\mathrm{frag}}$. A two-level software pipeline overlaps load and decode latency with ongoing `mma.sync` execution. While one fragment is multiplying on Tensor Cores, integer and bitwise units decompress the next FragTile.

Fragment-level decoding uses a 64-bit occupancy mask
$$
M=B_0\lor B_1\lor B_2,
$$
with
$$
\mathrm{mask}_i=(1\ll i)-1,\quad
\mathrm{popH}_i=\mathrm{popcount}(M\ \&\ \mathrm{mask}_i),\quad
\mathrm{popAll}_i=i,\quad
\mathrm{popL}_i=\mathrm{popAll}_i-\mathrm{popH}_i.
$$
If $M_i=1$, the decoder reads an 8-bit sign-plus-mantissa from $H$, reconstructs the exponent from the bitmaps and $\mathrm{BaseExp}$, and forms the BF16 value. Otherwise it reads a full BF16 from $L$. Each pair $(w_{2l},w_{2l+1})$ is packed into one 32-bit register [2603.17435].

## 4. ZipGEMM in SparseZipper: zipper encoding and CPU SpGEMM

In the SparseZipper literature, ZipGEMM operates on sparse vectors viewed as streams of `(key,value)` tuples. The index stream is a sequence of 32-bit column indices, and the value stream is a matching sequence of 32-bit values. Up to $R$ concurrent streams are mapped into the $R$ rows of two tile registers, $\mathrm{TR}_{\mathit{key}}$ and $\mathrm{TR}_{\mathit{val}}$, while vector registers hold chunk lengths and base addresses. Special counter vector registers $\mathrm{IC}_0,\mathrm{IC}_1$ and $\mathrm{OC}_0,\mathrm{OC}_1$ track how many entries have been consumed or produced [2502.11353].

The representation is organized around fixed-size chunks of $R$ tuples. With key cost $B_k$ bits and value cost $B_v$ bits, the raw storage for $W$ nonzeros is
$$
S_{\mathrm{raw}}=W(B_k+B_v).
$$
The zipped size is given as
$$
S_{\mathrm{zip}}=\Bigl\lceil \frac{W}{R}\Bigr\rceil\,(2R(B_k+B_v)+\delta),
$$
where $\delta$ is small per-chunk metadata overhead for routing, combine, or skip decisions. The effective compression ratio is
$$
\mathrm{CR}=\frac{W(B_k+B_v)}{\lceil W/R\rceil\,(2R(B_k+B_v)+\delta)}.
$$
In practice, the summary states that $R=16$ and $\delta\ll 2R(B_k+B_v)$.

The ISA additions are central to the method. The new instructions include indexed load `mlxe.t`, indexed store `msxe.t`, stream-sort keys `mssortk.tt`, stream-sort values `mssortv.tt`, stream-merge keys `mszipk.tt`, stream-merge values `mszipv.tt`, and counter-move `mmv.vi` and `mmv.vo`. The baseline target is a RISC-V-style matrix ISA with a 16×16 systolic array datapath. Per-PE changes include “two 1-bit tags per datum,” a comparator with 2:1 multiplexers, four popcount-based counters, and two loopback multiplexers. The textual datapath description states that equal keys are summed into a single south output, while routing state is recorded for the next compress pass.

The end-to-end algorithm follows Gustavson’s row-wise product. For each row of $A$ in CSR, the algorithm expands streams from products $A[i,*]\times B$ rows, splits the stream into partitions of at most $R$ tuples, sorts and compresses each partition with the zipper instructions, then tree-merges sorted partitions until a final CSR row of $C$ is produced. Dynamic zero bypassing is handled without explicit branching: invalid or skipped slots are carried as sentinel keys and excluded during the compress pass. Pointer control is maintained by reading the $\mathrm{IC}$ and $\mathrm{OC}$ counters [2502.11353].

## 5. Arithmetic intensity, throughput, and empirical results

For the ZipServ GPU kernel, the paper gives arithmetic-intensity expressions for standard BF16 GEMM, a decoupled lossless pipeline, and the fused ZipGEMM pipeline:
$$
CI_{\mathrm{GEMM}}=\frac{MNK}{MK+KN+MN},
$$
$$
CI_{\mathrm{Decoupled}}\approx \frac{MNK}{2.66MK+KN+MN},
$$
$$
CI_{\mathrm{ZipServ}}\approx \frac{MNK}{0.66MK+KN+MN}.
$$
The text states that the fused pipeline can exceed even the uncompressed baseline arithmetic intensity in memory-bound regimes, and that the Roofline in §3.3 on an RTX4090 moves many practical GEMMs out of the memory-bound region toward a higher performance ceiling [2603.17435].

Kernel-level performance is reported on RTX4090 and L40S, averaged over all linear layers and batch sizes $\{8,16,32\}$. On RTX4090, “ZipGEMM vs. cuBLAS_TensorCore” is reported as “1.31× avg, 1.71× peak,” while DietGPU (rANS), nvCOMP (rANS), and DFloat11 (Huffman) are listed as `0.17×`, `0.19×`, and `0.28×`. On L40S, ZipGEMM versus cuBLAS_TensorCore is “1.36× avg, 2.21× peak,” with DietGPU, nvCOMP, and DFloat11 at `0.20×`, `0.23×`, and `0.34×`. On an RTX5090, ZipGEMM is “1.34× faster for LLaMA-3.1-8B and 1.87× for Mistral-24B vs. native cuBLAS.” The paper also states that “a single RTX4090+ZipGEMM outperforms a DGX-A100’s cuBLAS by 9.3 % (LLaMA8B) and is within 2.7 % (Mistral)” [2603.17435].

For end-to-end inference, ZipServ reports that across all tested configurations it “reduced end-to-end latency by ≈17.6 % vs. vLLM and boosted throughput by 1.22×,” while “weight memory footprints fell by ≈29 % (e.g. 14.96 GB→10.83 GB for LLaMA8B)” [2603.17435].

For the SparseZipper CPU design, the sorting or merging latency per micro-operation is
$$
L_{\mathrm{sz}}=2N+2
$$
cycles, and sorting or merging a stream of $W$ tuples in $P=\lceil W/R\rceil$ chunks takes
$$
T_{\mathrm{zip}}=P\times L_{\mathrm{sz}}+T_{\mathrm{mem}}+T_{\mathrm{ovh}},
$$
where $T_{\mathrm{mem}}$ is approximately $4P$ vector micro-operations for indexed loads and stores. The throughput is written as
$$
\mathrm{Throughput}_{\mathrm{ZipGEMM}}=\frac{\sum_i \mathrm{nnz}(C_{i,*})}{T_{\mathrm{total}}},
$$
and the speedup over a scalar hash baseline is approximated by
$$
\mathrm{Speedup}_{\mathrm{hash}}\approx \frac{L_{\mathrm{hash}}}{L_{\mathrm{sz}}/R},
$$
with $L_{\mathrm{hash}}\approx 30$ cycles per nonzero [2502.11353].

Representative SuiteSparse results are reported for `p2p`, `wiki`, `scircuit`, and `p3d`. The summary gives average speedups of “5.98× over scalar-hash and 2.61× over ESC,” with highest gains when average per-row work is low and irregular, and with “31% fewer L1D accesses vs ESC” [2502.11353].

## 6. Hardware costs, limitations, and interpretive comparison

The ZipServ paper states several trade-offs. Small layers, such as `O_proj` in 8 B, can be “marginally slower (down to 0.79×)” because of fine-grain tiling overheads. In the compute-bound prefill stage with large $N$, on-the-fly fusion can slightly outweigh memory savings, and the system therefore switches to “a lightweight decoupled decompress+cuBLAS path” with “only 2–4 % overhead.” On “HBM-rich datacenter GPUs (A100, H800),” the very high memory bandwidth reduces ZipGEMM’s relative benefit, although the text states that it “remains the fastest decompression and fused GEMM implementation” [2603.17435].

The SparseZipper paper provides area and energy characterizations. Post-synthesis of the extended 16×16, 512-bit-wide systolic array in 12 nm yields a total area increase from `140.16 k μm²` to `158.00 k μm²`, corresponding to “12.7% area overhead on the 16×16 array.” Amortized over a full CPU plus caches, this is described as only “a few percent extra area.” Dynamic energy per nonzero is reduced by “∼30% versus ESC and ∼50% versus scalar,” and the extra comparator-plus-mux logic adds “<5 mW at full toggle,” which is characterized as negligible against the “500 mW array baseline.” Limitations include under-utilization for “very low-density matrices (work ≪ R)” and idle lanes under “highly skewed sparsity (work variation ≫1),” with row reordering as a partial mitigation [2502.11353].

A plausible implication is that the two ZipGEMM variants address different forms of wasted data movement. ZipServ targets redundant DRAM traffic arising from materialized decompression of dense BF16 weights, whereas SparseZipper targets wasted compute and memory activity arising from zeros and irregular sparse access. The common theme is close coupling between representation, kernel structure, and hardware datapath, but the mechanisms are not interchangeable: one depends on triple-bitmaps, Tensor Core fragments, and BF16 exponent skew; the other depends on keyed sparse streams, zipper instructions, and systolic-array sort/merge behavior.

## 7. Position within the broader GEMM landscape

Within the GPU LLM-serving setting, ZipServ positions ZipGEMM against cuBLAS, vLLM, DietGPU, nvCOMP, and DFloat11. Its distinguishing claim is that it is “the first lossless compression system that provides both storage savings and substantial acceleration for LLM inference on GPUs” [2603.17435]. The system is explicitly lossless and bit-exact, and its compression format is designed to avoid the SIMT divergence and redundant memory traffic associated with traditional entropy-coded pipelines.

Within the CPU sparse-linear-algebra setting, SparseZipper positions ZipGEMM against scalar hash-based SpGEMM and a vectorized ESC baseline. Its distinguishing claim is not bit-exact dense inference, but the ability to accelerate sparse-sparse GEMM “operating on highly sparse matrices with unstructured sparsity structures” while requiring only minimal modifications to existing matrix extensions and systolic-array micro-architectures [2502.11353].

A common misconception would be to treat “ZipGEMM” as a single canonical method. The published record instead supports a narrower and more precise view: the name denotes at least two separate research constructs, each embedded in its own hardware-software stack. For readers following the term in arXiv-era systems work, disambiguation by paper and workload is therefore essential.

Source: https://www.emergentmind.com/topics/zipgemm