bundle-course-cloud-platforms<\/a><\/h3>\n2. Global Memory: The High-Bandwidth Reservoir<\/b><\/h2>\n
Global Memory represents the largest, most persistent, and most accessible level of the memory hierarchy. It is the only memory space visible to the host CPU (via PCIe or NVLink) and all threads across all blocks on the GPU. Physically, Global Memory consists of the VRAM (Video RAM) soldered onto the graphics card PCB\u2014typically GDDR6X in consumer\/workstation cards (e.g., RTX 4090, RTX 6000 Ada) or HBM (High Bandwidth Memory) in data center accelerators (e.g., A100, H100).<\/span>14<\/span><\/p>\n2.1 The Physics of Bandwidth: GDDR vs. HBM<\/b><\/h3>\n
The performance characteristics of Global Memory are defined by the physical interface.<\/span><\/p>\n\n- GDDR6X:<\/b> Utilized in architectures like Ada Lovelace (e.g., RTX 4090), this technology relies on high clock speeds and narrow buses (e.g., 384-bit) to achieve bandwidths approaching 1 TB\/s. It uses PAM4 signaling to transmit two bits per clock cycle, increasing throughput but also sensitivity to signal integrity.<\/span>14<\/span><\/li>\n
- HBM2e\/HBM3:<\/b> Utilized in the Ampere A100 and Hopper H100, HBM stacks memory dies directly on the GPU interposer. This allows for an ultra-wide bus (e.g., 4096-bit or higher) running at lower clocks, delivering massive bandwidths of 1.5 TB\/s to over 3 TB\/s.<\/span>2<\/span><\/li>\n<\/ul>\n
Despite these massive numbers, Global Memory remains the bottleneck. The <\/span>Computational Intensity<\/b> (or Arithmetic Intensity) of a kernel\u2014defined as FLOPs performed per byte transferred\u2014must be sufficiently high to overcome the limitation of the memory bus. For an A100 GPU with 19.5 TFLOPS (FP32) and 1.6 TB\/s bandwidth, a kernel must perform roughly 12 floating-point operations for every byte loaded just to saturate the compute units. Most “simple” kernels (like vector addition) are strictly memory-bound, meaning their performance is purely a function of how efficiently they utilize Global Memory bandwidth.<\/span>2<\/span><\/p>\n2.2 Memory Coalescing: The Critical Optimization<\/b><\/h3>\n
The memory controller does not interact with DRAM at the granularity of individual bytes or floats. It operates on <\/span>transactions<\/b> (cache lines), typically 32 bytes, 64 bytes, or 128 bytes in size.<\/span>3<\/span> When a warp executes a global memory load instruction, the hardware inspects the addresses requested by the 32 threads.<\/span><\/p>\nThe Coalescing Mechanism:<\/span><\/p>\nIf the addresses are contiguous and aligned\u2014for example, Thread $k$ accesses address $Base + k \\times 4$\u2014the hardware coalesces these 32 requests into a single or minimum number of transactions. For 32-bit words (4 bytes), a full warp requests $32 \\times 4 = 128$ bytes. If aligned, this results in exactly one 128-byte transaction, achieving 100% bus utilization efficiency.3<\/span><\/p>\nThe Penalty of Uncoalesced Access:<\/span><\/p>\nIf threads access memory with a stride\u2014for example, Thread $k$ accesses $Base + k \\times 8$\u2014the requested addresses span 256 bytes. The memory controller must issue two 128-byte transactions (or more, depending on alignment) to fetch the data. However, only half of the data in those transactions is actually used by the threads. This reduces effective bandwidth by 50%. In random access patterns (e.g., pointer chasing or indirect indexing A[i]]), the efficiency can drop to 3-4%, as a full 128-byte line is fetched to satisfy a request for a single 4-byte value.3<\/span><\/p>\n\n\n\n| Access Pattern<\/b><\/td>\n | Description<\/b><\/td>\n | Transactions per Warp (approx)<\/b><\/td>\n | Bus Efficiency<\/b><\/td>\n<\/tr>\n\n| Sequential Aligned<\/b><\/td>\n | $Address = Base + tid$<\/span><\/td>\n | 1 (128 bytes)<\/span><\/td>\n | 100%<\/span><\/td>\n<\/tr>\n\n| Sequential Misaligned<\/b><\/td>\n | $Address = Base + tid + Offset$<\/span><\/td>\n | 2 (128 bytes)<\/span><\/td>\n | ~50-80%<\/span><\/td>\n<\/tr>\n\n| Strided (Stride 2)<\/b><\/td>\n | $Address = Base + tid \\times 2$<\/span><\/td>\n | 2 (128 bytes)<\/span><\/td>\n | 50%<\/span><\/td>\n<\/tr>\n\n| Strided (Large)<\/b><\/td>\n | $Address = Base + tid \\times 32$<\/span><\/td>\n | 32 (32 bytes each)<\/span><\/td>\n | ~12.5%<\/span><\/td>\n<\/tr>\n\n| Random<\/b><\/td>\n | Indirect access<\/span><\/td>\n | up to 32<\/span><\/td>\n | < 10%<\/span><\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n2.3 The Evolution of Caching: From Fermi to Ada Lovelace<\/b><\/h3>\nGlobal memory access is mediated by a multi-level cache hierarchy that has evolved significantly.<\/span><\/p>\n\n- Fermi\/Kepler:<\/b> Relied heavily on a small L1 cache and a relatively small L2. L1 could be configured to prefer Shared Memory or Cache.<\/span><\/li>\n
- Maxwell\/Pascal:<\/b> Global memory loads typically bypassed L1 and went straight to L2, using L1 primarily for register spills and local memory.<\/span><\/li>\n
- Volta\/Turing\/Ampere:<\/b> Re-introduced a strong, unified L1 Data Cache and Shared Memory architecture. In the NVIDIA A100 (Ampere), each SM contains 192 KB of on-chip memory that can be partitioned between L1 Cache and Shared Memory (e.g., 164 KB Shared \/ 28 KB L1). This allows the hardware to cache global loads in L1, providing a lower latency path for frequently accessed global data.<\/span>7<\/span><\/li>\n
- Ada Lovelace (RTX 40 Series):<\/b> Marked a paradigm shift by massively expanding the L2 Cache. The AD102 chip features up to 96 MB of L2 cache (compared to ~6 MB in Ampere GA102). This massive Last-Level Cache (LLC) allows entire working sets (e.g., intermediate activation layers in neural networks, ray tracing BVH structures) to reside on-chip, drastically reducing traffic to the slow GDDR6X memory.<\/span>19<\/span> This architectural change effectively turns Global Memory into a backing store for many workloads, mitigating the impact of non-coalesced access patterns if the working set fits in L2.<\/span><\/li>\n<\/ul>\n
2.4 Asynchronous Copy and the Compute-Data Overlap<\/b><\/h3>\nA critical bottleneck in legacy architectures was the utilization of execution cores for data movement. Loading data from Global to Shared memory required threads to issue load instructions, wait for data to arrive in registers, and then issue store instructions to Shared Memory.<\/span><\/p>\nThe Ampere architecture introduced the Asynchronous Copy (cp.async) instruction. This hardware feature allows the SM to offload the transfer of data from Global Memory to Shared Memory directly to the DMA (Direct Memory Access) engine, bypassing the Register File entirely.<\/span><\/p>\n\n- Mechanism:<\/b> Threads issue the copy command and then are free to execute other independent instructions (e.g., FP32 math) while the data is in flight.<\/span><\/li>\n
- Latency Hiding:<\/b> This explicitly overlaps compute and data transfer at the instruction level, rather than just the warp level.<\/span><\/li>\n
- Register Relief:<\/b> Because data does not pass through registers, register pressure is reduced, potentially allowing for higher occupancy.<\/span>7<\/span><\/li>\n<\/ul>\n
3. Shared Memory: The Programmer-Managed L1<\/b><\/h2>\nIf Global Memory is the warehouse, Shared Memory is the workbench. It is a block of high-speed SRAM located physically within the SM, offering low-latency (20-50 cycles) and high-bandwidth access comparable to the register file.<\/span>5<\/span> Unlike the L1 cache, which is managed by hardware eviction policies (LRU, etc.), Shared Memory is explicitly allocated and managed by the CUDA kernel code. This determinism makes it the most powerful tool for optimizing data reuse.<\/span><\/p>\n3.1 Use Cases: Tiling and Inter-Thread Communication<\/b><\/h3>\nThe primary application of Shared Memory is Tiling (or Blocking). In algorithms such as dense Matrix Multiplication ($C = A \\times B$), a naive implementation requires every thread to load a full row of $A$ and column of $B$ from global memory. For a matrix of size $N$, this results in $2N$ global loads per thread, or $2N^3$ loads total.<\/span><\/p>\nBy loading a small “tile” (e.g., $16 \\times 16$) of $A$ and $B$ into Shared Memory, threads can perform operations on this cached data. Each data element loaded from Global Memory is reused $Tile\\_Width$ times. This reduces Global Memory bandwidth pressure by a factor of the tile width, often transforming a bandwidth-bound kernel into a compute-bound one.2<\/span><\/p>\nFurthermore, Shared Memory is the only high-speed medium for communication between threads in a block. It enables parallel reduction algorithms (summing an array), prefix sums (scan), and sorting networks, where threads must exchange partial results.<\/span>21<\/span><\/p>\n3.2 The Mathematics of Bank Conflicts<\/b><\/h3>\nShared Memory is not a monolithic block; it is divided into <\/span>32 Banks<\/b> (in modern architectures). Memory addresses are mapped to banks in a round-robin fashion: Address $A$ maps to Bank $(A \/ 4 \\text{ bytes}) \\% 32$. Ideally, the 32 threads in a warp access 32 distinct banks simultaneously, yielding full bandwidth (e.g., 32 words per cycle).<\/span>23<\/span><\/p>\nBank Conflicts<\/b> arise when multiple threads in a warp request addresses that map to the <\/span>same<\/span><\/i> bank.<\/span><\/p>\n\n- Serialization:<\/b> If $N$ threads access the same bank, the hardware splits the request into $N$ separate serialized transactions. A 2-way conflict halves the throughput; a 32-way conflict reduces it to 1\/32nd of the peak.<\/span>17<\/span><\/li>\n
- Broadcast Exception:<\/b> If all threads (or a subset) access the <\/span>exact same address<\/span><\/i>, the hardware recognizes this and performs a broadcast, serving the data in a single cycle. Multicast involves serving multiple threads reading the same address once, then moving to the next unique address.<\/span>6<\/span><\/li>\n<\/ul>\n
3.2.1 Case Study: Stride and Padding<\/b><\/h4>\nConsider a matrix declared as __shared__ float A. If threads access a column (e.g., A[tid]), the access stride is 32 floats (128 bytes).<\/span><\/p>\n\n- Thread 0 accesses A $\\rightarrow$ Bank 0.<\/span><\/li>\n
- Thread 1 accesses A $\\rightarrow$ A. Since 32 words wrap around the 32 banks exactly, this address also maps to Bank 0.<\/span><\/li>\n
- Result: 32-way bank conflict. The column read is serialized.<\/span><\/li>\n<\/ul>\n
The Solution (Padding):<\/b> Declare the array as __shared__ float A.<\/span><\/p>\n\n- Stride is now 33 words.<\/span><\/li>\n
- Thread 0 accesses Bank 0.<\/span><\/li>\n
- Thread 1 accesses index 33, which maps to Bank 1 ($33 \\% 32 = 1$).<\/span><\/li>\n
- Result: Conflict-free access. This technique is standard in Matrix Transpose kernels.<\/span>26<\/span><\/li>\n<\/ul>\n
3.3 Dynamic vs. Static Allocation<\/b><\/h3>\nShared Memory can be declared statically or dynamically:<\/span><\/p>\n\n- Static:<\/b> __shared__ float data; – Size is fixed at compile time. Faster to implement but inflexible.<\/span><\/li>\n
- Dynamic: extern __shared__ float data; – Size is specified at kernel launch: kernel<<<grid, block, size>>>(…).<\/span>
\n<\/span>Dynamic allocation is crucial for creating portable code that maximizes utility across different GPU generations with varying shared memory capacities (e.g., 48KB on Kepler vs. 100KB+ on Ada).2 The kernel code must manually calculate pointers\/offsets into this single extern array if multiple data structures are needed.24<\/span><\/li>\n<\/ul>\n3.4 Hardware Acceleration: Async Barriers (Ampere+)<\/b><\/h3>\nWith the introduction of asynchronous copies (cp.async), synchronization became complex. Standard __syncthreads() is a heavy-handed barrier that waits for all threads to reach a point. Ampere introduced <\/span>Asynchronous Barriers<\/b> (mbarrier), which split the “arrive” and “wait” phases.<\/span><\/p>\n\n- Arrive:<\/b> Threads signal they have reached a point (e.g., issued copy commands).<\/span><\/li>\n
- Compute:<\/b> Threads execute independent math while waiting for memory.<\/span><\/li>\n
- Wait: Threads wait for the barrier (memory transfer) to complete.<\/span>
\n<\/span>This fine-grained control allows for software pipelining (double buffering) within shared memory, keeping the ALU pipeline full while loading the next tile of data.7<\/span><\/li>\n<\/ul>\n4. Local Memory: The Misunderstood Abstraction<\/b><\/h2>\n“Local Memory” is a term that frequently confuses newcomers because it refers to <\/span>scope<\/b>, not speed or location. Logically, it is private to a thread. Physically, it resides in <\/span>Global Memory<\/b> (DRAM). Consequently, it suffers from the same high latency and bandwidth constraints as Global Memory, although it benefits from L1\/L2 caching.<\/span>11<\/span><\/p>\n4.1 Triggering Local Memory Usage<\/b><\/h3>\nThe CUDA compiler (NVCC) resorts to Local Memory only when it cannot fit data into the Register File. This happens in three primary scenarios:<\/span><\/p>\n\n- Register Spilling:<\/b> This is the most common cause. If a kernel is complex and uses more registers than the hardware limit (e.g., 255 per thread) or the launch bounds limit, the compiler “spills” the excess variables to Local Memory. This manifests as “Local Load\/Store” instructions in profiling tools (Nsight Compute) and usually signals a severe performance degradation.<\/span>8<\/span><\/li>\n
- Dynamic Indexing of Arrays:<\/b> If a thread declares a small array float arr and accesses it with a variable index arr[i] where i is not known at compile time, the compiler <\/span>must<\/span><\/i> place arr in Local Memory. Registers cannot be addressed dynamically by the hardware (there is no “register indirect” addressing mode). If the index is constant (arr), it stays in registers.<\/span>13<\/span><\/li>\n
- Large Structures:<\/b> Structures or arrays too large to fit in the register budget are placed in Local Memory.<\/span>12<\/span><\/li>\n<\/ol>\n
4.2 Architectural Impact of Spilling<\/b><\/h3>\nSpilling registers to Local Memory increases the traffic on the L1\/L2 caches and memory bus. Since Local Memory is interleaved in Global Memory, heavy spilling can saturate the memory controller, starving explicit global memory loads.<\/span><\/p>\nMitigation via Shared Memory:<\/span><\/p>\nA recently introduced optimization in CUDA 13.0 (and available via research techniques like “RegDem” earlier) allows the compiler to spill registers to Shared Memory instead of Local Memory. Since Shared Memory is on-chip and faster, this reduces the penalty of spilling. This is particularly effective for kernels that have low Shared Memory occupancy but high register pressure.9<\/span><\/p>\n5. Texture Memory: The Graphic Legacy’s Gift to Compute<\/b><\/h2>\nTexture Memory is a specialized read-only path to Global Memory, utilizing dedicated hardware units designed originally for rendering graphics (mapping images onto 3D geometry).<\/span><\/p>\n5.1 Spatial Locality and Z-Order Curves<\/b><\/h3>\nStandard Global Memory is linear. A[y][x] and A[y+1][x] are separated by the width of the row in memory addresses. A thread block reading a 2D patch of data might trigger many cache lines for the row y and many distinct cache lines for row y+1, leading to poor cache reuse.<\/span><\/p>\nTexture Memory stores data in a Block Linear format (often using Morton Codes or Z-order curves). This layout maps 2D coordinates to 1D addresses such that pixels that are spatially close in 2D (neighbors in X and Y) are also close in linear memory address.<\/span><\/p>\nBenefit: For kernels accessing 2D\/3D neighborhoods (e.g., image convolution, stencil codes, fluid dynamics), Texture Memory dramatically improves cache hit rates compared to linear Global Memory.29<\/span><\/p>\n5.2 Hardware Features: Filtering and Boundary Handling<\/b><\/h3>\nTexture units provide “free” operations that would otherwise cost ALU cycles:<\/span><\/p>\n\n- Addressing Modes:<\/b> Handling boundary conditions (e.g., accessing pixel -1) usually requires if statements in code (if x < 0: x = 0). Texture units handle this in hardware via “Clamp”, “Wrap” (modulo), or “Mirror” modes, zeroing the instruction cost.<\/span>30<\/span><\/li>\n
- Linear Interpolation:<\/b> When fetching a coordinate like (1.5, 1.5), the texture unit can return the bilinear interpolation of the four surrounding pixels. This is fundamental for image resizing or volume rendering and provides massive speedups over software implementation.<\/span>10<\/span><\/li>\n
- Format Conversion:<\/b>
| | | | | |
![\"\"](\"https:\/\/uplatz.com\/blog\/wp-content\/uploads\/2025\/12\/The-CUDA-Memory-Hierarchy-Architectural-Analysis-Performance-Characteristics-and-Optimization-Strategies-1024x576.jpg\")
<\/p>\n