Tensor-Core-Aware Triple Bitmap Encoding

• TCA-TBE is a lossless BF16 weight compression format that encodes 8×8 tiles using three 64-bit bitmap planes to enable branch-light, register-resident decoding.
• It employs a fixed-length, tile-structured design paired with fused decompression-GEMM (ZipGEMM) to overcome memory bandwidth bottlenecks in GPU inference.
• Empirical results show kernel speedups of up to 1.71× and significant DRAM read reductions, highlighting its practical benefits for LLM inference.

Searching arXiv for the cited papers to ground the article in current records. {"queries":[{"query":"arXiv (Fan et al., 18 Mar 2026) ZipServ Fast and Memory-Efficient LLM Inference with Hardware-Aware Lossless Compression"},{"query":"arXiv (Wang et al., 2021) 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 (Fan et al., 18 Mar 2026, Wang et al., 2021).

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 (Fan et al., 18 Mar 2026).

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 (Fan et al., 18 Mar 2026).

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 (Fan et al., 18 Mar 2026). 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 (Fan et al., 18 Mar 2026).

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 (Fan et al., 18 Mar 2026). 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 (Fan et al., 18 Mar 2026).

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 (Fan et al., 18 Mar 2026).

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 (Fan et al., 18 Mar 2026).

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 (Fan et al., 18 Mar 2026).

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 (Fan et al., 18 Mar 2026).

A key property of the representation is that the three bitmaps can be OR-reduced to produce a spatial indicator

$\{e^\star, \dots, e^\star + 6\}.$0

where $\{e^\star, \dots, e^\star + 6\}.$1 iff position $\{e^\star, \dots, e^\star + 6\}.$2 is a covered high-frequency element and $\{e^\star, \dots, e^\star + 6\}.$3 iff it is a fallback element. This makes state classification constant-time and register-resident during decoding (Fan et al., 18 Mar 2026).

3. Compression and decoding mechanics

ZipServ formalizes TCA-TBE with an offline compressor. Compression begins with a per-matrix exponent histogram:

1. $\{e^\star, \dots, e^\star + 6\}.$4
2. $\{e^\star, \dots, e^\star + 6\}.$5
3. $\{e^\star, \dots, e^\star + 6\}.$6

The matrix is then partitioned into $\{e^\star, \dots, e^\star + 6\}.$7 tiles, and for each element $\{e^\star, \dots, e^\star + 6\}.$8 in a tile, the compressor extracts exponent $\{e^\star, \dots, e^\star + 6\}.$9. If $1.51\times$0, it computes $1.51\times$1, writes the three code bits into the corresponding local bitmaps, and appends $1.51\times$2 to the high-frequency buffer $1.51\times$3. Otherwise it appends the full BF16 value $1.51\times$4 to the fallback buffer $1.51\times$5 (Fan et al., 18 Mar 2026).

The decoding procedure is equally structured. For Tensor Core operand layout, each thread in a warp reconstructs two BF16 values from an $1.51\times$6 tile. Lane $1.51\times$7 owns positions

$1.51\times$8

within the tile, so 32 threads reconstruct 64 values, packing two BF16 values into one .bf16x2 register (Fan et al., 18 Mar 2026).

Decoding proceeds in three stages. First, the thread computes

$1.51\times$9

Second, it computes addresses into $16 / 10.6,$0 or $16 / 10.6,$1 by population count. For position $16 / 10.6,$2, define

$16 / 10.6,$3

Then the number of prior high-frequency elements is

$16 / 10.6,$4

If the current element is fallback, then the number of earlier fallback elements is

$16 / 10.6,$5

The paper explicitly cites GPU-native instructions such as __popc() and __shfl_sync() for these warp-local operations (Fan et al., 18 Mar 2026).

Third, exponent reconstruction depends on the state bit. If $16 / 10.6,$6, the decoder loads sign and mantissa from $16 / 10.6,$7, reconstructs the code

$16 / 10.6,$8

then computes

$16 / 10.6,$9

and assembles the original BF16 value with $8\times 8$0. If $8\times 8$1, it loads the full BF16 value from $8\times 8$2 (Fan et al., 18 Mar 2026).

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 (Fan et al., 18 Mar 2026). 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 $8\times 8$3, $8\times 8$4, and FP32 accumulator $8\times 8$5, with the intended operation

$8\times 8$6

Although some source typesetting is damaged, the intended mapping is standard (Fan et al., 18 Mar 2026).

TCA-TBE uses a three-level tiling hierarchy:

1. FragTile (FT): $8\times 8$7
2. TensorCoreTile (TT): $8\times 8$8, formed as a $8\times 8$9 grid of FragTiles
3. BlockTile (BT): $$E_{top} = \{e_{\min}, e_{\min}+1, \dots, e_{\min}+6\},$$0, processed cooperatively by a thread block (Fan et al., 18 Mar 2026)

This hierarchy aligns with execution in two ways. The $$E_{top} = \{e_{\min}, e_{\min}+1, \dots, e_{\min}+6\},$$1 level matches the $$E_{top} = \{e_{\min}, e_{\min}+1, \dots, e_{\min}+6\},$$2 operand dimensions for Tensor Core MMA, and the $$E_{top} = \{e_{\min}, e_{\min}+1, \dots, e_{\min}+6\},$$3 level matches a warp-local reconstruction pattern of 64 BF16 values with 32 lanes and 2 values per lane (Fan et al., 18 Mar 2026). 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 (Fan et al., 18 Mar 2026).

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 (Fan et al., 18 Mar 2026).

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 (Fan et al., 18 Mar 2026).

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 (Fan et al., 18 Mar 2026).

The paper formulates this in terms of compute intensity. For baseline GEMM,

$$E_{top} = \{e_{\min}, e_{\min}+1, \dots, e_{\min}+6\},$$4

For a decoupled compression pipeline with average compression ratio $$E_{top} = \{e_{\min}, e_{\min}+1, \dots, e_{\min}+6\},$$5, the paper gives a corrupted but interpretable expression implying that compute intensity is sharply reduced, and states that for $$E_{top} = \{e_{\min}, e_{\min}+1, \dots, e_{\min}+6\},$$6, the decoupled pipeline loses about 62% of compute intensity relative to standard GEMM across batch sizes 8–64 (Fan et al., 18 Mar 2026). 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 (Fan et al., 18 Mar 2026).

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 $$E_{top} = \{e_{\min}, e_{\min}+1, \dots, e_{\min}+6\},$$7 10.83 GB (72.4%)
• Mistral-24B: 43.92 GB $$E_{top} = \{e_{\min}, e_{\min}+1, \dots, e_{\min}+6\},$$8 31.30 GB (71.3%)
• LLaMA3.1-70B: 131.56 GB $$E_{top} = \{e_{\min}, e_{\min}+1, \dots, e_{\min}+6\},$$9 93.52 GB (71.1%) (Fan et al., 18 Mar 2026)

Kernel-level fused GEMM performance against cuBLAS Tensor Core BF16 GEMM is reported as:

• RTX4090: average $e_{base} = \min(E_{top}) - 1,$0, peak $e_{base} = \min(E_{top}) - 1,$1
• L40S: average $e_{base} = \min(E_{top}) - 1,$2, peak $e_{base} = \min(E_{top}) - 1,$3 (Fan et al., 18 Mar 2026)

Against decoupled compressed baselines, ZipGEMM is reported to achieve:

• DietGPU: $e_{base} = \min(E_{top}) - 1,$4 on RTX4090 / L40S
• nvCOMP: $e_{base} = \min(E_{top}) - 1,$5
• DFloat11: $e_{base} = \min(E_{top}) - 1,$6 (Fan et al., 18 Mar 2026)

The paper further reports layer-wise L40S averages on LLaMA3.1:

• GateUp_proj: $e_{base} = \min(E_{top}) - 1,$7
• Down_proj: $e_{base} = \min(E_{top}) - 1,$8

Transformer-block speedups are given as:

• LLaMA3.1-8B block: $e_{base} = \min(E_{top}) - 1,$9
• LLaMA3.1-405B block: $e$0 (Fan et al., 18 Mar 2026)

A standalone decompression kernel using the same TCA-TBE format, ZipServ-Decomp, achieves average speedups over:

• DietGPU: $e$1
• nvCOMP: $e$2
• DFloat11: $e$3 (Fan et al., 18 Mar 2026)

The paper attributes this to the format’s fixed-length structure, warp-aligned design, and elimination of control divergence. Microarchitecturally, on RTX4090 for $e$4, the paper reports:

• DRAM reads drop by 29.3%
• ALU utilization rises to 66.0%
• Tensor Core utilization remains 71.6% of cuBLAS baseline (Fan et al., 18 Mar 2026)

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 ($e$5) versus millions in baselines like DietGPU, crediting TCA-TBE’s data layout (Fan et al., 18 Mar 2026).

End-to-end serving results are likewise stage-specific. Compared to vLLM, Transformers, and DFloat11, ZipServ improves throughput on average by:

• $e$6 over vLLM
• $e$7 over Transformers
• $e$8 over DFloat11

Latency reductions average:

• 17.60% vs vLLM
• 60.79% vs Transformers
• 82.13% vs DFloat11 (Fan et al., 18 Mar 2026)

For LLaMA3.1-8B at batch 32 and output 2048 tokens, the paper reports 1105 tokens/s and $e$9 speedup over vLLM (Fan et al., 18 Mar 2026). 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 $c = e - e_{base}, \quad c \in [1,7].$0 makes GEMM compute-bound and on-the-fly decompression is less beneficial; decompression overhead in that mode is reported as only about $c = e - e_{base}, \quad c \in [1,7].$1 at $c = e - e_{base}, \quad c \in [1,7].$2 (Fan et al., 18 Mar 2026).

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 (Fan et al., 18 Mar 2026). 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 $c = e - e_{base}, \quad c \in [1,7].$3 layers, specifically LLaMA3.1-8B $c = e - e_{base}, \quad c \in [1,7].$4 on L40S at $c = e - e_{base}, \quad c \in [1,7].$5 relative to cuBLAS (Fan et al., 18 Mar 2026).

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 (Wang et al., 2021).

That earlier paper is not a literal triple-bitmap encoding scheme. Its core representation for sparse matrices is a two-tuple,

$c = e - e_{base}, \quad c \in [1,7].$6

where the bitmaps encode zero versus nonzero positions and the values arrays store condensed nonzeros (Wang et al., 2021). It also defines a bitmap outer product,

$c = e - e_{base}, \quad c \in [1,7].$7

with

$c = e - e_{base}, \quad c \in [1,7].$8

to derive the sparsity structure of outer-product partial results (Wang et al., 2021).

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 (Wang et al., 2021). For sparse convolution, it introduces a three-field representation,

$c = e - e_{base}, \quad c \in [1,7].$9

which is explicitly not three bitmaps, but rather one bitmap plus offsets and values (Wang et al., 2021).

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 (Wang et al., 2021).

The distinction is therefore precise. TCA-TBE is a lossless compression format for BF16 weights with three explicit bitmap planes per $e = e_{base} + c.$0 tile and direct register-local decoding for Tensor Core MMA (Fan et al., 18 Mar 2026). “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 (Wang et al., 2021). 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 (Fan et al., 18 Mar 2026, Wang et al., 2021).