diff --git a/sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp b/sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp index 96b6e7bbfbcdb..28ba0a0e3231a 100644 --- a/sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp +++ b/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 -ESIMD_INLINE ESIMD_NODEBUG - typename sycl::detail::enable_if_t<(n == 16 || n == 32), simd> - slm_load(simd offsets, simd_mask Pred = 1) { - return __esimd_slm_read(offsets.data(), Pred.data()); +ESIMD_INLINE ESIMD_NODEBUG std::enable_if_t<(n == 16 || n == 32), simd> +slm_gather(simd offsets, simd_mask pred = 1) { + return __esimd_slm_read(offsets.data(), pred.data()); +} + +/// SLM gather (deprecated version). +template +__SYCL_DEPRECATED("use slm_gather.") +ESIMD_INLINE + ESIMD_NODEBUG std::enable_if_t<(n == 16 || n == 32), simd> slm_load( + simd offsets, simd pred = 1) { + return slm_gather(offsets, pred); } /// SLM scatter. template -ESIMD_INLINE ESIMD_NODEBUG - typename sycl::detail::enable_if_t<(n == 16 || n == 32), void> - slm_store(simd vals, simd offsets, - simd_mask pred = 1) { +ESIMD_INLINE ESIMD_NODEBUG std::enable_if_t<(n == 16 || n == 32)> +slm_scatter(simd vals, simd offsets, simd_mask pred = 1) { __esimd_slm_write(offsets.data(), vals.data(), pred.data()); } +/// SLM scatter (deprecated version). +template +__SYCL_DEPRECATED("use slm_scatter.") +ESIMD_INLINE ESIMD_NODEBUG std::enable_if_t<(n == 16 || n == 32)> slm_store( + simd vals, simd offsets, simd pred = 1) { + slm_scatter(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 +ESIMD_INLINE ESIMD_NODEBUG + std::enable_if_t<(N == 8 || N == 16 || N == 32) && (sizeof(T) == 4), + simd> + slm_gather_rgba(simd offsets, simd pred = 1) { + return __esimd_slm_read4(offsets.data(), pred.data()); +} + /// SLM gather4. /// /// Only allow simd-8, simd-16 and simd-32. template -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> -slm_load4(simd offsets, simd_mask pred = 1) { - return __esimd_slm_read4(offsets.data(), pred.data()); + simd> slm_load4(simd + offsets, + simd_mask pred = + 1) { + return slm_gather_rgba(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 +ESIMD_INLINE ESIMD_NODEBUG + std::enable_if_t<(N == 8 || N == 16 || N == 32) && (sizeof(T) == 4)> + slm_scatter_rgba(simd vals, + simd offsets, simd_mask pred = 1) { + __esimd_slm_write4(offsets.data(), vals.data(), pred.data()); } /// SLM scatter4. template -ESIMD_INLINE ESIMD_NODEBUG typename sycl::detail::enable_if_t< - (n == 8 || n == 16 || n == 32) && (sizeof(T) == 4), void> -slm_store4(simd vals, - simd offsets, simd_mask pred = 1) { - __esimd_slm_write4(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 vals, + simd offsets, simd_mask pred = 1) { + slm_scatter_rgba(vals, offsets, pred); } /// SLM block-load. diff --git a/sycl/test/esimd/slm_load.cpp b/sycl/test/esimd/slm_gather_scatter.cpp similarity index 85% rename from sycl/test/esimd/slm_load.cpp rename to sycl/test/esimd/slm_gather_scatter.cpp index 4aa64bb559cd1..f9318d8ddafe4 100644 --- a/sycl/test/esimd/slm_load.cpp +++ b/sycl/test/esimd/slm_gather_scatter.cpp @@ -13,12 +13,12 @@ void kernel() __attribute__((sycl_device)) { simd offsets(0, 1); simd v1(0, 1); - auto v0 = slm_load(offsets); + auto v0 = slm_gather(offsets); esimd_fence(3); esimd_barrier(); v0 = v0 + v1; - slm_store(v0, offsets); + slm_scatter(v0, offsets); } diff --git a/sycl/test/esimd/slm_load4.cpp b/sycl/test/esimd/slm_gather_scatter_rgba.cpp similarity index 50% rename from sycl/test/esimd/slm_load4.cpp rename to sycl/test/esimd/slm_gather_scatter_rgba.cpp index a93a00752ead9..5bf51b95e3c21 100644 --- a/sycl/test/esimd/slm_load4.cpp +++ b/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 @@ -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(offsets); - v0 = slm_load4(offsets); + auto v0 = slm_gather_rgba(offsets); + v0 = slm_gather_rgba(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(v0, offsets); - slm_store4(v0, offsets); + slm_scatter_rgba(v0, offsets); + slm_scatter_rgba(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: