17 inline namespace _V1 {
22 if constexpr (std::is_same_v<T, half>)
23 return static_cast<float>(
x);
28 #if defined(__GNUC__) && !defined(__clang__)
32 #pragma GCC diagnostic push
33 #pragma GCC diagnostic ignored "-Wmaybe-uninitialized"
36 #define REL_BUILTIN_CUSTOM(NUM_ARGS, NAME, ...) \
37 template <typename... Ts> static auto NAME##_host_impl(Ts... xs) { \
38 using namespace detail; \
39 if constexpr ((... || is_vec_v<Ts>)) { \
40 return builtin_delegate_rel_impl( \
41 [](auto... xs) { return NAME##_host_impl(xs...); }, xs...); \
43 return __VA_ARGS__(xs...); \
46 EXPORT_SCALAR_AND_VEC_1_16(NUM_ARGS, NAME, FP_TYPES)
47 #define REL_BUILTIN(NUM_ARGS, NAME) \
48 REL_BUILTIN_CUSTOM(NUM_ARGS, NAME, [](auto... xs) { \
49 return std::NAME(process_arg_for_macos(xs)...); \
52 #if defined(__GNUC__) && !defined(__clang__)
53 #pragma GCC diagnostic pop
58 ([](auto
x, auto
y) {
return x !=
y; }))
61 ([](auto
x, auto
y) {
return x >=
y; }))
64 ([](auto
x, auto
y) {
return x <=
y; }))
66 ([](
auto x,
auto y) {
return x < y || x >
y; }))
73 ([](auto
x, auto
y) {
return !sycl::isunordered(
x,
y); }))
75 #define _SYCL_BUILTINS_GCC_VER \
76 (__GNUC__ * 10000 + __GNUC_MINOR__ * 100 + __GNUC_PATCHLEVEL__)
78 #if defined(__GNUC__) && !defined(__clang__) && \
79 ((_SYCL_BUILTINS_GCC_VER >= 100000 && _SYCL_BUILTINS_GCC_VER < 110000) || \
80 (_SYCL_BUILTINS_GCC_VER >= 110000 && _SYCL_BUILTINS_GCC_VER < 110500) || \
81 (_SYCL_BUILTINS_GCC_VER >= 120000 && _SYCL_BUILTINS_GCC_VER < 120400) || \
82 (_SYCL_BUILTINS_GCC_VER >= 130000 && _SYCL_BUILTINS_GCC_VER < 130300))
87 #define GCC_PR112816_DISABLE_OPT \
88 _Pragma("GCC push_options") _Pragma("GCC optimize(\"-O1\")")
89 #define GCC_PR112816_RESTORE_OPT _Pragma("GCC pop_options")
91 #define GCC_PR112816_DISABLE_OPT
92 #define GCC_PR112816_RESTORE_OPT
101 #undef GCC_PR112816_RESTORE_OPT
102 #undef GCC_PR112816_DISABLE_OPT
103 #undef _SYCL_BUILTINS_GCC_VER
106 using T0 = decltype(
x);
107 using T1 = decltype(
y);
108 using T2 = decltype(
z);
109 constexpr
size_t N =
sizeof(T0) * 8;
110 using bitset = std::bitset<N>;
112 static_assert(std::is_same_v<T0, T1> && std::is_same_v<T1, T2> &&
113 detail::is_scalar_arithmetic_v<T0>);
117 unsigned long,
unsigned long long>>;
118 static_assert(
sizeof(utype) ==
sizeof(T0));
119 bitset bx(bit_cast<utype>(
x)), by(bit_cast<utype>(
y)), bz(bit_cast<utype>(
z));
120 bitset res = (bz & by) | (~bz & bx);
121 unsigned long long ures = res.to_ullong();
123 return bit_cast<T0>(
static_cast<utype
>(ures));
#define FIXED_WIDTH_INTEGER_TYPES
#define FOR_EACH2(BASE_CASE, FIXED1, FIXED2,...)
#define EXPORT_SCALAR(NUM_ARGS, NAME, TYPE)
#define HOST_IMPL(NAME,...)
#define EXPORT_VEC_1_16(NUM_ARGS, NAME,...)
boost::mp11::mp_list< T... > type_list
typename make_type_impl< T, TL >::type make_type_t
std::enable_if_t< std::is_same_v< T, bfloat16 >, bool > isnan(T x)
auto auto autodecltype(x) z
static auto process_arg_for_macos(T x)
REL_BUILTIN_CUSTOM(TWO_ARGS, isequal,([](auto x, auto y) { return x==y;})) REL_BUILTIN_CUSTOM(TWO_ARGS
#define REL_BUILTIN(NUM_ARGS, NAME)
#define GCC_PR112816_RESTORE_OPT
#define GCC_PR112816_DISABLE_OPT