net: compute recv sameDevice against self, not the PXN proxyRank

[EylonKrause]

Description

recvSetup connects its proxy to the local rank — ncclProxyConnect(comm, TRANSPORT_NET, 0, myInfo->rank, &recv->proxyConn) — and the adjacent comment notes "We don't support PXN on receive yet", so the recv proxy always runs on this rank (self). But req.sameDevice is computed from proxyRank:

req.sameDevice = (comm->peerInfo[proxyRank].cudaDev == comm->cudaDev);

When graph != NULL, ncclTopoGetNetDev sets proxyRank to a PXN intermediate rank, which can be a different rank on a different GPU. So for the receiver sameDevice may be computed against the wrong device and come out false. The recv-side gate if (ncclGdrCopy && map->sameProcess && resources->sameDevice) then skips the GDRCopy SYNC/flush mapping, defeating the optimization on PXN-capable topologies even though the recv proxy is always on the kernel's own device. (The send path legitimately connects its proxy to proxyRank, so comparing against proxyRank is correct there.)

Related Issues

None.

Changes & Impact

• src/transport/net.cc (recvSetup): compute sameDevice against self (comm->peerInfo[myInfo->rank].cudaDev), matching where the recv proxy actually runs. The send path is unchanged. Behavior-preserving on non-PXN / single-NIC setups where proxyRank == myInfo->rank already; on PXN topologies it stops spuriously disabling recv-side GDRCopy. Non-breaking.

Performance Impact

Re-enables recv-side GDRCopy SYNC/flush on PXN-capable multi-GPU nodes where it was being spuriously disabled — a latency improvement on that path. No change elsewhere.

Testing

• Builds clean with make src.build; all_reduce_perf regression: Out of bounds values : 0 OK.
• Verified by inspection: the recv proxy is bound to myInfo->rank (and recvProxyConnect rejects a remote recv proxy), so self is the correct comparison; proxyRank can differ only via the graph/PXN path. The misgating manifests on multi-GPU / multi-NIC PXN nodes, which cannot be reproduced on a single-GPU machine.

[EylonKrause]

Disclosure: this contribution was authored with an AI coding assistant (Claude) and reviewed before submission.