Skip to content

Commit 3dae81e

Browse files
aelovikov-intelaelovikov-intel
committed
[NFC][SYCL] Add is_op_available_for_type helper for sycl::vec impl
Continuing refactoring before implementing functional changes under preview guard. Not that this availability would need to be shared between vec/swizzle, so it makes sense to outline it outside of macros.
1 parent 36ce10e commit 3dae81e

File tree

2 files changed

+95
-44
lines changed

2 files changed

+95
-44
lines changed

‎sycl/include/sycl/detail/vector_arith.hpp

Lines changed: 91 additions & 40 deletions
Original file line numberDiff line numberDiff line change
@@ -62,6 +62,52 @@ struct UnaryPlus {
6262
// Tag to map/templatize the mixin for prefix/postfix inc/dec operators.
6363
struct IncDec {};
6464

65+
template <class T> static constexpr bool not_fp = !is_vgenfloat_v<T>;
66+
67+
// To provide information about operators availability depending on vec/swizzle
68+
// element type.
69+
template <typename Op, typename T>
70+
inline constexpr bool is_op_available_for_type = false;
71+
72+
#define __SYCL_OP_AVAILABILITY(OP, COND) \
73+
template <typename T> \
74+
inline constexpr bool is_op_available_for_type<OP, T> = COND;
75+
76+
// clang-format off
77+
__SYCL_OP_AVAILABILITY(std::plus<void> , true)
78+
__SYCL_OP_AVAILABILITY(std::minus<void> , true)
79+
__SYCL_OP_AVAILABILITY(std::multiplies<void> , true)
80+
__SYCL_OP_AVAILABILITY(std::divides<void> , true)
81+
__SYCL_OP_AVAILABILITY(std::modulus<void> , not_fp<T>)
82+
83+
__SYCL_OP_AVAILABILITY(std::bit_and<void> , not_fp<T>)
84+
__SYCL_OP_AVAILABILITY(std::bit_or<void> , not_fp<T>)
85+
__SYCL_OP_AVAILABILITY(std::bit_xor<void> , not_fp<T>)
86+
87+
__SYCL_OP_AVAILABILITY(std::equal_to<void> , true)
88+
__SYCL_OP_AVAILABILITY(std::not_equal_to<void> , true)
89+
__SYCL_OP_AVAILABILITY(std::less<void> , true)
90+
__SYCL_OP_AVAILABILITY(std::greater<void> , true)
91+
__SYCL_OP_AVAILABILITY(std::less_equal<void> , true)
92+
__SYCL_OP_AVAILABILITY(std::greater_equal<void> , true)
93+
94+
__SYCL_OP_AVAILABILITY(std::logical_and<void> , not_fp<T>)
95+
__SYCL_OP_AVAILABILITY(std::logical_or<void> , not_fp<T>)
96+
97+
__SYCL_OP_AVAILABILITY(ShiftLeft , not_fp<T>)
98+
__SYCL_OP_AVAILABILITY(ShiftRight , not_fp<T>)
99+
100+
// Unary
101+
__SYCL_OP_AVAILABILITY(std::negate<void> , true)
102+
__SYCL_OP_AVAILABILITY(std::logical_not<void> , true)
103+
__SYCL_OP_AVAILABILITY(std::bit_not<void> , not_fp<T>)
104+
__SYCL_OP_AVAILABILITY(UnaryPlus , true)
105+
106+
__SYCL_OP_AVAILABILITY(IncDec , true)
107+
// clang-format on
108+
109+
#undef __SYCL_OP_AVAILABILITY
110+
65111
template <typename SelfOperandTy> struct IncDecImpl {
66112
using element_type = typename from_incomplete<SelfOperandTy>::element_type;
67113
using vec_t = simplify_if_swizzle_t<std::remove_const_t<SelfOperandTy>>;
@@ -213,7 +259,8 @@ template <typename Self> struct VecOperators {
213259
template <typename Op>
214260
struct OpMixin<Op, std::enable_if_t<std::is_same_v<Op, std::bit_not<void>>>> {
215261
template <typename T = typename from_incomplete<Self>::element_type>
216-
friend std::enable_if_t<!is_vgenfloat_v<T>, Self> operator~(const Self &v) {
262+
friend std::enable_if_t<is_op_available_for_type<Op, T>, Self>
263+
operator~(const Self &v) {
217264
return apply<std::bit_not<void>>(v);
218265
}
219266
};
@@ -233,32 +280,33 @@ template <typename Self> struct VecOperators {
233280
#error "Undefine __SYCL_BINOP macro"
234281
#endif
235282

236-
#define __SYCL_BINOP(BINOP, OPASSIGN, COND, FUNCTOR) \
283+
#define __SYCL_BINOP(BINOP, OPASSIGN, FUNCTOR) \
237284
template <typename T = DataT> \
238-
friend std::enable_if_t<(COND), vec_t> operator BINOP(const vec_t & Lhs, \
239-
const vec_t & Rhs) { \
285+
friend std::enable_if_t<is_op_available_for_type<FUNCTOR, T>, vec_t> \
286+
operator BINOP(const vec_t & Lhs, const vec_t & Rhs) { \
240287
return VecOperators<vec_t>::template apply<FUNCTOR>(Lhs, Rhs); \
241288
} \
242289
\
243290
template <typename T = DataT> \
244-
friend std::enable_if_t<(COND), vec_t> operator BINOP(const vec_t & Lhs, \
245-
const DataT & Rhs) { \
291+
friend std::enable_if_t<is_op_available_for_type<FUNCTOR, T>, vec_t> \
292+
operator BINOP(const vec_t & Lhs, const DataT & Rhs) { \
246293
return Lhs BINOP vec_t(Rhs); \
247294
} \
248295
template <typename T = DataT> \
249-
friend std::enable_if_t<(COND), vec_t> operator BINOP(const DataT & Lhs, \
250-
const vec_t & Rhs) { \
296+
friend std::enable_if_t<is_op_available_for_type<FUNCTOR, T>, vec_t> \
297+
operator BINOP(const DataT & Lhs, const vec_t & Rhs) { \
251298
return vec_t(Lhs) BINOP Rhs; \
252299
} \
253300
template <typename T = DataT> \
254-
friend std::enable_if_t<(COND), vec_t> &operator OPASSIGN( \
255-
vec_t & Lhs, const vec_t & Rhs) { \
301+
friend std::enable_if_t<is_op_available_for_type<FUNCTOR, T>, vec_t> \
302+
&operator OPASSIGN(vec_t & Lhs, const vec_t & Rhs) { \
256303
Lhs = Lhs BINOP Rhs; \
257304
return Lhs; \
258305
} \
259306
template <int Num = NumElements, typename T = DataT> \
260-
friend std::enable_if_t<(Num != 1) && (COND), vec_t &> operator OPASSIGN( \
261-
vec_t & Lhs, const DataT & Rhs) { \
307+
friend std::enable_if_t< \
308+
(Num != 1) && (is_op_available_for_type<FUNCTOR, T>), vec_t &> \
309+
operator OPASSIGN(vec_t & Lhs, const DataT & Rhs) { \
262310
Lhs = Lhs BINOP vec_t(Rhs); \
263311
return Lhs; \
264312
}
@@ -277,55 +325,58 @@ class vec_arith : public VecOperators<vec<DataT, NumElements>>::Combined {
277325
#error "Undefine __SYCL_RELLOGOP macro."
278326
#endif
279327

280-
#define __SYCL_RELLOGOP(RELLOGOP, COND, FUNCTOR) \
328+
#define __SYCL_RELLOGOP(RELLOGOP, FUNCTOR) \
281329
template <typename T = DataT> \
282-
friend std::enable_if_t<(COND), vec<ocl_t, NumElements>> operator RELLOGOP( \
283-
const vec_t & Lhs, const vec_t & Rhs) { \
330+
friend std::enable_if_t<is_op_available_for_type<FUNCTOR, T>, \
331+
vec<ocl_t, NumElements>> \
332+
operator RELLOGOP(const vec_t & Lhs, const vec_t & Rhs) { \
284333
return VecOperators<vec_t>::template apply<FUNCTOR>(Lhs, Rhs); \
285334
} \
286335
\
287336
template <typename T = DataT> \
288-
friend std::enable_if_t<(COND), vec<ocl_t, NumElements>> operator RELLOGOP( \
289-
const vec_t & Lhs, const DataT & Rhs) { \
337+
friend std::enable_if_t<is_op_available_for_type<FUNCTOR, T>, \
338+
vec<ocl_t, NumElements>> \
339+
operator RELLOGOP(const vec_t & Lhs, const DataT & Rhs) { \
290340
return Lhs RELLOGOP vec_t(Rhs); \
291341
} \
292342
template <typename T = DataT> \
293-
friend std::enable_if_t<(COND), vec<ocl_t, NumElements>> operator RELLOGOP( \
294-
const DataT & Lhs, const vec_t & Rhs) { \
343+
friend std::enable_if_t<is_op_available_for_type<FUNCTOR, T>, \
344+
vec<ocl_t, NumElements>> \
345+
operator RELLOGOP(const DataT & Lhs, const vec_t & Rhs) { \
295346
return vec_t(Lhs) RELLOGOP Rhs; \
296347
}
297348

298349
// OP is: ==, !=, <, >, <=, >=, &&, ||
299350
// vec<RET, NumElements> operatorOP(const vec<DataT, NumElements> &Rhs) const;
300351
// vec<RET, NumElements> operatorOP(const DataT &Rhs) const;
301-
__SYCL_RELLOGOP(==, true, std::equal_to<void>)
302-
__SYCL_RELLOGOP(!=, true, std::not_equal_to<void>)
303-
__SYCL_RELLOGOP(>, true, std::greater<void>)
304-
__SYCL_RELLOGOP(<, true, std::less<void>)
305-
__SYCL_RELLOGOP(>=, true, std::greater_equal<void>)
306-
__SYCL_RELLOGOP(<=, true, std::less_equal<void>)
352+
__SYCL_RELLOGOP(==, std::equal_to<void>)
353+
__SYCL_RELLOGOP(!=, std::not_equal_to<void>)
354+
__SYCL_RELLOGOP(>, std::greater<void>)
355+
__SYCL_RELLOGOP(<, std::less<void>)
356+
__SYCL_RELLOGOP(>=, std::greater_equal<void>)
357+
__SYCL_RELLOGOP(<=, std::less_equal<void>)
307358

308359
// Only available to integral types.
309-
__SYCL_RELLOGOP(&&, (!detail::is_vgenfloat_v<T>), std::logical_and<void>)
310-
__SYCL_RELLOGOP(||, (!detail::is_vgenfloat_v<T>), std::logical_or<void>)
360+
__SYCL_RELLOGOP(&&, std::logical_and<void>)
361+
__SYCL_RELLOGOP(||, std::logical_or<void>)
311362
#undef __SYCL_RELLOGOP
312363
#undef RELLOGOP_BASE
313364

314365
// Binary operations on sycl::vec<> for all types except std::byte.
315-
__SYCL_BINOP(+, +=, true, std::plus<void>)
316-
__SYCL_BINOP(-, -=, true, std::minus<void>)
317-
__SYCL_BINOP(*, *=, true, std::multiplies<void>)
318-
__SYCL_BINOP(/, /=, true, std::divides<void>)
366+
__SYCL_BINOP(+, +=, std::plus<void>)
367+
__SYCL_BINOP(-, -=, std::minus<void>)
368+
__SYCL_BINOP(*, *=, std::multiplies<void>)
369+
__SYCL_BINOP(/, /=, std::divides<void>)
319370

320371
// The following OPs are available only when: DataT != cl_float &&
321372
// DataT != cl_double && DataT != cl_half && DataT != BF16.
322-
__SYCL_BINOP(%, %=, (!detail::is_vgenfloat_v<T>), std::modulus<void>)
373+
__SYCL_BINOP(%, %=, std::modulus<void>)
323374
// Bitwise operations are allowed for std::byte.
324-
__SYCL_BINOP(|, |=, (!detail::is_vgenfloat_v<DataT>), std::bit_or<void>)
325-
__SYCL_BINOP(&, &=, (!detail::is_vgenfloat_v<DataT>), std::bit_and<void>)
326-
__SYCL_BINOP(^, ^=, (!detail::is_vgenfloat_v<DataT>), std::bit_xor<void>)
327-
__SYCL_BINOP(>>, >>=, (!detail::is_vgenfloat_v<DataT>), ShiftRight)
328-
__SYCL_BINOP(<<, <<=, (!detail::is_vgenfloat_v<DataT>), ShiftLeft)
375+
__SYCL_BINOP(|, |=, std::bit_or<void>)
376+
__SYCL_BINOP(&, &=, std::bit_and<void>)
377+
__SYCL_BINOP(^, ^=, std::bit_xor<void>)
378+
__SYCL_BINOP(>>, >>=, ShiftRight)
379+
__SYCL_BINOP(<<, <<=, ShiftLeft)
329380

330381
// friends
331382
template <typename T1, int T2> friend class __SYCL_EBO vec;
@@ -376,9 +427,9 @@ class vec_arith<std::byte, NumElements>
376427
return Lhs;
377428
}
378429

379-
__SYCL_BINOP(|, |=, true, std::bit_or<void>)
380-
__SYCL_BINOP(&, &=, true, std::bit_and<void>)
381-
__SYCL_BINOP(^, ^=, true, std::bit_xor<void>)
430+
__SYCL_BINOP(|, |=, std::bit_or<void>)
431+
__SYCL_BINOP(&, &=, std::bit_and<void>)
432+
__SYCL_BINOP(^, ^=, std::bit_xor<void>)
382433

383434
// friends
384435
template <typename T1, int T2> friend class __SYCL_EBO vec;

‎sycl/test/check_device_code/vector/vector_math_ops.cpp

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -119,7 +119,7 @@ SYCL_EXTERNAL auto TestAdd(vec<half, 3> a, vec<half, 3> b) { return a + b; }
119119
// CHECK: for.cond.i.i:
120120
// CHECK-NEXT: [[I_0_I_I:%.*]] = phi i64 [ 0, [[ENTRY:%.*]] ], [ [[INC_I_I:%.*]], [[FOR_BODY_I_I:%.*]] ]
121121
// CHECK-NEXT: [[CMP_I_I:%.*]] = icmp samesign ult i64 [[I_0_I_I]], 3
122-
// CHECK-NEXT: br i1 [[CMP_I_I]], label [[FOR_BODY_I_I]], label [[_ZN4SYCL3_V16DETAILPLINS0_3EXT6ONEAPI8BFLOAT16EEENS0_3VECIS5_LI3EEERKS7_S9__EXIT:%.*]]
122+
// CHECK-NEXT: br i1 [[CMP_I_I]], label [[FOR_BODY_I_I]], label [[_ZN4SYCL3_V16DETAILPLINS0_3EXT6ONEAPI8BFLOAT16EEENST9ENABLE_IFIX24IS_OP_AVAILABLE_FOR_TYPEIST4PLUSIVET_EENS0_3VECIS5_LI3EEEE4TYPEERKSB_SF__EXIT:%.*]]
123123
// CHECK: for.body.i.i:
124124
// CHECK-NEXT: [[ARRAYIDX_I_I_I_I_I:%.*]] = getelementptr inbounds nuw [4 x %"class.sycl::_V1::ext::oneapi::bfloat16"], ptr addrspace(4) [[A_ASCAST]], i64 0, i64 [[I_0_I_I]]
125125
// CHECK-NEXT: [[ARRAYIDX_I_I_I12_I_I:%.*]] = getelementptr inbounds nuw [4 x %"class.sycl::_V1::ext::oneapi::bfloat16"], ptr addrspace(4) [[B_ASCAST]], i64 0, i64 [[I_0_I_I]]
@@ -134,7 +134,7 @@ SYCL_EXTERNAL auto TestAdd(vec<half, 3> a, vec<half, 3> b) { return a + b; }
134134
// CHECK-NEXT: store i16 [[CALL_I_I3_I_I_I_I]], ptr [[ARRAYIDX_I_I_I14_I_I]], align 2, !tbaa [[TBAA88:![0-9]+]], !noalias [[META79]]
135135
// CHECK-NEXT: [[INC_I_I]] = add nuw nsw i64 [[I_0_I_I]], 1
136136
// CHECK-NEXT: br label [[FOR_COND_I_I]], !llvm.loop [[LOOP90:![0-9]+]]
137-
// CHECK: _ZN4sycl3_V16detailplINS0_3ext6oneapi8bfloat16EEENS0_3vecIS5_Li3EEERKS7_S9_.exit:
137+
// CHECK: _ZN4sycl3_V16detailplINS0_3ext6oneapi8bfloat16EEENSt9enable_ifIX24is_op_available_for_typeISt4plusIvET_EENS0_3vecIS5_Li3EEEE4typeERKSB_SF_.exit:
138138
// CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[RES_I_I]], align 8, !noalias [[META79]]
139139
// CHECK-NEXT: store i64 [[TMP0]], ptr addrspace(4) [[AGG_RESULT]], align 8, !alias.scope [[META79]]
140140
// CHECK-NEXT: call void @llvm.lifetime.end.p0(i64 8, ptr nonnull [[RES_I_I]]), !noalias [[META73]]
@@ -223,7 +223,7 @@ SYCL_EXTERNAL auto TestGreaterThan(vec<half, 8> a, vec<half, 8> b) {
223223
// CHECK: for.cond.i.i:
224224
// CHECK-NEXT: [[I_0_I_I:%.*]] = phi i64 [ 0, [[ENTRY:%.*]] ], [ [[INC_I_I:%.*]], [[FOR_BODY_I_I:%.*]] ]
225225
// CHECK-NEXT: [[CMP_I_I:%.*]] = icmp samesign ult i64 [[I_0_I_I]], 4
226-
// CHECK-NEXT: br i1 [[CMP_I_I]], label [[FOR_BODY_I_I]], label [[_ZN4SYCL3_V16DETAILGTINS0_3EXT6ONEAPI8BFLOAT16EEENS0_3VECISLI4EEERKNS6_IS5_LI4EEESA__EXIT:%.*]]
226+
// CHECK-NEXT: br i1 [[CMP_I_I]], label [[FOR_BODY_I_I]], label [[_ZN4SYCL3_V16DETAILGTINS0_3EXT6ONEAPI8BFLOAT16EEENST9ENABLE_IFIX24IS_OP_AVAILABLE_FOR_TYPEIST7GREATERIVET_EENS0_3VECISLI4EEEE4TYPEERKNSA_IS5_LI4EEESG__EXIT:%.*]]
227227
// CHECK: for.body.i.i:
228228
// CHECK-NEXT: [[ARRAYIDX_I_I_I_I_I:%.*]] = getelementptr inbounds nuw [4 x %"class.sycl::_V1::ext::oneapi::bfloat16"], ptr addrspace(4) [[A_ASCAST]], i64 0, i64 [[I_0_I_I]]
229229
// CHECK-NEXT: [[ARRAYIDX_I_I_I14_I_I:%.*]] = getelementptr inbounds nuw [4 x %"class.sycl::_V1::ext::oneapi::bfloat16"], ptr addrspace(4) [[B_ASCAST]], i64 0, i64 [[I_0_I_I]]
@@ -235,7 +235,7 @@ SYCL_EXTERNAL auto TestGreaterThan(vec<half, 8> a, vec<half, 8> b) {
235235
// CHECK-NEXT: store i16 [[CONV6_I_I]], ptr [[ARRAYIDX_I_I_I16_I_I]], align 2, !tbaa [[TBAA88]], !noalias [[META127]]
236236
// CHECK-NEXT: [[INC_I_I]] = add nuw nsw i64 [[I_0_I_I]], 1
237237
// CHECK-NEXT: br label [[FOR_COND_I_I]], !llvm.loop [[LOOP128:![0-9]+]]
238-
// CHECK: _ZN4sycl3_V16detailgtINS0_3ext6oneapi8bfloat16EEENS0_3vecIsLi4EEERKNS6_IS5_Li4EEESA_.exit:
238+
// CHECK: _ZN4sycl3_V16detailgtINS0_3ext6oneapi8bfloat16EEENSt9enable_ifIX24is_op_available_for_typeISt7greaterIvET_EENS0_3vecIsLi4EEEE4typeERKNSA_IS5_Li4EEESG_.exit:
239239
// CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[RES_I_I]], align 8, !noalias [[META127]]
240240
// CHECK-NEXT: store i64 [[TMP0]], ptr addrspace(4) [[AGG_RESULT]], align 8, !alias.scope [[META127]]
241241
// CHECK-NEXT: call void @llvm.lifetime.end.p0(i64 8, ptr nonnull [[RES_I_I]]), !noalias [[META121]]

0 commit comments

Comments
 (0)