nccl_helper.h 4.3 KB
Newer Older
Y
Yu Yang 已提交
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16
//   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.

#pragma once

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

T
typhoonzero 已提交
24 25
#define NCCL_ID_VARNAME "NCCLID"

Y
Yu Yang 已提交
26 27 28 29 30 31 32 33 34 35
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;
36 37
  } else if (type == typeid(int64_t)) {  // NOLINT
    return ncclInt64;
Y
Yu Yang 已提交
38 39 40 41 42
  } else {
    PADDLE_THROW("Not supported");
  }
}

Y
Yu Yang 已提交
43 44
class NCCLGroupGuard {
 public:
Y
Yu Yang 已提交
45 46 47 48 49
  static std::mutex &NCCLMutex() {
    static std::mutex mtx;
    return mtx;
  }

Y
Yu Yang 已提交
50
  inline NCCLGroupGuard() {
Y
Yu Yang 已提交
51
    NCCLMutex().lock();
Y
Yu Yang 已提交
52 53
    PADDLE_ENFORCE(dynload::ncclGroupStart());
  }
Y
Yu Yang 已提交
54 55

  inline ~NCCLGroupGuard() {
Y
yuyang18 已提交
56
    CHECK_EQ(dynload::ncclGroupEnd(), ncclSuccess);
Y
Yu Yang 已提交
57
    NCCLMutex().unlock();
Y
Yu Yang 已提交
58 59 60
  }
};

Y
Yu Yang 已提交
61 62 63 64 65
struct NCCLContext {
  std::unique_ptr<CUDADeviceContext> ctx_;
  ncclComm_t comm_;

  explicit NCCLContext(int dev_id)
Y
Yu Yang 已提交
66
      : ctx_(new CUDADeviceContext(CUDAPlace(dev_id))), comm_{nullptr} {}
Y
Yu Yang 已提交
67 68 69 70 71 72 73 74

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

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

Y
Yu Yang 已提交
75 76 77 78
struct NCCLContextMap {
  std::unordered_map<int, NCCLContext> contexts_;
  std::vector<int> order_;

T
typhoonzero 已提交
79 80
  explicit NCCLContextMap(const std::vector<platform::Place> &places,
                          ncclUniqueId *nccl_id = nullptr,
T
typhoonzero 已提交
81
                          size_t num_trainers = 1, size_t trainer_id = 0) {
Y
Yu Yang 已提交
82
    PADDLE_ENFORCE(!places.empty());
Y
Yu Yang 已提交
83 84 85 86 87 88 89 90 91 92
    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 已提交
93 94 95 96 97 98
    if (places.size() <= 1) {
      return;
    }
    std::unique_ptr<ncclComm_t[]> comms(new ncclComm_t[order_.size()]);
    // if pass nccl_id here, can assume we are doing multi node training
    if (nccl_id == nullptr) {
T
typhoonzero 已提交
99 100 101
      std::lock_guard<std::mutex> guard(NCCLGroupGuard::NCCLMutex());
      PADDLE_ENFORCE(platform::dynload::ncclCommInitAll(
          comms.get(), static_cast<int>(order_.size()), order_.data()));
T
typhoonzero 已提交
102
    } else {
T
typhoonzero 已提交
103
      PADDLE_ENFORCE_GT(num_trainers, 1);
T
typhoonzero 已提交
104
      // TODO(wuyi): need to ensure each node have same number of GPUs
Y
Yu Yang 已提交
105
      {
T
typhoonzero 已提交
106
        int nranks = num_trainers * order_.size();
T
typhoonzero 已提交
107
        NCCLGroupGuard gurad;
T
typhoonzero 已提交
108 109
        for (auto &gpu_id : order_) {
          int rank = trainer_id * order_.size() + gpu_id;
T
typhoonzero 已提交
110
          VLOG(3) << "init nccl rank: " << rank << " nranks: " << nranks;
T
typhoonzero 已提交
111
          PADDLE_ENFORCE(cudaSetDevice(gpu_id));
T
testing  
typhoonzero 已提交
112 113
          PADDLE_ENFORCE(platform::dynload::ncclCommInitRank(
              comms.get() + gpu_id, nranks, *nccl_id, rank));
T
typhoonzero 已提交
114
        }
Y
Yu Yang 已提交
115
      }
Y
Yu Yang 已提交
116
    }
T
typhoonzero 已提交
117 118 119 120
    int i = 0;
    for (auto &dev_id : order_) {
      contexts_.at(dev_id).comm_ = comms[i++];
    }
Y
Yu Yang 已提交
121 122
  }

Y
Yu Yang 已提交
123 124 125
  NCCLContextMap(const NCCLContextMap &other) = delete;
  NCCLContextMap &operator=(const NCCLContextMap &other) = delete;

Y
Yu Yang 已提交
126 127 128 129 130 131 132 133 134 135 136 137 138 139 140 141 142 143 144
  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 已提交
145 146
}  // namespace platform
}  // namespace paddle