From 03fba66e719cf5a2ba70bf581af4c568411d37ba Mon Sep 17 00:00:00 2001 From: corey-derochie-amd <161367113+corey-derochie-amd@users.noreply.github.com> Date: Wed, 11 Jun 2025 17:52:09 -0600 Subject: [PATCH] Deprecated MSCCL API functions (#1740) --- src/include/api_trace.h | 2 +- src/misc/api_trace.cc | 8 ++--- src/misc/msccl/msccl_lifecycle.cc | 57 ++++++++++++++++++++++++++++++- src/msccl.cc | 53 ++++------------------------ src/nccl.h.in | 10 +++--- 5 files changed, 74 insertions(+), 56 deletions(-) diff --git a/src/include/api_trace.h b/src/include/api_trace.h index a33d7e8d6a..8329718671 100644 --- a/src/include/api_trace.h +++ b/src/include/api_trace.h @@ -139,7 +139,7 @@ typedef ncclResult_t (*ncclMemAlloc_fn_t)(void** ptr, size_t size); typedef ncclResult_t (*ncclMemFree_fn_t)(void* ptr); typedef ncclResult_t (*mscclLoadAlgo_fn_t)(const char* mscclAlgoFilePath, - mscclAlgoHandle_t* mscclAlgoHandle, const ncclComm_t comm); + mscclAlgoHandle_t* mscclAlgoHandle, int rank); typedef ncclResult_t (*mscclRunAlgo_fn_t)( const void* sendBuff, const size_t sendCounts[], const size_t sDisPls[], diff --git a/src/misc/api_trace.cc b/src/misc/api_trace.cc index 5ae708c0c1..92e09eb149 100644 --- a/src/misc/api_trace.cc +++ b/src/misc/api_trace.cc @@ -135,7 +135,7 @@ ncclMemFree_impl(void* ptr); ncclResult_t mscclLoadAlgo_impl(const char* mscclAlgoFilePath, mscclAlgoHandle_t* mscclAlgoHandle, - const ncclComm_t comm); + int rank); ncclResult_t mscclRunAlgo_impl(const void* sendBuff, const size_t sendCounts[], const size_t sDisPls[], @@ -380,7 +380,7 @@ NCCL_API(ncclResult_t, ncclMemAlloc, void** ptr, size_t size); NCCL_API(ncclResult_t, ncclMemFree, void* ptr); NCCL_API(ncclResult_t, mscclLoadAlgo, const char* mscclAlgoFilePath, - mscclAlgoHandle_t* mscclAlgoHandle, const ncclComm_t comm); + mscclAlgoHandle_t* mscclAlgoHandle, int rank); NCCL_API(ncclResult_t, mscclRunAlgo, const void* sendBuff, const size_t sendCounts[], const size_t sDisPls[], void* recvBuff, const size_t recvCounts[], @@ -620,10 +620,10 @@ ncclMemFree(void* ptr) } ncclResult_t -mscclLoadAlgo(const char* mscclAlgoFilePath, mscclAlgoHandle_t* mscclAlgoHandle, const ncclComm_t comm) +mscclLoadAlgo(const char* mscclAlgoFilePath, mscclAlgoHandle_t* mscclAlgoHandle, int rank) { return ::rccl::RcclGetFunctionTable()->mscclLoadAlgo_fn(mscclAlgoFilePath, - mscclAlgoHandle, comm); + mscclAlgoHandle, rank); } ncclResult_t diff --git a/src/misc/msccl/msccl_lifecycle.cc b/src/misc/msccl/msccl_lifecycle.cc index 15062f6acf..38d2661b42 100644 --- a/src/misc/msccl/msccl_lifecycle.cc +++ b/src/misc/msccl/msccl_lifecycle.cc @@ -244,6 +244,29 @@ ncclResult_t mscclSchedulerInit(ncclComm_t comm, int* numChannelsRequired) { return ncclSuccess; } +static ncclResult_t mscclLoadAlgo(const char *mscclAlgoFilePath, mscclAlgoHandle_t *mscclAlgoHandle, const ncclComm_t comm) { + mscclStatus& status = mscclGetStatus(comm); + + if (status.freeAlgoHandles.size() == 0) { + WARN("MSCCL: MSCCL_MAX_NUM_ALGOS (%d) limit reached", MSCCL_MAX_NUM_ALGOS); + return ncclInvalidUsage; + } + *mscclAlgoHandle = *status.freeAlgoHandles.rbegin(); + status.freeAlgoHandles.pop_back(); + + struct mscclAlgo* hostAlgo; + NCCLCHECK(ncclCalloc(&hostAlgo, 1)); + NCCLCHECK(mscclGetAlgoFromXmlFile(mscclAlgoFilePath, hostAlgo, comm->rank)); + status.hostAlgos[*mscclAlgoHandle] = hostAlgo; + + struct mscclAlgo* devAlgo; + NCCLCHECK(ncclCudaMalloc(&devAlgo, 1)); + CUDACHECK(hipMemcpy(devAlgo, hostAlgo, sizeof(struct mscclAlgo), hipMemcpyHostToDevice)); + status.devAlgos[*mscclAlgoHandle] = devAlgo; + + return ncclSuccess; +} + ncclResult_t mscclInit(const ncclComm_t comm) { { mscclStatus& status = mscclGetStatus(comm); @@ -413,6 +436,38 @@ static ncclResult_t mscclSaveCountsAndDispls(struct mscclSavedSchedulerParam* pa return ncclSuccess; } +static ncclResult_t _mscclRunAlgo( + const void* sendBuff, const size_t sendCounts[], const size_t sDisPls[], + void* recvBuff, const size_t recvCounts[], const size_t rDisPls[], + size_t count, ncclDataType_t dataType, int root, int peer, ncclRedOp_t op, + mscclAlgoHandle_t mscclAlgoHandle, ncclComm_t comm, hipStream_t stream) { + mscclStatus& status = mscclGetStatus(comm); + struct mscclAlgo* hostAlgo = status.hostAlgos[mscclAlgoHandle]; + struct mscclAlgo* devAlgo = status.devAlgos[mscclAlgoHandle]; + + // NCCL adds a lot of guarantees that target device is getting used + // in its group management code, which we entirely skip when MSCCL is used + // Therefore, in single thread multiGPU mode + // setting the device is critical to be sure + // communication is done on the intended device + + CUDACHECK(hipSetDevice(comm->cudaDev)); + + NCCLCHECK(mscclGetCaptureStatus(comm, stream)); + + NCCLCHECK(mscclSetupCount(hostAlgo, comm, count, dataType)); + + NCCLCHECK(mscclSetupScratch(hostAlgo, stream)); + + NCCLCHECK(mscclSetupSyncFlags(comm, stream)); + + NCCLCHECK(mscclSetupProxy(hostAlgo, comm, stream)); + + NCCLCHECK(mscclSetupKernel(sendBuff, recvBuff, count, dataType, op, hostAlgo, devAlgo, comm, stream)); + + return ncclSuccess; +} + static ncclResult_t mscclRunSavedParams() { mscclThreadLocalStatus& threadLocalStatus = mscclGetThreadLocalStatus(); for (auto& param : threadLocalStatus.savedSchedulerParams) { @@ -420,7 +475,7 @@ static ncclResult_t mscclRunSavedParams() { mscclFuncNames[param.p.func], param.p.opCount, param.p.sendBuff, param.p.recvBuff, param.p.count, param.p.dataType, param.p.op, param.p.root, param.comm, param.p.nRanks, param.stream, param.comm->planner.nTasksP2p + param.comm->planner.nTasksColl, param.comm->localRankToRank[param.comm->localRank]); - NCCLCHECK(mscclRunAlgo( + NCCLCHECK(_mscclRunAlgo( param.p.sendBuff, param.p.sendCounts, param.p.sDisPls, param.p.recvBuff, param.p.recvCounts, param.p.rDisPls, param.p.count, param.p.dataType, param.p.root, param.p.peer, param.p.op, param.p.handle, param.comm, param.stream)); diff --git a/src/msccl.cc b/src/msccl.cc index 706c971808..51a9dd6470 100644 --- a/src/msccl.cc +++ b/src/msccl.cc @@ -14,28 +14,11 @@ using namespace rccl; -NCCL_API(ncclResult_t, mscclLoadAlgo, const char *mscclAlgoFilePath, mscclAlgoHandle_t *mscclAlgoHandle, const ncclComm_t comm); -ncclResult_t mscclLoadAlgo_impl(const char *mscclAlgoFilePath, mscclAlgoHandle_t *mscclAlgoHandle, const ncclComm_t comm) { +NCCL_API(ncclResult_t, mscclLoadAlgo, const char *mscclAlgoFilePath, mscclAlgoHandle_t *mscclAlgoHandle, int rank); +ncclResult_t mscclLoadAlgo_impl(const char *mscclAlgoFilePath, mscclAlgoHandle_t *mscclAlgoHandle, int rank) { + // deprecated Recorder::instance().record("mscclLoadAlgo"); - mscclStatus& status = mscclGetStatus(comm); - - if (status.freeAlgoHandles.size() == 0) { - WARN("MSCCL: MSCCL_MAX_NUM_ALGOS (%d) limit reached", MSCCL_MAX_NUM_ALGOS); - return ncclInvalidUsage; - } - *mscclAlgoHandle = *status.freeAlgoHandles.rbegin(); - status.freeAlgoHandles.pop_back(); - - struct mscclAlgo* hostAlgo; - NCCLCHECK(ncclCalloc(&hostAlgo, 1)); - NCCLCHECK(mscclGetAlgoFromXmlFile(mscclAlgoFilePath, hostAlgo, comm->rank)); - status.hostAlgos[*mscclAlgoHandle] = hostAlgo; - - struct mscclAlgo* devAlgo; - NCCLCHECK(ncclCudaMalloc(&devAlgo, 1)); - CUDACHECK(hipMemcpy(devAlgo, hostAlgo, sizeof(struct mscclAlgo), hipMemcpyHostToDevice)); - status.devAlgos[*mscclAlgoHandle] = devAlgo; - + WARN("mscclLoadAlgo is deprecated. Function call has no effect."); return ncclSuccess; } @@ -49,34 +32,11 @@ ncclResult_t mscclRunAlgo_impl( void* recvBuff, const size_t recvCounts[], const size_t rDisPls[], size_t count, ncclDataType_t dataType, int root, int peer, ncclRedOp_t op, mscclAlgoHandle_t mscclAlgoHandle, ncclComm_t comm, hipStream_t stream) { + // deprecated Recorder::instance().record("mscclRunAlgo"); NVTX3_FUNC_WITH_PARAMS(MSCCL, NcclNvtxParamsMSCCL, NVTX3_PAYLOAD(comm ? comm->commHash : 0, count * ncclTypeSize(dataType), op, dataType)); - - mscclStatus& status = mscclGetStatus(comm); - struct mscclAlgo* hostAlgo = status.hostAlgos[mscclAlgoHandle]; - struct mscclAlgo* devAlgo = status.devAlgos[mscclAlgoHandle]; - - // NCCL adds a lot of guarantees that target device is getting used - // in its group management code, which we entirely skip when MSCCL is used - // Therefore, in single thread multiGPU mode - // setting the device is critical to be sure - // communication is done on the intended device - - CUDACHECK(hipSetDevice(comm->cudaDev)); - - NCCLCHECK(mscclGetCaptureStatus(comm, stream)); - - NCCLCHECK(mscclSetupCount(hostAlgo, comm, count, dataType)); - - NCCLCHECK(mscclSetupScratch(hostAlgo, stream)); - - NCCLCHECK(mscclSetupSyncFlags(comm, stream)); - - NCCLCHECK(mscclSetupProxy(hostAlgo, comm, stream)); - - NCCLCHECK(mscclSetupKernel(sendBuff, recvBuff, count, dataType, op, hostAlgo, devAlgo, comm, stream)); - + WARN("mscclRunAlgo is deprecated. Function call has no effect."); return ncclSuccess; } @@ -84,5 +44,6 @@ NCCL_API(ncclResult_t, mscclUnloadAlgo, mscclAlgoHandle_t mscclAlgoHandle); ncclResult_t mscclUnloadAlgo_impl(mscclAlgoHandle_t mscclAlgoHandle) { // deprecated Recorder::instance().record("mscclUnloadAlgo"); + WARN("mscclUnloadAlgo is deprecated. Function call has no effect."); return ncclSuccess; } diff --git a/src/nccl.h.in b/src/nccl.h.in index 3700667437..87581aade3 100644 --- a/src/nccl.h.in +++ b/src/nccl.h.in @@ -756,6 +756,7 @@ ncclResult_t pncclAllToAllv(const void *sendbuff, const size_t sendcounts[], typedef int mscclAlgoHandle_t; /*! @brief MSCCL Load Algorithm + @deprecated This function has been removed from the public API. @details Load MSCCL algorithm file specified in mscclAlgoFilePath and return its handle via mscclAlgoHandle. This API is expected to be called by MSCCL scheduler instead of end users. @@ -763,13 +764,14 @@ typedef int mscclAlgoHandle_t; @param[in] mscclAlgoFilePath Path to MSCCL algorithm file @param[out] mscclAlgoHandle Returned handle to MSCCL algorithm - @param[in] comm Current rank's communicator */ -ncclResult_t mscclLoadAlgo(const char *mscclAlgoFilePath, mscclAlgoHandle_t *mscclAlgoHandle, const ncclComm_t comm); + @param[in] rank Current rank */ +ncclResult_t mscclLoadAlgo(const char *mscclAlgoFilePath, mscclAlgoHandle_t *mscclAlgoHandle, int rank); /*! @cond include_hidden */ -ncclResult_t pmscclLoadAlgo(const char *mscclAlgoFilePath, mscclAlgoHandle_t *mscclAlgoHandle, const ncclComm_t comm); +ncclResult_t pmscclLoadAlgo(const char *mscclAlgoFilePath, mscclAlgoHandle_t *mscclAlgoHandle, int rank); /*! @endcond */ /*! @brief MSCCL Run Algorithm + @deprecated This function has been removed from the public API. @details Run MSCCL algorithm specified by mscclAlgoHandle. The parameter list merges all possible parameters required by different operations as this is a general-purposed API. This API is expected to be called by MSCCL @@ -871,4 +873,4 @@ ncclResult_t pncclGroupSimulateEnd(ncclSimInfo_t* simInfo); } // end extern "C" #endif -#endif // end include guard \ No newline at end of file +#endif // end include guard