CUDA Memory Model — infra.bea.ai

A kernel’s performance is almost always determined by how it accesses memory, not how it computes. This chapter maps every CUDA memory type to hardware, explains access rules, and shows how the right memory choice can make a kernel 10–100× faster.

1. Memory Hierarchy Overview

CUDA Memory Hierarchy

CUDA exposes a software memory model that maps to hardware memory:

CUDA memory type	Declaration	Hardware location	Scope	Lifetime	Latency	Bandwidth
Registers	Local variables	SM register file	Per-thread	Kernel	~0 cycles	Highest
Local memory	Spilled registers	HBM (!)	Per-thread	Kernel	~600 cycles	Low
Shared memory	`__shared__`	SM on-chip SRAM	Per-block	Kernel	~30 cycles	~19 TB/s aggregate
L1 cache	Automatic	SM on-chip SRAM	Per-SM	Dynamic	~30 cycles	Shared with smem
L2 cache	Automatic	On-chip	Global	Dynamic	~200 cycles	~9 TB/s
Global memory	`__device__` / `cudaMalloc`	HBM	Global	Application	~600 cycles	3.35 TB/s (H100)
Constant memory	`__constant__`	HBM + cache	Global	Application	~1-600 cycles	Cached
Texture memory	Texture objects	HBM + cache	Global	Application	~1-600 cycles	Spatial locality

⚠️ CRITICAL: “Local memory” is the most misleading name in CUDA. It is NOT local to the SM — it lives in HBM and is as slow as global memory. It’s called “local” because it’s per-thread scope. When your kernel uses too many registers, the compiler spills to local memory — a silent performance killer.

2. Registers: The Fastest Storage

Every thread has its own registers — the fastest storage available.

__global__ void kernel() {
    float a = 1.0f;  // stored in a register
    float b = 2.0f;  // another register
    float c = a + b;  // computed using register ALU, result in register
}

Register Budget

GPU	Registers per SM	Registers per thread (max)
H100	65,536 (256 KB)	255
A100	65,536 (256 KB)	255

The tradeoff: more registers per thread → fewer threads can run simultaneously (lower occupancy). The compiler decides how many registers each kernel uses.

# Check register usage at compile time
nvcc -Xptxas -v kernel.cu
# Output: "Used 32 registers, 0 bytes shared memory"

Register Spilling

When a kernel needs more registers than available, the compiler spills to local memory (HBM):

High register usage → spill to local memory → 600-cycle penalty per access

How to reduce register pressure:

Simplify expressions (fewer temporaries)
Use __launch_bounds__(maxThreadsPerBlock, minBlocksPerSM) to hint the compiler
Trade registers for shared memory (store intermediates explicitly)

3. Shared Memory: Programmer-Managed Cache

Shared memory is fast on-chip SRAM shared by all threads in a block. It’s like an L1 cache that you manage.

Declaration and Usage

__global__ void kernel() {
    // Static allocation (size known at compile time)
    __shared__ float tile[256];

    // Dynamic allocation (size specified at launch)
    extern __shared__ float dynamicTile[];

    // Load data cooperatively
    tile[threadIdx.x] = globalData[blockIdx.x * 256 + threadIdx.x];
    __syncthreads();  // MUST sync before reading others' writes

    // Use shared data (fast — ~30 cycles)
    float val = tile[threadIdx.x] + tile[255 - threadIdx.x];
}

// Launch with dynamic shared memory
kernel<<<grid, block, sharedMemBytes>>>(args);

When to Use Shared Memory

Use case	Why
Data reuse within a block	Load from HBM once, use multiple times from shared memory
Inter-thread communication	Threads in a block can read each other’s data
Tiling for matrix operations	Load a tile, compute, repeat — the basis of fast GEMM
Reduction within a block	Accumulate values from all threads

Bank Conflicts

Shared memory is divided into 32 banks (one per warp lane). If multiple threads in a warp access the same bank (but different addresses), accesses are serialized:

// NO conflict: each thread accesses a different bank
shared[threadIdx.x] = val;  // thread 0→bank 0, thread 1→bank 1, ...

// 2-way conflict: threads 0 and 16 both access bank 0
shared[threadIdx.x * 2] = val;  // thread 0→bank 0, thread 1→bank 2, ..., thread 16→bank 0!

// BROADCAST (no conflict): all threads read the SAME address
float val = shared[0];  // all threads read bank 0, address 0 — broadcast

Bank index = (address / 4 bytes) % 32

Fix: Pad shared memory arrays to shift bank alignment:

__shared__ float tile[32][33];  // 33 instead of 32 — avoids column-access conflicts

Shared Memory vs L1 Cache

On modern GPUs (Ampere+), shared memory and L1 cache share the same physical SRAM. You can configure the split:

// Prefer more shared memory
cudaFuncSetAttribute(kernel, cudaFuncAttributePreferredSharedMemoryCarveout, 100);

// Or set at device level
cudaDeviceSetCacheConfig(cudaFuncCachePreferShared);

H100: 228 KB per SM, configurable between shared memory and L1.

L2 Cache

The L2 cache (50 MB on H100) sits between all SMs and HBM. It’s hardware-managed — you don’t control what gets cached.

L2 matters for:

Inter-block communication: if one block writes data that another block reads shortly after, it may hit L2 instead of going to HBM
Repeated global reads: data accessed by multiple blocks may stay in L2
Residency control (Hopper+): cudaAccessPolicyWindow lets you hint which data should stay in L2

// Hopper: set L2 persistence for frequently accessed data
cudaAccessPolicyWindow window;
window.base_ptr = d_weights;
window.num_bytes = weightSize;
window.hitRatio = 1.0f;
window.hitProp = cudaAccessPropertyPersisting;
window.missProp = cudaAccessPropertyStreaming;
cudaCtxSetAccessPolicyWindow(&window);

4. Global Memory: The Main Storage

Global memory is HBM — large, slow, and the bottleneck for most kernels.

Allocation and Access

float* d_data;
cudaMalloc(&d_data, N * sizeof(float));       // allocate
cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice);  // copy to GPU
cudaFree(d_data);                              // free

Coalesced Access: The #1 Optimization

Coalesced Access

When threads in a warp access consecutive memory addresses, the hardware combines them into a single transaction:

// COALESCED: threads 0-31 access addresses 0-31 (one 128-byte transaction)
float val = data[threadIdx.x];  // ✅ Fast

// STRIDED: threads 0-31 access addresses 0, stride, 2*stride, ...
float val = data[threadIdx.x * stride];  // ❌ Multiple transactions if stride > 1

// RANDOM: each thread accesses a random address
float val = data[indices[threadIdx.x]];  // ❌ Worst case: 32 separate transactions

Coalesced: 1 transaction for 32 threads (128 bytes). Uncoalesced: up to 32 transactions. That’s 32× slower.

Structure of Arrays (SoA) vs Array of Structures (AoS)

// AoS: BAD for GPU (strided access within a warp)
struct Particle { float x, y, z, w; };
Particle particles[N];
// Thread i accesses particles[i].x → stride of 16 bytes between threads

// SoA: GOOD for GPU (coalesced access)
float x[N], y[N], z[N], w[N];
// Thread i accesses x[i] → consecutive addresses → coalesced

Rule: Always prefer SoA layout for GPU data.

5. Constant and Texture Memory

Constant Memory

__constant__ float coefficients[256];  // 64 KB max total

// Set from host
cudaMemcpyToSymbol(coefficients, hostData, 256 * sizeof(float));

Read-only from device, written from host
Cached — very fast when all threads read the same address (broadcast)
Slow when threads read different addresses (serialized)

Best for: model hyperparameters, lookup tables, filter weights.

Texture Memory

Optimized for 2D spatial locality (nearby pixels). Rarely used in AI workloads — shared memory and global memory with coalescing are usually better.

6. Memory Access Patterns Cheat Sheet

Pattern	Memory type	When to use
Each thread computes with its own data	Registers	Always (compiler default)
Threads in a block share data	Shared memory	Tiling, reduction, transpose
All threads read same constant	Constant memory	Hyperparameters, small lookup tables
Sequential global access	Global (coalesced)	Input/output arrays, SoA layout
Random global access	Global (uncoalesced)	Avoid! Restructure data or use shared memory

7. Unified Memory (Managed Memory)

Unified Memory simplifies GPU programming by providing a single address space accessible from both CPU and GPU:

float* data;
cudaMallocManaged(&data, N * sizeof(float));

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

// GPU reads — driver automatically migrates pages
kernel<<<grid, block>>>(data, N);
cudaDeviceSynchronize();

// CPU reads result — driver migrates back
printf("%f\n", data[0]);

cudaFree(data);

How it works: The CUDA driver uses page faults to migrate data between CPU and GPU on demand. When the GPU accesses a page that’s in CPU memory, a page fault triggers migration.

Performance: Page fault migration adds latency. For performance-critical code, use explicit cudaMemcpy or prefetch hints:

cudaMemPrefetchAsync(data, size, deviceId, stream);  // pre-migrate before kernel

When to use:

Prototyping and debugging (simpler code)
Complex data structures with irregular access (hard to predict which data GPU needs)
Oversubscription (data larger than GPU memory — driver pages in/out)

When NOT to use:

Production AI kernels (explicit memory management is faster)
Latency-sensitive inference (page faults add unpredictable latency)

cudaMemcpy Patterns

Direction	Constant	Use
`cudaMemcpyHostToDevice`	CPU → GPU	Send input data before kernel
`cudaMemcpyDeviceToHost`	GPU → CPU	Retrieve results after kernel
`cudaMemcpyDeviceToDevice`	GPU → GPU	Copy between GPU buffers
`cudaMemcpyDefault`	Auto-detect	Works with unified memory

For async transfers (requires pinned host memory):

cudaMallocHost(&h_data, size);  // pinned memory
cudaMemcpyAsync(d_data, h_data, size, cudaMemcpyHostToDevice, stream);

Self-Check

Name all CUDA memory types and their hardware locations
Explain why “local memory” is misleadingly named (it’s in HBM)
Explain register spilling and how to detect it
Write a kernel that uses shared memory with __syncthreads()
Explain bank conflicts and how padding fixes them
Explain coalesced vs uncoalesced access and the performance difference
Explain SoA vs AoS and why SoA is preferred for GPUs
Explain Unified Memory and when to use it vs explicit cudaMemcpy
Name the four cudaMemcpy direction constants