SpInfer

02-04-2026

Background and Motivation

Weight pruning emerged as a promising method to reduce both memory consumption and computational burden in large language models. The idea is straightforward: introduce sparsity into the model by removing weights. Pruning methods come in different flavors (structured, semi-structured, or unstructured), each with different trade-offs between flexibility and hardware efficiency.

The challenge that remains unsolved is exploiting unstructured sparsity during LLM inference. State-of-the-art pruning techniques typically achieve around 50% sparsity before you start seeing adverse effects on model accuracy, but actually taking advantage of that sparsity on GPUs is where things get complicated.

SpInfer is a framework designed to accelerate LLM inference by leveraging this unstructured sparsity on GPUs. It introduces two key techniques: TCA-BME (Tensor-Core-Aware Bitmap Encoding), which encodes sparsity indexing information using bitmaps, and SMBD (Shared Memory Bitmap Decoding), which decodes these sparse matrices directly within shared memory while using asynchronous pipelining.

Before diving into the techniques themselves, it's worth understanding what they're optimizing. The paper identifies that weight storage occupies 87.6% of memory and GEMM operations consume 61.6% of execution time in LLMs. Given the roofline model, the more information we can pack into a given memory location, the better. We're always bounded by memory bandwidth. This is also why quantization works: it reduces the memory footprint by using smaller data types.

The Core Problem: Indexing Overhead

The recurring theme with existing sparse approaches is that they overlook the cost of indexing. When you store a sparse matrix, you need metadata to track where the nonzero elements actually are. This metadata has a cost, both in storage and in the time it takes to decode.

The paper introduces a metric called CR (compression ratio) that reveals the ratio between the size of sparsity metadata versus the actual data size. It turns out that common approaches see a CR close to or less than 1 at sparsity levels below 50%. In other words, the metadata can be as large as or larger than the actual data you're storing.

This has real implications for computational intensity. For sparse matrix multiplication (SpMM), the computational intensity becomes:

CI_{SpMM} = \frac{M \cdot N}{\frac{M}{CR} + N}

Compare this to the optimal case where you only store nonzeros:

CI_{Optimal} = \frac{M \cdot N}{M \cdot (1 - s) + N}

The takeaway is that traditional GEMM is typically memory-bounded, but once you factor in the compression ratio CR, it becomes an additional limiting factor, and potentially a large one.

TCA-BME: The Encoding Scheme

So how do you efficiently store sparse weight matrices while keeping indexing overhead low? SpInfer uses a multi-level tiling design that partitions the weight matrix into three hierarchical abstraction levels: BitmapTile, TCTile, and GroupTile. Each level corresponds to distinct computational units within the GPU hardware.

The BitmapTile is the fundamental building block, sized at $BT_{h} \times BT_{w} = 8 \times 8$ . This size lets them exploit tensor core operations and encode everything needed for a single tile in a uint64_t. That's 64 bits, one per element in the tile. It's elegant: each bit indicates whether that position holds a nonzero value.

The TCTile is $TT_{h} \times TT_{w} = 2 \times 2$ BitmapTiles, corresponding to the mma instruction requirements. The mma instruction requires matrix shapes of either $16 \times 16 \times 8$ or $16 \times 8 \times 8$ Tensor core specific constraints. BitmapTiles within a TCTile are laid out in column-major format for reasons relating to mma. It reduces online overhead during computation.

The GroupTile is a $GT_{h} \times GT_{w}$ collection of TCTiles, also stored in column-major order. This gives you three levels of hierarchy, each designed around the GPU's computational units.

The actual storage implementation uses two arrays. GTileOffset records the starting offset position of each GroupTile, acting as an index into the main data. The Values array stores all nonzero elements arranged in nested order: GroupTile → TCTile → BitmapTile. Each BitmapTile is represented as one 64-bit integer in this array.

The total storage overhead works out to:

4B \times (NGT + 1) + 8B \times NBT + 2B \times NNZ

where the number of GroupTiles is $NGT = (\frac{M}{GT_h}) \times (\frac{K}{GT_w})$ and the number of BitmapTiles is $NBT = (\frac{M}{BT_h}) \times (\frac{K}{BT_w})$ .

SMBD: Decoding and Computation

SpInfer leverages CUTLASS splitK parallelism, where each thread processes a portion of the K-dimension independently. Each thread block executes five operations in sequence:

Load GTile from global memory
Decode Wtile via SMBD into register files
Load XTile from global memory
Transfer XTile data into registers for tensor core computations
Execute tensor core computations

Data movement uses LDGSTS for vectorized loads of 128 bytes from global memory to shared memory. The LDSM.M88 instruction loads the XTile from shared memory. This instruction allows a warp to load a $16 \times 16$ matrix tile from shared memory and automatically arrange it into registers in the layout tensor cores expect.

The SMBD process handles decompressing the WTile into the register file. Since they use mma.m16n8k16 for half-precision computations, the warp calculates on $16 \times 16$ chunks of the matrix. Each thread needs to hold 8 elements, storing 2 in each of four registers: Ra0, Ra1, Ra2, Ra3. These registers get populated through bitmap decoding.

The decoding process itself is clever. PopCount (__popcll) counts the number of 1 bits in an integer, telling you how many nonzeros exist in a BitmapTile. MaskedPopCount determines the number of nonzeros before the current position to get the accurate offset into the Values array.

The implementation uses a two-phase approach to amortize costs. For the first element, you compute the base offset as i × 2i corresponds to the thread's "lane", its value [0...31] in a warp. You compute a mask as (1 << offset) - 1, which sets all bits before the offset bit, then run PopCount(b & mask) to find the number of nonzero values before that offset. For the second element at position 2i + 1, you no longer need to do the MaskedPopCount since you already have the true offset from the first element. You're essentially amortizing the cost of that operation across both values the thread needs.

Asynchronous Pipeline Design

The real performance gains come from overlapping memory operations with computation. SpInfer uses two separate shared memory buffers for GTiles and XTiles, allowing you to load data into one buffer while computing on the otherThis is known as latency hiding. The current data resides in one buffer while the other asynchronously fetches the next set via cp.async.

There are actually two cp.async groups working in concert. Once the GTile loading completes, SMBD begins, running concurrently with the XTile loading. This hides the latency of the SMBD decoding process. Even further, after the current tile begins its tensor core computation, SMBD on the next tile begins immediatelyThis is possible because SMBD relies on CUDA cores, not tensor cores. They're separate execution units.

Limitations

One notable limitation: SpInfer doesn't actually support dynamic activation sparsity. The framework is designed around weight sparsity, where the sparsity pattern is fixed after pruning. Dynamic sparsity in activations would require runtime pattern detection and adaptive indexing, which is a different problem entirely.