Does calling ncclGroupEnd imply that kernels are launched on multiple devices atomically? If so, we can remove the mutex. If not, we have to keep it around. If cudaFree races with the NCCL internals that are responsible for launching kernels on multiple devices, we reintroduce the potential deadlock.

I believe this PR is related to #14870.

NCCL 2.x documentation concerning Group Calls: https://docs.nvidia.com/deeplearning/sdk/nccl-developer-guide/docs/usage/groups.html

From the NCCL sources, in ncclGroupEnd:

  /* Collectives are done in three steps :
   * 1. Barrier Check In. Only the last call may call cudaLaunchKernel[cooperative]
   * 2. Barrier Wait. No CUDA call is permitted
   * 3. Enqueue Events. CUDA event wait/enqueue.
   * This is needed because step 2 cannot call any CUDA primitive, otherwise if
   * cudaFree happens between 1 and 3, it could block that CUDA call and
   * prevent some ranks from launching their network threads, which would
   * prevent the NCCL call from completing, blocking the cudaFree call.
   */

In light of the above, it's not clear to me whether NCCL 2 fixed the race with cudaFree or not. However, the NCCL 2 code guarantees that only the last CPU thread to arrive at the collective for a given communicator will perform the launching for all devices. For sufficiently new (CUDART_VERSION >= 9000) and capable hardware, it will use cudaCooperativeLaunchMultiDevice.

An alternative to this PR would be to remove the cuda_free_mutex and replace it with the recursive_mutex in the same file. This would also prevent deadlock, but without removing the cudaFree race protection. But it is not clear what the performance impact would be relative to just removing the cuda_free_mutex when NCCL v2 is in use.

Additional related issues/discussions:

• https://discuss.pytorch.org/t/distributed-data-parallel-freezes-without-error-message/8009
• https://discuss.pytorch.org/t/distributeddataparallel-deadlock/14809

Thanks for the detailed analysis, @jeffdaily. I went and looked at the NCCL code and did some more investigation myself today and found the following:

• Cooperative multi device launch is used if the CUDA version is >= 9.0 and the devices have cudaDevAttrCooperativeMultiDeviceLaunch set. I wrote a quick program to check this attribute on a few classes of GPUs and found this is supported since the Pascal generation. For older GPUs, it will loop over the participating devices and launch the kernels sequentially. We still support Kepler and Maxwell generation GPUs, so have to deal with the non-atomic multi device launch case.
• The cudaFree race is not solved. If during sequential launch (either in non-atomic group mode or parallel launch mode) any process runs a cudaFree on another thread, it risks deadlock. If I recall correctly, calling cudaFree triggers device synchronization and waits for all kernels to exit. If any of these are NCCL kernels, and not all of them have been launched, they will hang.

Regarding the lock ordering, this is indeed an issue. I like the idea of using the AutoNCCLGroup here. While we can't remove the free lock, we can narrow its scope to avoid locking the main allocator lock while holding the free lock. Then we'd have 2 loops: one to launch the NCCL kernels, and one to record streams. They don't depend on each other so we can sequence them one after the other to fix this. This doesn't depend on AutoNCCLGroup of course and can be done independently.

Hi @pietern . I have rebased the PR and revised it based on your comments. The free lock remains and its lifetime is managed by the AutoNcclGroup. There are now two loops, one to record the streams and one to launch the NCCL kernels. The stream recording happens first and does so without holding the free lock. Then, the NCCL kernel launches are within an AutoNcclGroup scope block.

Thanks for updating the PR @jeffdaily. This looks good to me. @mrshenli can you take a look as well?

@pietern I am looking now

Thanks @jeffdaily and @mrshenli!

@pietern merged this pull request in 29347cc.

This PR breaks master CUDA builds with:

Jul 18 14:59:18 In file included from /var/lib/jenkins/workspace/torch/lib/c10d/../c10d/ProcessGroupNCCL.hpp:6:0,
Jul 18 14:59:18 from /var/lib/jenkins/workspace/torch/lib/c10d/ProcessGroupNCCL.cpp:1:
Jul 18 14:59:18 /var/lib/jenkins/workspace/torch/lib/c10d/ProcessGroupNCCL.cpp: In destructor 'c10d::{anonymous}::AutoNcclGroup::~AutoNcclGroup()':
Jul 18 14:59:18 /var/lib/jenkins/workspace/torch/lib/c10d/../c10d/NCCLUtils.hpp:14:35: error: throw will always call terminate() [-Werror=terminate]
Jul 18 14:59:18 throw std::runtime_error(err);
Jul 18 14:59:18 ^
Jul 18 14:59:18 /var/lib/jenkins/workspace/torch/lib/c10d/ProcessGroupNCCL.cpp:27:5: note: in expansion of macro 'C10D_NCCL_CHECK'
Jul 18 14:59:18 C10D_NCCL_CHECK(ncclGroupEnd());
Jul 18 14:59:18 ^
Jul 18 14:59:18 /var/lib/jenkins/workspace/torch/lib/c10d/../c10d/NCCLUtils.hpp:14:35: note: in C++11 destructors default to noexcept
Jul 18 14:59:18 throw std::runtime_error(err);
Jul 18 14:59:18 ^
Jul 18 14:59:18 /var/lib/jenkins/workspace/torch/lib/c10d/ProcessGroupNCCL.cpp:27:5: note: in expansion of macro 'C10D_NCCL_CHECK'
Jul 18 14:59:18 C10D_NCCL_CHECK(ncclGroupEnd());
Jul 18 14:59:18 ^
Jul 18 14:59:18 cc1plus: all warnings being treated as errors
Jul 18 14:59:18 caffe2/lib_c10d/CMakeFiles/c10d.dir/build.make:206: recipe for target 'caffe2/lib_c10d/CMakeFiles/c10d.dir/ProcessGroupNCCL.cpp.o' failed
Jul 18 14:59:18 make[2]: *** [caffe2/lib_c10d/CMakeFiles/c10d.dir/ProcessGroupNCCL.cpp.o] Error 1
Jul 18 14:59:18 CMakeFiles/Makefile2:11636: recipe for target 'caffe2/lib_c10d/CMakeFiles/c10d.dir/all' failed
Jul 18 14:59:18 make[1]: *** [caffe2/lib_c10d/CMakeFiles/c10d.dir/all] Error 2

I am reverting it.

How did CI not catch this earlier? This AutoNcclGroup code is basically copied from torch/csrc/cuda/nccl.h, how does it work there?

@yf225 can we fix this with a new PR rather than revert the entire PR?

@jeffdaily If the fix PR takes longer, it would mean that all other PRs based on current master have their CUDA builds broken, which I don't think is a good idea.

Since most of this PR has been reviewed, opening a new PR with the fix added shouldn't take too long to review, and we can make sure CI is green before merging it.

@yf225 ok to revert. I will start the new PR now, with the fix added.

@ezyang do you know why our CI didn't catch this error?

@mrshenli I suspect that it's because the last CI run was too old (16 days ago)

@ezyang @mrshenli I have submitted PR #23040 to hopefully address the new CI failure.