What is a bank conflict?

A bank conflict happens when multiple threads in the same warp simultaneously access different addresses that happen to land in the same bank of shared memory. Where memory coalescing is about getting the most out of global memory accesses, bank conflicts are the equivalent efficiency concern for shared memory, and the fix looks nothing alike.

When a conflict occurs, the hardware has no choice but to serialize the colliding accesses, handling them one at a time instead of in parallel. That serialization costs memory throughput in whole-number multiples, so a two-way conflict roughly doubles the time an access takes, which keeps the Streaming Multiprocessor from ever saturating its shared memory bandwidth.

Shared memory, like other SRAM-based caches, is physically organized into banks, and banks can be read or written simultaneously, which is exactly what gives shared memory its bandwidth advantage over global memory in the first place. On current GPUs there are 32 banks, each 4 bytes wide, and consecutive 32-bit words map to consecutive banks. Addresses 128 bytes apart (32 banks times 4 bytes) land back on the same bank, and because a shared memory array is typically only kilobytes in size, it's common for several addresses in one array to alias onto a single bank.

Accessing shared memory sequentially avoids the problem entirely:

__shared__ float data[1024];

int tid = threadIdx.x;
float value = data[tid]; // addresses 0x00, 0x04, 0x08, ... one per bank

Every thread in the warp lands on a different bank here, so all 32 reads complete as a single transaction.

Trouble starts with a stride, for instance walking down a column of a row-major array with 32 elements per row:

float value = data[tid * 32]; // addresses 0x000, 0x080, 0x100, ...

Every thread now lands on bank 0, so all 32 accesses serialize, turning what should be one fast transaction into roughly 32 times the latency: on the order of hundreds of cycles instead of ten.

The usual fix is to change the shared memory layout so the stride no longer aligns with the bank count, commonly by transposing the array or padding each row with one extra element so successive rows no longer start on the same bank. One case never conflicts regardless of stride: if every thread in the collision reads the exact same address, the hardware broadcasts that single value to all of them instead of serializing.

Bank conflicts are invisible from the outside: a kernel that's conflict-bound still returns correct results, just slower than its arithmetic intensity would predict, and a profiler is usually the only way to catch one. Aquanode's GPU metrics page shows achieved versus theoretical shared memory throughput for jobs on rented H100 or B200 capacity, which tends to be the fastest way to notice a conflict is costing you anything at all.