Memory Coalescing

Memory coalescing transforms multiple, concurrent memory requests from parallel threads into a single, contiguous memory transaction. On architectures like GPUs and NPUs, memory is accessed in aligned segments (e.g., 32, 64, or 128 bytes).

• Key Insight: If Thread 0 accesses address A, Thread 1 accesses A+4, Thread 2 accesses A+8, and so on, the hardware can combine these into one fetch for the entire aligned block containing addresses A through A+128.
• Contrast with Uncoalesced Access: If threads access random, unaligned addresses (e.g., scattered across memory), each request may trigger a separate, inefficient memory transaction, crippling bandwidth.
• Primary Goal: Minimize the number of memory operations required to service all threads in a warp or wavefront, moving data in large, efficient bursts.

Coalescing is not software magic; it requires specific hardware support and careful data layout. The memory controller and cache hierarchy are designed to recognize and merge adjacent requests.

• Aligned Accesses: Threads must access addresses within the same aligned memory segment. For a 128-byte cache line, ideal access is to consecutive 4-byte words (e.g., float values) starting at a 128-byte boundary.
• Memory Transaction Sizes: Modern accelerators have configurable transaction sizes (32B, 64B, 128B). Coalescing effectiveness depends on matching access patterns to these sizes.
• Global Memory: Coalescing is most critical for global memory (high-latency DRAM). It is less relevant for on-chip memories like shared memory or scratchpad memory, which have lower latency and different access patterns.

Coalescing happens at the level of thread groups. On NVIDIA GPUs, a warp (32 threads) is the fundamental unit of execution and memory coalescing.

• Ideal Pattern (Sequential, Unit Stride): Threads in a warp accessing contiguous, increasing addresses (e.g., array[threadIdx.x]). This creates a single, contiguous block request.
• Strided Access: Threads accessing memory with a constant stride (e.g., every k elements). If the stride is 1, it's ideal. Large strides (>1) can reduce or prevent coalescing, as accesses fall into different memory segments.
• Worst Case (Random or Unaligned): Threads accessing completely unrelated addresses. This results in 32 separate memory transactions, the worst possible scenario for bandwidth.
• Compiler Role: Compilers for frameworks like CUDA or OpenCL attempt to structure loops and array indices to promote unit-stride access within warps.

Compilers for parallel architectures apply specific loop and data layout transformations to create coalesce-friendly access patterns.

• Loop Interchange: Swapping nested loops to ensure the innermost loop iterates over contiguous memory. For a 2D array in row-major order, the column index should be the inner loop variable.
• Array Padding: Adding unused elements to the end of array rows to ensure that the start of each row is aligned to a memory segment boundary, preventing misaligned accesses that break coalescing.
• Data Layout Transformation: Changing from Array of Structures (AoS) to Structure of Arrays (SoA). AoS (struct {float x, y, z;} points[N];) causes threads to access non-contiguous x fields. SoA (struct {float x[N], y[N], z[N];}) ensures all x values are contiguous, enabling perfect coalescing.

Memory coalescing does not operate in isolation; it interacts with and enables other critical performance techniques.

• Kernel Tiling: Tiling loads a block of data into fast shared memory. The load from global memory into the tile must be coalesced to be efficient. The tile is then accessed from shared memory with less stringent pattern requirements.
• Vectorized Load/Store Instructions: Using wider data types (e.g., float4) or vector intrinsics can explicitly request larger, contiguous memory transactions, making coalescing more explicit and robust.
• Bank Conflicts: After coalesced data is in shared memory, access must avoid bank conflicts. Coalescing optimizes the transfer to on-chip memory; other optimizations handle the usage.
• Prefetching: Coalesced access patterns are highly predictable, allowing hardware and software prefetchers to effectively bring data into caches before it is needed.

Performance engineers use specific tools and metrics to diagnose poor memory coalescing, which is a primary cause of memory-bound kernels.

• Key Metric: Global Memory Load/Store Efficiency: Profilers (e.g., NVIDIA Nsight Compute) report this as the ratio of requested memory transactions to actual transactions executed. A low efficiency (e.g., 25% or 12.5%) indicates uncoalesced access.
• Memory Throughput: Compare achieved DRAM bandwidth to the hardware's peak theoretical bandwidth. A large gap often points to coalescing or other memory access problems.
• The Roofline Model: Coalescing directly improves arithmetic intensity (operations per byte) by reducing the denominator (bytes fetched). A kernel stuck far below the memory-bound roofline likely suffers from inefficient access patterns.
• Visualization: Profilers can visualize memory addresses accessed by warps, showing scattered versus contiguous patterns.

A compiler optimization that merges multiple, separate computational kernels into a single, larger kernel. This reduces the overhead of repeated kernel launches and, crucially, eliminates intermediate writes and reads to global memory. By keeping data in faster memory hierarchies (like registers or shared memory) between fused operations, it inherently creates opportunities for improved memory access patterns, including coalescing.

The strategic orchestration of data movement across different levels of memory (e.g., DRAM, shared memory, registers) based on their latency, bandwidth, and size. Coalescing optimizes access at the DRAM-to-cache level. Effective hierarchy management ensures data resides in the fastest possible memory (registers/L1) for computation, minimizing the frequency and cost of coalesced transactions to slower global memory.

A key performance metric defined as the number of arithmetic operations performed per byte of data transferred from main memory. It determines if a kernel is compute-bound or memory-bound. Memory coalescing directly attacks the memory-bound problem by maximizing the useful data retrieved per memory transaction, thereby effectively increasing the operational intensity of a kernel and moving it closer to the hardware's compute roof.

A compiler transformation that swaps the order of nested loops to improve data locality. This is a foundational optimization that enables memory coalescing. For example, transforming a loop that accesses a column-major array in the inner loop (causing strided, non-coalesced accesses) to have the row-major loop innermost creates contiguous, coalesceable memory access patterns for parallel threads.

A loop transformation that partitions a large iteration space into smaller, regular blocks or tiles. The primary goal is to fit a working set of data into a faster memory hierarchy like shared memory or registers. Once data is in shared memory, subsequent accesses by threads in a block can be carefully orchestrated to ensure coalesced reads and writes to this on-chip memory, which is also subject to bank conflict issues analogous to coalescing.

The execution model used by GPUs and similar NPUs, where a single instruction is issued to a group of threads (a warp or wavefront). Memory coalescing is a hardware feature designed for the SIMT model. It relies on the fact that threads in a warp execute in lockstep, allowing their concurrent memory requests—if properly aligned and contiguous—to be serviced by a minimal number of wide memory transactions.