GPU Shared Memory

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.

GPU Shared Memory

Take gpu shared memory further