diff --git a/sycl/include/CL/sycl/builtins.hpp b/sycl/include/CL/sycl/builtins.hpp index 5843dc1562815..c268009216ddc 100644 --- a/sycl/include/CL/sycl/builtins.hpp +++ b/sycl/include/CL/sycl/builtins.hpp @@ -1109,9 +1109,12 @@ fast_normalize(T p) __NOEXC { return __sycl_std::__invoke_fast_normalize(p); } -/* SYCL 1.2.1 ---- 4.13.7 Relational functions. -----------------------------*/ -/* SYCL 2020 ---- 4.17.9 Relational functions. -----------------------------*/ - +/* --------------- 4.13.7 Relational functions. Device version --------------*/ +// int isequal (half x, half y) +// shortn isequal (halfn x, halfn y) +// igeninteger32bit isequal (genfloatf x, genfloatf y) +// int isequal (double x,double y); +// longn isequal (doublen x, doublen y) template ::value, T>> detail::common_rel_ret_t isequal(T x, T y) __NOEXC { @@ -1119,6 +1122,11 @@ detail::common_rel_ret_t isequal(T x, T y) __NOEXC { __sycl_std::__invoke_FOrdEqual>(x, y)); } +// int isnotequal (half x, half y) +// shortn isnotequal (halfn x, halfn y) +// igeninteger32bit isnotequal (genfloatf x, genfloatf y) +// int isnotequal (double x, double y) +// longn isnotequal (doublen x, doublen y) template ::value, T>> detail::common_rel_ret_t isnotequal(T x, T y) __NOEXC { @@ -1126,6 +1134,11 @@ detail::common_rel_ret_t isnotequal(T x, T y) __NOEXC { __sycl_std::__invoke_FUnordNotEqual>(x, y)); } +// int isgreater (half x, half y) +// shortn isgreater (halfn x, halfn y) +// igeninteger32bit isgreater (genfloatf x, genfloatf y) +// int isgreater (double x, double y) +// longn isgreater (doublen x, doublen y) template ::value, T>> detail::common_rel_ret_t isgreater(T x, T y) __NOEXC { @@ -1133,6 +1146,11 @@ detail::common_rel_ret_t isgreater(T x, T y) __NOEXC { __sycl_std::__invoke_FOrdGreaterThan>(x, y)); } +// int isgreaterequal (half x, half y) +// shortn isgreaterequal (halfn x, halfn y) +// igeninteger32bit isgreaterequal (genfloatf x, genfloatf y) +// int isgreaterequal (double x, double y) +// longn isgreaterequal (doublen x, doublen y) template ::value, T>> detail::common_rel_ret_t isgreaterequal(T x, T y) __NOEXC { @@ -1140,6 +1158,11 @@ detail::common_rel_ret_t isgreaterequal(T x, T y) __NOEXC { __sycl_std::__invoke_FOrdGreaterThanEqual>(x, y)); } +// int isless (half x, half y) +// shortn isless (halfn x, halfn y) +// igeninteger32bit isless (genfloatf x, genfloatf y) +// int isless (long x, long y) +// longn isless (doublen x, doublen y) template ::value, T>> detail::common_rel_ret_t isless(T x, T y) __NOEXC { @@ -1147,6 +1170,11 @@ detail::common_rel_ret_t isless(T x, T y) __NOEXC { __sycl_std::__invoke_FOrdLessThan>(x, y)); } +// int islessequal (half x, half y) +// shortn islessequal (halfn x, halfn y) +// igeninteger32bit islessequal (genfloatf x, genfloatf y) +// int islessequal (double x, double y) +// longn islessequal (doublen x, doublen y) template ::value, T>> detail::common_rel_ret_t islessequal(T x, T y) __NOEXC { @@ -1154,6 +1182,11 @@ detail::common_rel_ret_t islessequal(T x, T y) __NOEXC { __sycl_std::__invoke_FOrdLessThanEqual>(x, y)); } +// int islessgreater (half x, half y) +// shortn islessgreater (halfn x, halfn y) +// igeninteger32bit islessgreater (genfloatf x, genfloatf y) +// int islessgreater (double x, double y) +// longn islessgreater (doublen x, doublen y) template ::value, T>> detail::common_rel_ret_t islessgreater(T x, T y) __NOEXC { @@ -1161,6 +1194,11 @@ detail::common_rel_ret_t islessgreater(T x, T y) __NOEXC { __sycl_std::__invoke_FOrdNotEqual>(x, y)); } +// int isfinite (half x) +// shortn isfinite (halfn x) +// igeninteger32bit isfinite (genfloatf x) +// int isfinite (double x) +// longn isfinite (doublen x) template ::value, T>> detail::common_rel_ret_t isfinite(T x) __NOEXC { @@ -1168,6 +1206,11 @@ detail::common_rel_ret_t isfinite(T x) __NOEXC { __sycl_std::__invoke_IsFinite>(x)); } +// int isinf (half x) +// shortn isinf (halfn x) +// igeninteger32bit isinf (genfloatf x) +// int isinf (double x) +// longn isinf (doublen x) template ::value, T>> detail::common_rel_ret_t isinf(T x) __NOEXC { @@ -1175,6 +1218,11 @@ detail::common_rel_ret_t isinf(T x) __NOEXC { __sycl_std::__invoke_IsInf>(x)); } +// int isnan (half x) +// shortn isnan (halfn x) +// igeninteger32bit isnan (genfloatf x) +// int isnan (double x) +// longn isnan (doublen x) template ::value, T>> detail::common_rel_ret_t isnan(T x) __NOEXC { @@ -1182,6 +1230,11 @@ detail::common_rel_ret_t isnan(T x) __NOEXC { __sycl_std::__invoke_IsNan>(x)); } +// int isnormal (half x) +// shortn isnormal (halfn x) +// igeninteger32bit isnormal (genfloatf x) +// int isnormal (double x) +// longn isnormal (doublen x) template ::value, T>> detail::common_rel_ret_t isnormal(T x) __NOEXC { @@ -1189,6 +1242,11 @@ detail::common_rel_ret_t isnormal(T x) __NOEXC { __sycl_std::__invoke_IsNormal>(x)); } +// int isordered (half x) +// shortn isordered (halfn x, halfn y) +// igeninteger32bit isordered (genfloatf x, genfloatf y) +// int isordered (double x, double y) +// longn isordered (doublen x, doublen y) template ::value, T>> detail::common_rel_ret_t isordered(T x, T y) __NOEXC { @@ -1196,6 +1254,11 @@ detail::common_rel_ret_t isordered(T x, T y) __NOEXC { __sycl_std::__invoke_Ordered>(x, y)); } +// int isunordered (half x, half y) +// shortn isunordered (halfn x, halfn y) +// igeninteger32bit isunordered (genfloatf x, genfloatf y) +// int isunordered (double x, double y) +// longn isunordered (doublen x, doublen y) template ::value, T>> detail::common_rel_ret_t isunordered(T x, T y) __NOEXC { @@ -1203,6 +1266,11 @@ detail::common_rel_ret_t isunordered(T x, T y) __NOEXC { __sycl_std::__invoke_Unordered>(x, y)); } +// int signbit (half x) +// shortn signbit (halfn x) +// igeninteger32bit signbit (genfloatf x) +// int signbit (double) +// longn signbit (doublen x) template ::value, T>> detail::common_rel_ret_t signbit(T x) __NOEXC { @@ -1210,37 +1278,29 @@ detail::common_rel_ret_t signbit(T x) __NOEXC { __sycl_std::__invoke_SignBitSet>(x)); } -namespace detail { -#if defined(SYCL2020_CONFORMANT_APIS) && SYCL_LANGUAGE_VERSION >= 202001 -using anyall_ret_t = bool; -#else -using anyall_ret_t = int; -#endif -} // namespace detail - +// int any (sigeninteger x) template -detail::enable_if_t::value, detail::anyall_ret_t> -any(T x) __NOEXC { +detail::enable_if_t::value, int> any(T x) __NOEXC { return detail::Boolean<1>(int(detail::msbIsSet(x))); } +// int any (vigeninteger x) template -detail::enable_if_t::value, detail::anyall_ret_t> -any(T x) __NOEXC { +detail::enable_if_t::value, int> any(T x) __NOEXC { return detail::rel_sign_bit_test_ret_t( __sycl_std::__invoke_Any>( detail::rel_sign_bit_test_arg_t(x))); } +// int all (sigeninteger x) template -detail::enable_if_t::value, detail::anyall_ret_t> -all(T x) __NOEXC { +detail::enable_if_t::value, int> all(T x) __NOEXC { return detail::Boolean<1>(int(detail::msbIsSet(x))); } +// int all (vigeninteger x) template -detail::enable_if_t::value, detail::anyall_ret_t> -all(T x) __NOEXC { +detail::enable_if_t::value, int> all(T x) __NOEXC { return detail::rel_sign_bit_test_ret_t( __sycl_std::__invoke_All>( detail::rel_sign_bit_test_arg_t(x))); diff --git a/sycl/include/CL/sycl/detail/boolean.hpp b/sycl/include/CL/sycl/detail/boolean.hpp index bf75096c353d9..5f2c8776eca6d 100644 --- a/sycl/include/CL/sycl/detail/boolean.hpp +++ b/sycl/include/CL/sycl/detail/boolean.hpp @@ -110,8 +110,7 @@ template <> struct Boolean<1> { // Cast to a signed interger type template operator T() const { - static_assert(std::is_same::value || is_sgeninteger::value, - "Invalid conversion"); + static_assert(is_sgeninteger::value, "Invalid conversion"); return value; } diff --git a/sycl/include/CL/sycl/detail/generic_type_traits.hpp b/sycl/include/CL/sycl/detail/generic_type_traits.hpp index cb67dbfba1991..670fb49104a60 100644 --- a/sycl/include/CL/sycl/detail/generic_type_traits.hpp +++ b/sycl/include/CL/sycl/detail/generic_type_traits.hpp @@ -501,31 +501,9 @@ template inline constexpr bool msbIsSet(const T x) { return (x & msbMask(x)); } -#if defined(SYCL2020_CONFORMANT_APIS) && SYCL_LANGUAGE_VERSION >= 202001 -// SYCL 2020 4.17.9 (Relation functions), e.g. table 178 -// -// genbool isequal (genfloatf x, genfloatf y) -// genbool isequal (genfloatd x, genfloatd y) -// -// TODO: marray support isn't implemented yet. -template -using common_rel_ret_t = - conditional_t::value, make_singed_integer_t, bool>; -#else -// SYCL 1.2.1 4.13.7 (Relation functions), e.g. -// -// igeninteger32bit isequal (genfloatf x, genfloatf y) -// igeninteger64bit isequal (genfloatd x, genfloatd y) -// -// However, we have pre-existing bug so -// -// igeninteger32bit isequal (genfloatd x, genfloatd y) -// -// Fixing it would be an ABI-breaking change so isn't done. template using common_rel_ret_t = conditional_t::value, make_singed_integer_t, int>; -#endif // forward declaration template struct Boolean; diff --git a/sycl/test/basic_tests/relational_builtins.cpp b/sycl/test/basic_tests/relational_builtins.cpp deleted file mode 100644 index e0a74c260999e..0000000000000 --- a/sycl/test/basic_tests/relational_builtins.cpp +++ /dev/null @@ -1,353 +0,0 @@ -// RUN: %clangxx -DSYCL2020_CONFORMANT_APIS -fsycl %s -// RUN: %clangxx -sycl-std=121 -fsycl %s - -#include - -// Some helper macros to verify return type of the builtins. To be used like -// this -// -// CHECK(Expected return type in SYCL 1.2.1, -// Expected return type in SYCL 2020, -// builtin name, -// parameters' types...) -// -// C++17 doesn't allow lambdas in unevaluated context. Could be simplified -// further in C++20 including more std::declval usage. -template struct CheckHelper { - template static auto call(F f) { return f(Args()...); } -}; - -#if defined(SYCL2020_CONFORMANT_APIS) && SYCL_LANGUAGE_VERSION >= 202001 -#define CHECK(EXPECTED121, EXPECTED2020, FUNC, ...) \ - { \ - auto ret = CheckHelper<__VA_ARGS__>::call( \ - [](auto... args) { return cl::sycl::FUNC(args...); }); \ - static_assert(std::is_same_v); \ - } -// To be used for marray tests. Not yet implemented -// #define CHECK2020(...) CHECK(__VA_ARGS__) -#define CHECK2020(...) -#else -#define CHECK(EXPECTED121, EXPECTED2020, FUNC, ...) \ - { \ - auto ret = CheckHelper<__VA_ARGS__>::call( \ - [](auto... args) { return cl::sycl::FUNC(args...); }); \ - static_assert(std::is_same_v); \ - } -#define CHECK2020(...) -#endif - -void foo() { - using namespace cl::sycl; - using boolm = marray; - - using int16v = vec; - using int16m = marray; - - using uint16v = vec; - using uint16m = marray; - - using halfv = vec; - using halfm = marray; - - using int32v = vec; - using int32m = marray; - - using uint32v = vec; - using uint32m = marray; - - using floatv = vec; - using floatm = marray; - - using int64v = vec; - using int64m = marray; - - using uint64v = vec; - using uint64m = marray; - - using doublev = vec; - using doublem = marray; - - // isequal - CHECK(int32_t, bool, isequal, half, half); - CHECK(int16v, int16v, isequal, halfv, halfv); - CHECK2020(_, boolm, isequal, halfm, halfm); - - CHECK(int32_t, bool, isequal, float, float); - CHECK(int32v, int32v, isequal, floatv, floatv); - CHECK2020(_, boolm, isequal, floatm, floatm); - - // SYCL 1.2.1 has an ABI-affecting bug here (int32_t instead of int64_t for - // scalar case). - CHECK(int32_t, bool, isequal, double, double); - CHECK(int64v, int64v, isequal, doublev, doublev); - CHECK2020(_, boolm, isequal, doublem, doublem); - - // isnotequal - CHECK(int32_t, bool, isnotequal, half, half); - CHECK(int16v, int16v, isnotequal, halfv, halfv); - CHECK2020(_, boolm, isnotequal, halfm, halfm); - - CHECK(int32_t, bool, isnotequal, float, float); - CHECK(int32v, int32v, isnotequal, floatv, floatv); - CHECK2020(_, boolm, isnotequal, floatm, floatm); - - // SYCL 1.2.1 has an ABI-affecting bug here (int32_t instead of int64_t for - // scalar case). - CHECK(int32_t, bool, isnotequal, double, double); - CHECK(int64v, int64v, isnotequal, doublev, doublev); - CHECK2020(_, boolm, isnotequal, doublem, doublem); - - // isgreater - CHECK(int32_t, bool, isgreater, half, half); - CHECK(int16v, int16v, isgreater, halfv, halfv); - CHECK2020(_, boolm, isgreater, halfm, halfm); - - CHECK(int32_t, bool, isgreater, float, float); - CHECK(int32v, int32v, isgreater, floatv, floatv); - CHECK2020(_, boolm, isgreater, floatm, floatm); - - // SYCL 1.2.1 has an ABI-affecting bug here (int32_t instead of int64_t for - // scalar case). - CHECK(int32_t, bool, isgreater, double, double); - CHECK(int64v, int64v, isgreater, doublev, doublev); - CHECK2020(_, boolm, isgreater, doublem, doublem); - - // isgreaterequal - CHECK(int32_t, bool, isgreaterequal, half, half); - CHECK(int16v, int16v, isgreaterequal, halfv, halfv); - CHECK2020(_, boolm, isgreaterequal, halfm, halfm); - - CHECK(int32_t, bool, isgreaterequal, float, float); - CHECK(int32v, int32v, isgreaterequal, floatv, floatv); - CHECK2020(_, boolm, isgreaterequal, floatm, floatm); - - // SYCL 1.2.1 has an ABI-affecting bug here (int32_t instead of int64_t for - // scalar case). - CHECK(int32_t, bool, isgreaterequal, double, double); - CHECK(int64v, int64v, isgreaterequal, doublev, doublev); - CHECK2020(_, boolm, isgreaterequal, doublem, doublem); - - // isless - CHECK(int32_t, bool, isless, half, half); - CHECK(int16v, int16v, isless, halfv, halfv); - CHECK2020(_, boolm, isless, halfm, halfm); - - CHECK(int32_t, bool, isless, float, float); - CHECK(int32v, int32v, isless, floatv, floatv); - CHECK2020(_, boolm, isless, floatm, floatm); - - // SYCL 1.2.1 has an ABI-affecting bug here (int32_t instead of int64_t for - // scalar case). - CHECK(int32_t, bool, isless, double, double); - CHECK(int64v, int64v, isless, doublev, doublev); - CHECK2020(_, boolm, isless, doublem, doublem); - - // islessequal - CHECK(int32_t, bool, islessequal, half, half); - CHECK(int16v, int16v, islessequal, halfv, halfv); - CHECK2020(_, boolm, islessequal, halfm, halfm); - - CHECK(int32_t, bool, islessequal, float, float); - CHECK(int32v, int32v, islessequal, floatv, floatv); - CHECK2020(_, boolm, islessequal, floatm, floatm); - - // SYCL 1.2.1 has an ABI-affecting bug here (int32_t instead of int64_t for - // scalar case). - CHECK(int32_t, bool, islessequal, double, double); - CHECK(int64v, int64v, islessequal, doublev, doublev); - CHECK2020(_, boolm, islessequal, doublem, doublem); - - // islessgreater - CHECK(int32_t, bool, islessgreater, half, half); - CHECK(int16v, int16v, islessgreater, halfv, halfv); - CHECK2020(_, boolm, islessgreater, halfm, halfm); - - CHECK(int32_t, bool, islessgreater, float, float); - CHECK(int32v, int32v, islessgreater, floatv, floatv); - CHECK2020(_, boolm, islessgreater, floatm, floatm); - - // SYCL 1.2.1 has an ABI-affecting bug here (int32_t instead of int64_t for - // scalar case). - CHECK(int32_t, bool, islessgreater, double, double); - CHECK(int64v, int64v, islessgreater, doublev, doublev); - CHECK2020(_, boolm, islessgreater, doublem, doublem); - - // isfinite - CHECK(int32_t, bool, isfinite, half); - CHECK(int16v, int16v, isfinite, halfv); - CHECK2020(_, boolm, isfinite, halfm); - - CHECK(int32_t, bool, isfinite, float); - CHECK(int32v, int32v, isfinite, floatv); - CHECK2020(_, boolm, isfinite, floatm); - - // SYCL 1.2.1 has an ABI-affecting bug here (int32_t instead of int64_t for - // scalar case). - CHECK(int32_t, bool, isfinite, double); - CHECK(int64v, int64v, isfinite, doublev); - CHECK2020(_, boolm, isfinite, doublem); - - // isinf - CHECK(int32_t, bool, isinf, half); - CHECK(int16v, int16v, isinf, halfv); - CHECK2020(_, boolm, isinf, halfm); - - CHECK(int32_t, bool, isinf, float); - CHECK(int32v, int32v, isinf, floatv); - CHECK2020(_, boolm, isinf, floatm); - - // SYCL 1.2.1 has an ABI-affecting bug here (int32_t instead of int64_t for - // scalar case). - CHECK(int32_t, bool, isinf, double); - CHECK(int64v, int64v, isinf, doublev); - CHECK2020(_, boolm, isinf, doublem); - - // isnan - CHECK(int32_t, bool, isnan, half); - CHECK(int16v, int16v, isnan, halfv); - CHECK2020(_, boolm, isnan, halfm); - - CHECK(int32_t, bool, isnan, float); - CHECK(int32v, int32v, isnan, floatv); - CHECK2020(_, boolm, isnan, floatm); - - // SYCL 1.2.1 has an ABI-affecting bug here (int32_t instead of int64_t for - // scalar case). - CHECK(int32_t, bool, isnan, double); - CHECK(int64v, int64v, isnan, doublev); - CHECK2020(_, boolm, isnan, doublem); - - // isnormal - CHECK(int32_t, bool, isnormal, half); - CHECK(int16v, int16v, isnormal, halfv); - CHECK2020(_, boolm, isnormal, halfm); - - CHECK(int32_t, bool, isnormal, float); - CHECK(int32v, int32v, isnormal, floatv); - CHECK2020(_, boolm, isnormal, floatm); - - // SYCL 1.2.1 has an ABI-affecting bug here (int32_t instead of int64_t for - // scalar case). - CHECK(int32_t, bool, isnormal, double); - CHECK(int64v, int64v, isnormal, doublev); - CHECK2020(_, boolm, isnormal, doublem); - - // isordered - CHECK(int32_t, bool, isordered, half, half); - CHECK(int16v, int16v, isordered, halfv, halfv); - CHECK2020(_, boolm, isordered, halfm, halfm); - - CHECK(int32_t, bool, isordered, float, float); - CHECK(int32v, int32v, isordered, floatv, floatv); - CHECK2020(_, boolm, isordered, floatm, floatm); - - // SYCL 1.2.1 has an ABI-affecting bug here (int32_t instead of int64_t for - // scalar case). - CHECK(int32_t, bool, isordered, double, double); - CHECK(int64v, int64v, isordered, doublev, doublev); - CHECK2020(_, boolm, isordered, doublem, doublem); - - // isunordered - CHECK(int32_t, bool, isunordered, half, half); - CHECK(int16v, int16v, isunordered, halfv, halfv); - CHECK2020(_, boolm, isunordered, halfm, halfm); - - CHECK(int32_t, bool, isunordered, float, float); - CHECK(int32v, int32v, isunordered, floatv, floatv); - CHECK2020(_, boolm, isunordered, floatm, floatm); - - // SYCL 1.2.1 has an ABI-affecting bug here (int32_t instead of int64_t for - // scalar case). - CHECK(int32_t, bool, isunordered, double, double); - CHECK(int64v, int64v, isunordered, doublev, doublev); - CHECK2020(_, boolm, isunordered, doublem, doublem); - - // signbit - CHECK(int32_t, bool, signbit, half); - CHECK(int16v, int16v, signbit, halfv); - CHECK2020(_, boolm, signbit, halfm); - - CHECK(int32_t, bool, signbit, float); - CHECK(int32v, int32v, signbit, floatv); - CHECK2020(_, boolm, signbit, floatm); - - // SYCL 1.2.1 has an ABI-affecting bug here (int32_t instead of int64_t for - // scalar case). - CHECK(int32_t, bool, signbit, double); - CHECK(int64v, int64v, signbit, doublev); - CHECK2020(_, boolm, signbit, doublem); - - // any - CHECK(int, bool, any, int16_t) - CHECK(int, bool, any, int16v) - CHECK2020(_, bool, any, int16m) - - CHECK(int, bool, any, int32_t) - CHECK(int, bool, any, int32v) - CHECK2020(_, bool, any, int32m) - - CHECK(int, bool, any, int64_t) - CHECK(int, bool, any, int64v) - CHECK2020(_, bool, any, int64m) - - // all - CHECK(int, bool, all, int16_t) - CHECK(int, bool, all, int16v) - CHECK2020(_, bool, all, int16m) - - CHECK(int, bool, all, int32_t) - CHECK(int, bool, all, int32v) - CHECK2020(_, bool, all, int32m) - - CHECK(int, bool, all, int64_t) - CHECK(int, bool, all, int64v) - CHECK2020(_, bool, all, int64m) - - // bitselect - CHECK(int16_t, int16_t, bitselect, int16_t, int16_t, int16_t) - CHECK(int16v, int16v, bitselect, int16v, int16v, int16v) - CHECK2020(int16m, int16m, bitselect, int16m, int16m, int16m) - - CHECK(uint16_t, uint16_t, bitselect, uint16_t, uint16_t, uint16_t) - CHECK(uint16v, uint16v, bitselect, uint16v, uint16v, uint16v) - CHECK2020(uint16m, uint16m, bitselect, uint16m, uint16m, uint16m) - - CHECK(half, half, bitselect, half, half, half) - CHECK(halfv, halfv, bitselect, halfv, halfv, halfv) - - CHECK(int32_t, int32_t, bitselect, int32_t, int32_t, int32_t) - CHECK(int32v, int32v, bitselect, int32v, int32v, int32v) - CHECK2020(int32m, int32m, bitselect, int32m, int32m, int32m) - - CHECK(uint32_t, uint32_t, bitselect, uint32_t, uint32_t, uint32_t) - CHECK(uint32v, uint32v, bitselect, uint32v, uint32v, uint32v) - CHECK2020(uint32m, uint32m, bitselect, uint32m, uint32m, uint32m) - - CHECK(float, float, bitselect, float, float, float) - CHECK(floatv, floatv, bitselect, floatv, floatv, floatv) - CHECK2020(floatm, floatm, bitselect, floatm, floatm, floatm) - CHECK2020(floatm, floatm, bitselect, floatm, floatm, floatm) - - CHECK(int64_t, int64_t, bitselect, int64_t, int64_t, int64_t) - CHECK(int64v, int64v, bitselect, int64v, int64v, int64v) - CHECK2020(int64m, int64m, bitselect, int64m, int64m, int64m) - - CHECK(uint64_t, uint64_t, bitselect, uint64_t, uint64_t, uint64_t) - CHECK(uint64v, uint64v, bitselect, uint64v, uint64v, uint64v) - CHECK2020(uint64m, uint64m, bitselect, uint64m, uint64m, uint64m) - - CHECK(double, double, bitselect, double, double, double) - CHECK(doublev, doublev, bitselect, doublev, doublev, doublev) - CHECK2020(doublem, doublem, bitselect, doublem, doublem, doublem) -} - -int main() { - cl::sycl::queue q; - foo(); // Verify host. - q.submit([&](cl::sycl::handler &cgh) { - cgh.single_task([]() { - foo(); // verify device - }); - }); -}