Skip to content

[ESIMD] Add writeable subsript operator for simd_view #4384

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 6 commits into from
Aug 25, 2021
Merged

[ESIMD] Add writeable subsript operator for simd_view #4384

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
50d7285
[ESIMD] Add writeable subsript operator for simd_view
DenisBakhvalov Aug 13, 2021
1c420a5
clang-format
DenisBakhvalov Aug 20, 2021
0a8beb4
Fixed bitwise and relational ops
DenisBakhvalov Aug 23, 2021
8d9664c
Added new test case for nested view of len1
DenisBakhvalov Aug 23, 2021
2755b82
Clang-format
DenisBakhvalov Aug 23, 2021
0977675
renamed region_base_1 to region1d_scalar_t
DenisBakhvalov Aug 25, 2021
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
3 changes: 2 additions & 1 deletion sycl/include/sycl/ext/intel/experimental/esimd/detail/region.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -57,7 +57,8 @@ template <typename T, int SizeY, int StrideY, int SizeX, int StrideX>
using region2d_t = region_base<true, T, SizeY, StrideY, SizeX, StrideX>;

// A region with a single element.
template <typename T> using region_base_1 = region_base<false, T, 1, 0, 1, 0>;
template <typename T, int StrideY, int StrideX>
using region1d_scalar_t = region_base<false, T, 1, StrideY, 1, StrideX>;

// simd_view forward declaration.
template <typename BaseTy, typename RegionTy> class simd_view;
Expand Down
25 changes: 25 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 @@ -213,6 +213,11 @@ class simd_view_impl {
return ComputeTy(V2); \
} \
template <class T1 = Derived, class = std::enable_if_t<T1::length != 1>> \
ESIMD_INLINE friend auto operator BINOP(const Derived &X, \
const element_type &Y) { \
return X BINOP(value_type) Y; \
} \
template <class T1 = Derived, class = std::enable_if_t<T1::length != 1>> \
ESIMD_INLINE friend auto operator BINOP(const value_type &X, \
const Derived &Y) { \
using ComputeTy = detail::compute_type_t<value_type>; \
Expand Down Expand Up @@ -256,6 +261,11 @@ class simd_view_impl {
return simd<element_type, length>(V2); \
} \
template <class T1 = Derived, class = std::enable_if_t<T1::length != 1>> \
ESIMD_INLINE friend auto operator BITWISE_OP(const Derived &X, \
const element_type &Y) { \
return X BITWISE_OP(value_type) Y; \
} \
template <class T1 = Derived, class = std::enable_if_t<T1::length != 1>> \
ESIMD_INLINE friend auto operator BITWISE_OP(const value_type &X, \
const Derived &Y) { \
static_assert(std::is_integral<element_type>(), "not integral type"); \
Expand Down Expand Up @@ -352,6 +362,21 @@ class simd_view_impl {
return v[i];
}

/// Return a writeable view of a single element.
template <typename T = Derived,
typename = sycl::detail::enable_if_t<T::is1D()>>
auto operator[](int i) {
return select<1, 0>(i);
}

/// Return a writeable view of a single element.
template <typename T = Derived,
typename = sycl::detail::enable_if_t<T::is1D()>>
__SYCL_DEPRECATED("use operator[] form.")
auto operator()(int i) {
return select<1, 0>(i);
}

/// \name Replicate
/// Replicate simd instance given a simd_view
/// @{
Expand Down
100 changes: 92 additions & 8 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 @@ -36,8 +36,10 @@ class simd_view : public detail::simd_view_impl<BaseTy, RegionTy,
using ShapeTy = typename shape_type<RegionTy>::type;
static constexpr int length = ShapeTy::Size_x * ShapeTy::Size_y;

using element_type = typename ShapeTy::element_type;

/// The simd type if reading this simd_view object.
using value_type = simd<typename ShapeTy::element_type, length>;
using value_type = simd<element_type, length>;

private:
simd_view(BaseTy &Base, RegionTy Region) : BaseClass(Base, Region) {}
Expand Down Expand Up @@ -79,6 +81,10 @@ class simd_view : public detail::simd_view_impl<BaseTy, RegionTy,
mask_type_t<length> M(1); \
return M & detail::convert<mask_type_t<length>>(R); \
} \
ESIMD_INLINE friend simd<uint16_t, length> operator RELOP( \
const simd_view &X, const element_type &Y) { \
return X RELOP(value_type) Y; \
} \
ESIMD_INLINE friend simd<uint16_t, length> operator RELOP( \
const simd_view &X, const simd_view &Y) { \
return (X RELOP Y.read()); \
Expand All @@ -95,6 +101,9 @@ class simd_view : public detail::simd_view_impl<BaseTy, RegionTy,
};

/// This is a specialization of simd_view class with a single element.
/// Objects of such a class are created in the following situation:
/// simd<int, 4> v = 1;
/// auto v1 = v[0];
/// We allow implicit conversion to underlying type, e.g.:
/// simd<int, 4> v = 1;
/// int i = v[0];
Expand All @@ -103,38 +112,113 @@ class simd_view : public detail::simd_view_impl<BaseTy, RegionTy,
/// bool b = v[0] > v[1] && v[2] < 42;
///
/// \ingroup sycl_esimd
template <typename BaseTy>
class simd_view<BaseTy, region_base_1<typename BaseTy::element_type>>
template <typename BaseTy, typename T, int StrideY, int StrideX>
class simd_view<BaseTy, region1d_scalar_t<T, StrideY, StrideX>>
: public detail::simd_view_impl<
BaseTy, region_base_1<typename BaseTy::element_type>,
simd_view<BaseTy, region_base_1<typename BaseTy::element_type>>> {
BaseTy, region1d_scalar_t<T, StrideY, StrideX>,
simd_view<BaseTy, region1d_scalar_t<T, StrideY, StrideX>>> {
template <typename, int> friend class simd;
template <typename, typename, typename> friend class detail::simd_view_impl;

public:
using RegionTy = region_base_1<typename BaseTy::element_type>;
using RegionTy = region1d_scalar_t<T, StrideY, StrideX>;
using BaseClass =
detail::simd_view_impl<BaseTy, RegionTy, simd_view<BaseTy, RegionTy>>;
using ShapeTy = typename shape_type<RegionTy>::type;
static constexpr int length = ShapeTy::Size_x * ShapeTy::Size_y;
static_assert(1 == length, "length of this view is not equal to 1");
/// The element type of this class, which could be different from the element
/// type of the base object type.
using element_type = typename ShapeTy::element_type;
using element_type = T;

private:
simd_view(BaseTy &Base, RegionTy Region) : BaseClass(Base, Region) {}
simd_view(BaseTy &&Base, RegionTy Region) : BaseClass(Base, Region) {}

public:
operator element_type() const { return (*this)[0]; }
operator element_type() const {
const auto v = BaseClass::read();
return v[0];
}

using BaseClass::operator=;

#define DEF_RELOP(RELOP) \
ESIMD_INLINE friend bool operator RELOP(const simd_view &X, \
const simd_view &Y) { \
return (element_type)X RELOP(element_type) Y; \
} \
template <typename T1, typename = sycl::detail::enable_if_t< \
detail::is_esimd_scalar<T1>::value && \
detail::is_vectorizable_v<T1>::value>> \
ESIMD_INLINE friend bool operator RELOP(const simd_view &X, T1 Y) { \
return (element_type)X RELOP Y; \
}

DEF_RELOP(>)
DEF_RELOP(>=)
DEF_RELOP(<)
DEF_RELOP(<=)
DEF_RELOP(==)
DEF_RELOP(!=)

#undef DEF_RELOP
};

// TODO: remove code duplication in two class specializations for a simd_view
// with a single element

/// This is a specialization of nested simd_view class with a single element.
/// Objects of such a class are created in the following situation:
/// simd<int, 4> v = 1;
/// auto v1 = v.select<2, 1>(0);
/// auto v2 = v1[0]; // simd_view of a nested region for a single element
template <typename BaseTy, typename T, int StrideY, int StrideX,
typename NestedRegion>
class simd_view<BaseTy,
std::pair<region1d_scalar_t<T, StrideY, StrideX>, NestedRegion>>
: public detail::simd_view_impl<
BaseTy,
std::pair<region1d_scalar_t<T, StrideY, StrideX>, NestedRegion>,
simd_view<BaseTy, std::pair<region1d_scalar_t<T, StrideY, StrideX>,
NestedRegion>>> {
template <typename, int> friend class simd;
template <typename, typename, typename> friend class detail::simd_view_impl;

public:
using RegionTy =
std::pair<region1d_scalar_t<T, StrideY, StrideX>, NestedRegion>;
using BaseClass =
detail::simd_view_impl<BaseTy, RegionTy, simd_view<BaseTy, RegionTy>>;
using ShapeTy = typename shape_type<RegionTy>::type;
static constexpr int length = ShapeTy::Size_x * ShapeTy::Size_y;
static_assert(1 == length, "length of this view is not equal to 1");
/// The element type of this class, which could be different from the element
/// type of the base object type.
using element_type = T;

private:
simd_view(BaseTy &Base, RegionTy Region) : BaseClass(Base, Region) {}
simd_view(BaseTy &&Base, RegionTy Region) : BaseClass(Base, Region) {}

public:
operator element_type() const {
const auto v = BaseClass::read();
return v[0];
}

using BaseClass::operator=;

#define DEF_RELOP(RELOP) \
ESIMD_INLINE friend bool operator RELOP(const simd_view &X, \
const simd_view &Y) { \
return (element_type)X RELOP(element_type) Y; \
} \
template <typename T1, typename = sycl::detail::enable_if_t< \
detail::is_esimd_scalar<T1>::value && \
detail::is_vectorizable_v<T1>::value>> \
ESIMD_INLINE friend bool operator RELOP(const simd_view &X, T1 Y) { \
return (element_type)X RELOP Y; \
}

DEF_RELOP(>)
Expand Down
2 changes: 1 addition & 1 deletion sycl/test/esimd/esimd-util-compiler-eval.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -19,7 +19,7 @@ static_assert(log2<1024 * 1024>() == 20, "");

using BaseTy = simd<float, 4>;
using RegionTy = region1d_t<float, 2, 1>;
using RegionTy1 = region_base_1<float>;
using RegionTy1 = region1d_scalar_t<float, 0, 0>;
static_assert(
!is_simd_view_v<
simd_view_impl<BaseTy, RegionTy, simd_view<BaseTy, RegionTy>>>::value,
Expand Down
47 changes: 46 additions & 1 deletion sycl/test/esimd/simd_view.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -142,10 +142,55 @@ void test_simd_view_impl_api_ret_types() SYCL_ESIMD_FUNCTION {

void test_simd_view_subscript() SYCL_ESIMD_FUNCTION {
simd<int, 4> v = 1;
auto vv = v.select<2, 1>(0);
const auto vv = v.select<2, 1>(0);

int x = vv[1];
// expected-warning@+2 2 {{deprecated}}
// expected-note@sycl/ext/intel/experimental/esimd/detail/simd_view_impl.hpp:* 2 {{has been explicitly marked deprecated here}}
int y = vv(1);
}

void test_simd_view_writeable_subscript() SYCL_ESIMD_FUNCTION {
simd<int, 4> v = 1;
auto vv1 = v.select<2, 1>(0);
auto vv2 = v.select<2, 1>(0);
auto x = vv1 == vv2; // test relational operations
vv1[1] = 0; // returns writeable simd_view
int y = vv1[1]; // nested simd_view -> int

// expected-warning@+2 2 {{deprecated}}
// expected-note@sycl/ext/intel/experimental/esimd/detail/simd_view_impl.hpp:* 2 {{has been explicitly marked deprecated here}}
vv1(1) = 1;
}

// In this test `g.row(1)` return simd_view and `(g.row(1))[0]` returns
// simd_view of size 1. To avoid two implicit conversions (which is not
// allowed), simd_view_impl needs explicit definition of BINOP with a scalar.
void test_simd_view_binop_with_conv_to_scalar() SYCL_ESIMD_FUNCTION {
simd<ushort, 64> s = 0;
auto g = s.bit_cast_view<ushort, 4, 16>();
auto x = g.row(1) - (g.row(1))[0]; // binary op
auto y = g.row(1) & (g.row(1))[0]; // bitwise op
auto z = g.row(1) < (g.row(1))[0]; // relational op
}

// This code is OK. The result of bit_cast_view should be mapped
// to specialization of simd_view with length 1, so conversion
// to scalar is allowed.
void test_simd_view_len1_bitcast() SYCL_ESIMD_FUNCTION {
simd<int, 1> s = 0;
float f = s.bit_cast_view<float>();
}

// This test checks the same thing as previous one but for the
// nested simd_view specialization with length 1.
void test_nested_simd_view_len1_bitcast() SYCL_ESIMD_FUNCTION {
simd<int, 2> s = 0;
auto v = s.bit_cast_view<float>(); // generic simd_view
float f = v[0]; // result of v[0] is a specialized nested simd_view
// with length 1, which then converts to a scalar.

// checking nested views with several bitcasts.
simd<double, 4> s2;
((s2[0].bit_cast_view<float>())[1].bit_cast_view<int>())[0] = 1;
}