Host-Device Transfers and the Cost of Crossing the Bus

A discrete accelerator is reached over an interconnect that was designed for device enumeration and general I/O rather than for bulk streamed data. A PCIe lane carries data at a signalling rate that yields a fraction of it after encoding, and a slot of sixteen lanes, the common configuration for a GPU, moves data at a rate that is one to two orders of magnitude below the accelerator’s own memory bandwidth. Every byte that crosses the bus is therefore expensive in a way that no kernel optimisation can recover.

The interface’s transfer functions look like memcpy. They are not: the copy path depends on whether the source pages can be locked, whether the transfer is synchronous, and whether the runtime is silently performing two copies instead of one.

Pageable memory forces a staging copy

Host memory returned by malloc is pageable. The operating system may move any page of it at any time, and it may satisfy a page fault by reading from disk. A device that performs direct memory access has physical addresses baked into its descriptor; if the page moves mid-transfer, the device writes to whatever now occupies that physical frame.

The runtime cannot prevent the move, so for pageable memory it does something else. It allocates a page-locked staging buffer inside the driver, copies the data from the caller’s buffer into it on the host, and issues the DMA from the staging buffer. For a device-to-host transfer the order reverses: DMA into the staging buffer, then a host copy out. The observable cost is bandwidth roughly halved, since the data moves across the host memory bus twice, and the copy consumes host CPU time cyclically as well.

Pinned memory removes the intermediate. cudaHostAlloc and CL_MEM_ALLOC_HOST_PTR produce pages that are locked into physical memory for their lifetime, so the device can DMA directly to and from them. The transfer then runs at the interconnect’s rate with no host-side copy in the critical path.

The price is paid elsewhere and is real. Locked pages cannot be evicted, so they reduce the memory available to every other process, and a large pinned allocation degrades the whole system. Allocating and freeing pinned memory is expensive, because it manipulates page tables over a range, so the pattern that works is a pool of pinned buffers allocated once and recycled, not a pinned allocation per transfer.

// Pinned staging buffer, allocated once, reused across iterations.
float *h_in;
cudaHostAlloc((void **)&h_in, BYTES, cudaHostAllocDefault);

for (int it = 0; it < iterations; ++it) {
    memset(h_in, 0, BYTES);                       // host fills the pinned pages
    cudaMemcpyAsync(d_in, h_in, BYTES, cudaMemcpyHostToDevice, stream);
    scale<<<grid, block, 0, stream>>>(d_in, d_out, n);
    cudaMemcpyAsync(h_in, d_out, BYTES, cudaMemcpyDeviceToHost, stream);
    cudaStreamSynchronize(stream);                // one wait per iteration
}
cudaFreeHost(h_in);

The kernel between the two copies touches the same buffer that was just transferred in, so the pair of transfers and the launch serialise on one stream; overlapping them requires the segmented structure described below.

Note that the async copy from pageable memory is legal but degenerate: the runtime must first synchronise with the host copy it performs internally, so the call blocks until the staging copy is done. Asynchrony over pageable memory is therefore partial: the DMA overlaps, the staging copy does not.

Bandwidth is a product, not a constant

The interconnect publishes a per-lane rate and a lane count. PCIe 3.0 transfers 8 gigatransfers per second per lane with 128b/130b encoding, giving roughly 985 megabytes per second per lane of usable bandwidth, so a sixteen-lane slot yields about 15.75 GB/s in each direction. PCIe 4.0 doubles the per-lane rate to 16 gigatransfers per second and 5.0 doubles it again to 32, both keeping the 128b/130b encoding, so usable bandwidth scales with the signalling rate: about 2 GB/s per lane at 4.0 and 4 GB/s at 5.0. The numbers are directional: a slot’s aggregate figure is the sum of both directions, and a copy in one direction gets half of it.

Against device memory bandwidth of the order of hundreds of gigabytes per second on current accelerators, the interconnect is the binding constraint by a wide margin. A kernel whose arithmetic intensity is high enough can hide the transfer; a kernel that streams through a large array and performs a handful of operations per element cannot, because the bus cannot feed it.

The consequence for interface design is that the API’s transfer functions are not where performance comes from; they are a tax that the program must minimise by moving less data and overlapping what it does move.

Overlapping transfer with computation

Two transfers and a kernel on one in-order stream serialise. With two streams and pinned memory, the pattern is to split the data into segments and pipeline: while segment n is being computed, segment n+1 is being transferred in and segment n-1 out.

The overlap is not free of contention, because the copy engine and the compute units both read device memory. When the kernel is bandwidth-bound, a concurrent transfer competing for the same memory controllers reduces the kernel’s throughput; the total wall time improves but by less than a full overlap would suggest. Copy engines are separate hardware from the compute units, which is what makes the overlap possible at all.

A fully asynchronous pipeline needs the host not to block, which brings back the queue semantics treated in Command Queues, Events, and Cross-Stream Ordering: the transfer must be enqueued with a stream or queue argument or it is synchronous by default, and cudaMemcpy without a stream is a blocking call for pageable sources in most cases.

Mapped host memory and zero-copy

A third path removes the explicit transfer. cudaHostAllocMapped (or clCreateBuffer with CL_MEM_USE_HOST_PTR and a mapped pointer) makes a pinned buffer addressable from the device. A kernel can then read and write host memory directly over the interconnect.

The performance is exactly what the bus permits: every access to the mapped region is a bus transaction, and since the access granularity is a cache line, a kernel with poor spatial locality in the mapped region wastes most of each transaction. Zero-copy is useful when the data is read once, in a streaming pattern, and the alternative would be a full copy of a large buffer that would evict useful device data. It is a poor choice for anything accessed more than once.

Mapped memory also requires the device to support it: CUDA exposes the check as cudaDeviceCanMapHostMemory, while in OpenCL the property is the device’s support for host-visible buffers rather than the CL_DEVICE_HOST_UNIFIED_MEMORY query, which reports whether host and device share one physical memory, a different question.

Unified memory changes the mechanism, not the arithmetic

Unified memory presents one address space to host and device, and the driver migrates pages on access faults instead of the program issuing copies. The programming model benefit is that a pointer can be passed anywhere without an explicit transfer, and the pointer’s value is valid on both sides.

The arithmetic of the bus does not change. A page that is migrated from host to device still crosses the interconnect, and a fault-driven migration is more expensive than an explicit copy because it is serialised against the faulting access and involves the page tables. Access patterns that thrash, alternating between a device-resident array and a host-resident neighbour, produce repeated migrations, and the cost is the bus cost multiplied by the number of round trips.

The mitigations are the same as they always were: keep the working set device-resident for the duration of the kernels, prefetch where the API permits it, and use the memory advice interfaces to tell the driver which device will touch a region next. Prefetching converts a fault into an asynchronous copy ahead of the access; advising converts a migration into a no-op by placing the page where it will stay.

Querying peer access before relying on it

On a multi-GPU system, a device-to-device copy between two accelerators on different PCIe root complexes traverses the host bridge, and the bandwidth is the slower of the two links, possibly with an additional hop. Peer-to-peer access over the interconnect bypasses host memory where the topology allows it, which requires both devices to be behind the same switch and the platform to permit it; the runtime reports availability through a capability query rather than assuming it from the device count.

The general form of the lesson applies to every function in this group: the transfer API reports a capability, the program queries it, and a program that assumes a fast path exists is correct on the machine it was written on and silently slow elsewhere. That pattern is the whole subject of Capability Negotiation Across Architectures, and transfers are the place where getting it wrong costs the most, because a wrong assumption here costs a factor rather than a few percent.