Skip to content

[ESIMD] Fix 'ambiguous operator' error with length 1 simd operands #4149

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Jul 21, 2021
Merged

[ESIMD] Fix 'ambiguous operator' error with length 1 simd operands #4149

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
8 changes: 8 additions & 0 deletions sycl/include/sycl/ext/intel/experimental/esimd/detail/simd_view_impl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -193,6 +193,8 @@ template <typename BaseTy, typename RegionTy> class simd_view_impl {
}

#define DEF_BINOP(BINOP, OPASSIGN) \
template <class T1 = simd_view_impl, \
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

As I understood a[0] + x was ambiguous when the underlying types were different, right? E.g., int and uint32_t.
You resolve it by disabling one of the candidates and rely on implicit conversion to underlying type and built-in operator+ for primitive types.

Can we sink those two BINOP overloads into the derived class of generic simd_view implementation? That way I think we don't need this SFINAE trick.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

As I understood a[0] + x was ambiguous when the underlying types were different, right? E.g., int and uint32_t.

not really. int - int causes the failure too. In the test I just modelled the user test case where the types were different.

You resolve it by disabling one of the candidates and rely on implicit conversion to underlying type and built-in operator+ for primitive types.

yes, with length == 1, simd_view_impl and value_type are implicitly convertible to the element type, thus compiler will always be able to resolve BINOP(const simd_view_impl &, const value_type &) to the built-in operation over primitive types after implicitly converting both sides.

Can we sink those two BINOP overloads into the derived class of generic simd_view implementation?

not sure I understand. Do you mean moving BINOP from simd_view_impl to simd_view? That won't affect resolution in any way, since simd_view already binds to simd_view_impl w/o any conversions as simd_view extends simd_view_impl.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Per discussion with @DenisBakhvalov: his suggestion was to move BINOP to simd_view to make them unavailable for the length=1 simd_view specialization. That might make sense, but (1) would be pretty intrusive change at this point (2) my upcoming simd_mask introduction patch refactors operators anyway (none will be friends), so this can be re-visited when reviewing that patch. So we decided to go with the current fix at this point.

class = std::enable_if_t<T1::length != 1>> \
ESIMD_INLINE friend auto operator BINOP(const simd_view_impl &X, \
const value_type &Y) { \
using ComputeTy = detail::compute_type_t<value_type>; \
Expand All @@ -202,6 +204,8 @@ template <typename BaseTy, typename RegionTy> class simd_view_impl {
auto V2 = V0 BINOP V1; \
return ComputeTy(V2); \
} \
template <class T1 = simd_view_impl, \
class = std::enable_if_t<T1::length != 1>> \
ESIMD_INLINE friend auto operator BINOP(const value_type &X, \
const simd_view_impl &Y) { \
using ComputeTy = detail::compute_type_t<value_type>; \
Expand Down Expand Up @@ -237,12 +241,16 @@ template <typename BaseTy, typename RegionTy> class simd_view_impl {
#undef DEF_BINOP

#define DEF_BITWISE_OP(BITWISE_OP, OPASSIGN) \
template <class T1 = simd_view_impl, \
class = std::enable_if_t<T1::length != 1>> \
ESIMD_INLINE friend auto operator BITWISE_OP(const simd_view_impl &X, \
const value_type &Y) { \
static_assert(std::is_integral<element_type>(), "not integral type"); \
auto V2 = X.read().data() BITWISE_OP Y.data(); \
return simd<element_type, length>(V2); \
} \
template <class T1 = simd_view_impl, \
class = std::enable_if_t<T1::length != 1>> \
ESIMD_INLINE friend auto operator BITWISE_OP(const value_type &X, \
const simd_view_impl &Y) { \
static_assert(std::is_integral<element_type>(), "not integral type"); \
Expand Down
18 changes: 10 additions & 8 deletions sycl/include/sycl/ext/intel/experimental/esimd/simd.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -233,10 +233,11 @@ template <typename Ty, int N> class simd {
auto V2 = V0 BINOP V1; \
return ComputeTy(V2); \
} \
template <typename T = simd, \
typename = sycl::detail::enable_if_t<T::length == 1>> \
ESIMD_INLINE friend auto operator BINOP(const simd &X, const Ty &Y) { \
return X BINOP simd(Y); \
template <typename T1, typename T = simd, \
typename = sycl::detail::enable_if_t<T::length == 1 && \
std::is_arithmetic_v<T1>>> \
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<simd>; \
Expand Down Expand Up @@ -272,10 +273,11 @@ template <typename Ty, int N> class simd {
mask_type_t<N> M(1); \
return M & detail::convert<mask_type_t<N>>(R); \
} \
template <typename T = simd, \
typename = sycl::detail::enable_if_t<T::length == 1>> \
ESIMD_INLINE friend bool operator RELOP(const simd &X, const Ty &Y) { \
return (Ty)X RELOP Y; \
template <typename T1, typename T = simd, \
typename = sycl::detail::enable_if_t<(T::length == 1) && \
std::is_arithmetic_v<T1>>> \
ESIMD_INLINE friend bool operator RELOP(const simd &X, T1 Y) { \
return (Ty)X RELOP(Ty) Y; \
}

DEF_RELOP(>)
Expand Down
4 changes: 4 additions & 0 deletions sycl/include/sycl/ext/intel/experimental/esimd/simd_view.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -82,6 +82,8 @@ class simd_view : public detail::simd_view_impl<BaseTy, RegionTy> {
DEF_RELOP(<=)
DEF_RELOP(==)
DEF_RELOP(!=)

#undef DEF_RELOP
};

/// This is a specialization of simd_view class with a single element.
Expand Down Expand Up @@ -132,6 +134,8 @@ class simd_view<BaseTy, region_base_1<Is2D, T, StrideY, StrideX>>
DEF_RELOP(<=)
DEF_RELOP(==)
DEF_RELOP(!=)

#undef DEF_RELOP
};

} // namespace esimd
Expand Down
36 changes: 36 additions & 0 deletions sycl/test/esimd/regression/len1_operators.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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 <sycl/ext/intel/experimental/esimd.hpp>

#include <cstdint>

using namespace sycl::ext::intel::experimental::esimd;

template <typename T1, typename T2>
void test_esimd_ops(simd<T1, 1> 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<T1, 1> 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<uint32_t, 1>(3), (int)1, (int)9);
test_esimd_ops(simd<int, 1>(3), (uint32_t)1, (uint32_t)9);
test_esimd_ops(simd<uint16_t, 1>(3), 1, 9);
test_esimd_ops(simd<int16_t, 1>(3), 1, 9);
}