2.3 GPU Memory Hierarchy & Bandwidth

GPU Memory Hierarchy & Bandwidth Optimization

Deep learning training and inference workloads are frequently memory-bandwidth bound. In large language models (LLMs) and transformer architectures, memory bandwidth often limits token generation speed and computational efficiency more than raw FLOPs.

To optimize AI workloads, operations and infrastructure teams must understand the distinct tiers of the GPU memory hierarchy, memory coalescing rules, and Unified Memory data movement mechanisms.

1. The GPU Memory Hierarchy Tiers

NVIDIA GPUs employ a hierarchical memory subsystem ranging from ultra-fast, per-thread on-chip registers down to off-chip High-Bandwidth Memory (HBM) and host system RAM.

+-------------------------------------------------------------------------+
|                    PER-THREAD REGISTERS (On-Chip SRAM)                  |
|                    Latency: ~1 Cycle | Bandwidth: ~30-50 TB/s           |
+-------------------------------------------------------------------------+
                                     |
+-------------------------------------------------------------------------+
|              SHARED MEMORY / L1 DATA CACHE (Per-SM SRAM)                |
|                    Latency: ~20-30 Cycles | Bandwidth: ~10-20 TB/s      |
+-------------------------------------------------------------------------+
                                     |
+-------------------------------------------------------------------------+
|                     L2 CACHE (Chip-Wide Unified SRAM)                   |
|                    Latency: ~150-200 Cycles | Bandwidth: ~3-6 TB/s      |
+-------------------------------------------------------------------------+
                                     |
+-------------------------------------------------------------------------+
|             GLOBAL MEMORY (Device HBM3 / HBM3e / GDDR6 DRAM)            |
|                    Latency: ~400-800 Cycles | Bandwidth: 1-8 TB/s       |
+-------------------------------------------------------------------------+
                                     |
+-------------------------------------------------------------------------+
|            HOST SYSTEM RAM (DDR4 / DDR5 via PCIe / NVLink-C2C)          |
|                    Latency: ~1000+ Cycles | Bandwidth: 32-900 GB/s      |
+-------------------------------------------------------------------------+

1. Registers (Per-Thread)

• Location & Scope: Located directly inside the SM core. Private to each individual thread.
• Performance: Single-cycle access latency; aggregate register file bandwidth exceeds 30 to 50 TB/s across the chip.
• Capacity: Typically 64K 32-bit registers (256 KB) per SM (e.g., ~34 MB of total register storage across an entire H100 GPU).
• Register Spilling: If a kernel uses more registers than allocated by the compiler, excess variables "spill" to Local Memory (which is stored in off-chip DRAM and cached in L1/L2), causing significant performance degradation.

2. Shared Memory / L1 Data Cache (Per-SM)

• Location & Scope: On-chip SRAM residing inside each SM. Shared among all threads within a single thread block.
• Performance: ~20 to 30 clock cycles latency; aggregate bandwidth of 10 to 20 TB/s.
• Configurable Partition: On modern architectures (NVIDIA Ampere, Hopper, Blackwell), L1 Data Cache and Shared Memory share a unified on-chip physical SRAM pool. The partition can be configured dynamically per kernel launch (e.g., allocating up to 228 KB to Shared Memory and 28 KB to L1 cache on Hopper H100).
• Programmer Scratchpad: Explicitly declared in CUDA code using the __shared__ keyword to cache repeatedly accessed data tiles, eliminating redundant global memory loads.

3. L2 Cache (Chip-Wide Unified)

• Location & Scope: On-chip SRAM shared by all SMs across the entire GPU.
• Performance: ~150 to 200 clock cycles latency; several TB/s bandwidth.
• Capacity: Massively expanded in recent generations—from 6 MB on Volta V100 to 40 MB on Ampere A100 and 50 MB on Hopper H100.
• Function: Automatically caches all Global and Local memory reads and writes. Hopper also introduced L2 Cache Residency Controls, allowing developers to persist critical data structures (such as KV caches) in L2 SRAM.

4. Global Memory (Device VRAM / HBM)

• Location & Scope: Off-chip DRAM (High-Bandwidth Memory HBM2e/HBM3/HBM3e or GDDR6). Accessible by all threads in all thread blocks across all kernels.
• Performance: 400 to 800 clock cycles latency; bandwidth ranges from 2.0 TB/s (A100 HBM2e) to 3.35 TB/s (H100 HBM3) and 8.0 TB/s (B200 HBM3e).
• Lifetime: Persists for the lifetime of the CUDA application context.

GPU Memory Hierarchy Specifications Summary

2. Memory Access Patterns: Coalesced vs. Uncoalesced Access

Global memory transactions are executed at the granularity of warps (32 threads). When threads in a warp request data from global memory, the hardware memory controller analyzes the memory addresses requested across all 32 lanes.

Coalesced Memory Access

Memory coalescing occurs when the 32 threads in a warp access a contiguous, naturally aligned block of memory addresses (aligned to 32-byte, 64-byte, or 128-byte boundaries).

• If 32 threads in a warp simultaneously request consecutive 4-byte floating-point values ($32 \times 4\text{ bytes} = 128\text{ bytes}$), the hardware memory controller fulfills the entire warp's request in a single 128-byte memory transaction.
• Bus Efficiency: 100% of the data transferred across the memory bus is utilized by the kernel.

Warp Threads:     T0   T1   T2   T3   T4   ...  T31
Requested Addr:  [00] [04] [08] [12] [16]  ...  [124] (Contiguous & Aligned)
Memory Bus:      +---------------------------------------------------------+
                 | Single 128-Byte Coalesced Transaction (100% Efficient)  |
                 +---------------------------------------------------------+

Uncoalesced and Strided Memory Access

If threads access non-contiguous or strided memory addresses (such as traversing a matrix column-by-column in row-major layout, or indexing fields in an Array of Structures AoS):

• For a stride of 32 (where Thread $k$ accesses address $k \times 32$), each thread's requested 4-byte value falls into a completely different 32-byte memory segment.
• The memory controller is forced to issue 32 separate 32-byte transactions ($32 \times 32 = 1,024\text{ bytes}$ transferred) to retrieve only 128 bytes of useful data.
• Bus Efficiency: Only $128 / 1024 = 12.5%$ of the transferred bandwidth contains actual payload; 87.5% of memory bus bandwidth is wasted on unused padding.

Warp Threads:     T0          T1          T2          ...  T31
Requested Addr:  [00]        [32]        [64]         ...  [992] (Strided)
Memory Bus:      [32B Trans] [32B Trans] [32B Trans]  ...  [32B Trans]
Total Transferred: 1024 Bytes | Useful Data: 128 Bytes | Efficiency: 12.5%

Shared Memory Bank Conflicts

Shared memory is organized into 32 independent, equal-width memory banks (each 4 bytes / 32 bits wide):

• Consecutive 32-bit words are assigned to consecutive banks: Word 0 $\rightarrow$ Bank 0, Word 1 $\rightarrow$ Bank 1, ..., Word 31 $\rightarrow$ Bank 31, Word 32 $\rightarrow$ Bank 0.
• No Conflict: If all 32 threads access distinct banks, access is fully parallelized in a single cycle.
• Bank Conflict: If two or more threads in a warp access different memory addresses within the same bank, the hardware serializes the requests (2-way, 4-way, up to 32-way conflict), multiplying access latency.
• Broadcast Exception: If all threads in a warp read the exact same address in a bank, the value is broadcast in a single cycle without a bank conflict.

3. Unified Memory (UM), Page Migration & Pinned Memory

Managing complex distributed memory across host CPUs and GPU accelerators requires understanding virtual memory management mechanisms.

Unified Memory (UM) Architecture

Unified Memory (UM) creates a single, managed, coherent virtual address space that is accessible by both CPU host code and GPU device code using a single unified pointer allocated via cudaMallocManaged(&ptr, size).

// Allocate Unified Memory accessible by both CPU and GPU
float *data;
cudaMallocManaged(&data, N * sizeof(float));

// CPU initializes data
for (int i = 0; i < N; i++) data[i] = 1.0f;

// GPU executes kernel using the SAME pointer
processKernel<<<grid, block>>>(data, N);
cudaDeviceSynchronize();

// CPU reads back results directly
printf("Result: %f\n", data[0]);
cudaFree(data);

Hardware Page Faulting & The Page Migration Engine

Prior to the Pascal architecture, Unified Memory required bulk synchronization before kernel launches. Modern architectures (Pascal through Blackwell) utilize Hardware Page Faulting:

1. On-Demand Migration: When the GPU accesses a Unified Memory address that physically resides in host system RAM, the GPU Memory Management Unit (MMU) raises a hardware page fault.
2. Page Migration Engine: The GPU's hardware Page Migration Engine catches the fault, halts the requesting warp, and initiates a Direct Memory Access (DMA) transfer to migrate the enclosing 4 KB or 2 MB virtual memory page across PCIe or NVLink into GPU HBM.
3. Resumption: Once the page is resident in GPU HBM and page tables are updated, the warp resumes execution transparently.
4. Overcoming GPU Memory Limits (Oversubscription): Unified Memory allows GPUs to process datasets that exceed physical GPU VRAM capacity by automatically evicting cold pages back to host RAM.

Optimization Tip (cudaMemPrefetchAsync): While page faulting simplifies programming, runtime page fault stalls can degrade training performance. Developers use cudaMemPrefetchAsync(ptr, size, device, stream) to proactively stream pages to GPU HBM before kernel execution begins.

Pinned (Page-Locked) Host Memory

Standard host memory allocated via C malloc() or C++ new is pageable: the host operating system kernel may move or swap these pages to disk at any time without warning.

• The GPU DMA engine cannot safely transfer data directly from pageable host memory because physical memory addresses could shift during the transfer.
• When performing a standard cudaMemcpy() from pageable memory, the CUDA driver must first allocate an internal pinned "staging buffer", copy the data from pageable RAM to the staging buffer using the CPU, and then issue the DMA transfer to the GPU.
• Pinned Memory (cudaHostAlloc / cudaMallocHost): Locks the host physical pages in RAM, preventing the OS from paging them out. This enables the GPU DMA engine to access host memory directly (Zero-Copy access) and achieves maximum bidirectional PCIe Gen 5 transfer rates (up to 64 GB/s).