Bank conflict happens when threads in a warp access different addresses that fall in the same shared memory bank, and the hardware serialises them. Shared memory is divided into 32 banks, each 4 bytes wide and interleaved so that consecutive 4-byte words land in consecutive banks.

__shared__ float tile[32][32];
 
tile[threadIdx.x][0];   // 32-way conflict: whole column is in bank 0
tile[0][threadIdx.x];   // conflict-free: one word per bank

Word k lives in bank k % 32. In a [32][32] array, every element of a column is 32 words apart, so the whole column maps to one bank and a warp reading it takes 32 serialised transactions instead of one.

The standard fix is to pad the row length so the stride is no longer a multiple of 32:

__shared__ float tile[32][33];   // one extra column of padding
 
tile[threadIdx.x][0];            // now conflict-free

With a row length of 33, consecutive rows are offset by one bank, so a column walk touches all 32 banks. The padding wastes a little shared memory and nothing else.

Broadcasts are not conflicts. When every thread in a warp reads the same address, the hardware broadcasts one value at full speed, which is why tile[0][threadIdx.x] in a stencil and a uniform lookup both behave well.

This is the GPU counterpart of False sharing on a CPU. Both are cases where a memory system's granularity, cache lines in one case and banks in the other, turns an access pattern that looks fine in source into serialised traffic. Nsight Compute reports shared memory conflicts directly, so it is worth checking rather than guessing.