---
title: Tensor-Core-Aware Triple Bitmap Encoding
url: https://www.emergentmind.com/topics/tensor-core-aware-triple-bitmap-encoding-tca-tbe
type: topic
---

# Tensor-Core-Aware Triple Bitmap Encoding

Searching arXiv for the cited papers to ground the article in current records.
{"queries":[{"query":"arXiv 2603.17435 ZipServ Fast and Memory-Efficient LLM Inference with Hardware-Aware Lossless Compression"},{"query":"arXiv 2105.09564 Dual-side Sparse Tensor Core"}]}
Tensor-Core-Aware Triple Bitmap Encoding (TCA-TBE) is a fixed-length, tile-structured, lossless BF16 weight compression format introduced in ZipServ for GPU-based LLM inference. It is designed to make compression operationally useful at inference time rather than merely reducing storage footprint: each \(8\times 8\) tile is encoded with three 64-bit bitmap planes plus compact and fallback value buffers, and the resulting representation is decoded in a branch-light, warp-synchronous manner directly into Tensor Core operand registers. In ZipServ, TCA-TBE underpins a fused decompression-GEMM kernel, ZipGEMM, implementing a “load-compressed, compute-decompressed” execution model that targets the memory-bandwidth bottlenecks of autoregressive decode. The format is motivated by empirical regularities in BF16 exponent distributions and is conceptually related to earlier bitmap-driven Tensor Core designs for sparse computation, most notably “Dual-side Sparse Tensor Core,” although that earlier work is a bitmap-based precursor or alternative rather than a literal triple-bitmap encoding [2603.17435] [2105.09564].

## 1. Definition and motivating observations

TCA-TBE stands for **Tensor-Core-Aware Triple Bitmap Encoding**. In ZipServ, it is the central representation for lossless compression of BF16 model weights, and its objective is to preserve bit-exact values while enabling efficient GPU execution during inference [2603.17435].

The immediate motivation is twofold. At the kernel level, conventional entropy coders such as Huffman and ANS produce variable-length bitstreams that are poorly matched to SIMT execution because decoding becomes serialized, data-dependent, and divergent across threads. At the system level, decoupled decompression pipelines materialize full weights back into global memory before GEMM, which destroys much of the bandwidth benefit of compression. ZipServ addresses the first issue with TCA-TBE and the second with ZipGEMM’s fused execution model [2603.17435].

The format exploits a specific empirical property of BF16 LLM weights. A BF16 value has 1 sign bit, 8 exponent bits, and 7 mantissa bits, and the exponent field is reported to be highly redundant across major LLMs: the top-3 most frequent exponents cover more than **67%** of all weights, the top-7 most frequent exponents cover more than **95%**, specifically **96.4%** in Llama-3 and **97.4%** in Mistral-24B, and the exponent entropy is only **2.57–2.74 bits** [2603.17435]. The paper further reports that in **99.6%** of the 3,875 weight matrices examined across Gemma-3, Mistral, Qwen2.5, and LLaMA-3.1, the top-7 exponents form a numerically contiguous sequence,
\[
\{e^\star, \dots, e^\star + 6\}.
\]
This permits a base-plus-offset exponent representation rather than general entropy-coded symbols [2603.17435].

The paper interprets these observations as implying a theoretical lossless compression ratio of about **\(1.51\times\)** for BF16 values, written approximately as
\[
16 / 10.6,
\]
because a BF16 value uses 16 bits in total while the exponent carries about 2.6 bits of entropy [2603.17435]. This suggests that the compression opportunity is structurally concentrated in the exponent field, and that a fixed-length code can approach the information-theoretic limit closely enough to be preferable on GPUs.

## 2. Encoding format and data layout

The defining mechanism of TCA-TBE is the “triple bitmap” organization. Each weight element is assigned a fixed **3-bit codeword**, but instead of densely packing these 3-bit codes into a conventional bitstream, the format stores them as **three separate 64-bit bitmaps**, one bit-plane per bitmap, over an \(8\times 8\) tile called a **FragTile** [2603.17435].

For BF16 weights, the high-frequency path encodes values whose exponents lie in a selected top-7 contiguous window. For such elements, the format stores only the **sign bit** and **7-bit mantissa** in a compact 8-bit payload and reconstructs the exponent from a 3-bit code and a per-matrix **BaseExp**. Weights whose exponent lies outside that range are stored in a fallback buffer in full BF16 precision. Each element is therefore either on a **high-frequency path** or a **fallback path** [2603.17435].

The code space is fixed and exhaustive. The paper defines:
- `001`–`111` for the seven frequent exponent classes,
- `000` for fallback or outlier values.

If the selected exponent window is
\[
E_{top} = \{e_{\min}, e_{\min}+1, \dots, e_{\min}+6\},
\]
then the compressor stores
\[
e_{base} = \min(E_{top}) - 1,
\]
and for a covered exponent \(e\), the code is
\[
c = e - e_{base}, \quad c \in [1,7].
\]
Exponent reconstruction is therefore
\[
e = e_{base} + c.
\]
The paper describes this as an “implicit lookup” mechanism [2603.17435].

At the tile level, each FragTile contains 64 weights and uses five buffers:
1. **Bitmap 1**: a 64-bit bit-plane for the least-significant code bit.
2. **Bitmap 2**: a 64-bit bit-plane for the middle code bit.
3. **Bitmap 3**: a 64-bit bit-plane for the most-significant code bit.
4. **PackedSignMantissa buffer**: compact 8-bit payloads for high-frequency weights.
5. **FullValue buffer**: full BF16 values for fallback weights [2603.17435].

At the matrix level, the paper states that TCA-TBE organizes these buffers into four contiguous global arrays, nested according to the tiling hierarchy, and adds an Offset array recording the starting offset of each GroupTile within the PackedSignMantissa and FullValue arrays. The paper contains a minor inconsistency between “five buffers per tile” and “four contiguous global arrays” at matrix level; the directly supported point is that each tile has three bitmap planes plus two value buffers, and matrix-level metadata includes offsets into the compact and fallback streams [2603.17435].

A key property of the representation is that the three bitmaps can be OR-reduced to produce a spatial indicator
\[
\mathcal M = \mathcal B_1 \lor \mathcal B_2 \lor \mathcal B_3,
\]
where \(\mathcal M[p]=1\) iff position \(p\) is a covered high-frequency element and \(\mathcal M[p]=0\) iff it is a fallback element. This makes state classification constant-time and register-resident during decoding [2603.17435].

## 3. Compression and decoding mechanics

ZipServ formalizes TCA-TBE with an offline compressor. Compression begins with a per-matrix exponent histogram:
1. \(Hist \leftarrow ComputeExponentHistogram(\mathcal{W})\)
2. \(E_{top} \leftarrow SelectTop7ConsecutiveExponents(Hist)\)
3. \(e_{base} \leftarrow \min(E_{top}) - 1\)

The matrix is then partitioned into \(8\times 8\) tiles, and for each element \(w\) in a tile, the compressor extracts exponent \(e\). If \(e \in E_{top}\), it computes \(c=e-e_{base}\), writes the three code bits into the corresponding local bitmaps, and appends \(Pack(w.sign,w.mantissa)\) to the high-frequency buffer \(\mathcal H\). Otherwise it appends the full BF16 value \(w\) to the fallback buffer \(\mathcal L\) [2603.17435].

The decoding procedure is equally structured. For Tensor Core operand layout, each thread in a warp reconstructs **two BF16 values** from an \(8\times 8\) tile. Lane \(i\) owns positions
\[
2i \quad \text{and} \quad 2i+1
\]
within the tile, so 32 threads reconstruct 64 values, packing two BF16 values into one `.bf16x2` register [2603.17435].

Decoding proceeds in three stages. First, the thread computes
\[
\mathcal M \leftarrow \mathcal B_1 \lor \mathcal B_2 \lor \mathcal B_3.
\]
Second, it computes addresses into \(\mathcal H\) or \(\mathcal L\) by population count. For position \(p\), define
\[
mask = (1 \ll p) - 1.
\]
Then the number of prior high-frequency elements is
\[
idx_{\mathcal H} = Popc(\mathcal M \ \& \ mask).
\]
If the current element is fallback, then the number of earlier fallback elements is
\[
idx_{\mathcal L} = p - idx_{\mathcal H}.
\]
The paper explicitly cites GPU-native instructions such as `__popc()` and `__shfl_sync()` for these warp-local operations [2603.17435].

Third, exponent reconstruction depends on the state bit. If \(\mathcal M[p]=1\), the decoder loads sign and mantissa from \(\mathcal H\), reconstructs the code
\[
c \leftarrow (\mathcal B_3[p] \ll 2) \lor (\mathcal B_2[p] \ll 1) \lor \mathcal B_1[p],
\]
then computes
\[
e \leftarrow e_{base} + c,
\]
and assembles the original BF16 value with \(MakeBF16(val.sign, e, val.mantissa)\). If \(\mathcal M[p]=0\), it loads the full BF16 value from \(\mathcal L[\text{start}_{\mathcal L}+idx_{\mathcal L}]\) [2603.17435].

The paper repeatedly characterizes this decoding as **constant-time**. In context, this means that the per-element decode consists of a fixed small number of ORs, shifts, masks, population counts, direct loads, arithmetic reconstruction, and register packing, without variable-length parsing loops, hierarchical decode trees, or serialized bit-pointer advancement [2603.17435]. A plausible implication is that TCA-TBE’s computational regularity is at least as important as its compression ratio in determining end-to-end utility on GPUs.

## 4. Tensor-Core awareness and fused ZipGEMM execution

The “Tensor-Core-Aware” qualifier refers to an explicit co-design with NVIDIA Tensor Core fragment shapes, register layouts, and warp mapping. The target instruction is `mma.sync.m16n8k16`, and the paper describes a typical BF16 Tensor Core operation in terms of fragments \(A_{\text{frag}}\), \(B_{\text{frag}}\), and FP32 accumulator \(C_{\text{frag}}\), with the intended operation
\[
D_{\text{frag}} = A_{\text{frag}} \times B_{\text{frag}} + C_{\text{frag}}.
\]
Although some source typesetting is damaged, the intended mapping is standard [2603.17435].

TCA-TBE uses a three-level tiling hierarchy:
1. **FragTile (FT):** \(8\times 8\)
2. **TensorCoreTile (TT):** \(16\times 16\), formed as a \(2\times 2\) grid of FragTiles
3. **BlockTile (BT):** \(64\times 64\), processed cooperatively by a thread block [2603.17435]

This hierarchy aligns with execution in two ways. The \(16\times 16\) level matches the \(m=16, k=16\) operand dimensions for Tensor Core MMA, and the \(8\times 8\) level matches a warp-local reconstruction pattern of 64 BF16 values with 32 lanes and 2 values per lane [2603.17435]. The paper also states that the FragTiles within a TensorCoreTile are stored in **column-major order**, mirroring the register layout expected by Tensor Cores, such as `Ra0–Ra3`, so decompression produces values already arranged for `mma.sync` consumption [2603.17435].

This layout is paired with ZipGEMM, the fused decompression-GEMM kernel. Its runtime model is:
- load compressed weights from DRAM,
- stage them in shared memory,
- decode them on the fly in registers,
- feed the registers directly into Tensor Core MMA [2603.17435].

The paper calls this “load-compressed, compute-decompressed.” The workflow is:
1. compressed weight tile and activation tile are loaded from global memory to shared memory using asynchronous vectorized instructions such as `LDGSTS.128`, with compressed value arrays padded offline to ensure **128-bit alignment**;
2. each warp decodes the weight tile from shared memory directly into registers;
3. activations are moved from shared memory to registers via `LDSM.M88`;
4. MMA executes using the decoded weight registers and loaded activation registers [2603.17435].

This placement of decompression between shared memory and the register file is central to the system-level design. The paper argues that materializing decompressed weights in global memory would negate a large part of the bandwidth benefit by reintroducing decompressed write and read traffic. ZipServ’s fused path therefore avoids intermediate decompressed global-memory buffers and extra latency between decompression and GEMM [2603.17435].

The paper formulates this in terms of compute intensity. For baseline GEMM,
\[
CI_{GEMM} = \frac{MNK}{MK + KN + MN}.
\]
For a decoupled compression pipeline with average compression ratio \(CR = 1.51\), the paper gives a corrupted but interpretable expression implying that compute intensity is sharply reduced, and states that for \(M=K=4096\), the decoupled pipeline loses about **62%** of compute intensity relative to standard GEMM across batch sizes 8–64 [2603.17435]. For ZipServ’s fused design, the source equation is likewise partially garbled, but the reported conclusion is that fused execution can exceed the effective compute intensity of standard uncompressed GEMM by about **50%** in memory-bound regimes [2603.17435].

## 5. Quantitative behavior, benefits, and operating regime

The paper reports system-level weight footprint reductions of approximately 29–30%:
- **LLaMA3.1-8B:** 14.96 GB \(\rightarrow\) 10.83 GB (72.4%)
- **Mistral-24B:** 43.92 GB \(\rightarrow\) 31.30 GB (71.3%)
- **LLaMA3.1-70B:** 131.56 GB \(\rightarrow\) 93.52 GB (71.1%) [2603.17435]

Kernel-level fused GEMM performance against cuBLAS Tensor Core BF16 GEMM is reported as:
- **RTX4090:** average **\(1.31\times\)**, peak **\(1.71\times\)**
- **L40S:** average **\(1.36\times\)**, peak **\(2.21\times\)** [2603.17435]

Against decoupled compressed baselines, ZipGEMM is reported to achieve:
- DietGPU: **\(0.17\times / 0.20\times\)** on RTX4090 / L40S
- nvCOMP: **\(0.19\times / 0.23\times\)**
- DFloat11: **\(0.28\times / 0.34\times\)** [2603.17435]

The paper further reports layer-wise L40S averages on LLaMA3.1:
- GateUp\_proj: **\(1.39\times\)**
- Down\_proj: **\(1.64\times\)**

Transformer-block speedups are given as:
- LLaMA3.1-8B block: **\(1.35\times\)**
- LLaMA3.1-405B block: **\(1.48\times\)** [2603.17435]

A standalone decompression kernel using the same TCA-TBE format, ZipServ-Decomp, achieves average speedups over:
- DietGPU: **\(2.14\times\)**
- nvCOMP: **\(1.83\times\)**
- DFloat11: **\(1.10\times\)** [2603.17435]

The paper attributes this to the format’s fixed-length structure, warp-aligned design, and elimination of control divergence. Microarchitecturally, on RTX4090 for \(M=28672, K=4096, N=32\), the paper reports:
- DRAM reads drop by **29.3%**
- ALU utilization rises to **66.0%**
- Tensor Core utilization remains **71.6%** of cuBLAS baseline [2603.17435]

This is presented as the intended tradeoff: spend predictable integer-ALU work to save bandwidth. The paper also reports that shared-memory bank conflicts are nearly eliminated (\(\sim 4.7\text{K}\)) versus millions in baselines like DietGPU, crediting TCA-TBE’s data layout [2603.17435].

End-to-end serving results are likewise stage-specific. Compared to vLLM, Transformers, and DFloat11, ZipServ improves throughput on average by:
- **\(1.22\times\)** over vLLM
- **\(3.18\times\)** over Transformers
- **\(8.52\times\)** over DFloat11

Latency reductions average:
- **17.60%** vs vLLM
- **60.79%** vs Transformers
- **82.13%** vs DFloat11 [2603.17435]

For LLaMA3.1-8B at batch 32 and output 2048 tokens, the paper reports **1105 tokens/s** and **\(1.66\times\)** speedup over vLLM [2603.17435]. These gains are situated primarily in the decode stage. The paper explicitly states that ZipGEMM is used only in the **decode stage**, whereas the **prefill stage** uses decoupled decompression plus GEMM because large \(N\) makes GEMM compute-bound and on-the-fly decompression is less beneficial; decompression overhead in that mode is reported as only about **\(\sim 4\% / 2\%\)** at \(N=8192/16384\) [2603.17435].

The operating regime is therefore clear. TCA-TBE is most effective when weights are BF16, exponent distributions are strongly skewed, top-7 exponents are contiguous, and inference is memory-bound, especially in decode on consumer or inference GPUs [2603.17435]. It is weaker when exponent distributions are less concentrated, layers are small, or available bandwidth is sufficiently high that additional integer work is harder to hide. The paper notes a slowdown on some small \(O\_proj\) layers, specifically LLaMA3.1-8B \(O\_proj\) on L40S at **\(0.79\times\)** relative to cuBLAS [2603.17435].

## 6. Relation to prior bitmap-driven Tensor Core research

TCA-TBE belongs to a broader class of hardware-aware metadata encodings for Tensor Core execution, but its role is distinct. The most relevant precursor in the provided literature is “Dual-side Sparse Tensor Core,” which proposes a bitmap-driven sparse Tensor Core architecture to exploit both weight sparsity and activation sparsity in DNN inference [2105.09564].

That earlier paper is not a literal triple-bitmap encoding scheme. Its core representation for sparse matrices is a two-tuple,
\[
\mathbb{A} \rightarrow (\mathbb{A}_b, \mathbb{A}_v), \qquad \mathbb{B} \rightarrow (\mathbb{B}_b, \mathbb{B}_v),
\]
where the bitmaps encode zero versus nonzero positions and the values arrays store condensed nonzeros [2105.09564]. It also defines a bitmap outer product,
\[
\mathbb{D}^{(k)}_b = \mathbb{A}_b(:,k) \otimes \mathbb{B}_b(k,:),
\]
with
\[
\mathbb{D}^{(k)}_b(i,j) = \mathbb{A}_b(i,k) \land \mathbb{B}_b(k,j),
\]
to derive the sparsity structure of outer-product partial results [2105.09564].

At the device level, “Dual-side Sparse Tensor Core” uses a two-level bitmap hierarchy comprising an **element-bitmap** for fine-grained sparsity within a tile and a **warp-bitmap** for coarse tile occupancy, enabling entire warp tiles to be skipped when empty [2105.09564]. For sparse convolution, it introduces a three-field representation,
\[
(\text{bitmap}, \text{row offset}, \text{value}),
\]
which is explicitly not three bitmaps, but rather one bitmap plus offsets and values [2105.09564].

The architectural affinity to TCA-TBE lies elsewhere: the earlier work is also Tensor-Core-aware, also makes metadata active in execution rather than passive in storage, also uses bitmaps to guide predication and skipping, and also aligns encoding with Tensor Core tile shapes and instruction structure. It replaces inner-product Tensor Core computation with an outer-product primitive because dual-side sparse inner products require unpredictable matching of nonzero positions, whereas sparse vectors can be condensed independently for outer-product execution and later reconstructed through bitmap-guided accumulation [2105.09564].

The distinction is therefore precise. TCA-TBE is a **lossless compression format** for BF16 weights with three explicit bitmap planes per \(8\times 8\) tile and direct register-local decoding for Tensor Core MMA [2603.17435]. “Dual-side Sparse Tensor Core” is a **bitmap-based sparse execution architecture** with hierarchical metadata, bitmap outer products, OHMMA/BOHMMA-style ISA support, and a minimal-hardware-change argument for exploiting dual-side sparsity [2105.09564]. The latter is best understood as a bitmap-driven sparse Tensor Core precursor or alternative rather than an instance of TCA-TBE.

A plausible implication is that TCA-TBE extends the broader Tensor-Core-aware metadata design philosophy from sparsity exploitation to lossless compression: instead of encoding nonzero structure, it encodes exponent classes in a fixed tile geometry that is directly consumable by Tensor Core execution. In that sense, both works exemplify a common principle—co-designing metadata representation, decode logic, and Tensor Core dataflow—while solving different problems [2603.17435] [2105.09564].

Source: https://www.emergentmind.com/topics/tensor-core-aware-triple-bitmap-encoding-tca-tbe