GPU Page Faults and Demand Migration in Unified Memory

Unified memory is usually described as a convenience: one allocation, usable by host and device, with the driver moving the data where it is needed. The interesting part is the mechanism underneath, because the convenience is bought with page faults on the device’s critical path, and every performance characteristic of unified memory follows from how those faults are batched, serviced, and resolved.

The device-side memory management unit is the central actor. It translates every address a thread issues, using page tables the driver maintains. When a translation is absent, the MMU raises a fault rather than returning garbage, and the fault is delivered to the driver through a channel that has to be drained by software. That delivery path is orders of magnitude slower than a cache miss, and it happens on a unit of work, typically a warp, that is stalled until it resolves.

The fault path from MMU to driver

A device access to an address with no valid translation stops the issuing warp and records a fault entry. The entry carries the faulting address, the access type, and an identifier for the context and the engine that produced it. The device then signals the host, typically by writing to a queue that the driver polls or by raising an interrupt, depending on platform and driver configuration.

The driver reads the fault, determines which managed allocation the address falls in, and decides what to do. Three outcomes are possible. If the page is resident in host memory and the device is permitted to access host memory directly, the driver can install a mapping to the existing host page and resume. If the page needs to be moved, the driver allocates a device page, copies the contents, updates the device page table, and resumes. If the allocation is not permitted on this device at all, the fault is fatal and the context is destroyed.

Resuming is not a single operation. The stalled warp must be restarted at the faulting instruction, and the runtime must ensure the fault has actually been resolved before the instruction re-executes, or the warp faults again in a loop. On hardware that supports replay, the driver marks the fault as handled and the scheduler reissues the warp. The cost of one round trip through this path is on the order of microseconds on the host side, which is enormous compared with the few hundred nanoseconds an access costs when the translation is present.

This is why the first rule of unified memory performance is that faults are tolerable in proportion to how rarely they occur. A program that faults once per page of a large streaming array pays one fault per 4 KiB or per 2 MiB, depending on granularity, and the amortized cost can be small. A program that faults repeatedly on the same pages has a structural problem that no amount of tuning fixes.

Granularity decides the arithmetic

The fault is per page, and the page size is not fixed. A device commonly uses 64 KiB pages for its own memory, 2 MiB pages where it can obtain physically contiguous memory, and the host’s base page size when a range stays mapped to host memory. The driver prefers large pages for managed allocations when it can get them, because a 2 MiB mapping covers 512 times the address range of a 4 KiB mapping and requires one page table entry instead of 512.

The difference in fault behaviour is dramatic. A 1 GiB array covered by 4 KiB pages is 262,144 pages; covered by 2 MiB pages it is 512. If every page faults once, the first case involves a quarter of a million driver round trips and the second involves five hundred. At a microsecond or so per fault, the first is hundreds of milliseconds of pure overhead and the second is under a millisecond.

Large pages are not always available. Physical memory fragmentation can prevent the driver from finding a contiguous 2 MiB region, and the allocation falls back to small pages. Pinning a large managed buffer for a long period, or mixing managed and ordinary allocations in a pattern that fragments the device’s memory, makes the fallback more likely. The observable symptom is a program whose fault count is far higher than the page count would suggest, and the fix is usually to allocate managed memory earlier and in fewer, larger chunks.

#include <cuda_runtime.h>
#include <stdio.h>

__global__ void touch(int *p, size_t n)
{
    /* Stride by one element so consecutive threads cover consecutive
       addresses: the warp's accesses coalesce into one or two pages. */
    size_t i = blockIdx.x * (size_t)blockDim.x + threadIdx.x;
    if (i < n) p[i] = (int)i;
}

int main(void)
{
    const size_t n = 1u << 26;   /* 256 MiB of ints */
    int *p = NULL;

    /* Managed allocation: the driver places it and migrates on demand. */
    if (cudaMallocManaged((void **)&p, n * sizeof *p) != cudaSuccess)
        return 1;

    /* Prefetch removes the faults from the kernel's critical path; the
       migration still happens, but ahead of the accesses. */
    int dev = 0;
    cudaMemPrefetchAsync(p, n * sizeof *p, dev, 0);
    cudaDeviceSynchronize();

    touch<<<(n + 255) / 256, 256>>>(p, n);
    cudaDeviceSynchronize();

    cudaFree(p);
    return 0;
}

Prefetching moves the cost off the critical path

A prefetch call tells the driver to migrate a range before anything accesses it. The migration still costs the same copy, but it happens in bulk rather than in response to faults, so the driver can issue large transfers, batch page table updates, and overlap the copy with other work in a different stream.

Bulk migration is much cheaper per byte than fault-driven migration for a structural reason: a fault serializes. The device stops a warp, signals the host, waits for the driver to run, waits for a copy to complete, waits for a page table update, and only then resumes. A prefetch of the same range issues one copy and one set of page table updates, and the device can keep executing other warps throughout. The difference is not a constant factor on the copy; it is the elimination of a per-page round trip.

Prefetching has a failure mode of its own: it can be wrong. Prefetching a range the kernel will not touch migrates data that has to come back later, and the return trip is another copy. Prefetching too far ahead evicts pages that are still in use. The useful pattern is prefetch by tile, immediately before the tile is consumed, in the same stream that will consume it, so the prefetch is ordered against the compute without a host-side barrier.

Oversubscription and the eviction question

Unified memory allows an allocation larger than device memory. When the device needs a page and its memory is full, something must be evicted. The driver chooses a victim, typically by a least-recently-used approximation maintained from access counters that the hardware records per page, and migrates the victim back to host memory.

Two properties of that mechanism matter. First, eviction is not free even when the victim is clean, because the mapping must be torn down and the page table updated. Second, the access counters are a coarse approximation: hardware samples accesses at a granularity of pages, and the driver’s view of recency is a summary, not a trace. A working set that oscillates around the device’s capacity produces thrashing, where nearly every access evicts a page that will be needed again shortly.

Thrashing is not a tuning problem; it is a capacity problem. The fix is to restructure the computation so the working set fits, which usually means tiling: process the array in chunks that fit in device memory, prefetch each chunk, compute, and let the previous chunk’s pages be evicted cleanly or prefetched back to the host if they will be needed. The alternative, relying on the driver to sort out a working set that does not fit, produces a program whose throughput is dominated by the interconnect rather than by the device.

There is a related trap in the other direction. A managed allocation that fits comfortably in device memory but is also touched by host code on every iteration will migrate back and forth, because each side’s accesses invalidate the other’s residency. The pattern is common in code that computes a scalar on the host from an array the device just produced. The fix is to keep the array resident on one side and copy only the scalar, or to keep the scalar in host-mapped memory so no migration is involved; see host-mapped memory over PCIe.

Concurrency and the limits of fault handling

Fault handling serializes at the device level in ways that are easy to overlook. A fault from one warp stalls that warp, but the driver’s response involves the whole context: page table updates must be made visible to all translation caches, which on some architectures requires invalidating entries across the device. During that window, other warps that would hit the same stale translation may also fault, and the driver may coalesce them or handle them one at a time.

The coalescing behaviour is the reason access patterns matter so much. A warp whose 32 threads touch 32 consecutive elements of an array touches one or two pages, so one fault serves all of them. A warp whose threads are scattered across 32 pages produces up to 32 distinct faults for a single instruction, and the driver may serialize them. The difference is not a factor of two in bandwidth but a factor of 32 in faults for that one instruction, and it shows up entirely in fault handling rather than in memory bandwidth.

Multi-device systems add a dimension. A managed page can be resident on device 0 while device 1 accesses it, and the driver may either migrate it or, on a peer-capable topology, map it remotely. The choice depends on whether the topology supports direct peer access and on how the driver weights migration against remote access. On a system without peer support, a page shared between two devices ping-pongs, and the resulting traffic goes through host memory twice per round trip. The topology considerations are covered in peer-to-peer transfers and link topology, and the kernel-side machinery that makes host memory migratable at all is described in page table mirroring between CPU and device.

The ordering guarantees that determine when a migrated page’s contents are safe to read are not part of the migration mechanism; they come from the stream and the memory model, covered in stream ordering and submission semantics and system-scope atomics and memory ordering.