Memory Spaces
A CUDA kernel does not have one undifferentiated pool of memory to work with — it has six, each with its own scope, lifetime, and performance profile, and picking the wrong one for a given piece of data is one of the most common ways a kernel ends up an order of magnitude slower than it should be. This page is the map: what each space is, who can see it, how long it lives, and the rough latency and bandwidth numbers that make the choice matter. The pages that follow work through each space in depth.
Global Memory & Coalescing
Thread Indexing establishes that the fastest-varying array index should track threadIdx.x, so that adjacent threads in a warp touch adjacent addresses. This page is the mechanism that makes that rule matter: how the hardware actually turns a warp's 32 addresses into memory transactions, and why the difference between "adjacent" and "scattered" can be an 8x difference in delivered bandwidth for the exact same amount of useful data.
Shared Memory
Global memory is fast in aggregate but every access still pays a round trip through the memory system; when several threads in a block need the same piece of data, or a thread needs to hand a value to another thread in the same block, routing through global memory to do it wastes bandwidth on traffic that never needed to leave the chip. Shared memory exists for exactly that: a small, explicitly-managed, on-chip scratchpad that every thread in a block can read and write, fast enough to use as a staging area rather than just a cache.
Bank Conflicts
Shared memory earns its speed by serving a whole warp in one cycle, but that promise depends on the warp's 32 addresses landing in 32 different pieces of hardware. When they don't, shared memory — normally close behind register speed — degrades to a fraction of it, and the kernel doesn't fail, it just quietly runs slower with no compiler warning to explain why.
Registers & Local Memory
The Register File and Occupancy covers how a kernel's register usage feeds directly into the occupancy calculation. This page is the other side of that same fact: what registers actually hold, what forces a value out of the register file and into memory instead, and how to tell when that's happened to your kernel.
Constant & Texture
Global memory's performance rules assume threads in a warp want different addresses and reward spreading them out into distinct sectors. Constant memory and the read-only data path invert that assumption: they're fast precisely when a warp's threads all want the same data, and the further a kernel's access pattern gets from that, the less either buys over an ordinary global read.
Unified Memory
Every allocation covered so far draws a hard line between host and device memory: a pointer is valid on one side or the other, and moving data across the line is an explicit cudaMemcpy the programmer writes and pays for. Unified Memory erases that line for the source code — one pointer, usable from both host and device — while the underlying hardware and driver still have to physically move bytes between two separate memory systems whenever the data is touched from the "wrong" side. Understanding when that migration happens, and how to steer it, is the difference between Unified Memory being a convenience and Unified Memory being a performance trap.
Pinned Memory
The SAXPY program in Your First Kernel paid three costs without the source code drawing attention to any of them ordinary malloc'd host memory is pageable, and pageable memory cannot be the source or destination of an asynchronous transfer. Pinned memory removes that restriction, and once transfers are asynchronous, they can be scheduled to overlap with compute instead of paying for it serially.
Distributed Shared Memory
Shared Memory is scoped to a single block the cluster's combined on-chip shared memory, addressable across block boundaries, without routing through global memory at all.
Async Data Movement
The classic tiled kernel loop looks like load → syncthreads() → compute → syncthreads() while stage i computes, stage i+1's load can already be in flight.
Consistency & Fences
Every earlier page in this section has quietly leaned on one synchronization primitive or another — syncthreads(), a cluster.sync(), a pipeline's consumer_wait() — to make writes from one thread visible to reads from another. This page states the rule underneath all of them explicitly: CUDA's memory model is weakly ordered, and without an explicit fence or atomic, there is no guarantee about when, or even whether, one thread's writes become visible to another thread at all. Getting this wrong doesn't usually crash a kernel; it produces a result that's correct most of the time and silently wrong occasionally, which is far worse.