DPC++ Runtime
Runtime libraries for oneAPI DPC++
range.hpp
Go to the documentation of this file.
1 //==----------- range.hpp --- SYCL iteration range -------------------------==//
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/helpers.hpp> // for Builder
13 
14 #include <array> // for array
15 #include <stddef.h> // for size_t
16 #include <type_traits> // for enable_if_t
17 
18 namespace sycl {
19 inline namespace _V1 {
20 template <int Dimensions> class id;
21 
26 template <int Dimensions = 1> class range : public detail::array<Dimensions> {
27 public:
28  static constexpr int dimensions = Dimensions;
29 
30 private:
31  static_assert(Dimensions >= 1 && Dimensions <= 3,
32  "range can only be 1, 2, or 3 Dimensional.");
33  using base = detail::array<Dimensions>;
34  template <typename N, typename T>
35  using IntegralType = std::enable_if_t<std::is_integral_v<N>, T>;
36 
37 public:
38  /* The following constructor is only available in the range class
39  specialization where: Dimensions==1 */
40  template <int N = Dimensions>
41  range(typename std::enable_if_t<(N == 1), size_t> dim0) : base(dim0) {}
42 
43  /* The following constructor is only available in the range class
44  specialization where: Dimensions==2 */
45  template <int N = Dimensions>
46  range(typename std::enable_if_t<(N == 2), size_t> dim0, size_t dim1)
47  : base(dim0, dim1) {}
48 
49  /* The following constructor is only available in the range class
50  specialization where: Dimensions==3 */
51  template <int N = Dimensions>
52  range(typename std::enable_if_t<(N == 3), size_t> dim0, size_t dim1,
53  size_t dim2)
54  : base(dim0, dim1, dim2) {}
55 
56  size_t size() const {
57  size_t size = 1;
58  for (int i = 0; i < Dimensions; ++i) {
59  size *= this->get(i);
60  }
61  return size;
62  }
63 
64  range(const range<Dimensions> &rhs) = default;
65  range(range<Dimensions> &&rhs) = default;
66  range<Dimensions> &operator=(const range<Dimensions> &rhs) = default;
67  range<Dimensions> &operator=(range<Dimensions> &&rhs) = default;
68  range() = default;
69 
70 // OP is: +, -, *, /, %, <<, >>, &, |, ^, &&, ||, <, >, <=, >=
71 #define __SYCL_GEN_OPT_BASE(op) \
72  friend range<Dimensions> operator op(const range<Dimensions> &lhs, \
73  const range<Dimensions> &rhs) { \
74  range<Dimensions> result(lhs); \
75  for (int i = 0; i < Dimensions; ++i) { \
76  result.common_array[i] = lhs.common_array[i] op rhs.common_array[i]; \
77  } \
78  return result; \
79  }
80 
81 #ifndef __SYCL_DISABLE_ID_TO_INT_CONV__
82  // Enable operators with integral types only
83 #define __SYCL_GEN_OPT(op) \
84  __SYCL_GEN_OPT_BASE(op) \
85  template <typename T> \
86  friend IntegralType<T, range<Dimensions>> operator op( \
87  const range<Dimensions> &lhs, const T &rhs) { \
88  range<Dimensions> result(lhs); \
89  for (int i = 0; i < Dimensions; ++i) { \
90  result.common_array[i] = lhs.common_array[i] op rhs; \
91  } \
92  return result; \
93  } \
94  template <typename T> \
95  friend IntegralType<T, range<Dimensions>> operator op( \
96  const T &lhs, const range<Dimensions> &rhs) { \
97  range<Dimensions> result(rhs); \
98  for (int i = 0; i < Dimensions; ++i) { \
99  result.common_array[i] = lhs op rhs.common_array[i]; \
100  } \
101  return result; \
102  }
103 #else
104 #define __SYCL_GEN_OPT(op) \
105  __SYCL_GEN_OPT_BASE(op) \
106  friend range<Dimensions> operator op(const range<Dimensions> &lhs, \
107  const size_t &rhs) { \
108  range<Dimensions> result(lhs); \
109  for (int i = 0; i < Dimensions; ++i) { \
110  result.common_array[i] = lhs.common_array[i] op rhs; \
111  } \
112  return result; \
113  } \
114  friend range<Dimensions> operator op(const size_t &lhs, \
115  const range<Dimensions> &rhs) { \
116  range<Dimensions> result(rhs); \
117  for (int i = 0; i < Dimensions; ++i) { \
118  result.common_array[i] = lhs op rhs.common_array[i]; \
119  } \
120  return result; \
121  }
122 #endif // __SYCL_DISABLE_ID_TO_INT_CONV__
123 
124  __SYCL_GEN_OPT(+)
125  __SYCL_GEN_OPT(-)
126  __SYCL_GEN_OPT(*)
127  __SYCL_GEN_OPT(/)
128  __SYCL_GEN_OPT(%)
129  __SYCL_GEN_OPT(<<)
130  __SYCL_GEN_OPT(>>)
131  __SYCL_GEN_OPT(&)
132  __SYCL_GEN_OPT(|)
133  __SYCL_GEN_OPT(^)
134  __SYCL_GEN_OPT(&&)
135  __SYCL_GEN_OPT(||)
136  __SYCL_GEN_OPT(<)
137  __SYCL_GEN_OPT(>)
138  __SYCL_GEN_OPT(<=)
139  __SYCL_GEN_OPT(>=)
140 
141 #undef __SYCL_GEN_OPT
142 #undef __SYCL_GEN_OPT_BASE
143 
144 // OP is: +=, -=, *=, /=, %=, <<=, >>=, &=, |=, ^=
145 #define __SYCL_GEN_OPT(op) \
146  friend range<Dimensions> &operator op(range<Dimensions> &lhs, \
147  const range<Dimensions> &rhs) { \
148  for (int i = 0; i < Dimensions; ++i) { \
149  lhs.common_array[i] op rhs[i]; \
150  } \
151  return lhs; \
152  } \
153  friend range<Dimensions> &operator op(range<Dimensions> &lhs, \
154  const size_t &rhs) { \
155  for (int i = 0; i < Dimensions; ++i) { \
156  lhs.common_array[i] op rhs; \
157  } \
158  return lhs; \
159  }
160 
161  __SYCL_GEN_OPT(+=)
162  __SYCL_GEN_OPT(-=)
163  __SYCL_GEN_OPT(*=)
164  __SYCL_GEN_OPT(/=)
165  __SYCL_GEN_OPT(%=)
166  __SYCL_GEN_OPT(<<=)
167  __SYCL_GEN_OPT(>>=)
168  __SYCL_GEN_OPT(&=)
169  __SYCL_GEN_OPT(|=)
170  __SYCL_GEN_OPT(^=)
171 
172 #undef __SYCL_GEN_OPT
173 
174 // OP is unary +, -
175 #define __SYCL_GEN_OPT(op) \
176  friend range<Dimensions> operator op(const range<Dimensions> &rhs) { \
177  range<Dimensions> result(rhs); \
178  for (int i = 0; i < Dimensions; ++i) { \
179  result.common_array[i] = (op rhs.common_array[i]); \
180  } \
181  return result; \
182  }
183 
184  __SYCL_GEN_OPT(+)
185  __SYCL_GEN_OPT(-)
186 
187 #undef __SYCL_GEN_OPT
188 
189 // OP is prefix ++, --
190 #define __SYCL_GEN_OPT(op) \
191  friend range<Dimensions> &operator op(range<Dimensions> &rhs) { \
192  for (int i = 0; i < Dimensions; ++i) { \
193  op rhs.common_array[i]; \
194  } \
195  return rhs; \
196  }
197 
198  __SYCL_GEN_OPT(++)
199  __SYCL_GEN_OPT(--)
200 
201 #undef __SYCL_GEN_OPT
202 
203 // OP is postfix ++, --
204 #define __SYCL_GEN_OPT(op) \
205  friend range<Dimensions> operator op(range<Dimensions> &lhs, int) { \
206  range<Dimensions> old_lhs(lhs); \
207  for (int i = 0; i < Dimensions; ++i) { \
208  op lhs.common_array[i]; \
209  } \
210  return old_lhs; \
211  }
212 
213  __SYCL_GEN_OPT(++)
214  __SYCL_GEN_OPT(--)
215 
216 #undef __SYCL_GEN_OPT
217 
218 private:
219  friend class handler;
220  friend class detail::Builder;
221 
222  // Adjust the first dim of the range
223  void set_range_dim0(const size_t dim0) { this->common_array[0] = dim0; }
224 };
225 
226 #ifdef __cpp_deduction_guides
227 range(size_t)->range<1>;
228 range(size_t, size_t)->range<2>;
229 range(size_t, size_t, size_t)->range<3>;
230 #endif
231 
232 namespace detail {
233 // XPTI helpers for creating array from a range.
234 inline std::array<size_t, 3> rangeToArray(const range<3> &r) {
235  return {r[0], r[1], r[2]};
236 }
237 inline std::array<size_t, 3> rangeToArray(const range<2> &r) {
238  return {r[0], r[1], 0};
239 }
240 inline std::array<size_t, 3> rangeToArray(const range<1> &r) {
241  return {r[0], 0, 0};
242 }
243 } // namespace detail
244 
245 } // namespace _V1
246 } // namespace sycl
sycl::_V1::detail::Builder
Definition: helpers.hpp:68
sycl::_V1::detail::array
Definition: array.hpp:24
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::handler
Command group handler class.
Definition: handler.hpp:458
sycl::_V1::range
Defines the iteration domain of either a single work-group in a parallel dispatch,...
Definition: range.hpp:26
sycl::_V1::range::operator=
range< Dimensions > & operator=(range< Dimensions > &&rhs)=default
sycl::_V1::range::size
size_t size() const
Definition: range.hpp:56
sycl::_V1::range::range
range(typename std::enable_if_t<(N==3), size_t > dim0, size_t dim1, size_t dim2)
Definition: range.hpp:52
sycl::_V1::range::range
range(typename std::enable_if_t<(N==2), size_t > dim0, size_t dim1)
Definition: range.hpp:46
sycl::_V1::range::range
range(typename std::enable_if_t<(N==1), size_t > dim0)
Definition: range.hpp:41
sycl::_V1::range::range
range()=default
sycl::_V1::range::range
range(range< Dimensions > &&rhs)=default
sycl::_V1::range::operator=
range< Dimensions > & operator=(const range< Dimensions > &rhs)=default
sycl::_V1::range::range
range(const range< Dimensions > &rhs)=default
sycl::_V1::range::dimensions
static constexpr int dimensions
Definition: range.hpp:28
sycl::_V1::detail::rangeToArray
std::array< size_t, 3 > rangeToArray(const range< 3 > &r)
Definition: range.hpp:234
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::_V1::image_channel_order::r
@ r
sycl
Definition: access.hpp:18
__SYCL_GEN_OPT
#define __SYCL_GEN_OPT(op)
Definition: range.hpp:204