device_context.cc 16.3 KB
Newer Older
1
/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserved.
Q
qijun 已提交
2 3 4 5
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
6

Q
qijun 已提交
7 8 9 10 11
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. */
Y
Yi Wang 已提交
12
#include "paddle/fluid/platform/device_context.h"
13
#include <set>
14
#include <string>
Y
Yu Yang 已提交
15
#include <unordered_set>
16 17
#include <vector>

Y
Yi Wang 已提交
18
#include "paddle/fluid/memory/memory.h"
19 20
#ifdef PADDLE_WITH_CUDA
#include "paddle/fluid/framework/rw_lock.h"
21
#include "paddle/fluid/memory/allocation/cuda_device_context_allocator.h"
S
sneaxiy 已提交
22
#include "paddle/fluid/platform/cuda_device_guard.h"
23
#endif
24

25 26
#include "glog/logging.h"

27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53
namespace paddle {
namespace memory {

AllocationPtr Alloc(const platform::DeviceContext& dev_ctx, size_t size) {
  auto place = dev_ctx.GetPlace();
#ifdef PADDLE_WITH_CUDA
  if (size == 0 || !platform::is_gpu_place(place)) {
    return Alloc(place, size);
  }
  auto* default_dev_ctx = static_cast<platform::CUDADeviceContext*>(
      platform::DeviceContextPool::Instance().Get(place));
  auto& desired_dev_ctx =
      static_cast<const platform::CUDADeviceContext&>(dev_ctx);
  if (default_dev_ctx->stream() == desired_dev_ctx.stream()) {
    return Alloc(place, size);
  } else {
    return allocation::CUDADeviceContextAllocatorPool::Instance().Alloc(
        desired_dev_ctx, size);
  }
#else
  return Alloc(place, size);
#endif
}

}  // namespace memory
}  // namespace paddle

Q
qijun 已提交
54 55 56
namespace paddle {
namespace platform {

D
dzhwinter 已提交
57 58
DeviceContextPool* DeviceContextPool::pool = nullptr;

Y
Yu Yang 已提交
59
platform::DeviceContext* DeviceContextPool::Get(const platform::Place& place) {
D
dzhwinter 已提交
60 61 62
  auto it = device_contexts_.find(place);
  if (it == device_contexts_.end()) {
    PADDLE_THROW(
63 64 65 66
        "Place %s is not supported, Please check that your paddle compiles "
        "with WITH_GPU "
        "option or check that your train process hold the correct gpu_id if "
        "you use Executor",
M
minqiyang 已提交
67
        place);
D
dzhwinter 已提交
68
  }
69
  return it->second.get().get();
D
dzhwinter 已提交
70 71
}

72 73 74 75 76 77 78 79 80 81 82
template <typename DevCtx, typename PlaceType>
inline void EmplaceDeviceContext(
    std::map<Place, std::shared_future<std::unique_ptr<DeviceContext>>>*
        map_ptr,
    platform::Place p) {
  using PtrType = std::unique_ptr<DeviceContext>;
  map_ptr->emplace(p, std::async(std::launch::deferred, [=] {
                     // lazy evaluation. i.e., only create device context at
                     // first `Get`
                     return PtrType(new DevCtx(boost::get<PlaceType>(p)));
                   }));
C
chengduozh 已提交
83 84
}

D
dzhwinter 已提交
85 86 87
DeviceContextPool::DeviceContextPool(
    const std::vector<platform::Place>& places) {
  PADDLE_ENFORCE_GT(places.size(), 0);
88
  std::set<Place> set;
Y
Yu Yang 已提交
89 90 91 92 93
  for (auto& p : places) {
    set.insert(p);
  }
  for (auto& p : set) {
    if (platform::is_cpu_place(p)) {
94
#ifdef PADDLE_WITH_MKLDNN
95
      EmplaceDeviceContext<MKLDNNDeviceContext, CPUPlace>(&device_contexts_, p);
96
#else
97
      EmplaceDeviceContext<CPUDeviceContext, CPUPlace>(&device_contexts_, p);
98
#endif
Y
Yu Yang 已提交
99
    } else if (platform::is_gpu_place(p)) {
D
dzhwinter 已提交
100
#ifdef PADDLE_WITH_CUDA
101
      EmplaceDeviceContext<CUDADeviceContext, CUDAPlace>(&device_contexts_, p);
D
dzhwinter 已提交
102 103
#else
      PADDLE_THROW(
D
dzhwinter 已提交
104
          "'CUDAPlace' is not supported, Please re-compile with WITH_GPU "
D
dzhwinter 已提交
105
          "option");
C
chengduoZH 已提交
106 107 108
#endif
    } else if (platform::is_cuda_pinned_place(p)) {
#ifdef PADDLE_WITH_CUDA
109 110
      EmplaceDeviceContext<CUDAPinnedDeviceContext, CUDAPinnedPlace>(
          &device_contexts_, p);
C
chengduoZH 已提交
111 112 113 114
#else
      PADDLE_THROW(
          "'CUDAPlace' is not supported, Please re-compile with WITH_GPU "
          "option");
D
dzhwinter 已提交
115 116 117 118 119
#endif
    }
  }
}

120 121 122 123
CPUDeviceContext::CPUDeviceContext() {
  eigen_device_.reset(new Eigen::DefaultDevice());
}

D
dzhwinter 已提交
124
CPUDeviceContext::CPUDeviceContext(CPUPlace place) : place_(place) {
125 126 127 128 129 130 131
  eigen_device_.reset(new Eigen::DefaultDevice());
}

Eigen::DefaultDevice* CPUDeviceContext::eigen_device() const {
  return eigen_device_.get();
}

D
dzhwinter 已提交
132
Place CPUDeviceContext::GetPlace() const { return place_; }
133

134
#ifdef PADDLE_WITH_CUDA
135

Q
init  
qijun 已提交
136 137 138 139 140 141 142
class EigenCudaStreamDevice : public Eigen::StreamInterface {
 public:
  EigenCudaStreamDevice() : scratch_(nullptr), semaphore_(nullptr) {
    Eigen::initializeDeviceProp();
  }
  ~EigenCudaStreamDevice() override {}

D
dzhwinter 已提交
143
  void Reinitialize(const cudaStream_t* cuda_stream, CUDAPlace place) {
Q
init  
qijun 已提交
144 145 146 147 148 149 150 151 152 153 154 155
    stream_ = cuda_stream;
    place_ = place;
    device_prop_ = &Eigen::m_deviceProperties[place.device];
  }

  const cudaStream_t& stream() const override { return *stream_; }

  const cudaDeviceProp& deviceProperties() const override {
    return *device_prop_;
  }

  void* allocate(size_t num_bytes) const override {
S
sneaxiy 已提交
156 157 158
    if (UNLIKELY(num_bytes == 0)) {
      return nullptr;
    }
159 160 161
    auto buf = memory::Alloc(place_, num_bytes);
    VLOG(4) << "Eigen allocated at " << buf->ptr() << ", size" << buf->size()
            << " requested " << num_bytes;
162
    void* retv = buf->ptr();
S
sneaxiy 已提交
163 164 165 166
    {
      std::lock_guard<std::mutex> lock(mtx_);
      allocations_.emplace(retv, std::move(buf));
    }
167
    return retv;
Q
init  
qijun 已提交
168 169
  }

S
sneaxiy 已提交
170 171 172 173 174 175
  void deallocate(void* buffer) const override {
    if (LIKELY(buffer)) {
      std::lock_guard<std::mutex> lock(mtx_);
      allocations_.erase(buffer);
    }
  }
Q
init  
qijun 已提交
176 177 178

  void* scratchpad() const override {
    if (scratch_ == NULL) {
Z
Zhang Ting 已提交
179 180 181 182
// windows use an old version of eigen that uses kCudaScratchSize,
// once windows updates eigen to a recent version, the following code
// can use kGpuScratchSize uniformly
#ifdef _WIN32
Q
init  
qijun 已提交
183
      scratch_ = allocate(Eigen::kCudaScratchSize + sizeof(unsigned int));
Z
Zhang Ting 已提交
184 185 186
#else
      scratch_ = allocate(Eigen::kGpuScratchSize + sizeof(unsigned int));
#endif
Q
init  
qijun 已提交
187 188 189 190 191 192
    }
    return scratch_;
  }

  unsigned int* semaphore() const override {
    if (semaphore_ == NULL) {
Z
Zhang Ting 已提交
193
#ifdef _WIN32
Q
init  
qijun 已提交
194 195
      char* scratch =
          static_cast<char*>(scratchpad()) + Eigen::kCudaScratchSize;
Z
Zhang Ting 已提交
196 197 198
#else
      char* scratch = static_cast<char*>(scratchpad()) + Eigen::kGpuScratchSize;
#endif
Q
init  
qijun 已提交
199
      semaphore_ = reinterpret_cast<unsigned int*>(scratch);
200
      PADDLE_ENFORCE_CUDA_SUCCESS(
Q
init  
qijun 已提交
201 202 203 204 205 206
          cudaMemsetAsync(semaphore_, 0, sizeof(unsigned int), *stream_));
    }
    return semaphore_;
  }

 private:
D
dzhwinter 已提交
207
  CUDAPlace place_;
Q
init  
qijun 已提交
208 209
  const cudaStream_t* stream_;         // not owned;
  const cudaDeviceProp* device_prop_;  // not owned;
Q
qijun 已提交
210
  mutable void* scratch_;
Q
init  
qijun 已提交
211
  mutable unsigned int* semaphore_;
S
sneaxiy 已提交
212
  mutable std::mutex mtx_;  // to protect allocations_
Y
Yu Yang 已提交
213
  mutable std::unordered_map<void*, memory::AllocationPtr> allocations_;
Q
init  
qijun 已提交
214 215
};

216 217 218 219 220 221 222 223 224
void CudnnWorkspaceHandle::ReallocWorkspace(size_t required_workspace_bytes) {
  if (required_workspace_bytes <= WorkspaceSize()) {
    return;
  }
  // reset allocation first before re-allocate to save memory
  allocation_.reset();
  allocation_ = memory::Alloc(device_context_, required_workspace_bytes);
}

225 226 227 228 229 230 231 232 233 234 235 236
thread_local std::unordered_map<const CUDADeviceContext*,
                                std::shared_ptr<CUDAContext>>
    CUDADeviceContext::thread_ctx_;
thread_local std::mutex CUDADeviceContext::ctx_mtx_;

void CUDAContext::InitEigenContext() {
  eigen_stream_.reset(new EigenCudaStreamDevice());
  eigen_stream_->Reinitialize(&RawStream(), place_);
  eigen_device_.reset(new Eigen::GpuDevice(eigen_stream_.get()));
}

CUDAContext::CUDAContext(const CUDAPlace& place,
237
                         const stream::Priority& priority) {
238 239 240 241 242 243 244 245 246 247 248 249 250 251
  place_ = place;
  CUDADeviceGuard guard(place_.device);
  stream_.reset(new stream::CUDAStream(place, priority));
  InitEigenContext();
  InitCuBlasContext();
  InitCuDNNContext();
}

CUDAContext::~CUDAContext() {
  CUDADeviceGuard guard(place_.device);
  DestoryCuDNNContext();
  DestoryCuBlasContext();
}

252
CUDADeviceContext::CUDADeviceContext(CUDAPlace place) : place_(place) {
Y
Yu Yang 已提交
253
  CUDADeviceGuard guard(place_.device);
C
chengduo 已提交
254 255 256
  compute_capability_ = GetCUDAComputeCapability(place_.device);
  multi_process_ = GetCUDAMultiProcessors(place_.device);
  max_threads_per_mp_ = GetCUDAMaxThreadsPerMultiProcessor(place_.device);
257
  max_grid_dim_size_ = GetGpuMaxGridDimSize(place_.device);
258
  max_threads_per_block_ = GetCUDAMaxThreadsPerBlock(place_.device);
259

C
chengduo 已提交
260 261 262
  driver_version_ = GetCUDADriverVersion(place_.device);
  runtime_version_ = GetCUDARuntimeVersion(place_.device);

263 264
  LOG_FIRST_N(WARNING, 1) << "Please NOTE: device: " << place_.device
                          << ", CUDA Capability: " << compute_capability_
C
chengduo 已提交
265
                          << ", Driver API Version: " << driver_version_ / 1000
266
                          << "." << (driver_version_ % 100) / 10
C
chengduo 已提交
267 268 269
                          << ", Runtime API Version: "
                          << runtime_version_ / 1000 << "."
                          << (runtime_version_ % 100) / 10;
270 271 272
  size_t cudnn_dso_ver = dynload::cudnnGetVersion();
  LOG_FIRST_N(WARNING, 1) << "device: " << place_.device
                          << ", cuDNN Version: " << cudnn_dso_ver / 1000 << "."
273
                          << (cudnn_dso_ver % 1000) / 100 << ".";
S
sneaxiy 已提交
274 275 276

  {
    // Check CUDA/CUDNN version compatiblity
277 278 279 280
    auto local_cuda_version =
        (driver_version_ / 1000) * 10 + (driver_version_ % 100) / 10;
    auto compile_cuda_version =
        (CUDA_VERSION / 1000) * 10 + (CUDA_VERSION % 100) / 10;
S
sneaxiy 已提交
281 282 283 284 285 286 287 288 289 290 291 292
    if (local_cuda_version < compile_cuda_version) {
      LOG_FIRST_N(WARNING, 1)
          << "WARNING: device: " << place_.device
          << ". The installed Paddle is compiled with CUDA "
          << compile_cuda_version / 10 << "." << compile_cuda_version % 10
          << ", but CUDA runtime version in your machine is "
          << local_cuda_version / 10 << "." << local_cuda_version % 10
          << ", which may cause serious incompatible bug. "
          << "Please recompile or reinstall Paddle with compatible CUDA "
             "version.";
    }
  }
293
  default_ctx_.reset(new CUDAContext(place_));
294 295 296 297
}

CUDADeviceContext::~CUDADeviceContext() {
  SetDeviceId(place_.device);
298 299 300 301 302
#if defined(PADDLE_WITH_NCCL)
  if (nccl_comm_) {
    PADDLE_ENFORCE_CUDA_SUCCESS(dynload::ncclCommDestroy(nccl_comm_));
  }
#endif
303 304
}

L
liaogang 已提交
305
Place CUDADeviceContext::GetPlace() const { return place_; }
306

307
void CUDADeviceContext::Wait() const { context()->Stream()->Wait(); }
308

K
Kexin Zhao 已提交
309
int CUDADeviceContext::GetComputeCapability() const {
C
chengduo 已提交
310
  return compute_capability_;
K
Kexin Zhao 已提交
311 312
}

313
int CUDADeviceContext::GetMaxPhysicalThreadCount() const {
C
chengduo 已提交
314
  return multi_process_ * max_threads_per_mp_;
315 316
}

317 318 319 320 321 322
int CUDADeviceContext::GetSMCount() const { return multi_process_; }

int CUDADeviceContext::GetMaxThreadsPerBlock() const {
  return max_threads_per_block_;
}

323
Eigen::GpuDevice* CUDADeviceContext::eigen_device() const {
324
  return context()->EigenDevice().get();
325 326
}

327
bool CUDADeviceContext::tensor_core_available() const {
328
  return context()->CublasTensorCoreHandle() != nullptr;
S
sneaxiy 已提交
329 330
}

331 332 333 334
dim3 CUDADeviceContext::GetCUDAMaxGridDimSize() const {
  return max_grid_dim_size_;
}

335 336 337
cudnnHandle_t CUDADeviceContext::cudnn_handle() const {
  return context()->CudnnHandle();
}
338

S
sneaxiy 已提交
339
CudnnWorkspaceHandle CUDADeviceContext::cudnn_workspace_handle() const {
340
  return CudnnWorkspaceHandle(*this, &cudnn_handle_mtx_);
341
}
342

343 344 345
cudaStream_t CUDADeviceContext::stream() const {
  return context()->RawStream();
}
Q
qijun 已提交
346

C
chengduoZH 已提交
347 348 349 350 351 352 353 354 355 356 357 358 359 360
CUDAPinnedDeviceContext::CUDAPinnedDeviceContext() {
  eigen_device_.reset(new Eigen::DefaultDevice());
}

CUDAPinnedDeviceContext::CUDAPinnedDeviceContext(CUDAPinnedPlace place)
    : place_(place) {
  eigen_device_.reset(new Eigen::DefaultDevice());
}

Eigen::DefaultDevice* CUDAPinnedDeviceContext::eigen_device() const {
  return eigen_device_.get();
}

Place CUDAPinnedDeviceContext::GetPlace() const { return place_; }
L
Luo Tao 已提交
361
#endif
Q
qijun 已提交
362

T
tensor-tang 已提交
363 364
#ifdef PADDLE_WITH_MKLDNN
MKLDNNDeviceContext::MKLDNNDeviceContext(CPUPlace place)
A
Adam 已提交
365 366 367
    : CPUDeviceContext(place),
      engine_(mkldnn::engine::kind::cpu, 0),
      p_blobmap_() {
368 369
  p_blobmap_.reset(new BlobMap());
  p_mutex_.reset(new std::mutex());
T
tensor-tang 已提交
370 371
}

S
Sylwester Fraczek 已提交
372
namespace {
373 374
// Current mkldnn session id.
thread_local size_t cur_mkldnn_session_id = kMKLDNNSessionID_Default;
375 376 377 378
// Current data input shape string.
// - For fixed-shape, it's a null string in default.
// - For dynamic-shape, it's user specific.
thread_local std::string cur_input_shape_str = "";
379 380 381
// the cache capacity of different input shapes for MKLDNN.
// Default 1 means fixed input shape, not dynamic shape.
thread_local int cur_input_shape_cache_capacity = 1;
382 383 384 385
// Recently registered data_format. This is needed to
// know for converting MKL-DNN Tensor to non MKL-DNN
thread_local paddle::framework::DataLayout cur_paddle_data_layout =
    paddle::framework::DataLayout::kNCHW;
386
}  // namespace
S
Sylwester Fraczek 已提交
387

388 389
void set_cur_mkldnn_session_id(size_t sid) { cur_mkldnn_session_id = sid; }
size_t get_cur_mkldnn_session_id(void) { return cur_mkldnn_session_id; }
390 391 392
void set_cur_input_shape_str(std::string input_shape_str) {
  cur_input_shape_str = input_shape_str;
}
393 394 395
void set_cur_input_shape_cache_capacity(int input_shape_cache_capacity) {
  cur_input_shape_cache_capacity = input_shape_cache_capacity;
}
S
Sylwester Fraczek 已提交
396

397 398 399 400 401 402 403 404
void set_cur_paddle_data_layout(framework::DataLayout dl) {
  cur_paddle_data_layout = dl;
}

framework::DataLayout get_cur_paddle_data_layout(void) {
  return cur_paddle_data_layout;
}

405 406
void MKLDNNDeviceContext::ResetBlobMap() const { p_blobmap_->clear(); }

407 408 409 410 411 412 413 414 415 416 417
size_t MKLDNNDeviceContext::GetShapeBlobSize() const {
  std::lock_guard<std::mutex> lock(*p_mutex_);
  BlobMap* pMap = p_blobmap_.get();
  auto map_it = pMap->find(cur_mkldnn_session_id);
  if (map_it == pMap->end()) {
    LOG(FATAL) << "MKLDNNDeviceContext don't find cur_mkldnn_session_id : "
               << cur_mkldnn_session_id;
  }
  return map_it->second->size();
}

418 419
void MKLDNNDeviceContext::SetBlob(const std::string& name,
                                  std::shared_ptr<void> data) const {
420
  BlobMap* pMap = p_blobmap_.get();
421
  std::shared_ptr<ShapeBlob> sBlob = nullptr;
422 423
  std::shared_ptr<KeyBlob> pBlob = nullptr;

424
  int sid = platform::get_cur_mkldnn_session_id();
T
tensor-tang 已提交
425

426
  std::lock_guard<std::mutex> lock(*p_mutex_);
T
tensor-tang 已提交
427

428 429
  // Find ShapeBlob for current mkldnn session id.
  auto map_it = pMap->find(sid);
430 431 432

  if (map_it == pMap->end()) {
    // 1st time to set blob in current thread
433
    sBlob = std::shared_ptr<ShapeBlob>(new ShapeBlob());
434 435
    (*pMap)[sid] = sBlob;
    VLOG(2) << "SetBlob: sid=" << sid << ", add new sid\n";
436
  } else {
437
    sBlob = map_it->second;
438
  }
T
tensor-tang 已提交
439

440 441
  // Find KeyBlob for current input shape
  auto key_it = sBlob->find(cur_input_shape_str);
442

443
  if (key_it == sBlob->end()) {
444 445
    // In cache clearing mode, cur_input_shape_cache_capacity defines
    // max pblob capacity
446 447
    if ((static_cast<size_t>(sid) == kMKLDNNSessionID_CacheClearing) &&
        sBlob->size() &&
448 449 450 451 452 453
        (sBlob->size() >=
         static_cast<size_t>(cur_input_shape_cache_capacity))) {
      VLOG(2) << "sid=" << sid
              << ", remove all blobs of shape: " << sBlob->begin()->first;
      sBlob->erase(sBlob->begin()->first);
    }
454 455
    pBlob = std::shared_ptr<KeyBlob>(new KeyBlob());
    (*sBlob)[cur_input_shape_str] = pBlob;
456
  } else {
457
    pBlob = key_it->second;
458 459
  }

460 461 462 463 464 465 466
  // Find Blob via name
  auto blob_it = pBlob->find(name);
  if (blob_it == pBlob->end()) {
    (*pBlob)[name] = data;
  } else {
    blob_it->second = data;  // set data to existing blob
  }
467
  VLOG(2) << "SetBlob: sid=" << sid << ", add blob=" << name << "\n";
468
  // lock will be automatically released when out of scope
469
  return;
T
tensor-tang 已提交
470 471
}

472 473
std::shared_ptr<void> MKLDNNDeviceContext::GetBlob(
    const std::string& name) const {
474
  BlobMap* pMap = p_blobmap_.get();
475
  std::shared_ptr<ShapeBlob> sBlob = nullptr;
476
  std::shared_ptr<KeyBlob> pBlob = nullptr;
T
tensor-tang 已提交
477

478
  int sid = platform::get_cur_mkldnn_session_id();
T
tensor-tang 已提交
479

480
  std::lock_guard<std::mutex> lock(*p_mutex_);
481

482 483
  // Find ShapeBlob for current mkldnn session id firstly
  auto map_it = pMap->find(sid);
484
  if (map_it == pMap->end()) {
485
    VLOG(2) << "GetBlob: sid=" << sid << ", miss sid\n";
486 487 488 489 490 491 492
    return nullptr;
  }
  sBlob = map_it->second;

  // Find KeyBlob for current input shape secondly
  auto sBlob_it = sBlob->find(cur_input_shape_str);
  if (sBlob_it == sBlob->end()) {
493
    VLOG(2) << "GetBlob: sid=" << cur_input_shape_str
494 495 496 497
            << ", miss input_shape_str\n";
    return nullptr;
  }
  pBlob = sBlob_it->second;
498 499 500 501

  // Find Blob via name
  auto key_it = pBlob->find(name);

502
  if (key_it == pBlob->end()) {
503
    VLOG(2) << "GetBlob sid=" << sid << ", miss blob=" << name << "\n";
504 505
    return nullptr;
  }
506

507
  VLOG(2) << "GetBlob sid=" << sid << ", get blob=" << name << "\n";
508 509
  // lock will be automatically released when out of scope
  return key_it->second;
T
tensor-tang 已提交
510 511 512 513
}

#endif

Q
qijun 已提交
514
}  // namespace platform
Q
qijun 已提交
515
}  // namespace paddle