macro.cuh 3.34 KB
Newer Older
sangwzh's avatar
sangwzh committed
1
// !!! This is a file automatically generated by hipify!!!
2
/**
3
 *  Copyright (c) 2020 by Contributors
4
5
 * @file array/cuda/macro.cuh
 * @brief Macro to call SPMM/SDDMM cuda kernels.
6
7
8
9
10
11
12
13
14
15
 */
#ifndef DGL_ARRAY_CUDA_MACRO_CUH_
#define DGL_ARRAY_CUDA_MACRO_CUH_

///////////////////////// Dispatchers //////////////////////////

/* Macro used for switching between broadcasting and non-broadcasting kernels.
 * It also copies the auxiliary information for calculating broadcasting offsets
 * to GPU.
 */
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
#define BCAST_IDX_CTX_SWITCH(BCAST, EDGE_MAP, CTX, LHS_OFF, RHS_OFF, ...)     \
  do {                                                                        \
    const BcastOff &info = (BCAST);                                           \
    if (!info.use_bcast) {                                                    \
      constexpr bool UseBcast = false;                                        \
      if ((EDGE_MAP)) {                                                       \
        constexpr bool UseIdx = true;                                         \
        { __VA_ARGS__ }                                                       \
      } else {                                                                \
        constexpr bool UseIdx = false;                                        \
        { __VA_ARGS__ }                                                       \
      }                                                                       \
    } else {                                                                  \
      constexpr bool UseBcast = true;                                         \
      const DGLContext ctx = (CTX);                                           \
      const auto device = runtime::DeviceAPI::Get(ctx);                       \
      (LHS_OFF) = static_cast<int64_t *>(device->AllocWorkspace(              \
          ctx, sizeof(int64_t) * info.lhs_offset.size()));                    \
sangwzh's avatar
sangwzh committed
34
      CUDA_CALL(hipMemcpy(                                                   \
35
          (LHS_OFF), &info.lhs_offset[0],                                     \
sangwzh's avatar
sangwzh committed
36
          sizeof(int64_t) * info.lhs_offset.size(), hipMemcpyHostToDevice)); \
37
38
      (RHS_OFF) = static_cast<int64_t *>(device->AllocWorkspace(              \
          ctx, sizeof(int64_t) * info.rhs_offset.size()));                    \
sangwzh's avatar
sangwzh committed
39
      CUDA_CALL(hipMemcpy(                                                   \
40
          (RHS_OFF), &info.rhs_offset[0],                                     \
sangwzh's avatar
sangwzh committed
41
          sizeof(int64_t) * info.rhs_offset.size(), hipMemcpyHostToDevice)); \
42
43
44
45
46
47
48
49
50
51
52
      if ((EDGE_MAP)) {                                                       \
        constexpr bool UseIdx = true;                                         \
        { __VA_ARGS__ }                                                       \
      } else {                                                                \
        constexpr bool UseIdx = false;                                        \
        { __VA_ARGS__ }                                                       \
      }                                                                       \
      device->FreeWorkspace(ctx, (LHS_OFF));                                  \
      device->FreeWorkspace(ctx, (RHS_OFF));                                  \
    }                                                                         \
  } while (0)
53

54
#endif  // DGL_ARRAY_CUDA_MACRO_CUH_