Skip to content

Commit 5eebd1e

Browse files
iburyliburyl
authored
[SYCL] Add marray class as defined by SYCL 2020 provisional (#2928)
Signed-off-by: iburylov <[email protected]>
1 parent 6327221 commit 5eebd1e

File tree

3 files changed

+437
-0
lines changed

3 files changed

+437
-0
lines changed

‎sycl/include/CL/sycl.hpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -33,6 +33,7 @@
3333
#include <CL/sycl/image.hpp>
3434
#include <CL/sycl/item.hpp>
3535
#include <CL/sycl/kernel.hpp>
36+
#include <CL/sycl/marray.hpp>
3637
#include <CL/sycl/multi_ptr.hpp>
3738
#include <CL/sycl/nd_item.hpp>
3839
#include <CL/sycl/nd_range.hpp>

‎sycl/include/CL/sycl/marray.hpp

Lines changed: 349 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,349 @@
1+
//==----------------- marray.hpp --- Implements marray classes -------------==//
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/aliases.hpp>
12+
#include <CL/sycl/detail/common.hpp>
13+
#include <CL/sycl/detail/generic_type_traits.hpp>
14+
#include <CL/sycl/detail/type_traits.hpp>
15+
#include <CL/sycl/half_type.hpp>
16+
17+
__SYCL_INLINE_NAMESPACE(cl) {
18+
namespace sycl {
19+
20+
/// Provides a cross-patform math array class template that works on
21+
/// SYCL devices as well as in host C++ code.
22+
///
23+
/// \ingroup sycl_api
24+
template <typename Type, std::size_t NumElements> class marray {
25+
using DataT = Type;
26+
27+
public:
28+
using value_type = Type;
29+
using reference = Type &;
30+
using const_reference = const Type &;
31+
using iterator = Type *;
32+
using const_iterator = const Type *;
33+
34+
private:
35+
value_type MData[NumElements];
36+
37+
template <class...> struct conjunction : std::true_type {};
38+
template <class B1, class... tail>
39+
struct conjunction<B1, tail...>
40+
: std::conditional<bool(B1::value), conjunction<tail...>, B1>::type {};
41+
42+
// TypeChecker is needed for (const ArgTN &... Args) ctor to validate Args.
43+
template <typename T, typename DataT_>
44+
struct TypeChecker : std::is_convertible<T, DataT_> {};
45+
46+
// Shortcuts for Args validation in (const ArgTN &... Args) ctor.
47+
template <typename... ArgTN>
48+
using EnableIfSuitableTypes = typename std::enable_if<
49+
conjunction<TypeChecker<ArgTN, DataT>...>::value>::type;
50+
51+
public:
52+
marray() : MData{} {}
53+
54+
explicit marray(const Type &Arg) {
55+
for (std::size_t I = 0; I < NumElements; ++I) {
56+
MData[I] = Arg;
57+
}
58+
}
59+
60+
template <
61+
typename... ArgTN, typename = EnableIfSuitableTypes<ArgTN...>,
62+
typename = typename std::enable_if<sizeof...(ArgTN) == NumElements>::type>
63+
marray(const ArgTN &... Args) : MData{Args...} {}
64+
65+
marray(const marray<Type, NumElements> &Rhs) {
66+
for (std::size_t I = 0; I < NumElements; ++I) {
67+
MData[I] = Rhs.MData[I];
68+
}
69+
}
70+
71+
marray(marray<Type, NumElements> &&Rhs) {
72+
for (std::size_t I = 0; I < NumElements; ++I) {
73+
MData[I] = Rhs.MData[I];
74+
}
75+
}
76+
77+
// Available only when: NumElements == 1
78+
template <std::size_t Size = NumElements,
79+
typename = typename std::enable_if<Size == 1>>
80+
operator Type() const {
81+
return MData[0];
82+
}
83+
84+
static constexpr std::size_t size() noexcept { return NumElements; }
85+
86+
// subscript operator
87+
reference operator[](std::size_t index) { return MData[index]; }
88+
89+
const_reference operator[](std::size_t index) const { return MData[index]; }
90+
91+
marray &operator=(const marray<Type, NumElements> &Rhs) {
92+
for (std::size_t I = 0; I < NumElements; ++I) {
93+
MData[I] = Rhs.MData[I];
94+
}
95+
return *this;
96+
}
97+
98+
// broadcasting operator
99+
marray &operator=(const Type &Rhs) {
100+
for (std::size_t I = 0; I < NumElements; ++I) {
101+
MData[I] = Rhs;
102+
}
103+
return *this;
104+
}
105+
106+
// iterator functions
107+
iterator begin() { return MData; }
108+
109+
const_iterator begin() const { return MData; }
110+
111+
iterator end() { return MData + NumElements; }
112+
113+
const_iterator end() const { return MData + NumElements; }
114+
115+
#ifdef __SYCL_BINOP
116+
#error "Undefine __SYCL_BINOP macro"
117+
#endif
118+
119+
#ifdef __SYCL_BINOP_INTEGRAL
120+
#error "Undefine __SYCL_BINOP_INTEGRAL macro"
121+
#endif
122+
123+
#define __SYCL_BINOP(BINOP, OPASSIGN) \
124+
friend marray operator BINOP(const marray &Lhs, const marray &Rhs) { \
125+
marray Ret; \
126+
for (size_t I = 0; I < NumElements; ++I) { \
127+
Ret[I] = Lhs[I] BINOP Rhs[I]; \
128+
} \
129+
return Ret; \
130+
} \
131+
template <typename T> \
132+
friend typename std::enable_if< \
133+
std::is_convertible<DataT, T>::value && \
134+
(std::is_fundamental<T>::value || \
135+
std::is_same<typename std::remove_const<T>::type, half>::value), \
136+
marray>::type \
137+
operator BINOP(const marray &Lhs, const T &Rhs) { \
138+
return Lhs BINOP marray(static_cast<DataT>(Rhs)); \
139+
} \
140+
friend marray &operator OPASSIGN(marray &Lhs, const marray &Rhs) { \
141+
Lhs = Lhs BINOP Rhs; \
142+
return Lhs; \
143+
} \
144+
template <std::size_t Num = NumElements> \
145+
friend typename std::enable_if<Num != 1, marray &>::type operator OPASSIGN( \
146+
marray &Lhs, const DataT &Rhs) { \
147+
Lhs = Lhs BINOP marray(Rhs); \
148+
return Lhs; \
149+
}
150+
151+
#define __SYCL_BINOP_INTEGRAL(BINOP, OPASSIGN) \
152+
template <typename T = DataT> \
153+
friend typename std::enable_if<std::is_integral<T>::value, marray> \
154+
operator BINOP(const marray &Lhs, const marray &Rhs) { \
155+
marray Ret; \
156+
for (size_t I = 0; I < NumElements; ++I) { \
157+
Ret[I] = Lhs[I] BINOP Rhs[I]; \
158+
} \
159+
return Ret; \
160+
} \
161+
template <typename T, typename BaseT = DataT> \
162+
friend typename std::enable_if<std::is_convertible<T, DataT>::value && \
163+
std::is_integral<T>::value && \
164+
std::is_integral<BaseT>::value, \
165+
marray>::type \
166+
operator BINOP(const marray &Lhs, const T &Rhs) { \
167+
return Lhs BINOP marray(static_cast<DataT>(Rhs)); \
168+
} \
169+
template <typename T = DataT> \
170+
friend typename std::enable_if<std::is_integral<T>::value, marray> \
171+
&operator OPASSIGN(marray &Lhs, const marray &Rhs) { \
172+
Lhs = Lhs BINOP Rhs; \
173+
return Lhs; \
174+
} \
175+
template <std::size_t Num = NumElements, typename T = DataT> \
176+
friend typename std::enable_if<Num != 1 && std::is_integral<T>::value, \
177+
marray &>::type \
178+
operator OPASSIGN(marray &Lhs, const DataT &Rhs) { \
179+
Lhs = Lhs BINOP marray(Rhs); \
180+
return Lhs; \
181+
}
182+
183+
__SYCL_BINOP(+, +=)
184+
__SYCL_BINOP(-, -=)
185+
__SYCL_BINOP(*, *=)
186+
__SYCL_BINOP(/, /=)
187+
188+
__SYCL_BINOP_INTEGRAL(%, %=)
189+
__SYCL_BINOP_INTEGRAL(|, |=)
190+
__SYCL_BINOP_INTEGRAL(&, &=)
191+
__SYCL_BINOP_INTEGRAL(^, ^=)
192+
__SYCL_BINOP_INTEGRAL(>>, >>=)
193+
__SYCL_BINOP_INTEGRAL(<<, <<=)
194+
#undef __SYCL_BINOP
195+
#undef __SYCL_BINOP_INTEGRAL
196+
197+
#ifdef __SYCL_RELLOGOP
198+
#error "Undefine __SYCL_RELLOGOP macro"
199+
#endif
200+
201+
#ifdef __SYCL_RELLOGOP_INTEGRAL
202+
#error "Undefine __SYCL_RELLOGOP_INTEGRAL macro"
203+
#endif
204+
205+
#define __SYCL_RELLOGOP(RELLOGOP) \
206+
friend marray<bool, NumElements> operator RELLOGOP(const marray &Lhs, \
207+
const marray &Rhs) { \
208+
marray<bool, NumElements> Ret; \
209+
for (size_t I = 0; I < NumElements; ++I) { \
210+
Ret[I] = Lhs[I] RELLOGOP Rhs[I]; \
211+
} \
212+
return Ret; \
213+
} \
214+
template <typename T> \
215+
friend typename std::enable_if<std::is_convertible<T, DataT>::value && \
216+
(std::is_fundamental<T>::value || \
217+
std::is_same<T, half>::value), \
218+
marray<bool, NumElements>>::type \
219+
operator RELLOGOP(const marray &Lhs, const T &Rhs) { \
220+
return Lhs RELLOGOP marray(static_cast<const DataT &>(Rhs)); \
221+
}
222+
223+
#define __SYCL_RELLOGOP_INTEGRAL(RELLOGOP) \
224+
template <typename T = DataT> \
225+
friend typename std::enable_if<std::is_integral<T>::value, \
226+
marray<bool, NumElements>>::type \
227+
operator RELLOGOP(const marray &Lhs, const marray &Rhs) { \
228+
marray<bool, NumElements> Ret; \
229+
for (size_t I = 0; I < NumElements; ++I) { \
230+
Ret[I] = Lhs[I] RELLOGOP Rhs[I]; \
231+
} \
232+
return Ret; \
233+
} \
234+
template <typename T, typename BaseT = DataT> \
235+
friend typename std::enable_if<std::is_convertible<T, DataT>::value && \
236+
std::is_integral<T>::value && \
237+
std::is_integral<BaseT>::value, \
238+
marray<bool, NumElements>>::type \
239+
operator RELLOGOP(const marray &Lhs, const T &Rhs) { \
240+
return Lhs RELLOGOP marray(static_cast<const DataT &>(Rhs)); \
241+
}
242+
243+
__SYCL_RELLOGOP(==)
244+
__SYCL_RELLOGOP(!=)
245+
__SYCL_RELLOGOP(>)
246+
__SYCL_RELLOGOP(<)
247+
__SYCL_RELLOGOP(>=)
248+
__SYCL_RELLOGOP(<=)
249+
250+
__SYCL_RELLOGOP_INTEGRAL(&&)
251+
__SYCL_RELLOGOP_INTEGRAL(||)
252+
253+
#undef __SYCL_RELLOGOP
254+
#undef __SYCL_RELLOGOP_INTEGRAL
255+
256+
#ifdef __SYCL_UOP
257+
#error "Undefine __SYCL_UOP macro"
258+
#endif
259+
260+
#define __SYCL_UOP(UOP, OPASSIGN) \
261+
friend marray &operator UOP(marray &Lhs) { \
262+
Lhs OPASSIGN 1; \
263+
return Lhs; \
264+
} \
265+
friend marray operator UOP(marray &Lhs, int) { \
266+
marray Ret(Lhs); \
267+
Lhs OPASSIGN 1; \
268+
return Ret; \
269+
}
270+
271+
__SYCL_UOP(++, +=)
272+
__SYCL_UOP(--, -=)
273+
#undef __SYCL_UOP
274+
275+
// Available only when: dataT != cl_float && dataT != cl_double
276+
// && dataT != cl_half
277+
template <typename T = DataT>
278+
friend typename std::enable_if<std::is_integral<T>::value, marray>::type
279+
operator~(marray &Lhs) {
280+
marray Ret;
281+
for (size_t I = 0; I < NumElements; ++I) {
282+
Ret[I] = ~Lhs[I];
283+
}
284+
return Ret;
285+
}
286+
287+
friend marray<bool, NumElements> operator!(marray &Lhs) {
288+
marray<bool, NumElements> Ret;
289+
for (size_t I = 0; I < NumElements; ++I) {
290+
Ret[I] = !Lhs[I];
291+
}
292+
return Ret;
293+
}
294+
295+
friend marray operator+(marray &Lhs) {
296+
marray Ret;
297+
for (size_t I = 0; I < NumElements; ++I) {
298+
Ret[I] = +Lhs[I];
299+
}
300+
return Ret;
301+
}
302+
303+
friend marray operator-(marray &Lhs) {
304+
marray Ret;
305+
for (size_t I = 0; I < NumElements; ++I) {
306+
Ret[I] = -Lhs[I];
307+
}
308+
return Ret;
309+
}
310+
};
311+
312+
#define __SYCL_MAKE_MARRAY_ALIAS(ALIAS, TYPE, N) \
313+
using ALIAS##N = cl::sycl::marray<TYPE, N>;
314+
315+
#define __SYCL_MAKE_MARRAY_ALIASES_FOR_ARITHMETIC_TYPES(N) \
316+
__SYCL_MAKE_MARRAY_ALIAS(mchar, char, N) \
317+
__SYCL_MAKE_MARRAY_ALIAS(mshort, short, N) \
318+
__SYCL_MAKE_MARRAY_ALIAS(mint, int, N) \
319+
__SYCL_MAKE_MARRAY_ALIAS(mlong, long, N) \
320+
__SYCL_MAKE_MARRAY_ALIAS(mfloat, float, N) \
321+
__SYCL_MAKE_MARRAY_ALIAS(mdouble, double, N) \
322+
__SYCL_MAKE_MARRAY_ALIAS(mhalf, half, N)
323+
324+
#define __SYCL_MAKE_MARRAY_ALIASES_FOR_SIGNED_AND_UNSIGNED_TYPES(N) \
325+
__SYCL_MAKE_MARRAY_ALIAS(mschar, signed char, N) \
326+
__SYCL_MAKE_MARRAY_ALIAS(muchar, unsigned char, N) \
327+
__SYCL_MAKE_MARRAY_ALIAS(mushort, unsigned short, N) \
328+
__SYCL_MAKE_MARRAY_ALIAS(muint, unsigned int, N) \
329+
__SYCL_MAKE_MARRAY_ALIAS(mulong, unsigned long, N) \
330+
__SYCL_MAKE_MARRAY_ALIAS(mlonglong, long long, N) \
331+
__SYCL_MAKE_MARRAY_ALIAS(mulonglong, unsigned long long, N)
332+
333+
#define __SYCL_MAKE_MARRAY_ALIASES_FOR_MARRAY_LENGTH(N) \
334+
__SYCL_MAKE_MARRAY_ALIASES_FOR_ARITHMETIC_TYPES(N) \
335+
__SYCL_MAKE_MARRAY_ALIASES_FOR_SIGNED_AND_UNSIGNED_TYPES(N)
336+
337+
__SYCL_MAKE_MARRAY_ALIASES_FOR_MARRAY_LENGTH(2)
338+
__SYCL_MAKE_MARRAY_ALIASES_FOR_MARRAY_LENGTH(3)
339+
__SYCL_MAKE_MARRAY_ALIASES_FOR_MARRAY_LENGTH(4)
340+
__SYCL_MAKE_MARRAY_ALIASES_FOR_MARRAY_LENGTH(8)
341+
__SYCL_MAKE_MARRAY_ALIASES_FOR_MARRAY_LENGTH(16)
342+
343+
#undef __SYCL_MAKE_MARRAY_ALIAS
344+
#undef __SYCL_MAKE_MARRAY_ALIASES_FOR_ARITHMETIC_TYPES
345+
#undef __SYCL_MAKE_MARRAY_ALIASES_FOR_SIGNED_AND_UNSIGNED_TYPES
346+
#undef __SYCL_MAKE_MARRAY_ALIASES_FOR_MARRAY_LENGTH
347+
348+
} // namespace sycl
349+
} // __SYCL_INLINE_NAMESPACE(cl)

0 commit comments

Comments
 (0)