CUDA Shared Memory Optimization

CUDA Shared Memory Optimization is the technique of utilizing the fast, programmer-managed on-chip memory (shared memory) within each GPU streaming multiprocessor to cache frequently accessed data, enable inter-thread communication, and reduce costly global memory accesses — achieving 10-100× lower latency than global memory when properly utilized.

Shared Memory Architecture:
- Physical Implementation: shared memory is a low-latency SRAM bank array integrated within each SM — 48-164 KB per SM depending on GPU generation, configurable split with L1 cache
- Access Latency: ~20-30 cycles compared to ~200-800 cycles for global memory (L2 miss) — bandwidth of ~128 bytes/cycle per SM when accessed without bank conflicts
- Bank Organization: shared memory divided into 32 banks (one per warp lane), each 4 bytes wide — consecutive 4-byte addresses map to consecutive banks, enabling conflict-free access when each thread accesses a different bank
- Lifetime and Scope: allocated per thread block, accessible by all threads in the block, deallocated when the block completes — persists across all kernel phases within the block lifetime

Bank Conflicts:
- N-Way Conflict: when N threads in a warp access different addresses in the same bank, accesses are serialized into N sequential transactions — worst case 32-way conflict reduces bandwidth to 1/32
- Broadcast: when multiple threads read the same address in the same bank, hardware broadcasts the value to all requesting threads at no additional cost — reads are free from conflicts when all threads read the same word
- Padding Technique: adding one dummy element per row in a 2D shared memory array shifts column accesses to different banks — e.g., float smem[32][33] instead of float smem[32][32] eliminates column-access conflicts
- Access Pattern Analysis: NVIDIA Nsight Compute reports shared memory bank conflicts per warp — target zero conflicts for performance-critical kernels

Tiling Patterns:
- Matrix Multiply Tiling: load tiles of input matrices A and B into shared memory cooperatively, then compute partial results from shared memory — reduces global memory accesses by factor of tile_size (typically 16-32×)
- Stencil Computations: load neighborhood (halo cells) into shared memory, then each thread reads its complete stencil from shared memory — critical for image processing and PDE solvers where each output depends on neighboring inputs
- Cooperative Loading: threads in a block collectively load data from global memory to shared memory using coalesced access patterns — __syncthreads() barrier ensures all data is loaded before any thread reads from shared memory
- Double Buffering: overlap loading next tile from global memory with computation on current tile — two shared memory buffers alternate between load and compute phases

CUDA shared memory optimization is one of the most impactful GPU programming techniques — properly tiled algorithms with conflict-free shared memory access can approach the theoretical compute throughput of the GPU, which is impossible when limited by global memory bandwidth.