↓Skip to main content

SpInfer: How bitmaps make sparse LLMs actually fast

·8 mins

Preface #

I’ve been going through SpInfer as part of my work at The Burns Lab and it’s one of those papers where the core idea is elegant but takes a while to fully click. The paper is called SpInfer: Leveraging Low-Level Sparsity for Efficient Large Language Model Inference on GPUs and it was published at EuroSys ‘25.

The short version: unstructured pruning makes LLMs sparse but nobody’s been able to turn that sparsity into real performance gains at low sparsity levels. SpInfer fixes that. Here’s how I understand it.

The problem with being sparse #

When you prune an LLM, you zero out the less important weights. This introduces sparsity into the weight matrices. In theory, fewer non-zero values means less memory and less compute. In practice, it’s been surprisingly hard to make this work on GPUs, especially when sparsity is low (30-50%), which is exactly where LLMs sit without unacceptable accuracy loss.

The reason is indexing overhead. Most sparse formats like CSR (Compressed Sparse Row) need to store a 32-bit integer for every non-zero element to track its position. At 70%+ sparsity, there are few enough non-zeros that this overhead is worth it. But at 30-50% sparsity, the majority of your elements are still non-zero, so you end up storing:

  • The non-zero values themselves (2 bytes each, FP16)
  • A 32-bit index for every single one of them

That’s 6 bytes per non-zero vs 2 bytes in the dense format. At 30% sparsity with 70% of elements being non-zero, your “compressed” format is actually larger than the original dense matrix.

The paper formalizes this with a Compression Ratio:

$$CR = \frac{2B \times M \times K}{Stor_{Format}}$$

If CR < 1, your sparse format costs more than the dense matrix. CSR and Tiled-CSL (used by Flash-LLM) both have CR < 1 below roughly 50% sparsity. SparTA barely gets above 1. This is why SpMM kernels can’t beat cuBLAS at low sparsity — you’re not actually saving memory, so there’s no benefit to offset the computational overhead of dealing with a sparse format.

There’s also a second problem: even if you fix the memory side, GPU SpMM kernels are slow at low sparsity because fetching indices from global memory causes non-coalesced memory access. When 32 threads in a warp each need an index from a different location in memory, you get 32 separate memory transactions instead of 1. This kills bandwidth utilization.

How TCA-BME addresses the indexing overhead #

SpInfer introduces the Tensor-Core-Aware Bitmap Encoding (TCA-BME) format. Instead of storing a 32-bit index per non-zero, it stores 1 bit per element — regardless of whether that element is zero or non-zero. 64 elements fit into a single 64-bit integer. That’s it.

This works through a three-level tiling design:

  • BitmapTile (8×8): The smallest unit. 64 elements → one uint64_t. Dimensions chosen to match the minimum computational unit of Tensor Cores.
  • TCTile (16×16): A 2×2 arrangement of BitmapTiles. Matches the shape of the mma.m16n8k16 PTX instruction used for Tensor Core computation.
  • GroupTile: Multiple TCTiles, corresponding to the thread block level.

The storage for TCA-BME is:

$$Stor_{TCA-BME} = 4B \times (NGT+1) + 8B \times NBT + 2B \times NNZ$$

where NGT is the number of GroupTiles, NBT the number of BitmapTiles, and NNZ the number of non-zeros. At 30% sparsity this gives CR > 1, which no other format achieves.

Alongside the Bitmap array, there’s a Values array that holds only the non-zero elements, packed tightly in the order they appear in the matrix (GroupTile → TCTile → BitmapTile order). The dense matrix is discarded after compression. The Values array IS the storage — there’s no pointer back to anything.

The Values array and why it’s confusing at first #

This is the part that took me the longest to fully get.

Say your bitmap looks like this (using 8 bits for simplicity):

Bit position:  7  6  5  4  3  2  1  0
Bitmap:        1  0  0  1  0  1  1  0

Values: [ 0.73, -1.24, 0.05, 0.91 ]
           [0]    [1]   [2]   [3]

Positions 1, 2, 4, 7 are non-zero. The Values array stores them tightly packed, in order of appearance. No zeros, no gaps. The index of your value in Values is not your position in the matrix — it’s how many non-zeros appeared before you in the matrix.

If you’re at position 4, two non-zeros appeared before you (at positions 1 and 2), so your value lives at Values[2]. Not Values[4] — there’s no slot for zeros in the Values array.

This also means the dense matrix is gone for good. When a thread needs the value at position 4, it doesn’t go back to the original matrix. It uses the bitmap to figure out which slot in Values is its value, then fetches from there.

Shared Memory Bitmap Decoding (SMBD) #

The Tensor Core’s mma instruction requires a dense matrix fragment in a specific register layout. You can’t feed it sparse data. So before TC computation, SMBD reconstructs a dense 16×16 fragment in registers using the bitmap.

Each thread in a warp has a lane ID and is responsible for two elements (two FP16 values packed in a 32-bit register). Thread with lane ID l owns positions 2l and 2l+1.

To find its value in the Values array, the thread runs MaskedPopCount:

offset = lane_id × 2
mask   = (1 << offset) - 1    // all bits below my position
count  = PopCount(bitmap & mask)

(1 << offset) - 1 in binary gives all 1s below bit offset. AND-ing with the bitmap isolates only the non-zeros that precede the thread’s position. PopCount counts them. That count is the thread’s index into Values.

For a zero position (bitmap bit = 0), the thread places 0.0 in its register. For a non-zero, it loads Values[count].

After all 32 threads do this, the register file holds a complete dense 16×16 fragment — zeros and non-zeros alike — ready for the mma instruction.

Why is this fast? The bitmap is loaded into shared memory via a single coalesced transaction (one uint64_t, shared across the warp). The MaskedPopCount is a few register-level operations: one shift, one subtract, one AND, and one hardware __popcll instruction — single clock cycle. No global memory access after the initial load. No index arrays. Compare this to CSR where every thread needs to fetch its index from global memory (hundreds of cycles, non-coalesced).

SMBD also runs on CUDA cores, which are physically separate from Tensor Cores. This is key for the async pipeline.

The async pipeline #

SpInfer uses double buffering: two shared memory buffers for GTiles (weight tiles) and XTiles (input tiles). While the TC computes the current tile, the next tile is being fetched asynchronously.

The loading uses LDGSTS.128 — an Ampere-era instruction that moves data directly from global memory to shared memory, bypassing the L1 cache and register file entirely. Flash-LLM loads through the register file (global → registers → shared memory), consuming register space and adding latency. LDGSTS skips that roundtrip.

Two separate cp.async groups manage GTile and XTile loading independently, so they can proceed at their own rate. Once GTile loading completes, SMBD begins immediately on CUDA cores — running concurrently with XTile loading on the memory subsystem. After TC computation fires, SMBD for the next tile begins immediately. Three hardware units, all busy simultaneously:

Cycle N:   TC computes tile K        (Tensor Cores)
           SMBD decodes tile K+1     (CUDA cores)
           Load GTile/XTile K+2      (Memory subsystem)

The ablation study in the paper confirms this isn’t just theoretical — removing SMBD causes a 68.78% drop in bandwidth utilization and 78.41% drop in TC utilization. Removing the async pipeline drops TC utilization by 2%.

Results #

At the kernel level on an RTX 4090, SpInfer is the only method that consistently beats cuBLAS at 30% sparsity (94.44% of matrices). At 50% it wins on 96.30% of test cases. At 70%, 100%.

For comparison, Flash-LLM and SparTA only start reliably beating cuBLAS around 70%.

End-to-end on OPT-13B at 60% sparsity, SpInfer achieves 47.5% memory reduction (27.4 GB → 14.4 GB) and 1.35× speedup over Flash-LLM.

Limitations and what’s next #

SpInfer is fundamentally a memory-bound optimization. During prefill (when N = batch_size × seq_length is large), the CI formula shows why this becomes a problem:

$$CI_{SpMM} = \frac{M \times N}{\frac{M}{CR} + N}$$

As N grows large, the N term dominates and CI_SpMM converges to CI_GEMM — SpInfer’s CR advantage disappears. The operation moves into the compute-bound region where SMBD’s overhead costs without delivering memory savings. SpInfer can be up to 11.8% slower than cuBLAS during prefill at large sequence lengths.

Fixing this properly requires Sparse Tensor Cores — hardware that can natively skip zero multiplications without needing SMBD decoding at all.

SpInfer also doesn’t handle dynamic activation sparsity (sparsity in the input matrix X that varies per token at inference time). This is where something like a GPU-adapted Masked Matrix Multiplication could complement SpInfer — SpInfer owns static weight sparsity in W, a GPU-MMM handles dynamic activation sparsity in X. SpSpMM on GPU remains an open problem.

↑