RCCL Multinode DMA Buffer crash fix (#1682)
This commit handles DMABUF initialization and call appropriate handling function. This fixes crash in OS with no peermem support and relying on only DMABUF. * Initial test commit * Handling Dmabuf_fd opening and closing * Cleanup * Use DMABuff or Peermem as needed * Using user input for ibDmaBufSupportInitOnce * Revert all changes to rocmwrap.cc * Revert all changes to rocmwrap.cc * Changing to func definition braces * Reverting line removal in utils.h * useDmaBuf to calculate flushEnabled
This commit is contained in:
Avinash
GitHub
@@ -545,4 +545,25 @@ T* ncclIntruQueueMpscAbandon(ncclIntruQueueMpsc<T,next>* 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
|
||||
|
||||
@@ -8,7 +8,7 @@
|
||||
#include "core.h"
|
||||
|
||||
#include "nvmlwrap.h"
|
||||
|
||||
#include <unistd.h>
|
||||
#include <stdlib.h>
|
||||
|
||||
// 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;
|
||||
}
|
||||
|
||||
@@ -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));
|
||||
|
||||
Reference in New Issue
Block a user