CUDA Memory Spaces, Caches, and Consistency Behavior

Hello! Last lesson, you calculated how many blocks and warps an SM can keep resident. That capacity matters because some of those warps will wait for data. To interpret the wait, you need to distinguish where a CUDA value can be accessed, where it is physically stored or cached, and what makes another thread’s update safe to observe.

This lesson maps the principal CUDA memory spaces onto registers, shared memory, caches, and device DRAM. It then connects that map to the consistency rules that matter when threads, blocks, kernels, or the host exchange data.

A memory space is not necessarily a piece of hardware

Start with a hardware view, but treat the pictured capacities as specific to the GPU discussed—not as constants for every device.

The following diagram emphasizes CUDA’s named spaces rather than a complete cache hierarchy. In particular, it does not draw the ordinary L1 and L2 cache path between an SM and DRAM.

Now compare that picture with the software-facing classification:

Here is the working map. “DRAM-backed” describes ordinary device allocations and the usual backing store; a particular load can be served from a cache instead.

The first important correction to intuition is local memory: local means thread-private, not physically close to the execution units. Conversely, shared memory is physically close to an SM but is not shared across all blocks by default. On supported architectures, a block-cluster feature can extend shared-memory access beyond one block; the ordinary block-scoped model remains the right starting point here.

The second correction is that the table’s “host” access is not a promise that a CPU may directly dereference every device allocation or safely read it while a kernel writes it. Allocation type, transfer or mapping APIs, and synchronization determine how the host participates.

Follow an access, not just a variable declaration

For an ordinary global load, a useful performance model is: it may be satisfied by an applicable cache near the SM, by the device-wide L2 cache, or ultimately by its backing memory. This is a model of possible service locations, not a requirement that every instruction traverse identical levels. Architecture, instructions, and cache policies matter. L2 is shared among SMs; L1 behavior is closer to an individual SM.

Shared memory takes a different route. Threads explicitly place data there and read it back. That gives a block control over reuse, but also responsibility for coordinating producers and consumers. Because shared memory and registers consume finite per-SM resources, using more of either can lower the resident-block limit you calculated last lesson. A faster individual access is not, by itself, a faster kernel.

Suppose a kernel has a per-thread array indexed by a value known only at runtime. The compiler may place that array in local memory because individual elements cannot readily live in fixed registers. A load from it is then a memory operation, even though the source declares an automatic, thread-private variable. It need not reach DRAM on every access—a cache hit can serve it—but treating it as a register access would give the wrong performance model. Spills caused by register pressure create the same kind of local-memory traffic; you will inspect generated code for this in the compiler module.

A load of a global array has a different question attached: how much data do neighboring lanes request, and how much of what is fetched gets used? Caches can reward reuse, but they do not turn scattered requests into the same traffic pattern as compact ones. Later lessons will quantify transactions and data movement; for now, keep access pattern, cache service location, and backing storage separate.

Read-only paths: constant and texture

Constant and texture access illustrate why “stored in device memory” does not imply “behaves like an ordinary global load.” They provide specialized read paths, with different strengths and restrictions.

Constant memory is useful when many lanes of a warp request the same small set of unchanged values: its cache supports efficient delivery of those reads. Its advantage is access-pattern-dependent. If lanes request many different constant addresses, those requests can serialize; “constant” does not mean “always fast.” Kernel code reads this space rather than updating it as ordinary writable global memory.

Texture access is also cached and read-oriented, but its access facilities are suited to patterns such as spatial locality. A crucial correctness restriction accompanies that path: do not write underlying data through global memory and expect a texture fetch of that data within the same kernel invocation to observe the write. The guide says that such fetches return undefined data because the texture cache is not invalidated by those stores. A later, properly ordered kernel invocation or copy gives a different situation.

These are not merely performance footnotes. Choosing an access path affects what observations are permitted as well as where a read might be served.

Visibility needs an ordering relationship

A cache diagram tells you where bytes might be. It cannot, on its own, tell you whether one thread is allowed to observe another thread’s update. For that, distinguish three questions:

1. Access: Can both participants address the location? Two blocks can address global memory; only threads entitled to a block’s shared memory can use that block’s ordinary shared data.
2. Atomicity: Can an operation be observed as an indivisible update where required?
3. Ordering and visibility: What synchronization makes a consumer’s observation of a producer’s earlier writes well-defined?

For the simplest handoff, a block barrier supplies the needed relationship:

__global__ void publish_within_block(int* out) {
    __shared__ int value;

    if (threadIdx.x == 0) value = 42;
    __syncthreads();

    if (threadIdx.x == 1) out[0] = value;
}

Assume one block with at least two threads. Every thread reaches the barrier. Writes made before __syncthreads() are completed and visible to participating threads in that block after it, so thread 1 can read value. Removing the barrier would leave the producer–consumer ordering unspecified; putting it only inside thread 0’s conditional would not be a valid block-wide barrier. A warp-scoped synchronization can be appropriate when all relevant participants are in the same warp, but its participation and scope must actually cover them.

Across different blocks, an ordinary __syncthreads() cannot establish a grid-wide handoff. If one block writes global payload data and another waits on a flag, a plain global flag load/store is not a safe publication protocol. Use an appropriate atomic release operation to publish the flag and an acquire operation to observe it, with a scope that includes both blocks, or separate the producer and consumer into properly ordered kernels. A fence orders the calling thread’s memory operations at its specified scope; a fence alone is neither a block barrier nor a substitute for a correctly synchronized flag. A spinning consumer can also occupy residency needed by a producer, so an in-kernel protocol deserves more scrutiny than its flag code alone suggests.

The formal PTX view makes the role of scope explicit:

Kernel boundaries often provide a cleaner handoff. If a producing kernel completes before a consuming kernel starts in the same stream, the consumer can use the producer’s global-memory results. Work in different streams needs an explicit dependency when the order matters, such as an event recorded by the producer stream and waited on by the consumer stream. On the host, an asynchronous launch is not proof that results are ready: synchronize the relevant work before copying or reading them through a supported host-accessible path. CPU–GPU concurrent sharing has additional allocation and system-scope requirements; do not infer safety merely because both sides have an address for the data.

The practical rule is: cache coherence is not a replacement for a synchronization protocol. Equally, synchronization does not make an inappropriate access path—such as the same-invocation texture read after an underlying global write—safe.

Takeaways

CUDA memory names primarily describe an address space and its access rules. Registers and shared memory are on-SM resources; ordinary global and local accesses have device-memory backing and may be served by caches. Constant and texture reads use specialized cached paths. The exact cache route depends on the device and the access, so physical location alone is an incomplete latency prediction.

For correctness, ask who shares the location and what establishes ordering: a suitable block or warp synchronization, appropriately scoped atomics, an ordered kernel boundary, or host-side synchronization. This map will help you interpret memory stalls without confusing a cache hit with a visibility guarantee.

Next, you will examine branch divergence and predication: even when warps access the same memory hierarchy, the paths taken by their lanes can change how efficiently they execute.