[ESIMD] rename slm_load4/slm_store4 to slm_load_rgba/slm_store_rgba

Author: DenisBakhvalov

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

@@ -668,24 +668,63 @@ ESIMD_INLINE ESIMD_NODEBUG
   __esimd_slm_write<T, n>(offsets.data(), vals.data(), pred.data());
 }
 
+/// 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 typename sycl::detail::enable_if_t<
+    (N == 8 || N == 16 || N == 32) && (sizeof(T) == 4),
+    simd<T, N * get_num_channels_enabled(Mask)>>
+slm_load_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>
+__SYCL_DEPRECATED("use slm_load_rgba.")
 ESIMD_INLINE ESIMD_NODEBUG typename sycl::detail::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<uint16_t, 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<uint16_t, n>
+                                                               pred = 1) {
+  return slm_load_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 typename sycl::detail::enable_if_t<
+    (N == 8 || N == 16 || N == 32) && (sizeof(T) == 4), void>
+slm_store_rgba(simd<T, N * get_num_channels_enabled(Mask)> vals,
+               simd<uint32_t, N> offsets, simd<uint16_t, 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>
+__SYCL_DEPRECATED("use slm_store_rgba.")
 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<uint16_t, n> pred = 1) {
-  __esimd_slm_write4<T, n, Mask>(offsets.data(), vals.data(), pred.data());
+    (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<uint16_t, n> pred = 1) {
+  slm_store_rgba<T, n, Mask>(vals, offsets, pred);
 }
 
 /// SLM block-load.

sycl/test/esimd/slm_load4.cpp renamed to sycl/test/esimd/slm_load_store_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_load_rgba/slm_store_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_load_store_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_load_rgba<int, 32, ESIMD_ABGR_ENABLE>(offsets);
+  v0 = slm_load_rgba<int, 32, rgba_channel_mask::ABGR>(offsets);
 
   v0 = v0 + v1;
 
-  // CHECK: slm_load4.cpp:26{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated
+  // CHECK: slm_load_store_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_store_rgba<int, 32, ESIMD_ABGR_ENABLE>(v0, offsets);
+  slm_store_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_load_store_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_load_store_rgba.cpp:26{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated
 // CHECK: sycl/ext/intel/experimental/esimd/common.hpp:{{.*}}note: