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:
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  __SYCL_DEPRECATED("range() conversion is deprecated")
103  explicit operator range<Dimensions>() const {
104  range<Dimensions> result(
106  for (int i = 0; i < Dimensions; ++i) {
107  result[i] = this->get(i);
108  }
109  return result;
110  }
111 
112 #ifndef __SYCL_DISABLE_ID_TO_INT_CONV__
113  /* Template operator is not allowed because it disables further type
114  * conversion. For example, the next code will not work in case of template
115  * conversion:
116  * int a = id<1>(value); */
117 
118  __SYCL_ALWAYS_INLINE operator EnableIfT<(Dimensions == 1), size_t>() const {
119  size_t Result = this->common_array[0];
120  __SYCL_ASSUME_INT(Result);
121  return Result;
122  }
123 #endif // __SYCL_DISABLE_ID_TO_INT_CONV__
124 
125 // OP is: ==, !=
126 #ifndef __SYCL_DISABLE_ID_TO_INT_CONV__
128  // Needed for clang in C++20 mode as the above operator== would be ambigious
129  // between regular/reversed call for "Id == Id" case.
130  bool operator==(const id<Dimensions> &rhs) const {
131  return this->detail::array<Dimensions>::operator==(rhs);
132  }
133 #if __cpp_impl_three_way_comparison < 201907
135 #endif
136 
137  /* Enable operators with integral types.
138  * Template operators take precedence than type conversion. In the case of
139  * non-template operators, ambiguity appears: "id op size_t" may refer
140  * "size_t op size_t" and "id op size_t". In case of template operators it
141  * will be "id op size_t"*/
142 #define __SYCL_GEN_OPT(op) \
143  template <typename T> \
144  EnableIfIntegral<T, bool> operator op(const T &rhs) const { \
145  if (this->common_array[0] != rhs) \
146  return false op true; \
147  return true op true; \
148  } \
149  template <typename T> \
150  friend EnableIfIntegral<T, bool> operator op(const T &lhs, \
151  const id<Dimensions> &rhs) { \
152  if (lhs != rhs.common_array[0]) \
153  return false op true; \
154  return true op true; \
155  }
156 
157  __SYCL_GEN_OPT(==)
158  __SYCL_GEN_OPT(!=)
159 
160 #undef __SYCL_GEN_OPT
161 
162 #endif // __SYCL_DISABLE_ID_TO_INT_CONV__
163 
164 // OP is: +, -, *, /, %, <<, >>, &, |, ^, &&, ||, <, >, <=, >=
165 #define __SYCL_GEN_OPT_BASE(op) \
166  friend id<Dimensions> operator op(const id<Dimensions> &lhs, \
167  const id<Dimensions> &rhs) { \
168  id<Dimensions> result; \
169  for (int i = 0; i < Dimensions; ++i) { \
170  result.common_array[i] = lhs.common_array[i] op rhs.common_array[i]; \
171  } \
172  return result; \
173  }
174 
175 #ifndef __SYCL_DISABLE_ID_TO_INT_CONV__
176 // Enable operators with integral types only
177 #define __SYCL_GEN_OPT(op) \
178  __SYCL_GEN_OPT_BASE(op) \
179  template <typename T> \
180  friend EnableIfIntegral<T, id<Dimensions>> operator op( \
181  const id<Dimensions> &lhs, const T &rhs) { \
182  id<Dimensions> result; \
183  for (int i = 0; i < Dimensions; ++i) { \
184  result.common_array[i] = lhs.common_array[i] op rhs; \
185  } \
186  return result; \
187  } \
188  template <typename T> \
189  friend EnableIfIntegral<T, id<Dimensions>> operator op( \
190  const T &lhs, const id<Dimensions> &rhs) { \
191  id<Dimensions> result; \
192  for (int i = 0; i < Dimensions; ++i) { \
193  result.common_array[i] = lhs op rhs.common_array[i]; \
194  } \
195  return result; \
196  }
197 #else
198 #define __SYCL_GEN_OPT(op) \
199  __SYCL_GEN_OPT_BASE(op) \
200  friend id<Dimensions> operator op(const id<Dimensions> &lhs, \
201  const size_t &rhs) { \
202  id<Dimensions> result; \
203  for (int i = 0; i < Dimensions; ++i) { \
204  result.common_array[i] = lhs.common_array[i] op rhs; \
205  } \
206  return result; \
207  } \
208  friend id<Dimensions> operator op(const size_t &lhs, \
209  const id<Dimensions> &rhs) { \
210  id<Dimensions> result; \
211  for (int i = 0; i < Dimensions; ++i) { \
212  result.common_array[i] = lhs op rhs.common_array[i]; \
213  } \
214  return result; \
215  }
216 #endif // __SYCL_DISABLE_ID_TO_INT_CONV__
217 
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  __SYCL_GEN_OPT(>>)
225  __SYCL_GEN_OPT(&)
226  __SYCL_GEN_OPT(|)
227  __SYCL_GEN_OPT(^)
228  __SYCL_GEN_OPT(&&)
229  __SYCL_GEN_OPT(||)
230  __SYCL_GEN_OPT(<)
231  __SYCL_GEN_OPT(>)
232  __SYCL_GEN_OPT(<=)
233  __SYCL_GEN_OPT(>=)
234 
235 #undef __SYCL_GEN_OPT
236 #undef __SYCL_GEN_OPT_BASE
237 
238 // OP is: +=, -=, *=, /=, %=, <<=, >>=, &=, |=, ^=
239 #define __SYCL_GEN_OPT(op) \
240  friend id<Dimensions> &operator op(id<Dimensions> &lhs, \
241  const id<Dimensions> &rhs) { \
242  for (int i = 0; i < Dimensions; ++i) { \
243  lhs.common_array[i] op rhs.common_array[i]; \
244  } \
245  return lhs; \
246  } \
247  friend id<Dimensions> &operator op(id<Dimensions> &lhs, const size_t &rhs) { \
248  for (int i = 0; i < Dimensions; ++i) { \
249  lhs.common_array[i] op rhs; \
250  } \
251  return lhs; \
252  }
253 
254  __SYCL_GEN_OPT(+=)
255  __SYCL_GEN_OPT(-=)
256  __SYCL_GEN_OPT(*=)
257  __SYCL_GEN_OPT(/=)
258  __SYCL_GEN_OPT(%=)
259  __SYCL_GEN_OPT(<<=)
260  __SYCL_GEN_OPT(>>=)
261  __SYCL_GEN_OPT(&=)
262  __SYCL_GEN_OPT(|=)
263  __SYCL_GEN_OPT(^=)
264 
265 #undef __SYCL_GEN_OPT
266 
267 // OP is unary +, -
268 #define __SYCL_GEN_OPT(op) \
269  friend id<Dimensions> operator op(const id<Dimensions> &rhs) { \
270  id<Dimensions> result; \
271  for (int i = 0; i < Dimensions; ++i) { \
272  result.common_array[i] = (op rhs.common_array[i]); \
273  } \
274  return result; \
275  }
276 
277  __SYCL_GEN_OPT(+)
278  __SYCL_GEN_OPT(-)
279 
280 #undef __SYCL_GEN_OPT
281 
282 // OP is prefix ++, --
283 #define __SYCL_GEN_OPT(op) \
284  friend id<Dimensions> &operator op(id<Dimensions> &rhs) { \
285  for (int i = 0; i < Dimensions; ++i) { \
286  op rhs.common_array[i]; \
287  } \
288  return rhs; \
289  }
290 
291  __SYCL_GEN_OPT(++)
292  __SYCL_GEN_OPT(--)
293 
294 #undef __SYCL_GEN_OPT
295 
296 // OP is postfix ++, --
297 #define __SYCL_GEN_OPT(op) \
298  friend id<Dimensions> operator op(id<Dimensions> &lhs, int) { \
299  id<Dimensions> old_lhs; \
300  for (int i = 0; i < Dimensions; ++i) { \
301  old_lhs.common_array[i] = lhs.common_array[i]; \
302  op lhs.common_array[i]; \
303  } \
304  return old_lhs; \
305  }
306 
307  __SYCL_GEN_OPT(++)
308  __SYCL_GEN_OPT(--)
309 
310 #undef __SYCL_GEN_OPT
311 };
312 
313 namespace detail {
314 template <int Dimensions>
316  id<Dimensions> Offset) {
317  size_t offset = 0;
318  for (int i = 0; i < Dimensions; ++i)
319  offset = offset * Range[i] + Offset[i] + Id[i];
320  return offset;
321 }
322 
323 inline id<1> getDelinearizedId(const range<1> &, size_t Index) {
324  return {Index};
325 }
326 
327 inline id<2> getDelinearizedId(const range<2> &Range, size_t Index) {
328  size_t X = Index % Range[1];
329  size_t Y = Index / Range[1];
330  return {Y, X};
331 }
332 
333 inline id<3> getDelinearizedId(const range<3> &Range, size_t Index) {
334  size_t D1D2 = Range[1] * Range[2];
335  size_t Z = Index / D1D2;
336  size_t ZRest = Index % D1D2;
337  size_t Y = ZRest / Range[2];
338  size_t X = ZRest % Range[2];
339  return {Z, Y, X};
340 }
341 } // namespace detail
342 
343 // C++ feature test macros are supported by all supported compilers
344 // with the exception of MSVC 1914. It doesn't support deduction guides.
345 #ifdef __cpp_deduction_guides
346 id(size_t)->id<1>;
347 id(size_t, size_t)->id<2>;
348 id(size_t, size_t, size_t)->id<3>;
349 #endif
350 
351 template <int Dims>
352 __SYCL_DEPRECATED("use sycl::ext::oneapi::experimental::this_id() instead")
353 id<Dims> this_id() {
354 #ifdef __SYCL_DEVICE_ONLY__
355  return detail::Builder::getElement(detail::declptr<id<Dims>>());
356 #else
357  throw sycl::exception(
358  sycl::make_error_code(sycl::errc::feature_not_supported),
359  "Free function calls are not supported on host");
360 #endif
361 }
362 
363 namespace ext::oneapi::experimental {
364 template <int Dims> id<Dims> this_id() {
365 #ifdef __SYCL_DEVICE_ONLY__
366  return sycl::detail::Builder::getElement(sycl::detail::declptr<id<Dims>>());
367 #else
368  throw sycl::exception(
369  sycl::make_error_code(sycl::errc::feature_not_supported),
370  "Free function calls are not supported on host");
371 #endif
372 }
373 } // namespace ext::oneapi::experimental
374 } // namespace _V1
375 } // namespace sycl
bool operator==(const array< dimensions > &rhs) const
Definition: array.hpp:92
size_t common_array[dimensions]
Definition: array.hpp:113
size_t get(int dimension) const
Definition: array.hpp:70
A unique identifier of an item in an index space.
Definition: id.hpp:36
id(ParamTy< N, 2, size_t > dim0, size_t dim1)
Definition: id.hpp:78
id(ParamTy< N, 3, const range< Dimensions >> &range_size)
Definition: id.hpp:95
id(ParamTy< N, 3, const item< Dimensions, with_offset >> &item)
Definition: id.hpp:99
id(ParamTy< N, 2, const range< Dimensions >> &range_size)
Definition: id.hpp:81
id(ParamTy< N, 3, size_t > dim0, size_t dim1, size_t dim2)
Definition: id.hpp:91
static constexpr int dimensions
Definition: id.hpp:38
id(ParamTy< N, 1, const range< Dimensions >> &range_size)
Definition: id.hpp:68
id(ParamTy< N, 1, size_t > dim0)
Definition: id.hpp:65
id()=default
id(ParamTy< N, 1, const item< Dimensions, with_offset >> &item)
Definition: id.hpp:72
id(ParamTy< N, 2, const item< Dimensions, with_offset >> &item)
Definition: id.hpp:85
bool operator==(const id< Dimensions > &rhs) const
Definition: id.hpp:130
Identifies an instance of the function object executing at each point in a range.
Definition: item.hpp:37
#define __SYCL_ASSUME_INT(x)
Definition: defines.hpp:17
#define __SYCL_ALWAYS_INLINE
#define __SYCL_GEN_OPT(op)
Definition: id.hpp:297
id< 1 > getDelinearizedId(const range< 1 > &, size_t Index)
Definition: id.hpp:323
size_t getOffsetForId(range< Dimensions > Range, id< Dimensions > Id, id< Dimensions > Offset)
Definition: id.hpp:315
id< Dims > this_id()
Definition: id.hpp:353
class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor class __SYCL_EBO __SYCL_SPECIAL_CLASS Dimensions
Definition: accessor.hpp:3233
__SYCL_DEPRECATED("This is a deprecated argument type for SYCL nan built-in function.") std
std::error_code make_error_code(sycl::errc E) noexcept
Constructs an error code using e and sycl_category()
Definition: exception.cpp:94
Definition: access.hpp:18