DPC++ Runtime
Runtime libraries for oneAPI DPC++
id.hpp
Go to the documentation of this file.
1 //==-------------------- id.hpp --- SYCL iteration id ----------------------==//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 
9 #pragma once
10 
11 #include <sycl/detail/array.hpp> // for array
12 #include <sycl/detail/common.hpp> // for InitializedVal
13 #include <sycl/detail/defines.hpp> // for __SYCL_ASSUME_INT
14 #include <sycl/detail/defines_elementary.hpp> // for __SYCL_DEPRECATED, __SYCL_A...
15 #include <sycl/exception.hpp> // for make_error_code, errc, exce...
16 #include <sycl/range.hpp> // for range
17 
18 #include <stddef.h> // for size_t
19 #include <type_traits> // for enable_if_t, conditional_t
20 
21 namespace sycl {
22 inline namespace _V1 {
23 // Forward declarations
24 namespace detail {
25 template <typename TransformedArgType, int Dims, typename KernelType>
26 class RoundedRangeKernel;
27 template <typename TransformedArgType, int Dims, typename KernelType>
28 class RoundedRangeKernelWithKH;
29 } // namespace detail
30 template <int Dimensions> class range;
31 template <int Dimensions, bool with_offset> class item;
32 
36 template <int Dimensions = 1> class id : public detail::array<Dimensions> {
37 public:
38  static constexpr int dimensions = Dimensions;
39 
40 private:
41  using base = detail::array<Dimensions>;
42  static_assert(Dimensions >= 1 && Dimensions <= 3,
43  "id can only be 1, 2, or 3 Dimensional.");
44  template <int N, int val, typename T>
45  using ParamTy = std::enable_if_t<(N == val), T>;
46 
47 #ifndef __SYCL_DISABLE_ID_TO_INT_CONV__
48  /* Helper class for conversion operator. Void type is not suitable. User
49  * cannot even try to get address of the operator __private_class(). User
50  * may try to get an address of operator void() and will get the
51  * compile-time error */
52  class __private_class;
53 
54  template <typename N, typename T>
55  using EnableIfIntegral = std::enable_if_t<std::is_integral_v<N>, T>;
56  template <bool B, typename T>
57  using EnableIfT = std::conditional_t<B, T, __private_class>;
58 #endif // __SYCL_DISABLE_ID_TO_INT_CONV__
59 
60 public:
61  id() = default;
62 
63  /* The following constructor is only available in the id struct
64  * specialization where: Dimensions==1 */
65  template <int N = Dimensions> id(ParamTy<N, 1, size_t> dim0) : base(dim0) {}
66 
67  template <int N = Dimensions>
68  id(ParamTy<N, 1, const range<Dimensions>> &range_size)
69  : base(range_size.get(0)) {}
70 
71  template <int N = Dimensions, bool with_offset = true>
72  id(ParamTy<N, 1, const item<Dimensions, with_offset>> &item)
73  : base(item.get_id(0)) {}
74 
75  /* The following constructor is only available in the id struct
76  * specialization where: Dimensions==2 */
77  template <int N = Dimensions>
78  id(ParamTy<N, 2, size_t> dim0, size_t dim1) : base(dim0, dim1) {}
79 
80  template <int N = Dimensions>
81  id(ParamTy<N, 2, const range<Dimensions>> &range_size)
82  : base(range_size.get(0), range_size.get(1)) {}
83 
84  template <int N = Dimensions, bool with_offset = true>
85  id(ParamTy<N, 2, const item<Dimensions, with_offset>> &item)
86  : base(item.get_id(0), item.get_id(1)) {}
87 
88  /* The following constructor is only available in the id struct
89  * specialization where: Dimensions==3 */
90  template <int N = Dimensions>
91  id(ParamTy<N, 3, size_t> dim0, size_t dim1, size_t dim2)
92  : base(dim0, dim1, dim2) {}
93 
94  template <int N = Dimensions>
95  id(ParamTy<N, 3, const range<Dimensions>> &range_size)
96  : base(range_size.get(0), range_size.get(1), range_size.get(2)) {}
97 
98  template <int N = Dimensions, bool with_offset = true>
99  id(ParamTy<N, 3, const item<Dimensions, with_offset>> &item)
100  : base(item.get_id(0), item.get_id(1), item.get_id(2)) {}
101 
102 #ifndef __SYCL_DISABLE_ID_TO_INT_CONV__
103  /* Template operator is not allowed because it disables further type
104  * conversion. For example, the next code will not work in case of template
105  * conversion:
106  * int a = id<1>(value); */
107 
108  __SYCL_ALWAYS_INLINE operator EnableIfT<(Dimensions == 1), size_t>() const {
109  size_t Result = this->common_array[0];
110  __SYCL_ASSUME_INT(Result);
111  return Result;
112  }
113 #endif // __SYCL_DISABLE_ID_TO_INT_CONV__
114 
115 // OP is: ==, !=
116 #ifndef __SYCL_DISABLE_ID_TO_INT_CONV__
117  using detail::array<Dimensions>::operator==;
118  // Needed for clang in C++20 mode as the above operator== would be ambigious
119  // between regular/reversed call for "Id == Id" case.
120  bool operator==(const id<Dimensions> &rhs) const {
121  return this->detail::array<Dimensions>::operator==(rhs);
122  }
123 #if __cpp_impl_three_way_comparison < 201907
124  using detail::array<Dimensions>::operator!=;
125 #endif
126 
127  /* Enable operators with integral types.
128  * Template operators take precedence than type conversion. In the case of
129  * non-template operators, ambiguity appears: "id op size_t" may refer
130  * "size_t op size_t" and "id op size_t". In case of template operators it
131  * will be "id op size_t"*/
132 #define __SYCL_GEN_OPT(op) \
133  template <typename T> \
134  EnableIfIntegral<T, bool> operator op(const T &rhs) const { \
135  if (this->common_array[0] != rhs) \
136  return false op true; \
137  return true op true; \
138  } \
139  template <typename T> \
140  friend EnableIfIntegral<T, bool> operator op(const T &lhs, \
141  const id<Dimensions> &rhs) { \
142  if (lhs != rhs.common_array[0]) \
143  return false op true; \
144  return true op true; \
145  }
146 
147  __SYCL_GEN_OPT(==)
148  __SYCL_GEN_OPT(!=)
149 
150 #undef __SYCL_GEN_OPT
151 
152 #endif // __SYCL_DISABLE_ID_TO_INT_CONV__
153 
154 // OP is: +, -, *, /, %, <<, >>, &, |, ^, &&, ||, <, >, <=, >=
155 #define __SYCL_GEN_OPT_BASE(op) \
156  friend id<Dimensions> operator op(const id<Dimensions> &lhs, \
157  const id<Dimensions> &rhs) { \
158  id<Dimensions> result; \
159  for (int i = 0; i < Dimensions; ++i) { \
160  result.common_array[i] = lhs.common_array[i] op rhs.common_array[i]; \
161  } \
162  return result; \
163  }
164 
165 #ifndef __SYCL_DISABLE_ID_TO_INT_CONV__
166 // Enable operators with integral types only
167 #define __SYCL_GEN_OPT(op) \
168  __SYCL_GEN_OPT_BASE(op) \
169  template <typename T> \
170  friend EnableIfIntegral<T, id<Dimensions>> operator op( \
171  const id<Dimensions> &lhs, const T &rhs) { \
172  id<Dimensions> result; \
173  for (int i = 0; i < Dimensions; ++i) { \
174  result.common_array[i] = lhs.common_array[i] op rhs; \
175  } \
176  return result; \
177  } \
178  template <typename T> \
179  friend EnableIfIntegral<T, id<Dimensions>> operator op( \
180  const T &lhs, const id<Dimensions> &rhs) { \
181  id<Dimensions> result; \
182  for (int i = 0; i < Dimensions; ++i) { \
183  result.common_array[i] = lhs op rhs.common_array[i]; \
184  } \
185  return result; \
186  }
187 #else
188 #define __SYCL_GEN_OPT(op) \
189  __SYCL_GEN_OPT_BASE(op) \
190  friend id<Dimensions> operator op(const id<Dimensions> &lhs, \
191  const size_t &rhs) { \
192  id<Dimensions> result; \
193  for (int i = 0; i < Dimensions; ++i) { \
194  result.common_array[i] = lhs.common_array[i] op rhs; \
195  } \
196  return result; \
197  } \
198  friend id<Dimensions> operator op(const size_t &lhs, \
199  const id<Dimensions> &rhs) { \
200  id<Dimensions> result; \
201  for (int i = 0; i < Dimensions; ++i) { \
202  result.common_array[i] = lhs op rhs.common_array[i]; \
203  } \
204  return result; \
205  }
206 #endif // __SYCL_DISABLE_ID_TO_INT_CONV__
207 
208  __SYCL_GEN_OPT(+)
209  __SYCL_GEN_OPT(-)
210  __SYCL_GEN_OPT(*)
211  __SYCL_GEN_OPT(/)
212  __SYCL_GEN_OPT(%)
213  __SYCL_GEN_OPT(<<)
214  __SYCL_GEN_OPT(>>)
215  __SYCL_GEN_OPT(&)
216  __SYCL_GEN_OPT(|)
217  __SYCL_GEN_OPT(^)
218  __SYCL_GEN_OPT(&&)
219  __SYCL_GEN_OPT(||)
220  __SYCL_GEN_OPT(<)
221  __SYCL_GEN_OPT(>)
222  __SYCL_GEN_OPT(<=)
223  __SYCL_GEN_OPT(>=)
224 
225 #undef __SYCL_GEN_OPT
226 #undef __SYCL_GEN_OPT_BASE
227 
228 // OP is: +=, -=, *=, /=, %=, <<=, >>=, &=, |=, ^=
229 #define __SYCL_GEN_OPT(op) \
230  friend id<Dimensions> &operator op(id<Dimensions> &lhs, \
231  const id<Dimensions> &rhs) { \
232  for (int i = 0; i < Dimensions; ++i) { \
233  lhs.common_array[i] op rhs.common_array[i]; \
234  } \
235  return lhs; \
236  } \
237  friend id<Dimensions> &operator op(id<Dimensions> &lhs, const size_t &rhs) { \
238  for (int i = 0; i < Dimensions; ++i) { \
239  lhs.common_array[i] op rhs; \
240  } \
241  return lhs; \
242  }
243 
244  __SYCL_GEN_OPT(+=)
245  __SYCL_GEN_OPT(-=)
246  __SYCL_GEN_OPT(*=)
247  __SYCL_GEN_OPT(/=)
248  __SYCL_GEN_OPT(%=)
249  __SYCL_GEN_OPT(<<=)
250  __SYCL_GEN_OPT(>>=)
251  __SYCL_GEN_OPT(&=)
252  __SYCL_GEN_OPT(|=)
253  __SYCL_GEN_OPT(^=)
254 
255 #undef __SYCL_GEN_OPT
256 
257 // OP is unary +, -
258 #define __SYCL_GEN_OPT(op) \
259  friend id<Dimensions> operator op(const id<Dimensions> &rhs) { \
260  id<Dimensions> result; \
261  for (int i = 0; i < Dimensions; ++i) { \
262  result.common_array[i] = (op rhs.common_array[i]); \
263  } \
264  return result; \
265  }
266 
267  __SYCL_GEN_OPT(+)
268  __SYCL_GEN_OPT(-)
269 
270 #undef __SYCL_GEN_OPT
271 
272 // OP is prefix ++, --
273 #define __SYCL_GEN_OPT(op) \
274  friend id<Dimensions> &operator op(id<Dimensions> &rhs) { \
275  for (int i = 0; i < Dimensions; ++i) { \
276  op rhs.common_array[i]; \
277  } \
278  return rhs; \
279  }
280 
281  __SYCL_GEN_OPT(++)
282  __SYCL_GEN_OPT(--)
283 
284 #undef __SYCL_GEN_OPT
285 
286 // OP is postfix ++, --
287 #define __SYCL_GEN_OPT(op) \
288  friend id<Dimensions> operator op(id<Dimensions> &lhs, int) { \
289  id<Dimensions> old_lhs; \
290  for (int i = 0; i < Dimensions; ++i) { \
291  old_lhs.common_array[i] = lhs.common_array[i]; \
292  op lhs.common_array[i]; \
293  } \
294  return old_lhs; \
295  }
296 
297  __SYCL_GEN_OPT(++)
298  __SYCL_GEN_OPT(--)
299 
300 #undef __SYCL_GEN_OPT
301 };
302 
303 namespace detail {
304 template <int Dimensions>
305 size_t getOffsetForId(range<Dimensions> Range, id<Dimensions> Id,
306  id<Dimensions> Offset) {
307  size_t offset = 0;
308  for (int i = 0; i < Dimensions; ++i)
309  offset = offset * Range[i] + Offset[i] + Id[i];
310  return offset;
311 }
312 
313 inline id<1> getDelinearizedId(const range<1> &, size_t Index) {
314  return {Index};
315 }
316 
317 inline id<2> getDelinearizedId(const range<2> &Range, size_t Index) {
318  size_t X = Index % Range[1];
319  size_t Y = Index / Range[1];
320  return {Y, X};
321 }
322 
323 inline id<3> getDelinearizedId(const range<3> &Range, size_t Index) {
324  size_t D1D2 = Range[1] * Range[2];
325  size_t Z = Index / D1D2;
326  size_t ZRest = Index % D1D2;
327  size_t Y = ZRest / Range[2];
328  size_t X = ZRest % Range[2];
329  return {Z, Y, X};
330 }
331 } // namespace detail
332 
333 // C++ feature test macros are supported by all supported compilers
334 // with the exception of MSVC 1914. It doesn't support deduction guides.
335 #ifdef __cpp_deduction_guides
336 id(size_t)->id<1>;
337 id(size_t, size_t)->id<2>;
338 id(size_t, size_t, size_t)->id<3>;
339 #endif
340 } // namespace _V1
341 } // namespace sycl
sycl::_V1::detail::array
Definition: array.hpp:24
sycl::_V1::detail::array::operator==
bool operator==(const array< dimensions > &rhs) const
Definition: array.hpp:84
sycl::_V1::detail::array< 1 >::common_array
size_t common_array[dimensions]
Definition: array.hpp:105
sycl::_V1::detail::array< 1 >::get
size_t get(int dimension) const
Definition: array.hpp:62
sycl::_V1::id
A unique identifier of an item in an index space.
Definition: id.hpp:36
sycl::_V1::id::id
id(ParamTy< N, 2, size_t > dim0, size_t dim1)
Definition: id.hpp:78
sycl::_V1::id::id
id(ParamTy< N, 3, const range< Dimensions >> &range_size)
Definition: id.hpp:95
sycl::_V1::id::id
id(ParamTy< N, 3, const item< Dimensions, with_offset >> &item)
Definition: id.hpp:99
sycl::_V1::id::id
id(ParamTy< N, 2, const range< Dimensions >> &range_size)
Definition: id.hpp:81
sycl::_V1::id::id
id(ParamTy< N, 3, size_t > dim0, size_t dim1, size_t dim2)
Definition: id.hpp:91
sycl::_V1::id::dimensions
static constexpr int dimensions
Definition: id.hpp:38
sycl::_V1::id::id
id(ParamTy< N, 1, const range< Dimensions >> &range_size)
Definition: id.hpp:68
sycl::_V1::id::id
id(ParamTy< N, 1, size_t > dim0)
Definition: id.hpp:65
sycl::_V1::id::id
id()=default
sycl::_V1::id::id
id(ParamTy< N, 1, const item< Dimensions, with_offset >> &item)
Definition: id.hpp:72
sycl::_V1::id::id
id(ParamTy< N, 2, const item< Dimensions, with_offset >> &item)
Definition: id.hpp:85
sycl::_V1::id::operator==
bool operator==(const id< Dimensions > &rhs) const
Definition: id.hpp:120
sycl::_V1::item
Identifies an instance of the function object executing at each point in a range.
Definition: item.hpp:37
sycl::_V1::range< Dimensions >
__SYCL_ASSUME_INT
#define __SYCL_ASSUME_INT(x)
Definition: defines.hpp:17
defines_elementary.hpp
__SYCL_ALWAYS_INLINE
#define __SYCL_ALWAYS_INLINE
Definition: defines_elementary.hpp:23
__SYCL_GEN_OPT
#define __SYCL_GEN_OPT(op)
Definition: id.hpp:287
sycl::_V1::detail::getDelinearizedId
id< 1 > getDelinearizedId(const range< 1 > &, size_t Index)
Definition: id.hpp:313
sycl::_V1::detail::getOffsetForId
size_t getOffsetForId(range< Dimensions > Range, id< Dimensions > Id, id< Dimensions > Offset)
Definition: id.hpp:305
sycl::_V1::Dimensions
class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor class __SYCL_EBO __SYCL_SPECIAL_CLASS Dimensions
Definition: accessor.hpp:3233
sycl
Definition: access.hpp:18