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 #include <CL/sycl/detail/array.hpp>
11 #include <CL/sycl/detail/helpers.hpp>
12 
13 #include <stdexcept>
14 #include <type_traits>
15 
16 __SYCL_INLINE_NAMESPACE(cl) {
17 namespace sycl {
18 template <int dimensions> class id;
19 
24 template <int dimensions = 1> class range : public detail::array<dimensions> {
25  static_assert(dimensions >= 1 && dimensions <= 3,
26  "range can only be 1, 2, or 3 dimensional.");
27  using base = detail::array<dimensions>;
28  template <typename N, typename T>
29  using IntegralType = detail::enable_if_t<std::is_integral<N>::value, T>;
30 
31 public:
32  /* The following constructor is only available in the range class
33  specialization where: dimensions==1 */
34  template <int N = dimensions>
35  range(typename detail::enable_if_t<(N == 1), size_t> dim0) : base(dim0) {}
36 
37  /* The following constructor is only available in the range class
38  specialization where: dimensions==2 */
39  template <int N = dimensions>
40  range(typename detail::enable_if_t<(N == 2), size_t> dim0, size_t dim1)
41  : base(dim0, dim1) {}
42 
43  /* The following constructor is only available in the range class
44  specialization where: dimensions==3 */
45  template <int N = dimensions>
46  range(typename detail::enable_if_t<(N == 3), size_t> dim0, size_t dim1,
47  size_t dim2)
48  : base(dim0, dim1, dim2) {}
49 
50  size_t size() const {
51  size_t size = 1;
52  for (int i = 0; i < dimensions; ++i) {
53  size *= this->get(i);
54  }
55  return size;
56  }
57 
58  range(const range<dimensions> &rhs) = default;
59  range(range<dimensions> &&rhs) = default;
60  range<dimensions> &operator=(const range<dimensions> &rhs) = default;
61  range<dimensions> &operator=(range<dimensions> &&rhs) = default;
62  range() = delete;
63 
64 // OP is: +, -, *, /, %, <<, >>, &, |, ^, &&, ||, <, >, <=, >=
65 #define __SYCL_GEN_OPT_BASE(op) \
66  friend range<dimensions> operator op(const range<dimensions> &lhs, \
67  const range<dimensions> &rhs) { \
68  range<dimensions> result(lhs); \
69  for (int i = 0; i < dimensions; ++i) { \
70  result.common_array[i] = lhs.common_array[i] op rhs.common_array[i]; \
71  } \
72  return result; \
73  }
74 
75 #ifndef __SYCL_DISABLE_ID_TO_INT_CONV__
76  // Enable operators with integral types only
77 #define __SYCL_GEN_OPT(op) \
78  __SYCL_GEN_OPT_BASE(op) \
79  template <typename T> \
80  friend IntegralType<T, range<dimensions>> operator op( \
81  const range<dimensions> &lhs, const T &rhs) { \
82  range<dimensions> result(lhs); \
83  for (int i = 0; i < dimensions; ++i) { \
84  result.common_array[i] = lhs.common_array[i] op rhs; \
85  } \
86  return result; \
87  } \
88  template <typename T> \
89  friend IntegralType<T, range<dimensions>> operator op( \
90  const T &lhs, const range<dimensions> &rhs) { \
91  range<dimensions> result(rhs); \
92  for (int i = 0; i < dimensions; ++i) { \
93  result.common_array[i] = lhs op rhs.common_array[i]; \
94  } \
95  return result; \
96  }
97 #else
98 #define __SYCL_GEN_OPT(op) \
99  __SYCL_GEN_OPT_BASE(op) \
100  friend range<dimensions> operator op(const range<dimensions> &lhs, \
101  const size_t &rhs) { \
102  range<dimensions> result(lhs); \
103  for (int i = 0; i < dimensions; ++i) { \
104  result.common_array[i] = lhs.common_array[i] op rhs; \
105  } \
106  return result; \
107  } \
108  friend range<dimensions> operator op(const size_t &lhs, \
109  const range<dimensions> &rhs) { \
110  range<dimensions> result(rhs); \
111  for (int i = 0; i < dimensions; ++i) { \
112  result.common_array[i] = lhs op rhs.common_array[i]; \
113  } \
114  return result; \
115  }
116 #endif // __SYCL_DISABLE_ID_TO_INT_CONV__
117 
118  __SYCL_GEN_OPT(+)
119  __SYCL_GEN_OPT(-)
120  __SYCL_GEN_OPT(*)
121  __SYCL_GEN_OPT(/)
122  __SYCL_GEN_OPT(%)
123  __SYCL_GEN_OPT(<<)
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 
135 #undef __SYCL_GEN_OPT
136 #undef __SYCL_GEN_OPT_BASE
137 
138 // OP is: +=, -=, *=, /=, %=, <<=, >>=, &=, |=, ^=
139 #define __SYCL_GEN_OPT(op) \
140  friend range<dimensions> &operator op(range<dimensions> &lhs, \
141  const range<dimensions> &rhs) { \
142  for (int i = 0; i < dimensions; ++i) { \
143  lhs.common_array[i] op rhs[i]; \
144  } \
145  return lhs; \
146  } \
147  friend range<dimensions> &operator op(range<dimensions> &lhs, \
148  const size_t &rhs) { \
149  for (int i = 0; i < dimensions; ++i) { \
150  lhs.common_array[i] op rhs; \
151  } \
152  return lhs; \
153  }
154 
155  __SYCL_GEN_OPT(+=)
156  __SYCL_GEN_OPT(-=)
157  __SYCL_GEN_OPT(*=)
158  __SYCL_GEN_OPT(/=)
159  __SYCL_GEN_OPT(%=)
160  __SYCL_GEN_OPT(<<=)
161  __SYCL_GEN_OPT(>>=)
162  __SYCL_GEN_OPT(&=)
163  __SYCL_GEN_OPT(|=)
164  __SYCL_GEN_OPT(^=)
165 
166 #undef __SYCL_GEN_OPT
167 
168 // OP is unary +, -
169 #define __SYCL_GEN_OPT(op) \
170  friend range<dimensions> operator op(const range<dimensions> &rhs) { \
171  range<dimensions> result(rhs); \
172  for (int i = 0; i < dimensions; ++i) { \
173  result.common_array[i] = (op rhs.common_array[i]); \
174  } \
175  return result; \
176  }
177 
178  __SYCL_GEN_OPT(+)
179  __SYCL_GEN_OPT(-)
180 
181 #undef __SYCL_GEN_OPT
182 
183 // OP is prefix ++, --
184 #define __SYCL_GEN_OPT(op) \
185  friend range<dimensions> &operator op(range<dimensions> &rhs) { \
186  for (int i = 0; i < dimensions; ++i) { \
187  op rhs.common_array[i]; \
188  } \
189  return rhs; \
190  }
191 
192  __SYCL_GEN_OPT(++)
193  __SYCL_GEN_OPT(--)
194 
195 #undef __SYCL_GEN_OPT
196 
197 // OP is postfix ++, --
198 #define __SYCL_GEN_OPT(op) \
199  friend range<dimensions> operator op(range<dimensions> &lhs, int) { \
200  range<dimensions> old_lhs(lhs); \
201  for (int i = 0; i < dimensions; ++i) { \
202  op lhs.common_array[i]; \
203  } \
204  return old_lhs; \
205  }
206 
207  __SYCL_GEN_OPT(++)
208  __SYCL_GEN_OPT(--)
209 
210 #undef __SYCL_GEN_OPT
211 
212 private:
213  friend class handler;
214  friend class detail::Builder;
215 
216  // Adjust the first dim of the range
217  void set_range_dim0(const size_t dim0) { this->common_array[0] = dim0; }
218 };
219 
220 #ifdef __cpp_deduction_guides
221 range(size_t)->range<1>;
222 range(size_t, size_t)->range<2>;
223 range(size_t, size_t, size_t)->range<3>;
224 #endif
225 
226 } // namespace sycl
227 } // __SYCL_INLINE_NAMESPACE(cl)
cl::sycl::range::range
range(typename detail::enable_if_t<(N==3), size_t > dim0, size_t dim1, size_t dim2)
Definition: range.hpp:46
cl::sycl::detail::Builder
Definition: helpers.hpp:68
T
cl::sycl::range::range
range(typename detail::enable_if_t<(N==2), size_t > dim0, size_t dim1)
Definition: range.hpp:40
array.hpp
helpers.hpp
sycl
Definition: invoke_simd.hpp:68
cl::sycl::range
Defines the iteration domain of either a single work-group in a parallel dispatch,...
Definition: buffer.hpp:24
cl::sycl::range::size
size_t size() const
Definition: range.hpp:50
__SYCL_GEN_OPT
#define __SYCL_GEN_OPT(op)
Definition: range.hpp:198
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
cl::sycl::detail::array
Definition: array.hpp:21
cl::sycl::handler
Command group handler class.
Definition: handler.hpp:362
cl::sycl::range::range
range(typename detail::enable_if_t<(N==1), size_t > dim0)
Definition: range.hpp:35
std::get
constexpr tuple_element< I, tuple< Types... > >::type & get(cl::sycl::detail::tuple< Types... > &Arg) noexcept
Definition: tuple.hpp:199
cl::sycl::detail::enable_if_t
typename std::enable_if< B, T >::type enable_if_t
Definition: stl_type_traits.hpp:24
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12