Peer-to-Peer Transfers and the Topology That Decides Whether They Work

A copy from device 0 to device 1 through the host is two transfers: device 0 to host memory, then host memory to device 1. The data crosses the link twice, is written into host DRAM once and read out once, and both halves are scheduled on the copy engines of different devices, which means the two halves are serialized by the host-side synchronization that joins them. Peer-to-peer access replaces that with a single transfer whose descriptor names device 0’s memory as the source and device 1’s as the destination, with the data flowing directly over whatever path exists between them.

Whether that is possible, and how much it helps, depends entirely on the physical topology. The feature is often treated as a capability flag, supported or not, when it is really a property of the path between a specific pair of devices, and the interesting cases are the pairs where the path is not what the software assumes.

What peer access establishes

Enabling peer access from device 1 to device 0 asks the driver to install mappings in device 1’s page tables that translate addresses in device 0’s unified-address window to device 0’s physical memory, reachable over a path the hardware supports. Once installed, a kernel running on device 1 can dereference a pointer into device 0’s memory and the load is served remotely, without any host involvement and without any explicit copy.

Two properties follow. The first is that the mapping is directional and per pair: enabling 1 to access 0 does not enable 0 to access 1. The second is that the mapping is a page table entry, so the access is a remote load, not a local one. A kernel reading device 0’s memory from device 1 pays the interconnect latency for every access, exactly as a kernel reading host memory does. Peer access removes the copy; it does not make the memory local.

For copies rather than kernel accesses, the benefit is different. A copy engine on device 0 can be programmed with a descriptor whose destination is device 1’s memory, and the engine issues writes that travel over the peer path. The host is not involved, the data does not land in host DRAM, and the transfer is one hop rather than two. The reduction in link traffic is the point: on a PCIe topology where both devices hang from the same switch, a host-mediated copy crosses each device’s uplink once and the switch twice, while a peer copy crosses the switch once.

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

int main(void)
{
    int n = 0;
    cudaGetDeviceCount(&n);
    if (n < 2) return 0;

    int can = 0;
    /* Directional: can device 1 access device 0's memory? */
    if (cudaDeviceCanAccessPeer(&can, 1, 0) != cudaSuccess || !can)
        return 0;

    cudaSetDevice(1);
    /* Enabling is per pair and persists for the context. */
    if (cudaDeviceEnablePeerAccess(0, 0) != cudaSuccess) {
        cudaGetLastError();   /* already enabled, or unsupported */
    }

    float *a = NULL, *b = NULL;
    cudaSetDevice(0);
    cudaMalloc((void **)&a, 1 << 24);
    cudaSetDevice(1);
    cudaMalloc((void **)&b, 1 << 24);

    /* a is valid on device 1 because peer access mapped it there.
       With unified addressing the pointer value is unchanged. */
    cudaMemcpyPeerAsync(b, 1, a, 0, 1 << 24, 0);
    cudaDeviceSynchronize();

    cudaFree(b);
    cudaSetDevice(0);
    cudaFree(a);
    return 0;
}

Reading the topology before trusting it

The path between two devices is described by a topology query that reports the type of each link in the chain: a PCIe host bridge, a PCIe switch, a direct interconnect, or a NUMA node hop. The important distinction is between a path that stays inside one PCIe switch and a path that goes up to the root complex and back down.

Two devices under the same switch can exchange peer traffic without the traffic reaching the CPU’s root complex, which means it does not consume the host’s memory bandwidth and does not contend with host traffic. Two devices under different root complexes must traverse the inter-socket link, which is often the narrowest point in the machine, and peer traffic there competes with every other cross-socket access. The capability flag may be set in both cases; the performance is not comparable.

The worst case is a topology where peer access is reported as possible but the path is so indirect that the transfer is slower than the host-mediated alternative. This happens on virtualized platforms, where the hypervisor’s IOMMU configuration may present a peer path that is emulated or that routes through host memory anyway. The flag says the mapping can be installed; it does not say the mapping is fast.

Direct device interconnects change the calculus entirely. A link designed for device-to-device traffic carries far more bandwidth than a PCIe lane group and has lower latency, so peer access over it approaches local memory performance for bandwidth-bound work. The same code that is a modest win over PCIe becomes the obviously correct design over such a link, which is why topology-derived decisions should be made at run time rather than compiled in.

The failure modes of remote access

The first is assuming that peer memory is cached locally. It is not. Every access is a remote transaction, and the local caches hold copies that must be kept coherent with the remote memory. On some architectures, remote writes invalidate local cached lines, which means a kernel that reads the same remote line repeatedly pays the invalidation cost on each write from the other side. A loop over remote data that the other device is concurrently modifying is dramatically slower than the same loop over local data, and the slowdown is not explained by bandwidth.

The second is the access pattern problem, amplified. A coalesced read of 128 bytes from remote memory is one transaction; a scattered read of 32 separate 4-byte values is 32 transactions. On local memory the scattered case is several times slower; on remote memory it is far worse, because each transaction pays the interconnect latency and the link’s header overhead. Peer access rewards exactly the patterns that local memory rewards, and punishes the rest more severely.

The third is a partial failure that looks like success. Enabling peer access can succeed for a pair and still leave some allocations unmapped, if the allocation was made before the enable or by a different context. A kernel that follows a pointer into an unmapped peer region faults, and the fault is reported against the accessing device, not the one that owns the memory. The attribution is confusing enough that the first diagnosis is usually wrong.

The fourth is a deadlock built from symmetry. Two kernels on two devices, each waiting for a flag the other will set, is a deadlock regardless of the interconnect, and peer access makes it possible to write that pattern without any host-side transfer to break the symmetry. The system-scope atomics needed to express the flags correctly are covered in system-scope atomics and memory ordering; the deadlock is a design error that no memory ordering fixes.

Deciding between peer access and a host round trip

The comparison has three terms. A host-mediated copy costs two transfers over two links, plus a host memory write and read, plus the synchronization that joins the halves. A peer copy costs one transfer over the peer path. If the peer path has at least half the bandwidth of the host path and comparable latency, the peer copy wins on volume alone; if the peer path is narrow, the two-transfer route can be faster in wall-clock terms despite moving the data twice.

For kernel access rather than copies, the comparison is different. A host-mediated design copies the data to the device once and then accesses it locally at full speed. A peer design accesses it remotely on every touch. The crossover depends on how many times each byte is touched: once or twice, remote access is competitive; many times, the copy pays for itself. The same arithmetic governs the choice between direct host access and a copy, described in host-mapped memory over PCIe, and the reasoning transfers directly because both cases are remote access to memory another processor owns.

There is a middle option worth naming: a peer copy into a local buffer, followed by local access. That pays one transfer and then gets local bandwidth for the compute, which is usually the best of both when the data will be reused. It requires enough local memory to hold the working set, which brings back the capacity question that governs managed memory, covered in residency hints and thrashing under oversubscription.

Coordination when work spans devices

Peer access is a data path, not a scheduling mechanism. Deciding which device runs which part of a computation, and how the parts synchronize, is a separate problem with its own costs: every cross-device dependency is a synchronization point that may idle one device while the other works. The decomposition question is covered in partitioning work across devices.

The transport-level detail that makes any of this expressible, one address space in which a device’s memory has a name other devices can use, is the subject of unified virtual addressing and pointer provenance. And the ordering rules that decide when a peer write is visible to a peer reader are those of stream ordering and submission semantics, extended across devices only by explicit events.