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
| Scenario | Problem | Fix |
|---|---|---|
| Matrix column access | Stride = matrix width (power of 2) | Pad shared array: shared[N][N+1] |
| Struct array | Struct size = power of 2 bytes | Pad struct or use SoA layout |
| Reduction tree | Half-warp accesses same bank | Use sequential addressing, not interleaved |
| Histogram | Multiple threads update same bin | Use privatization, then merge |
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.
Related Topics
Explore 500+ Semiconductor & AI Topics
From EUV lithography to CUDA optimization — search the full knowledge base or chat with our AI assistant.