TL;DR: Shared memory has 32 banks, each 4 bytes wide, so word
wlives in bankw mod 32. Reading a column of afloat tile[32][32]puts all 32 lanes in bankc, a 32-way conflict that serializes the read into 32 wavefronts. Padding the row to 33 words makes the bank(r + c) mod 32, distinct for every lane, so the read takes one. Padding breaks once rows must stay 16-byte aligned for vector or TMA stores (a 4-float pad leaves a 4-way conflict), which is why production kernels XOR-swizzle the column index instead.
How to approach it
Write the address of the element each lane touches, convert it to a 4-byte word index, and take it mod 32. Ask the element type and the access width per lane, because both change the answer, and ask whether the tile is filled by plain stores, vector stores or TMA. Count conflicts as wavefronts beyond the minimum the request size forces.
A strong answer
A naive tiled transpose loads a 32 by 32 tile from global memory with coalesced row reads, writes it into shared memory row-wise, then reads it column-wise so the global store is also coalesced. The row-wise write is fine. The column-wise read is not:
element (r, c) of float tile[32][32] -> word r*32 + c
bank = (r*32 + c) mod 32 = c
lanes r = 0..31 read column c -> all in bank c, 32 distinct words
cost = 32-way conflict = 32 wavefronts
The one-column pad changes the row stride from 32 words to 33, which is odd and therefore coprime with 32:
float tile[32][33]: element (r, c) -> word r*33 + c
bank = (r + c) mod 32
lanes r = 0..31, fixed c -> 32 different banks, 1 wavefront
row-wise write, fixed r, c = 0..31 -> (r + c) mod 32, also distinct
cost of the pad = 32 words = 128 B per tile (3%)
template <int TILE = 32>
__global__ void transpose(const float* __restrict__ in, float* __restrict__ out,
int rows, int cols) {
__shared__ float tile[TILE][TILE + 1]; // +1: column reads hit 32 banks
int x = blockIdx.x * TILE + threadIdx.x; // blockDim = (32, 8)
for (int j = threadIdx.y; j < TILE; j += blockDim.y) {
int y = blockIdx.y * TILE + j;
if (x < cols && y < rows) tile[j][threadIdx.x] = in[(size_t)y * cols + x];
}
__syncthreads();
x = blockIdx.y * TILE + threadIdx.x;
for (int j = threadIdx.y; j < TILE; j += blockDim.y) {
int y = blockIdx.x * TILE + j;
if (x < rows && y < cols) out[(size_t)y * rows + x] = tile[threadIdx.x][j];
}
}
Now change the element type, which is what inference kernels actually move. A BF16 tile [32][32] packs two elements per word, so its rows are 16 words long:
bf16 (r, c) -> word r*16 + c/2 -> bank (16r + c/2) mod 32
even r -> bank c/2, odd r -> bank 16 + c/2: 2 banks for 32 lanes, 16-way conflict
pad to [32][33] (66 B rows) -> scalar 2-byte column reads conflict-free
pad to [32][34] (17-word rows) -> also conflict-free, and rows stay 4-byte aligned
The 33-element pad works for scalar reads but leaves rows on 2-byte boundaries, so any packed bf16x2 store into a row is misaligned. Pad by two elements.
Where padding fails. Fast kernels fill shared memory with 16-byte vector stores or with TMA, and both need 16-byte-aligned rows. A 33-float row is 132 bytes, not a multiple of 16. The aligned pad is 4 floats, a 36-word stride:
float tile[32][36]: bank (36r + c) mod 32 = (4r + c) mod 32
lanes r = 0..31 -> only 8 distinct banks -> 4-way conflict
The fix that keeps alignment is a swizzle: permute where data lands inside each row instead of shifting the rows. For a scalar FP32 tile, store element (r, c) at column c ^ r. Each row is still a permutation of its own 32 words, so row writes stay conflict-free, the tile has no wasted bytes, and a column read lands in 32 different banks. Tensor-core kernels apply the same idea at 16-byte granularity: with 128-byte rows split into eight 16-byte chunks, chunk q of row r is stored at chunk q ^ (r mod 8). Eight lanes that each read one 16-byte chunk from eight consecutive rows (the ldmatrix and WGMMA access pattern) then cover all 32 banks in one wavefront instead of piling 8 deep on 4 banks. TMA writes tiles in exactly these layouts with its 32-, 64- and 128-byte swizzle modes, and WGMMA reads its shared-memory operands in them, so a padded tile cannot feed it. Padding is a fix for hand-written scalar code; swizzling is the fix for tensor-core pipelines.
Every conflict-free claim above was checked by computing distinct words per bank for all 32 lanes.
What it costs. Shared memory delivers 32 banks x 4 B = 128 B per clock per SM. A conflict-free 4-byte read by a full warp is one wavefront; a 32-way conflict is 32, so the read phase of the transpose runs at 1/32 of shared-memory bandwidth. A warp-wide 16-byte load requests 512 B and needs at least 4 wavefronts by bandwidth alone, which is why Nsight Compute reports conflicts as excess wavefronts over that ideal rather than as a raw count.
The reversal condition: a transpose of a large matrix is bound by HBM, not by shared memory, so a moderate conflict may never show in end-to-end time:
shared memory, whole GPU = 132 SMs x 128 B/clock = 16,896 B per clock
at an assumed 1.7 GHz = 16,896 B x 1.7e9 = about 29 TB/s
HBM3, H100 SXM = 3.35 TB/s, and a transpose reads and writes every byte
A 32-way conflict cuts the shared-memory side to under 1 TB/s and becomes the bottleneck; a 2-way conflict leaves it near 14 TB/s, still well above HBM. Fix the 32-way case, measure, and stop when DRAM throughput is near peak.
What interviewers probe next
"Two lanes read the same word. Conflict?" No. Same-word accesses are a broadcast; conflicts are distinct words in one bank.
"Do this for FP64." A column of double tile[32][32] puts every lane on banks 2c and 2c+1, 32 deep. Padding by one double (a 66-word stride) spreads them so only lanes r and r+16 share banks: 2 wavefronts, which is the floor anyway, because the warp requests 256 bytes.
"Why not avoid shared memory and transpose with warp shuffles?" For a 32 by 32 FP32 tile per warp, shuffles work without any bank analysis, at the cost of more instructions per element; it pays off for small tiles held entirely in registers.
How it's scored
Weak: knows "add one to the column dimension" but cannot say which access conflicts or why 33 works.
Hire: writes word index mod 32 for both access patterns, gets 32-way and conflict-free, and explains coprimality.
Strong Hire: redoes the arithmetic for BF16 and 16-byte stores, shows the aligned pad leaves a 4-way conflict, and explains why TMA and WGMMA use swizzled layouts instead of padding.
Common mistakes
- Calling a broadcast (lanes reading one word) a bank conflict.
- Padding a BF16 tile by one element and then storing rows as packed pairs, which faults or silently splits into narrower accesses.
- Computing banks from element indices instead of 4-byte word indices.
- Chasing a 4-way conflict in a kernel whose profile shows HBM at 90% of peak.