Gather-Scatter Operations

A gather operation collects data elements from a source array using a list of indices, assembling them into a contiguous destination vector. This is essential for loading the non-zero weights and corresponding activations required for a sparse matrix multiplication.

• Mechanism: dest[i] = src[indices[i]]
• Use Case: Loading a dense vector of values from a sparse matrix's non-zero entries, as defined by its index arrays (e.g., CSR column indices).
• Performance Impact: Efficiency is dominated by the random read access pattern from the source memory, which can cause cache thrashing.

A scatter operation writes elements from a contiguous source vector to a destination array at specified indices. This is used to accumulate partial results, such as output activations, to their correct locations after a sparse computation.

• Mechanism: dest[indices[i]] = src[i]
• Use Case: Writing the computed results of a sparse matrix-vector multiplication to the correct rows of the output tensor.
• Performance Impact: Random write access can be even more costly than gathers due to write-back policies and potential for write conflicts, requiring atomic operations in parallel contexts.

The core purpose of gather-scatter is to bridge the gap between indexed data (sparse, irregular) and contiguous data (dense, regular).

• Sparse Tensor Formats (like CSR, COO) store data as (value, index) pairs.
• Compute Kernels (like SpMM) require dense vectors of operands for efficient SIMD/vector processing.
• Gather-Scatter Cycle: 1) Gather non-zero operands into dense buffers. 2) Perform dense compute on buffers. 3) Scatter results back via indices. This cycle is the heart of sparse inference.

Modern CPU and GPU ISAs include explicit instructions to accelerate these patterns, reducing overhead versus software-managed loops.

• CPU: Intel AVX-512 provides VGATHERDPS and VSCATTERDPS instructions.
• GPU: NVIDIA GPUs execute gather-scatter via their load/store units and the ld.global.nc/st.global.wb instructions, often coalesced within a warp.
• Specialized Accelerators: NPUs and TPUs often integrate dedicated gather-scatter units into their memory hierarchies to minimize latency for sparse workloads.

The primary challenge for gather-scatter is memory bandwidth utilization and latency, not compute.

• Irregular Access: Non-sequential memory reads/writes defeat spatial prefetchers and reduce cache line utilization.
• Metadata Overhead: Processing index arrays consumes additional memory bandwidth and instruction cycles.
• Load Imbalance: In parallel execution, threads may have varying numbers of indices to process, causing some cores to stall waiting for others.
• The Sparse Efficiency Gap: The theoretical FLOP reduction from sparsity is often not fully realized due to these memory system overheads.

Gather-scatter operations are the enabling primitives for higher-level sparse kernels like Sparse Matrix-Matrix Multiplication (SpMM) and Sparse Convolution.

• SpMM Kernel Phases: 1) Gather input feature map slices. 2) Perform dense sub-matrix multiplications. 3) Scatter-add results to the output matrix.
• Blocked Sparsity (e.g., N:M): Uses structured indices to make gather-scatter patterns more regular and efficient, improving cache locality and enabling use of Sparse Tensor Cores.
• Compiler Role: Advanced compilers fuse operations to minimize intermediate gather-scatter steps, a key optimization called sparse operator fusion.

A data structure for efficiently storing tensors where most elements are zero. Instead of a dense array, it encodes only non-zero values and their indices. Common formats include:

• CSR (Compressed Sparse Row): Optimized for row-wise operations.
• CSC (Compressed Sparse Column): Optimized for column-wise access.
• COO (Coordinate Format): Simple list of (index, value) tuples. The choice of format directly dictates the efficiency of the gather (reading) and scatter (writing) patterns required by sparse kernels.

The core computational kernel that multiplies a sparse matrix by a dense matrix. It is the workhorse for sparse linear layers in neural networks. Efficient SpMM implementation requires:

• Gathering rows/columns of the dense matrix based on the sparse matrix's indices.
• Performing reduced, skipped multiplications for zero entries.
• Scattering or accumulating partial results to the output. Performance is dominated by memory bandwidth and the overhead of index processing, not raw FLOPs.

A specialized hardware unit in modern GPUs (e.g., NVIDIA Ampere/Ada/Hopper architectures) designed to accelerate sparse matrix operations. It leverages structured sparsity patterns (like 2:4, where 2 of every 4 elements are non-zero).

• These cores natively understand pruning masks or bitmask encodings.
• They can effectively skip zero-valued computations at the hardware level, potentially doubling theoretical compute throughput for eligible operations.
• They require software to format data in a compatible sparse data layout.

A major performance challenge in parallel sparse computation. Because non-zero elements are distributed irregularly, different processor threads or cores are assigned vastly different amounts of work.

• One warp may process a dense row with many non-zeros, while another handles a nearly empty row.
• This leads to thread divergence and underutilization of parallel hardware, as faster threads wait for slower ones to finish.
• Advanced sparse kernel design focuses on workload partitioning and dynamic scheduling to mitigate this imbalance.

The additional computational cost incurred during sparse execution beyond the floating-point operations. This overhead can diminish the theoretical gains from zero-skipping and includes:

• Index Decoding: Reading and interpreting CSR row pointers or COO coordinate lists.
• Pointer Chasing: Irregular memory accesses to gather input data.
• Conditional Branching: Instructions to check for loop boundaries or zero values.
• Metadata Processing: Handling the pruning mask or bitmask. The sparse efficiency gap is often the result of this overhead outweighing FLOPs reduction.

A form of structured pruning where, in every block of M consecutive weights (e.g., within a 1x4 vector), only N are allowed to be non-zero. A common pattern is 2:4 sparsity.

• This creates a regular, predictable sparsity pattern that is hardware-friendly.
• It enables efficient gather-scatter using vectorized loads/stores and minimal metadata (a simple 2-bit mask per block).
• Modern Sparse Tensor Cores are explicitly designed to exploit this pattern, offering significant speedups over dense execution for compliant models.