Silent data corruption in GPU-to-GPU (P2P) copies between two RTX 6000 Ada (open kernel module 580.142)

[mvilucchio]

Description

Direct device-to-device copies between two RTX 6000 Ada GPUs silently return corrupted data. cudaDeviceCanAccessPeer reports 1 and every API call returns cudaSuccess, but the copied buffer does not match the source. The same data staged through host memory is bit-exact. No error, warning, or dmesg output — the corruption is completely silent, which caused multi-GPU training (PyTorch DataParallel weight broadcasts) to diverge with no error message.

Environment

• 2× NVIDIA RTX 6000 Ada Generation (no NVLink; nvidia-smi topo -m reports NODE — PCIe across host bridges within one NUMA node)
• Driver: 580.142, Open Kernel Module
• CUDA runtime: 13.0
• CPU: Intel Xeon w7-3465X (Sapphire Rapids), single socket
• OS: Ubuntu 22.04, kernel 6.8.0-111-generic
• Kernel cmdline has no IOMMU flags (quiet splash only)

Reproduction (raw CUDA runtime API, no frameworks)

cudaMemcpyPeer with peer access explicitly enabled, via ctypes:

import ctypes
import numpy as np

rt = ctypes.CDLL("libcudart.so")
rt.cudaGetErrorString.restype = ctypes.c_char_p


def check(code, what, tolerate=()):
    if code != 0 and code not in tolerate:
        raise RuntimeError(f"{what}: {rt.cudaGetErrorString(code).decode()} ({code})")


N = 256 * 256
NBYTES = N * 4
host_src = np.random.default_rng(0).standard_normal(N).astype(np.float32)
host_dst = np.empty_like(host_src)

can = ctypes.c_int(0)
check(rt.cudaDeviceCanAccessPeer(ctypes.byref(can), 0, 1), "canAccessPeer")
print(f"cudaDeviceCanAccessPeer(0 -> 1) = {can.value}")

check(rt.cudaSetDevice(0), "setDevice 0")
check(rt.cudaDeviceEnablePeerAccess(1, 0), "enablePeer 0->1", tolerate=(704, 601))
check(rt.cudaSetDevice(1), "setDevice 1")
check(rt.cudaDeviceEnablePeerAccess(0, 0), "enablePeer 1->0", tolerate=(704, 601))

d0 = ctypes.c_void_p()
d1 = ctypes.c_void_p()
check(rt.cudaSetDevice(0), "setDevice 0")
check(rt.cudaMalloc(ctypes.byref(d0), NBYTES), "malloc dev0")
check(rt.cudaMemcpy(d0, host_src.ctypes.data_as(ctypes.c_void_p), NBYTES, 1), "H2D dev0")
check(rt.cudaSetDevice(1), "setDevice 1")
check(rt.cudaMalloc(ctypes.byref(d1), NBYTES), "malloc dev1")

check(rt.cudaMemcpyPeer(d1, 1, d0, 0, NBYTES), "memcpyPeer 0->1")
check(rt.cudaDeviceSynchronize(), "sync")
check(rt.cudaMemcpy(host_dst.ctypes.data_as(ctypes.c_void_p), d1, NBYTES, 2), "D2H dev1")

exact = bool(np.array_equal(host_src, host_dst))
print(f"raw cudaMemcpyPeer 0 -> 1 bit-exact: {exact}")
if not exact:
    bad = np.flatnonzero(host_src != host_dst)
    print(f"  {bad.size}/{N} elements differ, first at index {bad[0]}")

Output on this system:

cudaDeviceCanAccessPeer(0 -> 1) = 1
raw cudaMemcpyPeer 0 -> 1 bit-exact: False
  49152/65536 elements differ, first at index 0

Also reproducible in one line from PyTorch 2.10.0+cu130 (x.to("cuda:1") of a tensor on cuda:0 differs from x in all 65536 elements, max abs error ~5.1, while x.cpu().to("cuda:1") is exact), so the corruption is independent of the client application.

Expected behavior

Either the copy is bit-exact, or P2P is reported as unsupported (cudaDeviceCanAccessPeer = 0) so the runtime falls back to host staging. Silent corruption with all-success return codes is the worst possible failure mode.

Notes

• Happy to attach nvidia-bug-report.log.gz, lspci -vvv (ACS state), or run nvbandwidth / cuda-samples simpleP2P on request.