What is tiling in CUDA?
Loading a block-sized piece of the input into shared memory once, then reusing it from there instead of going back to global memory.
That definition is the textbook one and it describes half the cases. A tile pays for itself two different ways, and only one of them is reuse. The reuse case is the matrix multiply everybody teaches: an element pulled into the tile is used once per row of the output tile, so a tile of side 32 cuts global traffic by roughly 32 and lifts arithmetic intensity by the same factor. The other case saves no traffic whatsoever, and it is the one that catches people out.
A transpose is that second case. It reads every element once and writes it once, tile or no tile, so the tiled kernel moves exactly the bytes the naive kernel moved. What the tile changes is which addresses arrive together. out[c][r] = in[r][c] means one of the two ends must walk down a column, and a column walk on the memory bus costs a sector per lane instead of a sector per eight lanes. Put a square in shared memory, read the input along its rows into the tile, then write the output along its rows out of the tile, and the column walk happens on the SM where there is no sector to waste. Same bytes, different order, and day 13 watched 65.2 GB/s become 118.6.
This is why a checker that only compares numbers cannot teach the pattern. Someone working through GPU Puzzles put it exactly: "I can pass the two tests but I haven't use the shared memory. I feel that my solution is not correct but I don't know why the shared memory is needed here." They were right on both counts. Correct-and-slow and correct-and-fast are two different outcomes and only a stopwatch tells them apart.
Two details decide whether a tile survives contact with real data. The first is the ragged edge. Tutorial code sizes the matrix as a multiple of the tile and then breaks on anything else, so day 13 uses 2003 by 3001 on purpose: neither side divides by 32, both loops guard the load and the store, and the guard never sits above the barrier, because a thread that returns early stops participating in __syncthreads() and leaves holes in the tile for someone else to read. The second is the tile's own shape. A [32][32] float tile read down a column puts all 32 lanes on one bank, which is a 32-way bank conflict that costs real time and is fixed by making the tile one column wider.
Measured
Day 13 ran a 2003 by 3001 float transpose with and without a 32 by 32 tile, the tile costing 4,096 bytes per block. Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), built with nvcc -O3 -arch=sm_75, captured 2026-08-30 on the project's verification node.
| Kernel | Time (ms) | GB/s |
|---|---|---|
copyRowMajor, the ceiling |
0.228 | 211.2 |
transposeNaive |
0.737 | 65.2 |
transposeTiledStatic |
0.406 | 118.6 |
The tiled kernel does not reach the copy and was never going to, because a copy has no column walk to place anywhere. What it does is cut most of the distance to that ceiling while moving the same bytes through the same bus. Reuse was not available on this kernel at all, so every bit of that came from address order.
What is left over is the tile's own shape, and day 15 measured the rest on a 4096 by 4096 transpose on the same card: a copy at 236.6 GB/s, tile[32][32] at 135.1 GB/s (57.1 percent of the copy), and tile[32][33] at 195.2 GB/s (82.5 percent). One extra column moved that kernel from 135.1 to 195.2. Tiling and the tile's layout are two decisions, and taking the first without the second leaves most of the win behind.
Diagram
tile-lifecycle, preset tile-correct with the tile size slider set to a 32 by 32 square: the five phases of global load, barrier, compute, barrier, global store, with the load phase drawn as rows arriving in parallel and the store phase drawn as a column being read out.
Alt text: "A tile's life inside one block. Thirty-two lanes load a row of the tile from global memory in one coalesced request, the barrier makes those writes visible to the block, and the store phase reads the tile down a column so the global write is coalesced too."
Code
From code/day13-shared-memory/shared_memory.cu. The two accesses that make it a transpose rather than a copy, with the barrier between them:
tile[threadIdx.y + j][threadIdx.x] = in[row * cols + col]; // rows in
__syncthreads();
out[row * rows + col] = tile[threadIdx.x][threadIdx.y + j]; // columns out
The index pair is swapped on the read, and the block indices behind row and col are recomputed from the other axis after the barrier, because the block that read the square at (x, y) writes the square at (y, x). Reuse the input's indices there and you get a program that compiles, runs, and transposes nothing.
Related terms
Where you meet this
- Day 13, CUDA shared memory and tiling, the lesson that owns this term and measured the table above.
- Day 15, shared memory bank conflicts, which changes only the width of the tile.
- Day 16, CUDA tiled matrix multiplication, where a tile finally buys reuse instead of only order, and where one barrier per loop iteration stops being enough.
- Day 20, CUDA image convolution, where the tile grows a halo because each output needs its neighbours.
invalid argument, what a launch returns when the tile you asked for is larger than the per-block default.
Sources
- CUDA Programming Guide, shared memory and the tiled transpose: https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/writing-cuda-kernels.html (checked 2026-08-30)
- CUDA C++ Best Practices Guide, on staging global accesses through shared memory: https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html (checked 2026-08-30)
- "I can pass the two tests but I haven't use the shared memory": https://github.com/srush/GPU-Puzzles/issues/18 (checked 2026-08-29)
cuda-samples, NVIDIA's transpose ladder, which is the same kernel with the same tile: https://github.com/NVIDIA/cuda-samples/tree/master/cpp/6_Performance/transpose (checked 2026-08-29)- Programming Massively Parallel Processors, 4th edition, chapter 5, on tiling and the memory hierarchy: https://shop.elsevier.com/books/programming-massively-parallel-processors/hwu/978-0-323-91231-0 (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.