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

Modify comm_split to avoid ucp #1649

Merged
merged 4 commits into from
Jul 24, 2023
Merged

Modify comm_split to avoid ucp #1649

Changes from 2 commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
bcf32b0
Modify comm_split to do allgather instead of requiring UCP
ChuckHastings Jul 14, 2023
77e16b3
fix code formatting
ChuckHastings Jul 14, 2023
5546d6f
Use sync_stream instead of cudaStreamSynchronize in the new code
ChuckHastings Jul 24, 2023
40031ac
Merge branch 'branch-23.08' into modify_split_to_avoid_ucp
ChuckHastings Jul 24, 2023
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
69 changes: 30 additions & 39 deletions cpp/include/raft/comms/detail/std_comms.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -28,6 +28,8 @@

#include <raft/util/cudart_utils.hpp>

#include <thrust/iterator/zip_iterator.h>

#include <cuda_runtime.h>

#include <ucp/api/ucp.h>
Expand Down Expand Up @@ -140,48 +142,37 @@ class std_comms : public comms_iface {

RAFT_CUDA_TRY(cudaStreamSynchronize(stream_));
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Now that we are using nccl instead of ucp, we should probably avoid any potential deadlocks from dying nccl ranks by using sync_stream here instead of the standard cuda synchronization.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Updated in latest push


std::vector<int> subcomm_ranks{};
std::vector<ucp_ep_h> new_ucx_ptrs{};

for (int i = 0; i < get_size(); ++i) {
if (h_colors[i] == color) {
subcomm_ranks.push_back(i);
if (ucp_worker_ != nullptr && subcomms_ucp_) { new_ucx_ptrs.push_back((*ucp_eps_)[i]); }
}
}
ncclComm_t nccl_comm;

// Create a structure to allgather...
ncclUniqueId id{};
if (get_rank() == subcomm_ranks[0]) { // root of the new subcommunicator
RAFT_NCCL_TRY(ncclGetUniqueId(&id));
std::vector<request_t> requests(subcomm_ranks.size() - 1);
for (size_t i = 1; i < subcomm_ranks.size(); ++i) {
isend(&id, sizeof(ncclUniqueId), subcomm_ranks[i], color, requests.data() + (i - 1));
}
waitall(requests.size(), requests.data());
} else {
request_t request{};
irecv(&id, sizeof(ncclUniqueId), subcomm_ranks[0], color, &request);
waitall(1, &request);
}
// FIXME: this seems unnecessary, do more testing and remove this
barrier();
rmm::device_uvector<ncclUniqueId> d_nccl_ids(get_size(), stream_);

ncclComm_t nccl_comm;
RAFT_NCCL_TRY(ncclCommInitRank(&nccl_comm, subcomm_ranks.size(), id, key));

if (ucp_worker_ != nullptr && subcomms_ucp_) {
auto eps_sp = std::make_shared<ucp_ep_h*>(new_ucx_ptrs.data());
return std::unique_ptr<comms_iface>(new std_comms(nccl_comm,
(ucp_worker_h)ucp_worker_,
eps_sp,
subcomm_ranks.size(),
key,
stream_,
subcomms_ucp_));
} else {
return std::unique_ptr<comms_iface>(
new std_comms(nccl_comm, subcomm_ranks.size(), key, stream_));
}
if (key == 0) { RAFT_NCCL_TRY(ncclGetUniqueId(&id)); }

update_device(d_nccl_ids.data() + get_rank(), &id, 1, stream_);

allgather(d_nccl_ids.data() + get_rank(),
d_nccl_ids.data(),
sizeof(ncclUniqueId),
datatype_t::UINT8,
stream_);

auto offset =
std::distance(thrust::make_zip_iterator(h_colors.begin(), h_keys.begin()),
std::find_if(thrust::make_zip_iterator(h_colors.begin(), h_keys.begin()),
thrust::make_zip_iterator(h_colors.end(), h_keys.end()),
[color](auto tuple) { return thrust::get<0>(tuple) == color; }));

auto subcomm_size = std::count(h_colors.begin(), h_colors.end(), color);

update_host(&id, d_nccl_ids.data() + offset, 1, stream_);

RAFT_CUDA_TRY(cudaStreamSynchronize(stream_));
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We should probably use sync_stream() here so that this is more tolerant to failed ranks.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Updated in latest push


RAFT_NCCL_TRY(ncclCommInitRank(&nccl_comm, subcomm_size, id, key));

return std::unique_ptr<comms_iface>(new std_comms(nccl_comm, subcomm_size, key, stream_));
}

void barrier() const
Expand Down