Coalesced Memory Access

Coalescing is defined at the warp level, the fundamental unit of execution on a GPU (typically 32 threads). The hardware fetches memory for all threads in a warp simultaneously. For coalescing to occur, the memory addresses requested by these 32 threads must form a contiguous, aligned block. If threads access scattered addresses, the hardware may issue multiple, smaller transactions, wasting bandwidth.

The ideal pattern is for thread T in a warp to access memory address A + T, where A is a base address aligned to a cache line boundary (e.g., 128 bytes). This allows the GPU's memory controller to service all requests with a single, wide transaction. Key patterns include:

• Unit stride: Consecutive threads access consecutive elements (e.g., array[threadIdx.x]).
• Misaligned access: Threads access a contiguous block, but starting at an unaligned address, may still be coalesced but can require two transactions.
• Strided or random access: Threads access non-consecutive addresses, causing serialization and poor performance.

GPUs access global memory in fixed-size transactions (e.g., 32, 64, or 128 bytes) aligned to cache lines. When a warp requests data, the hardware issues the minimum number of transactions to cover all requested addresses. Coalescing maximizes the utilization of each transaction. For example, if a warp of 32 threads each requests a 4-byte float, the ideal coalesced read fetches all 128 bytes in one 128-byte transaction. Non-coalesced access could result in 32 separate 32-byte transactions, a 32x bandwidth waste.

The primary metric for coalescing is effective memory bandwidth, calculated as (Bytes Needed by Application) / (Bytes Fetched by Hardware). Perfect coalescing approaches 100%. Poor patterns can drop this to low single digits. This directly impacts kernel performance, as memory latency often dominates execution time. Optimizing for coalesced access is frequently the most significant performance improvement for memory-bound GPU kernels.

Achieving coalescing often requires designing data structures for efficient access by parallel threads. The two primary strategies are:

• Array of Structures (AoS): struct {float x, y, z;} points[N]; Threads accessing points[tid].x are coalesced, but accessing all x, then all y leads to strided access.
• Structure of Arrays (SoA): struct {float x[N], y[N], z[N];} points; Threads accessing points.x[tid] are perfectly coalesced for sequential processing of each field. SoA is generally preferred for GPU computation.

Coalescing primarily optimizes access to global memory (DRAM). Its benefits are amplified by the L2 and L1 cache hierarchy. A coalesced access that hits in cache has ultra-low latency. Furthermore, techniques like prefetching into shared memory rely on first performing a coalesced load from global memory. Non-coalesced access wastes cache space by bringing in unnecessary data, causing thrashing and reducing cache effectiveness for all threads on the GPU.

The classic GEMM (General Matrix Multiply) kernel is a prime example. Optimal implementations load blocks of matrices into shared memory using coalesced reads:

• Threads in a warp load a contiguous row (or column) segment from global memory in a single transaction.
• This pattern minimizes the number of memory requests, saturating the memory bus.
• Non-coalesced access in naive implementations can reduce throughput by an order of magnitude.

Operations like convolution filters (e.g., Gaussian blur, edge detection) apply a kernel across an image. Coalescing is achieved by ensuring threads processing adjacent pixels read adjacent memory addresses.

• A warp processing a horizontal strip of pixels reads a contiguous block of pixel values.
• Padding and memory alignment are used to avoid misaligned accesses that break coalescing.
• This is critical for real-time video processing where latency is bounded.

Parallel reduction algorithms sum an array of values. The optimal pattern uses a strided index that changes with each reduction step to maintain coalescing.

• Initial step: Threads in a warp load contiguous data segments coalescedly.
• Subsequent steps: Use shared memory for intermediate results to avoid global memory access divergence.
• Poor indexing leads to scattered reads, causing 32 separate memory transactions per warp instead of 1 or 2.

Libraries like cuDNN, PyTorch, and TensorFlow rely on coalesced access in their low-level CUDA kernels for layers like Convolution and Fully Connected.

• Tensor layouts (e.g., NHWC vs NCHW) are chosen to ensure that the most rapidly changing dimension in memory is accessed by contiguous threads.
• Kernel auto-tuners experiment with thread block sizes and memory access patterns to maximize coalescing.
• This directly impacts training speed and inference latency.

Simulations involving particles or grid-based methods (e.g., fluid dynamics, n-body problems) store state in Structure of Arrays (SoA) format instead of Array of Structures (AoS).

• SoA stores all particle X coordinates contiguously, then all Y coordinates, etc.
• A warp processing a force calculation reads 32 consecutive X values in one coalesced transaction.
• AoS would cause non-coalesced access as threads read interleaved struct members.

GPU-accelerated databases use coalesced access for primitives like scan, sort, and hash table probes.

• Radix sort partitions data into buckets; writing to buckets uses coalesced patterns to avoid bank conflicts in shared memory and global memory.
• Merge paths ensure output phases maintain coalesced writes.
• Bandwidth utilization from coalescing determines the speedup over CPU implementations for large-scale data analytics.

The memory hierarchy organizes storage into tiers with different capacities, latencies, and bandwidths. On a GPU, this typically includes:

• Registers: Fastest, private to each thread.
• Shared Memory: Low-latency, shared within a thread block.
• L1/L2 Cache: Hardware-managed, automatic caching.
• Global Memory (DRAM): High-capacity but higher latency. Coalescing optimizes access at the global memory level, which has the highest penalty for poor access patterns. Understanding this hierarchy is essential for placing data in the right tier to minimize latency.

A memory access pattern describes the order and structure in which threads in a warp access memory. It is the overarching category for coalesced access. Key patterns include:

• Coalesced: Consecutive threads access consecutive addresses (optimal).
• Strided: Threads access addresses with a constant, non-unit stride.
• Random: Threads access addresses with no discernible pattern (worst-case). The pattern directly determines the number of memory transactions required. Coalescing aims to minimize this to a single transaction per warp for a contiguous memory segment.

A bank conflict occurs in shared memory, which is divided into equally sized memory banks. When two or more threads within a warp attempt to access different data words stored in the same bank simultaneously, the accesses are serialized, causing a performance penalty. This is the shared memory analog to uncoalesced access in global memory. Avoiding bank conflicts involves designing data layouts and access patterns so that concurrent requests map to unique banks.

Global memory, or device memory, is the high-bandwidth DRAM physically located on the GPU. It is:

• The largest memory pool accessible by all threads.
• Relatively high latency (hundreds of cycles).
• Accessed via 32-, 64-, or 128-byte memory transactions. Coalesced access is the primary technique for hiding this latency and saturating the available bandwidth. Uncoalesced access results in wasted bandwidth as only a fraction of the data in each transaction is used.

A warp is the fundamental unit of execution and scheduling on an NVIDIA GPU, typically consisting of 32 threads. Threads in a warp execute in Single Instruction, Multiple Thread (SIMT) fashion. Memory accesses are issued and processed per warp. For coalescing to occur, the 32 threads must coordinate their accesses to contiguous global memory addresses within a defined segment size (e.g., 128 bytes). The warp's behavior defines the granularity at which memory coalescing is evaluated.

The memory controller is the hardware unit on the GPU that manages all communication with the DRAM chips (global memory). It:

• Receives read/write requests.
• Translates them into specific DRAM commands (activate, read, write, precharge).
• Optimizes command scheduling to maximize bus efficiency. Coalesced access patterns allow the memory controller to service an entire warp's request with a minimal number of efficient, burst-oriented DRAM transactions, reducing command overhead and improving overall bandwidth utilization.