GPU Shared Memory Bank Conflict

GPU Shared Memory Bank Conflicts are the performance penalties that occur when multiple threads in a warp simultaneously access different addresses that map to the same shared memory bank — forcing the accesses to be serialized rather than served simultaneously, reducing effective bandwidth by a factor equal to the degree of the conflict, and representing one of the most common and impactful GPU optimization targets.

Shared Memory Bank Architecture

• Shared memory is divided into 32 banks (matching warp size).
• Banks are interleaved: Address 0 → Bank 0, Address 4 → Bank 1, ..., Address 124 → Bank 31, Address 128 → Bank 0.
• (For 4-byte words: bank = (address / 4) % 32.)
• Each bank can service one address per cycle.
• No conflict: All 32 threads access 32 different banks → 1 cycle (full bandwidth).
• N-way conflict: N threads access same bank, different addresses → N cycles (serialized).
• Broadcast: Multiple threads access SAME address in same bank → 1 cycle (broadcast, no conflict).

Conflict Examples

// No conflict — stride 1 (consecutive access)
shared[threadIdx.x]           // thread 0→bank 0, thread 1→bank 1, ...

// 2-way conflict — stride 2
shared[threadIdx.x * 2]       // thread 0→bank 0, thread 16→bank 0 (conflict!)

// 32-way conflict — stride 32 (worst case)
shared[threadIdx.x * 32]      // ALL threads hit bank 0 → fully serialized

// No conflict — stride that is odd
shared[threadIdx.x * 3]       // Odd stride → all banks hit uniquely

Bank Conflict Rule

• Conflict occurs when: stride is a multiple of any power of 2 that divides 32.
• No conflict when: stride is odd (coprime with 32).
• Stride 1: No conflict. Stride 2: 2-way. Stride 4: 4-way. Stride 32: 32-way.
• Stride 3: No conflict. Stride 5: No conflict. Stride 7: No conflict.

Common Conflict Scenarios and Fixes

Padding Technique (Most Common Fix)

// Problem: 32x32 matrix, column access = stride 32 = 32-way conflict
__shared__ float tile[32][32];       // column access: 32-way conflict

// Fix: Pad each row by 1 element
__shared__ float tile[32][32 + 1];   // column access: stride 33 (odd) → no conflict!

Diagnosing Bank Conflicts

• NVIDIA Nsight Compute: Reports shared memory bank conflicts per kernel.
• Metric: `l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_{load,store}`.
• Target: 0 conflicts. Acceptable: < 1 conflict per instruction on average.

GPU shared memory bank conflicts are one of the most frequent micro-architectural performance pitfalls — a single line of code using a power-of-2 stride can reduce shared memory throughput by 32x, making bank conflict analysis and padding/layout optimization essential skills for GPU performance engineers.