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.