PCIe & GPUDirect — The Bus Every GPU Sits On

Topics We'll Cover

Every GPU sits on PCIe, even the ones that mostly talk to siblings over NVLink. The board layer underneath the driver decides whether peer-to-peer works, how fast cold loads run, and whether your eight-GPU box is actually one fabric or two halves stitched through a CPU socket.

• PCIe in One Page
• GPU Generation → PCIe Generation
• BAR, BAR1 & Resizable BAR
• IOMMU & ACS
• P2P — GPU↔GPU Without Host Memory
• NUMA & Multi-Socket Pinning
• Reading the Topology
• GPUDirect P2P, RDMA, Storage
• Persistent Mode, Power, Cooling
• Multi-GPU on Consumer Boards — Where it Goes Wrong
• Interactive: PCIe Topology Lint

PCIe in One Page

PCI Express is bidirectional, packet-switched, point-to-point serial. Each lane is one differential pair per direction; a link is a bundle of lanes (x1, x4, x8, x16). Every generation roughly doubles the per-lane raw rate; the encoding evolves alongside (8b/10b on Gen1/2, 128b/130b from Gen3, PAM4 on Gen6).

Lane configurations you'll actually meet

• x16 — the headline GPU slot. Modern datacenter and consumer GPUs are designed for this.
• x8 — cheap inference cards (L4, T4) and many "second slot" desktop boards. Half the bandwidth of x16; for decoding the difference is barely visible, for prefill and weight-streaming it matters.
• Dual x16 — HGX baseboards, DPU-paired servers, and high-end workstations. One x16 to the GPU, another to a NIC or DPU sitting next to it.
• x4 — "the other slot" on consumer ATX. If you put a second GPU here you've quartered your link.

GPU Generation → PCIe Generation

Each GPU architecture pins itself to one PCIe generation at the host edge. The interesting twist: NVLink bandwidth on the same card has been pulling away from PCIe ever since Volta — the bus is the cheap path, NVLink is the fast one.

What a PCIe-only GPU loses

A "PCIe edition" H100 (no NVLink bridge across cards) ganged with another via PCIe 5 x16 sees ~64 GB/s peer-to-peer. The same pair of H100 SXM5 cards on a HGX baseboard sees 900 GB/s. Tensor parallelism on the PCIe pair is bandwidth-starved on every all-reduce; that's why DGX/HGX exist.

BAR, BAR1 & Resizable BAR

A PCIe device exposes its memory to the host through Base Address Registers (BARs). Each BAR is a window into the device's address space that the host CPU can map into its own virtual memory. For an NVIDIA GPU, BAR1 is the window onto VRAM — the aperture through which the CPU and other devices see GPU memory.

What ReBAR actually delivers

• Cold weight load: 2-5× faster on large models. Loading a 70B FP8 weight set drops from tens of seconds of wall clock to single digits.
• Direct kernel-side DMA: drivers can DMA into any region of VRAM without going through a bounce-buffer scheme.
• P2P bandwidth ceiling: with ReBAR off, peer-to-peer on PCIe is throttled because the source GPU can only see 256 MB of the destination at a time.
• Game/inference engines: any workload that streams large textures or weights wins.

lspci -vvv -s 01:00.0 | grep -A1 'Region 1:'
# Expect: Region 1: Memory at ... (64-bit, prefetchable) [size=24G]
# Bad   : Region 1: Memory at ... (64-bit, prefetchable) [size=256M]

# Or via NVIDIA driver:
nvidia-smi -q | grep -i bar1
# BAR1 Memory Usage
#     Total                             : 24576 MiB   ← full VRAM = ReBAR on
#     Total                             : 256 MiB     ← legacy = ReBAR off

IOMMU & ACS

The IOMMU (Intel VT-d, AMD-Vi) is the device-side equivalent of the CPU's MMU. It virtualises DMA addresses: a device's view of memory is translated by the IOMMU before reaching the memory controller. Two reasons it exists: VM passthrough, and protection from buggy or malicious DMA.

The two operating modes that matter

ACS — Access Control Services

ACS is a PCIe feature implemented by switches and root ports. It tells the fabric whether peer-to-peer traffic between two downstream devices should be routed directly (peer-to-peer DMA) or upstream-redirected (forced through the host root complex). ACS exists because the IOMMU can only enforce isolation on traffic it actually sees — if two devices talk peer-to-peer behind a switch, the IOMMU is bypassed.

• ACS off / disabled at root port (typical consumer ATX): all GPUs hanging off that root port end up in one IOMMU group. Peer-to-peer DMA works without switch redirection. Good for P2P, bad for VM isolation.
• ACS enabled by switch / server: peer traffic is bounced upstream. Each device gets its own IOMMU group. Good for VM isolation, blocks GPU-to-GPU P2P.
• ACS override patch: an out-of-tree kernel patch that lies to ACS so groups split where the silicon doesn't allow them to. Common on enthusiast multi-GPU desktop builds. Don't ship it on production.

# Which IOMMU mode is active?
dmesg | grep -E 'IOMMU|DMAR|AMD-Vi' | head -20

# Which devices share an IOMMU group?
for g in /sys/kernel/iommu_groups/*; do
    echo "Group ${g##*/}:"
    ls $g/devices
done

# Is ACS enabled on root ports / switches?
sudo lspci -vvv | grep -E 'IOMMU|ACSCtl'
# ACSCtl: SrcValid+ TransBlk- ReqRedir+ CmpltRedir+ ...   ← redirect on = blocks P2P
# ACSCtl: SrcValid- TransBlk- ReqRedir- CmpltRedir- ...   ← clean for P2P

P2P — GPU↔GPU Without Host Memory

Peer-to-peer DMA lets one GPU read or write the memory of another GPU directly across PCIe (or NVLink, when present), without staging through host RAM. The path is: GPU A → PCIe root complex (or switch) → GPU B. CPU is uninvolved beyond setting up the mappings.

Requirements for P2P to actually work

• Same IOMMU group (or compatible groups via ACS). If cudaDeviceCanAccessPeer(0,1) returns false, this is usually why.
• ReBAR on: without it, the source GPU can only DMA 256 MB into the destination at a time.
• ACS allows P2P: no upstream redirection at the relevant PCIe switch / root port.
• Driver enables it: cudaDeviceEnablePeerAccess(peer, 0) is required after the capability check.
• Same root complex / NUMA node: cross-socket P2P traffic that hops UPI usually fails or runs at a fraction of speed (covered next slide).

// Check capability
int can_p2p;
cudaDeviceCanAccessPeer(&can_p2p, 0, 1);
if (!can_p2p) { /* fall back to host-staged */ }

// Enable both directions
cudaSetDevice(0); cudaDeviceEnablePeerAccess(1, 0);
cudaSetDevice(1); cudaDeviceEnablePeerAccess(0, 0);

// Direct copy
cudaSetDevice(0);
cudaMemcpyPeerAsync(dst1, 1, src0, 0, n, stream);

// Or use the simpleP2P sample for a calibrated bandwidth measurement:
//   cuda-samples/Samples/0_Introduction/simpleP2P

NUMA & Multi-Socket Pinning

Two-CPU servers (and the few four-socket beasts that survive) are NUMA — non-uniform memory access. Each socket owns half the DRAM channels and half the PCIe lanes. Each GPU is wired to one socket. Cross-socket GPU traffic hops the inter-socket link — UPI on Intel, Infinity Fabric on AMD — which is slower and adds latency.

What it costs you

• P2P degraded or broken. nvidia-smi topo -m labels the link SYS. Most drivers refuse to enable P2P across sockets; those that allow it run at a tiny fraction of intra-socket speed.
• Host-staged transfers slower. A GPU on socket 0 reading host memory pinned on socket 1 hops UPI on every read.
• 5-15% throughput on prefill if your serving process runs on the wrong socket and pages from far DRAM.

The fix — pin everything

# Inspect topology
numactl --hardware
nvidia-smi topo -m
nvidia-smi topo -p2p r          # P2P read capability matrix

# Pin a serving process to NUMA 0 with GPUs on socket 0
CUDA_VISIBLE_DEVICES=0,1 \
    numactl --cpunodebind=0 --membind=0 \
    vllm serve meta-llama/Llama-3.1-70B-Instruct --tensor-parallel-size 2

# For multi-process MPI / torchrun: rank-to-NUMA mapping by local rank
for rank in 0 1 2 3; do
    node=$((rank / 2))            # 2 GPUs per socket on this box
    numactl --cpunodebind=$node --membind=$node \
        python worker.py --rank $rank &
done

Reading the Topology

nvidia-smi topo -m is the single most useful diagnostic for multi-GPU performance debugging. It prints a matrix of how every GPU pair is connected, using a fixed vocabulary of link classes.

Real example — 8×H100 SXM5 in DGX H100

      GPU0  GPU1  GPU2  GPU3  GPU4  GPU5  GPU6  GPU7  CPU Affinity  NUMA
GPU0   X    NV18  NV18  NV18  NV18  NV18  NV18  NV18  0-55,112-167   0
GPU1  NV18   X    NV18  NV18  NV18  NV18  NV18  NV18  0-55,112-167   0
GPU2  NV18  NV18   X    NV18  NV18  NV18  NV18  NV18  0-55,112-167   0
GPU3  NV18  NV18  NV18   X    NV18  NV18  NV18  NV18  0-55,112-167   0
GPU4  NV18  NV18  NV18  NV18   X    NV18  NV18  NV18  56-111,168-223 1
GPU5  NV18  NV18  NV18  NV18  NV18   X    NV18  NV18  56-111,168-223 1
GPU6  NV18  NV18  NV18  NV18  NV18  NV18   X    NV18  56-111,168-223 1
GPU7  NV18  NV18  NV18  NV18  NV18  NV18  NV18   X    56-111,168-223 1

Legend:  X = self  NV# = N NVLinks  PIX = single PCIe bridge  ...

Every pair shows NV18 — the NVSwitch fabric makes the 8 GPUs look like a single uniform NVLink domain. CPU affinity columns split 4/4 across two sockets, but the cross-socket SYS hop is hidden because all GPU-to-GPU traffic goes over NVSwitch, not over PCIe + UPI.

Real example — 2×RTX 4090 on a desktop board

      GPU0  GPU1  CPU Affinity  NUMA
GPU0   X    PIX   0-15           0
GPU1  PIX    X    0-15           0

# PIX = both GPUs hang off the same PCIe root port. P2P works.
# If instead you saw PHB or NODE, the second slot is on PCH-attached lanes
# and P2P will be slow or unreliable.

GPUDirect P2P, RDMA, Storage

"GPUDirect" is NVIDIA's umbrella for any path that lets a non-CPU device DMA into GPU memory or vice versa, bypassing host buffers. Three flavours actually matter for LLM and HPC work.

The data path, side by side

// Without GDS (standard POSIX read + cudaMemcpy):
NVMe SSD  →  PCIe  →  host page cache  →  user buffer  →  cudaMemcpyHtoD  →  GPU VRAM
          ^                                                       ^
          |  one PCIe traversal upstream                          |  another PCIe traversal back down
          |  CPU touches every byte                                  CPU touches every byte again

// With GDS (cuFile API):
NVMe SSD  →  PCIe  →  GPU VRAM
          ^
          |  single PCIe traversal, NVMe DMAs straight into VRAM
          |  CPU only manages the queue
          |  ~10x faster on a 70 GB checkpoint, >90% of NVMe headline BW

Persistent Mode, Power, Cooling

The bus discussion has knock-on effects on chassis, power, and driver state. Three operational details worth covering because they intersect directly with PCIe trade-offs.

Multi-GPU on Consumer Boards — Where it Goes Wrong

"I bolted a second 4090 in and now everything's slow." Here is the standard menu of failure modes, in roughly the order you should check them.

What "the right hardware" actually buys you

• Server-class motherboards (Threadripper PRO, EPYC, Xeon-W) split CPU lanes properly: 2 or 4 full x16 slots all connected directly to the CPU.
• PLX/Microchip PCIe switches let four GPUs share two CPU x16 lanes via one switch chip — consistent peer routing, predictable bandwidth, optional SR-IOV.
• HGX baseboards integrate NVSwitch + PCIe switches + dedicated NIC slots in one engineered piece. Expensive, but every consumer-board failure mode in the table above is engineered out.

Interactive: PCIe Topology Lint

Configure a hypothetical box. The lint tool reports per-GPU effective lanes, whether P2P will work, an estimated TP scaling factor, and the warnings you'd hit at deployment time.