From 46a2e9e24e043b9223155027bf1d2fa660c49d65 Mon Sep 17 00:00:00 2001 From: Josh Romero Date: Wed, 21 Jan 2026 12:45:52 -0800 Subject: [PATCH 1/3] Fix potential deadlock during NCCL local communicator creation. Signed-off-by: Josh Romero --- src/cudecomp.cc | 9 +++++++-- 1 file changed, 7 insertions(+), 2 deletions(-) diff --git a/src/cudecomp.cc b/src/cudecomp.cc index dd82575..f5e55a5 100644 --- a/src/cudecomp.cc +++ b/src/cudecomp.cc @@ -655,8 +655,13 @@ cudecompResult_t cudecompGridDescCreate(cudecompHandle_t handle, cudecompGridDes setCommInfo(handle, grid_desc, col_comm, CUDECOMP_COMM_COL); // Create local NCCL communicator if row or column communicator uses it - if ((grid_desc->row_comm_info.ngroups == 1 && grid_desc->row_comm_info.nranks > 1) || - (grid_desc->col_comm_info.ngroups == 1 && grid_desc->col_comm_info.nranks > 1)) { + int need_local_nccl_comm = static_cast((grid_desc->row_comm_info.ngroups == 1 && grid_desc->row_comm_info.nranks > 1) || + (grid_desc->col_comm_info.ngroups == 1 && grid_desc->col_comm_info.nranks > 1)); + + // Local comm can include ranks in other rows/columns, need additional check for those cases. + CHECK_MPI(MPI_Allreduce(MPI_IN_PLACE, &need_local_nccl_comm, 1, MPI_INT, MPI_LOR, handle->mpi_local_comm)); + + if (need_local_nccl_comm) { handle->nccl_local_comm = ncclCommFromMPIComm( handle->mpi_clique_comm != MPI_COMM_NULL ? handle->mpi_clique_comm : handle->mpi_local_comm); } From f22dce0a71a7d2aba529a799c95d151d97b7f7bc Mon Sep 17 00:00:00 2001 From: Josh Romero Date: Wed, 21 Jan 2026 13:02:01 -0800 Subject: [PATCH 2/3] Use clique comm when needed. Signed-off-by: Josh Romero --- src/cudecomp.cc | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/src/cudecomp.cc b/src/cudecomp.cc index f5e55a5..08278cf 100644 --- a/src/cudecomp.cc +++ b/src/cudecomp.cc @@ -659,7 +659,8 @@ cudecompResult_t cudecompGridDescCreate(cudecompHandle_t handle, cudecompGridDes (grid_desc->col_comm_info.ngroups == 1 && grid_desc->col_comm_info.nranks > 1)); // Local comm can include ranks in other rows/columns, need additional check for those cases. - CHECK_MPI(MPI_Allreduce(MPI_IN_PLACE, &need_local_nccl_comm, 1, MPI_INT, MPI_LOR, handle->mpi_local_comm)); + CHECK_MPI(MPI_Allreduce(MPI_IN_PLACE, &need_local_nccl_comm, 1, MPI_INT, MPI_LOR, + handle->mpi_clique_comm != MPI_COMM_NULL ? handle->mpi_clique_comm : handle->mpi_local_comm)); if (need_local_nccl_comm) { handle->nccl_local_comm = ncclCommFromMPIComm( From 40eadb7b4e3a2039a2ba1081426b009fd7761069 Mon Sep 17 00:00:00 2001 From: Josh Romero Date: Wed, 21 Jan 2026 13:13:36 -0800 Subject: [PATCH 3/3] Formatting. Signed-off-by: Josh Romero --- src/cudecomp.cc | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/src/cudecomp.cc b/src/cudecomp.cc index 08278cf..fb43822 100644 --- a/src/cudecomp.cc +++ b/src/cudecomp.cc @@ -655,12 +655,14 @@ cudecompResult_t cudecompGridDescCreate(cudecompHandle_t handle, cudecompGridDes setCommInfo(handle, grid_desc, col_comm, CUDECOMP_COMM_COL); // Create local NCCL communicator if row or column communicator uses it - int need_local_nccl_comm = static_cast((grid_desc->row_comm_info.ngroups == 1 && grid_desc->row_comm_info.nranks > 1) || - (grid_desc->col_comm_info.ngroups == 1 && grid_desc->col_comm_info.nranks > 1)); + int need_local_nccl_comm = + static_cast((grid_desc->row_comm_info.ngroups == 1 && grid_desc->row_comm_info.nranks > 1) || + (grid_desc->col_comm_info.ngroups == 1 && grid_desc->col_comm_info.nranks > 1)); // Local comm can include ranks in other rows/columns, need additional check for those cases. CHECK_MPI(MPI_Allreduce(MPI_IN_PLACE, &need_local_nccl_comm, 1, MPI_INT, MPI_LOR, - handle->mpi_clique_comm != MPI_COMM_NULL ? handle->mpi_clique_comm : handle->mpi_local_comm)); + handle->mpi_clique_comm != MPI_COMM_NULL ? handle->mpi_clique_comm + : handle->mpi_local_comm)); if (need_local_nccl_comm) { handle->nccl_local_comm = ncclCommFromMPIComm(