[ESIMD] Allow copy and move constructors for simd_view

This patch allows copy and move constructors and now simd_view
class adheres to the "rule of 5". We also mark APIs of simd class
with [[clang::lifetimebound]] attirbute to warn users when
simd_view escapes the lifetime of the underlying simd object.

Author: DenisBakhvalov

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

@@ -59,9 +59,9 @@ template <typename BaseTy, typename RegionTy> class simd_view_impl {
       : M_base(Base), M_region(Region) {}
 
 public:
-  // Disallow copy and move constructors.
-  simd_view_impl(const simd_view_impl &Other) = delete;
-  simd_view_impl(simd_view_impl &&Other) = delete;
+  // Default copy and move constructors.
+  simd_view_impl(const simd_view_impl &Other) = default;
+  simd_view_impl(simd_view_impl &&Other) = default;
   /// @}
 
   /// Conversion to simd type.
@@ -80,6 +80,11 @@ template <typename BaseTy, typename RegionTy> class simd_view_impl {
   simd_view_impl &operator=(const value_type &Val) { return write(Val); }
   /// @}
 
+  /// Move assignment operator.
+  simd_view_impl &operator=(simd_view_impl &&Other) {
+    return write(Other.read());
+  }
+
   /// @{
   /// Region accessors.
   static constexpr bool is1D() { return !ShapeTy::Is_2D; }

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

@@ -139,7 +139,7 @@ template <typename Ty, int N> class simd {
   }
 
   /// View this simd object in a different element type.
-  template <typename EltTy> auto bit_cast_view() & {
+  template <typename EltTy> auto bit_cast_view() &[[clang::lifetimebound]] {
     using TopRegionTy = detail::compute_format_type_t<simd, EltTy>;
     using RetTy = simd_view<simd, TopRegionTy>;
     TopRegionTy R(0);
@@ -153,7 +153,8 @@ template <typename Ty, int N> class simd {
   }
 
   /// View as a 2-dimensional simd_view.
-  template <typename EltTy, int Height, int Width> auto bit_cast_view() & {
+  template <typename EltTy, int Height, int Width>
+  auto bit_cast_view() &[[clang::lifetimebound]] {
     using TopRegionTy =
         detail::compute_format_type_2d_t<simd, EltTy, Height, Width>;
     using RetTy = simd_view<simd, TopRegionTy>;
@@ -174,7 +175,8 @@ template <typename Ty, int N> class simd {
   /// \param Offset is the starting element offset.
   /// \return the representing region object.
   template <int Size, int Stride>
-  simd_view<simd, region1d_t<Ty, Size, Stride>> select(uint16_t Offset = 0) & {
+  simd_view<simd, region1d_t<Ty, Size, Stride>>
+  select(uint16_t Offset = 0) &[[clang::lifetimebound]] {
     region1d_t<Ty, Size, Stride> Reg(Offset);
     return {*this, Reg};
   }
@@ -196,7 +198,8 @@ template <typename Ty, int N> class simd {
   Ty operator[](int i) const { return data()[i]; }
 
   /// Return writable view of a single element.
-  simd_view<simd, region1d_t<Ty, 1, 0>> operator[](int i) {
+  simd_view<simd, region1d_t<Ty, 1, 0>> operator[](int i)
+      [[clang::lifetimebound]] {
     return select<1, 0>(i);
   }
 

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

@@ -42,9 +42,9 @@ class simd_view : public detail::simd_view_impl<BaseTy, RegionTy> {
   simd_view(BaseTy &&Base, RegionTy Region) : BaseClass(Base, Region) {}
 
 public:
-  // Disallow copy and move constructors for simd_view.
-  simd_view(const simd_view &Other) = delete;
-  simd_view(simd_view &&Other) = delete;
+  // Default copy and move constructors for simd_view.
+  simd_view(const simd_view &Other) = default;
+  simd_view(simd_view &&Other) = default;
 
   /// @{
   /// Assignment operators.
@@ -58,6 +58,12 @@ class simd_view : public detail::simd_view_impl<BaseTy, RegionTy> {
   }
   /// @}
 
+  /// Move assignment operator.
+  simd_view &operator=(simd_view &&Other) {
+    *this = Other.read();
+    return *this;
+  }
+
 #define DEF_RELOP(RELOP)                                                       \
   ESIMD_INLINE friend simd<uint16_t, length> operator RELOP(                   \
       const simd_view &X, const value_type &Y) {                               \

sycl/test/esimd/simd_view.cpp

@@ -7,7 +7,7 @@
 
 using namespace sycl::ext::intel::experimental::esimd;
 
-bool test_simd_view_bin_ops() __attribute__((sycl_device)) {
+SYCL_ESIMD_FUNCTION bool test_simd_view_bin_ops() {
   simd<int, 16> v0 = 1;
   simd<int, 16> v1 = 2;
   auto ref0 = v0.select<8, 2>(0);
@@ -24,7 +24,7 @@ bool test_simd_view_bin_ops() __attribute__((sycl_device)) {
   return v0[0] == 1;
 }
 
-bool test_simd_view_unary_ops() __attribute__((sycl_device)) {
+SYCL_ESIMD_FUNCTION bool test_simd_view_unary_ops() {
   simd<int, 16> v0 = 1;
   simd<int, 16> v1 = 2;
   auto ref0 = v0.select<8, 2>(0);
@@ -36,21 +36,21 @@ bool test_simd_view_unary_ops() __attribute__((sycl_device)) {
   return v1[0] == 1;
 }
 
-bool test_simd_view_assign1() __attribute__((sycl_device)) {
+SYCL_ESIMD_FUNCTION bool test_simd_view_assign1() {
   simd<int, 32> v0(0, 1);
   simd<int, 16> v1(0, 1);
   v0.select<8, 1>(0) = v1.select<8, 1>(0) + v1.select<8, 1>(1);
   return v0[8] == 8 + 9;
 }
 
-bool test_simd_view_assign2() __attribute__((sycl_device)) {
+SYCL_ESIMD_FUNCTION bool test_simd_view_assign2() {
   simd<int, 32> v0 = 0;
   simd<int, 16> v1 = 1;
   v0.select<8, 1>(0) = v1.select<8, 1>(0);
   return v0[0] == 1;
 }
 
-bool test_simd_view_assign3() __attribute__((sycl_device)) {
+SYCL_ESIMD_FUNCTION bool test_simd_view_assign3() {
   simd<int, 64> v0 = 0;
   simd<int, 64> v1 = 1;
   auto mask = (v0.select<16, 1>(0) > v1.select<16, 1>(0));
@@ -62,3 +62,58 @@ bool test_simd_view_assign3() __attribute__((sycl_device)) {
       (g4.row(2) & mask2.bit_cast_view<ushort, 4, 16>().row(0));
   return val[0] == 0 && val1[0] == 0;
 }
+
+// copy constructor creates the same view of the underlying data.
+SYCL_ESIMD_FUNCTION void test_simd_view_copy_ctor() {
+  simd<int, 16> v0 = 1;
+  auto v0_view = v0.select<8, 1>(0);
+  auto v0_view_copy(v0_view);
+}
+
+// move constructor transfers the same view of the underlying data.
+SYCL_ESIMD_FUNCTION void test_simd_view_move_ctor() {
+  simd<int, 16> v0 = 1;
+  auto v0_view = v0.select<8, 1>(0);
+  auto v0_view_move(std::move(v0_view));
+}
+
+// assignment operator copies the underlying data.
+SYCL_ESIMD_FUNCTION void test_simd_view_copy_assign() {
+  simd<int, 16> v0 = 0;
+  simd<int, 16> v1 = 1;
+  auto v0_view = v0.select<8, 1>(0);
+  auto v1_view = v1.select<8, 1>(0);
+  v0_view = v1_view;
+}
+
+// move assignment operator copies the underlying data.
+SYCL_ESIMD_FUNCTION void test_simd_view_move_assign() {
+  simd<int, 16> v0 = 0;
+  simd<int, 16> v1 = 1;
+  auto v0_view = v0.select<8, 1>(0);
+  auto v1_view = v1.select<8, 1>(0);
+  v0_view = std::move(v1_view);
+}
+
+// Check that the same holds for specialization with lenght==1.
+SYCL_ESIMD_FUNCTION void test_simd_view_ctors_length1() {
+  simd<int, 16> v = 1;
+  auto v0_view = v[0];
+  auto v0_view_copy(v0_view);            // copy ctor
+  auto v0_view_move(std::move(v0_view)); // move ctor
+  simd<int, 16> vv = 2;
+  auto vv0_view = vv[0];
+  v0_view_copy = vv0_view;            // copy assign
+  v0_view_move = std::move(vv0_view); // move assign
+}
+
+// Check that simd_view can be passed as a by-value argument
+template <class BaseTy, class RegionTy>
+SYCL_ESIMD_FUNCTION void foo(simd_view<BaseTy, RegionTy> view) {}
+
+SYCL_ESIMD_FUNCTION void bar() {
+  simd<int, 16> v0 = 0;
+  auto v0_view = v0.select<8, 1>(0);
+  foo(v0_view);            // lvalue
+  foo(v0.select<8, 1>(0)); // rvalue
+}

sycl/test/esimd/simd_view_ret_warn.cpp

@@ -0,0 +1,28 @@
+// RUN: %clangxx -fsycl -fsycl-device-only -fsyntax-only -Xclang -verify %s
+
+#include <sycl/ext/intel/experimental/esimd.hpp>
+using namespace sycl::ext::intel::experimental::esimd;
+
+// Explicitly returning simd_view for local objects is wrong,
+// and it should be programmers fault, similar to string_view.
+// However, sometimes we could return simd_view from a function
+// implicitly. This test checks that users will see a warning in such situation.
+simd_view<simd<float, 4>, region1d_t<float, 1, 0>> f1(simd<float, 4> x) {
+  // expected-warning@+1 {{address of stack memory associated with parameter 'x' returned}}
+  return x[0];
+}
+
+simd_view<simd<float, 4>, region1d_t<float, 2, 1>> f2(simd<float, 4> x) {
+  // expected-warning@+1 {{address of stack memory associated with parameter 'x' returned}}
+  return x.select<2, 1>(0);
+}
+
+auto f3(simd<float, 4> x) {
+  // expected-warning@+1 {{address of stack memory associated with parameter 'x' returned}}
+  return x.bit_cast_view<int>();
+}
+
+auto f4(simd<float, 4> x) {
+  // expected-warning@+1 {{address of stack memory associated with parameter 'x' returned}}
+  return x.bit_cast_view<int, 2, 2>();
+}