[ESIMD] Improve simd::replicate API

sycl/doc/extensions/ExplicitSIMD/dpcpp-explicit-simd.md

@@ -155,17 +155,17 @@ a.select<4, 2>(0) = b;  // selected elements of a are replaced
 
 
 Gen ISA provides powerful register region addressing modes to facilitate cross-lane
-SIMD vector operation. To exploit this feature Explicit SIMD provides ```replicate``` function
+SIMD vector operation. To exploit this feature Explicit SIMD provides ```replicate``` functions
 to allow programmer to implement any native Gen ISA region in the following forms:
 - ```replicate<REP>()```: replicate a simd vector object **REP** times and return a new simd
 vector of **REP** * Width, where Width specifies the original vector size.
-- ```replicate<REP, W>(uint16_t i)```: replicate **W** consecutive elements starting at the
+- ```replicate_w<REP, W>(uint16_t i)```: replicate **W** consecutive elements starting at the
 i-th element from the simd vector object **REP** times, and return a new simd vector of **REP** * **W** length.
-- ```replicate<REP, VS, W>(uint16_t i)```: replicate **REP** blocks of **W** consecutive
+- ```replicate_vs_w<REP, VS, W>(uint16_t i)```: replicate **REP** blocks of **W** consecutive
 elements starting at the i-th from the simd vector object with each block strided by **VS**
 elements, and return a new vector of **REP** * **W** length. Selected blocks of **W**
 elements will overlap if **VS** < **W**.
-- ```replicate<REP, VS, W, HS>(uint16_t i=0 )```: replicate **REP** blocks of **W** sequential
+- ```replicate_vs_w_hs<REP, VS, W, HS>(uint16_t i=0 )```: replicate **REP** blocks of **W** sequential
 elements with a stride of **HS** starting at the i-th element from the simd vector object with
 each block strided by **VS** elements, and return a new vector of **REP** * **W** length.
 Selected blocks of **W** elements will overlap if **VS** < **W**.

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

@@ -344,8 +344,18 @@ template <typename Ty, int N> class simd {
   /// \tparam W is width of src region to replicate.
   /// \param Offset is offset in number of elements in src region.
   /// \return replicated simd instance.
-  template <int Rep, int W> simd<Ty, Rep * W> replicate(uint16_t Offset) {
-    return replicate<Rep, W, W, 1>(Offset);
+  template <int Rep, int W>
+  __SYCL_DEPRECATED("use simd::replicate_w")
+  simd<Ty, Rep * W> replicate(uint16_t Offset) {
+    return replicate_w<Rep, W>(Offset);
+  }
+
+  /// \tparam Rep is number of times region has to be replicated.
+  /// \tparam W is width of src region to replicate.
+  /// \param Offset is offset in number of elements in src region.
+  /// \return replicated simd instance.
+  template <int Rep, int W> simd<Ty, Rep * W> replicate_w(uint16_t Offset) {
+    return replicate_vs_w_hs<Rep, W, W, 1>(Offset);
   }
 
   /// \tparam Rep is number of times region has to be replicated.
@@ -354,31 +364,41 @@ template <typename Ty, int N> class simd {
   /// \param Offset is offset in number of elements in src region.
   /// \return replicated simd instance.
   template <int Rep, int VS, int W>
+  __SYCL_DEPRECATED("use simd::replicate_vs_w")
   simd<Ty, Rep * W> replicate(uint16_t Offset) {
-    return replicate<Rep, VS, W, 1>(Offset);
+    return replicate_vs_w<Rep, VS, W>(Offset);
+  }
+
+  /// \tparam Rep is number of times region has to be replicated.
+  /// \tparam VS vertical stride of src region to replicate.
+  /// \tparam W width of src region to replicate.
+  /// \param Offset offset in number of elements in src region.
+  /// \return replicated simd instance.
+  template <int Rep, int VS, int W>
+  simd<Ty, Rep * W> replicate_vs_w(uint16_t Offset) {
+    return replicate_vs_w_hs<Rep, VS, W, 1>(Offset);
   }
 
-  // TODO
-  // @rolandschulz
-  // {quote}
-  // - Template function with that many arguments are really ugly.
-  //   Are you sure there isn't a better interface? And that users won't
-  //   constantly forget what the correct order of the argument is?
-  //   Some kind of templated builder pattern would be a bit more verbose but
-  //   much more readable.
-  //   ...
-  //   The user would use (any of the extra method calls are optional)
-  //   s.replicate<R>(i).width<W>().vstride<VS>().hstride<HS>()
-  // {/quote}
-  // @jasonsewall-intel +1 for this
   /// \tparam Rep is number of times region has to be replicated.
   /// \tparam VS vertical stride of src region to replicate.
   /// \tparam W is width of src region to replicate.
   /// \tparam HS horizontal stride of src region to replicate.
   /// \param Offset is offset in number of elements in src region.
   /// \return replicated simd instance.
   template <int Rep, int VS, int W, int HS>
+  __SYCL_DEPRECATED("use simd::replicate_vs_w_hs")
   simd<Ty, Rep * W> replicate(uint16_t Offset) {
+    return replicate_vs_w_hs<Rep, VS, W, HS>(Offset);
+  }
+
+  /// \tparam Rep is number of times region has to be replicated.
+  /// \tparam VS vertical stride of src region to replicate.
+  /// \tparam W is width of src region to replicate.
+  /// \tparam HS horizontal stride of src region to replicate.
+  /// \param Offset is offset in number of elements in src region.
+  /// \return replicated simd instance.
+  template <int Rep, int VS, int W, int HS>
+  simd<Ty, Rep * W> replicate_vs_w_hs(uint16_t Offset) {
     return __esimd_rdregion<element_type, N, Rep * W, VS, W, HS, N>(
         data(), Offset * sizeof(Ty));
   }

sycl/test/esimd/simd.cpp

@@ -212,21 +212,21 @@ bool test_replicate() SYCL_ESIMD_FUNCTION {
 
 bool test_replicate1() SYCL_ESIMD_FUNCTION {
   simd<int, 8> v0(0, 1);
-  auto v0_rep = v0.replicate<4, 2>(2);
+  auto v0_rep = v0.replicate_w<4, 2>(2);
 
   return v0[2] == v0_rep[2] && v0[3] == v0_rep[5];
 }
 
 bool test_replicate2() SYCL_ESIMD_FUNCTION {
   simd<int, 8> v0(0, 1);
-  auto v0_rep = v0.replicate<2, 4, 2>(1);
+  auto v0_rep = v0.replicate_vs_w<2, 4, 2>(1);
 
   return v0_rep[0] == v0[1] && v0_rep[1] == v0[2] && v0_rep[2] == v0[5];
 }
 
 bool test_replicate3() SYCL_ESIMD_FUNCTION {
   simd<int, 8> v0(0, 1);
-  auto v0_rep = v0.replicate<2, 4, 2, 2>(1);
+  auto v0_rep = v0.replicate_vs_w_hs<2, 4, 2, 2>(1);
 
   return v0_rep[0] == v0[1] && v0_rep[1] == v0[3] && v0_rep[2] == v0[5];
 }

sycl/test/esimd/simd_replicate_deprecated.cpp

@@ -0,0 +1,37 @@
+// RUN: %clangxx -fsycl -fsycl-device-only -fsyntax-only -Xclang -verify %s
+
+// This test checks deprecated replicate ESIMD APIs
+
+#include <sycl/ext/intel/experimental/esimd.hpp>
+
+using namespace sycl::ext::intel::experimental::esimd;
+
+bool test_replicate1() {
+  simd<int, 8> v0(0, 1);
+  // expected-warning@+3 2 {{deprecated}}
+  // expected-note@sycl/ext/intel/experimental/esimd/simd.hpp:* {{}}
+  // expected-note@sycl/ext/intel/experimental/esimd/simd.hpp:* {{has been explicitly marked deprecated here}}
+  auto v0_rep = v0.replicate<4, 2>(2);
+
+  return v0[2] == v0_rep[2] && v0[3] == v0_rep[5];
+}
+
+bool test_replicate2() {
+  simd<int, 8> v0(0, 1);
+  // expected-warning@+3 2 {{deprecated}}
+  // expected-note@sycl/ext/intel/experimental/esimd/simd.hpp:* {{}}
+  // expected-note@sycl/ext/intel/experimental/esimd/simd.hpp:* {{has been explicitly marked deprecated here}}
+  auto v0_rep = v0.replicate<2, 4, 2>(1);
+
+  return v0_rep[0] == v0[1] && v0_rep[1] == v0[2] && v0_rep[2] == v0[5];
+}
+
+bool test_replicate3() {
+  simd<int, 8> v0(0, 1);
+  // expected-warning@+3 2 {{deprecated}}
+  // expected-note@sycl/ext/intel/experimental/esimd/simd.hpp:* {{}}
+  // expected-note@sycl/ext/intel/experimental/esimd/simd.hpp:* {{has been explicitly marked deprecated here}}
+  auto v0_rep = v0.replicate<2, 4, 2, 2>(1);
+
+  return v0_rep[0] == v0[1] && v0_rep[1] == v0[3] && v0_rep[2] == v0[5];
+}