Shared memory and bank conflicts: using on-chip storage correctly
A full derivation of tiled matmul, and that mysterious [TILE][TILE + 1]
- See the value of shared memory as a cache you manage yourself
- Derive tiled matmul from scratch and explain why each line is written that way
- Understand the 32-bank mechanism and what causes a conflict
- Use padding and swizzle as two ways to kill conflicts
Shared memory is a block of fast SRAM on each SM, about 1/20 the latency of global memory, shared by every thread in the block. The big difference from a CPU L1 is that you manage it explicitly. That is both a burden and a weapon: you know your reuse pattern better than any hardware prefetcher.
Start from naive matmul
1__global__ void matmulNaive(const float* A, const float* B, float* C, int N) {2 int col = blockIdx.x * blockDim.x + threadIdx.x;3 int row = blockIdx.y * blockDim.y + threadIdx.y;4 if (row >= N || col >= N) return;5 6 float acc = 0.0f;7 for (int k = 0; k < N; ++k) {8 acc += A[row * N + k] * B[k * N + col]; // 2 global loads and 1 multiply-add per iteration9 }10 C[row * N + col] = acc;11}Do the arithmetic for this kernel: each inner-loop iteration reads 2 floats (8 bytes) and does 2 FLOPs, so arithmetic intensity is 0.25 FLOP/Byte, 1/50 of the hardware balance point. Worse: the same row of A is read once by each of the N threads on that row, and the same column of B is read N times too. Total: 2N³ global accesses for only 2N² distinct values.
B Each step: ┌───┬───┬───┐ 1. Threads cooperate: load an A tile and a B tile into shared mem │ │Bt │ │ 2. __syncthreads() wait until the tile is fully loaded ├───┼───┼───┤ 3. TILE multiply-adds from shared mem, accumulate in registers │ │ │ │ 4. __syncthreads() wait until all are done before the next tile └───┴───┴───┘ A global reads: 2 * N^3 / TILE ┌───┬───┬───┐ ┌───┬───┐ shared mem reads: 2 * N^3 (but ~20x faster) │At │ │ │ │ C │ │ accumulator: stays in registers the whole time ├───┼───┼───┤ ├───┼───┤ │ │ │ │ │ │ │ └───┴───┴───┘ └───┴───┘
1#define TILE 322 3__global__ void matmulTiled(const float* A, const float* B, float* C, int N) {4 __shared__ float sA[TILE][TILE];5 __shared__ float sB[TILE][TILE];6 7 int tx = threadIdx.x, ty = threadIdx.y;8 int row = blockIdx.y * TILE + ty;9 int col = blockIdx.x * TILE + tx;10 11 float acc = 0.0f;12 13 for (int t = 0; t < (N + TILE - 1) / TILE; ++t) {14 // Cooperative load: each thread owns one element. Out of range → 0, so the inner loop needs no bounds check15 int aCol = t * TILE + tx;16 int bRow = t * TILE + ty;17 sA[ty][tx] = (row < N && aCol < N) ? A[row * N + aCol] : 0.0f;18 sB[ty][tx] = (bRow < N && col < N) ? B[bRow * N + col] : 0.0f;19 20 __syncthreads(); // barrier 1: the whole tile is in place21 22 #pragma unroll23 for (int k = 0; k < TILE; ++k) {24 acc += sA[ty][k] * sB[k][tx];25 }26 27 __syncthreads(); // barrier 2: everyone has finished computing, so the next tile can overwrite shared mem28 }29 30 if (row < N && col < N) C[row * N + col] = acc;31}32 banks: the parallel structure of shared memory
Shared memory is split into 32 banks, each 4 bytes wide, interleaved by address: address 0 is bank 0, address 4 is bank 1, ... address 128 wraps back to bank 0. All 32 banks can serve 32 requests in one cycle. When several threads in a warp hit different addresses in the same bank, hardware serializes those requests. That is a bank conflict.
float index: 0 1 2 ... 31 32 33 ...
bank: 0 1 2 ... 31 0 1 ...
└──────── one round of 32 banks ──────┘
no conflict: 32 threads hit 32 different banks → 1 cycle
broadcast: 32 threads hit the same address in one bank → 1 cycle (hardware broadcast, not a conflict)
2-way: every 2 threads hit different addrs in one bank → 2 cycles
32-way: all threads hit different addrs in one bank → 32 cycles (worst case)What the +1 in transpose actually fixes
1__shared__ float tile[32][32];2 3// Column access: threadIdx.x = 0..31, different rows, same column4float v = tile[threadIdx.x][0];5 6// Thread i touches float index = i * 32 + 07// bank = (i * 32) % 32 = 0 ← all 32 threads pile into bank 08// Result: 32-way conflict. A 1-cycle access becomes 32 cycles1__shared__ float tile[32][33]; // one extra float per row2 3float v = tile[threadIdx.x][0];4 5// Thread i touches float index = i * 33 + 06// bank = (i * 33) % 32 = i ← 33 and 32 are coprime, so the accesses fan out across 32 banks7// Result: zero conflicts. Cost is only 32 * 4 = 128 extra bytes of shared memoryThe other path: swizzle
Padding wastes shared memory, and a large tile can squeeze occupancy. The more advanced move is swizzle: keep the storage size, XOR-permute the column index so different rows of the same column land in different banks. CUTLASS and other high-performance GEMM libraries all take this route.
1__shared__ float tile[32][32]; // no padding, not one extra byte2 3__device__ inline int swz(int row, int col) {4 return col ^ (row & 31); // each row's column order is XOR-scrambled5}6 7// Write8tile[ty][swz(ty, tx)] = value;9 10// On a column read, the same logical col in different rows maps to different physical banks11float v = tile[threadIdx.x][swz(threadIdx.x, 0)];Self-check
Answer in your head first, then open the solution. If you cannot, reread that section.