Memory hierarchy: the table that sets your performance ceiling
Latency and bandwidth of registers, shared memory, L1/L2, and HBM, plus arithmetic intensity as a diagnostic
- Remember the order-of-magnitude gaps in latency and bandwidth across the hierarchy
- Use arithmetic intensity to tell whether a kernel is memory-bound or compute-bound
- Recognize the symptoms and cost of register spilling
- See why most real kernels are memory-bound
If you remember one thing about CUDA performance, remember this table. GPU compute has been in surplus for years. The scarce resource is feeding data into the ALUs. An A100 has 19.5 TFLOPS of FP32, but only 1.55 TB/s of device-memory bandwidth. For every float (4 bytes) you read from device memory, you need 50 floating-point operations to keep the ALUs busy. Most kernels never come close to that ratio.
| Level | Scope | Latency (cycles) | Bandwidth class | Capacity |
|---|---|---|---|---|
| Registers | Private per thread | ~1 | ~20 TB/s | 256 KB per SM |
| Shared memory / L1 | Shared within a block | ~20-30 | ~10 TB/s | up to 164 KB per SM |
| L2 cache | Device-wide | ~200 | ~4 TB/s | 40-50 MB |
| Global memory (HBM) | Device + host | ~400-800 | 1-3 TB/s | 16-80 GB |
| Constant memory | Device-wide, read-only | ~1 (cache hit) | Broadcast-optimized | 64 KB |
| Local memory | Private per thread (physically device memory) | ~400-800 | Same as global | Limited by device memory |
1nvcc -O3 -arch=sm_80 -Xptxas -v -c kernel.cu2 3# Typical output:4# ptxas info : Used 38 registers, 8192 bytes smem, 380 bytes cmem[0]5# ↑ registers per thread ↑ shared mem per block6#7# This line means you spilled. Fix it:8# ptxas info : 24 bytes spill stores, 24 bytes spill loadsArithmetic intensity: name the bottleneck before you optimize
Arithmetic intensity = floating-point operations / bytes accessed (FLOP/Byte). Compare it to the hardware's balance point and you know where the kernel is stuck. Balance point = peak compute / peak bandwidth. For A100 FP32 that is about 19500 / 1555 ≈ 12.5 FLOP/Byte.
| Operation | Arithmetic intensity | Bottleneck | What to optimize |
|---|---|---|---|
Vector add c = a + b | 1 FLOP / 12 B ≈ 0.08 | Severely memory-bound | Tune the access pattern. Extra FLOPS will not help |
SAXPY y = a*x + y | 2 FLOP / 12 B ≈ 0.17 | Severely memory-bound | Same: the goal is to saturate bandwidth |
| Naive matmul (no tiling) | 2 FLOP / 8 B = 0.25 | Memory-bound | Reuse data in shared memory |
| Tiled matmul (tile=32) | ≈ 8 | Nearly balanced | Grow the tile, add register blocking |
| Large matmul (register blocking) | > 50 | Compute-bound | Use Tensor Cores, push FLOPS |
How each level is declared
1__constant__ float c_filter[256]; // constant memory: host writes with cudaMemcpyToSymbol2 3__global__ void demo(const float* __restrict__ g_in, // global memory4 float* g_out) {5 __shared__ float s_tile[32][33]; // static shared memory (33 pads out bank conflicts)6 extern __shared__ float s_dyn[]; // dynamic shared memory, size from <<<,,bytes>>>7 8 float acc = 0.0f; // register9 float buf[8]; // compile-time-unrolled index → register; otherwise → local memory (slow)10 11 #pragma unroll // force unroll so buf stays in registers12 for (int i = 0; i < 8; ++i) buf[i] = g_in[i];13 14 for (int i = 0; i < 8; ++i) acc += buf[i] * c_filter[i];15 g_out[threadIdx.x] = acc;16}Shared memory and L1 are the same SRAM physically, split by a ratio. The driver picks a default, but you can state a preference. A kernel that leans on shared memory (tiled matmul, for example) should ask for more.
1// Prefer: give shared memory as much as possible2cudaFuncSetCacheConfig(myKernel, cudaFuncCachePreferShared);3 4// Volta and later: set a precise per-block shared-memory cap (bytes)5cudaFuncSetAttribute(myKernel,6 cudaFuncAttributeMaxDynamicSharedMemorySize,7 96 * 1024);Self-check
Answer in your head first, then open the solution. If you cannot, reread that section.