nccl_helper.h 4.6 KB
Newer Older
Y
Yu Yang 已提交
1 2 3 4 5 6 7 8 9 10 11 12 13 14
//   Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
//     http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.

P
peizhilin 已提交
15
#ifndef _WIN32
Y
Yu Yang 已提交
16 17
#pragma once

T
typhoonzero 已提交
18
#include <stdio.h>
19
#include <string>
20
#include <thread>  // NOLINT
Y
Yu Yang 已提交
21
#include <typeindex>
22
#include <vector>
Y
Yu Yang 已提交
23 24 25
#include "paddle/fluid/platform/dynload/nccl.h"
#include "paddle/fluid/platform/enforce.h"

T
typhoonzero 已提交
26 27
#define NCCL_ID_VARNAME "NCCLID"

Y
Yu Yang 已提交
28 29 30 31 32 33 34 35 36 37
namespace paddle {
namespace platform {

inline ncclDataType_t ToNCCLDataType(std::type_index type) {
  if (type == typeid(float)) {  // NOLINT
    return ncclFloat;
  } else if (type == typeid(double)) {  // NOLINT
    return ncclDouble;
  } else if (type == typeid(int)) {  // NOLINT
    return ncclInt;
38 39
  } else if (type == typeid(int64_t)) {  // NOLINT
    return ncclInt64;
Y
Yu Yang 已提交
40 41 42 43 44
  } else {
    PADDLE_THROW("Not supported");
  }
}

45 46 47 48 49
// NOTE(minqiyang): according to the ncclGroupEnd documentations:
// https://docs.nvidia.com/deeplearning/sdk/nccl-api/ncclapidoc.html,
// ncclGroupEnd will wait for all communicators to be initialized, which will
// cause blocking problem when a runtime_error was thrown, so try only guard
// NCCL actions when use it.
Y
Yu Yang 已提交
50 51
class NCCLGroupGuard {
 public:
Y
Yu Yang 已提交
52 53 54 55 56
  static std::mutex &NCCLMutex() {
    static std::mutex mtx;
    return mtx;
  }

Y
Yu Yang 已提交
57
  inline NCCLGroupGuard() {
Y
Yu Yang 已提交
58
    NCCLMutex().lock();
Y
Yu Yang 已提交
59 60
    PADDLE_ENFORCE(dynload::ncclGroupStart());
  }
Y
Yu Yang 已提交
61 62

  inline ~NCCLGroupGuard() {
Y
yuyang18 已提交
63
    CHECK_EQ(dynload::ncclGroupEnd(), ncclSuccess);
Y
Yu Yang 已提交
64
    NCCLMutex().unlock();
Y
Yu Yang 已提交
65 66 67
  }
};

Y
Yu Yang 已提交
68 69 70 71 72
struct NCCLContext {
  std::unique_ptr<CUDADeviceContext> ctx_;
  ncclComm_t comm_;

  explicit NCCLContext(int dev_id)
Y
Yu Yang 已提交
73
      : ctx_(new CUDADeviceContext(CUDAPlace(dev_id))), comm_{nullptr} {}
Y
Yu Yang 已提交
74 75 76 77 78 79 80 81

  cudaStream_t stream() const { return ctx_->stream(); }

  int device_id() const {
    return boost::get<platform::CUDAPlace>(ctx_->GetPlace()).device;
  }
};

Y
Yu Yang 已提交
82 83 84 85
struct NCCLContextMap {
  std::unordered_map<int, NCCLContext> contexts_;
  std::vector<int> order_;

T
typhoonzero 已提交
86 87
  explicit NCCLContextMap(const std::vector<platform::Place> &places,
                          ncclUniqueId *nccl_id = nullptr,
T
typhoonzero 已提交
88
                          size_t num_trainers = 1, size_t trainer_id = 0) {
Y
Yu Yang 已提交
89
    PADDLE_ENFORCE(!places.empty());
Y
Yu Yang 已提交
90 91 92 93 94 95 96 97 98 99
    order_.reserve(places.size());
    for (auto &p : places) {
      int dev_id = boost::get<CUDAPlace>(p).device;
      order_.emplace_back(dev_id);
      contexts_.emplace(dev_id, NCCLContext(dev_id));
    }
    PADDLE_ENFORCE_EQ(
        order_.size(), contexts_.size(),
        "NCCL Context Map does not support contain two or more same device");

T
typhoonzero 已提交
100 101 102 103
    if (places.size() <= 1) {
      return;
    }
    std::unique_ptr<ncclComm_t[]> comms(new ncclComm_t[order_.size()]);
W
Wu Yi 已提交
104 105
    // if num_trainers == 1, should create a new nccl id for local comms.
    if (num_trainers == 1) {
T
typhoonzero 已提交
106 107 108
      std::lock_guard<std::mutex> guard(NCCLGroupGuard::NCCLMutex());
      PADDLE_ENFORCE(platform::dynload::ncclCommInitAll(
          comms.get(), static_cast<int>(order_.size()), order_.data()));
T
typhoonzero 已提交
109
    } else {
W
Wu Yi 已提交
110
      PADDLE_ENFORCE_NOT_NULL(nccl_id);
Y
Yu Yang 已提交
111
      {
T
typhoonzero 已提交
112
        int nranks = num_trainers * order_.size();
T
typhoonzero 已提交
113
        NCCLGroupGuard gurad;
T
typhoonzero 已提交
114 115
        for (auto &gpu_id : order_) {
          int rank = trainer_id * order_.size() + gpu_id;
T
typhoonzero 已提交
116
          VLOG(3) << "init nccl rank: " << rank << " nranks: " << nranks;
T
typhoonzero 已提交
117
          PADDLE_ENFORCE(cudaSetDevice(gpu_id));
T
testing  
typhoonzero 已提交
118 119
          PADDLE_ENFORCE(platform::dynload::ncclCommInitRank(
              comms.get() + gpu_id, nranks, *nccl_id, rank));
T
typhoonzero 已提交
120
        }
Y
Yu Yang 已提交
121
      }
Y
Yu Yang 已提交
122
    }
T
typhoonzero 已提交
123 124 125 126
    int i = 0;
    for (auto &dev_id : order_) {
      contexts_.at(dev_id).comm_ = comms[i++];
    }
Y
Yu Yang 已提交
127 128
  }

Y
Yu Yang 已提交
129 130 131
  NCCLContextMap(const NCCLContextMap &other) = delete;
  NCCLContextMap &operator=(const NCCLContextMap &other) = delete;

Y
Yu Yang 已提交
132 133 134 135 136 137 138 139 140 141 142 143 144 145 146 147 148 149 150
  CUDADeviceContext *DevCtx(int dev_id) const { return at(dev_id).ctx_.get(); }

  CUDADeviceContext *DevCtx(platform::Place p) const {
    return DevCtx(boost::get<CUDAPlace>(p).device);
  }

  const NCCLContext &at(platform::Place p) const {
    return this->at(boost::get<CUDAPlace>(p).device);
  }

  const NCCLContext &at(int dev_id) const { return contexts_.at(dev_id); }

  void WaitAll() {
    for (auto &p : contexts_) {
      p.second.ctx_->Wait();
    }
  }
};

Y
Yu Yang 已提交
151 152
}  // namespace platform
}  // namespace paddle
P
peizhilin 已提交
153
#endif