Skip to content

[SYCL][ESIMD] Add more stringent compile time checks to atomic_update API #11683

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 22 commits into from
Dec 14, 2023
Merged

[SYCL][ESIMD] Add more stringent compile time checks to atomic_update API #11683

Show file tree
Hide file tree
Changes from 4 commits
Commits
Show all changes
22 commits
Select commit Hold shift + click to select a range
a7f4098
Introduce more stringent checks for atomic_update API
fineg74 Oct 26, 2023
c9bfece
Merge remote-tracking branch 'origin/sycl' into accessorCheck3
fineg74 Oct 26, 2023
f2b0c38
Add additional tests for block_load/store API
fineg74 Oct 27, 2023
cc1ce4c
Merge remote-tracking branch 'origin/sycl' into accessorCheck3
fineg74 Dec 1, 2023
b4ca8fe
Update the new API with checks,
fineg74 Dec 4, 2023
9671245
Fix test failures
fineg74 Dec 4, 2023
885e963
Address PR comments
fineg74 Dec 5, 2023
c3b734a
Merge branches 'sycl', 'accessorCheck3', 'accessorCheck3' and 'access…
fineg74 Dec 6, 2023
098ac3e
Merge remote-tracking branch 'origin/sycl' into accessorCheck3
fineg74 Dec 6, 2023
b10017b
Simplify the compile time checks
fineg74 Dec 7, 2023
339e376
Merge remote-tracking branch 'origin/sycl' into accessorCheck3
fineg74 Dec 7, 2023
fea8a3d
Merge remote-tracking branch 'origin/sycl' into accessorCheck3
fineg74 Dec 8, 2023
d574160
Update sycl/include/sycl/ext/intel/esimd/memory.hpp
fineg74 Dec 8, 2023
134f11c
Update sycl/include/sycl/ext/intel/esimd/memory.hpp
fineg74 Dec 8, 2023
b55f7d4
Update sycl/include/sycl/ext/intel/esimd/memory.hpp
fineg74 Dec 8, 2023
02ce465
Update sycl/test/esimd/flat_atomic.cpp
fineg74 Dec 8, 2023
842426c
Address PR comments
fineg74 Dec 8, 2023
d8eee69
Merge branch 'accessorCheck3' of https://github.com/fineg74/llvm into…
fineg74 Dec 9, 2023
6a74b23
Revert some of the last changes
fineg74 Dec 9, 2023
15a25a4
Address PR comments
fineg74 Dec 13, 2023
c60e0a0
Fix build break
fineg74 Dec 13, 2023
0e4131c
Address PR comments
fineg74 Dec 14, 2023
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
81 changes: 51 additions & 30 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 @@ -5666,14 +5666,15 @@ atomic_update(AccessorTy acc, Toffset byte_offset, simd_mask<N> mask) {
/// @return A vector of the old values at the memory locations before the
/// update.
///

template <atomic_op Op, typename T, int N, typename Toffset,
typename AccessorTy,
typename PropertyListT =
ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
__ESIMD_DNS::get_num_args<Op>() == 1 && std::is_integral_v<Toffset> &&
sycl::detail::acc_properties::is_accessor_v<AccessorTy> &&
!sycl::detail::acc_properties::is_local_accessor_v<AccessorTy> &&
detail::is_device_accessor_with_v<
AccessorTy, detail::accessor_mode_cap::can_write> &&
ext::oneapi::experimental::is_property_list_v<PropertyListT>,
simd<T, N>>
atomic_update(AccessorTy acc, simd<Toffset, N> byte_offset, simd<T, N> src0,
Expand Down Expand Up @@ -5898,10 +5899,12 @@ atomic_update(AccessorTy acc, simd<Toffset, N> byte_offset,
/// update.
///
template <atomic_op Op, typename Tx, int N, typename AccessorTy>
__ESIMD_API std::enable_if_t<
sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>, simd<Tx, N>>
atomic_update(AccessorTy acc, simd<uint32_t, N> offset, simd<Tx, N> src0,
simd_mask<N> mask) {
__ESIMD_API
std::enable_if_t<detail::is_local_accessor_with_v<
AccessorTy, detail::accessor_mode_cap::can_write>,
simd<Tx, N>>
atomic_update(AccessorTy acc, simd<uint32_t, N> offset, simd<Tx, N> src0,
simd_mask<N> mask) {
if constexpr ((Op == atomic_op::fmin) || (Op == atomic_op::fmax) ||
(Op == atomic_op::fadd) || (Op == atomic_op::fsub)) {
// Auto-convert FP atomics to LSC version.
Expand Down Expand Up @@ -5958,7 +5961,9 @@ template <atomic_op Op, typename T, int N, typename Toffset,
typename PropertyListT =
ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
__ESIMD_DNS::get_num_args<Op>() == 1 && !std::is_pointer_v<AccessorTy> &&
__ESIMD_DNS::get_num_args<Op>() == 1 &&
detail::is_accessor_with_v<AccessorTy,
detail::accessor_mode_cap::can_write> &&
ext::oneapi::experimental::is_property_list_v<PropertyListT>,
simd<T, N>>
atomic_update(AccessorTy acc, simd_view<Toffset, RegionTy> byte_offset,
Expand Down Expand Up @@ -6000,7 +6005,9 @@ template <atomic_op Op, typename T, int N, typename Toffset,
typename PropertyListT =
ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
__ESIMD_DNS::get_num_args<Op>() == 1 && !std::is_pointer_v<AccessorTy> &&
__ESIMD_DNS::get_num_args<Op>() == 1 &&
detail::is_accessor_with_v<AccessorTy,
detail::accessor_mode_cap::can_write> &&
ext::oneapi::experimental::is_property_list_v<PropertyListT>,
simd<T, N>>
atomic_update(AccessorTy acc, simd_view<Toffset, RegionTy> byte_offset,
Expand Down Expand Up @@ -6047,7 +6054,9 @@ template <atomic_op Op, typename T, int N, typename Toffset,
typename PropertyListT =
ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
__ESIMD_DNS::get_num_args<Op>() == 1 && !std::is_pointer_v<AccessorTy> &&
__ESIMD_DNS::get_num_args<Op>() == 1 &&
detail::is_accessor_with_v<AccessorTy,
detail::accessor_mode_cap::can_write> &&
ext::oneapi::experimental::is_property_list_v<PropertyListT>,
simd<T, N>>
atomic_update(AccessorTy acc, simd_view<Toffset, OffsetRegionTy> byte_offset,
Expand Down Expand Up @@ -6093,7 +6102,9 @@ template <atomic_op Op, typename T, int N, typename Toffset,
typename PropertyListT =
ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
__ESIMD_DNS::get_num_args<Op>() == 1 && !std::is_pointer_v<AccessorTy> &&
__ESIMD_DNS::get_num_args<Op>() == 1 &&
detail::is_accessor_with_v<AccessorTy,
detail::accessor_mode_cap::can_write> &&
ext::oneapi::experimental::is_property_list_v<PropertyListT>,
simd<T, N>>
atomic_update(AccessorTy acc, simd_view<Toffset, OffsetRegionTy> byte_offset,
Expand Down Expand Up @@ -6127,7 +6138,9 @@ atomic_update(AccessorTy acc, simd_view<Toffset, OffsetRegionTy> byte_offset,
template <atomic_op Op, typename Tx, int N, typename Toffset,
typename AccessorTy>
__ESIMD_API std::enable_if_t<
std::is_integral_v<Toffset> && !std::is_pointer_v<AccessorTy> &&
std::is_integral_v<Toffset> &&
detail::is_accessor_with_v<AccessorTy,
detail::accessor_mode_cap::can_write> &&
((Op != atomic_op::store && Op != atomic_op::xchg) || N == 1),
simd<Tx, N>>
atomic_update(AccessorTy acc, Toffset offset, simd<Tx, N> src0,
Expand Down Expand Up @@ -6160,13 +6173,13 @@ atomic_update(AccessorTy acc, Toffset offset, simd<Tx, N> src0,
///
template <atomic_op Op, typename Tx, int N, typename Toffset,
typename AccessorTy>
__ESIMD_API std::enable_if_t<
std::is_integral_v<Toffset> &&
sycl::detail::acc_properties::is_accessor_v<AccessorTy> &&
!sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>,
simd<Tx, N>>
atomic_update(AccessorTy acc, simd<Toffset, N> offset, simd<Tx, N> src0,
simd<Tx, N> src1, simd_mask<N> mask) {
__ESIMD_API
std::enable_if_t<std::is_integral_v<Toffset> &&
detail::is_device_accessor_with_v<
AccessorTy, detail::accessor_mode_cap::can_write>,
simd<Tx, N>>
atomic_update(AccessorTy acc, simd<Toffset, N> offset, simd<Tx, N> src0,
simd<Tx, N> src1, simd_mask<N> mask) {
#ifdef __ESIMD_FORCE_STATELESS_MEM
return atomic_update<Op, Tx, N>(__ESIMD_DNS::accessorToPointer<Tx>(acc),
offset, src0, src1, mask);
Expand Down Expand Up @@ -6211,10 +6224,12 @@ atomic_update(AccessorTy acc, simd<Toffset, N> offset, simd<Tx, N> src0,
/// update.
///
template <atomic_op Op, typename Tx, int N, typename AccessorTy>
__ESIMD_API std::enable_if_t<
sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>, simd<Tx, N>>
atomic_update(AccessorTy acc, simd<uint32_t, N> offset, simd<Tx, N> src0,
simd<Tx, N> src1, simd_mask<N> mask) {
__ESIMD_API
std::enable_if_t<detail::is_local_accessor_with_v<
AccessorTy, detail::accessor_mode_cap::can_write>,
simd<Tx, N>>
atomic_update(AccessorTy acc, simd<uint32_t, N> offset, simd<Tx, N> src0,
simd<Tx, N> src1, simd_mask<N> mask) {
if constexpr (Op == atomic_op::fcmpxchg) {
// Auto-convert FP atomics to LSC version.
return atomic_update<detail::to_lsc_atomic_op<Op>(), Tx, N>(
Expand Down Expand Up @@ -6247,10 +6262,13 @@ atomic_update(AccessorTy acc, simd<uint32_t, N> offset, simd<Tx, N> src0,
///
template <atomic_op Op, typename Tx, int N, typename Toffset,
typename AccessorTy, typename RegionTy = region1d_t<Toffset, N, 1>>
__ESIMD_API std::enable_if_t<
std::is_integral_v<Toffset> && !std::is_pointer_v<AccessorTy>, simd<Tx, N>>
atomic_update(AccessorTy acc, simd_view<Toffset, RegionTy> offsets,
simd<Tx, N> src0, simd<Tx, N> src1, simd_mask<N> mask) {
__ESIMD_API
std::enable_if_t<std::is_integral_v<Toffset> &&
detail::is_accessor_with_v<
AccessorTy, detail::accessor_mode_cap::can_write>,
simd<Tx, N>>
atomic_update(AccessorTy acc, simd_view<Toffset, RegionTy> offsets,
simd<Tx, N> src0, simd<Tx, N> src1, simd_mask<N> mask) {
return atomic_update<Op, Tx, N>(acc, offsets.read(), src0, src1, mask);
}

Expand All @@ -6276,10 +6294,13 @@ atomic_update(AccessorTy acc, simd_view<Toffset, RegionTy> offsets,
///
template <atomic_op Op, typename Tx, int N, typename Toffset,
typename AccessorTy>
__ESIMD_API std::enable_if_t<
std::is_integral_v<Toffset> && !std::is_pointer_v<AccessorTy>, simd<Tx, N>>
atomic_update(AccessorTy acc, Toffset offset, simd<Tx, N> src0,
simd<Tx, N> src1, simd_mask<N> mask) {
__ESIMD_API
std::enable_if_t<std::is_integral_v<Toffset> &&
detail::is_accessor_with_v<
AccessorTy, detail::accessor_mode_cap::can_write>,
simd<Tx, N>>
atomic_update(AccessorTy acc, Toffset offset, simd<Tx, N> src0,
simd<Tx, N> src1, simd_mask<N> mask) {
return atomic_update<Op, Tx, N>(acc, simd<Toffset, N>(offset), src0, src1,
mask);
}
Expand Down
27 changes: 20 additions & 7 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 @@ -19,7 +19,20 @@ kernel1(accessor<int, 1, access::mode::read_write, access::target::device> &buf)
block_store<int, 32>(buf, 0, v0);
}

SYCL_EXTERNAL void kernel2(int *ptr) SYCL_ESIMD_FUNCTION {
SYCL_EXTERNAL void
kernel2(accessor<int, 1, access::mode::read, access::target::device> &buf)
SYCL_ESIMD_FUNCTION {
auto v0 = block_load<int, 32>(buf, 0);
}

SYCL_EXTERNAL void
kernel3(accessor<int, 1, access::mode::write, access::target::device> &buf)
SYCL_ESIMD_FUNCTION {
simd<int, 32> v1(0, 1);
block_store<int, 32>(buf, 0, v1);
}

SYCL_EXTERNAL void kernel4(int *ptr) SYCL_ESIMD_FUNCTION {
simd<int, 32> v1(0, 1);
auto v0 = block_load<int, 32>(ptr);
v0 = v0 + v1;
Expand All @@ -30,29 +43,29 @@ SYCL_EXTERNAL void kernel2(int *ptr) SYCL_ESIMD_FUNCTION {

// Incompatible mode (write).
SYCL_EXTERNAL void
kernel4(accessor<int, 1, access::mode::write, access::target::device> &buf)
kernel5(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
// CHECK: block_load_store.cpp:51{{.*}}error: no matching function
// function for call to 'block_load'
v = block_load<int, 32>(buf, 0);
}

// Incompatible mode (read).
SYCL_EXTERNAL void
kernel5(accessor<int, 1, access::mode::read, access::target::device> &buf)
kernel6(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
// CHECK: block_load_store.cpp:61{{.*}}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 {
kernel7(local_accessor<const int, 1> &buf) SYCL_ESIMD_FUNCTION {
simd<int, 32> v(0, 1);
// CHECK: block_load_store.cpp:57{{.*}}error: no matching function
// CHECK: block_load_store.cpp:70{{.*}}error: no matching function
// function for call to 'block_store'
block_store<int, 32>(buf, 0, v);
}
105 changes: 104 additions & 1 deletion sycl/test/esimd/flat_atomic.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,4 +1,5 @@
// RUN: %clangxx -fsycl -fsyntax-only -Wno-unused-command-line-argument %s
// RUN: not %clangxx -fsycl -fsycl-device-only -fsyntax-only -Wno-unused-command-line-argument %s 2>&1 | FileCheck %s --implicit-check-not="warning:" --implicit-check-not="error:"
// RUN: not %clangxx %fsycl-host-only -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 atomic APIs.

Expand All @@ -10,6 +11,8 @@
using namespace sycl::ext::intel::esimd;
using namespace sycl;

// --- Postive tests.

void kernel0(uint32_t *ptr) SYCL_ESIMD_FUNCTION {
simd<uint32_t, 32> offsets(0, 1);

Expand All @@ -31,3 +34,103 @@ template <class T> void kernel2(T *ptr) SYCL_ESIMD_FUNCTION {
}

template void kernel2<uint32_t>(uint32_t *) SYCL_ESIMD_FUNCTION;

void kernel3(accessor<uint32_t, 1, access::mode::read_write,
access::target::device> &buf) SYCL_ESIMD_FUNCTION {
simd<uint32_t, 32> offsets(0, 1);

atomic_update<atomic_op::inc, uint32_t, 32>(buf, offsets, 1);
}

void kernel4(accessor<uint32_t, 1, access::mode::read_write,
access::target::device> &buf) SYCL_ESIMD_FUNCTION {
simd<uint32_t, 32> offsets(0, 1);
simd<uint32_t, 32> v1(0, 1);

atomic_update<atomic_op::add, uint32_t, 32>(buf, offsets, v1, 1);
}

void kernel5(accessor<uint32_t, 1, access::mode::read_write,
access::target::device> &buf) SYCL_ESIMD_FUNCTION {
simd<uint32_t, 32> offsets(0, 1);
simd<uint32_t, 32> v1(0, 1);

atomic_update<atomic_op::cmpxchg, uint32_t, 32>(buf, offsets, v1, v1, 1);
}

void kernel6(local_accessor<uint32_t, 1> &buf) SYCL_ESIMD_FUNCTION {
simd<uint32_t, 32> offsets(0, 1);

atomic_update<atomic_op::inc, uint32_t, 32>(buf, offsets, 1);
}

void kernel7(local_accessor<uint32_t, 1> &buf) SYCL_ESIMD_FUNCTION {
simd<uint32_t, 32> offsets(0, 1);
simd<uint32_t, 32> v1(0, 1);

atomic_update<atomic_op::add, uint32_t, 32>(buf, offsets, v1, 1);
}

void kernel8(local_accessor<uint32_t, 1> &buf) SYCL_ESIMD_FUNCTION {
simd<uint32_t, 32> offsets(0, 1);
simd<uint32_t, 32> v1(0, 1);

atomic_update<atomic_op::cmpxchg, uint32_t, 32>(buf, offsets, v1, v1, 1);
}

// --- Negative tests.

// Incompatible mode (read).
void kernel9(accessor<uint32_t, 1, access::mode::read, access::target::device>
&buf) SYCL_ESIMD_FUNCTION {
simd<uint32_t, 32> offsets(0, 1);

// CHECK: flat_atomic.cpp:89{{.*}}error: no matching function for call to 'atomic_update'
atomic_update<atomic_op::inc, uint32_t, 32>(buf, offsets, 1);
}

// Incompatible mode (read).
void kernel10(accessor<uint32_t, 1, access::mode::read, access::target::device>
&buf) SYCL_ESIMD_FUNCTION {
simd<uint32_t, 32> offsets(0, 1);
simd<uint32_t, 32> v1(0, 1);

// CHECK: flat_atomic.cpp:99{{.*}}error: no matching function for call to 'atomic_update'
atomic_update<atomic_op::add, uint32_t, 32>(buf, offsets, v1, 1);
}

// Incompatible mode (read).
void kernel11(accessor<uint32_t, 1, access::mode::read, access::target::device>
&buf) SYCL_ESIMD_FUNCTION {
simd<uint32_t, 32> offsets(0, 1);
simd<uint32_t, 32> v1(0, 1);

// CHECK: flat_atomic.cpp:109{{.*}}error: no matching function for call to 'atomic_update'
atomic_update<atomic_op::cmpxchg, uint32_t, 32>(buf, offsets, v1, v1, 1);
}

// Incompatible mode (read).
void kernel12(local_accessor<const uint32_t, 1> &buf) SYCL_ESIMD_FUNCTION {
simd<uint32_t, 32> offsets(0, 1);

// CHECK: flat_atomic.cpp:117{{.*}}error: no matching function for call to 'atomic_update'
atomic_update<atomic_op::inc, uint32_t, 32>(buf, offsets, 1);
}

// Incompatible mode (read).
void kernel13(local_accessor<const uint32_t, 1> &buf) SYCL_ESIMD_FUNCTION {
simd<uint32_t, 32> offsets(0, 1);
simd<uint32_t, 32> v1(0, 1);

// CHECK: flat_atomic.cpp:126{{.*}}error: no matching function for call to 'atomic_update'
atomic_update<atomic_op::add, uint32_t, 32>(buf, offsets, v1, 1);
}

// Incompatible mode (read).
void kernel8(const local_accessor<const uint32_t, 1> &buf) SYCL_ESIMD_FUNCTION {
simd<uint32_t, 32> offsets(0, 1);
simd<uint32_t, 32> v1(0, 1);

// CHECK: flat_atomic.cpp:135{{.*}}error: no matching function for call to 'atomic_update'
atomic_update<atomic_op::cmpxchg, uint32_t, 32>(buf, offsets, v1, v1, 1);
}