What is shared memory in CUDA?
Fast per-block scratch memory on the SM, which you fill yourself and which vanishes when the block ends.
Declare an array __shared__ and every thread block gets its own copy, sitting on the SM the block landed on. Nobody else can see it. The block next door has a different copy, the host cannot touch it, and when the block retires the storage goes back to the pool for the next block. That last part is the piece people skip: shared memory is not a cache. Nothing loads it for you, nothing evicts from it, and it holds whatever the previous block on that SM left behind. A learner on GPU Puzzles hit exactly that and asked whether it was a leak, having watched a value "carried over from Block 0". It was not a leak. It was an array they had not written.
Capacity is per SM, and the number you are allowed to ask for per block is a different number. On a Tesla T4 the SM has 64 KiB, a single block gets 48 KiB (49,152 bytes) without asking, and 64 KiB (65,536 bytes) is available if you opt in per kernel with cudaFuncSetAttribute. Table 31 of the compute-capability appendix splits the same two columns for every architecture and the gap is where most pages go wrong: 228 KB per SM on 9.0 and 100 KB on 12.x, against per-block maxima of 227 KB and 99 KB. Quote the per-SM figure as a per-block budget and your tile is 1 KB too big on an H100. The T4 hides this, because its two figures happen to agree at 64 KiB.
Asking for a lot has a second cost that has nothing to do with the wall. The SM's 64 KiB is divided among the blocks resident on it, so a block that claims half of it caps that SM at two blocks and occupancy falls with it. Day 13's tile is 32 by 32 floats, 4,096 bytes, so twelve of them would fit one block's default allowance and the SM keeps room for every block it can hold.
What shared memory buys is not fewer bytes. It is the order in which you touch global memory. A transpose reads every element once and writes it once whether or not there is a tile in the middle, so the tiled kernel moves the identical traffic. What changes is that the strided walk moves off the memory bus, where it costs sectors, and onto the SM, where it costs banks instead. That trade is only worth it if you then look at the banks: day 13's unpadded tile[32][32] puts all 32 lanes of the column read on one bank, a 32-way bank conflict that day 15 removes with one extra column. And because two threads now touch the same cell, every tile needs __syncthreads() between the write and the read.
Measured
On a Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), built with nvcc -O3 -arch=sm_75, day 13 transposed a 2003 by 3001 float matrix, 22.9 MiB per buffer, four ways. Captured 2026-08-30 on the project's verification node.
| Kernel | Time (ms) | GB/s |
|---|---|---|
copyRowMajor, no transpose |
0.228 | 211.2 |
transposeNaive |
0.737 | 65.2 |
transposeTiledStatic, tile[32][32] |
0.406 | 118.6 |
transposeTiledDynamic, same tile, sized at launch |
0.407 | 118.2 |
The tile is 4,096 bytes per block. Every row read and wrote every element once, so the whole gap between the naive kernel and the tiled one was bought with address order and nothing else. Static and dynamic land on the same number, 118.6 against 118.2, which settles the only question people ask about the two forms: pick on whether the size is known at compile time, not on speed.
The device's own limits came out of the same run: 49,152 bytes per block by default, 65,536 opt-in, 65,536 per SM. Asking a kernel for 53,248 bytes of dynamic shared memory failed at launch with invalid argument, a string that names no shared memory at all, and the same request succeeded after cudaFuncSetAttribute with cudaFuncAttributeMaxDynamicSharedMemorySize. A launch that goes over the register budget instead reports too many resources requested for launch, so the two failures read differently.
Diagram
tile-lifecycle, preset tile-correct: one block, one tile, the five phases of global load, barrier, compute, barrier, global store, with the previous block's residue toggle switched on so the tile starts full of another block's leftovers.
Alt text: "One block's shared tile through five phases. The tile starts holding the previous block's values, each thread writes its own cell, the barrier makes those writes visible, and only then does any thread read a cell it did not write."
Code
From code/day13-shared-memory/shared_memory.cu. The two declarations, which differ only in when the size is fixed:
__shared__ float tile[32][32]; // size fixed when you compile
extern __shared__ float tile[]; // size fixed when you launch
The dynamic form takes its byte count from the third argument of the execution configuration, kernel<<<grid, block, bytes>>>(...), and gives you one unsized array per block, so a two-dimensional tile becomes an explicit multiply.
Related terms
- tiling
__syncthreads()- memory bank
- bank conflict
- occupancy
- memory hierarchy
- distributed shared memory
Where you meet this
- Day 12, why transposing a matrix is slow, the kernel this one is measured against.
- Day 13, CUDA shared memory and tiling, the lesson that owns this term and produced the table above.
- Day 14, what
__syncthreads()guarantees, where a tile without a barrier is right at one block size and wrong at three others. - Day 15, shared memory bank conflicts, which pads the same tile and takes back most of the gap to a plain copy.
invalid argument, the launch failure you get for asking past the per-block default.
Sources
- CUDA Programming Guide, shared memory and the dynamic form: https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/writing-cuda-kernels.html (checked 2026-08-30)
- Compute Capabilities appendix, Table 31, which gives shared memory per SM and per thread block as separate columns: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/compute-capabilities.html (checked 2026-08-29)
- CUDA C++ Best Practices Guide, on shared memory and occupancy: https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html (checked 2026-08-30)
- "the value got carried over from Block 0", a learner meeting uninitialized shared memory: https://github.com/srush/GPU-Puzzles/issues/20 (checked 2026-08-29)
cuda-samples, NVIDIA's own transpose ladder over this kernel: https://github.com/NVIDIA/cuda-samples/tree/master/cpp/6_Performance/transpose (checked 2026-08-29)
Byline
Written by: pending. Reviewed by: pending. Written on: pending. Last checked on: pending. This entry stays a draft until a named author and a different named reviewer sign it.