Skip to content

Commit d3cdb95

Browse files
steffenlarsensteffenlarsen
authored
[SYCL] Fix group algorithms for non-uniform groups, marray and vec (#14364)
This commit makes the following fixes to group algorithms: * Some GroupNonUniform SPIR-V builtins were incorrectly named as they were bundled together with their KHR Group SPIR-V equivalents. These have been renamed to map correctly to the right SPIR-V operations. * `sycl::marray` is now considered when checking for arithmetic types, making it usable in group broadcast operations and the functions that use them. * The representation of `bool` in `sycl::vec` has been changed to unsigned to match the representation picked by `ConvertToOpenCLType`. * The representation of `char` in `sycl::vec` has been changed to to match the representation picked by `ConvertToOpenCLType` to avoid cases where signedness would cause type mismatches. --------- Signed-off-by: Larsen, Steffen <[email protected]>
1 parent 5f418ab commit d3cdb95

File tree

11 files changed

+452
-130
lines changed

11 files changed

+452
-130
lines changed

‎sycl/include/CL/__spirv/spirv_ops.hpp‎

Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1117,6 +1117,14 @@ template <typename ValueT>
11171117
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
11181118
__spirv_GroupNonUniformBitwiseAnd(__spv::Scope::Flag, unsigned int, ValueT);
11191119

1120+
template <typename ValueT>
1121+
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
1122+
__spirv_GroupNonUniformLogicalOr(__spv::Scope::Flag, unsigned int, ValueT);
1123+
1124+
template <typename ValueT>
1125+
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
1126+
__spirv_GroupNonUniformLogicalAnd(__spv::Scope::Flag, unsigned int, ValueT);
1127+
11201128
template <typename ValueT>
11211129
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
11221130
__spirv_GroupNonUniformSMin(__spv::Scope::Flag, unsigned int, ValueT,
@@ -1182,6 +1190,16 @@ __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
11821190
__spirv_GroupNonUniformBitwiseAnd(__spv::Scope::Flag, unsigned int, ValueT,
11831191
unsigned int);
11841192

1193+
template <typename ValueT>
1194+
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
1195+
__spirv_GroupNonUniformLogicalOr(__spv::Scope::Flag, unsigned int, ValueT,
1196+
unsigned int);
1197+
1198+
template <typename ValueT>
1199+
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT
1200+
__spirv_GroupNonUniformLogicalAnd(__spv::Scope::Flag, unsigned int, ValueT,
1201+
unsigned int);
1202+
11851203
extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT void
11861204
__clc_BarrierInitialize(int64_t *state, int32_t expected_count) noexcept;
11871205

‎sycl/include/sycl/builtins_utils_vec.hpp‎

Lines changed: 1 addition & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -10,22 +10,13 @@
1010

1111
#include <sycl/builtins_utils_scalar.hpp>
1212

13+
#include <sycl/detail/type_traits.hpp>
1314
#include <sycl/marray.hpp> // for marray
1415
#include <sycl/types.hpp> // for vec
1516

1617
namespace sycl {
1718
inline namespace _V1 {
1819
namespace detail {
19-
template <typename> struct is_swizzle : std::false_type {};
20-
template <typename VecT, typename OperationLeftT, typename OperationRightT,
21-
template <typename> class OperationCurrentT, int... Indexes>
22-
struct is_swizzle<SwizzleOp<VecT, OperationLeftT, OperationRightT,
23-
OperationCurrentT, Indexes...>> : std::true_type {};
24-
25-
template <typename T> constexpr bool is_swizzle_v = is_swizzle<T>::value;
26-
27-
template <typename T>
28-
constexpr bool is_vec_or_swizzle_v = is_vec_v<T> || is_swizzle_v<T>;
2920

3021
// Utility trait for checking if T's element type is in Ts.
3122
template <typename T, size_t N, typename... Ts>

‎sycl/include/sycl/detail/spirv.hpp‎

Lines changed: 35 additions & 32 deletions
Original file line numberDiff line numberDiff line change
@@ -236,8 +236,8 @@ bool GroupAny(const ext::oneapi::experimental::opportunistic_group &,
236236
template <typename T>
237237
using is_native_broadcast =
238238
std::bool_constant<detail::is_arithmetic<T>::value &&
239-
!std::is_same<T, half>::value &&
240-
!detail::is_vec<T>::value>;
239+
!std::is_same<T, half>::value && !detail::is_vec_v<T> &&
240+
!detail::is_marray_v<T> && !std::is_pointer_v<T>>;
241241

242242
template <typename T, typename IdT = size_t>
243243
using EnableIfNativeBroadcast = std::enable_if_t<
@@ -747,12 +747,15 @@ struct VecTypeIsProhibitedForShuffleEmulation
747747
template <typename T>
748748
using EnableIfNativeShuffle =
749749
std::enable_if_t<detail::is_arithmetic<T>::value &&
750-
!VecTypeIsProhibitedForShuffleEmulation<T>::value,
750+
!VecTypeIsProhibitedForShuffleEmulation<T>::value &&
751+
!detail::is_marray_v<T>,
751752
T>;
752753

753754
template <typename T>
754-
using EnableIfVectorShuffle =
755-
std::enable_if_t<VecTypeIsProhibitedForShuffleEmulation<T>::value, T>;
755+
using EnableIfNonScalarShuffle =
756+
std::enable_if_t<VecTypeIsProhibitedForShuffleEmulation<T>::value ||
757+
detail::is_marray_v<T>,
758+
T>;
756759

757760
#else // ifndef __NVPTX__
758761

@@ -761,8 +764,8 @@ using EnableIfNativeShuffle = std::enable_if_t<
761764
std::is_integral<T>::value && (sizeof(T) <= sizeof(int32_t)), T>;
762765

763766
template <typename T>
764-
using EnableIfVectorShuffle =
765-
std::enable_if_t<detail::is_vector_arithmetic<T>::value, T>;
767+
using EnableIfNonScalarShuffle =
768+
std::enable_if_t<detail::is_nonscalar_arithmetic<T>::value, T>;
766769
#endif // ifndef __NVPTX__
767770

768771
// Bitcast shuffles can be implemented using a single SubgroupShuffle
@@ -780,7 +783,7 @@ template <typename T>
780783
using EnableIfBitcastShuffle =
781784
std::enable_if_t<!(std::is_integral_v<T> &&
782785
(sizeof(T) <= sizeof(int32_t))) &&
783-
!detail::is_vector_arithmetic<T>::value &&
786+
!detail::is_nonscalar_arithmetic<T>::value &&
784787
(std::is_trivially_copyable_v<T> &&
785788
(sizeof(T) == 1 || sizeof(T) == 2 || sizeof(T) == 4)),
786789
T>;
@@ -802,7 +805,7 @@ using EnableIfGenericShuffle =
802805
template <typename T>
803806
using EnableIfGenericShuffle = std::enable_if_t<
804807
!(std::is_integral<T>::value && (sizeof(T) <= sizeof(int32_t))) &&
805-
!detail::is_vector_arithmetic<T>::value &&
808+
!detail::is_nonscalar_arithmetic<T>::value &&
806809
!(std::is_trivially_copyable_v<T> &&
807810
(sizeof(T) == 1 || sizeof(T) == 2 || sizeof(T) == 4)),
808811
T>;
@@ -1024,7 +1027,7 @@ EnableIfNativeShuffle<T> ShuffleUp(GroupT g, T x, uint32_t delta) {
10241027
}
10251028

10261029
template <typename GroupT, typename T>
1027-
EnableIfVectorShuffle<T> Shuffle(GroupT g, T x, id<1> local_id) {
1030+
EnableIfNonScalarShuffle<T> Shuffle(GroupT g, T x, id<1> local_id) {
10281031
T result;
10291032
for (int s = 0; s < x.size(); ++s) {
10301033
result[s] = Shuffle(g, x[s], local_id);
@@ -1033,7 +1036,7 @@ EnableIfVectorShuffle<T> Shuffle(GroupT g, T x, id<1> local_id) {
10331036
}
10341037

10351038
template <typename GroupT, typename T>
1036-
EnableIfVectorShuffle<T> ShuffleXor(GroupT g, T x, id<1> local_id) {
1039+
EnableIfNonScalarShuffle<T> ShuffleXor(GroupT g, T x, id<1> local_id) {
10371040
T result;
10381041
for (int s = 0; s < x.size(); ++s) {
10391042
result[s] = ShuffleXor(g, x[s], local_id);
@@ -1042,7 +1045,7 @@ EnableIfVectorShuffle<T> ShuffleXor(GroupT g, T x, id<1> local_id) {
10421045
}
10431046

10441047
template <typename GroupT, typename T>
1045-
EnableIfVectorShuffle<T> ShuffleDown(GroupT g, T x, uint32_t delta) {
1048+
EnableIfNonScalarShuffle<T> ShuffleDown(GroupT g, T x, uint32_t delta) {
10461049
T result;
10471050
for (int s = 0; s < x.size(); ++s) {
10481051
result[s] = ShuffleDown(g, x[s], delta);
@@ -1051,7 +1054,7 @@ EnableIfVectorShuffle<T> ShuffleDown(GroupT g, T x, uint32_t delta) {
10511054
}
10521055

10531056
template <typename GroupT, typename T>
1054-
EnableIfVectorShuffle<T> ShuffleUp(GroupT g, T x, uint32_t delta) {
1057+
EnableIfNonScalarShuffle<T> ShuffleUp(GroupT g, T x, uint32_t delta) {
10551058
T result;
10561059
for (int s = 0; s < x.size(); ++s) {
10571060
result[s] = ShuffleUp(g, x[s], delta);
@@ -1186,7 +1189,7 @@ ControlBarrier(Group g, memory_scope FenceScope, memory_order Order) {
11861189
}
11871190

11881191
// TODO: Refactor to avoid duplication after design settles
1189-
#define __SYCL_GROUP_COLLECTIVE_OVERLOAD(Instruction) \
1192+
#define __SYCL_GROUP_COLLECTIVE_OVERLOAD(Instruction, GroupExt) \
11901193
template <__spv::GroupOperation Op, typename Group, typename T> \
11911194
inline typename std::enable_if_t< \
11921195
ext::oneapi::experimental::is_fixed_topology_group_v<Group>, T> \
@@ -1201,8 +1204,8 @@ ControlBarrier(Group g, memory_scope FenceScope, memory_order Order) {
12011204
std::is_same<ConvertedT, opencl::cl_ushort>(), \
12021205
opencl::cl_uint, ConvertedT>>; \
12031206
OCLT Arg = x; \
1204-
OCLT Ret = __spirv_Group##Instruction(group_scope<Group>::value, \
1205-
static_cast<unsigned int>(Op), Arg); \
1207+
OCLT Ret = __spirv_Group##Instruction##GroupExt( \
1208+
group_scope<Group>::value, static_cast<unsigned int>(Op), Arg); \
12061209
return Ret; \
12071210
} \
12081211
\
@@ -1286,27 +1289,27 @@ ControlBarrier(Group g, memory_scope FenceScope, memory_order Order) {
12861289
return Ret; \
12871290
}
12881291

1289-
__SYCL_GROUP_COLLECTIVE_OVERLOAD(SMin)
1290-
__SYCL_GROUP_COLLECTIVE_OVERLOAD(UMin)
1291-
__SYCL_GROUP_COLLECTIVE_OVERLOAD(FMin)
1292+
__SYCL_GROUP_COLLECTIVE_OVERLOAD(SMin, )
1293+
__SYCL_GROUP_COLLECTIVE_OVERLOAD(UMin, )
1294+
__SYCL_GROUP_COLLECTIVE_OVERLOAD(FMin, )
12921295

1293-
__SYCL_GROUP_COLLECTIVE_OVERLOAD(SMax)
1294-
__SYCL_GROUP_COLLECTIVE_OVERLOAD(UMax)
1295-
__SYCL_GROUP_COLLECTIVE_OVERLOAD(FMax)
1296+
__SYCL_GROUP_COLLECTIVE_OVERLOAD(SMax, )
1297+
__SYCL_GROUP_COLLECTIVE_OVERLOAD(UMax, )
1298+
__SYCL_GROUP_COLLECTIVE_OVERLOAD(FMax, )
12961299

1297-
__SYCL_GROUP_COLLECTIVE_OVERLOAD(IAdd)
1298-
__SYCL_GROUP_COLLECTIVE_OVERLOAD(FAdd)
1300+
__SYCL_GROUP_COLLECTIVE_OVERLOAD(IAdd, )
1301+
__SYCL_GROUP_COLLECTIVE_OVERLOAD(FAdd, )
12991302

1300-
__SYCL_GROUP_COLLECTIVE_OVERLOAD(IMulKHR)
1301-
__SYCL_GROUP_COLLECTIVE_OVERLOAD(FMulKHR)
1302-
__SYCL_GROUP_COLLECTIVE_OVERLOAD(CMulINTEL)
1303+
__SYCL_GROUP_COLLECTIVE_OVERLOAD(IMul, KHR)
1304+
__SYCL_GROUP_COLLECTIVE_OVERLOAD(FMul, KHR)
1305+
__SYCL_GROUP_COLLECTIVE_OVERLOAD(CMulINTEL, )
13031306

1304-
__SYCL_GROUP_COLLECTIVE_OVERLOAD(BitwiseOrKHR)
1305-
__SYCL_GROUP_COLLECTIVE_OVERLOAD(BitwiseXorKHR)
1306-
__SYCL_GROUP_COLLECTIVE_OVERLOAD(BitwiseAndKHR)
1307+
__SYCL_GROUP_COLLECTIVE_OVERLOAD(BitwiseOr, KHR)
1308+
__SYCL_GROUP_COLLECTIVE_OVERLOAD(BitwiseXor, KHR)
1309+
__SYCL_GROUP_COLLECTIVE_OVERLOAD(BitwiseAnd, KHR)
13071310

1308-
__SYCL_GROUP_COLLECTIVE_OVERLOAD(LogicalAndKHR)
1309-
__SYCL_GROUP_COLLECTIVE_OVERLOAD(LogicalOrKHR)
1311+
__SYCL_GROUP_COLLECTIVE_OVERLOAD(LogicalAnd, KHR)
1312+
__SYCL_GROUP_COLLECTIVE_OVERLOAD(LogicalOr, KHR)
13101313

13111314
template <access::address_space Space, typename T>
13121315
auto GenericCastToPtr(T *Ptr) ->

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

Lines changed: 33 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -156,29 +156,29 @@ template <class T> using marray_element_t = typename T::value_type;
156156
// get_elem_type
157157
// Get the element type of T. If T is a scalar, the element type is considered
158158
// the type of the scalar.
159-
template <typename T, typename = void> struct get_elem_type {
159+
template <typename T, typename = void> struct get_elem_type_unqual {
160160
using type = T;
161161
};
162-
template <typename T, size_t N> struct get_elem_type<marray<T, N>> {
162+
template <typename T, size_t N> struct get_elem_type_unqual<marray<T, N>> {
163163
using type = T;
164164
};
165-
template <typename T, int N> struct get_elem_type<vec<T, N>> {
165+
template <typename T, int N> struct get_elem_type_unqual<vec<T, N>> {
166166
using type = T;
167167
};
168168
template <typename VecT, typename OperationLeftT, typename OperationRightT,
169169
template <typename> class OperationCurrentT, int... Indexes>
170-
struct get_elem_type<SwizzleOp<VecT, OperationLeftT, OperationRightT,
170+
struct get_elem_type_unqual<SwizzleOp<VecT, OperationLeftT, OperationRightT,
171171
OperationCurrentT, Indexes...>> {
172-
using type = typename get_elem_type<std::remove_cv_t<VecT>>::type;
172+
using type = typename get_elem_type_unqual<std::remove_cv_t<VecT>>::type;
173173
};
174174

175175
template <typename ElementType, access::address_space Space,
176176
access::decorated DecorateAddress>
177-
struct get_elem_type<multi_ptr<ElementType, Space, DecorateAddress>> {
177+
struct get_elem_type_unqual<multi_ptr<ElementType, Space, DecorateAddress>> {
178178
using type = ElementType;
179179
};
180180

181-
template <typename ElementType> struct get_elem_type<ElementType *> {
181+
template <typename ElementType> struct get_elem_type_unqual<ElementType *> {
182182
using type = ElementType;
183183
};
184184

@@ -194,10 +194,13 @@ template <typename T>
194194
inline constexpr bool is_ext_vector_v = is_ext_vector<T>::value;
195195

196196
template <typename T>
197-
struct get_elem_type<T, std::enable_if_t<is_ext_vector_v<T>>> {
197+
struct get_elem_type_unqual<T, std::enable_if_t<is_ext_vector_v<T>>> {
198198
using type = decltype(__builtin_reduce_max(std::declval<T>()));
199199
};
200200

201+
template <typename T>
202+
struct get_elem_type : get_elem_type_unqual<std::remove_cv_t<T>> {};
203+
201204
template <typename T> using get_elem_type_t = typename get_elem_type<T>::type;
202205

203206
// change_base_type_t
@@ -295,6 +298,20 @@ template <typename T, int N> struct get_vec_size<sycl::vec<T, N>> {
295298
static constexpr int size = N;
296299
};
297300

301+
// is_swizzle
302+
template <typename> struct is_swizzle : std::false_type {};
303+
template <typename VecT, typename OperationLeftT, typename OperationRightT,
304+
template <typename> class OperationCurrentT, int... Indexes>
305+
struct is_swizzle<SwizzleOp<VecT, OperationLeftT, OperationRightT,
306+
OperationCurrentT, Indexes...>> : std::true_type {};
307+
308+
template <typename T> constexpr bool is_swizzle_v = is_swizzle<T>::value;
309+
310+
// is_swizzle_or_vec_v
311+
312+
template <typename T>
313+
constexpr bool is_vec_or_swizzle_v = is_vec_v<T> || is_swizzle_v<T>;
314+
298315
// is_marray
299316
template <typename> struct is_marray : std::false_type {};
300317
template <typename T, size_t N>
@@ -304,7 +321,7 @@ template <typename T> constexpr bool is_marray_v = is_marray<T>::value;
304321

305322
// is_integral
306323
template <typename T>
307-
struct is_integral : std::is_integral<vector_element_t<T>> {};
324+
struct is_integral : std::is_integral<get_elem_type_t<T>> {};
308325

309326
// is_floating_point
310327
template <typename T>
@@ -314,7 +331,7 @@ template <> struct is_floating_point_impl<half> : std::true_type {};
314331

315332
template <typename T>
316333
struct is_floating_point
317-
: is_floating_point_impl<std::remove_cv_t<vector_element_t<T>>> {};
334+
: is_floating_point_impl<std::remove_cv_t<get_elem_type_t<T>>> {};
318335

319336
// is_arithmetic
320337
template <typename T>
@@ -324,14 +341,17 @@ struct is_arithmetic
324341

325342
template <typename T>
326343
struct is_scalar_arithmetic
327-
: std::bool_constant<!is_vec<T>::value && is_arithmetic<T>::value> {};
344+
: std::bool_constant<!is_vec_or_swizzle_v<T> && !is_ext_vector_v<T> &&
345+
!is_marray_v<T> && is_arithmetic<T>::value> {};
328346

329347
template <typename T>
330348
inline constexpr bool is_scalar_arithmetic_v = is_scalar_arithmetic<T>::value;
331349

332350
template <typename T>
333-
struct is_vector_arithmetic
334-
: std::bool_constant<is_vec<T>::value && is_arithmetic<T>::value> {};
351+
struct is_nonscalar_arithmetic
352+
: std::bool_constant<(is_vec_or_swizzle_v<T> || is_ext_vector_v<T> ||
353+
is_marray_v<T>) &&
354+
is_arithmetic<T>::value> {};
335355

336356
// is_bool
337357
template <typename T>

‎sycl/include/sycl/ext/oneapi/functional.hpp‎

Lines changed: 19 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -79,27 +79,27 @@ __SYCL_CALC_OVERLOAD(GroupOpISigned, IAdd, sycl::plus<T>)
7979
__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, IAdd, sycl::plus<T>)
8080
__SYCL_CALC_OVERLOAD(GroupOpFP, FAdd, sycl::plus<T>)
8181

82-
__SYCL_CALC_OVERLOAD(GroupOpISigned, IMulKHR, sycl::multiplies<T>)
83-
__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, IMulKHR, sycl::multiplies<T>)
84-
__SYCL_CALC_OVERLOAD(GroupOpFP, FMulKHR, sycl::multiplies<T>)
82+
__SYCL_CALC_OVERLOAD(GroupOpISigned, IMul, sycl::multiplies<T>)
83+
__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, IMul, sycl::multiplies<T>)
84+
__SYCL_CALC_OVERLOAD(GroupOpFP, FMul, sycl::multiplies<T>)
8585
__SYCL_CALC_OVERLOAD(GroupOpC, CMulINTEL, sycl::multiplies<T>)
8686

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>)
93-
94-
__SYCL_CALC_OVERLOAD(GroupOpBool, LogicalAndKHR, sycl::logical_and<T>)
95-
__SYCL_CALC_OVERLOAD(GroupOpISigned, LogicalAndKHR, sycl::logical_and<T>)
96-
__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, LogicalAndKHR, sycl::logical_and<T>)
97-
__SYCL_CALC_OVERLOAD(GroupOpFP, LogicalAndKHR, sycl::logical_and<T>)
98-
99-
__SYCL_CALC_OVERLOAD(GroupOpBool, LogicalOrKHR, sycl::logical_or<T>)
100-
__SYCL_CALC_OVERLOAD(GroupOpISigned, LogicalOrKHR, sycl::logical_or<T>)
101-
__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, LogicalOrKHR, sycl::logical_or<T>)
102-
__SYCL_CALC_OVERLOAD(GroupOpFP, LogicalOrKHR, sycl::logical_or<T>)
87+
__SYCL_CALC_OVERLOAD(GroupOpISigned, BitwiseOr, sycl::bit_or<T>)
88+
__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, BitwiseOr, sycl::bit_or<T>)
89+
__SYCL_CALC_OVERLOAD(GroupOpISigned, BitwiseXor, sycl::bit_xor<T>)
90+
__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, BitwiseXor, sycl::bit_xor<T>)
91+
__SYCL_CALC_OVERLOAD(GroupOpISigned, BitwiseAnd, sycl::bit_and<T>)
92+
__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, BitwiseAnd, sycl::bit_and<T>)
93+
94+
__SYCL_CALC_OVERLOAD(GroupOpBool, LogicalAnd, sycl::logical_and<T>)
95+
__SYCL_CALC_OVERLOAD(GroupOpISigned, LogicalAnd, sycl::logical_and<T>)
96+
__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, LogicalAnd, sycl::logical_and<T>)
97+
__SYCL_CALC_OVERLOAD(GroupOpFP, LogicalAnd, sycl::logical_and<T>)
98+
99+
__SYCL_CALC_OVERLOAD(GroupOpBool, LogicalOr, sycl::logical_or<T>)
100+
__SYCL_CALC_OVERLOAD(GroupOpISigned, LogicalOr, sycl::logical_or<T>)
101+
__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, LogicalOr, sycl::logical_or<T>)
102+
__SYCL_CALC_OVERLOAD(GroupOpFP, LogicalOr, sycl::logical_or<T>)
103103

104104
#undef __SYCL_CALC_OVERLOAD
105105

0 commit comments

Comments
 (0)