Multi-GPU Peer-to-Peer Access, in depth: the server as a graph, NCCL transport choice, one-shot all-reduce over peer pointers and placing groups on P2P cliques

The server as a graph

Build the graph from the per-pair answers: an edge exists if the two GPUs can access each other's memory, and its weight is the measured bandwidth. On an NVSwitch system (DGX or HGX class) the graph is complete with uniform weights: any group of GPUs is equally good. On a PCIe server it splits into islands. GPUs under the same PCIe switch talk at full link speed, GPUs on different switches under one CPU route through the root complex, and GPUs on different sockets cross the CPU interconnect, where P2P is often unsupported or slow. NVLink bridges on workstation cards add a third shape: fast pairs inside an otherwise PCIe graph.

The property that matters for collectives is the clique: a set of GPUs in which every pair has a fast edge. Ring algorithms only need a fast cycle, but direct all-to-all patterns, including the custom all-reduce kernels used by inference engines, need a full clique. The following script derives cliques from what the runtime reports:

import itertools, torch

def p2p_graph():
    n = torch.cuda.device_count()
    return {(i, j) for i in range(n) for j in range(n)
            if i != j and torch.cuda.can_device_access_peer(i, j)}

def is_clique(group, edges):
    return all((a, b) in edges and (b, a) in edges
               for a, b in itertools.combinations(group, 2))

def cliques_of_size(k):
    edges, n = p2p_graph(), torch.cuda.device_count()
    return [g for g in itertools.combinations(range(n), k) if is_clique(g, edges)]

print("TP=2 candidates:", cliques_of_size(2))
print("TP=4 candidates:", cliques_of_size(4))

Capability is not bandwidth. Two GPUs can report P2P access across the root complex at a fraction of switch speed, so weight the edges with a measured copy rate (the pairwise page has a ready-made script) before trusting a clique.

How NCCL chooses a transport

NCCL builds its own version of this graph at communicator creation and chooses a transport per peer connection. The intra-node options are P2P (direct GPU-to-GPU over NVLink or PCIe), SHM (staging through host shared memory, used when P2P is unavailable or judged slower), and, across nodes, NET (InfiniBand or sockets, optionally with GPUDirect RDMA). With NCCL_DEBUG=INFO the log prints each channel's route, and lines that mention via P2P or via SHM show the choice directly.

Two environment variables control the decision. NCCL_P2P_DISABLE=1 turns the P2P transport off entirely. NCCL_P2P_LEVEL sets the maximum topological distance at which NCCL will still use P2P, and takes the same vocabulary as nvidia-smi topo -m:

Leave it unset by default; NCCL picks based on the detected platform. Set it when you have measured that its choice is wrong for your machine, for example forcing PIX on a box where cross-switch P2P is technically allowed but slower than host staging, or when P2P across sockets causes hangs. The algorithm-level picture of ring and tree all-reduce is in NCCL collectives.

All-reduce directly over peer pointers

NCCL's ring all-reduce is bandwidth-optimal: each GPU sends and receives about 2(N-1)/N times the buffer size. It pays for that with 2(N-1) sequential steps, each with its own synchronization. For the small all-reduces in tensor-parallel inference, often tens of kilobytes per layer per token, latency dominates and the ring's step count is the cost.

Peer pointers allow a different design. Each GPU writes its input into a buffer that every peer has mapped (via CUDA IPC handles exchanged once at start-up). In a one-shot all-reduce, every GPU then reads all N buffers directly and sums them: one step and one barrier, but each GPU reads (N-1) times the buffer, so bandwidth cost grows with N. In a two-shot all-reduce, each GPU reduces only its 1/N slice from all peers, then every GPU gathers the reduced slices: two steps, roughly ring-like traffic. Engines choose one-shot for the smallest messages, two-shot for medium ones and NCCL for large ones.

// One-shot all-reduce over peer pointers (simplified; real kernels vectorize and
// use flag-based barriers on mapped memory instead of a host sync).
__global__ void oneshot_allreduce(float* const* peer_bufs,   // N mapped input buffers
                                  float* out, int n_ranks, size_t count,
                                  volatile int* const* peer_flags, int rank, int epoch) {
    // 1. barrier: signal "my input is ready" to every peer, wait for theirs
    if (blockIdx.x == 0 && threadIdx.x < n_ranks) {
        peer_flags[threadIdx.x][rank] = epoch;               // remote write over P2P
        while (peer_flags[rank][threadIdx.x] != epoch) {}    // spin on local flag
    }
    __syncthreads();   // per-block only: production kernels barrier in every block
    // 2. every rank reads every peer's buffer and reduces locally
    for (size_t i = blockIdx.x * blockDim.x + threadIdx.x; i < count;
         i += gridDim.x * blockDim.x) {
        float acc = 0.f;
        for (int r = 0; r < n_ranks; ++r) acc += peer_bufs[r][i];   // P2P loads
        out[i] = acc;
    }
}

The design only works when every GPU can load from every other GPU at full speed, which is why it needs a clique. vLLM's custom all-reduce checks P2P support and, for more than two GPUs, whether they are fully NVLink-connected; if not, it disables itself and uses NCCL, logging a warning that suggests --disable-custom-all-reduce to silence it. Reports of hangs at start-up on some PCIe platforms are usually resolved with the same flag. PyTorch is also moving in this direction with symmetric-memory primitives, which give each rank peer-mapped buffers for custom collectives; those APIs are still evolving, so check your version's documentation before relying on them.

Placing processes on cliques

The graph only helps if processes land on the right GPUs. By default CUDA enumerates devices "fastest first", which can differ from the PCI bus order that nvidia-smi prints. Set CUDA_DEVICE_ORDER=PCI_BUS_ID so indices match the topology tools, then use CUDA_VISIBLE_DEVICES to hand each job a clique. For multi-process launches, map ranks so that each tensor-parallel group is one clique and data-parallel replicas span cliques, because data-parallel traffic is larger but far less latency-sensitive.

# 8-GPU PCIe box, two sockets, GPUs 0-3 under switches on socket 0, 4-7 on socket 1.
export CUDA_DEVICE_ORDER=PCI_BUS_ID

# Two independent inference replicas, each TP=4 inside one socket:
CUDA_VISIBLE_DEVICES=0,1,2,3 vllm serve $MODEL --tensor-parallel-size 4 --port 8000 &
CUDA_VISIBLE_DEVICES=4,5,6,7 vllm serve $MODEL --tensor-parallel-size 4 --port 8001 &

# Training: torchrun ranks are contiguous, so build TP groups from contiguous ranks
# (0-3, 4-7) and DP groups across them ({0,4}, {1,5}, ...).

Pin CPU threads and host memory to the same NUMA node as the GPUs too; the SHM fallback and the data loader both cross sockets otherwise. How tensor-parallel groups are formed is covered in tensor parallelism.

Containers and virtual machines

Containers share the host's PCIe topology, so P2P generally works if every GPU of a job is visible to the container. Two things break. First, NCCL's SHM transport and CUDA IPC need shared memory and process visibility between ranks; small default /dev/shm sizes cause failures, which is why inference images are usually run with --ipc=host or a large --shm-size. Second, splitting one job's GPUs across containers that cannot see each other removes the IPC path even when the hardware path exists.

Virtual machines are harder. With PCIe passthrough the hypervisor decides what topology the guest sees, and the guest often sees GPUs behind a flat virtual root rather than the real switches. Peer traffic may then be routed through the IOMMU or blocked, and the guest's P2P matrix can differ from the host's. NVLink-connected GPUs passed through together usually keep NVLink. Always rebuild the graph inside the guest; never assume the host's.

Worked example: an eight-GPU PCIe server

Take the box in panel B extended to eight GPUs: two PCIe switches per socket, two GPUs per switch. On this host, cliques_of_size(4) returns {0,1,2,3} and {4,5,6,7} because P2P is allowed through the root complex within a socket but not across sockets. Measured bandwidth tells a finer story: about 25 GB/s between GPUs on the same switch (PCIe Gen4 x16), and noticeably lower across switches under the same CPU.

Serving a model that needs TP=4, the team runs two replicas pinned per socket, as in the placement script. vLLM logs that custom all-reduce is disabled, because four PCIe GPUs are not fully NVLink-connected, and uses NCCL. NCCL's log shows P2P inside each switch and P2P or SHM across switches. A test run with NCCL_P2P_LEVEL=PIX shows whether host staging is faster than the cross-switch P2P path on this platform; the faster setting is kept and recorded. If TP=2 is enough for the model, four replicas on {0,1}, {2,3}, {4,5}, {6,7} each get the fast switch path and two-GPU custom all-reduce, which is usually the better throughput configuration on such a box.

Failure modes

• Group spans islands. One slow edge makes the whole TP group run at its speed. Place groups on cliques.
• Index mismatch. Without CUDA_DEVICE_ORDER=PCI_BUS_ID, GPU 2 in your code is not GPU 2 in nvidia-smi, and placement built from the topology table is wrong.
• Silent fallback. Custom all-reduce disables itself and latency rises; read start-up logs and alert on the warning.
• Forced levels that outlive the hardware. A NCCL_P2P_LEVEL or NCCL_P2P_DISABLE set for one machine is baked into an image and slows the next one.
• Container shared memory too small. SHM transport fails or hangs at initialization.
• Guest topology differs from host. A VM loses P2P that bare metal had.

Trade-offs

NVSwitch systems remove the placement problem at a high price per GPU; PCIe servers are cheaper and fine for inference and data parallelism, but force group sizes to match islands. Custom P2P collectives cut small-message latency but add kernels to maintain and depend on a full clique; NCCL is portable and tuned but pays more steps. Forcing transports by environment variable fixes a specific machine and can hurt others, so keep such settings per host type, not global. At rack scale the same reasoning moves to the network: see SuperPOD topology.

What to do next

1. Run the clique script on each host type, weight edges with measured bandwidth, and store the result with the machine image.
2. Set CUDA_DEVICE_ORDER=PCI_BUS_ID everywhere and place each tensor-parallel group on a clique.
3. Start one job with NCCL_DEBUG=INFO and confirm P2P versus SHM per channel.
4. Only if measurements justify it, set NCCL_P2P_LEVEL per host type.
5. For inference, grep start-up logs for custom all-reduce being disabled and decide whether a smaller TP degree on a clique serves better.
6. In containers use --ipc=host or a large shared-memory size; in VMs rebuild the graph inside the guest.