16 namespace ext::oneapi {
18 template <
typename T =
void>
using plus = std::plus<T>;
19 template <
typename T =
void>
using multiplies = std::multiplies<T>;
20 template <
typename T =
void>
using bit_or = std::bit_or<T>;
21 template <
typename T =
void>
using bit_xor = std::bit_xor<T>;
22 template <
typename T =
void>
using bit_and = std::bit_and<T>;
23 template <
typename T =
void>
using maximum = sycl::maximum<T>;
24 template <
typename T =
void>
using minimum = sycl::minimum<T>;
28 #ifdef __SYCL_DEVICE_ONLY__
31 struct GroupOpISigned {};
32 struct GroupOpIUnsigned {};
35 template <
typename T,
typename =
void>
struct GroupOpTag;
38 struct GroupOpTag<T, detail::
enable_if_t<detail::is_sigeninteger<T>::value>> {
39 using type = GroupOpISigned;
43 struct GroupOpTag<T, detail::
enable_if_t<detail::is_sugeninteger<T>::value>> {
44 using type = GroupOpIUnsigned;
48 struct GroupOpTag<T, detail::
enable_if_t<detail::is_sgenfloat<T>::value>> {
49 using type = GroupOpFP;
52 #define __SYCL_CALC_OVERLOAD(GroupTag, SPIRVOperation, BinaryOperation) \
53 template <typename T, __spv::GroupOperation O, __spv::Scope::Flag S> \
54 static T calc(GroupTag, T x, BinaryOperation) { \
55 using ConvertedT = detail::ConvertToOpenCLType_t<T>; \
58 conditional_t<std::is_same<ConvertedT, cl_char>() || \
59 std::is_same<ConvertedT, cl_short>(), \
61 conditional_t<std::is_same<ConvertedT, cl_uchar>() || \
62 std::is_same<ConvertedT, cl_ushort>(), \
63 cl_uint, ConvertedT>>; \
66 __spirv_Group##SPIRVOperation(S, static_cast<unsigned int>(O), Arg); \
71 __SYCL_CALC_OVERLOAD(GroupOpISigned, SMin, sycl::minimum<T>)
72 __SYCL_CALC_OVERLOAD(GroupOpIUnsigned, UMin, sycl::minimum<T>)
73 __SYCL_CALC_OVERLOAD(GroupOpFP, FMin, sycl::minimum<T>)
75 __SYCL_CALC_OVERLOAD(GroupOpISigned, SMax, sycl::maximum<T>)
76 __SYCL_CALC_OVERLOAD(GroupOpIUnsigned, UMax, sycl::maximum<T>)
77 __SYCL_CALC_OVERLOAD(GroupOpFP, FMax, sycl::maximum<T>)
79 __SYCL_CALC_OVERLOAD(GroupOpISigned, IAdd, sycl::plus<T>)
80 __SYCL_CALC_OVERLOAD(GroupOpIUnsigned, IAdd, sycl::plus<T>)
81 __SYCL_CALC_OVERLOAD(GroupOpFP, FAdd, sycl::plus<T>)
83 __SYCL_CALC_OVERLOAD(GroupOpISigned, IMulKHR, sycl::multiplies<T>)
84 __SYCL_CALC_OVERLOAD(GroupOpIUnsigned, IMulKHR, sycl::multiplies<T>)
85 __SYCL_CALC_OVERLOAD(GroupOpFP, FMulKHR, sycl::multiplies<T>)
87 __SYCL_CALC_OVERLOAD(GroupOpISigned, BitwiseOrKHR, sycl::bit_or<T>)
88 __SYCL_CALC_OVERLOAD(GroupOpIUnsigned, BitwiseOrKHR, sycl::bit_or<T>)
89 __SYCL_CALC_OVERLOAD(GroupOpISigned, BitwiseXorKHR, sycl::bit_xor<T>)
90 __SYCL_CALC_OVERLOAD(GroupOpIUnsigned, BitwiseXorKHR, sycl::bit_xor<T>)
91 __SYCL_CALC_OVERLOAD(GroupOpISigned, BitwiseAndKHR, sycl::bit_and<T>)
92 __SYCL_CALC_OVERLOAD(GroupOpIUnsigned, BitwiseAndKHR, sycl::bit_and<T>)
94 #undef __SYCL_CALC_OVERLOAD
97 template <
typename>
class BinaryOperation>
98 static T calc(
typename GroupOpTag<T>::type, T x, BinaryOperation<void>) {
99 return calc<T, O, S>(
typename GroupOpTag<T>::type(), x, BinaryOperation<T>());
#define __SYCL_INLINE_VER_NAMESPACE(X)
typename std::enable_if< B, T >::type enable_if_t
std::multiplies< T > multiplies
sycl::minimum< T > minimum
sycl::maximum< T > maximum
std::bit_xor< T > bit_xor
std::bit_and< T > bit_and
---— Error handling, matching OpenCL plugin semantics.