Bank Conflicts Inside GPU Shared Memory

Shared memory is not free of access patterns; it has its own, and they are different from the ones that govern global memory. There is no cache line, no sector, and no tag lookup: the array is a static SRAM that the hardware reads in a deterministic number of passes per warp, and the pass count is determined by an address property that is trivially computable from the source. The good news is that unlike DRAM behaviour, the cost of a shared-memory access is computable from the address pattern and the bank geometry. The bad news is that the natural way to lay out a tile is often exactly the pattern that maximizes the cost.

How an address selects a bank

The array is divided into 32 banks, each one word wide, and the bank holding a given address is determined by the address itself: with four-byte words, bank = (word address) modulo 32. Because 32 is a power of two, this is a selection of the low five bits of the word address, and it is not affected by anything else in the address. There is no hashing of the address. The consequence is that a stride of 32 words — 128 bytes — between lanes, or any other multiple of 32 words, maps every lane to the same bank.

A warp-wide access is served in a number of passes equal to the largest number of lanes that target any single bank. If all thirty-two lanes hit distinct banks, the access completes in one pass. If the lanes split evenly across two banks, it takes two passes, and the effective bandwidth of the scratchpad is halved for that access. Thirty-two lanes on one bank take thirty-two passes.

Lane addresses within a warp	Banks hit	Passes
Consecutive words	32 distinct	1
Stride of 2 words	16 distinct, 2 lanes each	2
Same word in every lane	1, broadcast	1
Stride of 32 words	1 bank, 32 lanes	32

The four-byte word size is the reference case, and it is the one the numbers above assume. Everything downstream, including how a wider load is split into phrases, follows from it.

The broadcast rule

Reading the same word from several lanes is not a conflict. The hardware recognizes that all the lanes want one value and returns it to the warp in a single pass, regardless of how many lanes participate. This makes a scratchpad well suited to delivering a scalar that every thread needs, like a loop bound or a scale factor, without any of the cost that the same pattern would incur in global memory.

The rule is per word, not per lane group. Sixteen lanes reading word w and sixteen lanes reading word w + 1 is a broadcast in two banks, one pass. Sixteen lanes reading w and sixteen reading w + 32 is two lane groups on one bank, two passes, because the two words are 32 words apart and therefore share a bank.

Broadcasts also interact with the conflict rule in a way that occasionally saves an otherwise bad layout. A column of a tile with 32 rows and 32 columns can be conflict-free if the warp is reading the same column entry in many lanes; the conflicts appear when each lane wants a different entry of the same bank.

When a whole warp lands on one bank

The canonical failure is a square tile accessed by column. Declaring a tile as a two-dimensional array of 32 rows and 32 columns and having each lane read one element down a column gives addresses 0, 32, 64, 96 and so on in words. Every one of those is congruent to zero modulo 32, so a full warp’s read becomes thirty-two serialized passes, and a kernel that reads the tile this way in a loop of 32 iterations pays 1024 passes for what should have been 32.

The same arithmetic explains a subtler variant. A warp scattering values from a tile whose row pitch happens to be a multiple of 128 bytes will conflict on every read where the lanes differ only in row. Row pitch is a layout constant chosen in the declaration, and it is not visible in the loop that causes the problem.

The remedies are all variations on changing the pitch.

#define N 32

// Launched with blockDim = (32, 8); one warp covers all 32 columns of one row.
__global__ void column_read(const float *in, float *out, int n) {
    __shared__ float flat[N][N];        // pitch 32 words
    __shared__ float skew[N][N + 1];    // pitch 33 words

    int tx = threadIdx.x, ty = threadIdx.y;
    for (int j = 0; j < N; j += 8) {    // rows: lanes vary tx, conflict-free
        flat[ty + j][tx] = in[(ty + j) * n + tx];
        skew[ty + j][tx] = in[(ty + j) * n + tx];
    }
    __syncthreads();

    float sum = 0.0f;
    for (int r = 0; r < N; r += 8) {
        sum += flat[tx][ty + r];   // word tx*32 + c: one bank, 32 passes
        sum += skew[tx][ty + r];   // word tx*33 + c: 32 banks, one pass
    }
    if (ty == 0) out[tx] = sum;
}

Both reads use the same row loop shape, but the word index differs: tx*32 + c is constant modulo 32 across the warp, while tx*33 + c takes a different value modulo 32 for every lane. The padding does not merely sprinkle the addresses around; it makes the mapping from lane to bank a bijection, which is the property that has to hold for a single pass.

The example also shows why the pitch must be considered a property of the declaration rather than of the loop. The write phase is conflict-free in both layouts because its index varies the fastest along the contiguous axis, and nothing in the loop text distinguishes the two cases.

Padding, skewing, and what they cost

Adding one column to a 32 by 32 float tile costs 32 words, which is 128 bytes, or a little over 3 percent of the tile. It removes a thirty-two-way conflict and replaces it with none. That is one of the cheapest trades available in GPU memory programming, and it explains why padding is the first thing to reach for.

It is not always available. A pitch of 33 makes the tile’s row start drift by one word relative to the bank boundary, which is exactly what removes the conflict, but it also means the tile no longer has a power-of-two pitch, so the compiler cannot turn a two-dimensional index into a shift and an add. The address arithmetic gets slightly more expensive for every access, which is usually a good trade and occasionally not. Vectorized access is the other constraint: a pitch of 33 words is not a multiple of an aligned sixteen-byte vector, so a tile accessed as quad-word vectors cannot use the odd pitch without splitting the load. In those cases a skew, in which each row starts a different distance into the array, achieves the same bank spreading while keeping vector-friendly granularity; the storage cost is one row’s worth of extra span at the top of the tile.

There is a third option that costs nothing in storage: change which axis the lanes vary along, so that the strided axis is handled by the loop rather than by the warp. This is the same fix as the row-major trap in global memory, and it is available whenever the algorithm does not require a specific assignment of lanes to tile coordinates.

Wider accesses change the arithmetic

A load that moves eight or sixteen bytes per lane is not analyzed as one word per bank. The hardware breaks a warp’s wider access into phases, and each phase is served against the banks under the same modulo-32 rule. A warp reading sixteen consecutive bytes per lane covers 512 bytes, which is 128 words, well beyond what the array can return in one pass; the access therefore proceeds in several phases even though every lane is contiguous with its neighbour, and the layout must divide those phases evenly across the banks for any of them to be conflict-free.

Some architectures have offered a selectable bank width, in which the array is organized as 32 banks of eight bytes instead of four, which changes the modulo arithmetic and can remove conflicts from double-precision tiles at the cost of introducing them elsewhere. Because the geometry is a property of the architecture, and on some parts a per-kernel launch setting, code that assumes one bank width may be conflict-free on one generation and conflicted on the next. The durable approach is to derive the bank index from the platform’s reported geometry where the platform reports it, and otherwise to prefer odd pitches and lane assignments that distribute lanes across all thirty-two banks by construction.

The scratchpad also interacts with occupancy, since the array is per block and the amount of it a block reserves limits how many blocks reside on an SM. A padded tile costs slightly more of that budget, and a kernel that is already near the limit may lose a resident block to the extra column. The trade is still usually positive, but it is a real one, and it sits alongside the larger question of how much of the tile should live in the scratchpad at all.