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 <CL/sycl/detail/array.hpp>
15 #include <CL/sycl/range.hpp>
16 
18 namespace sycl {
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(
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__
121 #if __cpp_impl_three_way_comparison < 201907
123 #endif
124 
125  /* Enable operators with integral types.
126  * Template operators take precedence than type conversion. In the case of
127  * non-template operators, ambiguity appears: "id op size_t" may refer
128  * "size_t op size_t" and "id op size_t". In case of template operators it
129  * will be "id op size_t"*/
130 #define __SYCL_GEN_OPT(op) \
131  template <typename T> \
132  EnableIfIntegral<T, bool> operator op(const T &rhs) const { \
133  if (this->common_array[0] != rhs) \
134  return false op true; \
135  return true op true; \
136  } \
137  template <typename T> \
138  friend EnableIfIntegral<T, bool> operator op(const T &lhs, \
139  const id<dimensions> &rhs) { \
140  if (lhs != rhs.common_array[0]) \
141  return false op true; \
142  return true op true; \
143  }
144 
145  __SYCL_GEN_OPT(==)
146  __SYCL_GEN_OPT(!=)
147 
148 #undef __SYCL_GEN_OPT
149 
150 #endif // __SYCL_DISABLE_ID_TO_INT_CONV__
151 
152 // OP is: +, -, *, /, %, <<, >>, &, |, ^, &&, ||, <, >, <=, >=
153 #define __SYCL_GEN_OPT_BASE(op) \
154  friend id<dimensions> operator op(const id<dimensions> &lhs, \
155  const id<dimensions> &rhs) { \
156  id<dimensions> result; \
157  for (int i = 0; i < dimensions; ++i) { \
158  result.common_array[i] = lhs.common_array[i] op rhs.common_array[i]; \
159  } \
160  return result; \
161  }
162 
163 #ifndef __SYCL_DISABLE_ID_TO_INT_CONV__
164 // Enable operators with integral types only
165 #define __SYCL_GEN_OPT(op) \
166  __SYCL_GEN_OPT_BASE(op) \
167  template <typename T> \
168  friend EnableIfIntegral<T, id<dimensions>> operator op( \
169  const id<dimensions> &lhs, const T &rhs) { \
170  id<dimensions> result; \
171  for (int i = 0; i < dimensions; ++i) { \
172  result.common_array[i] = lhs.common_array[i] op rhs; \
173  } \
174  return result; \
175  } \
176  template <typename T> \
177  friend EnableIfIntegral<T, id<dimensions>> operator op( \
178  const T &lhs, const id<dimensions> &rhs) { \
179  id<dimensions> result; \
180  for (int i = 0; i < dimensions; ++i) { \
181  result.common_array[i] = lhs op rhs.common_array[i]; \
182  } \
183  return result; \
184  }
185 #else
186 #define __SYCL_GEN_OPT(op) \
187  __SYCL_GEN_OPT_BASE(op) \
188  friend id<dimensions> operator op(const id<dimensions> &lhs, \
189  const size_t &rhs) { \
190  id<dimensions> result; \
191  for (int i = 0; i < dimensions; ++i) { \
192  result.common_array[i] = lhs.common_array[i] op rhs; \
193  } \
194  return result; \
195  } \
196  friend id<dimensions> operator op(const size_t &lhs, \
197  const id<dimensions> &rhs) { \
198  id<dimensions> result; \
199  for (int i = 0; i < dimensions; ++i) { \
200  result.common_array[i] = lhs op rhs.common_array[i]; \
201  } \
202  return result; \
203  }
204 #endif // __SYCL_DISABLE_ID_TO_INT_CONV__
205 
206  __SYCL_GEN_OPT(+)
207  __SYCL_GEN_OPT(-)
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 
223 #undef __SYCL_GEN_OPT
224 #undef __SYCL_GEN_OPT_BASE
225 
226 // OP is: +=, -=, *=, /=, %=, <<=, >>=, &=, |=, ^=
227 #define __SYCL_GEN_OPT(op) \
228  friend id<dimensions> &operator op(id<dimensions> &lhs, \
229  const id<dimensions> &rhs) { \
230  for (int i = 0; i < dimensions; ++i) { \
231  lhs.common_array[i] op rhs.common_array[i]; \
232  } \
233  return lhs; \
234  } \
235  friend id<dimensions> &operator op(id<dimensions> &lhs, const size_t &rhs) { \
236  for (int i = 0; i < dimensions; ++i) { \
237  lhs.common_array[i] op rhs; \
238  } \
239  return lhs; \
240  }
241 
242  __SYCL_GEN_OPT(+=)
243  __SYCL_GEN_OPT(-=)
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 
253 #undef __SYCL_GEN_OPT
254 
255 // OP is unary +, -
256 #define __SYCL_GEN_OPT(op) \
257  friend id<dimensions> operator op(const id<dimensions> &rhs) { \
258  id<dimensions> result; \
259  for (int i = 0; i < dimensions; ++i) { \
260  result.common_array[i] = (op rhs.common_array[i]); \
261  } \
262  return result; \
263  }
264 
265  __SYCL_GEN_OPT(+)
266  __SYCL_GEN_OPT(-)
267 
268 #undef __SYCL_GEN_OPT
269 
270 // OP is prefix ++, --
271 #define __SYCL_GEN_OPT(op) \
272  friend id<dimensions> &operator op(id<dimensions> &rhs) { \
273  for (int i = 0; i < dimensions; ++i) { \
274  op rhs.common_array[i]; \
275  } \
276  return rhs; \
277  }
278 
279  __SYCL_GEN_OPT(++)
280  __SYCL_GEN_OPT(--)
281 
282 #undef __SYCL_GEN_OPT
283 
284 // OP is postfix ++, --
285 #define __SYCL_GEN_OPT(op) \
286  friend id<dimensions> operator op(id<dimensions> &lhs, int) { \
287  id<dimensions> old_lhs; \
288  for (int i = 0; i < dimensions; ++i) { \
289  old_lhs.common_array[i] = lhs.common_array[i]; \
290  op lhs.common_array[i]; \
291  } \
292  return old_lhs; \
293  }
294 
295  __SYCL_GEN_OPT(++)
296  __SYCL_GEN_OPT(--)
297 
298 #undef __SYCL_GEN_OPT
299 
300 private:
301  // Friend to get access to private method set_allowed_range().
302  template <typename, int, typename> friend class detail::RoundedRangeKernel;
303  template <typename, int, typename>
305  void set_allowed_range(range<dimensions> rnwi) { (void)rnwi[0]; }
306 };
307 
308 namespace detail {
309 template <int dimensions>
311  id<dimensions> Offset) {
312  size_t offset = 0;
313  for (int i = 0; i < dimensions; ++i)
314  offset = offset * Range[i] + Offset[i] + Id[i];
315  return offset;
316 }
317 
318 inline id<1> getDelinearizedId(const range<1> &, size_t Index) {
319  return {Index};
320 }
321 
322 inline id<2> getDelinearizedId(const range<2> &Range, size_t Index) {
323  size_t X = Index % Range[1];
324  size_t Y = Index / Range[1];
325  return {Y, X};
326 }
327 
328 inline id<3> getDelinearizedId(const range<3> &Range, size_t Index) {
329  size_t D1D2 = Range[1] * Range[2];
330  size_t Z = Index / D1D2;
331  size_t ZRest = Index % D1D2;
332  size_t Y = ZRest / Range[2];
333  size_t X = ZRest % Range[2];
334  return {Z, Y, X};
335 }
336 } // namespace detail
337 
338 // C++ feature test macros are supported by all supported compilers
339 // with the exception of MSVC 1914. It doesn't support deduction guides.
340 #ifdef __cpp_deduction_guides
341 id(size_t)->id<1>;
342 id(size_t, size_t)->id<2>;
343 id(size_t, size_t, size_t)->id<3>;
344 #endif
345 
346 template <int Dims>
347 __SYCL_DEPRECATED("use sycl::ext::oneapi::experimental::this_id() instead")
348 id<Dims> this_id() {
349 #ifdef __SYCL_DEVICE_ONLY__
350  return detail::Builder::getElement(detail::declptr<id<Dims>>());
351 #else
352  throw sycl::exception(
353  sycl::make_error_code(sycl::errc::feature_not_supported),
354  "Free function calls are not supported on host device");
355 #endif
356 }
357 
358 namespace ext {
359 namespace oneapi {
360 namespace experimental {
361 template <int Dims> id<Dims> this_id() {
362 #ifdef __SYCL_DEVICE_ONLY__
363  return sycl::detail::Builder::getElement(sycl::detail::declptr<id<Dims>>());
364 #else
365  throw sycl::exception(
366  sycl::make_error_code(sycl::errc::feature_not_supported),
367  "Free function calls are not supported on host device");
368 #endif
369 }
370 } // namespace experimental
371 } // namespace oneapi
372 } // namespace ext
373 } // namespace sycl
374 } // __SYCL_INLINE_NAMESPACE(cl)
cl::sycl::id::id
id(ParamTy< N, 2, size_t > dim0, size_t dim1)
Definition: id.hpp:71
cl::sycl::id::id
id(ParamTy< N, 1, const item< dimensions, with_offset >> &item)
Definition: id.hpp:65
cl::sycl::detail::get< 0 >
Definition: tuple.hpp:75
T
type_traits.hpp
cl::sycl::item
Identifies an instance of the function object executing at each point in a range.
Definition: helpers.hpp:28
cl::sycl::id
A unique identifier of an item in an index space.
Definition: array.hpp:17
cl::sycl::id::id
id(ParamTy< N, 2, const item< dimensions, with_offset >> &item)
Definition: id.hpp:78
__SYCL_DEPRECATED
#define __SYCL_DEPRECATED(message)
Definition: defines_elementary.hpp:47
cl::sycl::detail::declptr
T * declptr()
Definition: helpers.hpp:56
cl::sycl::detail::InitializedVal
Definition: common.hpp:227
array.hpp
cl::sycl::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:121
helpers.hpp
sycl
Definition: invoke_simd.hpp:68
cl::sycl::id::id
id(ParamTy< N, 3, size_t > dim0, size_t dim1, size_t dim2)
Definition: id.hpp:84
cl::sycl::range
Defines the iteration domain of either a single work-group in a parallel dispatch,...
Definition: buffer.hpp:24
cl::sycl::detail::get
Definition: tuple.hpp:59
cl::sycl::id::id
id(ParamTy< N, 2, const range< dimensions >> &range_size)
Definition: id.hpp:74
__SYCL_ALWAYS_INLINE
#define __SYCL_ALWAYS_INLINE
Definition: defines_elementary.hpp:29
cl::sycl::detail::RoundedRangeKernel
Definition: handler.hpp:204
range.hpp
cl::sycl::detail::getOffsetForId
size_t getOffsetForId(range< dimensions > Range, id< dimensions > Id, id< dimensions > Offset)
Definition: id.hpp:310
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
cl::sycl::id::id
id(ParamTy< N, 3, const item< dimensions, with_offset >> &item)
Definition: id.hpp:92
cl::sycl::detail::array
Definition: array.hpp:21
__SYCL_ASSUME_INT
#define __SYCL_ASSUME_INT(x)
Definition: defines.hpp:18
cl::sycl::detail::RoundedRangeKernelWithKH
Definition: handler.hpp:222
std::get
constexpr tuple_element< I, tuple< Types... > >::type & get(cl::sycl::detail::tuple< Types... > &Arg) noexcept
Definition: tuple.hpp:199
cl::sycl::id::id
id(ParamTy< N, 1, size_t > dim0)
Definition: id.hpp:58
cl::sycl::ext::oneapi::experimental::this_id
id< Dims > this_id()
Definition: id.hpp:361
cl::sycl::id::id
id(ParamTy< N, 1, const range< dimensions >> &range_size)
Definition: id.hpp:61
common.hpp
cl::sycl::exception
Definition: exception.hpp:64
cl::sycl::detail::getDelinearizedId
id< 3 > getDelinearizedId(const range< 3 > &Range, size_t Index)
Definition: id.hpp:328
cl::sycl::detail::enable_if_t
typename std::enable_if< B, T >::type enable_if_t
Definition: stl_type_traits.hpp:24
cl::sycl::id::id
id(ParamTy< N, 3, const range< dimensions >> &range_size)
Definition: id.hpp:88
__SYCL_GEN_OPT
#define __SYCL_GEN_OPT(op)
Definition: id.hpp:285
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12