GPU Glossary
/
Performance·2 min read

Bank Conflict

When threads in a warp access different addresses in the same shared memory bank at once, forcing their accesses to be serialized.

A bank conflict happens when several threads in a warp simultaneously request different addresses within the same bank of shared memory. The conflicting accesses are serialized, cutting memory throughput by an integer factor and preventing the kernel from saturating memory bandwidth.

Threads accessing distinct shared memory banks are served in parallel (left). Threads accessing different addresses in the same bank are serialized (right).
Threads accessing distinct shared memory banks are served in parallel (left). Threads accessing different addresses in the same bank are serialized (right).

How shared memory banks work#

Like other SRAM cache memories, the shared memory in a Streaming Multiprocessor is split into groups called banks. Banks can be accessed simultaneously, which raises bandwidth.

On GPUs there are 32 banks, each 4 bytes wide. Consecutive 32-bit words (not 64-bit; GPUs were designed around 32-bit floats and integers) map to consecutive banks:

Address:  0x00  0x04  0x08  0x0C  0x10  0x14  0x18  0x1C  ...  0x7C
Bank:       0     1     2     3     4     5     6     7   ...    31

Address:  0x80  0x84  0x88  0x8C  0x90  0x94  0x98  0x9C  ...  0xFC
Bank:       0     1     2     3     4     5     6     7   ...    31

Addresses 32 × 4 = 128 bytes apart map to the same bank. Shared memory is roughly kilobyte scale, so many addresses share each bank.

Conflict-free access#

When threads read sequential elements, each thread in the warp hits a different bank:

__shared__ float data[1024];  // array in shared memory

// all 32 threads access consecutive elements of data
int tid = threadIdx.x;
float value = data[tid];  // address LSBs: 0x00, 0x04, 0x08, ...

All 32 accesses complete in one memory transaction, as on the left of the figure.

A 32-way conflict#

Now suppose the threads read a column of a row-major shared memory array with 32 elements per row:

float value = data[tid * 32];  // address LSBs: 0x000, 0x080, 0x100 ...
// recall: floats are 4 bytes wide

Every access lands in Bank 0, as on the right of the figure. The accesses are serialized, a 32× increase in latency, from on the order of ten cycles to on the order of hundreds. Transposing the shared memory array fixes this conflict. For more techniques, see the Introduction to CUDA Programming and Performance Optimization talk from GTC 2024.

Key idea Only different addresses in the same bank conflict. If threads read the exact same address, the data can be multicast or broadcast and no conflict need occur.