Shared memory is physically split into 32 parallel banks, each serving one 4-byte word per cycle. A warp's 32 lanes can be served in one cycle in two cases:

• each lane's address falls in a different bank
• every lane wants the exact same word (broadcast).

Any other pattern serializes: two lanes on the same bank means two cycles, not one.

Banks are 32 separate SRAM arrays on the SM, each with its own address decoder, sense amplifiers, and read port. That hardware is why one bank serves one word per cycle.

Bank of the word at byte address A is (A / 4) % 32.

• Contiguous 4-byte memory blocks in SRAM can be visualized in terms of banks they belong to.
  • Each column in the visualization is one bank.

__shared__ float smem[32 * 32];   // 1024 words, 4 bytes each

// word i  ->  bank = i % 32,  row = i / 32
// smem[0..31]   -> banks 0..31, row 0
// smem[32..63]  -> banks 0..31 again, row 1
// Bank 0 owns smem[0], smem[32], smem[64], smem[96] - one whole column and so on

Accessing SRAM without and with bank conflicts

• Thread tid reads smem[tid]: no-conflict request
  • Each lane asks for a different column, so no bank is asked for two words.
  • Total cycles needed: 1

__shared__ float smem[32 * 32];
float x = smem[tid];   // lane t -> bank t, row 0. 32 distinct banks.

• Thread tid reads smem[tid * 2]: bank conflict
  • Doubling the index skips every other column. Two words are read per column, requiring two cycles.
  • Cycle 1: every bank serves its first word
  • Cycle 2: every bank serves its second word
  • Total cycles needed: 2

__shared__ float smem[32 * 32];
float x = smem[tid * 2];   // 16 even banks, two words each: 2 cycles

• Each thread tid reads the same data smem[constant].
  • The value is broadcasted. It doesn't count as bank conflict.
  • Total cycles needed: 1

__shared__ float smem[32 * 32];
float x = smem[0];   // one bank, but the same word: fanned out in 1 cycle

• Thread tid reads smem[(tid % 4) * 32]: four words, one bank
  • All 32 lanes land on bank 0, but they only want 4 distinct words.
  • The bank serves one word per cycle, fanning each out to the 8 lanes that asked for it.
  • Total cycles needed: 4

__shared__ float smem[32 * 32];
float x = smem[(tid % 4) * 32];   // 4 distinct words on bank 0: 4 cycles

• Thread tid reads smem[tid * 32]: worst-case scenario
  • Each warp lane accesses the same bank. But different addresses in the same bank. No broadcasting is possible here.
  • Each cycle reads one value at a time.
  • Total cycles needed: 32

__shared__ float smem[32 * 32];
float x = smem[tid * 32];   // all 32 lanes on bank 0, 32 different words: 32 cycles

Wider loads (8/16-byte) have different collision math, not covered here. This uses the classic per-cycle bank model - the number that matters is 32 banks, one 4-byte word each, per cycle.

The padding trick in matrix transpose

A matrix transpose reads a tile down a column.

• With a [32][32] tile that is the worst case above.
  • Lane t reads tile[t][c], which is index t*32 + c, and (t*32 + c) % 32 = c for every t.
  • All 32 lanes hit bank c: 32 cycles are needed to read per warp.

Padding the row to 33 costs one unused word per row and fixes it completely.

• Reading a column of a [32][32] tile: worst possible bank conflict
  • Total cycles needed: 32
• The same read on a [32][33] tile: no bank conflict
  • The extra column shifts each row by one, so bank = (r + c) % 32 instead of c.
  • All 32 words are served together.
  • Total cycles needed: 1

__shared__ float tile[32][33];   // 33, not 32
float x = tile[tid][c];          // bank = (tid + c) % 32: all 32 distinct

The cost is 32 wasted words per tile. The gain is turning a 32-cycle access into a 1-cycle one. This trick shows up in essentially every hand-written transpose and tiled matmul.