CUDA Memory Hierarchy: Global, Shared, Constant, Local, Registers

On a GPU the word "memory" hides a hierarchy spanning ~3 orders of magnitude in latency. A kernel can see:

• Registers — per-thread, on-chip, ~1 cycle.
• Shared memory — per-block, on-chip, ~20-30 cycles.
• L1/L2 cache — automatic, ~30/200 cycles.
• Constant memory — read-only, broadcast-optimised, cached.
• Global memory — device-wide VRAM, ~400-800 cycles.
• Local memory — per-thread but physically in global memory.

The whole performance-tuning game on GPUs is: keep hot data in fast tiers, touch global memory as little as possible, and make the global accesses you can't avoid look the way the hardware likes.

Global memory is the GPU's VRAM — gigabytes large, accessible by every thread in every block, persists across kernel launches. It's also the slowest tier by a factor of ~20-40 over on-chip memory.

You allocate it from the host with cudaMalloc:

float* d_arr;
cudaMalloc(&d_arr, N * sizeof(float));
// ... kernels read/write d_arr ...
cudaFree(d_arr);

Every __global__ kernel argument that's a pointer typically points here. The catch: every uncached load is hundreds of cycles. A naive kernel that does one global load per arithmetic op will be memory-bound, sitting on its hands waiting for VRAM. The whole rest of this lesson is about avoiding that.

The hardware reads global memory in 128-byte transactions aligned to 128-byte boundaries. If the 32 threads of a warp all access adjacent 4-byte words, the hardware serves them with one transaction. That's coalesced access.

If those 32 threads access scattered addresses (say, each one a different 128-byte page), the hardware issues 32 separate transactions — 32x more memory traffic for the same compute.

// Coalesced: thread i reads element i
a[blockIdx.x * blockDim.x + threadIdx.x];

// Strided: thread i reads element i * stride — bad
a[(blockIdx.x * blockDim.x + threadIdx.x) * stride];

Lay out your data so consecutive threads touch consecutive addresses. This is why GPU code prefers struct-of-arrays over array-of-structs.

Each Streaming Multiprocessor (SM) has a small pool of on-chip memory (~48-100 KB) that you can carve into shared memory for each block. Threads in the same block can read and write it; threads in different blocks can't see each other's.

__global__ void blockSum(float* in, float* out) {
  __shared__ float tile[256];
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  tile[threadIdx.x] = in[i];   // pull from global
  __syncthreads();             // wait for all threads
  // ... now use tile[] as fast scratch ...
}

It's not a hardware cache — you decide what goes in. Used right, it cuts global memory traffic by 10-100x on stencils, matmul tiles, reductions. Used wrong, it just adds code and bank-conflict pain (see next step).

Shared memory is split into 32 banks (one per warp lane). If 32 threads in a warp each hit a different bank, you get full bandwidth. If two or more threads in a warp hit the same bank at different addresses, the hardware serialises them — an n-way bank conflict runs n times slower.

The classic offender is a column-major access into a 32×32 tile: every thread in a warp hits the same column → same bank → 32-way conflict. The textbook fix is padding the tile to 33 columns: __shared__ float tile[32][33];. The wasted column shifts each row by one bank, so column reads spread across all 32 banks again.

Constant memory is a small (64 KB) device region declared at file scope:

__constant__ float kernel[16];   // host writes once, device reads many

__global__ void conv(/*...*/) {
  float w = kernel[0];          // hits constant cache
}

The host writes it once with cudaMemcpyToSymbol. Kernels read it. Why bother? When every thread in a warp reads the same address, the constant cache broadcasts the value in a single cycle — almost free. Perfect for filter weights, simulation parameters, lookup tables. Bad for anything where different threads need different elements — broadcast collapses to serial reads.

Plain local variables in a kernel live in registers — the fastest storage on the chip. Each SM has a fixed register file (e.g. 65,536 32-bit registers); the compiler divides it among the threads resident on the SM.

The trap: if your kernel uses too many registers per thread, or you index into a local array with a runtime index, the compiler spills the array to local memory — which sounds local but is physically global memory. Latency jumps from 1 cycle to ~500.

__global__ void bad() {
  float tmp[64];       // probably spilled to global!
  int i = threadIdx.x % 64;
  tmp[i] = ...;        // dynamic index = no register allocation
}

Check nvcc --ptxas-options=-v to see register usage and any local memory spills.

The optimisation playbook flows directly from this table: keep hot data in registers, use shared memory as a manual cache for block-local reuse, keep global accesses coalesced and minimal.

Where each space lives, and roughly how fast.