01 · The memory you manage
Shared memory is a scratchpad you schedule yourself
Day 07 ended on a tax: every trip to global memory costs hundreds of cycles, and coalescing only makes each trip efficient — it doesn't remove the trip. Shared memory removes the trip. Each of the 4070's 46 SMs carries a block of on-chip SRAM — up to 100 KB per SM — that you allocate explicitly with __shared__.
Think of it as a small whiteboard in the stadium instead of the warehouse across town (day 01's analogy, continued). It's only visible to threads inside the same block, it lives exactly as long as that block, and it's yours: the hardware will not decide what goes there. That is the whole deal — you get control in exchange for doing the management.
LOCATION
on-chip SRAM, per SM
VISIBILITY
one block — invisible across blocks
LIFETIME
exactly the block's lifetime
CAPACITY (RTX 4070)
up to 100 KB per block, 128 KB per SM
LATENCY
~30 cycles — roughly 15× faster than DRAM
MANAGED BY
you, explicitly
02 · A tile in shared memory
Stage it coalesced, read it back any way you like
The canonical pattern, which day 15 will build a whole matmul out of: threads cooperatively load a tile — a small rectangular chunk — from global memory into shared memory using a coalesced pattern, synchronize, then read the tile back in whatever pattern the algorithm actually wants. Coalescing handled the highway; the whiteboard absorbs the awkward accesses.
__shared__ float tile[16][16];
int tx = threadIdx.x; // 0..15
int ty = threadIdx.y; // 0..15
// stage: consecutive tx → consecutive addresses → coalesced
tile[ty][tx] = input[row0 + ty][col0 + tx];
// the wall: nobody reads until everyone has written
__syncthreads();
// now read back in any pattern — it's on-chip now
float v = tile[tx][ty];__syncthreads() is the fence that makes the staging correct — a barrier every thread in the block must reach before any of them continue. Forget it and half your block reads values the other half hasn't written yet. That's a race that sometimes passes tests, which is worse than one that never does.
03 · Thirty-two banks
Shared memory is 32 lanes wide
Shared memory isn't one big array — it's physically sliced into 32 banks, each 4 bytes wide, matching the warp. A warp that reads 32 consecutive floats touches all 32 banks in parallel: one transaction, everyone served. A warp whose 32 addresses land on thesame bank forces the hardware to serve them one at a time: a 32-way bank conflict, and a ~32× slowdown on what should have been a free read.
The rule is arithmetic, not magic: address a (in 4-byte words) lives in bank (a % 32). Two extra cases matter:
BROADCAST — FREE
all 32 lanes read the same address: one fetch, replicated to everyone
MULTICAST — FREE
lanes read a few distinct addresses: each fetched once, fanned out
N-WAY CONFLICT — SLOW
n distinct addresses hit one bank: n serialized trips
THE FIX
pad or rearrange so a warp's addresses spread across banks
04 · The classic accident
Row access is free. Column access is a trap.
Declare __shared__ float tile[32][32] and read down a column — tile[tx][j] for tx = 0…31. The stride between neighbors is 32 floats, and 32 % 32 == 0, so every thread in the warp lands on the same bank. Thirty-two-way conflict. The layout that made row access coalesced made column access serialize.
// 32-way conflict: column stride is 32 words → same bank
__shared__ float bad[32][32];
float v = bad[tx][j];
// the pad: +1 column shifts every row's start by one bank
__shared__ float good[32][32 + 1];
float w = good[tx][j]; // stride 33 → conflict-free
The pad costs almost nothing — a few percent more SRAM — and turns a 32× serialized read back into a parallel one. You'll see [32][33] in real kernels everywhere (day 15's matmul includes it); now you know it isn't superstition. Transpose-style kernels exist almost entirely to exercise this trap, and production kernels pad unprompted.
05 · When it matters
Conflicts tax the fast path, not the slow one
Perspective: a 32-way conflict makes one shared read ~32× slower — but even the slowed read is still faster than a trip to DRAM. Bank conflicts rarely turn a fast kernel slow on their own; they eat the margin shared memory just bought you. In a kernel that reloads a tile hundreds of times (day 15's matmul, day 24's attention), that margin is the whole game.
The profiler sees this directly — Nsight reports shared bank conflicts per kernel, and we'll read that number on day 18. If a kernel is "using shared memory correctly but not getting faster," conflicts are the first suspect.
- 01Shared memory is per-block, on-chip, and yours to manage. Up to 100 KB per SM on the 4070, ~30-cycle latency. You allocate it, you synchronize it, you own its layout.
- 02Stage coalesced, read back freely. Load tiles with coalesced writes, __syncthreads(), then access in whatever pattern the algorithm needs. This is the workhorse pattern of days 15 and 24.
- 0332 banks, 4 bytes each — spread a warp's addresses across them. Consecutive floats: conflict-free. Same-address reads broadcast for free. n addresses on one bank: n-way serialization.
- 04Pad row-major arrays when you read columns. tile[32][33] instead of tile[32][32] shifts every row one bank over and erases the conflict for a few percent more SRAM.
Next: day 09 — the even faster memory inside shared memory: registers, the per-thread budget that decides your occupancy, and what spilling quietly does to you.