Shared Memory and Bank Conflicts
REFERENCE CONTEXT · CUDA 12.2; RTX 2070 SUPER; Turing sm_75
Core Concept
Shared memory is divided into banks. When lanes in a warp access different addresses in the same bank, the operation is split; broadcasts from the same address are a special case.
Why It Affects Performance
Shared memory is meant to enable reuse and reduce DRAM traffic, but severe bank conflicts can serialize access and erase that benefit.
Example or Pseudocode
__shared__ float tile[32][33];
tile[threadIdx.y][threadIdx.x] = in[...];
__syncthreads();
float v = tile[threadIdx.x][threadIdx.y];The extra column changes bank mapping and commonly removes conflicts during a 32x32 transpose.
Profiling and Performance Judgment
Compare shared load/store transactions, conflict metrics, short-scoreboard stalls, and runtime before and after padding.
Five-Minute Exercise
Why can column reads from float tile[32][32] conflict while [32][33] often does not?
Show answer
A 32-element row stride repeats the same bank mapping for a column. Adding one element rotates each row's starting bank.