Explain shared memory bank conflicts with the bank arithmetic, show a kernel that has them, and fix it with padding.

TL;DR: Shared memory is split into 32 banks, 4 bytes wide, interleaved: the bank of a 4-byte word at byte address A is (A / 4) mod 32. A warp's shared-memory access completes in one pass if each bank is touched by at most one distinct address (the same address from several lanes is a broadcast and free); otherwise it takes as many passes as the most-loaded bank, an n-way conflict. In a tile declared float tile[32][32], element [r][c] is word 32r + c, bank c; a warp reading a column (lanes vary r, c fixed) puts all 32 lanes in bank c: a 32-way conflict, 32 passes. Declaring tile[32][33] makes element [r][c] word 33r + c, bank (33r + c) mod 32 = (r + c) mod 32, so 32 lanes with different r hit 32 different banks: one pass. The cost is 1/32 more shared memory. The transpose kernel is the standard case, and Nsight Compute's "shared load/store bank conflicts" counter drops to zero when the fix is right.

How to approach it

Give the bank formula and the one-pass rule. Work the column-access example with the arithmetic, then the padded version. Show the transpose kernel before and after. Name the profiler counter, the cost of padding, and the alternative (swizzling). Close with the other common conflict patterns.

A strong answer

A typical situation: a candidate fixes uncoalesced global reads in a transpose by staging through shared memory, and the kernel gets no faster; the profiler shows 32-way bank conflicts on the shared reads, which cost exactly what the uncoalesced global reads did.

The arithmetic:

bank(A) = (A / 4) mod 32 for a 4-byte word at byte address A (banks are 4 bytes wide, 32 of them, so 128 bytes per
  "row" of banks, and the pattern repeats every 128 bytes)
one-pass rule: a warp's access is serviced in one pass if no bank sees two different addresses; k different
  addresses in one bank → k passes (k-way conflict); identical addresses across lanes → broadcast, no conflict

tile[32][32] of floats, row-major: element [r][c] at word 32r + c → bank (32r + c) mod 32 = c
  row access: lanes vary c (0..31), r fixed → banks 0..31 → 1 pass ✓
  column access: lanes vary r (0..31), c fixed → all in bank c → 32-way conflict → 32 passes ✗
tile[32][33]: element [r][c] at word 33r + c → bank (33r + c) mod 32 = (r + c) mod 32
  row access: r fixed, c varies → banks (r + 0..31) mod 32 → all distinct → 1 pass ✓
  column access: c fixed, r varies → banks (0..31 + c) mod 32 → all distinct → 1 pass ✓
  cost: 32 × 33 × 4 = 4,224 bytes instead of 4,096; 3% more shared memory
sanity: a 32-way conflict makes the shared read 32× slower, which turns the "fast" staging path into the same cost as
        the uncoalesced global access it replaced; the profiler is the only way to see it, since the code looks fine.

Shared Memory and Bank Conflicts has the bank model and the 8-byte and 16-byte access rules (which change the effective bank width); GPU Memory Hierarchy places shared memory in the SM.

The transpose kernel, before and after:

#define TILE 32
#define ROWS 8      // 32 × 8 = 256 threads; each thread handles 4 elements of the tile

// out[j][i] = in[i][j], N×N, row-major
__global__ void transpose_conflicts(const float* __restrict__ in, float* __restrict__ out, int N) {
    __shared__ float tile[TILE][TILE];                 // no padding
    int x = blockIdx.x * TILE + threadIdx.x;           // column in 'in'
    int y = blockIdx.y * TILE + threadIdx.y;           // row in 'in'
    for (int k = 0; k < TILE; k += ROWS)               // coalesced global read: warp reads 32 consecutive floats of a row
        tile[threadIdx.y + k][threadIdx.x] = in[(y + k) * N + x];
    __syncthreads();
    x = blockIdx.y * TILE + threadIdx.x;               // transposed block coordinates
    y = blockIdx.x * TILE + threadIdx.y;
    for (int k = 0; k < TILE; k += ROWS)               // coalesced global write; but the shared READ walks a column:
        out[(y + k) * N + x] = tile[threadIdx.x][threadIdx.y + k];   // lanes vary threadIdx.x = row → bank conflict 32-way
}

__global__ void transpose_padded(const float* __restrict__ in, float* __restrict__ out, int N) {
    __shared__ float tile[TILE][TILE + 1];             // +1 column of padding
    int x = blockIdx.x * TILE + threadIdx.x;
    int y = blockIdx.y * TILE + threadIdx.y;
    for (int k = 0; k < TILE; k += ROWS)
        tile[threadIdx.y + k][threadIdx.x] = in[(y + k) * N + x];
    __syncthreads();
    x = blockIdx.y * TILE + threadIdx.x;
    y = blockIdx.x * TILE + threadIdx.y;
    for (int k = 0; k < TILE; k += ROWS)
        out[(y + k) * N + x] = tile[threadIdx.x][threadIdx.y + k];   // word 33·x + (y+k): banks (x + y + k) mod 32, distinct across x
}

Both versions have coalesced global reads and writes (each warp touches 32 consecutive floats of a row in both phases); the difference is only the shared-memory read pattern in the second loop. The measured effect on an H100 for a 4,096 × 4,096 matrix: the unpadded version reaches roughly a third of copy bandwidth, the padded one roughly 90% of it. Transposing without Uncoalesced Writes (the companion question) covers the rest of the kernel's design.

The profiler counter and the diagnosis:

Nsight Compute → Memory Workload Analysis → Shared memory:
  "bank conflicts" (l1tex__data_bank_conflicts_pipe_lsu_mem_shared) per load and store; zero after the fix
  "shared memory wavefronts" vs "ideal wavefronts": a 32-way conflict shows as 32× the ideal
  the source view attributes them to the line; the column read in the second loop lights up

Profiling with Nsight walks the section.

Padding's cost and the alternative:

padding: +1 word per row costs 1/TILE of the tile (3% at 32); for 2D tiles in a GEMM with 8-byte or 16-byte
  vectorized shared loads, the padding must keep rows aligned for the vector width (pad by 4 or 8 words) and the
  waste grows; also breaks the alignment cp.async and TMA want
swizzling: store element [r][c] at column c XOR (r mod 32) (or a permutation of column groups) instead of c; the
  layout stays dense and aligned, and both row and column accesses hit distinct banks; this is what CUTLASS and
  the tensor-core paths use, because ldmatrix and wgmma need dense, aligned tiles
rule: padding for hand-written kernels with scalar shared accesses; swizzling when the tile feeds tensor cores or
  vector loads

CUTLASS and Tensor Core Kernels shows the swizzled layouts; Tiled Matrix Multiplication is the GEMM where they appear.

Other conflict patterns worth naming: a stride of 2 words across lanes (lane t reads word 2t) is a 2-way conflict (banks 0, 2, 4, ... repeat after 16 lanes); a stride of any even number conflicts, odd strides do not (the reason 33 works); 8-byte accesses use pairs of banks and 16 lanes per pass, so the arithmetic changes; and a structure-of-arrays layout in shared memory often turns an array-of-structures conflict into none.

The reversal condition: a kernel whose shared accesses are all row-wise (lanes vary the last index) has no conflicts and needs no padding; adding padding there wastes shared memory and can reduce occupancy for nothing, so the profiler counter decides.

What interviewers probe next

• "Why does 33 work and 34 not?" Bank = (33r + c) mod 32 = (r + c) mod 32 cycles through all banks as r varies; (34r + c) mod 32 = (2r + c) mod 32 repeats every 16 rows: 2-way conflict. Odd padding widths work; even ones do not fully.
• "What is the cost of a 2-way conflict?" Two passes instead of one on that instruction; often hidden by other warps' work; a 32-way one is not hidden.
• "Does padding help global memory?" No; global memory is sectors and lines, not banks; the fix there is coalescing. The two are different mechanisms with similar-looking symptoms.
• "How do you see conflicts without a profiler?" You mostly cannot; a micro-benchmark of the shared access pattern alone, timed with events, shows the 32× slowdown.

Common mistakes

• Staging through shared memory to fix coalescing and re-creating the cost as bank conflicts.
• Padding by an even number.
• Padding a tile that feeds tensor-core loads and breaking their alignment.
• Confusing bank conflicts (shared memory) with uncoalesced access (global memory).

Key takeaways

• bank = (word index) mod 32; a warp access takes as many passes as the most-loaded bank; broadcast is free.
• A column walk on a 32-wide tile is a 32-way conflict; [32][33] makes it one pass at 3% extra memory.
• Confirm with Nsight Compute's shared bank-conflict counter.
• Padding for scalar hand-written tiles; swizzling for tensor-core and vector-load tiles.