Shared Memory & Bank Conflicts
Using shared memory effectively — bank structure, conflict types, matrix transpose optimization, padding, and swizzle patterns.
Shared memory is the programmer’s L1 cache — fast, on-chip, and shared across threads in a block. But accessing it carelessly creates bank conflicts that serialize what should be parallel reads. This chapter covers the bank structure, common conflict patterns, and how to fix them.
1. Bank Structure
Shared memory is divided into 32 banks, each 4 bytes wide. Banks are interleaved:
Address 0- 3 → Bank 0
Address 4- 7 → Bank 1
Address 8-11 → Bank 2
...
Address 124-127 → Bank 31
Address 128-131 → Bank 0 (wraps around)
Address 132-135 → Bank 1Bank index = (byte_address / 4) % 32
When multiple threads in a warp access different addresses in the same bank, the accesses are serialized — a bank conflict.
2. Conflict Types
No Conflict: Different Banks
// Each thread accesses a unique bank
shared[threadIdx.x] = val; // thread 0→bank 0, thread 1→bank 1, ...32 threads, 32 banks, no conflicts → 1 cycle.
No Conflict: Broadcast
// All threads read the SAME address
float val = shared[0]; // all threads read bank 0, address 0Hardware broadcasts to all threads → 1 cycle.
2-Way Conflict
// Stride of 2: threads 0 and 16 both hit bank 0
shared[threadIdx.x * 2] = val;
// thread 0 → bank 0, thread 1 → bank 2, ..., thread 16 → bank 0!Two threads per bank → 2 cycles (2× slower).
N-Way Conflict
// Stride of 32: ALL threads hit bank 0
shared[threadIdx.x * 32] = val;
// Every thread accesses bank 0 → 32-way conflict → 32 cycles!The Column Access Problem
The most common conflict in practice — accessing a 2D array by columns:
__shared__ float tile[32][32];
// ROW access: thread i reads tile[row][i] — consecutive banks, no conflict ✅
float val = tile[row][threadIdx.x];
// COLUMN access: thread i reads tile[i][col] — stride 32, every thread hits same bank ❌
float val = tile[threadIdx.x][col]; // 32-way conflict!3. Classic Example: Matrix Transpose
Transpose reads rows (coalesced from global, no bank conflict in shared) but writes columns (bank conflicts):
Naive Transpose (Bank Conflicts)
__shared__ float tile[32][32];
// Load: coalesced global read, sequential shared write — no conflict
tile[threadIdx.y][threadIdx.x] = input[gy * N + gx];
__syncthreads();
// Store: transpose — column read from shared → 32-way bank conflict!
output[gx * N + gy] = tile[threadIdx.x][threadIdx.y]; // ❌Fixed with Padding
__shared__ float tile[32][33]; // 33 instead of 32!
// Now column access: thread i reads tile[i][col]
// tile[0][col] → bank (0*33 + col) % 32
// tile[1][col] → bank (1*33 + col) % 32
// Each thread hits a different bank because 33 and 32 are coprime → no conflict ✅
tile[threadIdx.y][threadIdx.x] = input[gy * N + gx];
__syncthreads();
output[gx * N + gy] = tile[threadIdx.x][threadIdx.y]; // no conflict!The padding trick: Adding 1 element to the row width shifts each row’s bank alignment, ensuring column accesses hit different banks.
Cost: 1 extra float per row × 32 rows = 32 floats = 128 bytes wasted. Negligible compared to the performance gain.
4. Swizzle Patterns (Advanced)
Padding wastes shared memory. Swizzle remaps addresses without wasting space:
// Instead of linear indexing:
int idx = row * 32 + col;
// Swizzle: XOR the row into the column index
int idx = row * 32 + (col ^ row); // conflict-free for many access patternsSwizzle is used in production GEMM kernels (cuBLAS, CUTLASS) where shared memory is precious. For most AI infra work, padding is sufficient.
5. Detecting Bank Conflicts
Nsight Compute
ncu --metrics l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum kernelThe metric shows total bank conflict events. Zero = no conflicts.
Quick Test
Compare kernel time with and without padding. If padding gives >10% speedup, you had significant conflicts.
6. When Shared Memory vs Registers vs Global
| Situation | Best choice | Why |
|---|---|---|
| Per-thread private data | Registers | Fastest, no sharing needed |
| Data reused by multiple threads in block | Shared memory | Load once from HBM, reuse many times |
| Data too large for shared memory | Global memory with L2 | Falls through to HBM |
| All threads read same constant | Constant memory | Cached broadcast |
| Reduction across threads | Warp shuffle first, then shared | Warp shuffle is faster for intra-warp |
Self-Check
- Calculate the bank index for a given byte address
- Explain why stride-2 access causes 2-way conflicts
- Explain the column access problem in a 32×32 shared array
- Fix bank conflicts using the padding trick (
[32][33]) - Implement conflict-free matrix transpose using shared memory + padding
- Use Nsight Compute to detect bank conflicts