From 7069fc936ff251b496bfadfb4ab20bd17e7d8b6b Mon Sep 17 00:00:00 2001 From: mberenjk <146776561+mberenjk@users.noreply.github.com> Date: Wed, 21 Jan 2026 11:08:12 -0800 Subject: [PATCH] Adding a check to respect DMABUF being disabled by the user (#2076) Co-authored-by: Marzieh Berenjkoub [ROCm/rccl commit: 9a443f305499e768095c4deabce075c5ad7927d5] --- projects/rccl/src/include/net.h | 1 + projects/rccl/src/transport/net_ib.cc | 2 +- 2 files changed, 2 insertions(+), 1 deletion(-) diff --git a/projects/rccl/src/include/net.h b/projects/rccl/src/include/net.h index f7f914b7cd..70f053e6f1 100644 --- a/projects/rccl/src/include/net.h +++ b/projects/rccl/src/include/net.h @@ -30,6 +30,7 @@ extern ncclNet_t ncclNetIb; extern ncclNet_t ncclNetSocket; extern ncclResult_t rcclNetP2pPolicy(void* handle, int isP2p); +extern int64_t ncclParamDmaBufEnable(); #if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__) extern ncclNet_t rocmNetIb; diff --git a/projects/rccl/src/transport/net_ib.cc b/projects/rccl/src/transport/net_ib.cc index deac53a3cf..7908db72ca 100644 --- a/projects/rccl/src/transport/net_ib.cc +++ b/projects/rccl/src/transport/net_ib.cc @@ -1990,7 +1990,7 @@ ib_recv: } } - useDmaBuf = (ncclIbDmaBufSupport(lComm->dev) == ncclSuccess); + useDmaBuf = (ncclIbDmaBufSupport(lComm->dev) == ncclSuccess && ncclParamDmaBufEnable()); rComm->flushEnabled = ((ncclIbGdrSupport() == ncclSuccess || useDmaBuf) && (ncclParamIbGdrFlushDisable() == 0)) ? 1 : 0; for (int i = 0; i < rComm->base.vProps.ndevs; i++) {