GPU-to-GPU P2P Bandwidth on Z890 (Dual RTX PRO 5000)

Two RTX PRO 5000s on an ASUS ProArt Z890-CREATOR WIFI: each card sits on its own CPU-attached Gen5 x8 link; the route between them is a PHB hop with no NVLink. Measured transfer rates across that route — direct peer copies and copies staged through pinned host memory. Current figures were taken with IOMMU translation off (intel_iommu=off iommu=off); August data was collected while VT-d was rejecting DMA from GPU 0 and is kept below, labeled, as a failure-mode record. The fault class itself: Fix NVIDIA Multi-GPU P2P Corruption on Intel Linux.

Test conditions

The spec blocks above carry the board, both GPUs with their PCI addresses, the driver, the current kernel, both power caps and the IOMMU state of each measurement era. Three conditions they do not carry:

• No LLM server or other CUDA workload was running during the probes.
• The boot used for the current measurements showed zero DMAR faults.
• The August fault-era runs were made on kernel 7.0.0-30-generic — see the fault-era section below.

Both cards are directly CPU-attached (separate root ports, not the chipset) and on the same NUMA node:

        GPU0    GPU1    CPU Affinity    NUMA Affinity   GPU NUMA ID
GPU0     X      PHB     0-17            0               N/A
GPU1    PHB      X     0-17            0               N/A

00:06.0 -> 02:00.0 -> GPU0
00:06.3 -> 03:00.0 -> GPU1

The cards idle at PCIe Gen1 x8 and ramp to Gen5 x8 under benchmark load — link readings taken after a run ends prove nothing; sample during transfers.

Current measurements (IOMMU disabled, 2026-09-14)

Custom probe, idle machine. Run record: /proc/cmdline with intel_iommu=off iommu=off, DMAR fault count 0 for that boot, peer matrices OK both directions.

GPU0 -> GPU1 (peer)      0.5 GB in 105.04 ms avg ->   5.11 GB/s
GPU1 -> GPU0 (peer)      0.5 GB in 103.97 ms avg ->   5.16 GB/s
GPU0 -> GPU1 repeat      0.5 GB in 105.60 ms avg ->   5.08 GB/s
32B copy x2048 avg latency: 4.05 us (launch-overhead-inclusive)

GPU0 -> GPU1 (peer)      1.2 GB in 235.67 ms avg ->   5.01 GB/s
GPU1 -> GPU0 (peer)      1.2 GB in 236.41 ms avg ->   4.99 GB/s
GPU0 -> GPU1 repeat      1.2 GB in 235.97 ms avg ->   5.00 GB/s
32B copy x2048 avg latency: 4.05 us (launch-overhead-inclusive)

The other transfer path — same binary, explicit staging mode (one DtoH + one HtoD through a pinned host buffer per iteration, timed over both hops):

GPU0 -> GPU1 (via pinned host)    0.5 GB in 37.51 ms avg ->  14.31 GB/s
GPU1 -> GPU0 (via pinned host)    0.5 GB in 37.42 ms avg ->  14.35 GB/s
GPU0 -> GPU1 repeat               0.5 GB in 37.52 ms avg ->  14.31 GB/s

GPU0 -> GPU1 (via pinned host)    1.2 GB in 82.34 ms avg ->  14.34 GB/s
GPU1 -> GPU0 (via pinned host)    1.2 GB in 82.25 ms avg ->  14.36 GB/s
GPU0 -> GPU1 repeat               1.2 GB in 82.39 ms avg ->  14.33 GB/s

Findings:

• Both paths are symmetric. The two peer directions differ by about 1%, and so do the staged ones.
• The ~5 GB/s peer ceiling is a property of this platform’s PHB route. It held across transfer sizes and repeats, with the links sampled during the transfers holding P1 at Gen5 ×8 on both cards (idle samples fall to P8/Gen1 — which is why post-run link readings prove nothing).
• Staging wins because it uses two root-port links instead of the inter-GPU route: each card talks to pinned RAM at near line rate, so a GPU→host→GPU copy finishes ahead of a direct peer cudaMemcpy on this board.
• Small-copy latency never looked impaired: 4.05 µs now versus 4.02–4.03 µs in August — consistent with layer-split inference being fast all along.
• CUDA_DISABLE_PEER_ACCESS=1 (the obvious “P2P off” switch) does nothing on driver 595.84: identical timings, access still reported. Staging figures come from explicit staging mode in the probe, not that env var.

What this means for inference here

• Layer-split inference moves tens of kilobytes per layer boundary, so it is latency-bound: Qwen 3.5 122B-A10B measures tg 95.8 t/s post-fix vs 96.0 pre-fix. Layer mode’s speed was never about peer bandwidth.
• Tensor parallelism talks across the inter-card link on every layer of every token, which is where peer behaviour bites. The post-fix control matrix — +61% generation from a 1/1 tensor split on a 27B model, tensor beating layer by +16% even at 122B — is on the RTX PRO 5000 page.
• Whether host staging helps an application depends on its copy patterns. At link level it wins on this box; llama.cpp’s staged build (-DGGML_CUDA_NO_PEER_COPY=ON) has not been measured.
• Judge a multi-GPU box by measured application-level behavior and a quiet kernel log — not by link labels, not by a capability matrix.

Reproduction commands

The host-staged variant is a flag, not a second program:

./gpu-p2p-bench 512 10          # peer DtoD (default)
./gpu-p2p-bench 512 10 host     # DtoH + HtoD through pinned RAM, both hops timed

# 1. Topology + NUMA, and the P2P capability matrices
nvidia-smi topo -m
nvidia-smi topo -p2p r && nvidia-smi topo -p2p w

# 2. Negotiated links while a transfer is in flight (should hold Gen5 x8)
watch -n 0.2 'nvidia-smi --query-gpu=index,pstate,pcie.link.gen.current,pcie.link.width.current --format=csv'

# 3. The custom probe these numbers came from (source in the collapsible block
#    at the bottom of this page)
nvcc -O3 -arch=native -o p2p_bench gpu-p2p-bench.cu
./p2p_bench 512 10 && ./p2p_bench 1126 10

# 4. Record the DMA state in the same sitting as the numbers
cat /proc/cmdline
sudo dmesg | grep -Ei 'DMAR|IOMMU' | head
journalctl -k --since boot | grep -ciE '\[DMA (Write|Read) NO_PASID\]'   # expect 0

NVIDIA’s p2pBandwidthLatencyTest, which produced the matrices in the historical section, no longer exists in cuda-samples master as of this check (only simpleP2P and streamOrderedAllocationP2P remain), so the embedded probe is the reference for both paths.

Always capture the kernel command line and DMAR fault count in the same record as the bandwidth numbers, and keep pre-fix and post-fix data in separate tables.

Fault-era measurements (August 2026, IOMMU translating)

Everything in this section was measured while VT-d was rejecting DMA writes from GPU 0 — it describes that failure mode, not this platform’s capability. Kernel evidence from the same window is just below.

DMAR: [DMA Write NO_PASID] Request device [02:00.0] fault addr 0xa000002000
      [fault reason 0x71] SM: Present bit in first-level paging entry is clear

The custom probe under translation:

GPU0 -> GPU1 (peer)      0.5 GB in 103.00 ms avg ->   5.21 GB/s
GPU1 -> GPU0 (peer)      0.5 GB in 276.99 ms avg ->   1.94 GB/s
GPU0 -> GPU1 repeat      0.5 GB in 102.61 ms avg ->   5.23 GB/s
32B copy x2048 avg latency: 4.03 us

GPU0 -> GPU1 (peer)      1.1 GB in 204.99 ms avg ->   5.24 GB/s
GPU1 -> GPU0 (peer)      1.1 GB in 566.05 ms avg ->   1.90 GB/s
GPU0 -> GPU1 repeat      1.1 GB in 207.08 ms avg ->   5.19 GB/s
32B copy x2048 avg latency: 4.02 us

NVIDIA’s p2pBandwidthLatencyTest reproduced the asymmetry almost exactly, and added the host-staged comparison from inside the fault window:

Unidirectional P2P=Disabled Bandwidth Matrix (GB/s)
   D\D     0      1
     0 1112.57  19.57
     1  14.35 1138.87

Unidirectional P2P=Enabled Bandwidth (P2P Writes) Matrix (GB/s)
   D\D     0      1
     0 1119.27   5.34
     1   2.05 1137.19

Bidirectional P2P=Disabled Bandwidth Matrix (GB/s)
   D\D     0      1
     0 1113.25  18.45
     1  18.38 1119.63

Bidirectional P2P=Enabled Bandwidth Matrix (GB/s)
   D\D     0      1
     0 1114.04   3.85
     1   3.89 1122.85

P2P=Disabled Latency Matrix (us)      P2P=Enabled Latency (us)
GPU0 -> GPU1:  ~14.39                 GPU0 -> GPU1:  ~0.39
GPU1 -> GPU0:  ~14.38                 GPU1 -> GPU0:  ~0.38

Trap: these tables are repeatable, which makes them look like a designed limit. They are not — the asymmetry (2 GB/s one direction) is the fault; the ~5 GB/s ceiling is real and survives the fix.

Application symptoms from the same window

Bandwidth tables only mislead; application output lying to you is harder to ignore. Multi-GPU llama.cpp and ik_llama.cpp runs using direct peer transfers produced output resembling random multilingual tokens and identifiers instead of text:

Glob headoinedeliverjudT head hint JPviewport infipersRated...

A longer run then failed inside CUDA near a 128K-token context position:

CUDA error: an internal operation failed
current device: 1, in function ggml_cuda_op_mul_mat_cublas
cublasGemmEx(...)

Controls, same weights and prompt:

Upstream llama.cpp documents this failure class: its multi-GPU guide notes that peer access “may cause crashes or corrupted outputs on some motherboards or BIOS configurations (e.g. when IOMMU is enabled)”, and that direct CUDA peer access is opt-in at runtime via GGML_CUDA_P2P. (See the multi-GPU document and the CUDA P2P options.)

The kernel evidence

The journal from that window carries the fault census, which is the part that settles whether DMAR NO_PASID lines are “harmless discovery messages” (they are not):

journalctl --no-pager | grep -ciE '\[DMA (Write|Read) NO_PASID\]'
journalctl --no-pager | grep -oiE 'Request device \[[0-9a-f:.]+\]' | sort | uniq -c
journalctl --no-pager | grep -oE 'dmar_fault: [0-9]+ callbacks suppressed' \
  | awk '{ if ($2 > m) m = $2 } END { print m }'

2980                                     # logged faults, first on 2026-08-14, last 2026-09-13
   2916 Request device [02:00.0]         # the 72 GB card
     64 Request device [03:00.0]         # the 48 GB card
largest single suppression burst: 3175687

Two things worth noting precisely:

• The window covers every multi-GPU measurement taken on this host up to then. The retained journal’s first fault is 2026-08-14; the P2P probes ran 2026-08-26.
• It was overwhelmingly one device, one fault reason. 2,916 of 2,980 logged faults name 02:00.0, and every one is reason 0x71 — “Present bit in first-level paging entry is clear”, i.e. the IOMMU’s page tables had no mapping for the address the GPU chose to write.

PCI inspection identifies the requestor as the 72 GB card:

lspci -nnk -s 02:00.0

02:00.0 VGA compatible controller: NVIDIA Corporation GB202GL …
Kernel driver in use: nvidia

A dropped write inside a tensor copy produces no CUDA error at the point of loss. The consumer of that buffer just reads whatever was there before — which is why the first symptom was nonsense tokens rather than an exception, and why the cuBLAS failure arrived much later, once a buffer that mattered had been corrupted badly enough.

IOMMU-disabled configuration (current state)

Booted with:

BOOT_IMAGE=/boot/vmlinuz-7.0.0-31-generic … intel_iommu=off iommu=off …

Verification, and the three things that confuse people about it:

sudo dmesg | grep -Ei 'DMAR|IOMMU'
readlink /sys/bus/pci/devices/0000:02:00.0/iommu_group
readlink /sys/bus/pci/devices/0000:03:00.0/iommu_group
ls /sys/class/iommu/

DMAR: IOMMU disabled
DMAR: Host address width 42                      # table parsing continues — this is
DMAR-IR: Enabled IRQ remapping in x2apic mode    # NOT "translation is still on"
(no readlink output — the per-device iommu_group links are gone)
(empty /sys/class/iommu/)

DMAR-IR: Enabled IRQ remapping is interrupt remapping, a separate feature from DMA isolation. The absent iommu_group symlinks are the practical confirmation — and also why VFIO/passthrough stops working here.

Capability re-checked after the change, both directions:

        GPU0    GPU1                    GPU0    GPU1
GPU0      X      OK            GPU0      X      OK
GPU1     OK       X           GPU1     OK       X
             (reads)                       (writes)

And the faults stopped: journalctl -k --since boot | grep -c NO_PASID → 0, over a boot that has run dual-GPU inference with both cards nearly full.

Security and virtualization tradeoffs

Disabling the IOMMU is not a free performance setting. Turning off DMA translation removes protection for:

• VFIO and PCIe device passthrough — needs DMA remapping; it will not work with intel_iommu=off.
• Virtual machines relying on IOMMU isolation of assigned devices.
• Defense against malicious or compromised DMA-capable devices — any device that can write arbitrary physical memory can compromise the host, and the IOMMU is what bounds that.
• Thunderbolt / PCIe-enclosure peripheral isolation, the case most likely to matter on a workstation with hot-pluggable ports.

For a dedicated, physically trusted inference box it can be an acceptable permanent setting; as a general “turn off VT-d for CUDA” tip it is not. If you need any of the above, investigate BIOS/kernel/driver fixes instead — and if this class of bug ever matters to you, note that upstream llama.cpp already treats GGML_CUDA_P2P as opt-in precisely because peer access is unreliable on some platforms.

Full procedure, including a temporary test at GRUB and revert steps: Fix NVIDIA Multi-GPU P2P Corruption on Intel Linux.

Short checklist

If your two-GPU box is slow or producing nonsense: check the kernel log for rejected DMA (sudo dmesg -T | grep -Ei 'DMAR|NO_PASID|Xid'), separate transport from model with a host-staged or RPC build on the same weights and prompt, and record /proc/cmdline plus the fault count in the same sitting as any numbers. The full procedure — including mapping the faulting device and reverting the change — is in the fix guide.

System:
GPU0: NVIDIA RTX PRO 5000 72GB Blackwell   (PCI 0000:02:00.0)
GPU1: NVIDIA RTX PRO 5000 Blackwell 48GB   (PCI 0000:03:00.0)
Driver: 595.84
CUDA reported by nvidia-smi: 13.2
Topology: PHB
NUMA: 0 for both GPUs
IOMMU: VT-d translation active (no intel_iommu=off); kernel 7.0.0-30-generic
NOTE (added 2026-09-14): DMAR [DMA Write NO_PASID] faults against 02:00.0 were
active on this host throughout this capture. Treat as a failure-mode record.

nvidia-smi topo -m:
        GPU0    GPU1    CPU Affinity    NUMA Affinity   GPU NUMA ID
GPU0     X      PHB     0-17    0               N/A
GPU1    PHB      X     0-17    0               N/A

PCIe capability / idle state:
GPU0 LnkCap: Speed 32GT/s, Width x16
GPU0 LnkSta: Speed 2.5GT/s, Width x8
GPU1 LnkCap: Speed 32GT/s, Width x16
GPU1 LnkSta: Speed 2.5GT/s, Width x8

Under benchmark load:
index, pstate, pcie.link.gen.current, pcie.link.width.current
0, P8, 5, 8
1, P8, 5, 8

Custom benchmark:
GPU0 -> GPU1 (peer)      0.5 GB in 103.00 ms avg ->   5.21 GB/s
GPU1 -> GPU0 (peer)      0.5 GB in 276.99 ms avg ->   1.94 GB/s
GPU0 -> GPU1 repeat      0.5 GB in 102.61 ms avg ->   5.23 GB/s
32B copy x2048 avg latency: 4.03 us

GPU0 -> GPU1 (peer)      1.1 GB in 204.99 ms avg ->   5.24 GB/s
GPU1 -> GPU0 (peer)      1.1 GB in 566.05 ms avg ->   1.90 GB/s
GPU0 -> GPU1 repeat      1.1 GB in 207.08 ms avg ->   5.19 GB/s
32B copy x2048 avg latency: 4.02 us

NVIDIA p2pBandwidthLatencyTest:
Unidirectional P2P=Disabled Bandwidth Matrix (GB/s)
   D\D     0      1
     0 1112.57  19.57
     1  14.35 1138.87

Unidirectional P2P=Enabled Bandwidth (P2P Writes) Matrix (GB/s)
   D\D     0      1
     0 1119.27   5.34
     1   2.05 1137.19

Bidirectional P2P=Disabled Bandwidth Matrix (GB/s)
   D\D     0      1
     0 1113.25  18.45
     1  18.38 1119.63

Bidirectional P2P=Enabled Bandwidth Matrix (GB/s)
   D\D     0      1
     0 1114.04   3.85
     1   3.89 1122.85

P2P=Disabled Latency Matrix (us)
GPU:
0 -> 1: 14.39
1 -> 0: 14.38

P2P=Enabled Latency (P2P Writes) Matrix (us)
GPU:
0 -> 1: 0.39
1 -> 0: 0.38

// GPU<->GPU P2P probe for this machine (PHB, no NVLink).
// Build: nvcc -O3 -arch=native -o p2p_bench gpu-p2p-bench.cu
// Run:   ./p2p_bench [sizeMB=512] [iters=10]
#include <cuda_runtime.h>
#include <cstdio>
#include <cstdlib>
#include <vector>

#define CK(x) do { cudaError_t e = (x); if (e != cudaSuccess) { \
  printf("CUDA error %s at line %d: %s\n", #x, __LINE__, cudaGetErrorString(e)); exit(1);} } while (0)

static double gbps(size_t bytes, float ms) { return (double)bytes / 1e9 / (ms / 1e3); }

int main(int argc, char** argv) {
    int ndev = 0; CK(cudaGetDeviceCount(&ndev));
    if (ndev < 2) { printf("need 2 GPUs\n"); return 1; }
    size_t bytes = (size_t)(argc > 1 ? atof(argv[1]) : 512.0) * 1024 * 1024;
    int iters = argc > 2 ? atoi(argv[2]) : 10;

    for (int d = 0; d < ndev; d++) {
        cudaDeviceProp p; CK(cudaGetDeviceProperties(&p, d));
        printf("GPU%d: %s (%zu GB)\n", d, p.name, p.totalGlobalMem >> 30);
    }

    int can01 = 0, can10 = 0;
    CK(cudaDeviceCanAccessPeer(&can01, 0, 1));
    CK(cudaDeviceCanAccessPeer(&can10, 1, 0));
    printf("peer access: 0->1 %s, 1->0 %s\n", can01 ? "yes" : "NO", can10 ? "yes" : "NO");
    if (!can01 || !can10) { printf("no P2P — copies fall back through host/system\n"); }
    else {
        CK(cudaSetDevice(0));
        CK(cudaDeviceEnablePeerAccess(1, 0)); // dev0 may read/write dev1
        CK(cudaSetDevice(1));
        CK(cudaDeviceEnablePeerAccess(0, 0)); // dev1 may read/write dev0
    }

    void *a0 = nullptr, *b0 = nullptr, *a1 = nullptr, *b1 = nullptr;
    CK(cudaSetDevice(0));
    CK(cudaMalloc(&a0, bytes));   // src on GPU0
    CK(cudaMalloc(&b0, bytes));   // dst on GPU0
    CK(cudaSetDevice(1));
    CK(cudaMalloc(&a1, bytes));   // src on GPU1
    CK(cudaMalloc(&b1, bytes));   // dst on GPU1

    auto bench = [&](int srcDev, int dstDev, void* src, void* dst, const char* tag) {
        CK(cudaSetDevice(dstDev));
        cudaEvent_t t0, t1;
        CK(cudaEventCreate(&t0)); CK(cudaEventCreate(&t1));
        // warmup
        CK(cudaMemcpy(dst, src, bytes, cudaMemcpyDefault));
        float total = 0.f;
        for (int i = 0; i < iters; i++) {
            CK(cudaEventRecord(t0));
            CK(cudaMemcpy(dst, src, bytes, cudaMemcpyDefault));
            CK(cudaEventRecord(t1));
            CK(cudaEventSynchronize(t1));
            float ms = 0.f; CK(cudaEventElapsedTime(&ms, t0, t1));
            total += ms;
        }
        printf("%-22s %5.1f GB in %.2f ms avg -> %6.2f GB/s\n", tag, bytes/1e9, total/iters, gbps(bytes, total/iters));
    };

    bench(0, 1, a0, b1, "GPU0 -> GPU1 (peer)");
    bench(1, 0, a1, b0, "GPU1 -> GPU0 (peer)");
    bench(0, 1, a0, b1, "GPU0 -> GPU1 repeat");

    // small-copy latency probe: 32-byte DtoD via peer, timed in a tight loop
    size_t small = 32;
    void *s0 = nullptr, *d1 = nullptr;
    CK(cudaSetDevice(0)); CK(cudaMalloc(&s0, small));
    CK(cudaSetDevice(1)); CK(cudaMalloc(&d1, small));
    const int nsmall = 2048;
    CK(cudaSetDevice(1));
    cudaEvent_t s0e, s1e; CK(cudaEventCreate(&s0e)); CK(cudaEventCreate(&s1e));
    CK(cudaMemcpy(d1, s0, small, cudaMemcpyDefault)); // warmup
    CK(cudaEventRecord(s0e));
    for (int i = 0; i < nsmall; i++) CK(cudaMemcpy(d1, s0, small, cudaMemcpyDefault));
    CK(cudaEventRecord(s1e));
    CK(cudaEventSynchronize(s1e));
    float ms = 0.f; CK(cudaEventElapsedTime(&ms, s0e, s1e));
    printf("32B copy x%d avg latency: %.2f us (launch-overhead-inclusive)\n", nsmall, ms*1000.f/nsmall);

    return 0;
}

Sources and upstream documentation

• llama.cpp multi-GPU guide — documents GGML_CUDA_P2P as opt-in and warns that peer access may corrupt output or crash with IOMMU enabled.
• Fix NVIDIA Multi-GPU P2P Corruption on Intel Linux — the diagnostic procedure, the fix, its tradeoffs and revert steps.
• NVIDIA p2pBandwidthLatencyTest — the CUDA sample used for the historical matrices; retired from cuda-samples master as of September 2026.
• cudaDeviceEnablePeerAccess — runtime API behind the custom probe’s peer setup.
• The Linux kernel command-line parameters — intel_iommu= and iommu=.
• Checking your GPU’s real PCIe bandwidth — link negotiation, LnkCap/LnkSta, and when downgrades matter.
• RTX PRO 5000 llama.cpp benchmarks — same host; the corrected tensor-split control matrix.