GPU Shared Memory Optimization — using the fast, programmer-managed on-chip memory (shared memory / SMEM) within each GPU Streaming Multiprocessor (SM) to drastically reduce global memory accesses.

Shared Memory Properties

• Location: On-chip SRAM within each SM
• Size: 48–228 KB per SM (configurable vs. L1 cache)
• Latency: ~20-30 cycles (vs. ~400 cycles for global memory)
• Bandwidth: ~10 TB/s aggregate (vs. ~2 TB/s for HBM)
• Scope: Shared among all threads in a thread block

Classic Pattern: Tiled Matrix Multiply

1. Load tile of A from global → shared memory
2. Load tile of B from global → shared memory
3. __syncthreads()  // All threads in block sync
4. Compute partial result using fast shared memory reads
5. Repeat for next tile

• Without shared memory: Each element read from slow global memory multiple times
• With shared memory: Each element loaded once from global, reused many times from SMEM
• Speedup: 10-20x for matrix multiply

Bank Conflicts

• Shared memory divided into 32 banks
• Threads in a warp accessing different banks → simultaneous (fast)
• Multiple threads accessing same bank → serialized (bank conflict, slow)
• Solution: Pad shared memory arrays to avoid conflict patterns

Best Practices

• Use shared memory for data reused across threads in a block
• Always __syncthreads() between write and read phases
• Avoid bank conflicts by careful indexing

Shared memory is the #1 optimization technique in CUDA programming — mastering it is what separates a 10x kernel from a 100x kernel.