AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

2.16.2-1

Browse Source 
Add support for CUDA 12.0, drop Kepler (sm_35).
Support for H100 features.
Make socket code more robust and protected. Solves #555.
Improve performance on large CUDA graphs, reducing dependencies.
Reduce inter-socket bandwidth on AMD CPUs to favor better paths.
Various fixes to ncclCommAbort.
Make service thread polling resistant to EINTR.
Compile with profiling API by default.
Extend NVTX instrumentation with call arguments.


[ROCm/rccl commit: 28189e2df8]
This commit is contained in:
Sylvain Jeaugey
2022-11-29 04:27:46 -08:00
parent 1bcb7a7cb6
commit 2ce8946622
46 changed files with 3325 additions and 1037 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rccl/src/misc/shmutils.cc
+133 -60
View File
@@ -15,79 +15,152 @@
#include <stdlib.h>
#include <unistd.h>
// Change functions behavior to match other SYS functions
static int shm_allocate(int fd, const int shmSize) {
int err = posix_fallocate(fd, 0, shmSize);
if (err) { errno = err; return -1; }
return 0;
}
static int shm_map(int fd, const int shmSize, void** ptr) {
*ptr = mmap(NULL, shmSize, PROT_READ | PROT_WRITE, MAP_SHARED, fd, 0);
return (*ptr == MAP_FAILED) ? -1 : 0;
struct shmHandleInternal {
int fd;
char* shmPath;
char* shmPtr;
void* devShmPtr;
size_t shmSize;
size_t realShmSize;
int* refcount;
};
static void shmHandleInit(int fd, char* shmPath, size_t shmSize, size_t realShmSize, char* hptr, void* dptr, bool create, struct shmHandleInternal* handle) {
handle->fd = fd;
handle->shmPtr = hptr;
handle->devShmPtr = dptr;
handle->shmSize = shmSize;
handle->realShmSize = realShmSize;
handle->refcount = (int*)(hptr + shmSize);
if (create) {
int slen = strlen(shmPath);
handle->shmPath = (char*)malloc(slen + 1);
memcpy(handle->shmPath, shmPath, slen + 1);
if (hptr) memset(hptr, 0, shmSize);
} else {
handle->shmPath = NULL;
}
return;
}
static ncclResult_t ncclShmSetup(char* shmPath, const int shmSize, int* fd, void** ptr, int create) {
ncclResult_t ncclShmOpen(char* shmPath, size_t shmSize, void** shmPtr, void** devShmPtr, int refcount, ncclShmHandle_t* handle) {
int fd = -1;
char* hptr = NULL;
void* dptr = NULL;
ncclResult_t ret = ncclSuccess;
struct shmHandleInternal* tmphandle;
bool create = refcount > 0 ? true : false;
const size_t refSize = sizeof(int); /* extra sizeof(int) bytes for reference count */
const size_t realShmSize = shmSize + refSize;
*handle = *shmPtr = NULL; /* assume shmPtr and handle always set correctly by users. */
EQCHECKGOTO(tmphandle = (struct shmHandleInternal*)malloc(sizeof(struct shmHandleInternal)), NULL, ret, fail);
if (create) {
/* refcount > 0 means the caller tries to allocate a shared memory. This shared memory segment will have
* refcount references; when the peer attaches, it should pass -1 to reduce one reference count. When it
* goes down to 0, unlink should be called in order to delete shared memory file. */
if (shmPath[0] == '\0') {
sprintf(shmPath, "/dev/shm/nccl-XXXXXX");
*fd = mkstemp(shmPath);
fd = mkstemp(shmPath);
} else {
SYSCHECKVAL(open(shmPath, O_CREAT | O_RDWR, S_IRUSR | S_IWUSR), "open", *fd);
SYSCHECKGOTO(fd = open(shmPath, O_CREAT | O_RDWR, S_IRUSR | S_IWUSR), ret, fail);
}
if (ftruncate(*fd, shmSize) != 0) {
WARN("Error: failed to extend %s to %d bytes", shmPath, shmSize);
return ncclSystemError;
if (ftruncate(fd, realShmSize) != 0) {
WARN("Error: failed to extend %s to %ld bytes", shmPath, realShmSize);
ret = ncclSystemError;
goto fail;
}
INFO(NCCL_ALLOC, "Allocated %d bytes of shared memory in %s\n", shmSize, shmPath);
INFO(NCCL_ALLOC, "Allocated %ld bytes of shared memory in %s", realShmSize, shmPath);
} else {
SYSCHECKVAL(open(shmPath, O_RDWR, S_IRUSR | S_IWUSR), "open", *fd);
}
*ptr = (char*)mmap(NULL, shmSize, PROT_READ|PROT_WRITE, MAP_SHARED, *fd, 0);
if (*ptr == NULL) {
WARN("Could not map %s\n", shmPath);
return ncclSystemError;
}
close(*fd);
*fd = -1;
if (create) memset(*ptr, 0, shmSize);
return ncclSuccess;
}
ncclResult_t ncclShmOpen(char* shmPath, const int shmSize, void** shmPtr, void** devShmPtr, int create) {
int fd = -1;
void* ptr = MAP_FAILED;
ncclResult_t res = ncclSuccess;
NCCLCHECKGOTO(ncclShmSetup(shmPath, shmSize, &fd, &ptr, create), res, sysError);
if (devShmPtr) {
CUDACHECKGOTO(cudaHostRegister(ptr, shmSize, cudaHostRegisterMapped), res, cudaError);
CUDACHECKGOTO(cudaHostGetDevicePointer(devShmPtr, ptr, 0), res, cudaError);
SYSCHECKGOTO(fd = open(shmPath, O_RDWR, S_IRUSR | S_IWUSR), ret, fail);
}
*shmPtr = ptr;
return ncclSuccess;
sysError:
WARN("Error while %s shared memory segment %s (size %d)", create ? "creating" : "attaching to", shmPath, shmSize);
cudaError:
if (fd != -1) close(fd);
if (create) shm_unlink(shmPath);
if (ptr != MAP_FAILED) munmap(ptr, shmSize);
*shmPtr = NULL;
return res;
}
hptr = (char*)mmap(NULL, realShmSize, PROT_READ | PROT_WRITE, MAP_SHARED, fd, 0);
if (hptr == MAP_FAILED) {
WARN("Could not map %s size %zi, error: %s", shmPath, realShmSize, strerror(errno));
ret = ncclSystemError;
goto fail;
}
ncclResult_t ncclShmUnlink(const char* shmPath) {
if (shmPath != NULL) SYSCHECK(unlink(shmPath), "unlink");
return ncclSuccess;
}
if (create) {
*(int*)(hptr + shmSize) = refcount;
} else {
int remref = __atomic_sub_fetch((int*)(hptr + shmSize), 1, __ATOMIC_RELAXED);
if (remref == 0) {
/* the last peer has completed attachment, it should unlink the shm mem file. */
if (unlink(shmPath) != 0) {
WARN("unlink shared memory %s failed, error: %s", shmPath, strerror(errno));
}
}
ncclResult_t ncclShmClose(void* shmPtr, void* devShmPtr, const int shmSize) {
if (shmPtr) {
if (devShmPtr) CUDACHECK(cudaHostUnregister(shmPtr));
if (munmap(shmPtr, shmSize) != 0) {
WARN("munmap of shared memory failed");
return ncclSystemError;
if (refcount != -1) {
WARN("attaching memory should only reduce refcount by 1 but %d is passed", refcount);
}
}
return ncclSuccess;
if (devShmPtr) {
CUDACHECKGOTO(cudaHostRegister((void*)hptr, realShmSize, cudaHostRegisterMapped), ret, fail);
CUDACHECKGOTO(cudaHostGetDevicePointer(&dptr, (void*)hptr, 0), ret, fail);
}
shmHandleInit(fd, shmPath, shmSize, realShmSize, hptr, dptr, create, tmphandle);
exit:
*shmPtr = hptr;
if (devShmPtr) *devShmPtr = dptr;
*handle = (ncclShmHandle_t)tmphandle;
return ret;
fail:
WARN("Error while %s shared memory segment %s (size %ld)", create ? "creating" : "attaching to", shmPath, shmSize);
if (tmphandle) {
shmHandleInit(fd, shmPath, shmSize, realShmSize, hptr, dptr, create, tmphandle);
ncclShmClose((ncclShmHandle_t)tmphandle);
tmphandle = NULL;
}
hptr = NULL;
dptr = NULL;
goto exit;
}
ncclResult_t ncclShmClose(ncclShmHandle_t handle) {
ncclResult_t ret = ncclSuccess;
struct shmHandleInternal* tmphandle = (struct shmHandleInternal*)handle;
if (tmphandle) {
if (tmphandle->fd >= 0) {
close(tmphandle->fd);
if (tmphandle->shmPath != NULL && *tmphandle->refcount > 0) {
if (unlink(tmphandle->shmPath) != 0) {
WARN("unlink shared memory %s failed, error: %s", tmphandle->shmPath, strerror(errno));
ret = ncclSystemError;
}
free(tmphandle->shmPath);
}
}
if (tmphandle->shmPtr) {
if (tmphandle->devShmPtr) CUDACHECK(cudaHostUnregister(tmphandle->shmPtr));
if (munmap(tmphandle->shmPtr, tmphandle->realShmSize) != 0) {
WARN("munmap of shared memory %p size %ld failed, error: %s", tmphandle->shmPtr, tmphandle->realShmSize, strerror(errno));
ret = ncclSystemError;
}
}
free(tmphandle);
}
return ret;
}
ncclResult_t ncclShmUnlink(ncclShmHandle_t handle) {
ncclResult_t ret = ncclSuccess;
struct shmHandleInternal* tmphandle = (struct shmHandleInternal*)handle;
if (tmphandle) {
if (tmphandle->shmPath != NULL) {
if (unlink(tmphandle->shmPath) != 0) {
WARN("unlink shared memory %s failed, error: %s", tmphandle->shmPath, strerror(errno));
ret = ncclSystemError;
}
free(tmphandle->shmPath);
tmphandle->shmPath = NULL;
}
}
return ret;
}