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.**