Unified Virtual Addressing and How a Pointer Becomes Usable on a Device

Before unified addressing, a device had its own address space with its own numbering, starting at zero. A pointer in a kernel was an offset into device memory, and a host pointer was an offset into host memory, and the two were indistinguishable at the type level. Every API that moved data had to be told which side a buffer lived on, and every kernel argument that pointed at data had to be accompanied by an implicit agreement about which space the pointer referred to. The failure mode was not subtle: a kernel given a host pointer treated it as a device offset, and the resulting access was somewhere in device memory rather than nowhere.

Unified virtual addressing removes the ambiguity by construction. The host and all devices on a platform share one virtual address space, with disjoint ranges carved out for each processor’s memory. Once the ranges are disjoint and known, the numeric value of a pointer determines where it points. That property is what allows a copy API to take only a pointer and a size, and it is also the property that makes cross-device pointer passing possible at all.

One address space, several physical memories

The layout is not a hardware feature so much as an agreement maintained by the driver. At startup the driver reserves regions of the process’s virtual address space, and each device is assigned a window within which device allocations are placed. Host memory occupies the ordinary range that malloc and mmap return. The reservation is large: on a 64-bit platform, device windows are sized so that a device with tens of gigabytes of memory has room to grow, and the reserved regions are far larger than the physical memory they map.

Nothing about that layout changes the fact that the memories are physically separate. A virtual address in a device’s window still resolves to device DRAM, reached over PCIe or a vendor interconnect. What changes is that the CPU can name it. Whether the CPU may actually load from it is a separate question answered by the mapping flags: some device windows are mapped for host access and some are not, and reading an unmapped one faults.

The device side of the agreement is the device MMU. Every device access, whether to its own memory or to a host window, goes through a translation. That translation table is populated by the driver, and it can cover host memory, device memory, or another device’s memory, depending on what has been established. Unified addressing therefore describes a namespace, not a guarantee of reachability. Reachability is per mapping, and mappings are created by allocation and by explicit peer or host-mapping operations.

There is a parallel arrangement on CPUs with an IOMMU, and the two are often confused. An IOMMU translates the addresses a device emits when it performs DMA, using an I/O page table that the kernel maintains for that device. Device-side translation, driven by the GPU’s own page tables, is a different mechanism with a different table. Both exist on a modern server, both must agree about which physical frame a given host page occupies, and a bug in either produces the same symptom: the device reads the wrong bytes or faults.

What the runtime can infer, and what it cannot

Given a pointer, the runtime compares its value against the registered windows. A value in a device window is a device pointer; a value outside all windows is a host pointer. That comparison is the entirety of the inference, and it has a direct consequence: it cannot distinguish a host pointer from a garbage pointer, and it cannot distinguish a device pointer from a stale device pointer. Provenance is not tracked. Only the numeric range is checked.

#include <cuda_runtime.h>

/* Unified addressing lets the runtime classify a pointer by value.
   cudaPointerGetAttributes reports the owning context and location. */
static void describe(const void *p)
{
    cudaPointerAttributes a;
    if (cudaPointerGetAttributes(&a, p) != cudaSuccess) {
        /* Not a registered allocation: treat as pageable host memory. */
        cudaGetLastError();
        return;
    }
    /* a.type distinguishes host, device, and managed memory; a.device
       is the owning device for device and managed allocations. */
}

int main(void)
{
    float *h = NULL, *d = NULL;
    cudaMallocHost((void **)&h, 4096);
    cudaMalloc((void **)&d, 4096);

    describe(h);   /* host allocation, pinned */
    describe(d);   /* device allocation on the current device */

    /* No direction argument is needed when the source is unambiguous. */
    cudaMemcpy(d, h, 4096, cudaMemcpyDefault);

    cudaFree(d);
    cudaFreeHost(h);
    return 0;
}

The cudaMemcpyDefault direction is the visible benefit of the namespace: the runtime inspects both pointers, determines the transfer class, and picks the path. That works because the classification is reliable for pointers that came from the runtime’s own allocators. It is not reliable for pointers manufactured by arithmetic, and it does not apply to allocations made by another process, another context, or a library that manages its own memory.

Two failure modes follow directly. A pointer to the interior of an allocation classifies correctly, since the window covers the whole range. A pointer that has been cast to an integer, hashed, stored, and reconstructed still classifies correctly, for the same reason. But a pointer from a different process has a different virtual address space and is meaningless here regardless of its value, and a pointer from a device in another context may fall in a window that this process has not registered, in which case it is misclassified as host memory.

Why a shared namespace is what makes peer access expressible

Consider two devices in one process. Each has a window. If device 1 can map device 0’s memory into its own page tables, then a kernel running on device 1 can dereference a pointer into device 0’s window, because the value is the same value device 0’s own kernels use. Without a shared namespace the pointer would have to be rebased, and every structure containing pointers, such as a linked list, a tree, or a struct with a pointer to its own payload, would need to be rewritten for each device that traverses it.

That property is what makes device-to-device copies and peer mappings expressible with a single pointer, and it is the foundation for the topology questions covered in peer-to-peer transfers and link topology. It also explains a constraint: peer mappings are established per device pair, and a kernel that follows a pointer into a device whose memory is not mapped into the current device’s tables faults rather than silently reading zeros.

Managed memory extends the same namespace with a different contract. A managed allocation is placed somewhere, in device memory, host memory, or both, and migrates on access. Its pointer value is valid everywhere because the driver can move the backing storage without changing the virtual address. That is precisely why migration is possible at all: the virtual-to-physical translation is what changes, not the pointer. The mechanics of that migration, including the page faults that trigger it, are the subject of GPU page faults and demand migration.

The cost of a flat namespace

A single namespace is not free. It constrains the driver’s ability to place allocations, because every allocation must land inside a reserved window and windows cannot overlap. On a platform with many devices the reservation grows with the device count, and address space exhaustion is a real limit on 32-bit platforms and on some virtualized configurations where the guest’s address space is smaller.

It also makes some errors undetectable. With separate address spaces, passing the wrong kind of pointer to a kernel usually produced an immediate, obvious failure. With unified addressing, a host pointer passed to a kernel on a system where host memory is not mapped into the device’s tables produces a fault at first access, which is fine, but on a system where it is mapped through a peer mapping or a managed allocation, the access succeeds and reads host memory at PCIe latency. A kernel that accidentally dereferences a host pointer can therefore run to completion, hundreds of times slower than intended, with no error at all. The only signal is the elapsed time.

The last constraint is about visibility. The namespace is per process. Two processes each get their own, and a pointer shared between them requires either an inter-process communication mechanism that maps the same physical memory into both, or a copy. Unified addressing does not create a system-wide namespace, and treating a device pointer as something that can be written to a file, sent over a socket, or stored in shared memory is a category error that fails in a way that depends entirely on what the receiving process happens to have mapped at that address.

The transfer mechanics that operate on this namespace are covered in pinned host memory and asynchronous transfers, and the ordering guarantees that govern when a device may observe a write made through one of these pointers are those of stream ordering and submission semantics.