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.