[ESIMD] Add writeable subsript operator for simd_view (#4384)

The primary goal for this patch is to allow writing through
simd_view subscript operator, e.g.:
```
  simd<int, 4> v = 1;
  auto vv = v.select<2, 1>(0);
  vv[0] = 42;
```

All the other changes are required to make that happen:

1. One more specialization for nested simd_view class with len=1.

2. Added explicit definition of BINOP for simd_view_impl with
a scalar value because of the following code:
```
  simd<ushort, 64> s = 0;
  auto g = s.bit_cast_view<ushort, 4, 16>();
  auto x = g.row(1) - (g.row(1))[0];
```
Here `g.row(1)` return simd_view and `(g.row(1))[0]` returns
simd_view of size 1. To use any binops defined in simd_view_impl
before this patch we would need to implicitly convert RHS to
scalar and then convert scalar to vector type. Two implicit
conversions are not allowed according to C++ language rules,
hence new BINOP operator definition for simd_view_impl.

3. I also added additional template parameter for simd_view
specialization for lenght=1 to support the following code:
```
  simd<int, 1> s = 0;
  float f = s.bit_cast_view<float>();
```
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. To enable such mapping, I fixed the definition
of region_base_1. Previously we had zeros for StrideY and StrideX
but in fact, we don't care about their values, they can take any
values.

Author: DenisBakhvalov

sycl/include/sycl/ext/intel/experimental/esimd/detail/region.hpp

@@ -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;

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

@@ -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>;                      \
@@ -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");      \
@@ -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
   /// @{

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

@@ -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) {}
@@ -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());                                                 \
@@ -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];
@@ -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(>)

sycl/test/esimd/esimd-util-compiler-eval.cpp

@@ -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,

sycl/test/esimd/simd_view.cpp

@@ -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;
+}