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**
| 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)**
```cuda
// 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.