thread_storage_scope.h 5.68 KB
Newer Older
1
/**
Minjie Wang's avatar
Minjie Wang committed
2
 *  Copyright (c) 2017 by Contributors
3
4
 * @file thread_storage_scope.h
 * @brief Extract thread axis configuration from DGLArgs.
Minjie Wang's avatar
Minjie Wang committed
5
 */
6
7
#ifndef DGL_RUNTIME_THREAD_STORAGE_SCOPE_H_
#define DGL_RUNTIME_THREAD_STORAGE_SCOPE_H_
Minjie Wang's avatar
Minjie Wang committed
8
9

#include <dgl/runtime/packed_func.h>
10

Minjie Wang's avatar
Minjie Wang committed
11
12
13
#include <string>
#include <vector>

14
namespace dgl {
Minjie Wang's avatar
Minjie Wang committed
15
16
namespace runtime {

17
/**
18
19
 * @brief Memory hierachy rank in the storage system
 * @note The global rank and shared rank have one to one
Minjie Wang's avatar
Minjie Wang committed
20
21
22
 *       correspondence to the thread rank.
 */
enum class StorageRank {
23
  /** @brief global memory */
Minjie Wang's avatar
Minjie Wang committed
24
  kGlobal = 0,
25
  /** @brief shared memory among thread group */
Minjie Wang's avatar
Minjie Wang committed
26
  kShared = 1,
27
  /**
28
   * @brief reserved for warp memory.
Minjie Wang's avatar
Minjie Wang committed
29
30
31
32
33
   *  This is only used by programming model.
   *  There is no such memory usually in GPU.
   *  Instead, we can simulate it by registers and shuffle.
   */
  kWarp = 2,
34
  /** @brief thread local memory */
Minjie Wang's avatar
Minjie Wang committed
35
36
37
  kLocal = 3
};

38
/**
39
40
 * @param thread_scope_rank The thread scope rank
 * @return default storage rank given the thread scope
Minjie Wang's avatar
Minjie Wang committed
41
42
43
 */
inline StorageRank DefaultStorageRank(int thread_scope_rank) {
  switch (thread_scope_rank) {
44
45
46
47
48
49
    case -1:
      return StorageRank::kGlobal;
    case 0:
      return StorageRank::kShared;
    case 1:
      return StorageRank::kLocal;
Minjie Wang's avatar
Minjie Wang committed
50
51
52
53
54
55
56
    default: {
      LOG(FATAL) << "unknown rank";
      return StorageRank::kGlobal;
    }
  }
}

57
/** @brief class to represent storage scope */
Minjie Wang's avatar
Minjie Wang committed
58
struct StorageScope {
59
  /** @brief The rank of the storage */
Minjie Wang's avatar
Minjie Wang committed
60
  StorageRank rank{StorageRank::kGlobal};
61
  /** @brief tag for special purpose memory. */
Minjie Wang's avatar
Minjie Wang committed
62
63
64
65
66
67
68
69
70
71
72
  std::string tag;
  // comparator
  inline bool operator==(const StorageScope& other) const {
    return rank == other.rank && tag == other.tag;
  }
  inline bool operator!=(const StorageScope& other) const {
    return !(*this == other);
  }
  inline std::string to_string() const {
    std::string ret;
    switch (rank) {
73
74
75
76
77
78
79
80
81
82
83
      case StorageRank::kGlobal:
        return "global" + tag;
      case StorageRank::kShared:
        return "shared" + tag;
      case StorageRank::kWarp:
        return "warp" + tag;
      case StorageRank::kLocal:
        return "local" + tag;
      default:
        LOG(FATAL) << "unknown storage scope";
        return "";
Minjie Wang's avatar
Minjie Wang committed
84
85
    }
  }
86
  /**
87
88
89
   * @brief make storage scope from string
   * @param s The string to be parsed.
   * @return The storage scope.
Minjie Wang's avatar
Minjie Wang committed
90
91
92
   */
  static StorageScope make(const std::string& s) {
    StorageScope r;
93
    if (s.compare(0, 6, "global") == 0) {
Minjie Wang's avatar
Minjie Wang committed
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
      r.rank = StorageRank::kGlobal;
      r.tag = s.substr(6, std::string::npos);
    } else if (s.compare(0, 6, "shared") == 0) {
      r.rank = StorageRank::kShared;
      r.tag = s.substr(6, std::string::npos);
    } else if (s.compare(0, 4, "warp") == 0) {
      r.rank = StorageRank::kWarp;
      r.tag = s.substr(4, std::string::npos);
    } else if (s.compare(0, 5, "local") == 0) {
      r.rank = StorageRank::kLocal;
      r.tag = s.substr(5, std::string::npos);
    } else {
      LOG(FATAL) << "unknown storage scope " << s;
    }
    return r;
  }
};

112
/** @brief class to represent thread scope */
Minjie Wang's avatar
Minjie Wang committed
113
struct ThreadScope {
114
  /** @brief The rank of thread scope */
Minjie Wang's avatar
Minjie Wang committed
115
  int rank{0};
116
  /** @brief the dimension index under the rank */
Minjie Wang's avatar
Minjie Wang committed
117
  int dim_index{0};
118
  /**
119
120
121
   * @brief make storage scope from string
   * @param s The string to be parsed.
   * @return The storage scope.
Minjie Wang's avatar
Minjie Wang committed
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
   */
  static ThreadScope make(const std::string& s) {
    ThreadScope r;
    if (s == "vthread" || s == "cthread") {
      // virtual thread at the same level as local
      r.rank = 1;
      r.dim_index = -1;
    } else if (s.compare(0, 9, "blockIdx.") == 0) {
      r.rank = 0;
      r.dim_index = static_cast<int>(s[9] - 'x');
    } else if (s.compare(0, 10, "threadIdx.") == 0) {
      r.rank = 1;
      r.dim_index = static_cast<int>(s[10] - 'x');
    } else {
      LOG(FATAL) << "Unknown threadscope " << s;
    }
    return r;
  }
};

142
/** @brief workload speccification */
Minjie Wang's avatar
Minjie Wang committed
143
144
145
struct ThreadWorkLoad {
  // array, first three are thread configuration.
  size_t work_size[6];
146
  /**
147
148
   * @param i The block dimension.
   * @return i-th block dim
Minjie Wang's avatar
Minjie Wang committed
149
   */
150
  inline size_t block_dim(size_t i) const { return work_size[i + 3]; }
151
  /**
152
153
   * @param i The grid dimension.
   * @return i-th grid dim
Minjie Wang's avatar
Minjie Wang committed
154
   */
155
  inline size_t grid_dim(size_t i) const { return work_size[i]; }
Minjie Wang's avatar
Minjie Wang committed
156
};
157
/** @brief Thread axis configuration */
Minjie Wang's avatar
Minjie Wang committed
158
159
class ThreadAxisConfig {
 public:
160
  void Init(size_t base, const std::vector<std::string>& thread_axis_tags) {
Minjie Wang's avatar
Minjie Wang committed
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
    base_ = base;
    std::vector<bool> filled(6, false);
    for (size_t i = 0; i < thread_axis_tags.size(); ++i) {
      const std::string& tag = thread_axis_tags[i];
      ThreadScope ts = ThreadScope::make(tag);
      arg_index_map_.push_back(ts.rank * 3 + ts.dim_index);
      filled[ts.rank * 3 + ts.dim_index] = true;
    }
    work_dim_ = 1;
    for (int i = 0; i < 3; ++i) {
      if (filled[i] || filled[i + 3]) {
        work_dim_ = i + 1;
      }
    }
  }
  // extract workload from arguments.
177
  ThreadWorkLoad Extract(DGLArgs x) const {
Minjie Wang's avatar
Minjie Wang committed
178
179
180
181
182
183
184
185
186
    ThreadWorkLoad w;
    std::fill(w.work_size, w.work_size + 6, 1);
    for (size_t i = 0; i < arg_index_map_.size(); ++i) {
      w.work_size[arg_index_map_[i]] =
          static_cast<size_t>(x.values[base_ + i].v_int64);
    }
    return w;
  }
  // return the work dim
187
  size_t work_dim() const { return work_dim_; }
Minjie Wang's avatar
Minjie Wang committed
188
189

 private:
190
  /** @brief base axis */
Minjie Wang's avatar
Minjie Wang committed
191
  size_t base_;
192
  /** @brief The worker dimension */
Minjie Wang's avatar
Minjie Wang committed
193
  size_t work_dim_;
194
  /** @brief The index mapping. */
Minjie Wang's avatar
Minjie Wang committed
195
196
197
198
  std::vector<uint32_t> arg_index_map_;
};

}  // namespace runtime
199
}  // namespace dgl
Minjie Wang's avatar
Minjie Wang committed
200
201
202

namespace std {
template <>
203
204
struct hash<::dgl::runtime::StorageScope> {
  std::size_t operator()(const ::dgl::runtime::StorageScope& k) const {
Minjie Wang's avatar
Minjie Wang committed
205
206
207
208
    return static_cast<size_t>(k.rank);
  }
};
}  // namespace std
209
#endif  // DGL_RUNTIME_THREAD_STORAGE_SCOPE_H_