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>
12 #include <sycl/detail/common.hpp>
13 #include <sycl/detail/helpers.hpp>
14 #include <sycl/detail/type_traits.hpp>
15 #include <sycl/range.hpp>
16 
17 namespace sycl {
18 __SYCL_INLINE_VER_NAMESPACE(_V1) {
19 // Forward declarations
20 namespace detail {
21 template <typename TransformedArgType, int Dims, typename KernelType>
22 class RoundedRangeKernel;
23 template <typename TransformedArgType, int Dims, typename KernelType>
24 class RoundedRangeKernelWithKH;
25 } // namespace detail
26 template <int dimensions> class range;
27 template <int dimensions, bool with_offset> class item;
28 
32 template <int dimensions = 1> class id : public detail::array<dimensions> {
33 private:
34  using base = detail::array<dimensions>;
35  static_assert(dimensions >= 1 && dimensions <= 3,
36  "id can only be 1, 2, or 3 dimensional.");
37  template <int N, int val, typename T>
38  using ParamTy = detail::enable_if_t<(N == val), T>;
39 
40 #ifndef __SYCL_DISABLE_ID_TO_INT_CONV__
41  /* Helper class for conversion operator. Void type is not suitable. User
42  * cannot even try to get address of the operator __private_class(). User
43  * may try to get an address of operator void() and will get the
44  * compile-time error */
45  class __private_class;
46 
47  template <typename N, typename T>
48  using EnableIfIntegral = detail::enable_if_t<std::is_integral<N>::value, T>;
49  template <bool B, typename T>
50  using EnableIfT = detail::conditional_t<B, T, __private_class>;
51 #endif // __SYCL_DISABLE_ID_TO_INT_CONV__
52 
53 public:
54  id() = default;
55 
56  /* The following constructor is only available in the id struct
57  * specialization where: dimensions==1 */
58  template <int N = dimensions> id(ParamTy<N, 1, size_t> dim0) : base(dim0) {}
59 
60  template <int N = dimensions>
61  id(ParamTy<N, 1, const range<dimensions>> &range_size)
62  : base(range_size.get(0)) {}
63 
64  template <int N = dimensions, bool with_offset = true>
65  id(ParamTy<N, 1, const item<dimensions, with_offset>> &item)
66  : base(item.get_id(0)) {}
67 
68  /* The following constructor is only available in the id struct
69  * specialization where: dimensions==2 */
70  template <int N = dimensions>
71  id(ParamTy<N, 2, size_t> dim0, size_t dim1) : base(dim0, dim1) {}
72 
73  template <int N = dimensions>
74  id(ParamTy<N, 2, const range<dimensions>> &range_size)
75  : base(range_size.get(0), range_size.get(1)) {}
76 
77  template <int N = dimensions, bool with_offset = true>
78  id(ParamTy<N, 2, const item<dimensions, with_offset>> &item)
79  : base(item.get_id(0), item.get_id(1)) {}
80 
81  /* The following constructor is only available in the id struct
82  * specialization where: dimensions==3 */
83  template <int N = dimensions>
84  id(ParamTy<N, 3, size_t> dim0, size_t dim1, size_t dim2)
85  : base(dim0, dim1, dim2) {}
86 
87  template <int N = dimensions>
88  id(ParamTy<N, 3, const range<dimensions>> &range_size)
89  : base(range_size.get(0), range_size.get(1), range_size.get(2)) {}
90 
91  template <int N = dimensions, bool with_offset = true>
92  id(ParamTy<N, 3, const item<dimensions, with_offset>> &item)
93  : base(item.get_id(0), item.get_id(1), item.get_id(2)) {}
94 
95  __SYCL_DEPRECATED("range() conversion is deprecated")
96  explicit operator range<dimensions>() const {
97  range<dimensions> result(
98  detail::InitializedVal<dimensions, range>::template get<0>());
99  for (int i = 0; i < dimensions; ++i) {
100  result[i] = this->get(i);
101  }
102  return result;
103  }
104 
105 #ifndef __SYCL_DISABLE_ID_TO_INT_CONV__
106  /* Template operator is not allowed because it disables further type
107  * conversion. For example, the next code will not work in case of template
108  * conversion:
109  * int a = id<1>(value); */
110 
111  __SYCL_ALWAYS_INLINE operator EnableIfT<(dimensions == 1), size_t>() const {
112  size_t Result = this->common_array[0];
113  __SYCL_ASSUME_INT(Result);
114  return Result;
115  }
116 #endif // __SYCL_DISABLE_ID_TO_INT_CONV__
117 
118 // OP is: ==, !=
119 #ifndef __SYCL_DISABLE_ID_TO_INT_CONV__
120  using detail::array<dimensions>::operator==;
121  // Needed for clang in C++20 mode as the above operator== would be ambigious
122  // between regular/reversed call for "Id == Id" case.
123  bool operator==(const id<dimensions> &rhs) const {
124  return this->detail::array<dimensions>::operator==(rhs);
125  }
126 #if __cpp_impl_three_way_comparison < 201907
127  using detail::array<dimensions>::operator!=;
128 #endif
129 
130  /* Enable operators with integral types.
131  * Template operators take precedence than type conversion. In the case of
132  * non-template operators, ambiguity appears: "id op size_t" may refer
133  * "size_t op size_t" and "id op size_t". In case of template operators it
134  * will be "id op size_t"*/
135 #define __SYCL_GEN_OPT(op) \
136  template <typename T> \
137  EnableIfIntegral<T, bool> operator op(const T &rhs) const { \
138  if (this->common_array[0] != rhs) \
139  return false op true; \
140  return true op true; \
141  } \
142  template <typename T> \
143  friend EnableIfIntegral<T, bool> operator op(const T &lhs, \
144  const id<dimensions> &rhs) { \
145  if (lhs != rhs.common_array[0]) \
146  return false op true; \
147  return true op true; \
148  }
149 
150  __SYCL_GEN_OPT(==)
151  __SYCL_GEN_OPT(!=)
152 
153 #undef __SYCL_GEN_OPT
154 
155 #endif // __SYCL_DISABLE_ID_TO_INT_CONV__
156 
157 // OP is: +, -, *, /, %, <<, >>, &, |, ^, &&, ||, <, >, <=, >=
158 #define __SYCL_GEN_OPT_BASE(op) \
159  friend id<dimensions> operator op(const id<dimensions> &lhs, \
160  const id<dimensions> &rhs) { \
161  id<dimensions> result; \
162  for (int i = 0; i < dimensions; ++i) { \
163  result.common_array[i] = lhs.common_array[i] op rhs.common_array[i]; \
164  } \
165  return result; \
166  }
167 
168 #ifndef __SYCL_DISABLE_ID_TO_INT_CONV__
169 // Enable operators with integral types only
170 #define __SYCL_GEN_OPT(op) \
171  __SYCL_GEN_OPT_BASE(op) \
172  template <typename T> \
173  friend EnableIfIntegral<T, id<dimensions>> operator op( \
174  const id<dimensions> &lhs, const T &rhs) { \
175  id<dimensions> result; \
176  for (int i = 0; i < dimensions; ++i) { \
177  result.common_array[i] = lhs.common_array[i] op rhs; \
178  } \
179  return result; \
180  } \
181  template <typename T> \
182  friend EnableIfIntegral<T, id<dimensions>> operator op( \
183  const T &lhs, const id<dimensions> &rhs) { \
184  id<dimensions> result; \
185  for (int i = 0; i < dimensions; ++i) { \
186  result.common_array[i] = lhs op rhs.common_array[i]; \
187  } \
188  return result; \
189  }
190 #else
191 #define __SYCL_GEN_OPT(op) \
192  __SYCL_GEN_OPT_BASE(op) \
193  friend id<dimensions> operator op(const id<dimensions> &lhs, \
194  const size_t &rhs) { \
195  id<dimensions> result; \
196  for (int i = 0; i < dimensions; ++i) { \
197  result.common_array[i] = lhs.common_array[i] op rhs; \
198  } \
199  return result; \
200  } \
201  friend id<dimensions> operator op(const size_t &lhs, \
202  const id<dimensions> &rhs) { \
203  id<dimensions> result; \
204  for (int i = 0; i < dimensions; ++i) { \
205  result.common_array[i] = lhs op rhs.common_array[i]; \
206  } \
207  return result; \
208  }
209 #endif // __SYCL_DISABLE_ID_TO_INT_CONV__
210 
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  __SYCL_GEN_OPT(>)
225  __SYCL_GEN_OPT(<=)
226  __SYCL_GEN_OPT(>=)
227 
228 #undef __SYCL_GEN_OPT
229 #undef __SYCL_GEN_OPT_BASE
230 
231 // OP is: +=, -=, *=, /=, %=, <<=, >>=, &=, |=, ^=
232 #define __SYCL_GEN_OPT(op) \
233  friend id<dimensions> &operator op(id<dimensions> &lhs, \
234  const id<dimensions> &rhs) { \
235  for (int i = 0; i < dimensions; ++i) { \
236  lhs.common_array[i] op rhs.common_array[i]; \
237  } \
238  return lhs; \
239  } \
240  friend id<dimensions> &operator op(id<dimensions> &lhs, const size_t &rhs) { \
241  for (int i = 0; i < dimensions; ++i) { \
242  lhs.common_array[i] op rhs; \
243  } \
244  return lhs; \
245  }
246 
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  __SYCL_GEN_OPT(&=)
255  __SYCL_GEN_OPT(|=)
256  __SYCL_GEN_OPT(^=)
257 
258 #undef __SYCL_GEN_OPT
259 
260 // OP is unary +, -
261 #define __SYCL_GEN_OPT(op) \
262  friend id<dimensions> operator op(const id<dimensions> &rhs) { \
263  id<dimensions> result; \
264  for (int i = 0; i < dimensions; ++i) { \
265  result.common_array[i] = (op rhs.common_array[i]); \
266  } \
267  return result; \
268  }
269 
270  __SYCL_GEN_OPT(+)
271  __SYCL_GEN_OPT(-)
272 
273 #undef __SYCL_GEN_OPT
274 
275 // OP is prefix ++, --
276 #define __SYCL_GEN_OPT(op) \
277  friend id<dimensions> &operator op(id<dimensions> &rhs) { \
278  for (int i = 0; i < dimensions; ++i) { \
279  op rhs.common_array[i]; \
280  } \
281  return rhs; \
282  }
283 
284  __SYCL_GEN_OPT(++)
285  __SYCL_GEN_OPT(--)
286 
287 #undef __SYCL_GEN_OPT
288 
289 // OP is postfix ++, --
290 #define __SYCL_GEN_OPT(op) \
291  friend id<dimensions> operator op(id<dimensions> &lhs, int) { \
292  id<dimensions> old_lhs; \
293  for (int i = 0; i < dimensions; ++i) { \
294  old_lhs.common_array[i] = lhs.common_array[i]; \
295  op lhs.common_array[i]; \
296  } \
297  return old_lhs; \
298  }
299 
300  __SYCL_GEN_OPT(++)
301  __SYCL_GEN_OPT(--)
302 
303 #undef __SYCL_GEN_OPT
304 
305 private:
306  // Friend to get access to private method set_allowed_range().
307  template <typename, int, typename> friend class detail::RoundedRangeKernel;
308  template <typename, int, typename>
309  friend class detail::RoundedRangeKernelWithKH;
310  void set_allowed_range(range<dimensions> rnwi) { (void)rnwi[0]; }
311 };
312 
313 namespace detail {
314 template <int dimensions>
315 size_t getOffsetForId(range<dimensions> Range, id<dimensions> Id,
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 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
375 } // namespace sycl
sycl::_V1::detail::InitializedVal
Definition: common.hpp:324
type_traits.hpp
sycl::_V1::make_error_code
std::error_code make_error_code(sycl::errc E) noexcept
Constructs an error code using e and sycl_category()
Definition: exception.cpp:92
sycl::_V1::id::id
id(ParamTy< N, 2, const item< dimensions, with_offset >> &item)
Definition: id.hpp:78
sycl::_V1::detail::declptr
T * declptr()
Definition: helpers.hpp:49
__SYCL_INLINE_VER_NAMESPACE
#define __SYCL_INLINE_VER_NAMESPACE(X)
Definition: defines_elementary.hpp:11
__SYCL_ALWAYS_INLINE
#define __SYCL_ALWAYS_INLINE
Definition: defines_elementary.hpp:25
array.hpp
helpers.hpp
sycl
---— Error handling, matching OpenCL plugin semantics.
Definition: access.hpp:14
operator==
bool operator==(const Slab &Lhs, const Slab &Rhs)
Definition: usm_allocator.cpp:313
sycl::_V1::id
A unique identifier of an item in an index space.
Definition: array.hpp:17
__SYCL_DEPRECATED
#define __SYCL_DEPRECATED(message)
Definition: defines_elementary.hpp:46
sycl::_V1::detail::enable_if_t
typename std::enable_if< B, T >::type enable_if_t
Definition: stl_type_traits.hpp:24
sycl::_V1::range
Defines the iteration domain of either a single work-group in a parallel dispatch,...
Definition: buffer.hpp:28
sycl::_V1::id::id
id(ParamTy< N, 1, size_t > dim0)
Definition: id.hpp:58
std::get
constexpr tuple_element< I, tuple< Types... > >::type & get(sycl::detail::tuple< Types... > &Arg) noexcept
Definition: tuple.hpp:199
sycl::_V1::item
Identifies an instance of the function object executing at each point in a range.
Definition: helpers.hpp:29
range.hpp
common.hpp
sycl::_V1::detail::getOffsetForId
size_t getOffsetForId(range< dimensions > Range, id< dimensions > Id, id< dimensions > Offset)
Definition: id.hpp:315
sycl::_V1::id::id
id(ParamTy< N, 1, const item< dimensions, with_offset >> &item)
Definition: id.hpp:65
sycl::_V1::id::id
id(ParamTy< N, 1, const range< dimensions >> &range_size)
Definition: id.hpp:61
sycl::_V1::id::id
id(ParamTy< N, 3, size_t > dim0, size_t dim1, size_t dim2)
Definition: id.hpp:84
__SYCL_ASSUME_INT
#define __SYCL_ASSUME_INT(x)
Definition: defines.hpp:18
sycl::_V1::id::id
id(ParamTy< N, 3, const item< dimensions, with_offset >> &item)
Definition: id.hpp:92
sycl::_V1::detail::getDelinearizedId
id< 3 > getDelinearizedId(const range< 3 > &Range, size_t Index)
Definition: id.hpp:333
sycl::_V1::detail::RoundedRangeKernel
Definition: handler.hpp:242
sycl::_V1::detail::RoundedRangeKernelWithKH
Definition: handler.hpp:260
sycl::_V1::id::id
id(ParamTy< N, 3, const range< dimensions >> &range_size)
Definition: id.hpp:88
sycl::_V1::ext::oneapi::experimental::this_id
id< Dims > this_id()
Definition: id.hpp:364
sycl::_V1::id::operator==
bool operator==(const id< dimensions > &rhs) const
Definition: id.hpp:123
sycl::_V1::id::id
id(ParamTy< N, 2, const range< dimensions >> &range_size)
Definition: id.hpp:74
sycl::_V1::id::id
id(ParamTy< N, 2, size_t > dim0, size_t dim1)
Definition: id.hpp:71
sycl::_V1::detail::array
Definition: array.hpp:21
__SYCL_GEN_OPT
#define __SYCL_GEN_OPT(op)
Definition: id.hpp:290