GPU Memory Hierarchy in CUDA

GPU performance is often limited less by "how many cores exist" and more by how fast data can reach those cores. CUDA exposes several memory spaces, each with a different location, scope, lifetime, and performance profile.^[1]

This note focuses on the CUDA programmer's view: registers, shared memory, local memory, global memory, and constant memory.

Table of contents

• On-chip vs off-chip memory
• CUDA memory spaces
• 1. Registers
• 2. Shared memory
• 3. Local memory
• 4. Global memory
• 5. Constant memory
• Comparison table
• Two useful mental models
• References

On-chip vs off-chip memory#

_{Figure 1. A simplified on-chip/off-chip memory hierarchy for GPU systems.}

On-chip memory is physically inside the GPU chip. It is smaller and more expensive in area, but much faster to access. Registers, shared memory, and L1 cache are on-chip resources.

Off-chip memory is outside the GPU chip, usually in the GPU's device DRAM/HBM/GDDR memory package. It is much larger, but access latency is higher. Global memory and most local memory traffic live here, although they can be cached.

The basic engineering tradeoff is simple:

$$closer\ memory = lower\ latency + smaller\ capacity$$

That is why high-performance CUDA kernels try to reuse data in registers and shared memory before going back to global memory.

CUDA memory spaces#

_{Figure 2. CUDA memory spaces organized by location, visibility, and programmer control.^[1]}

tldr:

• Registers are private to one thread.
• Shared memory is visible to threads in one block.
• Global memory is visible across the grid and can be managed by the host.
• Constant memory is read-only from the device and useful for broadcast-style reads.
• Local memory is private to one thread but physically backed by device memory.

1. Registers#

• Location: on-chip register file in the SM.
• Scope: one thread.
• Access: read/write.
• Lifetime: one thread's execution.
• Programmer control: indirect; the compiler allocates registers.

Registers are the closest (=fastest) memory resource to execution. It also has the smallest capacity. Usually the faster bandwidth it has, the smaller capacity it has. This is not something fundamental, but only using SRAMs in a chip is not that cost-efficient and very expensive. Therefore computer architecture, including GPUs, usually use memory hierarchy and cache to have similar speed as SRAM but large storage like DRAM (much more cost-efficient; hardware-software hack basically).

In CUDA C/C++, ordinary automatic variables often live in registers when the compiler can keep them there.^[1]

If a kernel needs too many registers per thread, the compiler may spill values into local memory, which is much slower to fetch. High register use can also reduce occupancy because each resident thread block consumes part of the SM's register file.

At the PTX level, registers are represented by the .reg state space. PTX registers are not generally addressable like normal memory which is one reason CUDA programmers cannot directly "take the address" of a register as if it were an array in global memory.^[2]

2. Shared memory#

• Location: on-chip, near the SM.
• Scope: one thread block.
• Access: read/write.
• Lifetime: one thread block.
• Declaration: __shared__ in a kernel, or dynamically via launch configuration.

Shared memory is fast(about ~100x faster than global memory) on-chip memory that all threads in the same thread block can access.^[1] It is one of the most important memory for CUDA performance because it lets a block explicitly stage and reuse data.

On many NVIDIA architectures, shared memory and L1 cache use the same physical on-chip resource, often called the unified data cache. The portion reserved for shared memory can be configured per kernel on supported devices.^[3]

Common examples include tiled matrix multiplication, stencil kernels, reductions, and data-layout transformations before Tensor Core operations.

3. Local memory#

• Location: device memory, usually off-chip, cached through GPU caches.
• Scope: one thread.
• Access: read/write.
• Lifetime: one thread.
• Programmer control: mostly compiler-managed.

Local memory is off-chip memory, which means it resides in DRAM. Usually large logical data structures or local arrays are saved here.

From a performance point of view, local memory should usually be treated like global-memory traffic with thread-private semantics.

4. Global memory#

• Location: off-chip device memory.
• Scope: visible to all threads in the grid and accessible through host/device APIs.
• Access: read/write.
• Lifetime: host-controlled; data persists across kernels until freed or overwritten.
• Declaration/allocation: cudaMalloc(), CUDA-managed allocations, or __device__ globals.

Global memory is the main device memory space, often interchangeably called as GPU VRAM or HBM.

Global memory is largest but the slowest. CUDA performance often depends on making global-memory accesses coalesced, reducing unnecessary transfers, and reusing loaded data in registers or shared memory.

5. Constant memory#

• Location: device memory, backed by an on-chip constant cache.
• Scope: visible to all threads in a grid.
• Access: read-only from device code.
• Lifetime: host-controlled.
• Declaration: __constant__, with host copies such as cudaMemcpyToSymbol().

Constant memory is a read-only device memory space. It is useful when many threads read the same values, especially when threads in a warp read the same address and the hardware can broadcast the value efficiently.^[1]

In short, constant memory is a specialized path for small read-only data with favorable access patterns.

Comparison table#

Two useful mental models#

_{Figure 3. Memory hierarchy diagram from Programming Massively Parallel Processors.^[4]}

First, memory is a hierarchy of capacity and latency. Registers and shared memory are precious because they are close to the SM. Global memory is plentiful, but every unnecessary trip to it costs time.

_{Figure 4. CUDA memory hierarchy summary from Modal's GPU glossary.^[5]}

Second, CUDA memory spaces are about visibility as much as speed. Registers are private to a thread. Shared memory is shared by a block. Global and constant memory are visible across the device. Local memory is thread-private but slow because of where it is physically stored.

References#

1. NVIDIA, "CUDA C++ Programming Guide: Writing CUDA SIMT Kernels." Link
2. NVIDIA, "Parallel Thread Execution ISA: Register State Space." Link
3. NVIDIA, "CUDA C++ Programming Guide: Advanced Kernel Programming." Link
4. Hwu, Kirk, and El Hajj, Programming Massively Parallel Processors. Link
5. Modal, "GPU glossary: Memory hierarchy." Link