Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Add comm related nccl APIs #430

Closed
wants to merge 3 commits into from
Closed

Add comm related nccl APIs #430

Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
7446c02
add ncclCommSplit
Binyang2014 Dec 20, 2024
f36a733
add ncclCommSplit
Binyang2014 Dec 21, 2024
38100e7
Merge branch 'main' into binyli/nccl-api-commsplit
Binyang2014 Dec 26, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
39 changes: 36 additions & 3 deletions apps/nccl/src/nccl.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -69,6 +69,12 @@ struct ChannelInfo {
std::shared_ptr<mscclpp::DeviceHandle<mscclpp::SmChannel>> smChannelDeviceHandles;
};

struct splitCommInfo {
int color;
int key;
int originalRank;
};

struct ncclComm {
std::shared_ptr<mscclpp::Communicator> comm;
std::vector<std::shared_ptr<mscclpp::Connection>> connections;
Expand Down Expand Up @@ -421,9 +427,36 @@ NCCL_API ncclResult_t ncclCommAbort(ncclComm_t) {
return ncclSuccess;
}

NCCL_API ncclResult_t ncclCommSplit(ncclComm_t, int, int, ncclComm_t*, ncclConfig_t*) {
// TODO: implement this function
return ncclInternalError;
NCCL_API ncclResult_t ncclCommSplit(ncclComm_t comm, int color, int key, ncclComm_t* newcomm, ncclConfig_t*) {
*newcomm = NCCL_COMM_NULL;
int nRanks = comm->comm->bootstrap()->getNranks();
int rank = comm->comm->bootstrap()->getRank();
splitCommInfo info{color, key, comm->comm->bootstrap()->getRank()};
std::vector<splitCommInfo> infos(nRanks);
infos[rank] = info;
comm->comm->bootstrap()->allGather(infos.data(), sizeof(splitCommInfo));
comm->comm->bootstrap()->barrier();
std::vector<splitCommInfo> group;
std::copy_if(infos.begin(), infos.end(), std::back_inserter(group),
[color](const splitCommInfo& info) { return info.color == color; });
std::sort(group.begin(), group.end(), [](const splitCommInfo& a, const splitCommInfo& b) { return a.key < b.key; });
int newRank = std::distance(group.begin(),
std::find_if(group.begin(), group.end(),
[rank](const splitCommInfo& info) { return info.originalRank == rank; }));
int groupSize = group.size();
ncclUniqueId uniqueId;
if (newRank == 0) {
ncclGetUniqueId(&uniqueId);
}
std::vector<ncclUniqueId> uniqueIds(nRanks);
uniqueIds[rank] = uniqueId;
comm->comm->bootstrap()->allGather(uniqueIds.data(), sizeof(ncclUniqueId));
comm->comm->bootstrap()->barrier();
uniqueId = uniqueIds[group.front().originalRank];
if (color == NCCL_SPLIT_NOCOLOR) {
return ncclSuccess;
}
return ncclCommInitRankConfig(newcomm, groupSize, uniqueId, newRank, nullptr);
}

NCCL_API const char* ncclGetErrorString(ncclResult_t result) {
Expand Down
Loading