hl_cuda_device.cc 20.3 KB
Newer Older
1
/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserved.
Z
zhangjinchao01 已提交
2 3 4 5 6 7 8 9 10 11 12 13 14

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. */

Y
Yu Yang 已提交
15
// clang-format off
16
// Because clang-format 4.X and clang-format 3.8+ format
Y
Yu Yang 已提交
17
// following lines in different. So disable clang-format.
Y
Yu Yang 已提交
18
#include "hl_cuda.h"
L
liaogang 已提交
19
#include <cuda_profiler_api.h>
Z
zhangjinchao01 已提交
20 21
#include <string.h>
#include <sys/syscall.h>
L
liaogang 已提交
22 23
#include <sys/time.h>
#include <unistd.h>
Z
zhangjinchao01 已提交
24
#include "hl_cuda.ph"
L
liaogang 已提交
25
#include "hl_thread.ph"
X
Xin Pan 已提交
26 27
#include "paddle/legacy/utils/Logging.h"
#include "paddle/legacy/utils/DynamicLoader.h"
Y
Yu Yang 已提交
28
// clang-format on
Z
zhangjinchao01 已提交
29 30 31 32

namespace dynload {

std::once_flag curand_dso_flag;
33
void *curand_dso_handle = nullptr;
Z
zhangjinchao01 已提交
34 35 36 37 38 39 40 41 42

/**
 * The following macro definition can generate structs
 * (for each function) to dynamic load curand routine
 * via operator overloading.
 *
 * note: default dynamic linked libs
 */
#ifdef PADDLE_USE_DSO
43 44 45 46 47 48 49 50 51 52
#define DYNAMIC_LOAD_CURAND_WRAP(__name)                                       \
  struct DynLoad__##__name {                                                   \
    template <typename... Args>                                                \
    curandStatus_t operator()(Args... args) {                                  \
      typedef curandStatus_t (*curandFunc)(Args...);                           \
      std::call_once(curand_dso_flag, GetCurandDsoHandle, &curand_dso_handle); \
      void *p_##__name = dlsym(curand_dso_handle, #__name);                    \
      return reinterpret_cast<curandFunc>(p_##__name)(args...);                \
    }                                                                          \
  } __name; /* struct DynLoad__##__name */
Z
zhangjinchao01 已提交
53
#else
54 55 56 57 58 59 60
#define DYNAMIC_LOAD_CURAND_WRAP(__name)      \
  struct DynLoad__##__name {                  \
    template <typename... Args>               \
    curandStatus_t operator()(Args... args) { \
      return __name(args...);                 \
    }                                         \
  } __name; /* struct DynLoad__##__name */
Z
zhangjinchao01 已提交
61 62 63
#endif

/* include all needed curand functions in HPPL */
L
Luo Tao 已提交
64
// clang-format off
Z
zhangjinchao01 已提交
65 66 67 68 69 70
#define CURAND_RAND_ROUTINE_EACH(__macro)    \
  __macro(curandCreateGenerator)             \
  __macro(curandSetStream)                   \
  __macro(curandSetPseudoRandomGeneratorSeed)\
  __macro(curandGenerateUniform)             \
  __macro(curandGenerateUniformDouble)
L
Luo Tao 已提交
71
// clang-format on
Z
zhangjinchao01 已提交
72 73 74 75 76 77

CURAND_RAND_ROUTINE_EACH(DYNAMIC_LOAD_CURAND_WRAP)

#undef CURAND_RAND_ROUTINE_EACH
#undef DYNAMIC_LOAD_CURAND_WRAP

78
} /* namespace dynload */
Z
zhangjinchao01 已提交
79 80 81 82

/**
 * @brief   global resource.
 */
83 84 85 86
int g_system_device_num = 0;                /* system device number */
int device_num = 0;                         /* use    device number */
hl_device_prop *g_device;                   /* device info table */
__thread thread_device_resources *t_device; /* device resources table */
Z
zhangjinchao01 已提交
87 88 89
int g_cuda_lib_version = 0;

/* number of global stream */
90
#define NUMBER_OF_GLOBAL_STREAM (HPPL_THREAD_STREAM_1)
Z
zhangjinchao01 已提交
91
/* number of thread stream */
92
#define NUMBER_OF_THREAD_STREAM (HPPL_STREAM_END - HPPL_THREAD_STREAM_1)
Z
zhangjinchao01 已提交
93
/* sizeof of device memory */
94
#define HPPL_GPU_MEMORY_SIZE (256 * 4)
Z
zhangjinchao01 已提交
95 96

/**
97
 * Check build-in cuda function using glog and it **does not**
Z
zhangjinchao01 已提交
98 99
 * support << operator for more details error info.
 */
L
liaogang 已提交
100 101 102 103 104
#define CHECK_CUDA(cudaFunc)                                         \
  do {                                                               \
    cudaError_t cudaStat = cudaFunc;                                 \
    CHECK_EQ(cudaSuccess, cudaStat) << "Cuda Error: "                \
                                    << cudaGetErrorString(cudaStat); \
105
  } while (0)
Z
zhangjinchao01 已提交
106 107 108 109

/**
 * @brief   thread resource.
 */
110 111 112 113 114 115 116 117 118 119 120 121
__thread _hl_thread_resource t_resource = {{0},    /* stream */
                                           0,      /* handle */
                                           0,      /* gen */
                                           0,      /* cudnn_handle */
                                           0,      /* cudnn_desc */
                                           NULL,   /* gen_mutex */
                                           NULL,   /* gpu_mem */
                                           NULL,   /* cpu_mem */
                                           0,      /* event */
                                           -1,     /* device */
                                           0,      /* major */
                                           false}; /* is_init */
Z
zhangjinchao01 已提交
122 123 124 125 126

__thread cudaStream_t default_stream = 0;
__thread bool g_sync_flag = true;
bool hl_start_flag = false;

L
liaogang 已提交
127 128
inline pid_t gettid() {
#if defined(__APPLE__) || defined(__OSX__)
G
gangliao 已提交
129 130 131 132 133
  // syscall is deprecated: first deprecated in macOS 10.12.
  // syscall is unsupported;
  // syscall pid_t tid = syscall(SYS_thread_selfid);
  uint64_t tid;
  pthread_threadid_np(NULL, &tid);
L
liaogang 已提交
134
#else
135 136 137
#ifndef __NR_gettid
#define __NR_gettid 224
#endif
L
liaogang 已提交
138 139
  pid_t tid = syscall(__NR_gettid);
#endif
140 141
  CHECK_NE((int)tid, -1);
  return tid;
L
liaogang 已提交
142
}
Z
zhangjinchao01 已提交
143 144

void hl_init(int device) {
145
  CHECK(hl_start_flag) << "[Init failed] hl_start() did not succeed.";
Z
zhangjinchao01 已提交
146 147 148 149 150 151 152 153 154 155

  /* thread has been initialized */
  if (true == t_resource.is_init) {
    hl_set_device(device);
    return;
  }

  /* create thread devcie resources */
  char *tmp;
  thread_device_resources device_res;
156 157
  tmp = (char *)malloc(g_system_device_num * sizeof(thread_device_resources *) +
                       device_num * sizeof(_thread_device_resources));
Z
zhangjinchao01 已提交
158
  CHECK_NOTNULL(tmp);
159 160 161 162
  t_device = (thread_device_resources *)tmp;
  device_res = (thread_device_resources)(
      (char *)tmp + g_system_device_num * sizeof(thread_device_resources *));
  memset(t_device, 0, g_system_device_num * sizeof(thread_device_resources *));
Z
zhangjinchao01 已提交
163

164 165
  char *tmp_stream = (char *)malloc(device_num * NUMBER_OF_THREAD_STREAM *
                                    sizeof(cudaStream_t));
Z
zhangjinchao01 已提交
166 167 168 169 170 171 172 173 174
  CHECK_NOTNULL(tmp_stream);

  int num = 0;
  for (int dev = 0; dev < g_system_device_num; dev++) {
    if (!g_device[dev]) {
      continue;
    }

    t_device[dev] = &device_res[num];
175 176 177
    t_device[dev]->stream =
        (cudaStream_t *)(tmp_stream +
                         num * NUMBER_OF_THREAD_STREAM * sizeof(cudaStream_t));
Z
zhangjinchao01 已提交
178 179 180 181 182 183 184 185 186 187 188 189 190 191 192 193 194 195 196 197 198 199 200 201 202

    hl_create_thread_resources(dev, t_device[dev]);
    num++;
  }

  hl_cudnn_desc_init(&t_resource.cudnn_desc);

  /* thread initialization is complete */
  t_resource.is_init = true;
  /* set device */
  t_resource.device = -1;
  hl_set_device(device);
}

void hl_fini() {
  if (false == t_resource.is_init) {
    return;
  }

  /* hppl stream fini */
  t_resource.device = -1;
  for (int i = NUMBER_OF_GLOBAL_STREAM; i < HPPL_STREAM_END; i++) {
    t_resource.stream[i] = 0;
  }

203 204
  char *tmp = (char *)t_device;
  char *tmp_stream = NULL;
Z
zhangjinchao01 已提交
205 206 207 208 209
  for (int dev = 0; dev < g_system_device_num; dev++) {
    if (!t_device[dev]) {
      continue;
    }
    if (!tmp_stream) {
210
      tmp_stream = (char *)t_device[dev]->stream;
Z
zhangjinchao01 已提交
211 212
    }
    for (int j = 0; j < NUMBER_OF_THREAD_STREAM; j++) {
L
liaogang 已提交
213
      CHECK_CUDA(cudaStreamDestroy(t_device[dev]->stream[j]));
Z
zhangjinchao01 已提交
214 215 216 217 218
    }

    /* free device memory */
    hl_free_mem_device(t_device[dev]->gpu_mem);
    hl_free_mem_host(t_device[dev]->cpu_mem);
L
liaogang 已提交
219
    CHECK_CUDA(cudaEventDestroy(t_device[dev]->mem_event));
Z
zhangjinchao01 已提交
220 221 222 223 224 225 226
  }

  free(tmp);
  free(tmp_stream);
  t_resource.is_init = false;
}

227
int hl_get_device_count() { return device_num; }
Z
zhangjinchao01 已提交
228 229 230 231 232 233 234

void hl_set_device(int device) {
  if (device == t_resource.device) {
    return;
  }

  CHECK(device >= 0 && device < g_system_device_num && g_device[device])
235
      << "Device: " << device << " is not specified in startup.";
Z
zhangjinchao01 已提交
236

L
liaogang 已提交
237
  CHECK_CUDA(cudaSetDevice(device));
Z
zhangjinchao01 已提交
238 239 240 241 242 243 244 245 246

  /* switch thread stream */
  for (int i = 0; i < NUMBER_OF_GLOBAL_STREAM; i++) {
    t_resource.stream[i] = g_device[device]->device_resources->stream[i];
  }

  if (true == t_resource.is_init) {
    for (int i = NUMBER_OF_GLOBAL_STREAM; i < HPPL_STREAM_END; i++) {
      t_resource.stream[i] =
247
          t_device[device]->stream[i - NUMBER_OF_GLOBAL_STREAM];
Z
zhangjinchao01 已提交
248 249 250
    }
    t_resource.gpu_mem = t_device[device]->gpu_mem;
    t_resource.cpu_mem = t_device[device]->cpu_mem;
251
    t_resource.event = t_device[device]->mem_event;
Z
zhangjinchao01 已提交
252 253 254 255 256 257 258 259 260 261 262 263 264
  }

  t_resource.handle = g_device[device]->device_resources->handle;
  t_resource.gen = g_device[device]->device_resources->gen;
  t_resource.cudnn_handle = g_device[device]->device_resources->cudnn_handle;
  t_resource.gen_mutex = g_device[device]->device_resources->gen_mutex;
  t_resource.device = device;
  t_resource.major = g_device[device]->major;
  default_stream = t_resource.stream[0];
}

int hl_get_device() {
  int device;
L
liaogang 已提交
265
  CHECK_CUDA(cudaGetDevice(&device));
Z
zhangjinchao01 已提交
266 267 268
  return device;
}

269
void *hl_malloc_device(size_t size) {
Z
zhangjinchao01 已提交
270 271 272
  void *dest_d;

  CHECK(size) << __func__ << ": the size for device memory is 0, please check.";
L
liaogang 已提交
273
  CHECK_CUDA(cudaMalloc((void **)&dest_d, size));
Z
zhangjinchao01 已提交
274 275 276 277 278 279 280

  return dest_d;
}

void hl_free_mem_device(void *dest_d) {
  CHECK_NOTNULL(dest_d);

L
liaogang 已提交
281
  cudaError_t err = cudaFree(dest_d);
Z
zhangjinchao01 已提交
282
  CHECK(cudaSuccess == err || cudaErrorCudartUnloading == err)
283
      << hl_get_device_error_string();
Z
zhangjinchao01 已提交
284 285
}

286
void *hl_malloc_host(size_t size) {
Z
zhangjinchao01 已提交
287 288 289
  void *dest_h;

  CHECK(size) << __func__ << ": the size for device memory is 0, please check.";
L
liaogang 已提交
290
  CHECK_CUDA(cudaHostAlloc((void **)&dest_h, size, cudaHostAllocDefault));
Z
zhangjinchao01 已提交
291 292 293 294 295 296 297

  return dest_h;
}

void hl_free_mem_host(void *dest_h) {
  CHECK_NOTNULL(dest_h);

L
liaogang 已提交
298
  cudaError_t err = cudaFreeHost(dest_h);
299
  CHECK(cudaSuccess == err || cudaErrorCudartUnloading == err)
300
      << hl_get_device_error_string();
Z
zhangjinchao01 已提交
301 302 303 304 305 306 307 308
}

void hl_memcpy(void *dst, void *src, size_t size) {
  if (0 == size) {
    return;
  }
  CHECK_NOTNULL(dst);
  CHECK_NOTNULL(src);
L
liaogang 已提交
309
  CHECK_CUDA(cudaMemcpy(dst, src, size, cudaMemcpyDefault));
Z
zhangjinchao01 已提交
310 311 312
}

void hl_memset_device(void *dest_d, int value, size_t size) {
L
liaogang 已提交
313
  CHECK_CUDA(cudaMemset(dest_d, value, size));
Z
zhangjinchao01 已提交
314 315 316 317 318 319 320 321
}

void hl_memcpy_host2device(void *dest_d, void *src_h, size_t size) {
  if (0 == size) {
    return;
  }
  CHECK_NOTNULL(src_h);
  CHECK_NOTNULL(dest_d);
L
liaogang 已提交
322
  CHECK_CUDA(cudaMemcpy(dest_d, src_h, size, cudaMemcpyHostToDevice));
Z
zhangjinchao01 已提交
323 324 325 326 327 328 329 330
}

void hl_memcpy_device2host(void *dest_h, void *src_d, size_t size) {
  if (0 == size) {
    return;
  }
  CHECK_NOTNULL(dest_h);
  CHECK_NOTNULL(src_d);
L
liaogang 已提交
331
  CHECK_CUDA(cudaMemcpy(dest_h, src_d, size, cudaMemcpyDeviceToHost));
Z
zhangjinchao01 已提交
332 333 334 335 336 337 338 339
}

void hl_memcpy_device2device(void *dest_d, void *src_d, size_t size) {
  if (0 == size) {
    return;
  }
  CHECK_NOTNULL(dest_d);
  CHECK_NOTNULL(src_d);
L
liaogang 已提交
340
  CHECK_CUDA(cudaMemcpy(dest_d, src_d, size, cudaMemcpyDeviceToDevice));
Z
zhangjinchao01 已提交
341 342 343 344 345 346 347 348 349 350 351 352 353
}

void hl_memcpy_async(void *dst, void *src, size_t size, hl_stream_t stream) {
  cudaStream_t cu_stream;

  if (0 == size) {
    return;
  }
  CHECK_NOTNULL(dst);
  CHECK_NOTNULL(src);
  CHECK_LT(stream, HPPL_STREAM_END);
  cu_stream = t_resource.stream[stream];

L
liaogang 已提交
354
  CHECK_CUDA(cudaMemcpyAsync(dst, src, size, cudaMemcpyDefault, cu_stream));
Z
zhangjinchao01 已提交
355 356 357 358 359 360 361 362 363 364
}

void hl_start() {
  hl_specify_devices_start(NULL, 0);
  /* set default device */
  hl_set_device(0);
}

bool hl_device_can_access_peer(int device, int peerDevice) {
  int canAccessPeer;
L
liaogang 已提交
365
  CHECK_CUDA(cudaDeviceCanAccessPeer(&canAccessPeer, device, peerDevice));
Z
zhangjinchao01 已提交
366 367 368 369 370 371 372 373 374

  if (canAccessPeer == 1) {
    return true;
  } else {
    return false;
  }
}

void hl_device_enable_peer_access(int peerDevice) {
L
liaogang 已提交
375
  cudaError_t err = cudaDeviceEnablePeerAccess(peerDevice, 0);
Z
zhangjinchao01 已提交
376
  if (cudaErrorPeerAccessAlreadyEnabled == err) {
L
liaogang 已提交
377
    cudaGetLastError();
Z
zhangjinchao01 已提交
378 379 380 381 382 383 384 385 386 387
  } else {
    CHECK_CUDA(err);
  }
}

void hl_create_global_resources(hl_device_prop device_prop) {
  struct cudaDeviceProp cu_prop;
  int device = device_prop->device;
  global_device_resources device_res = device_prop->device_resources;

L
liaogang 已提交
388
  CHECK_CUDA(cudaSetDevice(device));
Z
zhangjinchao01 已提交
389
  /* device properties */
L
liaogang 已提交
390
  CHECK_CUDA(cudaGetDeviceProperties(&cu_prop, device));
Z
zhangjinchao01 已提交
391 392 393 394 395 396 397 398

  device_prop->major = cu_prop.major;
  device_prop->minor = cu_prop.minor;
  strncpy(device_prop->device_name, cu_prop.name, 256);
  device_prop->device_mem = cu_prop.totalGlobalMem;

  /* create device stream */
  for (int j = 0; j < NUMBER_OF_GLOBAL_STREAM; j++) {
L
liaogang 已提交
399
    CHECK_CUDA(cudaStreamCreate(&device_res->stream[j]));
Z
zhangjinchao01 已提交
400 401 402 403 404 405 406
  }

  /* cublas init */
  hl_cublas_init(&device_res->handle, device_res->stream[0]);

  /* create curand gen */
  CHECK_EQ(dynload::curandCreateGenerator(&device_res->gen,
407 408 409
                                          CURAND_RNG_PSEUDO_DEFAULT),
           CURAND_STATUS_SUCCESS)
      << "[Start failed] Curand init failed.";
Z
zhangjinchao01 已提交
410

411 412 413
  CHECK_EQ(dynload::curandSetStream(device_res->gen, device_res->stream[0]),
           CURAND_STATUS_SUCCESS)
      << "[Start failed] Curand set stream failed!";
Z
zhangjinchao01 已提交
414 415 416 417 418

  /* create cudnn handle */
  hl_cudnn_init(&device_res->cudnn_handle, device_res->stream[0]);

  int seed = gettid();
419 420 421
  CHECK_EQ(dynload::curandSetPseudoRandomGeneratorSeed(device_res->gen,
                                                       seed + device),
           CURAND_STATUS_SUCCESS);
Z
zhangjinchao01 已提交
422

423
  device_res->gen_mutex = (pthread_mutex_t *)(malloc(sizeof(pthread_mutex_t)));
Z
zhangjinchao01 已提交
424 425
  pthread_mutex_init(device_res->gen_mutex, NULL);

L
liaogang 已提交
426
  CHECK_CUDA(cudaRuntimeGetVersion(&g_cuda_lib_version));
Z
zhangjinchao01 已提交
427 428
}

429
int hl_get_cuda_version() { return g_cuda_lib_version; }
Z
zhangjinchao01 已提交
430

431
void hl_create_thread_resources(int device,
432
                                thread_device_resources device_res) {
L
liaogang 已提交
433
  CHECK_CUDA(cudaSetDevice(device));
Z
zhangjinchao01 已提交
434 435 436

  /* create thread stream */
  for (int j = 0; j < NUMBER_OF_THREAD_STREAM; j++) {
L
liaogang 已提交
437
    CHECK_CUDA(cudaStreamCreate(&device_res->stream[j]));
Z
zhangjinchao01 已提交
438 439 440
  }

  /* allocation device memory */
441
  device_res->gpu_mem = (real *)hl_malloc_device(HPPL_GPU_MEMORY_SIZE);
Z
zhangjinchao01 已提交
442 443

  /* allocation host memory */
444
  device_res->cpu_mem = (real *)hl_malloc_host(HPPL_GPU_MEMORY_SIZE);
Z
zhangjinchao01 已提交
445

L
liaogang 已提交
446
  CHECK_CUDA(cudaEventCreate(&device_res->mem_event));
Z
zhangjinchao01 已提交
447 448
}

449
void hl_specify_devices_start(int *device, int number) {
Z
zhangjinchao01 已提交
450 451 452
  if (hl_start_flag) return;

  /* 1. get the number of devices */
L
liaogang 已提交
453
  CHECK_CUDA(cudaGetDeviceCount(&g_system_device_num));
Z
zhangjinchao01 已提交
454 455 456 457 458 459 460
  CHECK_NE(g_system_device_num, 0) << "[Start failed] there is no GPU device";
  if (device == NULL) {
    number = g_system_device_num;
  }

  /* 2. check device & create device property table */
  CHECK_LE(number, g_system_device_num)
461 462
      << "[Start failed] System does not have enough device. "
      << "Device number: " << g_system_device_num << "Input number: " << number;
Z
zhangjinchao01 已提交
463 464 465

  char *tmp;
  hl_device_prop device_prop;
466 467
  tmp = (char *)malloc(g_system_device_num * sizeof(hl_device_prop *) +
                       number * sizeof(_hl_device_prop));
Z
zhangjinchao01 已提交
468 469
  CHECK(tmp) << "[Start failed] System memory is not enough.";

470
  g_device = (hl_device_prop *)tmp;
L
liaogang 已提交
471 472
  device_prop = (hl_device_prop)(
      (char *)tmp + g_system_device_num * sizeof(hl_device_prop *));
473
  memset(g_device, 0, g_system_device_num * sizeof(hl_device_prop *));
Z
zhangjinchao01 已提交
474 475 476 477 478 479 480 481 482 483
  int num = 0;
  for (int i = 0; i < number; i++) {
    int dev;
    if (device == NULL) {
      dev = i;
    } else {
      dev = device[i];
    }

    CHECK_LT(dev, g_system_device_num)
484 485 486
        << "[Start failed] The specified device number is "
        << "out of range. Max device number: " << g_system_device_num - 1
        << " Specified devcie number: " << dev;
Z
zhangjinchao01 已提交
487 488 489

    if (g_device[dev]) {
      /* Warning */
490
      LOG(WARNING) << "[Warning] Repeat specify device: " << dev;
Z
zhangjinchao01 已提交
491 492 493 494 495 496 497 498 499 500
      continue;
    }

    g_device[dev] = &device_prop[num];
    g_device[dev]->device = dev;
    num++;
  }
  device_num = num;

  /* 3.  create global device resources */
501
  char *tmp_res = (char *)malloc(device_num * sizeof(_global_device_resources));
Z
zhangjinchao01 已提交
502 503
  CHECK_NOTNULL(tmp_res);

504 505
  char *tmp_stream = (char *)malloc(device_num * NUMBER_OF_GLOBAL_STREAM *
                                    sizeof(cudaStream_t));
Z
zhangjinchao01 已提交
506 507 508 509 510 511 512 513
  CHECK_NOTNULL(tmp_stream);

  num = 0;
  for (int i = 0; i < g_system_device_num; i++) {
    if (!g_device[i]) {
      continue;
    }

514 515 516 517 518
    g_device[i]->device_resources = (global_device_resources)(
        tmp_res + num * sizeof(_global_device_resources));
    g_device[i]->device_resources->stream =
        (cudaStream_t *)(tmp_stream +
                         num * NUMBER_OF_GLOBAL_STREAM * sizeof(cudaStream_t));
Z
zhangjinchao01 已提交
519 520 521 522 523 524 525 526 527

    hl_create_global_resources(g_device[i]);
    num++;
  }

  /* hl_start() is ok */
  hl_start_flag = true;
  /* set default device */
  if (device == NULL) {
528
    hl_set_device(0);
Z
zhangjinchao01 已提交
529
  } else {
530
    hl_set_device(device[0]);
Z
zhangjinchao01 已提交
531 532 533 534 535 536
  }
}

void hl_rand(real *dest_d, size_t num) {
  pthread_mutex_lock(t_resource.gen_mutex);
  CHECK_EQ(
537
#ifndef PADDLE_TYPE_DOUBLE
538
      dynload::curandGenerateUniform(t_resource.gen, dest_d, num),
Z
zhangjinchao01 已提交
539
#else
540
      dynload::curandGenerateUniformDouble(t_resource.gen, dest_d, num),
Z
zhangjinchao01 已提交
541
#endif
542
      CURAND_STATUS_SUCCESS);
Z
zhangjinchao01 已提交
543 544 545 546 547 548
  pthread_mutex_unlock(t_resource.gen_mutex);
  CHECK_SYNC("hl_rand failed");
}

void hl_srand(unsigned int seed) {
  pthread_mutex_lock(t_resource.gen_mutex);
549 550
  CHECK_EQ(dynload::curandSetPseudoRandomGeneratorSeed(t_resource.gen, seed),
           CURAND_STATUS_SUCCESS);
Z
zhangjinchao01 已提交
551 552 553
  pthread_mutex_unlock(t_resource.gen_mutex);
}

554
void hl_set_sync_flag(bool flag) { g_sync_flag = flag; }
Z
zhangjinchao01 已提交
555

556
bool hl_get_sync_flag() { return g_sync_flag; }
Z
zhangjinchao01 已提交
557 558 559 560

void hl_stream_synchronize(hl_stream_t stream) {
  cudaStream_t cu_stream;

L
liaogang 已提交
561 562
  CHECK_LT(stream, HPPL_STREAM_END) << __func__
                                    << ": the parameter stream is error.";
Z
zhangjinchao01 已提交
563 564

  cu_stream = t_resource.stream[stream];
L
liaogang 已提交
565
  CHECK_CUDA(cudaStreamSynchronize(cu_stream));
Z
zhangjinchao01 已提交
566 567 568 569 570
}

void hl_create_event(hl_event_t *event) {
  CHECK_NOTNULL(event);

571 572
  struct _hl_event_st *st_event =
      (struct _hl_event_st *)malloc(sizeof(struct _hl_event_st));
Z
zhangjinchao01 已提交
573

L
liaogang 已提交
574
  CHECK_CUDA(cudaEventCreate(&st_event->cu_event));
Z
zhangjinchao01 已提交
575 576 577 578 579 580 581 582 583

  *event = st_event;
}

float hl_event_elapsed_time(hl_event_t start, hl_event_t end) {
  float time;
  CHECK_NOTNULL(start);
  CHECK_NOTNULL(end);

L
liaogang 已提交
584
  CHECK_CUDA(cudaEventElapsedTime(&time, start->cu_event, end->cu_event));
Z
zhangjinchao01 已提交
585 586 587 588 589 590 591
  return time;
}

void hl_stream_record_event(hl_stream_t stream, hl_event_t event) {
  cudaStream_t cu_stream;

  CHECK_NOTNULL(event);
L
liaogang 已提交
592 593
  CHECK_LT(stream, HPPL_STREAM_END) << __func__
                                    << ": the parameter stream is error.";
Z
zhangjinchao01 已提交
594 595

  cu_stream = t_resource.stream[stream];
L
liaogang 已提交
596
  CHECK_CUDA(cudaEventRecord(event->cu_event, cu_stream));
Z
zhangjinchao01 已提交
597 598 599 600 601 602
}

void hl_stream_wait_event(hl_stream_t stream, hl_event_t event) {
  cudaStream_t cu_stream;

  CHECK_NOTNULL(event);
L
liaogang 已提交
603 604
  CHECK_LT(stream, HPPL_STREAM_END) << __func__
                                    << ": the parameter stream is error.";
Z
zhangjinchao01 已提交
605 606

  cu_stream = t_resource.stream[stream];
L
liaogang 已提交
607
  CHECK_CUDA(cudaStreamWaitEvent(cu_stream, event->cu_event, 0));
Z
zhangjinchao01 已提交
608 609 610 611
}

void hl_destroy_event(hl_event_t event) {
  CHECK_NOTNULL(event);
L
liaogang 已提交
612
  CHECK_CUDA(cudaEventDestroy(event->cu_event));
Z
zhangjinchao01 已提交
613 614 615 616 617 618 619

  free(event);
  event = NULL;
}

void hl_event_synchronize(hl_event_t event) {
  CHECK_NOTNULL(event);
L
liaogang 已提交
620
  CHECK_CUDA(cudaEventSynchronize(event->cu_event));
Z
zhangjinchao01 已提交
621 622 623 624 625
}

void hl_get_device_name(char *name, int len, int device) {
  CHECK_NOTNULL(name);
  CHECK(device >= 0 && device < g_system_device_num && g_device[device])
626
      << "Device(" << device << ") is not specified in startup.";
Z
zhangjinchao01 已提交
627

628
  strncpy(name, g_device[device]->device_name, len);
Z
zhangjinchao01 已提交
629 630 631 632 633
}

void hl_get_device_memory(size_t *mem_size, int device) {
  CHECK_NOTNULL(mem_size);
  CHECK(device >= 0 && device < g_system_device_num && g_device[device])
634
      << "Device(" << device << ") is not specified in startup.";
Z
zhangjinchao01 已提交
635 636 637 638 639 640 641 642

  *mem_size = g_device[device]->device_mem;
}

void hl_get_device_compute_capability(int *major, int *minor, int device) {
  CHECK_NOTNULL(major);
  CHECK_NOTNULL(minor);
  CHECK(device >= 0 && device < g_system_device_num && g_device[device])
643
      << "Device(" << device << ") is not specified in startup.";
Z
zhangjinchao01 已提交
644 645 646 647 648

  *major = g_device[device]->major;
  *minor = g_device[device]->minor;
}

L
liaogang 已提交
649
int hl_get_device_last_error() { return (int)cudaGetLastError(); }
Z
zhangjinchao01 已提交
650

651
const char *hl_get_device_error_string() {
L
liaogang 已提交
652 653
  cudaError_t err = cudaGetLastError();
  return cudaGetErrorString(err);
Z
zhangjinchao01 已提交
654 655
}

656
const char *hl_get_device_error_string(size_t err) {
L
liaogang 已提交
657
  return cudaGetErrorString((cudaError_t)err);
Z
zhangjinchao01 已提交
658 659
}

L
liaogang 已提交
660
void hl_device_synchronize() { CHECK_CUDA(cudaDeviceSynchronize()); }
Z
zhangjinchao01 已提交
661
void hl_set_device_flags_block() {
L
liaogang 已提交
662
  CHECK_CUDA(cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync));
Z
zhangjinchao01 已提交
663 664
}

L
liaogang 已提交
665
bool hl_cuda_event_is_ready(hl_event_t event) {
L
liaogang 已提交
666
  cudaError_t err = cudaEventQuery(event->cu_event);
Z
zhangjinchao01 已提交
667 668 669
  CHECK(cudaSuccess == err || cudaErrorNotReady == err);

  if (cudaErrorNotReady == err) {
L
liaogang 已提交
670
    return false;
Z
zhangjinchao01 已提交
671
  }
L
liaogang 已提交
672
  return true;
Z
zhangjinchao01 已提交
673
}
L
liaogang 已提交
674

L
liaogang 已提交
675
void hl_profiler_start() { CHECK_CUDA(cudaProfilerStart()); }
L
liaogang 已提交
676

L
liaogang 已提交
677
void hl_profiler_end() { CHECK_CUDA(cudaProfilerStop()); }