[ESIMD] Allow constructing simd_view from simd

SYCL/ESIMD/api/simd_view_copy_move_assign.cpp

@@ -24,45 +24,56 @@
 using namespace cl::sycl;
 using namespace sycl::ext::intel::experimental::esimd;
 
-template <class T> bool test(queue q, std::string str, T funcUnderTest) {
-  std::cout << "Testing " << str << " ...\n";
-  constexpr unsigned VL = 8;
-
+template <unsigned VL, class T>
+bool test(queue q, std::string str, T funcUnderTest) {
+  std::cout << "Testing " << str << ", VL = " << VL << " ...\n";
   int A[VL];
   int B[VL];
-  // As a result, A should have first 4 values from B, next 4 values from A.
-  int gold[VL] = {0, 1, 2, 3, -4, -5, -6, -7};
+  constexpr unsigned HalfVL = VL > 1 ? (VL / 2) : 1;
 
+  // The expected result gets the first half of values from B,
+  // the rest of values from A.
+  int gold[VL];
   for (unsigned i = 0; i < VL; ++i) {
     A[i] = -i;
     B[i] = i;
+    gold[i] = (i < HalfVL) ? B[i] : A[i];
   }
 
   try {
     buffer<int, 1> bufA(A, range<1>(VL));
     buffer<int, 1> bufB(B, range<1>(VL));
     range<1> glob_range{1};
 
-    auto e = q.submit([&](handler &cgh) {
-      auto PA = bufA.template get_access<access::mode::read_write>(cgh);
-      auto PB = bufB.template get_access<access::mode::read_write>(cgh);
-      cgh.parallel_for(glob_range, [=](id<1> i) SYCL_ESIMD_KERNEL {
-        using namespace sycl::ext::intel::experimental::esimd;
-        unsigned int offset = i * VL * sizeof(int);
-        simd<int, VL> va;
-        va.copy_from(PA, offset);
-        simd<int, VL> vb;
-        vb.copy_from(PB, offset);
-        auto va_view = va.select<4, 1>(0);
-        auto vb_view = vb.select<4, 1>(0);
-
-        funcUnderTest(va_view, vb_view);
-
-        va.copy_to(PA, offset);
-        vb.copy_to(PB, offset);
-      });
-    });
-    q.wait_and_throw();
+    q.submit([&](handler &cgh) {
+       auto PA = bufA.template get_access<access::mode::read_write>(cgh);
+       auto PB = bufB.template get_access<access::mode::read_write>(cgh);
+       cgh.parallel_for(glob_range, [=](id<1> i) SYCL_ESIMD_KERNEL {
+         using namespace sycl::ext::intel::experimental::esimd;
+         unsigned int offset = i * VL * sizeof(int);
+         simd<int, VL> va;
+         simd<int, VL> vb;
+         if constexpr (VL == 1) {
+           va[0] = PA[0];
+           vb[0] = PB[0];
+         } else {
+           va.copy_from(PA, offset);
+           vb.copy_from(PB, offset);
+         }
+
+         auto va_view = va.template select<HalfVL, 1>(0);
+         auto vb_view = vb.template select<HalfVL, 1>(0);
+         funcUnderTest(va_view, vb_view);
+
+         if constexpr (VL == 1) {
+           PA[0] = va[0];
+           PB[0] = vb[0];
+         } else {
+           va.copy_to(PA, offset);
+           vb.copy_to(PB, offset);
+         }
+       });
+     }).wait_and_throw();
   } catch (sycl::exception const &e) {
     std::cout << "SYCL exception caught: " << e.what() << '\n';
     return false; // not success
@@ -86,34 +97,67 @@ template <class T> bool test(queue q, std::string str, T funcUnderTest) {
   return err_cnt > 0 ? false : true;
 }
 
+template <class FuncT>
+bool tests(queue &Q, std::string Str, FuncT FuncUnderTest) {
+  bool Passed = true;
+  Passed &= test<8>(Q, Str, FuncUnderTest);
+  Passed &= test<1>(Q, Str, FuncUnderTest);
+  return Passed;
+}
+
 int main(void) {
   queue q(esimd_test::ESIMDSelector{}, esimd_test::createExceptionHandler());
   auto dev = q.get_device();
   std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
 
   bool passed = true;
   // copy constructor creates the same view of the underlying data.
-  passed &= test(q, "copy constructor",
-                 [](auto &va_view, auto &vb_view) SYCL_ESIMD_FUNCTION {
-                   auto va_view_copy(va_view);
-                   auto vb_view_copy(vb_view);
-                   va_view_copy = vb_view_copy;
-                 });
+  passed &= tests(q, "copy constructor",
+                  [](auto &va_view, auto &vb_view) SYCL_ESIMD_FUNCTION {
+                    auto va_view_copy(va_view);
+                    auto vb_view_copy(vb_view);
+                    va_view_copy = vb_view_copy;
+                  });
   // move constructor transfers the same view of the underlying data.
-  passed &= test(q, "move constructor",
-                 [](auto &va_view, auto &vb_view) SYCL_ESIMD_FUNCTION {
-                   auto va_view_move(std::move(va_view));
-                   auto vb_view_move(std::move(vb_view));
-                   va_view_move = vb_view_move;
-                 });
+  passed &= tests(q, "move constructor",
+                  [](auto &va_view, auto &vb_view) SYCL_ESIMD_FUNCTION {
+                    auto va_view_move(std::move(va_view));
+                    auto vb_view_move(std::move(vb_view));
+                    va_view_move = vb_view_move;
+                  });
   // assignment operator copies the underlying data.
-  passed &= test(q, "assignment operator",
-                 [](auto &va_view, auto &vb_view)
-                     SYCL_ESIMD_FUNCTION { va_view = vb_view; });
+  passed &= tests(q, "assignment operator",
+                  [](auto &va_view, auto &vb_view)
+                      SYCL_ESIMD_FUNCTION { va_view = vb_view; });
   // move assignment operator copies the underlying data.
-  passed &= test(q, "move assignment operator",
-                 [](auto &va_view, auto &vb_view)
-                     SYCL_ESIMD_FUNCTION { va_view = std::move(vb_view); });
-
+  passed &= tests(q, "move assignment operator",
+                  [](auto &va_view, auto &vb_view)
+                      SYCL_ESIMD_FUNCTION { va_view = std::move(vb_view); });
+
+  // construct complete view of a vector.
+#if 0
+  // TODO: When/if template arguments deducing is implemented for simd
+  // constructor accepting simd_view are implemented, the following
+  // shorter and more elegant version of the code may be used.
+  passed &= tests(q, "constructor from simd",
+                  [](auto &va_view, auto &vb_view) SYCL_ESIMD_FUNCTION {
+                    simd vb = vb_view;
+                    simd_view new_vb_view = vb; // ctor from simd
+                    va_view = new_vb_view;
+                  });
+#else
+  passed &= test<8>(q, "constructor from simd",
+                    [](auto &va_view, auto &vb_view) SYCL_ESIMD_FUNCTION {
+                      simd<int, 4> vb = vb_view;
+                      simd_view new_vb_view = vb; // ctor from simd
+                      va_view = new_vb_view;
+                    });
+  passed &= test<1>(q, "constructor from simd",
+                    [](auto &va_view, auto &vb_view) SYCL_ESIMD_FUNCTION {
+                      simd<int, 1> vb = vb_view;
+                      simd_view new_vb_view = vb; // ctor from simd
+                      va_view = new_vb_view;
+                    });
+#endif
   return passed ? 0 : 1;
 }