Skip to content

[SYCL][ESIMD] Add more stringent compile time checks to local_accessor version of block_load/block_store, gather/scatter API #11653

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 1 commit into from
Oct 26, 2023
Merged

[SYCL][ESIMD] Add more stringent compile time checks to local_accessor version of block_load/block_store, gather/scatter API #11653

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
49 changes: 27 additions & 22 deletions sycl/include/sycl/ext/intel/esimd/memory.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3049,11 +3049,12 @@ __ESIMD_API void media_block_store(AccessorTy acc, unsigned x, unsigned y,
///
template <typename Tx, int N, typename AccessorTy,
typename Flags = overaligned_tag<detail::OperandSize::OWORD>>
__ESIMD_API std::enable_if_t<
sycl::detail::acc_properties::is_local_accessor_v<AccessorTy> &&
is_simd_flag_type_v<Flags>,
simd<Tx, N>>
block_load(AccessorTy acc, uint32_t offset, Flags = {}) {
__ESIMD_API
std::enable_if_t<detail::is_local_accessor_with_v<
AccessorTy, detail::accessor_mode_cap::can_read> &&
is_simd_flag_type_v<Flags>,
simd<Tx, N>>
block_load(AccessorTy acc, uint32_t offset, Flags = {}) {
return slm_block_load<Tx, N, Flags>(offset +
__ESIMD_DNS::localAccessorToOffset(acc));
}
Expand All @@ -3077,10 +3078,11 @@ block_load(AccessorTy acc, uint32_t offset, Flags = {}) {
///
template <typename Tx, int N, typename AccessorTy,
typename Flags = overaligned_tag<detail::OperandSize::OWORD>>
__ESIMD_API std::enable_if_t<
sycl::detail::acc_properties::is_local_accessor_v<AccessorTy> &&
is_simd_flag_type_v<Flags>>
block_store(AccessorTy acc, uint32_t offset, simd<Tx, N> vals, Flags = {}) {
__ESIMD_API
std::enable_if_t<detail::is_local_accessor_with_v<
AccessorTy, detail::accessor_mode_cap::can_write> &&
is_simd_flag_type_v<Flags>>
block_store(AccessorTy acc, uint32_t offset, simd<Tx, N> vals, Flags = {}) {
slm_block_store<Tx, N, Flags>(
offset + __ESIMD_DNS::localAccessorToOffset(acc), vals);
}
Expand All @@ -3103,10 +3105,12 @@ block_store(AccessorTy acc, uint32_t offset, simd<Tx, N> vals, Flags = {}) {
/// undefined.
///
template <typename T, int N, typename AccessorTy>
__ESIMD_API std::enable_if_t<
sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>, simd<T, N>>
gather(AccessorTy acc, simd<uint32_t, N> offsets, uint32_t glob_offset = 0,
simd_mask<N> mask = 1) {
__ESIMD_API
std::enable_if_t<detail::is_local_accessor_with_v<
AccessorTy, detail::accessor_mode_cap::can_read>,
simd<T, N>>
gather(AccessorTy acc, simd<uint32_t, N> offsets, uint32_t glob_offset = 0,
simd_mask<N> mask = 1) {
return slm_gather<T, N>(
offsets + glob_offset + __ESIMD_DNS::localAccessorToOffset(acc), mask);
}
Expand All @@ -3130,8 +3134,8 @@ gather(AccessorTy acc, simd<uint32_t, N> offsets, uint32_t glob_offset = 0,
///
///
template <typename T, int N, typename AccessorTy>
__ESIMD_API std::enable_if_t<
sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>>
__ESIMD_API std::enable_if_t<detail::is_local_accessor_with_v<
AccessorTy, detail::accessor_mode_cap::can_write>>
scatter(AccessorTy acc, simd<uint32_t, N> offsets, simd<T, N> vals,
uint32_t glob_offset = 0, simd_mask<N> mask = 1) {
slm_scatter<T, N>(offsets + glob_offset +
Expand Down Expand Up @@ -3166,11 +3170,12 @@ scatter(AccessorTy acc, simd<uint32_t, N> offsets, simd<T, N> vals,
template <rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR,
typename AccessorT, int N,
typename T = typename AccessorT::value_type>
__ESIMD_API std::enable_if_t<
sycl::detail::acc_properties::is_local_accessor_v<AccessorT>,
simd<T, N * get_num_channels_enabled(RGBAMask)>>
gather_rgba(AccessorT acc, simd<uint32_t, N> offsets,
uint32_t global_offset = 0, simd_mask<N> mask = 1) {
__ESIMD_API
std::enable_if_t<detail::is_local_accessor_with_v<
AccessorT, detail::accessor_mode_cap::can_read>,
simd<T, N * get_num_channels_enabled(RGBAMask)>>
gather_rgba(AccessorT acc, simd<uint32_t, N> offsets,
uint32_t global_offset = 0, simd_mask<N> mask = 1) {
return slm_gather_rgba<T, N, RGBAMask>(
offsets + global_offset + __ESIMD_DNS::localAccessorToOffset(acc), mask);
}
Expand All @@ -3194,8 +3199,8 @@ gather_rgba(AccessorT acc, simd<uint32_t, N> offsets,
template <rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR,
typename AccessorT, int N,
typename T = typename AccessorT::value_type>
__ESIMD_API std::enable_if_t<
sycl::detail::acc_properties::is_local_accessor_v<AccessorT>>
__ESIMD_API std::enable_if_t<detail::is_local_accessor_with_v<
AccessorT, detail::accessor_mode_cap::can_write>>
scatter_rgba(AccessorT acc, simd<uint32_t, N> offsets,
simd<T, N * get_num_channels_enabled(RGBAMask)> vals,
uint32_t global_offset = 0, simd_mask<N> mask = 1) {
Expand Down
13 changes: 11 additions & 2 deletions sycl/test/esimd/block_load_store.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -30,7 +30,7 @@ SYCL_EXTERNAL void kernel2(int *ptr) SYCL_ESIMD_FUNCTION {

// Incompatible mode (write).
SYCL_EXTERNAL void
kernel3(accessor<int, 1, access::mode::write, access::target::device> &buf)
kernel4(accessor<int, 1, access::mode::write, access::target::device> &buf)
SYCL_ESIMD_FUNCTION {
simd<int, 32> v;
// CHECK: block_load_store.cpp:38{{.*}}error: no matching function
Expand All @@ -40,10 +40,19 @@ kernel3(accessor<int, 1, access::mode::write, access::target::device> &buf)

// Incompatible mode (read).
SYCL_EXTERNAL void
kernel4(accessor<int, 1, access::mode::read, access::target::device> &buf)
kernel5(accessor<int, 1, access::mode::read, access::target::device> &buf)
SYCL_ESIMD_FUNCTION {
simd<int, 32> v(0, 1);
// CHECK: block_load_store.cpp:48{{.*}}error: no matching function
// function for call to 'block_store'
block_store<int, 32>(buf, 0, v);
}

// Incompatible mode (read).
SYCL_EXTERNAL void
kernel6(local_accessor<const int, 1> &buf) SYCL_ESIMD_FUNCTION {
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

'const int' means it is read only. It is surprising that SYCL implementation allowed creating it because
a) it is odd to have read-only local accessor - that memory cannot be written and thus is not usable
b) SYCL 2020 SPEC says: "This specialization of accessor is only available for access modes access_mode::read_write and access_mode::atomic.".

simd<int, 32> v(0, 1);
// CHECK: block_load_store.cpp:57{{.*}}error: no matching function
// function for call to 'block_store'
block_store<int, 32>(buf, 0, v);
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

kernel1() in this test is the only positive test for load/store for read/write accessor.
Wouldn't it be useful to add 2 more cases - 1 load from read accesor and 1 store to write accessor?

}
20 changes: 20 additions & 0 deletions sycl/test/esimd/gather_scatter.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -104,3 +104,23 @@ kernel5(accessor<int, 1, access::mode::read, access::target::device> &buf)
// function for call to 'scatter'
scatter_rgba(buf, offset, v);
}

// Incompatible mode (read).
SYCL_EXTERNAL void
kernel6(local_accessor<const int, 1> &buf) SYCL_ESIMD_FUNCTION {
simd<int, 32> v(0, 1);
simd<uint32_t, 32> offset(0, 1);
// CHECK: gather_scatter.cpp:115{{.*}}error: no matching function
// function for call to 'scatter'
scatter<int, 32>(buf, offset, v);
}

// Incompatible mode (read).
SYCL_EXTERNAL void
kernel7(local_accessor<const int, 1> &buf) SYCL_ESIMD_FUNCTION {
simd<int, 32 * 4> v(0, 1);
simd<uint32_t, 32> offset(0, sizeof(int) * 4);
// CHECK: gather_scatter.cpp:125{{.*}}error: no matching function
// function for call to 'scatter'
scatter_rgba(buf, offset, v);
}