What is the GPU memory hierarchy?
The stack of storage a kernel can reach, from registers through shared memory and caches out to global memory, each one bigger and slower than the last.
The useful question is not which level is fastest. It is who can see a value and who put it there, because those two answers decide where a variable lands and you rarely get to override them. A plain local variable lives in registers, private to one thread, allocated by the compiler. An array that thread indexes with a runtime value spills to local memory, still private, now in DRAM. An array marked __shared__ is one copy per thread block, on the SM, and you fill it yourself. L1 and L2 are filled by the hardware and you never address them. Constant memory is written by the host and read by everyone. Global memory is the only level the host can copy into and the only one that survives the launch.
Two of the six names are lies and both cost people a day. Local memory is not local to anything on chip: it is global DRAM with per-thread addressing, which is why an indexed private array is not a cheap way to hold state. Shared memory is not shared with the grid, only within one block, which is why a barrier inside a block is enough to order it and nothing orders it between blocks.
The pyramid diagram everyone copies also gets the top two rungs backwards on real hardware. A Tesla T4 SM holds 65,536 32-bit registers and 64 KiB of shared memory, so the register file is the bigger of the two in bytes. Registers sit at the top because they are private and free to address, not because they are scarce. What is scarce is registers per thread: 65,536 across the 1,024 threads an SM can hold works out at 64 each, and a kernel wanting more gets fewer threads resident, which is the occupancy trade day 17 measures.
The reason to know the ladder is that most kernel tuning is choosing which level absorbs an awkward access pattern. A transpose has to walk a column somewhere. Do it in global memory and each lane of a warp pulls its own 32-byte sector. Do it in shared memory and the same walk costs bank conflicts instead, which are cheaper and, unlike sectors, fixable by padding. Nothing was cached, nothing was reused, and the same bytes crossed the same bus either way. Only the level changed.
Measured
Every figure below came off one Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), built with nvcc -O3 -arch=sm_75, on the project's verification node, captured 2026-08-30. Capacities are what the device reported to the runtime; the right-hand column is what a kernel in the named lesson actually got.
| Level | Visible to | Capacity on this T4 | Measured |
|---|---|---|---|
| Registers | one thread | 65,536 32-bit per SM, so 64 per thread to keep 1,024 threads resident | a 10-register kernel ran in 0.038 ms at 100% occupancy; a 72-register one in the same program ran at 75%, 0.164 ms (day 17) |
| Local memory | one thread, in DRAM | no separate budget | that 72-register kernel rewritten with its taps in an indexed private array spent 256 bytes per thread of local and 1.556 ms for the identical answer (day 17) |
| Shared memory | one block | 64 KiB per SM, 49,152 bytes per block by default, 65,536 opt-in | a 4,096-byte tile took a transpose from 65.2 to 118.6 GB/s (day 13); a strided read inside it cost 0.286 ms at stride 1 and 8.886 ms at stride 32 (day 15) |
| L2 cache | the whole device | 4,096 KiB | this card reports a persisting-L2 maximum of 0 bytes and a policy-window maximum of 0 bytes, so the access-policy controls are not available on it (day 18) |
| Constant memory | every thread, read-only | 65,536 bytes | a 32-tap filter using 128 of them ran at 0.68 times the global-memory version when all lanes read one address, and 15.08 times slower when each lane read its own (day 18) |
| Global memory | every thread, and the host | 14,912 MiB | 232.9 GB/s with consecutive lanes on consecutive floats, 9.6 GB/s at stride 32, identical bytes moved (day 11) |
There is no latency column, and that is deliberate rather than an omission. This course measures wall time and delivered bandwidth on kernels it ships, not cycle counts from a pointer-chase microbenchmark, so it has no honest per-level latency figure to publish and will not copy one. The table above is what the levels cost when real kernels used them.
Read the two shared-memory entries together. The level is faster than global memory and still has a 31.10x spread inside itself depending on the addresses. Picking the right level is the first decision; the access pattern within it is a second, separate one, and skipping the second gives back most of what the first won.
Diagram
memory-hierarchy-svg: registers and shared memory drawn inside one SM box with L1 beside them, the 4,096 KiB L2 as a single band spanning all 40 SMs, and global memory as an off-chip band below it. The register box is drawn wider than the shared-memory box, matching the T4's actual figures, and each band carries the capacity and the measured number from the table above.
Alt text: "The memory hierarchy on a Tesla T4. Registers and shared memory sit inside the SM, and on this card the 65,536-register file holds more bytes than the 64 KiB of shared memory beside it. The 4096 KiB L2 is shared by all 40 SMs, and 14,912 MiB of global memory sits off chip, delivering between 9.6 and 232.9 GB/s depending on the access pattern."
Related terms
- registers
- local memory
- shared memory
- L1 and L2 cache
- constant memory
- global memory
- distributed shared memory
Where you meet this
- Day 11, CUDA memory coalescing, for the bottom rung and how much its access pattern is worth.
- Day 13, CUDA shared memory and tiling, the lesson that owns this term and moves an access from one level to another.
- Day 17, registers, local memory and spills, for the top rung and the level that pretends to be it.
- Day 18, CUDA constant memory, for the read-only path and the caches behind it.
- How to set up CUDA, which prints your own card's register file, shared memory, L2 and global memory.
too many resources requested for launch, what a launch returns when a block asks the SM for more registers or shared memory than it has.
Sources
- CUDA Programming Guide, on the memory spaces a kernel can address: https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/writing-cuda-kernels.html (checked 2026-08-30)
- Compute Capabilities appendix, Tables 30 and 31, for register file and shared memory per SM and per block on every architecture: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/compute-capabilities.html (checked 2026-08-29)
- CUDA C++ Best Practices Guide, on picking a memory space: https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html (checked 2026-08-30)
- Turing Tuning Guide, for the L1, L2 and shared-memory arrangement on compute capability 7.5: https://docs.nvidia.com/cuda/turing-tuning-guide/index.html (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.