diff --git a/src/include/utils.h b/src/include/utils.h index 383f678c87..7094891bdb 100644 --- a/src/include/utils.h +++ b/src/include/utils.h @@ -545,4 +545,25 @@ T* ncclIntruQueueMpscAbandon(ncclIntruQueueMpsc* me) { return head; } } +/** + * @brief function to get page size of the system + */ +size_t get_sc_page_size(void); + +/** + * @brief function to get system's page size aligned memory address and buffersize + * + * Given a pointer `ptr` to a buffer of size `bufsize`, this function computes: + * 1. A new pointer `aligned_ptr` which points to the start of the page-aligned memory region that includes `ptr`. + * 2. A new size `aligned_size` that is the minimum number of bytes (aligned to page size) needed to cover the original buffer from `aligned_ptr`. + * + * This is useful, for example, when performing operations such as memory mapping or advising memory usage (e.g., with `mmap`, `madvise`, `mlock`, etc.), which often require page-aligned memory ranges. + * This function doesn't dereferece the input pointer + * + * @param[in] ptr Pointer to the start of the original memory buffer. + * @param[in] bufsize Size (in bytes) of the original buffer starting at `ptr`. + * @param[out] aligned_ptr Pointer to a variable that will be set to the aligned base address. + * @param[out] aligned_size Pointer to a variable that will be set to the aligned size. + */ +void get_aligned_ptr_and_size(const void *ptr, const size_t bufsize, void **aligned_ptr, size_t *aligned_size); #endif diff --git a/src/misc/utils.cc b/src/misc/utils.cc index bb59947e46..ee72bf60d3 100644 --- a/src/misc/utils.cc +++ b/src/misc/utils.cc @@ -8,7 +8,7 @@ #include "core.h" #include "nvmlwrap.h" - +#include #include // Get current Compute Capability @@ -289,3 +289,22 @@ void ncclMemoryStackDestruct(struct ncclMemoryStack* me) { h = h1; } } + +size_t get_sc_page_size() { + static size_t cached_page_size = 0; + size_t ps = __atomic_load_n(&cached_page_size,__ATOMIC_RELAXED); + if (ps == 0) { + ps = (size_t)sysconf(_SC_PAGESIZE); + __atomic_store_n(&cached_page_size, ps,__ATOMIC_RELAXED); + } + return ps; +} + +void get_aligned_ptr_and_size(const void *ptr, const size_t bufsize, void **aligned_ptr, size_t *aligned_size) { + if (!aligned_ptr || !aligned_size) return; + const size_t page_size = get_sc_page_size(); + uintptr_t aligned_ptr_local = (uintptr_t)ptr & ~(page_size - 1); + size_t local_offset = (size_t)((uintptr_t)ptr - aligned_ptr_local); + *aligned_size = (bufsize + local_offset + page_size - 1) & ~(page_size - 1); + *aligned_ptr = (void *)aligned_ptr_local; +} diff --git a/src/transport/net_ib.cc b/src/transport/net_ib.cc index 8b288bca77..d1fc9a9604 100644 --- a/src/transport/net_ib.cc +++ b/src/transport/net_ib.cc @@ -793,7 +793,8 @@ static void ibDmaBufSupportInitOnce(){ ncclIbDev* ibDev = ncclIbDevs + mergedDev->vProps.devs[0]; struct ibv_pd* pd; struct ibv_context* ctx = ibDev->context; - rocmLibraryInit(); + res = rocmLibraryInit(); + if (res != ncclSuccess) goto failure; NCCLCHECKGOTO(wrap_ibv_alloc_pd(&pd, ctx), res, failure); // Test kernel DMA-BUF support with a dummy call (fd=-1) (void)wrap_direct_ibv_reg_dmabuf_mr(pd, 0ULL /*offset*/, 0ULL /*len*/, 0ULL /*iova*/, -1 /*fd*/, 0 /*flags*/); @@ -1071,6 +1072,7 @@ struct ncclIbGpuFlush { int* gpuFlushGpuMem; struct ibv_sge sge; struct ncclIbQp qp; + int dmabuf_fd; }; struct ncclIbRemFifo { @@ -1620,6 +1622,7 @@ ib_recv: struct ncclIbRecvCommDev* rCommDev; struct ncclIbDevInfo* remDevInfo; struct ncclIbQp* qp; + bool useDmaBuf; mergedDev = ncclIbMergedDevs + lComm->dev; rComm->base.nRemDevs = remMeta.ndevs; @@ -1685,9 +1688,9 @@ ib_recv: NCCLCHECKGOTO(ncclIbRtsQp(qp->qp), ret, fail); } - rComm->flushEnabled = ((ncclIbGdrSupport() == ncclSuccess || ncclIbDmaBufSupport(lComm->dev) == ncclSuccess) - && (ncclParamIbGdrFlushDisable() == 0)) ? 1 : 0; - + useDmaBuf = (ncclIbDmaBufSupport(lComm->dev) == ncclSuccess); + rComm->flushEnabled = ((ncclIbGdrSupport() == ncclSuccess || useDmaBuf) + && (ncclParamIbGdrFlushDisable() == 0)) ? 1 : 0; for (int i = 0; i < rComm->base.vProps.ndevs; i++) { rCommDev = rComm->devs + i; ibDev = ncclIbDevs + rCommDev->base.ibDevN; @@ -1706,10 +1709,29 @@ ib_recv: #else NCCLCHECKGOTO(ncclCudaCalloc(&rCommDev->gpuFlush.gpuFlushGpuMem, sizeof(int), nullptr, hipDeviceMallocFinegrained), ret, fail); #endif - NCCLCHECKGOTO(wrap_ibv_reg_mr(&rCommDev->gpuFlush.gpuMr, rCommDev->base.pd, rCommDev->gpuFlush.gpuFlushGpuMem, sizeof(int), IBV_ACCESS_LOCAL_WRITE|IBV_ACCESS_REMOTE_WRITE|IBV_ACCESS_REMOTE_READ), ret, fail); + if (useDmaBuf) + { + uint64_t export_offset = 0; + void *aligned_ptr = NULL; + size_t aligned_size = 0; + get_aligned_ptr_and_size(rCommDev->gpuFlush.gpuFlushGpuMem, sizeof(int) /*devicebuffersize*/, &aligned_ptr, &aligned_size); + hsa_status_t export_status = pfn_hsa_amd_portable_export_dmabuf(aligned_ptr, aligned_size, &rCommDev->gpuFlush.dmabuf_fd, &export_offset); + if (rCommDev->gpuFlush.dmabuf_fd < 0 || export_status != HSA_STATUS_SUCCESS) + { + WARN("Failed to export DMA BUF"); + goto fail; + } + NCCLCHECKGOTO(wrap_ibv_reg_dmabuf_mr(&rCommDev->gpuFlush.gpuMr, rCommDev->base.pd, export_offset, sizeof(int), (uint64_t)rCommDev->gpuFlush.gpuFlushGpuMem /*iova*/, rCommDev->gpuFlush.dmabuf_fd, IBV_ACCESS_LOCAL_WRITE | IBV_ACCESS_REMOTE_WRITE | IBV_ACCESS_REMOTE_READ), ret, fail); + } + else + { + rCommDev->gpuFlush.dmabuf_fd = -1; + NCCLCHECKGOTO(wrap_ibv_reg_mr(&rCommDev->gpuFlush.gpuMr, rCommDev->base.pd, rCommDev->gpuFlush.gpuFlushGpuMem, sizeof(int), IBV_ACCESS_LOCAL_WRITE | IBV_ACCESS_REMOTE_WRITE | IBV_ACCESS_REMOTE_READ), ret, fail); + } } else { rCommDev->gpuFlush.gpuFlushGpuMem = nullptr; rCommDev->gpuFlush.gpuMr = nullptr; + rCommDev->gpuFlush.dmabuf_fd = -1; } NCCLCHECKGOTO(wrap_ibv_reg_mr(&rCommDev->gpuFlush.hostMr, rCommDev->base.pd, &rComm->gpuFlushHostMem, sizeof(int), IBV_ACCESS_LOCAL_WRITE), ret, fail); rCommDev->gpuFlush.sge.addr = (uint64_t)&rComm->gpuFlushHostMem; @@ -2446,6 +2468,7 @@ ncclResult_t ncclIbCloseRecv(void* recvComm) { commDev->gpuFlush.gpuFlushGpuMem = nullptr; if (commDev->gpuFlush.gpuMr != nullptr) NCCLCHECK(wrap_ibv_dereg_mr(commDev->gpuFlush.gpuMr)); commDev->gpuFlush.gpuMr = nullptr; + if(commDev->gpuFlush.dmabuf_fd > 0) { close(commDev->gpuFlush.dmabuf_fd);} } if (commDev->gpuFlush.qp.qp != NULL) NCCLCHECK(wrap_ibv_destroy_qp(commDev->gpuFlush.qp.qp)); if (commDev->gpuFlush.hostMr != NULL) NCCLCHECK(wrap_ibv_dereg_mr(commDev->gpuFlush.hostMr));