TL;DR: Coalescing is whether a warp's 32 global-memory addresses fall in a few contiguous sectors (fast) or scatter into up to 32 transactions (10x slower); bank conflicts are whether a warp's 32 shared-memory addresses hit 32 distinct banks or serialize k-way in one bank. The matrix transpose hits both, and the canonical fix is staging through a padded
tile[32][33]so column accesses spread across banks.
How to approach it
These are sibling problems at two levels of the hierarchy: coalescing is about how a warp touches global memory (HBM), bank conflicts about how it touches shared memory. Define each with its mechanical cause, give the canonical numbers (32-thread warp, 32 banks), then walk the matrix transpose, which exhibits both and has a famous two-line fix.
A strong answer
Coalescing. When a warp issues a load, the hardware combines the 32 addresses into as few memory transactions as possible. If thread i reads data[base + i], meaning 32 consecutive 4-byte words, that is one or a few 32-byte sectors and near-perfect bandwidth use. If threads stride (data[i * 1024], column access in a row-major matrix), each thread's word lands in a different sector and the warp triggers up to 32 separate transactions to deliver the same 128 useful bytes. You are reading the DRAM at full speed and throwing most of it away; effective bandwidth can drop by an order of magnitude. This is also the argument for structure-of-arrays over array-of-structs: with AoS, a warp reading particles[i].x skips over y, z, mass on every load.
Bank conflicts. Shared memory is divided into 32 banks, 4 bytes wide, interleaved by address. Each bank serves one access per cycle. If the 32 threads of a warp hit 32 different banks (or all read the same address, which is a broadcast and free), the access completes in one cycle. If k threads hit different addresses in the same bank, the access serializes k-way. The classic trigger is a 2D shared array tile[32][32] accessed by column: tile[threadIdx.x][col] walks addresses 32 floats apart, which all map to the same bank, a 32-way conflict.
The transpose. A naive transpose reads rows and writes columns: either the read or the write of global memory is uncoalesced. The standard fix stages through shared memory: read a 32×32 tile row-wise (coalesced), write it out transposed row-wise (also coalesced), doing the index swap inside shared memory. But that inside swap is exactly the column access that causes a 32-way bank conflict. The fix is one character:
__shared__ float tile[32][33]; // 33, not 32: padding skews columns across banks
The padding column shifts each row's bank mapping by one, so column accesses now touch 32 different banks. The proof is two lines of modular arithmetic, and being able to write it is what "be ready to say why padding works" means. An element tile[r][c] sits at linear address r * width + c, and its bank is that address mod 32. With width 32: bank = (r * 32 + c) mod 32 = c, so a column access (fixed c, r running 0 to 31) maps every thread to the one bank c, a 32-way conflict. With width 33: bank = (r * 33 + c) mod 32 = (r + c) mod 32, because 33 mod 32 is 1, so the same column access now walks all 32 banks exactly once (checked by enumerating it). The extra column is dead storage, 32 floats per tile, purchasing a change in the address-to-bank residue, which is also why the follow-up answer is "any width coprime with 32 works, 33 is just the cheapest." Naive, then tiled, then padded is typically a 3-8x progression on bandwidth-bound transposes, and reciting it credibly is worth a lot in an NVIDIA loop.
Diagnosis order in practice: Nsight Compute will show you global memory efficiency (requested vs transferred bytes) for coalescing and a shared-memory bank-conflict counter. Fix coalescing first, because HBM transactions cost hundreds of cycles versus tens for a conflict replay.
What interviewers probe next
- "Why 33 and not 34 or 64?" Any padding coprime with 32 breaks the alignment; 1 element is the cheapest in wasted space.
- "Do bank conflicts apply to global memory?" No; banks are a shared-memory concept. Global's analogue is sectors and coalescing. Mixing these up is an instant signal.
- "What if my data is inherently strided?" Stage it: do one coalesced pass to reorder into a friendly layout (or via shared memory per-tile), then run the compute, amortizing the shuffle.
- "How do modern attention kernels avoid these?" FlashAttention-style kernels are built around coalesced tile loads and conflict-free shared layouts (with swizzling rather than padding on Hopper, since TMA prefers unpadded tiles).
Common mistakes
- Conflating the two: "bank conflicts in global memory" or "coalescing shared memory" both reveal the mental model is missing a level.
- Forgetting broadcast: all 32 threads reading the same shared address is free, not a 32-way conflict, and candidates who do not know this over-engineer reductions.
- Quoting the transpose fix without being able to say why padding works (changing the address-to-bank mapping per row).
- Optimizing bank conflicts on a kernel that is globally uncoalesced. Wrong order; HBM waste dominates.
Key takeaways
- Coalescing lives in global memory (sectors); bank conflicts live in shared memory (32 banks). Never swap the vocabulary.
- The transpose needs both fixes: tile through shared memory for coalescing, pad to
[32][33]for conflicts. - Same-address shared reads broadcast for free; fix global coalescing before chasing conflicts.