The Address Spaces a GPU Kernel Can Touch

A pointer inside a kernel looks identical to a pointer in host code: eight bytes, arithmetic works, dereferencing produces a value. The resemblance ends there. Behind that single syntax sit four or five physically distinct stores with different owners, different lifetimes, and access latencies spanning two orders of magnitude. The compiler commits to one of them where it emits the load, because the instruction that reads a per-SM scratchpad is not the instruction that reads device DRAM. Getting a kernel to run at a fraction of peak is usually a matter of noticing that the data was placed in the wrong space, or that the compiler could not prove which space a pointer referred to.

One declaration syntax, several physical stores

The C-family GPU languages mark the space on the declaration rather than in the type. A variable declared at file scope for the device lives in global memory: device DRAM, visible to every thread of every block, alive for as long as the module is loaded. Buffers allocated by the runtime at the host’s request also live there and survive until the host frees them. A variable inside a kernel body is a register if the compiler can keep it in one, and otherwise a slot in local memory. A variable marked as shared lives in a scratchpad attached to the streaming multiprocessor (SM) and belongs to one block. A variable marked constant lives in a read-only bank with its own small cache.

__constant__ float coeff[64];        // read-only bank, device-wide

__global__ void scale(float *out, const float *in, int n) {
    extern __shared__ float tile[];  // size chosen at launch
    __shared__ int stamp[256];       // static, per block
    int t = threadIdx.x;
    float acc = 0.0f;                // register, or a local-memory slot
    for (int i = t; i < n; i += blockDim.x)
        acc += in[i] * coeff[blockIdx.x & 63];  // uniform within the block
    tile[t] = acc;                   // banked scratchpad
    __syncthreads();
    out[t] = tile[t ^ 1];            // exchange through shared memory
}

Those declarations cover four different memories — global, constant, shared, and register — and only the register file is not named in the source at all. That asymmetry runs through the whole model: the fastest storage in the machine is the one the language does not allow a kernel to request directly.

Space	Scope	Lifetime	Backing store	Typical size
Registers	One thread	Thread	Register file in the SM	Tens of 32-bit registers per thread
Shared	One block	Block	SRAM in the SM	Tens of KB, up to a couple of hundred kilobytes on recent parts
Local	One thread	Thread	Device DRAM, cached	Bounded by the same DRAM as global
Global	Grid and host	Until freed	Device DRAM	Gigabytes
Constant	Grid and host	Until freed	Device DRAM plus a small cache	64 KB total

The sizes are the design statement. Registers and shared memory are the two spaces a programmer budgets explicitly at compile time — the register count per thread and the scratchpad per block — while global and local memory are sized by the host and by the launch. Constant memory sits apart as a fixed 64 KB bank that neither the programmer nor the launch can resize.

Why local memory exists at all

The name is misleading. Local memory is not local in the physical sense; it is thread-private in the logical sense, and it is backed by the same DRAM as global memory. A thread that needs more registers than the architecture allows, or that declares an array it indexes dynamically, gets those objects placed in local memory. Loads and stores to them go out over the same memory path as global accesses, cached in L1 and L2 like anything else.

The reason for the indirection is that register allocation happens before the number of threads per block is known. A compiler that promised 128 registers per thread would make a 1024-thread block impossible on an SM with a 64K register file, so instead it caps the register count and spills the excess. Spill slots are laid out so that consecutive threads of a warp occupy consecutive addresses, which means a spill access by a full warp still coalesces into a small number of transactions even though each thread is touching private data. The layout is deliberate: a naive per-thread allocation would turn every spill into a scattered access, as described in register pressure and spilling to local memory. That traffic is invisible in the source and shows up only as a lower achieved bandwidth.

Shared memory is a scratchpad, not a cache

On a CPU, the level-one cache is invisible: the program names an address and the hardware decides whether the line is resident. GPU shared memory inverts that. The program names a location inside a small per-SM array, the hardware does not allocate or evict anything, and the contents persist exactly as long as the block does. Nothing is fetched on a miss because there are no misses.

That inversion is what makes the space useful. A blocked algorithm can load a tile of a matrix into shared memory once, then have every thread read it many times at a latency of a few dozen cycles rather than a few hundred, subject to the banking constraints on the shared-memory access pattern. It also makes the space dangerous: an array index past the declared size reads whatever another block left there, or faults, and the compiler cannot help. And because the array is per block, the size chosen at launch directly caps how many blocks the SM can host concurrently, which couples a memory decision to occupancy.

The per-SM array is finite and shared with the level-one cache on architectures that unify the two. The programming model exposes a static declaration for arrays whose size is known at compile time and an extern declaration for arrays sized at launch; the latter costs an extra launch parameter and, past a threshold, an explicit opt-in because the default per-block allowance is smaller than the hardware maximum.

The constant bank and its broadcast rule

Constant memory holds 64 KB for the whole grid, is written from the host between launches, and is read-only inside kernels. Its cache is small and sits close to the instruction pipeline, and it has one property the other spaces lack: when every thread of a warp reads the same address, the access is served in a single pass by broadcast. Uniform coefficients, filter weights, and configuration structures therefore cost almost nothing, even in a loop.

The rule reverses under divergence. If the threads of a warp read different addresses, the constant cache serializes the distinct addresses, and a warp that touches 32 different words pays 32 times. A table lookup indexed by thread ID is the canonical mistake: it turns the cheapest space in the machine into the most expensive one. The fix is to move the table to global memory, where divergent addresses at least coalesce when they happen to be near each other, or to restructure so the index is uniform across the warp.

Generic addressing and the cost of not knowing

Hardware needs to distinguish these spaces because they have different address decoders. The instruction set carries a separate load and store for each, and a generic form that accepts a pointer to any of them. Converting a space-specific address into a generic one and back is an explicit operation, and the generic load has to determine at run time which window the address falls into before it can issue.

The practical consequence is that a pointer whose space the compiler cannot prove costs more than one it can. Passing a shared-memory array to a function through a plain pointer, or storing a shared address into a global structure, forces generic addressing. Keeping the declaration visible at the point of use lets the compiler infer the space and emit the narrower instruction. This is the reason a great deal of GPU code carries qualifiers that look redundant: they are not type decorations, they are address-space information the backend needs before it can pick an opcode.

Where the boundaries move

The set of spaces is stable across architectures; their sizes and their relationships are not. The constant bank has stayed at 64 KB for many generations while the shared-memory array grew severalfold and, on some parts, became configurable between cache and scratchpad. The read-only path that once sat on a separate texture cache was later folded into the same physical array as the level-one data cache. Local memory has remained backed by device memory, but its spill traffic became cheaper as caches grew.

Code that hard-codes a shared-memory budget therefore ages badly. The durable approach is to query the device limits at run time and to size tiles from the reported per-block maximum, then let the launch configuration follow. The same caution applies to the per-thread register budget: the arithmetic that converts a register count into a maximum number of resident threads depends on the register file size of the SM, which is not the same on every part.

What does not change is the ownership rule underneath all of it. Every value a kernel touches is either in a register, in a per-SM array that the block owns, or in DRAM behind a cache hierarchy. Deciding which of those a given piece of data should occupy, and then making sure the compiler can see the decision, is most of the work of writing memory-efficient GPU code.