P2P over NVLink

[osayamenja]

Hey folks!

I tried running p2p benchmarks on a GCP VM and it fails with No such file or directory: '/sys/class/infiniband/' which is due to here. Since EFA is optional, can we wrap that with a try-catch?

thanks.

[osayamenja]

So I added the below and tried running but the benchmark still failed with E1016 23:07:35.057309 18812 rdma_io.cc:68] [Error] Unable to get device list.

is_efa = False
if os.path.isdir("/sys/class/infiniband"):
    is_efa = "rdmap" in " ".join(os.listdir("/sys/class/infiniband/"))

Is it a hard requirement for UCCL that the underlying hardware be RDMA-compliant?

[YangZhou1997]

Yes, UCCL-p2p only supports RDMA NICs for now.

…

On Thu, Oct 16, 2025 at 4:10 PM Osayamen Aimuyo ***@***.***> wrote: *osayamenja* left a comment (uccl-project/uccl#449) <#449 (comment)> So I added the below and tried running but the benchmark still failed with Unable to get device list. is_efa = Falseif os.path.isdir("/sys/class/infiniband"): is_efa = "rdmap" in " ".join(os.listdir("/sys/class/infiniband/")) Is it a hard requirement for UCCL that the underlying hardware be RDMA-compliant? — Reply to this email directly, view it on GitHub <#449 (comment)>, or unsubscribe <https://github.com/notifications/unsubscribe-auth/AJFTPQUFZXM675J5XV6A63D3YAQWNAVCNFSM6AAAAACJNMX3IGVHI2DSMVQWIX3LMV43OSLTON2WKQ3PNVWWK3TUHMZTIMJTGE4TANBZGQ> . You are receiving this because you are subscribed to this thread.Message ID: ***@***.***>

[osayamenja]

Thanks for getting back @YangZhou1997! I was hoping to use P2P over NVLink, is that not possible?

[YangZhou1997]

Oh I see. We indeed want to support P2P over NVLink. But our own implementation currently causes some value errors when using cuda/hip IPC for Pytorch tensors, which might relates to some tensor address alignment issues, cc'ing @monopodium.

In addition to this, we find it is usually very hard to beat NCCL/RCCL on NVLink P2P performance, so we are now planning to integrate NCCL/RCCL for intra-node P2P under the hood.

[YangZhou1997]

@osayamenja we just fixed our P2P over NVLink! Now for intranode P2P, it will go through cudaMemcpy()