diff --git a/sycl/include/sycl/ext/intel/experimental/esimd/detail/simd_view_impl.hpp b/sycl/include/sycl/ext/intel/experimental/esimd/detail/simd_view_impl.hpp index 2724b3f51ccb0..6678bf378c309 100644 --- a/sycl/include/sycl/ext/intel/experimental/esimd/detail/simd_view_impl.hpp +++ b/sycl/include/sycl/ext/intel/experimental/esimd/detail/simd_view_impl.hpp @@ -193,6 +193,8 @@ template class simd_view_impl { } #define DEF_BINOP(BINOP, OPASSIGN) \ + template > \ ESIMD_INLINE friend auto operator BINOP(const simd_view_impl &X, \ const value_type &Y) { \ using ComputeTy = detail::compute_type_t; \ @@ -202,6 +204,8 @@ template class simd_view_impl { auto V2 = V0 BINOP V1; \ return ComputeTy(V2); \ } \ + template > \ ESIMD_INLINE friend auto operator BINOP(const value_type &X, \ const simd_view_impl &Y) { \ using ComputeTy = detail::compute_type_t; \ @@ -237,12 +241,16 @@ template class simd_view_impl { #undef DEF_BINOP #define DEF_BITWISE_OP(BITWISE_OP, OPASSIGN) \ + template > \ ESIMD_INLINE friend auto operator BITWISE_OP(const simd_view_impl &X, \ const value_type &Y) { \ static_assert(std::is_integral(), "not integral type"); \ auto V2 = X.read().data() BITWISE_OP Y.data(); \ return simd(V2); \ } \ + template > \ ESIMD_INLINE friend auto operator BITWISE_OP(const value_type &X, \ const simd_view_impl &Y) { \ static_assert(std::is_integral(), "not integral type"); \ diff --git a/sycl/include/sycl/ext/intel/experimental/esimd/simd.hpp b/sycl/include/sycl/ext/intel/experimental/esimd/simd.hpp index 5b38a8d677111..4f64af0af6e7d 100644 --- a/sycl/include/sycl/ext/intel/experimental/esimd/simd.hpp +++ b/sycl/include/sycl/ext/intel/experimental/esimd/simd.hpp @@ -233,10 +233,11 @@ template class simd { auto V2 = V0 BINOP V1; \ return ComputeTy(V2); \ } \ - template > \ - ESIMD_INLINE friend auto operator BINOP(const simd &X, const Ty &Y) { \ - return X BINOP simd(Y); \ + template >> \ + ESIMD_INLINE friend auto operator BINOP(const simd &X, T1 Y) { \ + return X BINOP simd((Ty)Y); \ } \ ESIMD_INLINE friend simd &operator OPASSIGN(simd &LHS, const simd &RHS) { \ using ComputeTy = detail::compute_type_t; \ @@ -272,10 +273,11 @@ template class simd { mask_type_t M(1); \ return M & detail::convert>(R); \ } \ - template > \ - ESIMD_INLINE friend bool operator RELOP(const simd &X, const Ty &Y) { \ - return (Ty)X RELOP Y; \ + template >> \ + ESIMD_INLINE friend bool operator RELOP(const simd &X, T1 Y) { \ + return (Ty)X RELOP(Ty) Y; \ } DEF_RELOP(>) diff --git a/sycl/include/sycl/ext/intel/experimental/esimd/simd_view.hpp b/sycl/include/sycl/ext/intel/experimental/esimd/simd_view.hpp index c4d9aa5d32274..a2915c713d53f 100644 --- a/sycl/include/sycl/ext/intel/experimental/esimd/simd_view.hpp +++ b/sycl/include/sycl/ext/intel/experimental/esimd/simd_view.hpp @@ -82,6 +82,8 @@ class simd_view : public detail::simd_view_impl { DEF_RELOP(<=) DEF_RELOP(==) DEF_RELOP(!=) + +#undef DEF_RELOP }; /// This is a specialization of simd_view class with a single element. @@ -132,6 +134,8 @@ class simd_view> DEF_RELOP(<=) DEF_RELOP(==) DEF_RELOP(!=) + +#undef DEF_RELOP }; } // namespace esimd diff --git a/sycl/test/esimd/regression/len1_operators.cpp b/sycl/test/esimd/regression/len1_operators.cpp new file mode 100644 index 0000000000000..2ba42b2abaa01 --- /dev/null +++ b/sycl/test/esimd/regression/len1_operators.cpp @@ -0,0 +1,36 @@ +// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s +// expected-no-diagnostics + +// This test checks that compiler does not report 'ambiguous operator' error +// when compiling simd or simd_view operations with lenth = 1. + +#include + +#include + +using namespace sycl::ext::intel::experimental::esimd; + +template +void test_esimd_ops(simd a, T2 b, T2 w) SYCL_ESIMD_FUNCTION { + T2 c1 = a[0] * w + b; + T2 c2 = a[0] * T2{2} + b; + T2 c3 = T2{2} * a[0] + b; + T2 d1 = a[0] ^ w; + T2 d2 = a[0] ^ T2 { 2 }; + T2 d3 = T2{2} ^ a[0]; + auto e1 = a[0] < w; + auto e2 = a[0] < T2{2}; + auto e3 = T2{2} < a[0]; + simd z{4}; + auto f1 = a[0] ^ z; + auto f2 = z ^ a[0]; + auto f3 = a[0] < z; + auto f4 = z < a[0]; +} + +void foo() SYCL_ESIMD_FUNCTION { + test_esimd_ops(simd(3), (int)1, (int)9); + test_esimd_ops(simd(3), (uint32_t)1, (uint32_t)9); + test_esimd_ops(simd(3), 1, 9); + test_esimd_ops(simd(3), 1, 9); +}