# Bank conflicts (CUDA)
[[shared-memory-cuda|Shared memory]] is divided into 32 **banks**, each 4 bytes wide, interleaved so that consecutive 4-byte words land in consecutive banks. A **bank conflict** happens when threads in a [[warps-cuda|warp]] access different addresses that fall in the same bank, and the hardware serialises them.

```c
__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:

```c
__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.