Skip to content

Commit b5b0787

Browse files
kbobrovskbobrovs
authored
[ESIMD] Fix 'ambiguous operator' error with length 1 simd operands (#4149)
Signed-off-by: kbobrovs <[email protected]>
1 parent a4165b3 commit b5b0787

File tree

4 files changed

+58
-8
lines changed

4 files changed

+58
-8
lines changed

‎sycl/include/sycl/ext/intel/experimental/esimd/detail/simd_view_impl.hpp

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -193,6 +193,8 @@ template <typename BaseTy, typename RegionTy> class simd_view_impl {
193193
}
194194

195195
#define DEF_BINOP(BINOP, OPASSIGN) \
196+
template <class T1 = simd_view_impl, \
197+
class = std::enable_if_t<T1::length != 1>> \
196198
ESIMD_INLINE friend auto operator BINOP(const simd_view_impl &X, \
197199
const value_type &Y) { \
198200
using ComputeTy = detail::compute_type_t<value_type>; \
@@ -202,6 +204,8 @@ template <typename BaseTy, typename RegionTy> class simd_view_impl {
202204
auto V2 = V0 BINOP V1; \
203205
return ComputeTy(V2); \
204206
} \
207+
template <class T1 = simd_view_impl, \
208+
class = std::enable_if_t<T1::length != 1>> \
205209
ESIMD_INLINE friend auto operator BINOP(const value_type &X, \
206210
const simd_view_impl &Y) { \
207211
using ComputeTy = detail::compute_type_t<value_type>; \
@@ -237,12 +241,16 @@ template <typename BaseTy, typename RegionTy> class simd_view_impl {
237241
#undef DEF_BINOP
238242

239243
#define DEF_BITWISE_OP(BITWISE_OP, OPASSIGN) \
244+
template <class T1 = simd_view_impl, \
245+
class = std::enable_if_t<T1::length != 1>> \
240246
ESIMD_INLINE friend auto operator BITWISE_OP(const simd_view_impl &X, \
241247
const value_type &Y) { \
242248
static_assert(std::is_integral<element_type>(), "not integral type"); \
243249
auto V2 = X.read().data() BITWISE_OP Y.data(); \
244250
return simd<element_type, length>(V2); \
245251
} \
252+
template <class T1 = simd_view_impl, \
253+
class = std::enable_if_t<T1::length != 1>> \
246254
ESIMD_INLINE friend auto operator BITWISE_OP(const value_type &X, \
247255
const simd_view_impl &Y) { \
248256
static_assert(std::is_integral<element_type>(), "not integral type"); \

‎sycl/include/sycl/ext/intel/experimental/esimd/simd.hpp

Lines changed: 10 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -233,10 +233,11 @@ template <typename Ty, int N> class simd {
233233
auto V2 = V0 BINOP V1; \
234234
return ComputeTy(V2); \
235235
} \
236-
template <typename T = simd, \
237-
typename = sycl::detail::enable_if_t<T::length == 1>> \
238-
ESIMD_INLINE friend auto operator BINOP(const simd &X, const Ty &Y) { \
239-
return X BINOP simd(Y); \
236+
template <typename T1, typename T = simd, \
237+
typename = sycl::detail::enable_if_t<T::length == 1 && \
238+
std::is_arithmetic_v<T1>>> \
239+
ESIMD_INLINE friend auto operator BINOP(const simd &X, T1 Y) { \
240+
return X BINOP simd((Ty)Y); \
240241
} \
241242
ESIMD_INLINE friend simd &operator OPASSIGN(simd &LHS, const simd &RHS) { \
242243
using ComputeTy = detail::compute_type_t<simd>; \
@@ -272,10 +273,11 @@ template <typename Ty, int N> class simd {
272273
mask_type_t<N> M(1); \
273274
return M & detail::convert<mask_type_t<N>>(R); \
274275
} \
275-
template <typename T = simd, \
276-
typename = sycl::detail::enable_if_t<T::length == 1>> \
277-
ESIMD_INLINE friend bool operator RELOP(const simd &X, const Ty &Y) { \
278-
return (Ty)X RELOP Y; \
276+
template <typename T1, typename T = simd, \
277+
typename = sycl::detail::enable_if_t<(T::length == 1) && \
278+
std::is_arithmetic_v<T1>>> \
279+
ESIMD_INLINE friend bool operator RELOP(const simd &X, T1 Y) { \
280+
return (Ty)X RELOP(Ty) Y; \
279281
}
280282

281283
DEF_RELOP(>)

‎sycl/include/sycl/ext/intel/experimental/esimd/simd_view.hpp

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -82,6 +82,8 @@ class simd_view : public detail::simd_view_impl<BaseTy, RegionTy> {
8282
DEF_RELOP(<=)
8383
DEF_RELOP(==)
8484
DEF_RELOP(!=)
85+
86+
#undef DEF_RELOP
8587
};
8688

8789
/// This is a specialization of simd_view class with a single element.
@@ -132,6 +134,8 @@ class simd_view<BaseTy, region_base_1<Is2D, T, StrideY, StrideX>>
132134
DEF_RELOP(<=)
133135
DEF_RELOP(==)
134136
DEF_RELOP(!=)
137+
138+
#undef DEF_RELOP
135139
};
136140

137141
} // namespace esimd

‎sycl/test/esimd/regression/len1_operators.cpp

Lines changed: 36 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,36 @@
1+
// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s
2+
// expected-no-diagnostics
3+
4+
// This test checks that compiler does not report 'ambiguous operator' error
5+
// when compiling simd or simd_view operations with lenth = 1.
6+
7+
#include <sycl/ext/intel/experimental/esimd.hpp>
8+
9+
#include <cstdint>
10+
11+
using namespace sycl::ext::intel::experimental::esimd;
12+
13+
template <typename T1, typename T2>
14+
void test_esimd_ops(simd<T1, 1> a, T2 b, T2 w) SYCL_ESIMD_FUNCTION {
15+
T2 c1 = a[0] * w + b;
16+
T2 c2 = a[0] * T2{2} + b;
17+
T2 c3 = T2{2} * a[0] + b;
18+
T2 d1 = a[0] ^ w;
19+
T2 d2 = a[0] ^ T2 { 2 };
20+
T2 d3 = T2{2} ^ a[0];
21+
auto e1 = a[0] < w;
22+
auto e2 = a[0] < T2{2};
23+
auto e3 = T2{2} < a[0];
24+
simd<T1, 1> z{4};
25+
auto f1 = a[0] ^ z;
26+
auto f2 = z ^ a[0];
27+
auto f3 = a[0] < z;
28+
auto f4 = z < a[0];
29+
}
30+
31+
void foo() SYCL_ESIMD_FUNCTION {
32+
test_esimd_ops(simd<uint32_t, 1>(3), (int)1, (int)9);
33+
test_esimd_ops(simd<int, 1>(3), (uint32_t)1, (uint32_t)9);
34+
test_esimd_ops(simd<uint16_t, 1>(3), 1, 9);
35+
test_esimd_ops(simd<int16_t, 1>(3), 1, 9);
36+
}

0 commit comments

Comments
 (0)