Pinned Host Memory and Why Pageable Transfers Cannot Be Asynchronous

A device-side copy engine does not execute loads. It is programmed with a descriptor holding a source address, a destination address, a length, and a direction, and it walks that descriptor with hardware that knows nothing about the process that supplied the buffer. The addresses it consumes are device virtual addresses, translated by the GPU’s MMU and, behind it, by the IOMMU. What the engine cannot do is resolve an address that has no mapping: there is no path from a missing translation in a DMA descriptor to the operating system faulting in a pageable user page. That one property explains almost everything about how host memory reaches a device, and why the word “asynchronous” in a transfer API means less than it appears to.

The address a copy engine can actually use

A host virtual address is a promise made by the CPU’s memory management unit. Behind it sits a page table entry that maps a 4 KiB region, on the common architectures, to a physical frame. The kernel is free to change that entry at any time: NUMA balancing can migrate the page to another node, reclaim can write it out and drop it under memory pressure, swapping can move it to disk, and copy-on-write can replace the frame when another process writes to a shared mapping. None of these events are visible to a device that was handed the virtual address a moment earlier.

There are exactly two ways to make the address stable for the duration of a transfer. The first is to pin the pages: instruct the kernel that they may not be reclaimed, migrated, or written back until the pin is released, then translate the virtual range to physical addresses once and program the engine with the result. The second is to copy the data into a buffer that is already pinned and owned by the driver, and hand the engine the address of that buffer instead.

Everything else follows from which of these two paths a given API call takes. It also explains a persistent confusion: the device MMU makes device memory addressable, but a host pointer is not in that space until the driver has explicitly installed a mapping for it, which is the same pinning problem with different bookkeeping. An IOMMU does not remove the problem either. It lets the driver create an I/O page table entry for a host frame, but the driver still has to know which frame, and the frame still has to stay put.

What the driver does with a pageable pointer

When an application passes an ordinary malloc buffer to an asynchronous copy, the driver cannot program the engine with that address. It takes the second path. The sequence, as documented in the CUDA programming guide’s description of API synchronization behaviour, is:

1. Wait for all previously enqueued work in the target stream to complete.
2. Copy the contents of the user buffer into an internal pinned staging buffer using the CPU.
3. Submit a DMA descriptor from the staging buffer to device memory.
4. Return to the caller.

Step one is what makes the call synchronous in the way that matters. The driver cannot let earlier kernels keep running while it reads the user buffer, because those kernels might be writing to it. So the stream is drained, and any overlap the application hoped to achieve between the copy and prior compute is gone before the copy is even submitted. Step four returns once the data has reached staging, not once it has reached the device.

The arithmetic makes the penalty visible. A 1 GiB transfer over a PCIe Gen3 x16 link moves at roughly 16 GB/s in one direction, so the DMA alone takes on the order of 60 ms. The CPU-side copy into staging runs at host DRAM bandwidth, tens of GB/s on a server, and it adds a read of the user buffer plus a write to staging. The DMA then reads staging again. The data is read by the CPU, written into staging, and read once more by the engine, so it passes through host memory three times before any of it reaches the link, and the copy that the application believed was asynchronous has spent a substantial fraction of its time on a CPU core.

The reverse direction is worse. For a transfer from device memory to pageable host memory, the call returns only after the DMA has completed and the CPU has copied staging into the user buffer. There is no window in which the application can do anything at all.

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

int main(void)
{
    const size_t n = 1u << 24;          /* 16 MiB */
    float *h = NULL, *d = NULL;

    /* Pinned allocation: the pages are locked at allocation time. */
    if (cudaHostAlloc((void **)&h, n * sizeof *h, cudaHostAllocDefault) != cudaSuccess)
        return 1;
    if (cudaMalloc((void **)&d, n * sizeof *d) != cudaSuccess)
        return 1;

    for (size_t i = 0; i < n; i++)
        h[i] = (float)i;

    /* Because h is pinned, this call does not stage through an internal
       buffer and does not drain the stream before submitting. */
    cudaMemcpyAsync(d, h, n * sizeof *h, cudaMemcpyHostToDevice, 0);

    /* The copy is still in flight; the host buffer must stay alive. */
    cudaDeviceSynchronize();

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

What pinning costs, and where the cost lands

cudaHostAlloc and cudaMallocHost allocate and pin in one operation. cudaHostRegister pins a range the application already owns, without moving it. Both make the pages permanently resident until the corresponding free or unregister call.

Residency is the first cost. A 1 GiB pinned buffer is 262,144 pages of 4 KiB that the kernel can never reclaim, migrate, or swap. On a machine with 16 GiB of RAM that is a sixteenth of physical memory removed from the page cache’s working set for as long as the buffer lives. Pinning is also accounted against limits: the operating system may cap how much memory a process can lock, and a request beyond that cap fails rather than degrading gracefully.

The second cost is allocation latency. Pinning is per page, not per byte, so the work scales with the page count. Finding, locking, and building a scatter-gather list for a large region takes long enough that allocating a fresh pinned buffer inside an inner loop is a design error; the buffer should be allocated once and reused.

The third cost is fragmentation. Pinned pages cannot be moved, so a long-lived pinned buffer fragments the physical address space around it. A driver that later needs a large contiguous region may find one unavailable. This is the same problem the kernel faces with huge pages, and it is why the guidance for pinned memory is always to allocate early, allocate once, and allocate the smallest buffer that does the job.

Pinning does not make host memory faster for the CPU. It changes where the pages live and what the kernel may do with them; it does not change cache behaviour, and a pinned buffer read by a CPU core is neither faster nor slower than a pageable one.

Overlap requires more than an asynchronous call

Two pinned buffers and a stream that contains independent work are the minimum for overlap. The standard arrangement is double buffering: while the engine transfers chunk i+1, the device computes on chunk i. The per-iteration cost becomes the larger of the transfer time and the compute time rather than their sum, which is the whole point of the exercise.

There is a hard limit on how much of this can happen at once. A device exposes a small number of copy engines, and copies assigned to the same engine serialize with each other even when they are submitted on different streams. Whether a host-to-device copy and a device-to-host copy can proceed concurrently depends on whether the runtime assigns them to different engines; on many devices they can, because the link is full duplex, but two host-to-device copies assigned to the same engine will not. Sizing a pipeline around the assumption that many concurrent transfers will overlap is a common way to end up with a pipeline that does not.

The submission side has its own cost. Each copy or kernel launch is a command written into a queue and fetched by the device’s front end, and the host-side cost of submitting one is on the order of microseconds. A loop that submits a transfer per kilobyte spends more time in submission than in data movement. Chunk sizes in the hundreds of kilobytes to low megabytes amortize submission without making the pipeline so coarse that the tail dominates.

Failure modes that surface far from the transfer

The most common failure is passing a pageable pointer to a kernel rather than to a copy call. A kernel that dereferences an unmapped host address does not get a helpful error; it reads whatever the device MMU resolves, which is typically nothing, and the result is a fault or silent corruption depending on the platform.

The second is freeing a host buffer while a transfer from it is still in flight. With pageable memory the driver’s staging copy has already completed by the time the call returns, so the window is narrow. With pinned memory the engine may be reading the buffer long after the call returned, and releasing it is a use-after-free that the runtime has no way to detect. The buffer must outlive the synchronization point, not the submission point.

The third is over-pinning. Applications that pin every buffer they own, on the theory that pinned is faster, eventually starve the kernel of reclaimable memory and slow the entire machine down, including their own host-side code.

The address space questions that determine whether a given pointer is even meaningful on the device are covered in unified virtual addressing and pointer provenance. The alternative to copying at all, mapping host memory directly into the device’s address space, has its own cost model described in host-mapped memory over PCIe. And once transfers are genuinely asynchronous, the ordering rules that decide when the data is safe to read are those of stream ordering and submission semantics.