Host-Mapped Memory over PCIe: Zero-Copy Access and Its Latency Wall

Copying is the default way to get data to a device, and for bulk data it is usually the right way, because a copy engine moves bytes in large, well-formed bursts while the compute units do something else. There is a second path: map host memory into the device’s address space and let kernels read and write it directly across the interconnect. No copy, no staging buffer, no explicit transfer to schedule. The data is wherever it already was.

The path is real, it is supported on discrete GPUs over PCIe, and it is slower than copying for almost every access pattern that moves more than a small amount of data. Understanding why requires looking at what a single load from host memory costs when it is issued by a thread running on a device, and at how the memory system handles writes that will never be read by the CPU.

The mapping, and what it is not

A pinned host allocation can be registered as mapped, and the device receives a pointer into the same buffer that the device’s page tables resolve to the host physical frames. With unified addressing, that pointer is the host pointer itself; without it, the runtime returns a separate device-side address for the same memory.

The mapping is a page table entry, not a cache. Every load the device issues to that range becomes a PCIe transaction: a read request packet travels to the root complex, is routed to memory, and the response comes back with the data. A read of 4 bytes by one thread is a transaction carrying 4 bytes of payload plus the packet headers, and PCIe’s header overhead means the wire carries substantially more than the payload. The efficiency of the link is determined by how much useful data each transaction carries, which is why coalescing dominates everything else in this mode.

Latency is the other half. A device access to local memory misses to DRAM and costs on the order of hundreds of cycles. An access to host memory crosses the bus and the root complex, and its latency is on the order of a microsecond, several times the local figure, and that is the optimistic case with no contention. A warp that stalls on such a load stalls for a long time, and the latency hiding that normally works by switching to other warps requires far more resident warps to cover it.

#include <cuda_runtime.h>

/* Host buffer mapped into the device address space. Every load below
   travels over PCIe; the coalescing pattern is what determines cost. */
__global__ void sum_coalesced(const float *host, float *out, int n)
{
    float acc = 0.0f;
    for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < n;
         i += blockDim.x * gridDim.x)
        acc += host[i];          /* consecutive threads, consecutive floats */
    if (acc == 1234.5f) *out = acc;   /* keep the loop from being removed */
}

int main(void)
{
    const int n = 1 << 20;
    float *h = NULL;

    /* Mapped, so the device can reach it, but not write-combined: this
       kernel only reads the buffer, and write-combining would make the
       host loop below slow, since those writes are never cached. */
    if (cudaHostAlloc((void **)&h, n * sizeof *h, cudaHostAllocMapped) != cudaSuccess)
        return 1;
    for (int i = 0; i < n; i++)
        h[i] = (float)i;

    /* With unified addressing this returns the same value as h. */
    float *d_ptr = NULL;
    cudaHostGetDevicePointer((void **)&d_ptr, h, 0);

    float *out = NULL;
    cudaMalloc((void **)&out, sizeof *out);

    sum_coalesced<<<64, 256>>>(d_ptr, out, n);
    cudaDeviceSynchronize();

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

Why coalescing decides the outcome

A warp’s 32 threads issue one memory instruction that the coalescer turns into the smallest number of transactions covering the addresses. When the addresses are 32 consecutive 4-byte values, the coalescer produces one transaction of 128 bytes, a full cache line, and the link carries one header for 128 bytes of payload. Efficiency is high, and the latency is paid once for 128 bytes.

When the addresses are strided by, say, one 4 KiB page each, the coalescer produces 32 separate transactions of 4 bytes each. The link now carries 32 headers for 128 bytes of payload. The effective bandwidth drops by roughly the ratio of header-plus-overhead to payload, which on PCIe can be an order of magnitude, and 32 separate round trips are in flight where one would have done. The kernel has not changed; only the address pattern has.

This is the single most important fact about zero-copy access: it rewards exactly the access pattern that device code already prefers, and punishes scattered access far more severely than local memory does. A pointer-chasing structure such as a linked list, a tree, or a chained hash table is the worst case, because each hop is a dependent 4-byte or 8-byte read with no opportunity to coalesce and no opportunity to overlap. A structure of arrays read sequentially by consecutive threads is the best case, and even then it is slower than a copy for anything that will be read more than once.

Write-combining and the direction asymmetry

Host memory mapped for device access can be allocated with a write-combining attribute. The CPU side of the mapping becomes uncached: writes from the host are buffered and flushed in bursts rather than cached, and reads from the host are slow because each one goes to DRAM. The device side gains the ability to write without triggering cache-coherence traffic on the host’s caches, since there is no cached copy to invalidate.

The asymmetry is the point. For a buffer the device fills and the host later reads once, write-combining is the right attribute: the device’s writes stream out efficiently, and the host pays for one uncached read pass. For a buffer the host writes and the device reads, write-combining is wrong: the host’s writes are uncached, and a host loop that writes one element at a time to uncached memory is dramatically slower than the same loop to cached memory. The attribute should be chosen by who writes and who reads, and mixing both directions in one buffer gets the worst of both.

There is a subtler trap on the device side. Writes to host memory from a kernel are posted: the device does not wait for a completion before continuing, so a kernel can issue many writes without stalling. Reads are not posted; the thread must wait for the response. A kernel that writes results to mapped host memory therefore runs at close to full speed for the write traffic, while a kernel that reads inputs from mapped host memory stalls on every access. The asymmetry can make a write-only mapped buffer look like a good deal and a read-only mapped buffer look like a disaster, on the same link.

When direct access actually wins

Three situations justify it. The first is small, latency-dominated data: a few kilobytes of parameters, flags, or counters that the kernel needs and that would cost more to schedule a copy for than to read across the bus. The second is data touched once, in a pattern that coalesces well, where the copy would double the traffic because the data must be read from host memory into a staging buffer and then written to device memory, while direct access reads it once. The third is irregular, sparse access where only a small fraction of a large structure is needed: copying the whole structure costs bandwidth proportional to its size, while direct access costs bandwidth proportional to what is touched.

The third case deserves care, because the touched fraction is usually not as small as it looks. A structure where 1% of the elements are read still pays a full transaction per element, and if those elements are scattered, the transactions are 4 bytes of payload each. The comparison is not “1% of the copy” but “1% of the elements, each at the full per-transaction cost of a scattered access”, which is a different bill and usually a worse one.

The situation where the argument collapses entirely is an integrated device that shares physical memory with the CPU. There, “host memory” is device-accessible at local latency, the mapping is not a bus transaction, and the whole analysis above does not apply. Code written for a discrete GPU’s zero-copy path can be dramatically slower or faster on an integrated part depending on which assumption it encoded, which is why the attribute should be treated as a platform-specific tuning decision rather than a portable design.

Interaction with migration and residency

Mapped host memory and managed memory are different mechanisms with superficially similar goals. A managed page can be migrated to the device, and after migration the kernel’s accesses are local; a mapped host page cannot be migrated, because the mapping is to a specific physical frame that the host owns. The kernel’s accesses stay remote forever.

That makes mapped memory a residency decision stated in advance: this buffer lives on the host, always. Managed memory is the opposite: the driver decides, possibly changing its mind repeatedly. Choosing between them is choosing between a predictable slow path and an adaptive path that is fast when the working set fits and pathological when it does not. The adaptive path’s failure mode, repeated migration between host and device for a buffer both sides touch, is described in GPU page faults and demand migration.

The pinning cost is the same in both cases and is not reduced by mapping. A mapped buffer is pinned for its lifetime, and the same limits on locked memory apply. See pinned host memory and asynchronous transfers for what that costs the rest of the system.

Failure modes

The first is a kernel that writes to a mapped buffer and expects the host to see the data immediately. Posted writes may still be in flight when the kernel completes, and the host must synchronize before reading. The synchronization that drains the stream also flushes these writes, which is part of what makes it expensive.

The second is mapping a buffer on one device and accessing it from another. A mapping is established in a particular device’s page tables; a second device needs its own mapping, and without peer support it will fault or, worse, resolve to a different frame. Multi-device mapping is covered in peer-to-peer transfers and link topology.

The third is assuming that because the pointer is valid on both sides, concurrent access is safe. It is not. A host thread writing to a mapped buffer while a kernel reads it is a data race with no coherence protocol to resolve it at fine granularity; the only safe patterns are strict alternation with synchronization, or atomics with an explicitly system-wide scope, as described in system-scope atomics and memory ordering.