bank conflicts
**Bank Conflicts** are a **GPU performance bottleneck that occurs when multiple threads in a warp simultaneously access different addresses within the same shared memory bank** — causing memory accesses to be serialized rather than executed in parallel, potentially reducing shared memory throughput by up to 32× in the worst case, making bank conflict avoidance one of the most critical optimizations for high-performance CUDA kernels used in deep learning inference and training.
**What Are Bank Conflicts?**
- **Definition**: GPU shared memory is divided into 32 banks (on NVIDIA GPUs), each 4 bytes wide, with consecutive 4-byte words mapped to consecutive banks in a round-robin pattern — a bank conflict occurs when two or more threads in the same warp access different addresses that map to the same bank, forcing those accesses to be serialized.
- **Shared Memory Banks**: Bank 0 holds addresses 0-3, Bank 1 holds addresses 4-7, ..., Bank 31 holds addresses 124-127, then Bank 0 holds addresses 128-131, and so on — addresses that are 128 bytes apart (32 banks × 4 bytes) map to the same bank.
- **Conflict Example**: Thread 0 accesses address 0 (Bank 0) and Thread 1 accesses address 128 (also Bank 0) — both addresses are in Bank 0, so the accesses are serialized into two sequential transactions instead of one parallel transaction.
- **Broadcast Exception**: If all threads in a warp read the exact same address, there is no conflict — the hardware broadcasts the single read to all threads in one transaction.
**Bank Conflict Severity**
| Scenario | Threads Conflicting | Throughput Impact | Example |
|----------|-------------------|------------------|---------|
| No conflict | 0 | 100% (optimal) | Stride-1 access pattern |
| 2-way conflict | 2 per bank | 50% | Stride-2 access |
| 4-way conflict | 4 per bank | 25% | Stride-8 access |
| 32-way conflict | All 32 | 3% (worst case) | All threads same bank, different addr |
| Broadcast | All same address | 100% | All threads read same value |
**Common Causes in Deep Learning**
- **Matrix Transpose**: Naive shared memory transpose with stride equal to the tile width causes 32-way bank conflicts — the classic CUDA optimization example.
- **Reduction Operations**: Parallel reductions where threads access shared memory with power-of-2 strides create systematic bank conflicts.
- **Attention Kernels**: Custom attention implementations that load Q, K, V tiles into shared memory can suffer bank conflicts if tile dimensions align with bank boundaries.
**Avoidance Techniques**
- **Padding**: Add 1 element of padding per row in shared memory arrays — `__shared__ float tile[32][33]` instead of `[32][32]` shifts each row by one bank, eliminating stride-32 conflicts.
- **Access Pattern Redesign**: Rearrange data layout so that threads in a warp access consecutive banks — stride-1 access patterns are always conflict-free.
- **Swizzling**: XOR-based address swizzling remaps thread-to-bank assignments — used in CUTLASS and cuBLAS for high-performance matrix multiplication tiles.
**Bank conflicts are the hidden performance killer in GPU shared memory access** — causing up to 32× throughput reduction when multiple warp threads hit the same memory bank, making conflict-free access patterns through padding, swizzling, and layout optimization essential for achieving peak performance in CUDA kernels for deep learning.