Skip to content

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

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 6 commits into from
Sep 23, 2021
Merged

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

Show file tree
Hide file tree
Changes from 4 commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
ca15170
[ESIMD] rename slm_load4/slm_store4 to slm_load_rgba/slm_store_rgba
DenisBakhvalov Jul 21, 2021
8a54b66
Renamed slm_[load|store]_rgba -> slm_[gather|scatter]_rgba
DenisBakhvalov Jul 29, 2021
b672a97
Renamed slm_[load|store] to slm_[gather|scatter]
DenisBakhvalov Jul 29, 2021
2570610
Fix slm_store function
v-klochkov Sep 18, 2021
9d1b437
Merge remote-tracking branch 'intel_llvm/sycl' into slm_load4_rename
v-klochkov Sep 23, 2021
91ee441
Minor NFC fix
v-klochkov Sep 23, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
86 changes: 70 additions & 16 deletions sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -653,39 +653,93 @@ 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<uint16_t, n> pred = 1) {
ESIMD_INLINE ESIMD_NODEBUG std::enable_if_t<(n == 16 || n == 32), simd<T, n>>
slm_gather(simd<uint32_t, n> offsets, simd<uint16_t, 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<uint16_t, 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<uint16_t, 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<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_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<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>
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());
__SYCL_DEPRECATED("use slm_scatter_rgba.")
ESIMD_INLINE ESIMD_NODEBUG std::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_scatter_rgba<T, n, Mask>(vals, offsets, pred);
}

/// SLM block-load.
Expand Down
4 changes: 2 additions & 2 deletions sycl/test/esimd/slm_load.cpp → sycl/test/esimd/slm_gather_scatter.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
}
20 changes: 10 additions & 10 deletions sycl/test/esimd/slm_load4.cpp → sycl/test/esimd/slm_gather_scatter_rgba.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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>

Expand All @@ -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: