GPU P2P Memory Access, in depth: enabling peers, copies versus direct loads, cross-device ordering, CUDA IPC and proving the fast path

The model: UVA, a link and a mapping

Three pieces make P2P work. First, unified virtual addressing (UVA): on 64-bit platforms, host allocations and every GPU's allocations share one virtual address space, so a pointer identifies which device owns it. Second, the hardware path: two GPUs under the same PCIe switch, or connected by NVLink, can issue memory transactions to each other. Third, a mapping: the driver must map the peer's memory into the accessing GPU's page tables, which is what enabling peer access does.

With those in place, peer memory can be used two ways. A copy (cudaMemcpyPeerAsync, or cudaMemcpyAsync with cudaMemcpyDefault) is executed by a DMA copy engine; it moves a block efficiently and overlaps with compute. A direct access is a kernel on GPU 0 dereferencing a pointer that lives on GPU 1; each load or store becomes a remote transaction. Direct access avoids an intermediate buffer and lets you fuse the communication into a kernel, but it spends SM time waiting on a link with far higher latency than local memory, so it pays off for streaming, coalesced access and for small fine-grained exchanges, and loses badly for random access.

Discovering and enabling peers

Capability and enablement are separate steps, and access is per direction: enabling access from device 0 to device 1 does not let device 1 read device 0. The CUDA runtime documents that flags must be 0 and that a repeated call returns cudaErrorPeerAccessAlreadyEnabled, which callers should treat as success.

#include <cuda_runtime.h>
#include <cstdio>
#include <cstdlib>
#define CK(x) do { cudaError_t e_ = (x); if (e_ != cudaSuccess) { \
  fprintf(stderr, "%s:%d %s\n", __FILE__, __LINE__, cudaGetErrorString(e_)); \
  exit(1); } } while (0)

// Enable every possible peer mapping and print the capability matrix.
int main() {
  int n = 0;
  CK(cudaGetDeviceCount(&n));
  for (int i = 0; i < n; ++i) {
    CK(cudaSetDevice(i));                       // enablement applies to the current device
    for (int j = 0; j < n; ++j) {
      if (i == j) continue;
      int can = 0, rank = 0, atomics = 0;
      CK(cudaDeviceCanAccessPeer(&can, i, j));  // can i access j's memory?
      CK(cudaDeviceGetP2PAttribute(&rank, cudaDevP2PAttrPerformanceRank, i, j));
      CK(cudaDeviceGetP2PAttribute(&atomics, cudaDevP2PAttrNativeAtomicSupported, i, j));
      printf("%d -> %d  access=%d  perf_rank=%d  native_atomics=%d\n",
             i, j, can, rank, atomics);
      if (!can) continue;
      cudaError_t e = cudaDeviceEnablePeerAccess(j, 0);   // flags must be 0
      if (e == cudaErrorPeerAccessAlreadyEnabled) cudaGetLastError();  // clear, fine
      else CK(e);
    }
  }
  return 0;
}

Enable peers once at start-up, not per operation: mapping costs time and the mapping consumes resources on both devices. Log the matrix at start-up as well. It is the single most useful line in a multi-GPU incident, because the most common P2P bug is a machine that quietly lost it after a BIOS, driver or virtualisation change.

Copy or dereference

Copies are the safe default. The runtime picks the direct path when peer access is available and stages through host memory when it is not, so correctness never depends on the topology, only speed does. Direct access requires the mapping and fails with an illegal-address error if it is missing.

// Copy: DMA engine, overlaps with compute on other streams.
CK(cudaSetDevice(0));
CK(cudaMemcpyPeerAsync(dst_on_0, /*dstDevice=*/0, src_on_1, /*srcDevice=*/1, bytes, s0));

// Direct access: a kernel on GPU 0 reads GPU 1's buffer through UVA.
__global__ void add_from_peer(float* __restrict__ local,
                              const float* __restrict__ peer, size_t n) {
  size_t i = blockIdx.x * (size_t)blockDim.x + threadIdx.x;
  for (; i < n; i += (size_t)gridDim.x * blockDim.x)
    local[i] += peer[i];          // coalesced remote loads; requires 0 -> 1 enabled
}
add_from_peer<<<num_sms * 4, 256, 0, s0>>>(buf_on_0, buf_on_1, n);

The fused kernel saves a full write and re-read of an intermediate buffer, which is why custom all-reduce kernels for small messages in inference engines read peers directly. The copy version leaves the SMs free and usually reaches higher sustained bandwidth for large transfers. Measure both on your topology before choosing; the crossover depends on message size and on whether the link is NVLink or PCIe.

Ordering across devices

Streams order work within a device. Across devices you must add the edges yourself. The standard tool is an event: record it on the producer's stream on GPU 1, and make the consumer's stream on GPU 0 wait for it. cudaStreamWaitEvent accepts an event recorded on another device, and the wait happens on the GPU without blocking the host thread.

CK(cudaSetDevice(1));
produce<<<grid, block, 0, s1>>>(buf_on_1);
CK(cudaEventRecord(ready_1, s1));          // event created on device 1

CK(cudaSetDevice(0));
CK(cudaStreamWaitEvent(s0, ready_1, 0));   // s0 waits for GPU 1's kernel
add_from_peer<<<grid, block, 0, s0>>>(buf_on_0, buf_on_1, n);

Fine-grained protocols, where a kernel on one GPU writes data and then raises a flag that a kernel on another GPU polls, need memory fences, because writes over a link can become visible out of order. The writer issues __threadfence_system() between the data and the flag (or uses a release store with system scope via cuda::atomic_ref), and the reader polls with an acquire load or a volatile read followed by a fence. Get this wrong and the bug is a rare, load-dependent corruption. Prefer events unless you are writing a communication kernel on purpose, and check cudaDevP2PAttrNativeAtomicSupported before relying on atomics to peer memory.

Across processes: CUDA IPC

Training and serving stacks run one process per GPU, so the common case is not two devices in one process but two processes. CUDA IPC handles it: the owner calls cudaIpcGetMemHandle on the base of a cudaMalloc allocation and sends the resulting 64-byte handle to the other process over any channel; the receiver calls cudaIpcOpenMemHandle with cudaIpcMemLazyEnablePeerAccess and gets a device pointer it can copy from or dereference. The CUDA runtime reference restricts IPC to devices with unified addressing and, while it works on both Linux and Windows, describes the Windows support as a compatibility path that is not recommended because it costs performance. Check cudaDevAttrIpcEventSupport before relying on it.

// Process A (owner of the allocation)
cudaIpcMemHandle_t h;
CK(cudaIpcGetMemHandle(&h, base_ptr));       // base of the cudaMalloc allocation
send_bytes(sock, &h, sizeof h);              // plus the offset you care about

// Process B (importer)
cudaIpcMemHandle_t h;
recv_bytes(sock, &h, sizeof h);
void* base = nullptr;
CK(cudaIpcOpenMemHandle(&base, h, cudaIpcMemLazyEnablePeerAccess));
float* view = (float*)((char*)base + offset);
// ... use view ...
CK(cudaIpcCloseMemHandle(base));             // before the owner frees

Three rules prevent most IPC incidents. Ship offsets separately, because handles refer to whole allocations and caching allocators sub-allocate. Keep the owner's allocation alive until every importer has closed its mapping; freeing first leaves importers with dangling device pointers. Synchronise with IPC events, created with cudaEventInterprocess | cudaEventDisableTiming and shared via cudaIpcGetEventHandle. The newer virtual memory management driver API (cuMemCreate, cuMemExportToShareableHandle, cuMemImportFromShareableHandle, cuMemMap, cuMemSetAccess) exports a file descriptor instead, passed over a Unix socket, and gives explicit per-device access control; recent NCCL and framework versions use it.

Pools, allocators and frameworks

Stream-ordered allocations from cudaMallocAsync live in memory pools, and access to a pool from other devices is granted per pool with cudaMemPoolSetAccess rather than by the device-wide peer enable; pools also have their own export and import calls for sharing across processes. If a kernel faults reading a peer buffer that came from a pool even though peer access is enabled, this is the first thing to check.

Frameworks wrap all of this. PyTorch exposes torch.cuda.can_device_access_peer, routes cross-device copy_ through the peer path when it can, and shares CUDA tensors between processes in torch.multiprocessing via CUDA IPC, which is why the producer must keep the tensor alive while consumers hold it. Its caching allocator also explains why an IPC-shared tensor can keep a large block resident: the handle pins the whole underlying segment, not the tensor's slice. The site's NCCL collectives page covers the library that most training code uses instead of raw P2P. Within a node NCCL chooses between a P2P transport and shared-memory staging; NCCL_P2P_DISABLE=1 forces it off, and NCCL_P2P_LEVEL (for example NVL, PIX, PXB, PHB, SYS) sets the most distant topology at which it will still use P2P.

Worked example: the server that lost P2P

Suppose a new four-GPU PCIe server trains 30 percent slower than its predecessor and nobody changed the code. The diagnosis takes four commands. First, nvidia-smi topo -m prints the link type between each pair (NVLink, or PCIe distances such as PIX, PXB, PHB, SYS). Second, nvidia-smi topo -p2p r and -p2p w print read and write P2P status per pair, with codes such as OK, CNS (chipset not supported) and TNS (topology not supported). Third, the CUDA samples' p2pBandwidthLatencyTest measures bandwidth and latency with P2P disabled and enabled, side by side. Fourth, run the training job with NCCL_DEBUG=INFO and read the channel lines: a P2P route reports via P2P/..., a staged one via SHM.

A Python version of the measurement fits in a few lines and is worth keeping in the repository as a start-up self-test:

import time, torch

def copy_gbps(src: int, dst: int, nbytes: int = 1 << 30, iters: int = 20) -> float:
    a = torch.empty(nbytes, dtype=torch.uint8, device=f"cuda:{src}")
    b = torch.empty(nbytes, dtype=torch.uint8, device=f"cuda:{dst}")
    b.copy_(a)                                     # warm-up, maps peers if needed
    torch.cuda.synchronize(src); torch.cuda.synchronize(dst)
    t0 = time.perf_counter()
    for _ in range(iters):
        b.copy_(a, non_blocking=True)
    torch.cuda.synchronize(src); torch.cuda.synchronize(dst)
    return nbytes * iters / (time.perf_counter() - t0) / 1e9

n = torch.cuda.device_count()
for i in range(n):
    for j in range(n):
        if i != j:
            peer = torch.cuda.can_device_access_peer(i, j)
            print(f"{i}->{j} peer={peer} {copy_gbps(i, j):6.1f} GB/s")

In our hypothetical server the matrix shows access=0 on every pair, while the old machine showed 1. The usual culprit on bare metal is PCIe Access Control Services (ACS) or an IOMMU setting that redirects peer traffic up to the root complex, which the platform then refuses or routes slowly. On virtual machines, P2P depends on how the hypervisor passes the devices through. Consumer GeForce boards may report no P2P support with the stock driver regardless of topology. Fix the platform, re-run the four checks, and record the matrix as the machine's expected baseline.

Failure modes

• Silent host staging. Copies succeed at a fraction of the expected bandwidth. Assert the capability matrix at start-up.
• One direction enabled. GPU 0 reads GPU 1, but GPU 1's kernel faults reading GPU 0. Enable both directions explicitly.
• Illegal address on direct access. Peer access was never enabled, the pointer came from a pool without cudaMemPoolSetAccess, or it was freed.
• Hangs or poor bandwidth with ACS enabled. Peer transactions take the root-complex path; NCCL may hang or crawl. Fix in BIOS or kernel settings, or disable P2P for the job.
• Use-after-free across processes. The owner frees or the caching allocator reuses memory still mapped by an importer.
• Missing fences in flag protocols. The consumer sees the flag before the data; corruption appears only under load.
• Random-access kernels over the link. Direct access with scattered reads saturates latency, not bandwidth; copy first.

Trade-offs

Default to copies and NCCL, which degrade gracefully. Reach for direct access when profiling shows the intermediate buffer or kernel boundaries are the bottleneck, and when you control the topology the code will run on. For KV-cache hand-off between serving processes, the design questions are covered in P2P KV cache transfer, and for paths that include the network or storage see GPUDirect.

What to do next

1. Run nvidia-smi topo -m and nvidia-smi topo -p2p r on every machine type you operate, and save the output as the expected baseline.
2. Add a start-up check that logs cudaDeviceCanAccessPeer (or torch.cuda.can_device_access_peer) for all pairs and fails loudly on a mismatch with the baseline.
3. Run p2pBandwidthLatencyTest once per hardware generation and after BIOS or driver upgrades.
4. Read one NCCL_DEBUG=INFO log and confirm intra-node channels use P2P.
5. Audit IPC code for offset handling, owner lifetime and event-based synchronisation.
6. Only then prototype a direct-access kernel, and keep the copy path as its fallback.