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

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.