| 167 | } |
| 168 | |
| 169 | void Communicator::wait() { |
| 170 | if (!device_) { |
| 171 | // just return if it has not been synchronized |
| 172 | return; |
| 173 | } |
| 174 | |
| 175 | device_->Exec( |
| 176 | [this](Context *ctx) mutable { |
| 177 | // synchronizing on all the CUDA streams used by communicator |
| 178 | CUDA_CHECK(cudaEventRecord(event, ctx->s)); |
| 179 | CUDA_CHECK(cudaStreamWaitEvent(ctx->stream, event, 0)); |
| 180 | CUDA_CHECK(cudaEventRecord(event, ctx->c1)); |
| 181 | CUDA_CHECK(cudaStreamWaitEvent(ctx->stream, event, 0)); |
| 182 | CUDA_CHECK(cudaEventRecord(event, ctx->c2)); |
| 183 | CUDA_CHECK(cudaStreamWaitEvent(ctx->stream, event, 0)); |
| 184 | }, |
| 185 | blocks_, blocks_, "Waiting"); |
| 186 | } |
| 187 | |
| 188 | Communicator::~Communicator() { |
| 189 | // finalizing NCCL |