Renamed slm_[load|store]_rgba -> slm_[gather|scatter]_rgba

Author: DenisBakhvalov

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

@@ -681,22 +681,22 @@ 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) {
+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>
-__SYCL_DEPRECATED("use slm_load_rgba.")
+__SYCL_DEPRECATED("use slm_gather_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 slm_load_rgba<T, n, Mask>(offsets, pred);
+  return slm_gather_rgba<T, n, Mask>(offsets, pred);
 }
 
 /// Scatter write to the SLM given specified \p offsets.
@@ -712,19 +712,19 @@ ESIMD_INLINE ESIMD_NODEBUG typename sycl::detail::enable_if_t<
 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) {
+slm_scatter_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.")
+__SYCL_DEPRECATED("use slm_scatter_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) {
-  slm_store_rgba<T, n, Mask>(vals, offsets, pred);
+  slm_scatter_rgba<T, n, Mask>(vals, offsets, pred);
 }
 
 /// SLM block-load.

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

@@ -1,6 +1,6 @@
 // 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_load_rgba/slm_store_rgba APIs.
+// 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_load_store_rgba.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_load_rgba<int, 32, ESIMD_ABGR_ENABLE>(offsets);
-  v0 = slm_load_rgba<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_load_store_rgba.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_store_rgba<int, 32, ESIMD_ABGR_ENABLE>(v0, offsets);
-  slm_store_rgba<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_load_store_rgba.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_load_store_rgba.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: