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

Only enable MSCCL++ AllReduce for message sizes that are multiples 32 (#1253)

Browse Source 
* Only enable MSCCL++ AllReduce for message sizes that are multiples of 32. MSCCL++ does not handle these other sizes.

* Sanitized MSCCL++ logging.
This commit is contained in:
corey-derochie-amd
2024-07-12 18:04:23 -06:00
committed by GitHub
parent 6dc47eecd7
commit 9cbb3da224
2 changed files with 16 additions and 12 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/init.cc
+4 -4
View File
@@ -200,8 +200,8 @@ ncclResult_t ncclGetUniqueId(ncclUniqueId* out) {
hipDeviceProp_t devProp;
CUDACHECK(hipGetDeviceProperties(&devProp, dev));
if (IsArchMatch(devProp.gcnArchName, "gfx94")) {
INFO(NCCL_INIT, "MSCCL++: mscclpp_ncclGetUniqueId");
res = mscclpp_ncclGetUniqueId(&(mscclpp_uniqueIdMap[*out]));
TRACE_CALL("mscclpp_ncclGetUniqueId");
} else {
WARN("MSCCL++: Cannot enable MSCCL++ on %s architecture", devProp.gcnArchName);
}
@@ -1975,11 +1975,11 @@ static ncclResult_t ncclCommInitRankFunc(struct ncclAsyncJob* job_) {
comm->mscclppCompatible = IsArchMatch(devProp.gcnArchName, "gfx94");
if (comm->mscclppCompatible) {
NCCLCHECKGOTO(bootstrapIntraNodeBroadcast(comm->bootstrap, comm->localRankToRank, comm->localRank, comm->localRanks, 0, &(mscclpp_uniqueIdMap[job->commId]), sizeof(mscclpp_ncclUniqueId)), res, fail);
INFO(NCCL_INIT, "MSCCL++: Broadcast mscclpp_ncclUniqueId to %d ranks", (comm->localRanks - 1));
TRACE_CALL("bootstrapIntraNodeBroadcast(rank=%d, nranks=%d, root=%d, bcastData=<mscclpp_ncclUniqueId>)", comm->localRank, comm->localRanks, 0);
comm->mscclpp_threshold = rcclParamMscclppThreshold();
INFO(NCCL_INIT, "MSCCL++: Enabled! Msg size threshold=%zu", comm->mscclpp_threshold);
INFO(NCCL_INIT, "MSCCL++: mscclpp_ncclCommInitRank (nranks=%d)", job->nranks);
NCCLCHECKGOTO(mscclpp_ncclCommInitRank(&(comm->mscclpp_comm), job->nranks, mscclpp_uniqueIdMap[job->commId], job->myrank), res, fail);
TRACE_CALL("mscclpp_ncclCommInitRank (nranks=%d, myrank=%d)", job->nranks, job->myrank);
} else {
WARN("MSCCL++: Cannot enable MSCCL++ on %s architecture", devProp.gcnArchName);
}
@@ -2586,8 +2586,8 @@ ncclResult_t ncclCommDestroy(ncclComm_t comm) {
#ifdef ENABLE_MSCCLPP
if (comm->mscclppCompatible) {
INFO(NCCL_INIT, "MSCCL++: mscclpp_ncclCommDestroy");
ncclResult_t res = mscclpp_ncclCommDestroy(comm->mscclpp_comm);
TRACE_CALL("mscclpp_ncclCommDestroy");
if (res != ncclSuccess) {
WARN("MSCCL++: mscclpp_ncclCommDestroy failed (%s)", ncclGetErrorString(res));
}
src/misc/msccl/msccl_lifecycle.cc
+12 -8
View File
@@ -459,15 +459,17 @@ ncclResult_t mscclEnqueueCheck(
#ifdef ENABLE_MSCCLPP
if (comm->mscclppCompatible) {
/* check if one rank per GPU and graph mode is enabled */
if ((nBytes >= 32) && (threadLocalStatus.captureStatus != mscclNoCapture) && comm->mscclCompatible) {
if (func == mscclFuncAllReduce && nBytes <= comm->mscclpp_threshold) {
INFO(NCCL_INIT, "MSCCL++: mscclpp_ncclAllReduce (groupStatus=mscclNoGroup)");
if ((threadLocalStatus.captureStatus != mscclNoCapture) && comm->mscclCompatible) {
if (func == mscclFuncAllReduce && nBytes <= comm->mscclpp_threshold && (nBytes & 31) == 0) {
INFO(NCCL_COLL,"%s: opCount %lx sendbuff %p recvbuff %p count %zi datatype %d op %d root %d comm %p [nranks=%d] stream %p",
"mscclpp_ncclAllReduce", comm->opCount, sendBuff, recvBuff, count, dataType, op, root, comm, comm->nRanks, stream);
NCCLCHECK(mscclpp_ncclAllReduce(sendBuff, recvBuff, count, dataType, op, comm->mscclpp_comm, stream));
threadLocalStatus.savedSchedulerParams.clear();
break;
}
else if (func == mscclFuncAllGather && nBytes * comm->nRanks <= comm->mscclpp_threshold) {
INFO(NCCL_INIT, "MSCCL++: mscclpp_ncclAllGather (groupStatus=mscclNoGroup)");
INFO(NCCL_COLL,"%s: opCount %lx sendbuff %p recvbuff %p count %zi datatype %d op %d root %d comm %p [nranks=%d] stream %p",
"mscclpp_ncclAllGather", comm->opCount, sendBuff, recvBuff, count, dataType, op, root, comm, comm->nRanks, stream);
NCCLCHECK(mscclpp_ncclAllGather(sendBuff, recvBuff, count, dataType, comm->mscclpp_comm, stream));
threadLocalStatus.savedSchedulerParams.clear();
break;
@@ -488,15 +490,17 @@ ncclResult_t mscclEnqueueCheck(
#ifdef ENABLE_MSCCLPP
if (comm->mscclppCompatible) {
/* check if one rank per GPU and graph mode is enabled */
if ((nBytes >= 32) && (threadLocalStatus.captureStatus != mscclNoCapture) && comm->mscclCompatible) {
if (func == mscclFuncAllReduce && nBytes <= comm->mscclpp_threshold) {
INFO(NCCL_INIT, "MSCCL++: mscclpp_ncclAllReduce (groupStatus=mscclGroupSupportedOp)");
if ((threadLocalStatus.captureStatus != mscclNoCapture) && comm->mscclCompatible) {
if (func == mscclFuncAllReduce && nBytes <= comm->mscclpp_threshold && (nBytes & 31) == 0) {
INFO(NCCL_COLL,"%s: opCount %lx sendbuff %p recvbuff %p count %zi datatype %d op %d root %d comm %p [nranks=%d] stream %p",
"mscclpp_ncclAllReduce", comm->opCount, sendBuff, recvBuff, count, dataType, op, root, comm, comm->nRanks, stream);
NCCLCHECK(mscclpp_ncclAllReduce(sendBuff, recvBuff, count, dataType, op, comm->mscclpp_comm, stream));
threadLocalStatus.savedSchedulerParams.clear();
break;
}
else if (func == mscclFuncAllGather && nBytes * comm->nRanks <= comm->mscclpp_threshold) {
INFO(NCCL_INIT, "MSCCL++: mscclpp_ncclAllGather (groupStatus=mscclGroupSupportedOp)");
INFO(NCCL_COLL,"%s: opCount %lx sendbuff %p recvbuff %p count %zi datatype %d op %d root %d comm %p [nranks=%d] stream %p",
"mscclpp_ncclAllGather", comm->opCount, sendBuff, recvBuff, count, dataType, op, root, comm, comm->nRanks, stream);
NCCLCHECK(mscclpp_ncclAllGather(sendBuff, recvBuff, count, dataType, comm->mscclpp_comm, stream));
threadLocalStatus.savedSchedulerParams.clear();
break;