The CUDA Handbook ▸ Chapter 5. Memory

5.7 Managed Memory

Explicit memory management always has been perceived as CUDA’s Achilles’ heel. Every new CUDA programmer remembers learning cudaMalloc() and having to copy inputs to, and outputs from, memory attached to the GPU, and every superficial critique of the programming model seizes on that complexity first. Both NVIDIA and AMD long have aspired to make it obsolete: NVIDIA re-entered the data center CPU business with Grace and designed the Grace Hopper superchip around a coherent, unified memory model, while AMD built consumer APUs – CPU/GPU devices sharing a single memory pool – in the early 2010s and returned to the idea at data center scale with the MI300A in the early 2020s. All other things being equal, unified memory is unquestionably easier to use.

But for the best performance – the lowest-latency access by the CPU and the highest-bandwidth access by the GPU – the exigencies of physical design often favor separate memory pools and explicit allocation, if the application can tolerate them. And since the bulk of the market has moved to AI, where explicit memory management always has been the norm, the market has not favored unified memory architectures. For a fuller history of unified CPU/GPU designs and their fortunes, see APUs: When and Why They Work. In any case, managed memory was created to service this developer requirement for simpler memory management, where “a pointer is a pointer.”

Managed memory hides, but does not eliminate, the need to distinguish between host and device memory. Allocations made with cudaMallocManaged() (or cuMemAllocManaged() in the driver API) return a single pointer that is valid on both the CPU and the GPU: the same buffer can be initialized with ordinary host code, handed to a kernel, and examined again by the host, all without an intervening memcpy. The distinction between host and device memory has not gone away – the data still must reside somewhere, and performance still depends on where it resides when it is accessed. Managed memory shifts the burden of migrating data from the programmer to the underlying system.

When managed memory was introduced in 2014, no NVIDIA GPU could take a page fault, so coherence was maintained in software: the driver migrated every page the CPU had touched back to the device before any kernel launch, and host code was forbidden from touching managed buffers at all while a kernel might be running. Capacity was limited to physical device memory. As described in Section 2.4.2, the Pascal architecture (2016) added hardware demand paging, enabling pages to migrate when and where they are referenced; with this new virtual memory support, managed allocations may exceed physical device memory and the CPU and GPU can access managed buffers concurrently. Volta’s access counters (2017) further refined migration policy by moving pages based on observed access frequency rather than on first touch. Not every platform has crossed over, however: under Windows’ WDDM driver model, managed memory still follows the original software-coherence rules, even on the newest hardware. Applications can query whether hardware coherence is available with the cudaDevAttrConcurrentManagedAccess device attribute.

Because migration happens at page granularity in response to faults, the performance of managed memory is shaped by locality of reference: a page that stays where it is used is free after first touch, while a page that ping-pongs between processors incurs a fault and a migration on every move. CUDA provides tools to steer pages to where they belong: cudaMemPrefetchAsync() migrates a range to a specified processor ahead of use, and cudaMemAdvise() expresses intent – cudaMemAdviseSetReadMostly to duplicate read-shared pages on both processors, cudaMemAdviseSetPreferredLocation to anchor pages, and cudaMemAdviseSetAccessedBy to establish mappings that avoid faults altogether.

The costs described above can be measured using the managedOverhead.cu program in the concurrency/ directory of the source code. It times a null kernel launch followed by cudaDeviceSynchronize() – the same operation measured in Section 6.1 – while varying the number of managed pages allocated, from none to 4096, in two variants: one in which the CPU touches every page between launches, and one in which the allocation is left untouched. The two variants highlight a sharp distinction between the two regimes. Under software coherence, every page the CPU dirties must be migrated back to the device before the next kernel may launch, so launch time grows steeply with the number of touched pages even though the kernel never references them. Under hardware demand paging, memory the kernel never references does not migrate at all: on an Ampere-class GeForce RTX 3060 under Linux, the untouched variant holds at the null-launch floor of 3.9 μs whether the allocation spans zero pages or 4,096, and the touched variant grows only from 4.0 μs to 66.9 μs over the same range – about 15 ns per page touched, the signature of the CPU’s own cache-missing stores rather than of page migration. The original implementation was reverse-engineered with this same program in a 2017 blog post: under software coherence, the driver unmapped the CPU’s copy at every kernel launch, so each CPU touch afterward cost a page fault and a 4KB copy – an effective bandwidth under 500 MB/s on Maxwell, and roughly half a millisecond of overhead per launch with 16MB dirtied. Reader-submitted results for this benchmark, across GPU generations and operating systems, are collected at cudahandbook.com/benchmarks.

System-Allocated Memory and Hardware Coherence

Managed memory still asks the programmer for one concession: the allocation must come from cudaMallocManaged() rather than from the system’s own allocator. For GPUs that can access system-allocated memory, denoted by the cudaDevAttrPageableMemoryAccess device attribute, running kernels can dereference an ordinary pointer—memory returned by malloc(), mapped with mmap(), or taken from a global or the stack—with no cudaMallocManaged() and no cudaHostRegister(). For systems with this device attribute (e.g., Grace Blackwell), the aspiration that “a pointer is a pointer” becomes true: allocations performed by the operating system for the CPU are directly addressable by the GPU.

These capabilities were rolled out over several generations of hardware, with a steady increase in performance and function. The earlier mechanism, targeting discrete GPUs over PCIe, used Linux’s Heterogeneous Memory Management (HMM) to extend the managed-memory machinery to every allocation: the GPU signals page faults against system memory and migrates pages on demand, exactly as for managed memory. HMM needs a recent Linux kernel and driver; where it is absent, a system pointer must still be pinned with cudaHostRegister() or staged through an explicit copy.

On more recent hardware, such as the coherent superchips of Section 2.1.5, that software-based mechanism is replaced by coherency hardware. Address Translation Services (ATS) allow the GPU to walk the CPU’s own page tables, so a host virtual address resolves natively on the device, and the coherent NVLink-C2C link keeps the two memories consistent without any page migration. Systems with this capability set the cudaDevAttrPageableMemoryAccessUsesHostPageTables attribute.

For many applications, coherent CPU/GPU memory simplifies programming without any performance implications. But for others, affinity concerns may still loom. Like CPU systems with NUMA, nonlocal memory accesses incur a cost, and steering memory traffic to allocations that are local to the NUMA node can confer performance benefits. The difference is that in a NUMA configuration for CPUs, the peer devices are relatively similar. For coherent CPU/GPU systems, the two devices have different latency tolerance and bandwidth characteristics that make coherency harder to deliver.

On the Grace-generation superchips, that asymmetry is reflected in the coherency properties: the CPU may cache the GPU’s memory, but the GPU does not cache the CPU’s (Section 2.1.5). The coming Vera Rubin generation is said to relax this constraint, tightening the coherence between the two processors. Taken together, cudaMalloc(), cudaMallocManaged(), HMM, and hardware coherence form one trajectory: from two address spaces bridged by explicit copies, to a single address space with statically allocated address ranges corresponding to each device, to a single address space whose pages migrate on fault, to a single address space that only ‘migrates’ cache lines. Where a given application lands on that trajectory is a property of the platform, not of its source code—which is exactly the portability managed memory set out to provide.