AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

Deprecated MSCCL API functions (#1740)

Просмотр исходного кода
Этот коммит содержится в:
corey-derochie-amd
2025-06-11 17:52:09 -06:00
коммит произвёл GitHub
родитель 75c3c8215c
Коммит 03fba66e71
5 изменённых файлов: 74 добавлений и 56 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
src/include/api_trace.h
+1 -1
Просмотреть файл
@@ -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[],
src/misc/api_trace.cc
+4 -4
Просмотреть файл
@@ -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
src/misc/msccl/msccl_lifecycle.cc
+56 -1
Просмотреть файл
@@ -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));
src/msccl.cc
+7 -46
Просмотреть файл
@@ -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;
}
src/nccl.h.in
+6 -4
Просмотреть файл
@@ -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
#endif // end include guard