Skip to content

Commit 8b0b210

Browse files
dm-vodopyanovdm-vodopyanov
authored
[SYCL] Add missing relational functions + bunch of other fixes (#9110)
This patch makes math_builtin_api SYCL CTS test compilable and fixes most of the run failures. 1. Add missing marray relational functions: `any`, `all`, `bitselect` 2. Fix scalar `select` relational function: it didn't work for most of the types 3. Align scalar `abs` integer function with SYCL 2020: now it returns signed integer type instead of unsigned for signed argument 4. Fix math functions which take multi_ptr argument: now they accept multi_ptr without decoration in accordance with SYCL 2020
1 parent 50f1605 commit 8b0b210

File tree

7 files changed

+167
-39
lines changed

7 files changed

+167
-39
lines changed

‎sycl/include/sycl/builtins.hpp‎

Lines changed: 66 additions & 21 deletions
Original file line numberDiff line numberDiff line change
@@ -15,6 +15,8 @@
1515
#include <sycl/pointers.hpp>
1616
#include <sycl/types.hpp>
1717

18+
#include <algorithm>
19+
1820
// TODO Decide whether to mark functions with this attribute.
1921
#define __NOEXC /*noexcept*/
2022

@@ -1082,12 +1084,14 @@ detail::enable_if_t<detail::is_ugeninteger<T>::value, T> abs(T x) __NOEXC {
10821084
return __sycl_std::__invoke_u_abs<T>(x);
10831085
}
10841086

1085-
// ugeninteger abs (geninteger x)
1087+
// igeninteger abs (geninteger x)
10861088
template <typename T>
1087-
detail::enable_if_t<detail::is_igeninteger<T>::value,
1088-
detail::make_unsigned_t<T>>
1089-
abs(T x) __NOEXC {
1090-
return __sycl_std::__invoke_s_abs<detail::make_unsigned_t<T>>(x);
1089+
detail::enable_if_t<detail::is_igeninteger<T>::value, T> abs(T x) __NOEXC {
1090+
auto res = __sycl_std::__invoke_s_abs<detail::make_unsigned_t<T>>(x);
1091+
if constexpr (detail::is_vigeninteger<T>::value) {
1092+
return res.template convert<detail::vector_element_t<T>>();
1093+
} else
1094+
return detail::make_signed_t<decltype(res)>(res);
10911095
}
10921096

10931097
// ugeninteger abs_diff (geninteger x, geninteger y)
@@ -1434,9 +1438,8 @@ mul24(T x, T y) __NOEXC {
14341438

14351439
#define __SYCL_MARRAY_INTEGER_FUNCTION_ABS_I_OVERLOAD(NAME, ARG, ...) \
14361440
template <typename T, size_t N> \
1437-
std::enable_if_t<detail::is_igeninteger<T>::value, \
1438-
marray<detail::make_unsigned_t<T>, N>> \
1439-
NAME(marray<T, N> ARG) __NOEXC { \
1441+
std::enable_if_t<detail::is_igeninteger<T>::value, marray<T, N>> NAME( \
1442+
marray<T, N> ARG) __NOEXC { \
14401443
__SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
14411444
}
14421445

@@ -2073,19 +2076,27 @@ detail::enable_if_t<detail::is_gentype<T>::value, T> bitselect(T a, T b,
20732076
template <typename T>
20742077
detail::enable_if_t<detail::is_sgentype<T>::value, T> select(T a, T b,
20752078
bool c) __NOEXC {
2076-
return __sycl_std::__invoke_select<T>(a, b, static_cast<int>(c));
2077-
}
2078-
2079-
// mgentype select (mgentype a, mgentype b, marray<bool, { N }> c)
2080-
template <typename T,
2081-
typename = std::enable_if_t<detail::is_mgenfloat<T>::value>>
2082-
sycl::marray<detail::marray_element_t<T>, T::size()>
2083-
select(T a, T b, sycl::marray<bool, T::size()> c) __NOEXC {
2084-
sycl::marray<detail::marray_element_t<T>, T::size()> res;
2085-
for (int i = 0; i < a.size(); i++) {
2086-
res[i] = select(a[i], b[i], c[i]);
2087-
}
2088-
return res;
2079+
constexpr size_t SizeT = sizeof(T);
2080+
2081+
// sycl::select(sgentype a, sgentype b, bool c) calls OpenCL built-in
2082+
// select(sgentype a, sgentype b, igentype c). This type trait makes the
2083+
// proper conversion for argument c from bool to igentype, based on sgentype
2084+
// == T.
2085+
using get_select_opencl_builtin_c_arg_type = typename std::conditional_t<
2086+
SizeT == 1, char,
2087+
std::conditional_t<
2088+
SizeT == 2, short,
2089+
std::conditional_t<
2090+
(detail::is_contained<
2091+
T, detail::type_list<long, unsigned long>>::value &&
2092+
(SizeT == 4 || SizeT == 8)),
2093+
long, // long and ulong are 32-bit on
2094+
// Windows and 64-bit on Linux
2095+
std::conditional_t<SizeT == 4, int,
2096+
std::conditional_t<SizeT == 8, long, void>>>>>;
2097+
2098+
return __sycl_std::__invoke_select<T>(
2099+
a, b, static_cast<get_select_opencl_builtin_c_arg_type>(c));
20892100
}
20902101

20912102
// geninteger select (geninteger a, geninteger b, igeninteger c)
@@ -2164,6 +2175,40 @@ select(T a, T b, T2 c) __NOEXC {
21642175
return __sycl_std::__invoke_select<T>(a, b, c);
21652176
}
21662177

2178+
// other marray relational functions
2179+
2180+
template <typename T, size_t N>
2181+
detail::enable_if_t<detail::is_sigeninteger<T>::value, bool>
2182+
any(marray<T, N> x) __NOEXC {
2183+
return std::any_of(x.begin(), x.end(), [](T i) { return any(i); });
2184+
}
2185+
2186+
template <typename T, size_t N>
2187+
detail::enable_if_t<detail::is_sigeninteger<T>::value, bool>
2188+
all(marray<T, N> x) __NOEXC {
2189+
return std::all_of(x.begin(), x.end(), [](T i) { return all(i); });
2190+
}
2191+
2192+
template <typename T, size_t N>
2193+
detail::enable_if_t<detail::is_gentype<T>::value, marray<T, N>>
2194+
bitselect(marray<T, N> a, marray<T, N> b, marray<T, N> c) __NOEXC {
2195+
marray<T, N> res;
2196+
for (int i = 0; i < N; i++) {
2197+
res[i] = bitselect(a[i], b[i], c[i]);
2198+
}
2199+
return res;
2200+
}
2201+
2202+
template <typename T, size_t N>
2203+
detail::enable_if_t<detail::is_gentype<T>::value, marray<T, N>>
2204+
select(marray<T, N> a, marray<T, N> b, marray<bool, N> c) __NOEXC {
2205+
marray<T, N> res;
2206+
for (int i = 0; i < N; i++) {
2207+
res[i] = select(a[i], b[i], c[i]);
2208+
}
2209+
return res;
2210+
}
2211+
21672212
namespace native {
21682213
/* ----------------- 4.13.3 Math functions. ---------------------------------*/
21692214

‎sycl/include/sycl/detail/generic_type_traits.hpp‎

Lines changed: 9 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -386,9 +386,15 @@ template <typename To> struct PointerConverter {
386386
}
387387

388388
template <typename From> static To Convert(From &t) {
389-
// TODO find the better way to get the pointer to underlying data from vec
390-
// class
391-
return reinterpret_cast<To>(t.get());
389+
if constexpr (is_non_legacy_multi_ptr_v<From>) {
390+
return detail::cast_AS<To>(t.get_decorated());
391+
} else if constexpr (is_legacy_multi_ptr_v<From>) {
392+
return detail::cast_AS<To>(t.get());
393+
} else {
394+
// TODO find the better way to get the pointer to underlying data from vec
395+
// class
396+
return reinterpret_cast<To>(t.get());
397+
}
392398
}
393399
};
394400

‎sycl/include/sycl/detail/type_traits.hpp‎

Lines changed: 37 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -292,6 +292,43 @@ struct is_pointer_impl<multi_ptr<T, Space, DecorateAddress>> : std::true_type {
292292

293293
template <typename T> struct is_pointer : is_pointer_impl<remove_cv_t<T>> {};
294294

295+
// is_multi_ptr
296+
template <typename T> struct is_multi_ptr : std::false_type {};
297+
298+
template <typename ElementType, access::address_space Space,
299+
access::decorated IsDecorated>
300+
struct is_multi_ptr<multi_ptr<ElementType, Space, IsDecorated>>
301+
: std::true_type {};
302+
303+
template <class T>
304+
inline constexpr bool is_multi_ptr_v = is_multi_ptr<T>::value;
305+
306+
// is_non_legacy_multi_ptr
307+
template <typename T> struct is_non_legacy_multi_ptr : std::false_type {};
308+
309+
template <typename ElementType, access::address_space Space>
310+
struct is_non_legacy_multi_ptr<
311+
multi_ptr<ElementType, Space, access::decorated::yes>> : std::true_type {};
312+
313+
template <typename ElementType, access::address_space Space>
314+
struct is_non_legacy_multi_ptr<
315+
multi_ptr<ElementType, Space, access::decorated::no>> : std::true_type {};
316+
317+
template <class T>
318+
inline constexpr bool is_non_legacy_multi_ptr_v =
319+
is_non_legacy_multi_ptr<T>::value;
320+
321+
// is_legacy_multi_ptr
322+
template <typename T> struct is_legacy_multi_ptr : std::false_type {};
323+
324+
template <typename ElementType, access::address_space Space>
325+
struct is_legacy_multi_ptr<
326+
multi_ptr<ElementType, Space, access::decorated::legacy>> : std::true_type {
327+
};
328+
329+
template <class T>
330+
inline constexpr bool is_legacy_multi_ptr_v = is_legacy_multi_ptr<T>::value;
331+
295332
// remove_pointer_t
296333
template <typename T> struct remove_pointer_impl {
297334
using type = T;

‎sycl/test-e2e/Basic/built-ins/marray_math.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -44,7 +44,7 @@
4444
cgh.single_task([=]() { \
4545
sycl::multi_ptr<sycl::marray<PTR_TYPE, DIM>, \
4646
sycl::access::address_space::global_space, \
47-
sycl::access::decorated::yes> \
47+
sycl::access::decorated::no> \
4848
ptr(res_ptr_access); \
4949
sycl::marray<MARRAY_ELEM_TYPE, DIM> res = FUNC(__VA_ARGS__, ptr); \
5050
for (int i = 0; i < DIM; i++) \

‎sycl/test-e2e/DeviceLib/built-ins/marray_relational.cpp‎ renamed to ‎sycl/test-e2e/Basic/built-ins/marray_relational.cpp‎

Lines changed: 26 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -25,6 +25,25 @@
2525
} \
2626
}
2727

28+
#define TEST2(FUNC, TYPE, EXPECTED, N, ...) \
29+
{ \
30+
{ \
31+
TYPE result[1]; \
32+
{ \
33+
sycl::buffer<TYPE> b(result, sycl::range{1}); \
34+
deviceQueue.submit([&](sycl::handler &cgh) { \
35+
sycl::accessor res_access{b, cgh}; \
36+
cgh.single_task([=]() { \
37+
int res = FUNC(__VA_ARGS__); \
38+
for (int i = 0; i < N; i++) \
39+
res_access[0] = res; \
40+
}); \
41+
}); \
42+
} \
43+
assert(result[0] == EXPECTED[0]); \
44+
} \
45+
}
46+
2847
#define EXPECTED(TYPE, ...) ((TYPE[]){__VA_ARGS__})
2948

3049
int main() {
@@ -36,6 +55,10 @@ int main() {
3655
sycl::marray<float, 2> ma4{2.0, 2.0};
3756
sycl::marray<float, 3> ma5{2.0, 2.0, 1.0};
3857
sycl::marray<float, 3> ma6{1.0, 5.0, 8.0};
58+
sycl::marray<int, 3> ma7{50, 2, 31};
59+
sycl::marray<float, 2> ma8{1.0, 1.0};
60+
sycl::marray<float, 2> ma9{0.5, 0.5};
61+
sycl::marray<float, 2> ma10{2.0, 2.0};
3962
sycl::marray<bool, 3> c(1, 0, 1);
4063

4164
TEST(sycl::isequal, bool, EXPECTED(bool, 1, 1), 2, ma1, ma2);
@@ -52,6 +75,9 @@ int main() {
5275
TEST(sycl::isordered, bool, EXPECTED(bool, 1, 1), 2, ma1, ma2);
5376
TEST(sycl::isunordered, bool, EXPECTED(bool, 0, 0), 2, ma1, ma2);
5477
TEST(sycl::signbit, bool, EXPECTED(bool, 0, 0), 2, ma1);
78+
TEST2(sycl::all, int, EXPECTED(bool, false), 3, ma7);
79+
TEST2(sycl::any, int, EXPECTED(bool, false), 3, ma7);
80+
TEST(sycl::bitselect, float, EXPECTED(float, 1.0, 1.0), 2, ma8, ma9, ma10);
5581
TEST(sycl::select, float, EXPECTED(float, 1.0, 2.0, 8.0), 3, ma5, ma6, c);
5682

5783
return 0;

‎sycl/test-e2e/DeviceLib/built-ins/vector_integer.cpp‎

Lines changed: 4 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -224,11 +224,10 @@ int main() {
224224

225225
// abs (longlong)
226226
{
227-
using ulonglong2 = s::vec<unsigned long long, 2>;
228227
using longlong2 = s::vec<long long, 2>;
229-
ulonglong2 r{0};
228+
longlong2 r{0};
230229
{
231-
s::buffer<ulonglong2, 1> BufR(&r, s::range<1>(1));
230+
s::buffer<longlong2, 1> BufR(&r, s::range<1>(1));
232231
s::queue myQueue;
233232
myQueue.submit([&](s::handler &cgh) {
234233
auto AccR = BufR.get_access<s::access::mode::write>(cgh);
@@ -237,8 +236,8 @@ int main() {
237236
});
238237
});
239238
}
240-
unsigned long long r1 = r.x();
241-
unsigned long long r2 = r.y();
239+
long long r1 = r.x();
240+
long long r2 = r.y();
242241
assert(r1 == 5);
243242
assert(r2 == 2);
244243
}

‎sycl/test/regression/host_builtins_gcc.cpp‎

Lines changed: 24 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -51,19 +51,31 @@
5151
float f; \
5252
double d; \
5353
int *ip = nullptr; \
54-
TEST_CASE_2(OP, h, sycl::private_ptr<int>{ip}, sycl::half) \
55-
TEST_CASE_2(OP, f, sycl::private_ptr<int>{ip}, float) \
56-
TEST_CASE_2(OP, d, sycl::private_ptr<int>{ip}, double) \
54+
sycl::decorated_private_ptr<int> ptr = \
55+
sycl::address_space_cast<sycl::access::address_space::private_space, \
56+
sycl::access::decorated::yes>(ip); \
57+
TEST_CASE_2(OP, h, ptr, sycl::half) \
58+
TEST_CASE_2(OP, f, ptr, float) \
59+
TEST_CASE_2(OP, d, ptr, double) \
5760
}
5861

5962
#define TEST_CASES_F1FP1(OP) \
6063
{ \
6164
sycl::half h; \
6265
float f; \
6366
double d; \
64-
TEST_CASE_2(OP, h, sycl::private_ptr<sycl::half>{&h}, sycl::half) \
65-
TEST_CASE_2(OP, f, sycl::private_ptr<float>{&f}, float) \
66-
TEST_CASE_2(OP, d, sycl::private_ptr<double>{&d}, double) \
67+
sycl::decorated_private_ptr<sycl::half> ptr_h = \
68+
sycl::address_space_cast<sycl::access::address_space::private_space, \
69+
sycl::access::decorated::yes>(&h); \
70+
sycl::decorated_private_ptr<float> ptr_f = \
71+
sycl::address_space_cast<sycl::access::address_space::private_space, \
72+
sycl::access::decorated::yes>(&f); \
73+
sycl::decorated_private_ptr<double> ptr_d = \
74+
sycl::address_space_cast<sycl::access::address_space::private_space, \
75+
sycl::access::decorated::yes>(&d); \
76+
TEST_CASE_2(OP, h, ptr_h, sycl::half) \
77+
TEST_CASE_2(OP, f, ptr_f, float) \
78+
TEST_CASE_2(OP, d, ptr_d, double) \
6779
}
6880

6981
#define TEST_CASES_F2IP1(OP) \
@@ -72,9 +84,12 @@
7284
float f; \
7385
double d; \
7486
int *ip = nullptr; \
75-
TEST_CASE_3(OP, h, h, sycl::private_ptr<int>{ip}, sycl::half) \
76-
TEST_CASE_3(OP, f, f, sycl::private_ptr<int>{ip}, float) \
77-
TEST_CASE_3(OP, d, d, sycl::private_ptr<int>{ip}, double) \
87+
sycl::decorated_private_ptr<int> ptr = \
88+
sycl::address_space_cast<sycl::access::address_space::private_space, \
89+
sycl::access::decorated::yes>(ip); \
90+
TEST_CASE_3(OP, h, h, ptr, sycl::half) \
91+
TEST_CASE_3(OP, f, f, ptr, float) \
92+
TEST_CASE_3(OP, d, d, ptr, double) \
7893
}
7994

8095
int main() {

0 commit comments

Comments
 (0)