Global, shared, constant, texture, and register memory — access patterns, bank conflicts, coalescing rules, and when to use each level of the GPU memory hierarchy.
This tutorial maps out every level of the CUDA memory hierarchy — from the fastest registers to the slowest global DRAM — and teaches you when and how to use each.
Tutorials 01–03. You should be comfortable with kernel launches, the thread/block/grid hierarchy, and basic host–device data transfers.
GPU memory is organised in a hierarchy trading off size for speed. Data that lives closer to the compute units is faster to access but far more limited in capacity.
| Memory Level | Latency | Bandwidth | Scope |
|---|---|---|---|
| Registers | ~1 cycle | ~8 TB/s (aggregate) | Per-thread |
| Shared / L1 | ~5 cycles | ~12 TB/s (aggregate) | Per-block |
| L2 Cache | ~30 cycles | ~4 TB/s | Device-wide |
| Constant Cache | ~5 cycles (hit) | Broadcast to warp | Device-wide, read-only |
| Global (DRAM) | ~400–800 cycles | ~1–3 TB/s | Device-wide |
Global memory access is 100–800× slower than a register read. The entire art of CUDA optimisation is about moving data up the hierarchy — from global to shared to registers — before computing on it.
Global memory is the largest and slowest memory on the GPU. It is allocated on the host with cudaMalloc and persists for the lifetime of the allocation. Every thread in every block can read and write it.
float *d_data;
size_t bytes = N * sizeof(float);
cudaMalloc((void**)&d_data, bytes); // allocate on device
cudaMemcpy(d_data, h_data, bytes,
cudaMemcpyHostToDevice); // copy host → device
myKernel<<<grid, block>>>(d_data, N); // launch kernel
cudaMemcpy(h_data, d_data, bytes,
cudaMemcpyDeviceToHost); // copy device → host
cudaFree(d_data); // release memoryWhen all 32 threads in a warp access consecutive 4-byte addresses, the hardware combines them into a single 128-byte transaction. This is coalesced access — the fastest way to read global memory.
// COALESCED: thread i reads element i
float val = data[threadIdx.x]; // stride-1 access → fast
// STRIDED: thread i reads element i * stride
float val = data[threadIdx.x * 32]; // stride-32 access → slow
// RANDOM: thread i reads a random index
float val = data[indices[threadIdx.x]]; // scatter access → slowestAlways have consecutive threads access consecutive memory addresses. If your data layout forces strided access, consider transposing it or using shared memory as a staging buffer.
Shared memory is on-chip SRAM that is private to each thread block. It is roughly 100× faster than global memory and is the primary tool for inter-thread communication within a block.
__global__ void myKernel(float *data) {
// Static allocation — size known at compile time
__shared__ float tile[256];
// Dynamic allocation — size set at launch
extern __shared__ float dynTile[];
tile[threadIdx.x] = data[blockIdx.x * blockDim.x + threadIdx.x];
__syncthreads(); // barrier: all threads must finish writing
// Now all threads can safely read any element of tile[]
}Each block gets its own copy. Threads in different blocks cannot share data through shared memory.
Located on the SM die alongside compute units. ~5 cycle latency vs ~400+ cycles for global memory.
48–228 KB per SM (architecture-dependent). Shared across all active blocks on that SM.
Shared memory is divided into 32 banks (one per warp lane). Each bank can serve one address per cycle. When two threads in a warp access different addresses in the same bank, the accesses are serialised — this is a bank conflict.
Use stride-1 access patterns (consecutive threads access consecutive words). If you need a stride that is a power of 2, add 1 element of padding to each row: __shared__ float tile[32][32+1];
The tiling pattern is the most important shared memory technique. The idea: load a tile of data from slow global memory into fast shared memory, synchronise, then compute from the shared copy.
A 1D stencil computes each output element from a neighbourhood of input elements. This example uses a radius-3 stencil — each output depends on 7 input values. The halo cells are the extra elements needed at the tile boundaries.
#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>
#define RADIUS 3
#define BLOCK_SIZE 256
// ── 1D stencil kernel using shared memory with halo cells ──
__global__ void stencil1D(const float *in, float *out, int n) {
// Shared tile: BLOCK_SIZE elements + RADIUS on each side
__shared__ float tile[BLOCK_SIZE + 2 * RADIUS];
int gIdx = blockIdx.x * blockDim.x + threadIdx.x;
int lIdx = threadIdx.x + RADIUS; // offset for halo
// Load the main body of the tile
if (gIdx < n)
tile[lIdx] = in[gIdx];
else
tile[lIdx] = 0.0f;
// Load left halo
if (threadIdx.x < RADIUS) {
int haloIdx = gIdx - RADIUS;
tile[threadIdx.x] = (haloIdx >= 0) ? in[haloIdx] : 0.0f;
}
// Load right halo
if (threadIdx.x >= BLOCK_SIZE - RADIUS) {
int haloIdx = gIdx + RADIUS;
tile[lIdx + RADIUS] = (haloIdx < n) ? in[haloIdx] : 0.0f;
}
__syncthreads(); // all data in shared memory now
// Compute stencil: average of 2*RADIUS+1 neighbours
if (gIdx < n) {
float sum = 0.0f;
for (int offset = -RADIUS; offset <= RADIUS; offset++)
sum += tile[lIdx + offset];
out[gIdx] = sum / (2 * RADIUS + 1);
}
}
// ── Host code ──
int main() {
const int N = 1024;
size_t bytes = N * sizeof(float);
// Allocate host memory
float *h_in = (float*)malloc(bytes);
float *h_out = (float*)malloc(bytes);
// Initialise input: simple ramp
for (int i = 0; i < N; i++)
h_in[i] = (float)i;
// Allocate device memory
float *d_in, *d_out;
cudaMalloc(&d_in, bytes);
cudaMalloc(&d_out, bytes);
// Copy input to device
cudaMemcpy(d_in, h_in, bytes, cudaMemcpyHostToDevice);
// Launch kernel
int gridSize = (N + BLOCK_SIZE - 1) / BLOCK_SIZE;
stencil1D<<<gridSize, BLOCK_SIZE>>>(d_in, d_out, N);
// Copy result back
cudaMemcpy(h_out, d_out, bytes, cudaMemcpyDeviceToHost);
// Verify a few values
printf("Stencil results (radius=%d):\n", RADIUS);
for (int i = RADIUS; i < RADIUS + 5; i++)
printf(" out[%d] = %.2f (expected %.2f)\n",
i, h_out[i], (float)i);
// Cleanup
cudaFree(d_in);
cudaFree(d_out);
free(h_in);
free(h_out);
return 0;
}Without shared memory, each thread would read 7 values from global memory. With 256 threads, that is 1,792 global reads. With tiling, we load only 256 + 6 = 262 values from global memory into shared, then each thread reads its 7 neighbours from fast shared memory. 6.8× fewer global transactions.
Constant memory is a 64 KB read-only region that is cached on-chip. When all threads in a warp read the same address, the value is broadcast to all 32 threads in a single cycle — perfect for lookup tables, filter coefficients, and configuration parameters.
// Declare at file scope (not inside a function)
__constant__ float coeffs[7];
__global__ void applyFilter(const float *in, float *out, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) {
float sum = 0.0f;
for (int k = 0; k < 7; k++)
sum += coeffs[k] * in[i + k]; // all threads read same coeffs[k]
out[i] = sum;
}
}
int main() {
float h_coeffs[7] = {0.1f, 0.15f, 0.2f, 0.1f, 0.2f, 0.15f, 0.1f};
// Copy to constant memory — not cudaMemcpy!
cudaMemcpyToSymbol(coeffs, h_coeffs, 7 * sizeof(float));
// ... allocate, launch, etc.
}The constant cache can serve one address per cycle. If all 32 threads request the same address, the value is broadcast to the entire warp for free. If all 32 request different addresses, it takes 32 serialised cycles — worse than a global memory read.
Registers are the fastest storage on the GPU — zero additional latency beyond the instruction itself. Every automatic (local) variable in your kernel lives in a register, unless the compiler runs out of space.
__global__ void compute(float *data, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x; // register
float a = data[idx]; // register
float b = a * a; // register
float c = b + 1.0f; // register
data[idx] = c;
// Large arrays CANNOT fit in registers:
float bigArray[1024]; // SPILLS to local memory (slow!)
}When a thread needs more registers than available (typically 255 per thread), the compiler spills variables to local memory. Despite the name, local memory is actually stored in global DRAM — it is private to each thread but as slow as global memory.
// Method 1: Compiler flag
// nvcc --maxrregcount=32 kernel.cu
// Method 2: Per-kernel attribute
__global__ __launch_bounds__(256, 4) // max 256 threads, min 4 blocks/SM
void myKernel(...) {
// compiler will limit registers to allow 4 blocks of 256 threads
}Compile with nvcc --ptxas-options=-v to see per-kernel register, shared memory, and local memory usage. Target fewer than 32–64 registers per thread for good occupancy.
Unified Memory provides a single pointer that is accessible from both the CPU and GPU. The CUDA runtime automatically migrates pages between host and device memory as needed — no explicit cudaMemcpy calls required.
float *data;
cudaMallocManaged(&data, N * sizeof(float));
// Initialise on the CPU — no cudaMemcpy needed
for (int i = 0; i < N; i++)
data[i] = (float)i;
// Launch kernel — pages migrate to GPU automatically
myKernel<<<grid, block>>>(data, N);
cudaDeviceSynchronize();
// Read results on CPU — pages migrate back automatically
printf("Result: %f\n", data[0]);
cudaFree(data); // single free for both host and devicecudaMemcpy still winscudaMemPrefetchAsync to pre-migrateNVIDIA's Grace Hopper architecture features unified CPU+GPU memory with NVLink-C2C, where CPU and GPU share the same physical memory pool. On such systems, cudaMallocManaged incurs virtually zero migration overhead — the memory is already coherent. This makes unified memory the natural programming model for Grace Hopper and DGX Spark systems.
Quick reference for choosing the right memory type for your data.
| Memory Type | Location | Scope | Speed | Size | Typical Use |
|---|---|---|---|---|---|
| Registers | On-chip | Per-thread | ~1 cycle | ~255 per thread | Local variables, accumulators |
| Shared | On-chip | Per-block | ~5 cycles | 48–228 KB/SM | Tiling, reductions, inter-thread comms |
| L1 Cache | On-chip | Per-SM | ~30 cycles | Shared with smem | Automatic caching of global reads |
| L2 Cache | On-chip | Device | ~30 cycles | 1.5–126 MB | Automatic caching layer |
| Constant | DRAM + cache | Device (R/O) | ~5 cycles (hit) | 64 KB | Uniform reads: weights, LUTs |
| Texture | DRAM + cache | Device (R/O) | ~5 cycles (hit) | Varies | 2D spatial locality, interpolation |
| Local | DRAM | Per-thread | ~400 cycles | Up to 512 KB | Register spills, large arrays |
| Global | DRAM | Device | ~400 cycles | 8–80 GB | Main data store, I/O buffers |
cudaMallocPitch for 2D arrays__syncthreads() after writingIs the data read-only and small (≤64 KB)? → Constant memory. Is it reused by threads in the same block? → Shared memory. Does it fit in registers? → Keep it local. Otherwise → Global memory with coalesced access.
The single most impactful optimisation in CUDA is reducing global memory traffic. Use shared memory tiling whenever threads reuse data.
Coalesced global reads and conflict-free shared memory access can make the difference between 10% and 90% of peak bandwidth.
Matrix Multiplication — apply the tiling pattern to the classic GEMM problem. We'll build from a naive O(N³) kernel to an optimised tiled version, measuring the impact of shared memory at every step.