# Memory model (CUDA)
A GPU has several distinct memory spaces with very different sizes, speeds, and scopes. The **memory model** determines which threads can see a piece of data and how fast they can reach it. Choosing the right space is usually a bigger performance lever than anything done to the arithmetic.

^ Space ^ Scope ^ Latency ^ Declared as ^
| Registers | one thread | ~1 cycle | plain local variables |
| [[shared-memory-cuda|Shared]] | one block | ~30 cycles | `__shared__` |
| [[global-memory-cuda|Global]] | whole grid, plus host | 400-800 cycles | `cudaMalloc` |
| [[constant-memory-cuda|Constant]] | whole grid, read-only | ~1 cycle cached | `__constant__` |
| Local | one thread | same as global | spilled locals and indexed arrays |

Registers are the fastest and are allocated per thread by the compiler. A kernel that uses too many of them limits how many warps fit on a multiprocessor ([[occupancy-cuda]]), and past that limit the compiler spills the excess to local memory.

Local memory is a trap in the naming. It is not close to the thread and it is not fast. It lives in the same device DRAM as global memory, and it only exists because registers ran out or because an array was indexed with a value the compiler could not resolve at compile time.

```c
__global__ void k(float *out) {
    float tmp[8];        // stays in registers only if indexed with constants
    int i = threadIdx.x;
    tmp[i % 8] = 1.0f;   // dynamic index, so tmp is pushed to local memory
}
```

Scope also determines lifetime. Registers and shared memory die when the block finishes, so anything that has to survive a kernel launch lives in global memory. That constraint is why multi-phase GPU algorithms are written as a sequence of kernel launches, with global memory as the only thing carried between them.