Skip to main content

Shared Memory Bank Conflicts

Shared memory earns its speed by serving a whole warp in one cycle, but that promise depends on the warp's 32 addresses landing in 32 different pieces of hardware. When they don't, shared memory — normally close behind register speed — degrades to a fraction of it, and the kernel doesn't fail, it just quietly runs slower with no compiler warning to explain why.

32 banks​

Shared memory is physically organized into 32 banks, each 4 bytes wide, interleaved across consecutive addresses: bank = (address / 4) % 32. This layout exists so that 32 threads accessing 32 consecutive 4-byte words each land in a different bank and can all be serviced in the same cycle — the shared-memory analog of a coalesced global access, though the mechanism (banks, not sectors) is different, per the glossary entry.

What counts as a conflict​

Threads in a warp hitting distinct banks are serviced in one cycle, no matter how those addresses are otherwise scattered within the 32-bank space. Threads hitting distinct addresses within the same bank cannot be serviced together: N threads colliding on one bank serialize into N sequential cycles, one address at a time. This is a bank conflict — the glossary entry defines it as addresses that map to the same memory bank.

The broadcast exception​

There's one collision that isn't a conflict at all: if every thread in the warp reads the exact same address, the hardware detects it and serves all 32 threads from a single fetch in one cycle — a broadcast, not a conflict. The distinction is address, not bank: same bank, same address is free; same bank, different addresses is N-way serialized.

The padding fix​

The canonical conflict shows up in a shared-memory tile accessed column-wise. Take a square tile:

__shared__ float tile[32][32];

tile is laid out row-major, so tile[row][col] is at word offset row * 32 + col (the bank formula above takes a byte address; for a word offset like this one, dividing by the 4-byte word size is already done, so bank = word_offset % 32). Reading tile[threadIdx.x][k] for a fixed k across a warp — a column read, with threadIdx.x selecting the row — means thread t's word offset is t * 32 + k. Bank is (t * 32 + k) % 32; since 32 is a multiple of 32, every thread's t * 32 term is itself a multiple of 32 and vanishes under % 32, leaving k % 32 regardless of t. All 32 threads land in bank k % 32 — a full 32-way conflict, serializing what should be one cycle into 32.

Padding the row by one extra word breaks the arithmetic:

__shared__ float tile[32][33];

Now tile[row][col] is at offset row * 33 + col. Thread t's address for the same column read is t * 33 + k. Because 33 is not a multiple of 32, each successive t shifts the bank by one relative to the last (33 mod 32 == 1), so 32 threads land in 32 distinct banks — the conflict is gone, at the cost of one wasted word per row (32 × 4 bytes of padding per tile).

Swizzling​

Padding wastes shared memory — one word per row, which adds up when shared memory is the binding occupancy limiter. An XOR swizzle achieves the same bank-spreading effect without allocating any extra space, by permuting the column index instead of shifting the row stride:

// column index swizzled instead of padded
int col = k ^ threadIdx.x;

Indexing tile[threadIdx.x][col] into an unpadded tile[32][32] with this swizzled column spreads the same column read established above — tile[threadIdx.x][k] for a fixed k — across distinct banks the same way padding does, but keeps every row exactly 32 words wide. The tradeoff is complexity: the swizzle function has to be inverted consistently everywhere the tile is written and read, whereas padding only changes a declaration. Reach for swizzling specifically when shared memory is the occupancy limiter and the wasted padding column would drop a block below the next occupancy tier; otherwise padding is simpler and just as fast.

Measuring conflicts​

Nsight Compute reports conflicts directly

The l1tex__data_bank_conflicts_pipe_lsu_mem_shared metric counts shared-memory bank conflicts directly. A nonzero value on a tiled kernel almost always means a missing pad (or a swizzle that isn't actually spreading banks the way it's meant to) — check the tile's declared stride before looking anywhere else.

See also​

  • Shared Memory — allocation and the synchronization rule bank conflicts assume is already in place.
  • Shared Memory Tiling — applying the padded-tile pattern in a full tiled kernel.
  • Matrix Transpose — the canonical kernel where a column-wise shared-memory read causes this exact conflict.
  • GPU & Accelerators — the section index and its three learning paths.