GPUDirect / GDR: zero-copy GPU↔NIC↔Network↔NIC↔GPU

The single biggest perf foot-gun on a multi-node GPU training cluster: NCCL works, RDMA works, you hit ~30% of the bandwidth you expected, and nobody can figure out why. The answer is almost always that GPUDirect RDMA isn't engaged, so every byte is bouncing through host RAM. The fix is small (load a kernel module, enable a kernel param, fix ACS in BIOS) but the diagnosis takes hours if you don't know what to look for.

This page is about what GDR is, the PCIe / IOMMU / ACS plumbing it depends on, and how to confirm it's actually working.

The problem GDR solves

Without GDR, the data path for GPU 0 on node A → NIC 0 → wire → NIC 0 on node B → GPU 0 on node B looks like:

1. CPU on node A: cudaMemcpy(pinnedHostBuf, gpuBuf, len, DeviceToHost) — DMA over PCIe from GPU to host RAM.
2. App posts RDMA WRITE with pinnedHostBuf as source. NIC DMAs from host RAM to wire.
3. Wire transit.
4. NIC on node B writes into pinnedHostBuf on node B via DMA.
5. CPU on node B: cudaMemcpy(gpuBuf, pinnedHostBuf, len, HostToDevice) — DMA from host RAM to GPU.

Two extra trips through host RAM, two extra PCIe traversals, twice the host bandwidth consumed. Throughput is gated by host RAM bandwidth (~50-100 GB/s shared across all uses) and by the cudaMemcpy overhead (microseconds per call).

With GDR, the data path is:

1. App posts RDMA WRITE with gpuBuf as source. NIC reads gpuBuf directly from GPU memory via PCIe peer-to-peer.
2. Wire transit.
3. NIC on node B writes into gpuBuf on node B directly via PCIe peer-to-peer.

Zero host RAM, zero CPU. The NIC and GPU talk to each other over the PCIe fabric directly. This is what the literature calls "peer-to-peer DMA" or "P2P".

The performance delta is brutal: a tested 4-node H200 cluster moved from ~1.2 GB/s socket-fallback to 363 GB/s busbw on NCCL allreduce after enabling GDR. ~300x.

How GDR actually works at the PCIe level

The NIC issues PCIe Memory Read/Write TLPs targeting GPU BAR1 addresses. BAR1 is the GPU's PCIe-visible aperture — a window into device memory that the host can map. On H100 it's 64 GB; on B200 it's 128 GB; on B300 it's 256 GB. The framebuffer (HBM) is exposed through this window.

For the NIC to read GPU memory, two things must hold at the hardware level:

1. The NIC and the GPU must share a common PCIe path that allows P2P. If the path goes through the CPU root complex, the root complex must support P2P (most do, with caveats), or peer-to-peer is silently disallowed.
2. The PCIe path must not have ACS-enabled switches between the two devices, because ACS by default forces all traffic up to the root complex. If ACS is on, P2P "works" but bounces through the root complex, defeating the purpose.

For the kernel software side:

• The NIC driver (mlx5_ib) needs to know how to register memory regions whose virtual address points to GPU memory.
• The GPU driver needs to expose its physical addresses to other PCIe devices.
• The bridge between the two is the nvidia-peermem (formerly nv_peer_mem) kernel module. It plugs into the IB subsystem so that ibv_reg_mr(pd, gpu_va, len, ...) works.

When nvidia-peermem is loaded, the IB stack can recognize a CUDA pointer at MR registration time, ask the NVIDIA driver for the physical pages backing it, and pin them for the NIC. Without it, ibv_reg_mr either fails (modern HCAs) or falls back to host-bounce (older).

Requirements checklist

iommu=pt in detail

The IOMMU translates DMA addresses for security/virtualization. In normal mode it builds per-device translation tables, which costs memory and adds a small (~5-10%) latency per DMA. Pass-through mode (iommu=pt) keeps the IOMMU enabled (so VFIO/passthrough still works) but uses identity mapping for the host kernel's devices — DMA addresses == physical addresses for kernel users.

Why GDR cares: the GPU driver gives nvidia-peermem raw physical addresses. If the IOMMU is in default (translating) mode, the NIC's DMA needs IOMMU mapping for those addresses. The IB stack can do that, but it's slower and on some platforms it doesn't get done correctly for cross-device P2P. With iommu=pt, the addresses are used as-is and the path is fastest.

iommu=off works too but disables IOMMU entirely — bad for any virtualization, bad for protected DMA.

$ cat /proc/cmdline
BOOT_IMAGE=/boot/vmlinuz-6.8.0 root=UUID=... ro intel_iommu=on iommu=pt

ACS in detail

ACS = Access Control Services. PCIe switches use it to enforce that traffic between two endpoints below the switch goes UP to the root complex first, instead of being switched directly. This is required for some virtualization use cases (preventing one VM-assigned device from talking to another).

On a GPU node, ACS is the enemy of P2P. If two endpoints (GPU + NIC) are below the same PCIe switch, ACS forces their traffic through the root complex. Two consequences:

1. Latency: extra hops.
2. Bandwidth: depending on root complex P2P support, you might get bounce-through-host-RAM instead of true P2P.

$ lspci -vvv | grep -A 2 "Access Control"
Capabilities: [220 v1] Access Control Services
        ACSCap: SrcValid+ TransBlk+ ReqRedir+ CmpltRedir+ UpstreamFwd+ EgressCtrl+ DirectTrans+
        ACSCtl: SrcValid- TransBlk- ReqRedir- CmpltRedir- UpstreamFwd- EgressCtrl- DirectTrans-

ACSCtl (control) tells you what's actively enabled. All - means ACS is off — good for P2P. Any + in SrcValid / UpstreamFwd / RequestRedirect means traffic is being forced upward.

To disable, three options:

1. BIOS option (best): "ACS Disable" or "PCIe ACS" toggle. On Supermicro, Dell, HPE platforms the option is usually under "Advanced > PCIe Configuration."
2. Kernel cmdline pcie_acs_override=downstream,multifunction — overrides ACS for downstream-port and multifunction-device cases. Works at boot but logs a warning.
3. Per-device runtime via setpci — fragile, breaks on driver reload.

For mixed-tenant clusters, leaving ACS on is sometimes a security requirement (preventing cross-tenant DMA). In that case you trade ~30% perf for isolation. Most single-tenant ML clusters disable ACS.

BAR1 sizing

nvidia-smi -q | grep -A 3 BAR1:

BAR1 Memory Usage
    Total                 : 65536 MiB
    Used                  : 64 MiB
    Free                  : 65472 MiB

64 GB on H100, 128 GB on B100/B200, 256 GB on B300. If your workload pins more GPU memory than BAR1 allows, GDR registration fails and you fall back to bounce. With 80 GB H100 framebuffers and 64 GB BAR1, you can pin slightly less than the whole GPU memory at once — usually fine; NCCL's working set is much smaller than full HBM.

If Total is far smaller than expected (e.g., 256 MB on H100), the BIOS isn't exposing the resizable BAR — enable "Above 4G Decoding" and "Resizable BAR" in BIOS, then power-cycle. Without it, GDR works but is heavily restricted.

Validating that GDR is engaged

Method 1: count nvidia-peermem clients

$ cat /proc/driver/nvidia-peermem/clients
mlx5_0 (in_use=1)
mlx5_1 (in_use=1)
mlx5_2 (in_use=1)
mlx5_3 (in_use=1)
mlx5_8 (in_use=1)
mlx5_9 (in_use=1)
mlx5_10 (in_use=1)
mlx5_11 (in_use=1)

(The exact path varies by driver version: try /sys/kernel/debug/nvidia-peermem/clients, /proc/driver/nv_peer_mem/clients, or dmesg | grep peer_mem.)

If this is empty while a NCCL job is running, GDR isn't being used.

Method 2: NCCL_DEBUG=INFO

NCCL INFO Channel 00 : 0[0] -> 1[0] [send] via NET/IB/0/GDRDMA
NCCL INFO Channel 00 : 0[0] -> 1[0] [send] via NET/IB/0/Read    <-- no GDRDMA = bouncing

GDRDMA in the path = engaged. Just NET/IB/0 (no GDRDMA suffix) = not engaged. NCCL also logs:

NCCL INFO NET/IB : Using [0]mlx5_0:1/RoCE [1]mlx5_1:1/RoCE [2]mlx5_2:1/RoCE [3]mlx5_3:1/RoCE [4]mlx5_8:1/RoCE [5]mlx5_9:1/RoCE [6]mlx5_10:1/RoCE [7]mlx5_11:1/RoCE
NCCL INFO NET/IB : GPU Direct RDMA Enabled for HCA 0 'mlx5_0' (read 1, write 1)

GPU Direct RDMA Enabled is the sentence you want to see, eight times (or however many NICs).

If you see GPU Direct RDMA Disabled, NCCL probed for it and the kernel said no. Reasons: peermem not loaded, IOMMU not in pass-through, ACS forcing root-complex routing, or NCCL_NET_GDR_LEVEL set too restrictively (default is OK).

Method 3: perftest with --use_cuda

$ ib_write_bw -d mlx5_0 --use_cuda=0 -F --report_gbits
$ ib_write_bw -d mlx5_0 --use_cuda=0 -F --report_gbits server-host

--use_cuda=0 makes both sides allocate the buffer on GPU 0. If GDR works, you'll see line rate (~388 Gb/s on a 400G NIC). If it falls back to bounce, you'll see ~50 Gb/s (host RAM bandwidth shared).

Method 4: check for the mover

$ cat /proc/<nccl-pid>/maps | grep nvidia
7f1234567000-7f1234d67000 rw-s ... /dev/nvidia0

GPU memory is mmap'd into the process. NIC then registers it as MR. If NCCL/process is memory-mapping /dev/nvidia0 AND pinning it via verbs, GDR is set up.

Failure modes

nvidia-peermem missing

$ lsmod | grep peermem
$ # nothing

Load it:

$ modprobe nvidia-peermem
$ lsmod | grep peermem
nvidia_peermem         16384  0

Add to /etc/modules-load.d/nvidia-peermem.conf to persist.

If modprobe fails with Module not found: you have nvidia-driver but not the matching peermem package. On Ubuntu: apt install nvidia-modprobe is not the right one — install the matching nvidia-fabricmanager-XXX and nvidia-peermem-XXX packages, or use the open-gpu-kernel-modules build.

On older systems with nv_peer_mem (now deprecated), it was a separate dkms package. Check which one your driver expects with modinfo.

ACS enabled silently

You see GDR clients populated, but bandwidth is half line rate. ACS may be redirecting to root complex.

$ for d in $(lspci -d 15b3: -n | awk '{print $1}'); do
    echo "=== mlx5 at $d ==="
    sudo lspci -vvv -s $d | grep -A 1 "ACSCtl"
  done

Any SrcValid+ or UpstreamFwd+ = ACS is forcing redirection. Disable in BIOS or boot with pcie_acs_override=downstream,multifunction.

IOMMU not in pass-through

$ dmesg | grep -i iommu
[    1.234567] DMAR: IOMMU enabled
[    1.234599] DMAR: dmar0: reg_base_addr ...

Doesn't tell you about pt vs default. Use:

$ cat /sys/class/iommu/dmar0/intel-iommu/version       # exists on Intel
$ cat /proc/cmdline | grep -o 'iommu=[a-z]*'
iommu=pt

If iommu=pt is missing, add to GRUB and reboot. Verify after reboot.

BAR1 too small

$ nvidia-smi -q | grep -A 3 BAR1
BAR1 Memory Usage
    Total                 : 256 MiB

256 MiB on a B200 = BIOS isn't exposing resizable BAR. Enter BIOS → enable "Above 4G Decoding" + "Resizable BAR" + reboot. Some servers also need NVIDIA Fabric Manager bundle / VBIOS update.

GDR disabled by NCCL env

Check:

$ env | grep -E 'NCCL_(IB_DISABLE|NET_GDR_LEVEL|IB_GDR_DISABLE|GDR)'

Common offenders:

• NCCL_NET_GDR_LEVEL=0 disables GDR entirely.
• NCCL_IB_DISABLE=1 disables IB transport entirely → falls to socket → no GDR.
• NCCL_NET_GDR_READ=0 disables GDR-on-read (sender pulling) but keeps GDR-on-write.

Default NCCL_NET_GDR_LEVEL=PIX (or 5 in older NCCL) enables GDR when the GPU and NIC are within one PCIe switch. Lower values are more permissive.

See also

• PCIe topology + ACS
• RDMA fundamentals
• RoCE v2 + DCB/PFC/ECN
• NCCL multi-node tuning
• Network performance tuning