[ESIMD] Rename slm_load4/slm_store4 to slm_load_rgba/slm_store_rgba (#4158)

* [ESIMD] rename slm_load4/slm_store4 to slm_load_rgba/slm_store_rgba

Signed-off-by: Vyacheslav N Klochkov <vyacheslav.n.klochkov@intel.com>

Co-authored-by: Denis Bakhvalov

Author: DenisBakhvalov

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

@@ -655,39 +655,92 @@ SYCL_EXTERNAL SYCL_ESIMD_FUNCTION void slm_init(uint32_t size);
 ///
 /// Only allow simd-16 and simd-32.
 template <typename T, int n>
-ESIMD_INLINE ESIMD_NODEBUG
-    typename sycl::detail::enable_if_t<(n == 16 || n == 32), simd<T, n>>
-    slm_load(simd<uint32_t, n> offsets, simd_mask<n> Pred = 1) {
-  return __esimd_slm_read<T, n>(offsets.data(), Pred.data());
+ESIMD_INLINE ESIMD_NODEBUG std::enable_if_t<(n == 16 || n == 32), simd<T, n>>
+slm_gather(simd<uint32_t, n> offsets, simd_mask<n> pred = 1) {
+  return __esimd_slm_read<T, n>(offsets.data(), pred.data());
+}
+
+/// SLM gather (deprecated version).
+template <typename T, int n>
+__SYCL_DEPRECATED("use slm_gather.")
+ESIMD_INLINE
+    ESIMD_NODEBUG std::enable_if_t<(n == 16 || n == 32), simd<T, n>> slm_load(
+        simd<uint32_t, n> offsets, simd<uint16_t, n> pred = 1) {
+  return slm_gather<T, n>(offsets, pred);
 }
 
 /// SLM scatter.
 template <typename T, int n>
-ESIMD_INLINE ESIMD_NODEBUG
-    typename sycl::detail::enable_if_t<(n == 16 || n == 32), void>
-    slm_store(simd<T, n> vals, simd<uint32_t, n> offsets,
-              simd_mask<n> pred = 1) {
+ESIMD_INLINE ESIMD_NODEBUG std::enable_if_t<(n == 16 || n == 32)>
+slm_scatter(simd<T, n> vals, simd<uint32_t, n> offsets, simd_mask<n> pred = 1) {
   __esimd_slm_write<T, n>(offsets.data(), vals.data(), pred.data());
 }
 
+/// SLM scatter (deprecated version).
+template <typename T, int n>
+__SYCL_DEPRECATED("use slm_scatter.")
+ESIMD_INLINE ESIMD_NODEBUG std::enable_if_t<(n == 16 || n == 32)> slm_store(
+    simd<T, n> vals, simd<uint32_t, n> offsets, simd<uint16_t, n> pred = 1) {
+  slm_scatter<T, n>(vals, offsets, pred);
+}
+
+/// Gathering read from the SLM given specified \p offsets.
+/// Up to 4 data elements may be accessed at each address depending on the
+/// enabled channel \p Mask.
+/// \tparam T element type of the returned vector. Must be 4-byte.
+/// \tparam N size of the \p offsets vector. Must be 8, 16 or 32.
+/// \tparam Mask represents a pixel's channel mask.
+/// @param offsets byte-offsets within the SLM.
+/// @param pred predication control used for masking lanes.
+/// \ingroup sycl_esimd
+template <typename T, int N, rgba_channel_mask Mask>
+ESIMD_INLINE ESIMD_NODEBUG
+    std::enable_if_t<(N == 8 || N == 16 || N == 32) && (sizeof(T) == 4),
+                     simd<T, N * get_num_channels_enabled(Mask)>>
+    slm_gather_rgba(simd<uint32_t, N> offsets, simd<uint16_t, N> pred = 1) {
+  return __esimd_slm_read4<T, N, Mask>(offsets.data(), pred.data());
+}
+
 /// SLM gather4.
 ///
 /// Only allow simd-8, simd-16 and simd-32.
 template <typename T, int n, rgba_channel_mask Mask>
-ESIMD_INLINE ESIMD_NODEBUG typename sycl::detail::enable_if_t<
+__SYCL_DEPRECATED("use slm_gather_rgba.")
+ESIMD_INLINE ESIMD_NODEBUG std::enable_if_t<
     (n == 8 || n == 16 || n == 32) && (sizeof(T) == 4),
-    simd<T, n * get_num_channels_enabled(Mask)>>
-slm_load4(simd<uint32_t, n> offsets, simd_mask<n> pred = 1) {
-  return __esimd_slm_read4<T, n, Mask>(offsets.data(), pred.data());
+    simd<T, n * get_num_channels_enabled(Mask)>> slm_load4(simd<uint32_t, n>
+                                                               offsets,
+                                                           simd_mask<n> pred =
+                                                               1) {
+  return slm_gather_rgba<T, n, Mask>(offsets, pred);
+}
+
+/// Scatter write to the SLM given specified \p offsets.
+/// Up to 4 data elements may be written at each address depending on the
+/// enabled channel \p Mask.
+/// \tparam T element type of the input vector. Must be 4-byte.
+/// \tparam N size of the \p offsets vector. Must be 8, 16 or 32.
+/// \tparam Mask represents a pixel's channel mask.
+/// @param vals values to be written.
+/// @param offsets byte-offsets within the SLM.
+/// @param pred predication control used for masking lanes.
+/// \ingroup sycl_esimd
+template <typename T, int N, rgba_channel_mask Mask>
+ESIMD_INLINE ESIMD_NODEBUG
+    std::enable_if_t<(N == 8 || N == 16 || N == 32) && (sizeof(T) == 4)>
+    slm_scatter_rgba(simd<T, N * get_num_channels_enabled(Mask)> vals,
+                     simd<uint32_t, N> offsets, simd_mask<N> pred = 1) {
+  __esimd_slm_write4<T, N, Mask>(offsets.data(), vals.data(), pred.data());
 }
 
 /// SLM scatter4.
 template <typename T, int n, rgba_channel_mask Mask>
-ESIMD_INLINE ESIMD_NODEBUG typename sycl::detail::enable_if_t<
-    (n == 8 || n == 16 || n == 32) && (sizeof(T) == 4), void>
-slm_store4(simd<T, n * get_num_channels_enabled(Mask)> vals,
-           simd<uint32_t, n> offsets, simd_mask<n> pred = 1) {
-  __esimd_slm_write4<T, n, Mask>(offsets.data(), vals.data(), pred.data());
+__SYCL_DEPRECATED("use slm_scatter_rgba.")
+ESIMD_INLINE ESIMD_NODEBUG std::
+    enable_if_t<(n == 8 || n == 16 || n == 32) && (sizeof(T) == 4)> slm_store4(
+        simd<T, n * get_num_channels_enabled(Mask)> vals,
+        simd<uint32_t, n> offsets, simd_mask<n> pred = 1) {
+  slm_scatter_rgba<T, n, Mask>(vals, offsets, pred);
 }
 
 /// SLM block-load.

sycl/test/esimd/slm_load.cpp renamed to sycl/test/esimd/slm_gather_scatter.cpp

@@ -13,12 +13,12 @@ void kernel() __attribute__((sycl_device)) {
   simd<uint32_t, 32> offsets(0, 1);
   simd<int, 32> v1(0, 1);
 
-  auto v0 = slm_load<int, 32>(offsets);
+  auto v0 = slm_gather<int, 32>(offsets);
 
   esimd_fence(3);
   esimd_barrier();
 
   v0 = v0 + v1;
 
-  slm_store<int, 32>(v0, offsets);
+  slm_scatter<int, 32>(v0, offsets);
 }

sycl/test/esimd/slm_load4.cpp renamed to sycl/test/esimd/slm_gather_scatter_rgba.cpp

@@ -1,7 +1,7 @@
 // RUN: %clangxx -fsycl -fsyntax-only -Wno-unused-command-line-argument %s 2>&1 | FileCheck %s --implicit-check-not="warning:" --implicit-check-not="error:"
 
-// This test checks compilation of ESIMD slm load4/store4 APIs. Those which are
-// deprecated must produce deprecation messages.
+// This test checks compilation of ESIMD slm_gather_rgba/slm_scatter_rgba APIs.
+// Those which are deprecated must produce deprecation messages.
 
 #include <sycl/ext/intel/experimental/esimd.hpp>
 
@@ -14,22 +14,22 @@ void caller() SYCL_ESIMD_FUNCTION {
 
   slm_init(1024);
 
-  // CHECK: slm_load4.cpp:19{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated
+  // CHECK: slm_gather_scatter_rgba.cpp:19{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated
   // CHECK: sycl/ext/intel/experimental/esimd/common.hpp:{{.*}}note:
-  auto v0 = slm_load4<int, 32, ESIMD_ABGR_ENABLE>(offsets);
-  v0 = slm_load4<int, 32, rgba_channel_mask::ABGR>(offsets);
+  auto v0 = slm_gather_rgba<int, 32, ESIMD_ABGR_ENABLE>(offsets);
+  v0 = slm_gather_rgba<int, 32, rgba_channel_mask::ABGR>(offsets);
 
   v0 = v0 + v1;
 
-  // CHECK: slm_load4.cpp:26{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated
+  // CHECK: slm_gather_scatter_rgba.cpp:26{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated
   // CHECK: sycl/ext/intel/experimental/esimd/common.hpp:{{.*}}note:
-  slm_store4<int, 32, ESIMD_ABGR_ENABLE>(v0, offsets);
-  slm_store4<int, 32, rgba_channel_mask::ABGR>(v0, offsets);
+  slm_scatter_rgba<int, 32, ESIMD_ABGR_ENABLE>(v0, offsets);
+  slm_scatter_rgba<int, 32, rgba_channel_mask::ABGR>(v0, offsets);
 }
 
 // A "border" between host and device compilations
 // CHECK-LABEL: 2 warnings generated
-// CHECK: slm_load4.cpp:19{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated
+// CHECK: slm_gather_scatter_rgba.cpp:19{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated
 // CHECK: sycl/ext/intel/experimental/esimd/common.hpp:{{.*}}note:
-// CHECK: slm_load4.cpp:26{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated
+// CHECK: slm_gather_scatter_rgba.cpp:26{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated
 // CHECK: sycl/ext/intel/experimental/esimd/common.hpp:{{.*}}note: