Memory Access Pattern

Coalesced memory access is the optimal pattern for GPU global memory. It occurs when consecutive threads in a warp (e.g., thread 0, 1, 2...) access consecutive memory addresses in a single, aligned transaction.

• Mechanism: The memory subsystem can combine these individual requests into one or a minimal number of cache line transactions (e.g., 128 bytes).
• Impact: Maximizes the utilization of the memory bus, achieving near-peak bandwidth.
• Example: A warp of 32 threads reading 32 consecutive 4-byte floats results in one or two efficient memory transactions.

Strided access is a suboptimal pattern where threads in a warp access memory locations separated by a constant stride (distance).

• Small Stride: A stride of 2 or 4 may still be partially coalesced but reduces efficiency.
• Large Stride: A stride larger than the warp size (e.g., 1024) results in completely uncoalesced, scattered accesses. Each thread's request may require a separate memory transaction, devastating bandwidth.
• Common Cause: Often arises from poor data layout, such as accessing a column of a row-major matrix.

A bank conflict occurs in GPU shared memory, which is organized into independent memory banks (typically 32). When two or more threads within a warp attempt to access different data words stored in the same bank on the same clock cycle, the accesses are serialized.

• Mechanism: Shared memory has high bandwidth only if all 32 threads access 32 different banks concurrently.
• Impact: Causes serialized, sequential access, turning a potential 32-wide operation into 2-32 sequential steps.
• Resolution: Can often be avoided through data padding or access pattern restructuring.

These are fundamental principles for cache efficiency across the memory hierarchy.

• Spatial Locality: The tendency that if a memory location is accessed, nearby locations will likely be accessed soon. Coalesced access exploits spatial locality for global memory. Caches load data in blocks (cache lines) assuming spatial locality.
• Temporal Locality: The tendency that a recently accessed memory location will be accessed again in the near future. Efficient algorithms reuse loaded data from registers or shared memory to exploit temporal locality, minimizing trips to slower global memory.

GPU global memory accesses are serviced in fixed-size transactions (e.g., 32, 64, or 128 bytes) aligned to their size. Misalignment wastes bandwidth.

• Aligned Access: The first address accessed by a warp is a multiple of the transaction size. This ensures the minimal number of transactions are used.
• Misaligned Access: The warp's access pattern straddles a transaction boundary, requiring two transactions to fetch the same amount of useful data, effectively halving bandwidth.
• Optimization: Data structures should be aligned in memory to the natural transaction boundaries of the architecture.

For data that is strictly read-only and accessed uniformly by all threads (e.g., lookup tables, constant parameters), GPUs provide specialized cache paths.

• Constant Memory: A small, cached region (64KB) optimized for broadcast. When all threads in a warp read the same address, the data is fetched once and broadcast, achieving tremendous bandwidth.
• Texture Memory / Read-Only Cache: A dedicated cache path for global memory marked as read-only (via __ldg() or const __restrict__). It is optimized for spatial locality and 2D/3D access patterns, can reduce contention with the L1 cache, and is vital for specific workloads like image processing.

Coalesced memory access is the optimal pattern where consecutive threads in a GPU warp access consecutive memory addresses within a single, aligned transaction. This allows the memory controller to service all 32 threads' requests in one or a minimal number of operations, maximizing bandwidth utilization.

• Mechanism: When threads T0, T1, T2... T31 access addresses A, A+4, A+8... A+124 (for 4-byte data), the hardware can combine these into one 128-byte cache line request.
• Impact: Non-coalesced access, where threads read from scattered addresses, forces serialized transactions, reducing effective bandwidth by a factor of 32 or more.
• Implementation: Requires careful data structure design (e.g., Structure of Arrays vs. Array of Structures) and kernel indexing to ensure adjacent threads process adjacent data elements.

Shared memory is a programmer-managed, on-chip cache (typically 64-128KB per Streaming Multiprocessor) used to stage data from global memory for cooperative reuse within a thread block. Its ultra-low latency (≈1-2 cycles) makes it ideal for eliminating redundant global memory loads.

• Common Pattern: A thread block collectively loads a tile of data from global memory into shared memory. Threads then perform numerous computations accessing this fast, local cache.
• Key Challenge: Avoiding bank conflicts. Shared memory is divided into 32 banks (matching warp size). Simultaneous accesses to different addresses in the same bank cause serialization.
• Optimization: Use padding in data structures or memory access stride patterns to ensure concurrent accesses map to unique banks, enabling true parallel access.

Effective use of the GPU memory hierarchy—L1/L2 cache, shared memory, registers—reduces traffic to high-latency global memory. Prefetching is a key technique to hide latency by asynchronously loading data before it is needed.

• L1/L2 Cache: Hardware-managed, transparent to the programmer. Access patterns with spatial locality (accessing nearby addresses) and temporal locality (reusing data) improve cache hit rates.
• Software Prefetching: Kernels can be written to explicitly load data for the next computational iteration into registers or shared memory while processing the current iteration, keeping execution pipelines full.
• Constant Memory: For read-only data accessed uniformly by all threads, constant memory (cached) provides broadcast efficiency, where a single read can service an entire warp.

Bank conflicts in shared memory and partition camping in global memory are two major access pattern pathologies that serialize operations and destroy parallelism.

• Bank Conflict Resolution: Occurs when multiple threads in a warp access different words within the same shared memory bank. Solved by memory layout transformations (padding arrays) or modifying access strides.
• Partition Camping: In global memory, the DRAM is divided into partitions. When many concurrent memory requests target the same partition, a queue forms while other partitions sit idle. This is mitigated by memory address interleaving—ensuring sequential thread accesses are spread across multiple partitions via strategic base address offsets or data layout choices.

GPU global memory transactions are served in units of 32, 64, or 128 bytes (cache lines). Wasting these transactions leads to bandwidth inefficiency. Optimization involves aligning data and ensuring full utilization of each transaction.

• Alignment: Allocating memory on 128-byte boundaries ensures the first access by a warp is a single transaction. Misaligned accesses may require two transactions.
• Transaction Size: Using the smallest transaction size that satisfies the access pattern saves bandwidth. For example, if a warp only needs 32 consecutive bytes, enabling 32-byte transactions (via compiler flags or GPU architecture) is more efficient than a default 128-byte transaction.
• Structure Padding: In structures, arranging elements by access frequency and coalescing potential can reduce the number of cache lines touched per thread.

Identifying suboptimal memory patterns requires instrumentation. Profilers provide metrics to diagnose issues and guide optimization.

• Key Metrics:
  • Global Memory Load/Store Efficiency: Percentage of bytes requested in a transaction that are actually used by the kernel. Low efficiency indicates poor coalescing.
  • Shared Memory Bank Conflict Count: Direct measure of bank serialization events.
  • L1/L2 Cache Hit Rates: Low hit rates suggest poor locality.
• Essential Tools:
  • NVIDIA Nsight Compute: Provides detailed per-kernel analysis of memory access patterns, transaction sizes, and stall reasons.
  • nvprof / NVIDIA Nsight Systems: Offers system-wide timeline profiling to correlate kernel execution with memory transfer bottlenecks.
• Workflow: Profile, identify the primary bottleneck (e.g., low coalescing efficiency), apply a targeted optimization, and re-profile to quantify the gain.

An optimal GPU memory access pattern where consecutive threads in a warp access consecutive memory addresses in a single, aligned transaction. This pattern is fundamental for achieving peak bandwidth from GPU global memory.

• Mechanism: The GPU's memory controller can service a warp's request for 32 consecutive 4-byte words in one or two 128-byte cache line transactions.
• Impact: Non-coalesced access forces multiple smaller transactions, drastically reducing effective bandwidth and increasing latency.
• Example: A kernel where thread i accesses array element array[i] is typically coalesced.

The organization of memory subsystems in a GPU into multiple levels with differing capacities, latencies, and bandwidths, designed to exploit data locality.

Key levels include:

• Registers: Fastest, private to each thread.
• Shared Memory: Low-latency, software-managed cache shared by threads in a block.
• L1/L2 Cache: Hardware-managed caches for automatic data reuse.
• Global Memory: High-capacity, high-latency DRAM on the GPU board.

Optimizing access patterns involves promoting data to faster levels of this hierarchy.

A performance penalty that occurs in GPU shared memory when two or more threads within a warp attempt to access different data words stored in the same memory bank simultaneously.

• Cause: Shared memory is divided into equally sized banks (e.g., 32 banks). Concurrent access to the same bank causes serialized access.
• Avoidance: Techniques include memory padding or designing access patterns so threads access different banks (e.g., stride-1 access).
• Severity: A 32-way bank conflict can serialize a warp's access 32 times, negating the benefit of shared memory.

A memory management architecture that creates a single, contiguous virtual address space shared between a CPU and GPU, simplifying data sharing and enabling on-demand paging.

• Benefit: Eliminates the need for explicit cudaMemcpy calls; pointers can be accessed from either processor.
• Access Pattern Impact: UVM can hide the complexity of data movement, but poor access patterns (causing frequent page faults) can lead to severe performance degradation as pages are migrated between CPU and GPU memory.
• Use Case: Essential for workloads where data access patterns are sparse or unpredictable.

Host (CPU) memory that is prevented from being paged out to disk by the operating system, enabling high-bandwidth Direct Memory Access (DMA) transfers to and from GPU device memory.

• Performance: Critical for achieving maximum bandwidth for cudaMemcpy operations. Transfers from pageable host memory require a staging buffer, adding latency.
• Trade-off: Excessive allocation can reduce overall system performance by reducing available physical memory for other processes.
• Pattern Relevance: Using pinned memory for host buffers accessed via zero-copy or frequent H2D/D2H transfers is a key optimization for streaming data.

A performance optimization technique where a large block of memory (a pool) is pre-allocated, and smaller allocations are sub-allocated from this pool, reducing fragmentation and allocation overhead.

• Problem Solved: Frequent cudaMalloc and cudaFree calls are expensive and can lead to memory fragmentation, preventing allocation of large contiguous segments.
• Access Pattern Link: A memory pool ensures that allocations for tensors with similar lifecycles are physically closer, which can improve cache locality and access patterns.
• Implementation: Used by deep learning frameworks (e.g., PyTorch's Caching Allocator) and CUDA's cudaMallocAsync for stream-ordered allocations.