Shared Memory Bank Conflicts, in depth: the bank function, counting wavefronts, padding versus XOR swizzle, and proving the fix in Nsight Compute

The bank model from first principles

Since compute capability 5.x, shared memory has 32 banks, each 4 bytes wide. Successive 32-bit words go to successive banks, so the bank of a byte address is bank = (addr / 4) % 32. Byte 0 to 3 live in bank 0, bytes 4 to 7 in bank 1, and byte 128 wraps back to bank 0. A row of 32 floats therefore covers each bank exactly once, which is why 32 and its multiples are the dangerous strides.

When a warp issues a shared load, the hardware groups the 32 requested addresses by bank. Three cases follow. If every bank is asked for at most one distinct word, the request completes in a single pass, called a wavefront. If several threads ask for the same 32-bit word, that word is read once and broadcast to all of them, so there is no conflict; reading the same scalar from every lane is free. If a bank is asked for k distinct words, the request needs k wavefronts. The degree of the conflict is the maximum k over all banks, not the sum, because the banks still work in parallel within each pass.

The useful corollary is the stride rule. If lane t reads 32-bit word t * s, the number of lanes that collide in one bank is gcd(s, 32). Any odd stride is conflict-free, stride 2 costs two passes, stride 8 costs eight and stride 32 costs thirty-two. You can look at an index expression, read off the stride in words across threadIdx.x, and know the answer before you profile.

Wide accesses and a wavefront simulator

Wider accesses change the grouping. A 64-bit load moves 256 bytes for a warp, more than the 128 bytes the banks deliver per pass, so the hardware resolves it in two phases of 16 lanes; a 128-bit load such as float4 is resolved in four phases of 8 lanes. Conflicts are only counted inside a phase. Contiguous double or float4 accesses therefore cost the ideal two or four wavefronts, and are not conflicts at all. Sub-word types go the other way: lanes reading neighbouring half or char values share words and are served by broadcast.

These rules fit in a dozen lines of Python, which is the fastest way to check a layout before writing CUDA. The model below is not a cycle-accurate description of the arbiter, but it reproduces the wavefront counts Nsight Compute reports for the common patterns, and every wavefront count quoted here came from running it.

def wavefronts(addrs, width):
    """addrs: 32 byte addresses for one warp-wide shared access; width in bytes."""
    lanes_per_phase = {1: 32, 2: 32, 4: 32, 8: 16, 16: 8}[width]
    total = 0
    for start in range(0, 32, lanes_per_phase):
        words_in_bank = {}
        for a in addrs[start:start + lanes_per_phase]:
            for w in range(a // 4, (a + width - 1) // 4 + 1):
                words_in_bank.setdefault(w % 32, set()).add(w)
        total += max(len(s) for s in words_in_bank.values())   # worst bank sets the cost
    return total

lanes = range(32)
wavefronts([4 * t for t in lanes], 4)        # 1   row of floats
wavefronts([128 * t for t in lanes], 4)      # 32  column of tile[32][32]
wavefronts([132 * t for t in lanes], 4)      # 1   column of tile[32][33]
wavefronts([0 for t in lanes], 4)            # 1   broadcast
wavefronts([8 * t for t in lanes], 8)        # 2   doubles: ideal, two phases
wavefronts([16 * t for t in lanes], 16)      # 4   float4: ideal, four phases

Worked example: the transpose tile

The canonical case is a tiled matrix transpose, which also appears inside data layout conversions and some attention backward passes. A block of 32 by 8 threads loads a 32 by 32 tile from global memory row by row, which is coalesced, and writes it out column by column so that the global stores are coalesced too. The shared writes go along a row and are conflict-free. The shared reads go down a column: lane t reads tile[t][k], whose address is (t * 32 + k) * 4. The stride is 32 words, every lane lands in the same bank, and the simulator says 32 wavefronts for every read instruction.

Here is the kernel with the usual fix already applied. The only change from the conflicting version is TILE + 1 in the shared declaration.

#define TILE 32
__global__ void transpose(float* out, const float* in, int n) {
    __shared__ float tile[TILE][TILE + 1];            // +1 word of padding per row
    int x = blockIdx.x * TILE + threadIdx.x;
    int y = blockIdx.y * TILE + threadIdx.y;
    for (int j = 0; j < TILE; j += blockDim.y)        // blockDim = (32, 8)
        if (x < n && y + j < n)
            tile[threadIdx.y + j][threadIdx.x] = in[(y + j) * n + x];
    __syncthreads();
    x = blockIdx.y * TILE + threadIdx.x;              // swap the block indices
    y = blockIdx.x * TILE + threadIdx.y;
    for (int j = 0; j < TILE; j += blockDim.y)
        if (x < n && y + j < n)
            out[(y + j) * n + x] = tile[threadIdx.x][threadIdx.y + j];   // column read
}

Fix one: padding

Padding changes the row pitch from 32 words to 33. The column read now has stride 33, which is odd, so gcd(33, 32) = 1 and the conflict disappears. The cost is one wasted word per row, 128 bytes for this tile, about 3 percent. For a tile of half values, padding by two elements is the usual choice, since it keeps rows 4-byte aligned for half2 access; for 16-byte vectors you pad by a whole 16-byte chunk so the vector stays aligned.

Padding is the right first move when the kernel uses plain scalar loads and the shared buffer is private to your code. It stops being a good choice in three situations. First, when you want 16-byte vector loads, an odd pitch misaligns every row after the first. Second, when the shared tile is written by the Tensor Memory Accelerator, which copies a dense box with no notion of a row gap. Third, when occupancy is limited by shared memory and the extra bytes cost a resident block.

Fix two: XOR swizzling

Swizzling keeps the dense layout and permutes where each element lives inside its row. For the 32 by 32 float tile, store element (r, k) at column k ^ r. A row write still touches 32 distinct columns, because XOR with a fixed r is a permutation of 0 to 31. A column read for fixed k touches column k ^ t for t = 0 to 31, again a permutation, so every lane hits a different bank. No memory is wasted and alignment is untouched.

__shared__ float tile[TILE][TILE];                          // dense, no padding
tile[threadIdx.y + j][threadIdx.x ^ (threadIdx.y + j)] = in[(y + j) * n + x];
__syncthreads();
out[(y + j) * n + x] = tile[threadIdx.x][(threadIdx.y + j) ^ threadIdx.x];

Real kernels swizzle at the granularity of their widest access. Tensor core operand tiles are loaded with ldmatrix, where each lane supplies the address of one 16-byte row segment. If the tile is stored row-major with 128-byte rows, 64 halves, the eight segments of one 8 by 8 matrix are 128 bytes apart and all land in banks 0 to 3; the simulator reports 32 wavefronts where 4 is ideal, an 8-way conflict. XOR the 16-byte chunk index with row % 8 and the same load costs 4 wavefronts. This is exactly the pattern that CUTLASS expresses as a Swizzle<B, M, S> layout functor (XOR B bits taken from above bit M+S into the bits at M), that Triton chooses automatically for the shared operands of tl.dot, and that TMA applies in hardware through the CU_TENSOR_MAP_SWIZZLE_32B, _64B and _128B modes of cuTensorMapEncodeTiled. When TMA writes a tile with a swizzle mode, every reader must apply the identical XOR, or the kernel computes on scrambled data without any error.

Measuring conflicts in Nsight Compute

Confirm before and after every change. Nsight Compute exposes the counts directly. The most reliable signal is the ratio of wavefronts to instructions, compared with the ideal for the access width: about 1 for 32-bit accesses, 2 for 64-bit and 4 for 128-bit.

ncu --kernel-name transpose --metrics \
smsp__sass_inst_executed_op_shared_ld.sum,\
l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.sum,\
l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum,\
smsp__sass_inst_executed_op_shared_st.sum,\
l1tex__data_pipe_lsu_wavefronts_mem_shared_op_st.sum,\
l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum \
./transpose_bench

For the unpadded transpose, load wavefronts should come out at roughly 32 times the load instructions; with padding or swizzling they fall to roughly 1 times. The dedicated l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld counter is useful but is known to over-count, because it also records some arbitration stalls that are not conflicts in your access pattern, so do not chase a small non-zero value. To find the guilty line, collect a full report and open the Source page, which attributes an L1 Wavefronts Shared Excessive count to each SASS and source line. The Nsight guide for LLM workloads covers report collection on real training jobs.

Where conflicts hide in training kernels

Few people write a transpose for a living, but the same patterns sit inside the kernels that dominate training time. GEMM and attention kernels stage A, B, Q, K and V tiles in shared memory and feed tensor cores with ldmatrix or TMA plus warpgroup instructions; their tile layouts are swizzled precisely so that these loads run conflict-free, which is one reason FlashAttention and library GEMMs are hard to beat by hand. Global-to-shared copies should be coalesced on the global side and conflict-free on the shared side at the same time, and the swizzle is what reconciles the two orders.

The places where custom kernels lose time are more mundane. Reductions that store per-warp partials with a stride of 32 floats, softmax or layer norm kernels that keep a two-dimensional scratch array indexed column-first, radix and histogram kernels in top-k sampling that hammer a few hot bins with shared atomics, and mixture-of-experts permutation kernels that gather rows through an index list. Atomics to the same address serialise for a different reason, contention rather than banking, and a privatised per-warp histogram is the usual fix.

Failure modes

• Swizzle mismatch. Writer and reader use different XOR functions, or the TMA descriptor uses 128B mode while the consumer assumes 64B. There is no fault, just wrong numbers. Unit test each layout with an identity matrix.
• Padding that breaks vector loads. A pitch of 33 floats makes float4 loads misaligned on every other row; the compiler either falls back to scalar loads or the kernel faults on a misaligned address.
• Fixing the wrong bottleneck. A kernel limited by global bandwidth does not get faster when its shared loads drop from 2 to 1 wavefront. Check the memory workload analysis first.
• Occupancy loss from padding. Extra bytes per buffer, multiplied by multi-stage pipelines, can push a block over the per-SM shared budget and reduce resident blocks.
• Trusting a single counter. The bank-conflict counter over-counts; use the wavefront to instruction ratio and the Source page together.
• Layouts that change with the tile size. k ^ r only permutes cleanly when the tile width is a power of two; a 48-wide tile needs a different scheme.

Trade-offs

Padding is the easiest to read and review; swizzling is what production libraries use because it composes with vector loads and asynchronous copies. Sometimes the cheapest fix is not a layout change at all: interchange the loops so lanes walk along rows, or replace a shared-memory reduction with __shfl_down_sync and avoid the banks entirely.

What to do next

1. Write the address of every shared access in terms of threadIdx.x and compute its word stride; flag any even stride.
2. Paste the address expressions into the simulator and record the expected wavefronts per instruction.
3. Profile with the six metrics above and compare wavefronts to instructions against the ideal for each access width.
4. Open the Source page to find the lines with excessive shared wavefronts.
5. Apply padding for scalar kernels, or an XOR swizzle at the width of your widest access for vectorised and tensor core kernels.
6. If TMA fills the tile, pick the matching swizzle mode and test the reader with an identity matrix.
7. Re-measure end-to-end kernel time, not only the counter, and keep the change only if the kernel got faster.