From 6f2b50a8746ef2b99f8fb24708029c6914d35dd1 Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Tue, 22 Feb 2022 14:48:39 -0800 Subject: [PATCH 1/3] increase parallelism in allgatherv --- cpp/include/raft/comms/detail/mpi_comms.hpp | 2 ++ cpp/include/raft/comms/detail/std_comms.hpp | 2 ++ 2 files changed, 4 insertions(+) diff --git a/cpp/include/raft/comms/detail/mpi_comms.hpp b/cpp/include/raft/comms/detail/mpi_comms.hpp index ec1101032e..148566ca02 100644 --- a/cpp/include/raft/comms/detail/mpi_comms.hpp +++ b/cpp/include/raft/comms/detail/mpi_comms.hpp @@ -277,6 +277,7 @@ class mpi_comms : public comms_iface { { // From: "An Empirical Evaluation of Allgatherv on Multi-GPU Systems" - // https://arxiv.org/pdf/1812.05964.pdf Listing 1 on page 4. + RAFT_NCCL_TRY(ncclGroupStart()); for (int root = 0; root < size_; ++root) { RAFT_NCCL_TRY( ncclBroadcast(sendbuf, @@ -287,6 +288,7 @@ class mpi_comms : public comms_iface { nccl_comm_, stream)); } + RAFT_NCCL_TRY(ncclGroupEnd()); } void gather(const void* sendbuff, diff --git a/cpp/include/raft/comms/detail/std_comms.hpp b/cpp/include/raft/comms/detail/std_comms.hpp index 1a4cc2fcf9..30b58066c8 100644 --- a/cpp/include/raft/comms/detail/std_comms.hpp +++ b/cpp/include/raft/comms/detail/std_comms.hpp @@ -367,6 +367,7 @@ class std_comms : public comms_iface { { // From: "An Empirical Evaluation of Allgatherv on Multi-GPU Systems" - // https://arxiv.org/pdf/1812.05964.pdf Listing 1 on page 4. + RAFT_NCCL_TRY(ncclGroupStart()); for (int root = 0; root < num_ranks_; ++root) { size_t dtype_size = get_datatype_size(datatype); RAFT_NCCL_TRY(ncclBroadcast(sendbuf, @@ -377,6 +378,7 @@ class std_comms : public comms_iface { nccl_comm_, stream)); } + RAFT_NCCL_TRY(ncclGroupEnd()); } void gather(const void* sendbuff, From affb7df8a7965e00fd784e9fe3a603246262d27e Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Tue, 22 Feb 2022 16:59:34 -0800 Subject: [PATCH 2/3] add check for # operations between ncclGroupStart & ncclGroupEnd to be less than 2048 --- cpp/include/raft/comms/detail/mpi_comms.hpp | 2 ++ cpp/include/raft/comms/detail/std_comms.hpp | 2 ++ 2 files changed, 4 insertions(+) diff --git a/cpp/include/raft/comms/detail/mpi_comms.hpp b/cpp/include/raft/comms/detail/mpi_comms.hpp index 148566ca02..f06506888c 100644 --- a/cpp/include/raft/comms/detail/mpi_comms.hpp +++ b/cpp/include/raft/comms/detail/mpi_comms.hpp @@ -275,6 +275,8 @@ class mpi_comms : public comms_iface { datatype_t datatype, cudaStream_t stream) const { + RAFT_EXPECTS(size_ <= 2048, + "# NCCL operations between ncclGroupStart & ncclGroupEnd shouldn't exceed 2048."); // From: "An Empirical Evaluation of Allgatherv on Multi-GPU Systems" - // https://arxiv.org/pdf/1812.05964.pdf Listing 1 on page 4. RAFT_NCCL_TRY(ncclGroupStart()); diff --git a/cpp/include/raft/comms/detail/std_comms.hpp b/cpp/include/raft/comms/detail/std_comms.hpp index 30b58066c8..17a62d96ec 100644 --- a/cpp/include/raft/comms/detail/std_comms.hpp +++ b/cpp/include/raft/comms/detail/std_comms.hpp @@ -367,6 +367,8 @@ class std_comms : public comms_iface { { // From: "An Empirical Evaluation of Allgatherv on Multi-GPU Systems" - // https://arxiv.org/pdf/1812.05964.pdf Listing 1 on page 4. + RAFT_EXPECTS(size_ <= 2048, + "# NCCL operations between ncclGroupStart & ncclGroupEnd shouldn't exceed 2048."); RAFT_NCCL_TRY(ncclGroupStart()); for (int root = 0; root < num_ranks_; ++root) { size_t dtype_size = get_datatype_size(datatype); From a7ba3568c7918401979a2f6128d24891d0c88379 Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Tue, 22 Feb 2022 17:27:52 -0800 Subject: [PATCH 3/3] fix compile error --- cpp/include/raft/comms/detail/std_comms.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/include/raft/comms/detail/std_comms.hpp b/cpp/include/raft/comms/detail/std_comms.hpp index 17a62d96ec..0d54a7e55c 100644 --- a/cpp/include/raft/comms/detail/std_comms.hpp +++ b/cpp/include/raft/comms/detail/std_comms.hpp @@ -367,7 +367,7 @@ class std_comms : public comms_iface { { // From: "An Empirical Evaluation of Allgatherv on Multi-GPU Systems" - // https://arxiv.org/pdf/1812.05964.pdf Listing 1 on page 4. - RAFT_EXPECTS(size_ <= 2048, + RAFT_EXPECTS(num_ranks_ <= 2048, "# NCCL operations between ncclGroupStart & ncclGroupEnd shouldn't exceed 2048."); RAFT_NCCL_TRY(ncclGroupStart()); for (int root = 0; root < num_ranks_; ++root) {