Demystifying Shared Memory Bank Conflicts in CUDA

When I first started writing CUDA kernels, shared memory felt like a magic bullet. You take data from slow global memory, stage it into fast on-chip SRAM, and bypass high-latency memory buses. Because shared memory sits directly on the Streaming Multiprocessor next to the execution units, it delivers terabytes per second of bandwidth.

However, as you start writing 2D matrix operations or custom tile loaders, you often hit an invisible speed bump: your code is using shared memory, yet execution slows down and threads seem to stall.

More often than not, the culprit is a shared memory bank conflict.

In this post, we will break down what bank conflicts actually are, why the hardware is wired this way, and how a simple padding trick makes your code run at full speed.

---

Why Shared Memory Has Banks

To understand bank conflicts, picture what happens under the hood.

Threads execute in lockstep groups of 32 called warps. When a warp reads from shared memory, all 32 threads can ask for data at the exact same clock cycle.

Now imagine shared memory as a single notebook. If 32 people want to read 32 completely different lines from that same notebook at the exact same millisecond, the hardware faces a physical problem. For an on-chip memory cell to let multiple threads read from it at the same moment, you need dedicated physical wiring and sensing circuits for each simultaneous reader, known as a read port.

Building a single memory block with 32 physical read ports is impractical. The wiring complexity and silicon footprint blow up rapidly, consuming massive chip space, generating excess heat, and driving up costs.

Instead, hardware architects chose a practical engineering compromise: they split shared memory into 32 independent memory banks, each with a single, lightweight read port.

Because there are 32 banks and 32 threads in a warp, the hardware achieves an ideal match:

• If all 32 threads ask for data from 32 distinct banks, all 32 requests finish in parallel in a single clock cycle.
• If multiple threads ask for different words from the same bank, that bank can only serve one request at a time. The hardware must serialize them into multiple cycles.

That serialization is what we call a bank conflict.

---

The Mental Model: 32 Bank Tellers

Picture a bank branch with 32 teller windows, numbered from Window 0 to Window 31.

A group of 32 customers (the threads in a warp) walks into the lobby at the exact same moment:

1. The Ideal Scenario (Parallel Access): Each customer walks up to a different window. Customer 0 visits Window 0, Customer 1 visits Window 1, and so on. Every teller serves one customer simultaneously. Everyone finishes in 1 cycle.
2. The Broadcast (Same Address): Several customers walk up to Window 0 asking for the exact same public sheet of paper. The teller simply reads it aloud to all of them. In CUDA, when multiple threads read the exact same address from a bank, the value is broadcast in a single cycle with zero conflict.
3. The Bank Conflict (A Queue Forms): Several customers line up at Window 0, but each person asks for a different private document. The teller can only help one customer at a time. Customer 0 is served first, then Customer 1, then Customer 2.

If all 32 threads line up at Window 0 for distinct pieces of data, that single memory instruction gets split into 32 separate transactions. What should have taken 1 cycle now takes 32 cycles.

---

How Addresses Map to Banks

On modern NVIDIA GPUs, shared memory is divided into 32 banks, and each bank is 4 bytes (32 bits) wide, matching the size of a standard float or int.

Memory addresses are striped across the banks sequentially word by word:

• Word 0 (bytes 0 to 3) belongs to Bank 0
• Word 1 (bytes 4 to 7) belongs to Bank 1
• Word 2 (bytes 8 to 11) belongs to Bank 2
• ...
• Word 31 (bytes 124 to 127) belongs to Bank 31
• Word 32 (bytes 128 to 131) wraps back around to Bank 0

The bank assignment formula is:

bank_index = (byte_address / 4) % 32

Because 32 banks multiplied by 4 bytes equals 128 bytes, addresses wrap around every 128 bytes. Any two addresses whose byte difference is a multiple of 128 bytes map to the exact same bank.

---

The Classic Trap: Column Reads in a 2D Tile

Sequential access is naturally conflict-free:

__shared__ float buffer[1024];
float val = buffer[threadIdx.x]; // Thread 0 hits Bank 0, Thread 1 hits Bank 1...

Every thread accesses a unique bank, completing in a single transaction.

The problem strikes when threads read down a column in a 2D tile, such as during matrix transposition or when loading matrix columns during tiled matrix multiplication and 2D block tiling:

__shared__ float tile[32][32];

// Threads in a warp read down column 0
int row = threadIdx.x;
float val = tile[row][0];

Because C and CUDA arrange 2D arrays in row-major order:

• Thread 0 accesses index 0 * 32 = 0 (Bank 0)
• Thread 1 accesses index 1 * 32 = 32 (Bank 0)
• Thread 2 accesses index 2 * 32 = 64 (Bank 0)
• ...
• Thread 31 accesses index 31 * 32 = 992 (Bank 0)

All 32 threads target Bank 0. The hardware has to serialize this access into 32 successive clock cycles, cutting shared memory bandwidth down to roughly 3% of peak capability.

---

The Simple Fix: Array Padding

The most reliable, practical way to eliminate this conflict is adding a small amount of padding to the array definition.

Instead of declaring tile[32][32], add just one extra dummy element per row:

__shared__ float tile[32][33]; // One extra column of padding

Let us look at what happens to the bank assignments when threads read down column 0 now:

• Thread 0 reads index 0 * 33 + 0 = 0 (Bank 0)
• Thread 1 reads index 1 * 33 + 0 = 33 (Bank 1)
• Thread 2 reads index 2 * 33 + 0 = 66 (Bank 2)
• Thread 3 reads index 3 * 33 + 0 = 99 (Bank 3)
• ...
• Thread 31 reads index 31 * 33 + 0 = 1023 (Bank 31)

By shifting each successive row by 1 float, the column access rotates across all 32 distinct banks. What was previously a vertical collision down Bank 0 becomes a diagonal walk across all 32 banks.

The 32-way bank conflict vanishes completely, dropping latency from 32 cycles back down to 1 cycle.

---

Quick Summary

• 32 Banks: Shared memory is split into 32 banks, 4 bytes wide each.
• Different Banks = Fast: When all threads in a warp touch different banks, access finishes in 1 cycle.
• Same Address = Fast: When threads read the exact same word, the value is broadcast for free.
• Different Words, Same Bank = Slow: Accesses are serialized, forming a queue.
• The Stride-32 Danger: Column access in a [32][32] tile causes all threads to slam into Bank 0.
• Pad by 1: Changing [32][32] to [32][33] offsets each row by one bank and eliminates the conflict.