torch/lib/c10d/ProcessGroupNCCL.cpp
// Waiting on the work's corresponding CUDA events
Waiting on the work's corresponding CUDA events
Waiting on the work ' s corresponding CUDA events
void ProcessGroupNCCL::WorkNCCL::synchronize() {
  auto thcState = ::at::globalContext().lazyInitCUDA();
  at::DeviceGuard gpuGuard;
  for (size_t i = 0; i < devices_.size(); ++i) {
    gpuGuard.set_index(devices_[i].index());
    auto thcStream =
        THCState_getCurrentStreamOnDevice(thcState, devices_[i].index());
    auto& cudaEvent = cudaEvents_[i];
    C10D_CUDA_CHECK(cudaStreamWaitEvent(thcStream, cudaEvent.getEvent(), 0));
  }
}
void ProcessGroupNCCL :: WorkNCCL :: synchronize ( ) { auto thcState = :: at :: globalContext ( ) . lazyInitCUDA ( ) ; at :: DeviceGuard gpuGuard ; for ( size_t i = 0 ; i < devices_ . size ( ) ; ++ i ) { gpuGuard . set_index ( devices_ [ i ] . index ( ) ) ; auto thcStream = THCState_getCurrentStreamOnDevice ( thcState , devices_ [ i ] . index ( ) ) ; auto & cudaEvent = cudaEvents_ [ i ] ; C10D_CUDA_CHECK ( cudaStreamWaitEvent ( thcStream , cudaEvent . getEvent ( ) , 0 ) ) ; } }
Waiting on the work's corresponding CUDA events
Waiting on the work ' s corresponding CUDA events
// Let THC stream wait for the NCCL stream