Page-Locked (Pinned) Host Memory, in depth: why DMA needs it, the CUDA allocation and registration APIs, PyTorch pin_memory and non_blocking, and pinning without starving the node

Why a GPU cannot copy from ordinary memory

Every process sees virtual addresses. The kernel maps each 4 KiB page to a physical frame and is free to change that mapping at any time: it can swap a cold page to disk, migrate it between NUMA nodes, merge identical pages, or leave it unallocated until first touch. The CPU copes because every access goes through the page tables and a missing page simply faults and is brought back.

A DMA engine has no such safety net. It is handed bus addresses (physical, or translated by an IOMMU) when the transfer is programmed, and then it streams bytes for milliseconds without consulting the CPU. If the kernel moved a page in the middle, the engine would read whatever now occupies that frame. So a device may only DMA from memory whose pages are resident and pinned to fixed frames for the whole transfer. Page-locking is that guarantee: the driver locks the pages, records their bus addresses, and keeps them until you free or unregister the buffer.

When you hand cudaMemcpy an ordinary malloc pointer, the driver cannot DMA from it, so it copies chunks through a small pinned staging buffer it owns: CPU copy in, DMA out, repeat. That costs an extra pass over the data on the CPU, and because the driver must reuse its staging chunks, an asynchronous copy from pageable memory cannot overlap with host work the way the API name suggests.

The two transfer paths

Two practical effects follow. Bandwidth: the pinned path usually runs close to the measured link rate, while the pageable path is capped by the CPU copy and the chunking, and is noticeably lower on most systems. Concurrency: only the pinned path gives you a copy that is genuinely queued on a stream, so the copy engine can move batch n+1 while the SMs compute batch n. On many training jobs the second effect matters more than the first.

The CUDA API and a double-buffered pipeline

The rule that ties these together: allocate or register a bounded set of buffers once, reuse them for the life of the process, and never let the CPU overwrite a pinned buffer that a queued copy has not finished reading. The pipeline below enforces the last point with one event per buffer. Recording the event after the kernel, not after the copy, also protects the device buffer from being overwritten while the kernel still reads it.

// Double-buffered upload: the CPU fills one pinned buffer while the other is in flight.
// Wrap every call in your CHECK() macro; omitted here for length.
constexpr size_t CHUNK = 64u << 20;            // 64 MiB per chunk
float *h[2], *d[2];
cudaStream_t s[2];
cudaEvent_t  done[2];
for (int k = 0; k < 2; ++k) {
    cudaHostAlloc((void**)&h[k], CHUNK, cudaHostAllocDefault);   // pinned once, reused
    cudaMalloc((void**)&d[k], CHUNK);
    cudaStreamCreateWithFlags(&s[k], cudaStreamNonBlocking);
    cudaEventCreateWithFlags(&done[k], cudaEventDisableTiming);
}
for (size_t i = 0; i < nchunks; ++i) {
    int k = i & 1;
    cudaEventSynchronize(done[k]);             // wait until buffer k is free again
    fill_from_source(h[k], i);                 // CPU writes the next chunk
    cudaMemcpyAsync(d[k], h[k], CHUNK, cudaMemcpyHostToDevice, s[k]);
    process<<<grid, block, 0, s[k]>>>(d[k], CHUNK / sizeof(float));
    cudaEventRecord(done[k], s[k]);            // fires after copy AND kernel finish
}
cudaDeviceSynchronize();

With two buffers, the CPU fill of chunk i+1, the copy of chunk i and the kernel on chunk i-1 can all be in flight at once. Add a third buffer if the fill step is bursty. Allocation is deliberately outside the loop: pinning requires the driver to fault in and lock every page and update its mappings, which is far slower than a malloc, so pin-per-iteration code often ends up slower than the pageable code it replaced. Streams and events are covered in more depth in CUDA streams and graphs.

Portable, mapped and write-combined memory

Portable memory. The runtime documentation defines cudaHostAllocPortable as making the allocation count as pinned memory in every CUDA context, not just the one that allocated it. Set it in multi-GPU processes where one thread allocates buffers that several devices copy from; it costs nothing and removes a dependency on platform defaults.

Mapped (zero-copy) memory. cudaHostAllocMapped also maps the buffer into the device address space, and cudaHostGetDevicePointer returns a pointer a kernel can dereference. Every access then crosses the link at link latency, so it only pays off when each byte is read once: a kernel streaming a large input once, a small flag the host polls, or an integrated GPU that shares DRAM with the CPU. A kernel that re-reads mapped memory in a loop runs at PCIe speed instead of HBM speed. The CUDA best-practices guide notes that under unified virtual addressing, memory from cudaHostAlloc has identical host and device pointers; calling cudaHostGetDevicePointer anyway keeps code correct on every platform.

Write-combined memory. cudaHostAllocWriteCombined makes the pages uncached on the CPU and lets writes be combined into bursts. The CUDA documentation says it can transfer across PCIe more quickly on some system configurations but cannot be read efficiently by most CPUs. Use it only for buffers the CPU writes sequentially and never reads back, such as an upload ring; a CPU loop that reads it will be dramatically slow.

Pinning is also what GPUDirect generalises: instead of locking host pages for the GPU, it lets NICs and storage DMA into GPU memory itself, as described in GPUDirect. Do not confuse any of this with cp.async, which moves data from global to shared memory inside the GPU.

Pinned memory in PyTorch

PyTorch wraps all of this. tensor.pin_memory() returns a pinned copy, torch.empty(..., pin_memory=True) allocates pinned directly, and DataLoader(pin_memory=True) runs a dedicated thread in the main process that copies each collated batch from the workers into pinned memory, so the training loop never waits on that step. Pinned blocks come from a caching host allocator, which keeps freed blocks for reuse and records an event on each copy so a block is not handed out again while a transfer still reads it.

Then t.to('cuda', non_blocking=True) enqueues the copy on the current stream and returns immediately. Host-to-device is safe without extra synchronisation because later kernels on the same stream are ordered after the copy. The prefetcher below puts copies on a side stream so they overlap with compute on the main one:

import torch

loader = torch.utils.data.DataLoader(ds, batch_size=256, num_workers=8,
                                     pin_memory=True, persistent_workers=True)

class CudaPrefetcher:
    """Copies batch n+1 on a side stream while the model trains on batch n."""
    def __init__(self, loader, device="cuda"):
        self.it, self.device = iter(loader), device
        self.stream = torch.cuda.Stream(device)
        self._preload()

    def _preload(self):
        try:
            cpu_batch = next(self.it)          # already pinned by the loader
        except StopIteration:
            self.batch = None
            return
        with torch.cuda.stream(self.stream):
            self.batch = [t.to(self.device, non_blocking=True) for t in cpu_batch]

    def __iter__(self):
        return self

    def __next__(self):
        if self.batch is None:
            raise StopIteration
        main = torch.cuda.current_stream()
        main.wait_stream(self.stream)          # compute waits for the copy, CPU does not
        batch = self.batch
        for t in batch:
            t.record_stream(main)              # allocator must not recycle early
        self._preload()
        return batch

Two traps are documented in the official PyTorch guide to non_blocking and pin_memory(). First, device-to-host copies with non_blocking=True are not safe to read on the CPU until you synchronise: gpu_t.to('cpu', non_blocking=True).mean() can compute on bytes that have not arrived yet. Call torch.cuda.synchronize() or wait on an event first. Second, pinning on the fly does not help: pageable.pin_memory().to('cuda') pays for an allocation and a CPU copy that the driver would have done anyway, and the guide measures it slower than a plain transfer. Pin where the data is produced, in the loader, not just before the copy.

Worked example: measuring and sizing the win

Take an image model that consumes 256 MiB of input per step and computes for 40 ms. Measure, do not guess, the two copy rates on your machine with a script like this:

import torch

def gbps(src, dst, reps=20):
    torch.cuda.synchronize()
    start, end = torch.cuda.Event(enable_timing=True), torch.cuda.Event(enable_timing=True)
    start.record()
    for _ in range(reps):
        dst.copy_(src, non_blocking=True)
    end.record()
    end.synchronize()
    return reps * src.numel() * src.element_size() / (start.elapsed_time(end) / 1e3) / 1e9

n = 256 << 20                                          # 256 MiB
pageable = torch.empty(n, dtype=torch.uint8)
pinned = torch.empty(n, dtype=torch.uint8, pin_memory=True)
dev = torch.empty(n, dtype=torch.uint8, device="cuda")
print(f"H2D pageable {gbps(pageable, dev):6.1f} GB/s")
print(f"H2D pinned   {gbps(pinned, dev):6.1f} GB/s")
print(f"D2H pinned   {gbps(dev, pinned):6.1f} GB/s")

Suppose it reports 24 GB/s pinned and 10 GB/s pageable (illustrative figures; yours depend on link generation, CPU and NUMA placement). A 256 MiB batch then takes about 11 ms pinned and 27 ms pageable. With pageable memory the copy cannot overlap, so a step costs roughly 27 + 40 = 67 ms. With pinned memory and the prefetcher, the 11 ms copy hides under the 40 ms of compute and a step costs about 40 ms: step time falls by about 40 percent, roughly 1.7 times the throughput, from a flag and a few lines. If compute were only 5 ms, the job would be copy-bound either way, and the right fix would be fewer bytes: decode on the GPU, send uint8 instead of float32, or cache the dataset in device memory.

Confirm the overlap in a profiler such as Nsight Systems: the memcpy rows should sit underneath kernel rows, not between them. If a copy appears labelled as pageable, some tensor escaped the pinned path, often one created inside the training loop.

Operating pinned memory safely

• Budget pinned memory explicitly. Locked pages cannot be swapped or reclaimed, and the CUDA documentation warns that allocating too much reduces the memory the OS can page and degrades the whole system. A common budget is the staging rings plus loader prefetch: batch size times prefetch_factor times workers, not the dataset.
• Count it in container limits. Pinned allocations belong to your process and count against its memory limit; because the kernel cannot reclaim them, a cgroup that runs short reaches the OOM killer sooner than its pageable usage would suggest.
• Respect NUMA. Pin buffers on the socket attached to the GPU. Bind the loader and its memory with numactl --cpunodebind=N --membind=N or equivalent; a cross-socket copy adds an interconnect hop and loses bandwidth.
• Register foreign buffers once. cudaHostRegister is the tool for memory owned by another library, such as a decode buffer. Register at start-up, keep the range alive, and unregister before it is freed.
• Instrument it. Log total pinned bytes at start-up and expose transfer time per step as a metric, so a regression to pageable copies is visible.

Failure modes

• Reading a non-blocking D2H result early. Metrics or checkpoints computed from bytes still in flight. Synchronise on an event before any CPU read.
• Reusing a pinned buffer too soon. The CPU refills a buffer an earlier cudaMemcpyAsync is still reading, corrupting a batch silently. Gate reuse on an event, as in the pipeline above.
• Pinning per step. Allocation or registration inside the loop; throughput drops and a profiler shows large host API time.
• Over-pinning. Tens of gigabytes locked by one job; other processes swap, the node stalls, or the job is OOM-killed.
• Zero-copy hot loops. A kernel re-reading mapped host memory; it is correct and very slow.
• Reading write-combined memory on the CPU. Uncached reads make a trivial loop crawl.
• Forgetting record_stream. A tensor used on a stream other than the one it was allocated on can be recycled while that stream still uses it.

Trade-offs

What to do next

1. Run the bandwidth script on each machine type you train on and record pageable, pinned H2D and pinned D2H rates.
2. Turn on pin_memory=True in every DataLoader feeding a GPU, and add non_blocking=True to the host-to-device copies.
3. Profile one training step and check that copies overlap compute; add a side-stream prefetcher if they do not.
4. Search the code for device-to-host non_blocking=True copies and add a synchronisation before every CPU read.
5. Set a pinned-memory budget per job, bind loaders to the GPU NUMA node, and log pinned bytes at start-up.
6. If the step is still copy-bound, reduce bytes moved before buying a faster link.