CUDA Shared-Memory Bank Conflicts: what to measure and why

The mechanism

What cuda shared-memory bank conflicts means

For devices covered by current CUDA best-practice guidance, successive 32-bit words map to successive banks and a warp contains 32 threads. A row-major 32×32 tile therefore has a row stride of 32 words. If lanes in a warp read down one column, those addresses can map back to the same bank. Adding one padding element changes the stride to 33 words and rotates successive accesses across banks. Multiple lanes reading the exact same address are a broadcast case, not the same as requesting different conflicting addresses.

• Bank conflicts are evaluated for the shared-memory addresses requested by one warp instruction.
• A conflict serializes different addresses that map to the same bank.
• Broadcasting one shared-memory address to multiple lanes is handled specially.
• Padding is effective when it changes a problematic power-of-two stride, but it is not a universal fix.

A minimal pattern

Change one thing at a time

The snippets isolate the address or launch pattern being discussed. A real optimization still has to preserve output and be measured in the complete kernel.

__shared__ float tile[32][32];

// Lanes varying threadIdx.x read with a stride of 32 words.
float value = tile[threadIdx.x][threadIdx.y];

__shared__ float tile[32][33];

// A 33-word stride rotates successive rows across banks.
float value = tile[threadIdx.x][threadIdx.y];

Profiler evidence

What to measure

Compare the same input and access semantics while measuring kernel duration and shared-memory conflict or wavefront metrics available for the GPU. Use source or SASS correlation to identify the exact load or store. Confirm that the padded array does not push shared-memory use across a residency threshold that changes occupancy.

Reason about the result

How to interpret the measurements

If padding removes conflict serialization, the relevant shared-memory instruction should require fewer service steps and the kernel may run faster. The total speedup depends on how much time the kernel spent on that access. Padding also increases shared-memory allocation, so a conflict reduction can trade against block residency in resource-constrained kernels.

Avoid false conclusions

Common mistakes

• Calling a broadcast a bank conflict when all participating lanes request the same address.
• Adding padding without identifying which warp instruction has the conflicting pattern.
• Ignoring element width, vector access, layout, and architecture when reasoning about bank mapping.
• Removing conflicts while increasing shared-memory allocation enough to reduce useful residency.