[SYCL][ESIMD] Add simd constructor from simd_view (#15174)

intel · Aug 28, 2024 · 1f5e1dc · 1f5e1dc
1 parent 3ff3a28
commit 1f5e1dc
Show file tree

Hide file tree

Showing 2 changed files with 29 additions and 11 deletions.
diff --git a/sycl/include/sycl/ext/intel/esimd/simd.hpp b/sycl/include/sycl/ext/intel/esimd/simd.hpp
@@ -77,14 +77,19 @@ class simd : public detail::simd_obj_impl<
   }
 
   // Implicit conversion constructor from sycl::ext::oneapi::experimental::simd
-  template <
-      int N1 = N, class Ty1 = Ty,
-      class SFINAE = std::enable_if_t<
-          (N1 == N) && (N1 <= std::experimental::simd_abi::max_fixed_size<
-                                  Ty>)&&!detail::is_wrapper_elem_type_v<Ty1>>>
+  template <int N1 = N, class Ty1 = Ty,
+            class SFINAE = std::enable_if_t<
+                (N1 == N) &&
+                (N1 <= std::experimental::simd_abi::max_fixed_size<Ty>) &&
+                !detail::is_wrapper_elem_type_v<Ty1>>>
   simd(const sycl::ext::oneapi::experimental::simd<Ty, N1> &v)
       : simd(static_cast<raw_vector_type>(v)) {}
 
+  // Implicit conversion constructor from 1D simd_view
+  template <typename BaseTy, int Stride>
+  simd(simd_view<BaseTy, region_base<false, Ty, 1, 1, N, Stride>> &v)
+      : simd(v.read()) {}
+
   /// Broadcast constructor with conversion. Converts given value to
   /// #element_type and replicates it in all elements.
   /// Available when \c T1 is a valid simd element type.
@@ -113,11 +118,11 @@ class simd : public detail::simd_obj_impl<
   /// object. Available when the number of elements does not exceed maximum
   /// fixed size of the oneapi's simd_abi and (TODO, temporary limitation) the
   /// element type is a primitive type (e.g. can't be sycl::half).
-  template <
-      int N1, class Ty1 = Ty,
-      class SFINAE = std::enable_if_t<
-          (N1 == N) && (N1 <= std::experimental::simd_abi::max_fixed_size<
-                                  Ty>)&&!detail::is_wrapper_elem_type_v<Ty1>>>
+  template <int N1, class Ty1 = Ty,
+            class SFINAE = std::enable_if_t<
+                (N1 == N) &&
+                (N1 <= std::experimental::simd_abi::max_fixed_size<Ty>) &&
+                !detail::is_wrapper_elem_type_v<Ty1>>>
   operator sycl::ext::oneapi::experimental::simd<Ty, N1>() {
     return sycl::ext::oneapi::experimental::simd<Ty, N1>(base_type::data());
   }
@@ -210,7 +215,8 @@ template <int N> using simd_mask = detail::simd_mask_type<N>;
 template <typename Ty, int N>
 std::ostream &operator<<(std::ostream &OS, const __ESIMD_NS::simd<Ty, N> &V)
 #ifdef __SYCL_DEVICE_ONLY__
-    {}
+{
+}
 #else
 {
   __ESIMD_UNSUPPORTED_ON_HOST;

diff --git a/sycl/test/esimd/simd.cpp b/sycl/test/esimd/simd.cpp
@@ -376,3 +376,15 @@ void test_simd_binop_honor_int_promo() SYCL_ESIMD_FUNCTION {
   static_assert(std::is_same<decltype(c + c), simd<int, 32>>::value, "");
   static_assert(std::is_same<decltype(d + d), simd<int, 32>>::value, "");
 }
+
+template <class T> bool test_simd_view_ctors() SYCL_ESIMD_FUNCTION {
+  simd<T, 8> vec8;
+  auto view4 = vec8.template select<4, 1>(0);
+
+  simd<T, 4> vec4 = view4;
+  simd vec41 = view4;
+  return true;
+}
+
+template bool test_simd_view_ctors<int>() SYCL_ESIMD_FUNCTION;
+template bool test_simd_view_ctors<sycl::half>() SYCL_ESIMD_FUNCTION;