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.