From 45d516ccd65de0a0c3a1df68321975ee0be26acf Mon Sep 17 00:00:00 2001 From: Steffen Larsen Date: Fri, 18 Nov 2022 09:04:31 +0100 Subject: [PATCH] [SYCL] Fix return type of relational functions on scalars (#7414) This fix updates scalar versions of relation functions that now return "bool" instead of a signed integer in SYCL 2020 mode. The fix is ABI-breaking and is put under SYCL2020_CONFORMANT_APIS guard. This is a revival of the reverted patch https://github.com/intel/llvm/pull/5975. Signed-off-by: Larsen, Steffen Co-authored-by: aelovikov-intel --- sycl/include/sycl/builtins.hpp | 125 ++----- sycl/include/sycl/detail/boolean.hpp | 3 +- .../sycl/detail/generic_type_traits.hpp | 46 ++- sycl/test/basic_tests/relational_builtins.cpp | 353 ++++++++++++++++++ 4 files changed, 433 insertions(+), 94 deletions(-) create mode 100644 sycl/test/basic_tests/relational_builtins.cpp diff --git a/sycl/include/sycl/builtins.hpp b/sycl/include/sycl/builtins.hpp index 4dbbec61409c8..923deb1cd71c0 100644 --- a/sycl/include/sycl/builtins.hpp +++ b/sycl/include/sycl/builtins.hpp @@ -1109,184 +1109,129 @@ fast_normalize(T p) __NOEXC { return __sycl_std::__invoke_fast_normalize(p); } -/* --------------- 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) +/* SYCL 1.2.1 ---- 4.13.7 Relational functions. -----------------------------*/ +/* SYCL 2020 ---- 4.17.9 Relational functions. -----------------------------*/ + template ::value, T>> detail::common_rel_ret_t isequal(T x, T y) __NOEXC { return detail::RelConverter::apply( - __sycl_std::__invoke_FOrdEqual>(x, y)); + __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 { return detail::RelConverter::apply( - __sycl_std::__invoke_FUnordNotEqual>(x, y)); + __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 { return detail::RelConverter::apply( - __sycl_std::__invoke_FOrdGreaterThan>(x, y)); + __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 { return detail::RelConverter::apply( - __sycl_std::__invoke_FOrdGreaterThanEqual>(x, y)); + __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 { return detail::RelConverter::apply( - __sycl_std::__invoke_FOrdLessThan>(x, y)); + __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 { return detail::RelConverter::apply( - __sycl_std::__invoke_FOrdLessThanEqual>(x, y)); + __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 { return detail::RelConverter::apply( - __sycl_std::__invoke_FOrdNotEqual>(x, y)); + __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 { return detail::RelConverter::apply( - __sycl_std::__invoke_IsFinite>(x)); + __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 { return detail::RelConverter::apply( - __sycl_std::__invoke_IsInf>(x)); + __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 { return detail::RelConverter::apply( - __sycl_std::__invoke_IsNan>(x)); + __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 { return detail::RelConverter::apply( - __sycl_std::__invoke_IsNormal>(x)); + __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 { return detail::RelConverter::apply( - __sycl_std::__invoke_Ordered>(x, y)); + __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 { return detail::RelConverter::apply( - __sycl_std::__invoke_Unordered>(x, y)); + __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 { return detail::RelConverter::apply( - __sycl_std::__invoke_SignBitSet>(x)); + __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, int> any(T x) __NOEXC { +detail::enable_if_t::value, detail::anyall_ret_t> +any(T x) __NOEXC { return detail::Boolean<1>(int(detail::msbIsSet(x))); } // int any (vigeninteger x) template -detail::enable_if_t::value, int> any(T x) __NOEXC { +detail::enable_if_t::value, detail::anyall_ret_t> +any(T x) __NOEXC { return detail::rel_sign_bit_test_ret_t( __sycl_std::__invoke_Any>( detail::rel_sign_bit_test_arg_t(x))); @@ -1294,13 +1239,15 @@ detail::enable_if_t::value, int> any(T x) __NOEXC { // int all (sigeninteger x) template -detail::enable_if_t::value, int> all(T x) __NOEXC { +detail::enable_if_t::value, detail::anyall_ret_t> +all(T x) __NOEXC { return detail::Boolean<1>(int(detail::msbIsSet(x))); } // int all (vigeninteger x) template -detail::enable_if_t::value, int> all(T x) __NOEXC { +detail::enable_if_t::value, detail::anyall_ret_t> +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/sycl/detail/boolean.hpp b/sycl/include/sycl/detail/boolean.hpp index 60461596fcd7a..21ee47b0ac441 100644 --- a/sycl/include/sycl/detail/boolean.hpp +++ b/sycl/include/sycl/detail/boolean.hpp @@ -110,7 +110,8 @@ template <> struct Boolean<1> { // Cast to a signed interger type template operator T() const { - static_assert(is_sgeninteger::value, "Invalid conversion"); + static_assert(std::is_same::value || is_sgeninteger::value, + "Invalid conversion"); return value; } diff --git a/sycl/include/sycl/detail/generic_type_traits.hpp b/sycl/include/sycl/detail/generic_type_traits.hpp index 178a184d1ea7f..b81d7f49393ad 100644 --- a/sycl/include/sycl/detail/generic_type_traits.hpp +++ b/sycl/include/sycl/detail/generic_type_traits.hpp @@ -572,9 +572,37 @@ 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>; + +// TODO: Remove this when common_rel_ret_t is promoted. +template +using internal_host_rel_ret_t = + conditional_t::value, make_singed_integer_t, int>; +#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>; +template using internal_host_rel_ret_t = common_rel_ret_t; +#endif // forward declaration template struct Boolean; @@ -598,11 +626,21 @@ template struct RelationalReturnType { #ifdef __SYCL_DEVICE_ONLY__ using type = Boolean::value>; #else - using type = common_rel_ret_t; + // After changing the return type of scalar relational operations to boolean + // we keep the old representation of the internal implementation of the + // host-side builtins to avoid ABI-breaks. + // TODO: Use common_rel_ret_t when ABI break is allowed and the boolean return + // type for relationals are promoted out of SYCL2020_CONFORMANT_APIS. + // The scalar relational builtins in + // sycl/source/detail/builtins_relational.cpp should likewise be updated + // to return boolean values. + using type = internal_host_rel_ret_t; #endif }; -template using rel_ret_t = typename RelationalReturnType::type; +// Type representing the internal return type of relational builtins. +template +using internal_rel_ret_t = typename RelationalReturnType::type; // Used for any and all built-in functions template struct RelationalTestForSignBitType { @@ -634,7 +672,7 @@ struct RelConverter< using ret_t = common_rel_ret_t; #else using bool_t = Boolean; - using ret_t = rel_ret_t; + using ret_t = internal_rel_ret_t; #endif static ret_t apply(bool_t value) { @@ -653,7 +691,7 @@ struct RelConverter< template struct RelConverter< T, typename detail::enable_if_t::value>> { - using R = rel_ret_t; + using R = internal_rel_ret_t; #ifdef __SYCL_DEVICE_ONLY__ using value_t = bool; #else diff --git a/sycl/test/basic_tests/relational_builtins.cpp b/sycl/test/basic_tests/relational_builtins.cpp new file mode 100644 index 0000000000000..e0a74c260999e --- /dev/null +++ b/sycl/test/basic_tests/relational_builtins.cpp @@ -0,0 +1,353 @@ +// 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 + }); + }); +}