torch/lib/c10d/ProcessGroupNCCL.cpp
// Helper that lets the input ncclStreams to wait for the THC stream
Helper that lets the input ncclStreams to wait for the THC stream
Helper that lets the input ncclStreams to wait for the THC stream
void syncStreams(
    THCState* thcState,
    const std::vector<at::Device>& devices,
    std::vector<CUDAEvent>& ncclEvents,
    std::vector<CUDAStream>& ncclStreams) {
  at::DeviceGuard gpuGuard;
  for (size_t i = 0; i < devices.size(); ++i) {
    gpuGuard.set_index(devices[i].index());
    auto currentThcStream =
        THCState_getCurrentStreamOnDevice(thcState, devices[i].index());
    CUDAStream& ncclStream = ncclStreams[i];
    CUDAEvent& ncclEvent = ncclEvents[i];
    C10D_CUDA_CHECK(cudaEventRecord(ncclEvent.getEvent(), currentThcStream));
    C10D_CUDA_CHECK(
        cudaStreamWaitEvent(ncclStream.getStream(), ncclEvent.getEvent(), 0));
  }
}
void syncStreams ( THCState * thcState , const std :: vector < at :: Device > & devices , std :: vector < CUDAEvent > & ncclEvents , std :: vector < CUDAStream > & ncclStreams ) { at :: DeviceGuard gpuGuard ; for ( size_t i = 0 ; i < devices . size ( ) ; ++ i ) { gpuGuard . set_index ( devices [ i ] . index ( ) ) ; auto currentThcStream = THCState_getCurrentStreamOnDevice ( thcState , devices [ i ] . index ( ) ) ; CUDAStream & ncclStream = ncclStreams [ i ] ; CUDAEvent & ncclEvent = ncclEvents [ i ] ; C10D_CUDA_CHECK ( cudaEventRecord ( ncclEvent . getEvent ( ) , currentThcStream ) ) ; C10D_CUDA_CHECK ( cudaStreamWaitEvent ( ncclStream . getStream ( ) , ncclEvent . getEvent ( ) , 0 ) ) ; } }
Helper that lets the input ncclStreams to wait for the THC stream
Helper that lets the input ncclStreams to wait for the THC stream