[ESIMD] Add esimd::merge free function. (#5308)

* [ESIMD] Add esimd::merge free function.

Also, remove deprecation note from the 'simd_mask_impl(const simd<T, N> &Val)'
constructor, as it is used when loading a mask from memory, and is otherwise
useful constructor.

Signed-off-by: Konstantin S Bobrovsky <[email protected]>

Author: kbobrovs

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

@@ -12,6 +12,7 @@
 
 /// \defgroup sycl_esimd DPC++ Explicit SIMD API
 
+#include <sycl/ext/intel/experimental/esimd/alt_ui.hpp>
 #include <sycl/ext/intel/experimental/esimd/common.hpp>
 #include <sycl/ext/intel/experimental/esimd/math.hpp>
 #include <sycl/ext/intel/experimental/esimd/memory.hpp>

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

@@ -0,0 +1,61 @@
+//==-------------- alt_ui.hpp - DPC++ Explicit SIMD API   ------------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+// "Alternative" convenience Explicit SIMD APIs.
+//===----------------------------------------------------------------------===//
+
+#include <sycl/ext/intel/experimental/esimd/simd.hpp>
+#include <sycl/ext/intel/experimental/esimd/simd_view.hpp>
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace ext {
+namespace intel {
+namespace experimental {
+namespace esimd {
+
+/// "Merges" elements of the input vectors according to the merge mask.
+/// @param a the first vector
+/// @param b the second vector
+/// @param m the merge mask
+/// @return a vector, where each element equals to corresponding element from
+///    \c a (if corresponding mask element is zero) or \c b (otherwise)
+/// \ingroup sycl_esimd
+template <class T, int N>
+__ESIMD_API simd<T, N> merge(simd<T, N> a, simd<T, N> b, simd_mask<N> m) {
+  a.merge(b, m);
+  return a;
+}
+
+/// "Merges" elements of the input masks according to the merge mask.
+template <int N>
+__ESIMD_API simd_mask<N> merge(simd_mask<N> a, simd_mask<N> b, simd_mask<N> m) {
+  a.merge(b, m);
+  return a;
+}
+
+/// "Merges" elements of vectors referenced by the input views.
+/// Available only when all of the following holds:
+/// - the length and the element type of the sub-regions referenced by both
+///   input views are the same.
+template <class BaseT1, class BaseT2, class RegionT1, class RegionT2,
+          class = std::enable_if_t<
+              (shape_type<RegionT1>::length == shape_type<RegionT2>::length) &&
+              std::is_same_v<detail::element_type_t<BaseT1>,
+                             detail::element_type_t<BaseT2>>>>
+__ESIMD_API auto merge(simd_view<BaseT1, RegionT1> v1,
+                       simd_view<BaseT2, RegionT2> v2,
+                       simd_mask<shape_type<RegionT1>::length> m) {
+  return merge(v1.read(), v2.read(), m);
+}
+
+} // namespace esimd
+} // namespace experimental
+} // namespace intel
+} // namespace ext
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)

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

@@ -62,7 +62,6 @@ class simd_mask_impl
   }
 
   /// Implicit conversion from simd.
-  __SYCL_DEPRECATED(__ESIMD_MASK_DEPRECATION_MSG)
   simd_mask_impl(const simd<T, N> &Val) : base_type(Val.data()) {}
 
   /// Implicit conversion from simd_view<simd,...>.

sycl/test/esimd/esimd_merge.cpp

@@ -0,0 +1,50 @@
+// RUN: %clangxx -fsycl -fsycl-device-only -fsyntax-only -Xclang -verify %s
+// expected-no-diagnostics
+
+#include <sycl/ext/intel/experimental/esimd.hpp>
+
+using namespace sycl::ext::intel::experimental::esimd;
+using namespace sycl::ext::intel::experimental;
+using namespace cl::sycl;
+
+using simd_mask_elem_t = simd_mask<1>::element_type;
+
+template <class T, int N> struct SimdMergeTest {
+  inline void test(T *in_a, T *in_b, simd_mask_elem_t *in_mask, T *out)
+      __attribute__((sycl_device)) {
+    simd<T, N> a(in_a);
+    simd<T, N> b(in_b);
+    simd_mask<N> m(in_mask);
+    simd<T, N> c = esimd::merge(a, b, m);
+    c.copy_to(out);
+  }
+};
+
+template <class T, int N> struct SimdViewMergeTest {
+  inline void test(T *in_a, T *in_b, simd_mask_elem_t *in_mask, T *out)
+      __attribute__((sycl_device)) {
+    simd<T, N> a(in_a);
+    simd<T, N> b(in_b);
+    simd_mask<N / 2> m(in_mask);
+    simd<T, N / 2> c = esimd::merge(a.template select<N / 2, 1>(1),
+                                    b.template select<N / 2, 2>(0), m);
+    c.copy_to(out);
+  }
+};
+
+template <class T, int N> struct SimdView2MergeTest {
+  inline void test(T *in_a, T *in_b, simd_mask_elem_t *in_mask, T *out)
+      __attribute__((sycl_device)) {
+    simd<T, N> a(in_a);
+    simd<T, N> b(in_b);
+    simd_mask<N / 4> m(in_mask);
+    simd<T, N / 4> c = esimd::merge(
+        a.template select<N / 2, 1>(1).template select<N / 4, 1>(),
+        b.template select<N / 2, 2>(0).template select<N / 4, 1>(), m);
+    c.copy_to(out);
+  }
+};
+
+template struct SimdMergeTest<float, 8>;
+template struct SimdViewMergeTest<float, 8>;
+template struct SimdView2MergeTest<sycl::half, 8>;

sycl/test/esimd/simd_mask.cpp

@@ -63,16 +63,12 @@ SYCL_EXTERNAL SYCL_ESIMD_FUNCTION void compat_test(float *ptr) {
   simd<int, 8> pred1(1);
   auto pred2 = pred1.bit_cast_view<unsigned short>();
 
-  // expected-warning@+1 {{deprecated}}
   auto x1 = gather<float, 16>(ptr, offsets, pred);
   // expected-warning@+1 {{deprecated}}
   auto x11 = gather<float, 16>(ptr, offsets, pred2);
-  // expected-warning@+1 {{deprecated}}
   auto x2 = gather<float, 16>(ptr, offsets, simd<unsigned short, 16>{});
   simd_mask<16> m1(0);
-  // expected-warning@+1 {{deprecated}}
   m1 = pred;
   simd_mask<16> m2(0);
-  // expected-warning@+1 {{deprecated}}
   m2 = std::move(pred);
 }