GPU Shared Memory Optimization

GPU Shared Memory Optimization is the critical CUDA/GPU programming technique of using on-chip shared memory (32-228 KB per Streaming Multiprocessor) as a programmer-managed cache to reduce global memory accesses — where properly tiled algorithms using shared memory achieve 5-50x speedup over naive global memory implementations because shared memory provides ~20 cycle latency and ~100 TB/s aggregate bandwidth compared to global memory's ~400 cycle latency and ~2-8 TB/s bandwidth.

Shared Memory Architecture

• Location: On-chip SRAM within each SM/CU, shared among all threads in a thread block/workgroup.
• Size: 48-228 KB per SM (configurable split with L1 cache on NVIDIA GPUs). Ampere: up to 164 KB. Hopper: up to 228 KB.
• Bandwidth: 128 bytes per clock per bank. With 32 banks operating at ~1.5 GHz: ~6 TB/s per SM.
• Latency: ~20-30 cycles. Comparable to L1 cache, 10-20x faster than global memory.

Bank Conflicts

Shared memory is organized into 32 banks (NVIDIA). Consecutive 4-byte words map to consecutive banks. If multiple threads in a warp access different addresses in the same bank in the same cycle, the accesses serialize (bank conflict):
- No conflict: Each thread accesses a different bank. Full bandwidth.
- 2-way conflict: Two threads hit the same bank. Half bandwidth.
- 32-way conflict: All threads hit the same bank. 1/32 bandwidth (serial access).

Common conflict patterns:
- Stride-32 access: threads access every 32nd word — all map to the same bank. Worst case.
- Fix: Pad the shared memory array by one element per row: `__shared__ float tile[32][33];` — the extra column shifts each row's bank mapping, eliminating conflicts.

Tiling Pattern

The canonical optimization pattern for matrix operations:
1. Load tile: Threads cooperatively load a tile of input data from global memory into shared memory (coalesced global reads).
2. __syncthreads(): Barrier ensures all threads have completed loading.
3. Compute: Threads read from shared memory (fast, reusable) to compute their outputs. Each element loaded once from global memory but read multiple times from shared memory.
4. __syncthreads(): Barrier before the next tile load (prevent overwriting data still in use).
5. Repeat: Iterate over tiles until the full input is processed.

GEMM Example

Naive GEMM: each element of C reads an entire row of A and column of B from global memory — N³ global reads for an N×N matrix multiply. Tiled GEMM with shared memory: load a TILE_SIZE × TILE_SIZE block of A and B into shared memory, compute partial products, iterate over tiles. Global memory reads drop from N³ to N³/TILE_SIZE — a 16-32x reduction for typical tile sizes.

GPU Shared Memory is the key lever that transforms memory-bound GPU kernels into compute-bound ones — enabling the data reuse patterns that are essential to achieve a significant fraction of the GPU's peak computational throughput.