Skip to content

[SYCL][ESIMD] Add compile time checks for lsc_atomic_update accessor based API #11849

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
Dec 6, 2023
Merged

[SYCL][ESIMD] Add compile time checks for lsc_atomic_update accessor based API #11849

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
d1b3d07
Add more stringent checks to lsc_atomic_update
fineg74 Nov 9, 2023
de820e6
Merge remote-tracking branch 'origin/sycl' into accessorCheck5
fineg74 Dec 2, 2023
8f0ebd6
Add more checks to the API
fineg74 Dec 4, 2023
4812489
Fix formatting issues
fineg74 Dec 4, 2023
b3bb871
Remove incorrect checks
fineg74 Dec 5, 2023
652db82
Address PR comments
fineg74 Dec 5, 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
10 changes: 10 additions & 0 deletions sycl/include/sycl/ext/intel/esimd/detail/sycl_util.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -105,6 +105,16 @@ inline constexpr bool is_accessor_with_v =
is_device_accessor_with_v<T, Capability> ||
is_local_accessor_with_v<T, Capability>;

template <typename T>
inline constexpr bool is_rw_device_accessor_v =
is_device_accessor_with_v<T, accessor_mode_cap::can_read> &&
is_device_accessor_with_v<T, accessor_mode_cap::can_write>;

template <typename T>
inline constexpr bool is_rw_local_accessor_v =
is_local_accessor_with_v<T, accessor_mode_cap::can_read> &&
is_local_accessor_with_v<T, accessor_mode_cap::can_write>;

template <typename T, accessor_mode_cap_val_t Capability, typename RetT>
using EnableIfAccessor =
std::enable_if_t<detail::is_device_accessor_with_v<T, Capability>, RetT>;
Expand Down
15 changes: 11 additions & 4 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 @@ -3755,6 +3755,7 @@ template <atomic_op Op, typename T, int N, lsc_data_size DS, cache_hint L1H,
__ESIMD_API std::enable_if_t<get_num_args<Op>() == 0, simd<T, N>>
atomic_update_impl(T *p, simd<Toffset, N> offsets, simd_mask<N> pred) {
static_assert(sizeof(T) > 1, "Unsupported data type");
static_assert(std::is_integral_v<Toffset>, "Unsupported offset type");
check_atomic<Op, T, N, 0, /*IsLSC*/ true>();
check_lsc_data_size<T, DS>();
check_cache_hint<cache_action::atomic, L1H, L2H>();
Expand Down Expand Up @@ -3795,6 +3796,7 @@ __ESIMD_API std::enable_if_t<get_num_args<Op>() == 1, simd<T, N>>
atomic_update_impl(T *p, simd<Toffset, N> offsets, simd<T, N> src0,
simd_mask<N> pred) {
static_assert(sizeof(T) > 1, "Unsupported data type");
static_assert(std::is_integral_v<Toffset>, "Unsupported offset type");
check_lsc_data_size<T, DS>();
check_atomic<Op, T, N, 1, /*IsLSC*/ true>();
check_cache_hint<cache_action::atomic, L1H, L2H>();
Expand Down Expand Up @@ -3837,6 +3839,7 @@ __ESIMD_API std::enable_if_t<get_num_args<Op>() == 2, simd<T, N>>
atomic_update_impl(T *p, simd<Toffset, N> offsets, simd<T, N> src0,
simd<T, N> src1, simd_mask<N> pred) {
static_assert(sizeof(T) > 1, "Unsupported data type");
static_assert(std::is_integral_v<Toffset>, "Unsupported offset type");
check_lsc_data_size<T, DS>();
check_atomic<Op, T, N, 2, /*IsLSC*/ true>();
check_cache_hint<cache_action::atomic, L1H, L2H>();
Expand Down Expand Up @@ -3880,8 +3883,10 @@ template <atomic_op Op, typename T, int N,
typename AccessorTy, typename Toffset>
__ESIMD_API std::enable_if_t<
get_num_args<Op>() == 0 &&
sycl::detail::acc_properties::is_accessor_v<AccessorTy> &&
!sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>,
__ESIMD_DNS::is_device_accessor_with_v<
AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read> &&
__ESIMD_DNS::is_device_accessor_with_v<
AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_write>,
simd<T, N>>
atomic_update_impl(AccessorTy acc, simd<Toffset, N> byte_offsets,
simd_mask<N> pred) {
Expand Down Expand Up @@ -3933,8 +3938,10 @@ template <atomic_op Op, typename T, int N, lsc_data_size DS, cache_hint L1H,
cache_hint L2H, typename AccessorTy, typename Toffset>
__ESIMD_API std::enable_if_t<
get_num_args<Op>() == 1 &&
sycl::detail::acc_properties::is_accessor_v<AccessorTy> &&
!sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>,
__ESIMD_DNS::is_device_accessor_with_v<
AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read> &&
__ESIMD_DNS::is_device_accessor_with_v<
AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_write>,
simd<T, N>>
atomic_update_impl(AccessorTy acc, simd<Toffset, N> byte_offset,
simd<T, N> src0, simd_mask<N> pred) {
Expand Down
93 changes: 38 additions & 55 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 @@ -2854,8 +2854,11 @@ template <__ESIMD_NS::atomic_op Op, typename T, int N,
cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
typename AccessorTy, typename Toffset>
__ESIMD_API std::enable_if_t<
sycl::detail::acc_properties::is_accessor_v<AccessorTy> &&
!sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>,
__ESIMD_DNS::is_device_accessor_with_v<
AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read> &&
(Op == __ESIMD_NS::atomic_op::load ||
__ESIMD_DNS::is_device_accessor_with_v<
AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_write>),
__ESIMD_NS::simd<T, N>>
lsc_atomic_update(AccessorTy acc, __ESIMD_NS::simd<Toffset, N> offsets,
__ESIMD_NS::simd_mask<N> pred) {
Expand All @@ -2882,9 +2885,8 @@ template <__ESIMD_NS::atomic_op Op, typename T, int N,
lsc_data_size DS = lsc_data_size::default_size,
cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
typename AccessorTy>
__ESIMD_API std::enable_if_t<
sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>,
__ESIMD_NS::simd<T, N>>
__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_rw_local_accessor_v<AccessorTy>,
__ESIMD_NS::simd<T, N>>
lsc_atomic_update(AccessorTy acc, __ESIMD_NS::simd<uint32_t, N> offsets,
__ESIMD_NS::simd_mask<N> pred) {
return lsc_slm_atomic_update<Op, T, N, DS>(
Expand Down Expand Up @@ -2913,10 +2915,8 @@ template <__ESIMD_NS::atomic_op Op, typename T, int N,
lsc_data_size DS = lsc_data_size::default_size,
cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
typename AccessorTy, typename Toffset>
__ESIMD_API std::enable_if_t<
sycl::detail::acc_properties::is_accessor_v<AccessorTy> &&
!sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>,
__ESIMD_NS::simd<T, N>>
__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_rw_device_accessor_v<AccessorTy>,
__ESIMD_NS::simd<T, N>>
lsc_atomic_update(AccessorTy acc, __ESIMD_NS::simd<Toffset, N> offsets,
__ESIMD_NS::simd<T, N> src0, __ESIMD_NS::simd_mask<N> pred) {
return __ESIMD_DNS::atomic_update_impl<Op, T, N, DS, L1H, L3H>(acc, offsets,
Expand All @@ -2943,9 +2943,8 @@ template <__ESIMD_NS::atomic_op Op, typename T, int N,
lsc_data_size DS = lsc_data_size::default_size,
cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
typename AccessorTy>
__ESIMD_API std::enable_if_t<
sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>,
__ESIMD_NS::simd<T, N>>
__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_rw_local_accessor_v<AccessorTy>,
__ESIMD_NS::simd<T, N>>
lsc_atomic_update(AccessorTy acc, __ESIMD_NS::simd<uint32_t, N> offsets,
__ESIMD_NS::simd<T, N> src0, __ESIMD_NS::simd_mask<N> pred) {
return lsc_slm_atomic_update<Op, T, N, DS>(
Expand Down Expand Up @@ -2975,10 +2974,8 @@ template <__ESIMD_NS::atomic_op Op, typename T, int N,
lsc_data_size DS = lsc_data_size::default_size,
cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
typename AccessorTy, typename Toffset>
__ESIMD_API std::enable_if_t<
sycl::detail::acc_properties::is_accessor_v<AccessorTy> &&
!sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>,
__ESIMD_NS::simd<T, N>>
__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_rw_device_accessor_v<AccessorTy>,
__ESIMD_NS::simd<T, N>>
lsc_atomic_update(AccessorTy acc, __ESIMD_NS::simd<Toffset, N> offsets,
__ESIMD_NS::simd<T, N> src0, __ESIMD_NS::simd<T, N> src1,
__ESIMD_NS::simd_mask<N> pred) {
Expand Down Expand Up @@ -3033,9 +3030,8 @@ template <__ESIMD_NS::atomic_op Op, typename T, int N,
lsc_data_size DS = lsc_data_size::default_size,
cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
typename AccessorTy>
__ESIMD_API std::enable_if_t<
sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>,
__ESIMD_NS::simd<T, N>>
__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_rw_local_accessor_v<AccessorTy>,
__ESIMD_NS::simd<T, N>>
lsc_atomic_update(AccessorTy acc, __ESIMD_NS::simd<uint32_t, N> offsets,
__ESIMD_NS::simd<T, N> src0, __ESIMD_NS::simd<T, N> src1,
__ESIMD_NS::simd_mask<N> pred) {
Expand Down Expand Up @@ -3112,9 +3108,7 @@ atomic_update(T *p, simd<Toffset, N> offset, simd_mask<N> mask) {

template <native::lsc::atomic_op Op, typename T, int N, typename Toffset,
typename RegionTy = __ESIMD_NS::region1d_t<Toffset, N, 1>>
__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> &&
__ESIMD_DNS::get_num_args<Op>() == 0,
simd<T, N>>
__ESIMD_API std::enable_if_t<__ESIMD_DNS::get_num_args<Op>() == 0, simd<T, N>>
atomic_update(T *p, simd_view<Toffset, RegionTy> offsets,
simd_mask<N> mask = 1) {
return __ESIMD_ENS::lsc_atomic_update<detail::to_atomic_op<Op>(), T, N>(
Expand All @@ -3132,24 +3126,20 @@ atomic_update(T *p, Toffset offset, simd_mask<N> mask = 1) {

/// LSC version of the single-argument atomic update.
template <native::lsc::atomic_op Op, typename T, int N, typename Toffset>
__ESIMD_API
__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> &&
__ESIMD_DNS::get_num_args<Op>() == 1,
simd<T, N>>
atomic_update(T *p, simd<Toffset, N> offset, simd<T, N> src0,
simd_mask<N> mask) {
__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> &&
__ESIMD_DNS::get_num_args<Op>() == 1,
simd<T, N>>
atomic_update(T *p, simd<Toffset, N> offset, simd<T, N> src0,
simd_mask<N> mask) {
return __ESIMD_ENS::lsc_atomic_update<detail::to_atomic_op<Op>(), T, N>(
p, offset, src0, mask);
}

template <native::lsc::atomic_op Op, typename T, int N, typename Toffset,
typename RegionTy = __ESIMD_NS::region1d_t<Toffset, N, 1>>
__ESIMD_API
__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> &&
__ESIMD_DNS::get_num_args<Op>() == 1,
simd<T, N>>
atomic_update(T *p, simd_view<Toffset, RegionTy> offsets, simd<T, N> src0,
simd_mask<N> mask = 1) {
__ESIMD_API std::enable_if_t<__ESIMD_DNS::get_num_args<Op>() == 1, simd<T, N>>
atomic_update(T *p, simd_view<Toffset, RegionTy> offsets, simd<T, N> src0,
simd_mask<N> mask = 1) {
return __ESIMD_ENS::lsc_atomic_update<detail::to_atomic_op<Op>(), T, N>(
p, offsets, src0, mask);
}
Expand Down Expand Up @@ -3179,9 +3169,7 @@ atomic_update(T *p, simd<Toffset, N> offset, simd<T, N> src0, simd<T, N> src1,

template <native::lsc::atomic_op Op, typename T, int N, typename Toffset,
typename RegionTy = __ESIMD_NS::region1d_t<Toffset, N, 1>>
__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> &&
__ESIMD_DNS::get_num_args<Op>() == 2,
simd<T, N>>
__ESIMD_API std::enable_if_t<__ESIMD_DNS::get_num_args<Op>() == 2, simd<T, N>>
atomic_update(T *p, simd_view<Toffset, RegionTy> offsets, simd<T, N> src0,
simd<T, N> src1, simd_mask<N> mask = 1) {
return __ESIMD_ENS::lsc_atomic_update<detail::to_atomic_op<Op>(), T, N>(
Expand Down Expand Up @@ -3212,8 +3200,7 @@ atomic_update(AccessorTy acc, simd<Toffset, N> offset, simd_mask<N> mask) {
template <native::lsc::atomic_op Op, typename T, int N, typename Toffset,
typename RegionTy = __ESIMD_NS::region1d_t<Toffset, N, 1>,
typename AccessorTy>
__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> &&
__ESIMD_DNS::get_num_args<Op>() == 0 &&
__ESIMD_API std::enable_if_t<__ESIMD_DNS::get_num_args<Op>() == 0 &&
!std::is_pointer_v<AccessorTy>,
simd<T, N>>
atomic_update(AccessorTy acc, simd_view<Toffset, RegionTy> offsets,
Expand All @@ -3236,27 +3223,24 @@ atomic_update(AccessorTy acc, Toffset offset, simd_mask<N> mask) {
/// LSC version of the single-argument atomic update.
template <native::lsc::atomic_op Op, typename T, int N, typename Toffset,
typename AccessorTy>
__ESIMD_API
__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> &&
__ESIMD_DNS::get_num_args<Op>() == 1 &&
!std::is_pointer_v<AccessorTy>,
simd<T, N>>
atomic_update(AccessorTy acc, simd<Toffset, N> offset, simd<T, N> src0,
simd_mask<N> mask) {
__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> &&
__ESIMD_DNS::get_num_args<Op>() == 1 &&
!std::is_pointer_v<AccessorTy>,
simd<T, N>>
atomic_update(AccessorTy acc, simd<Toffset, N> offset, simd<T, N> src0,
simd_mask<N> mask) {
return __ESIMD_ENS::lsc_atomic_update<detail::to_atomic_op<Op>(), T, N>(
acc, offset, src0, mask);
}

template <native::lsc::atomic_op Op, typename T, int N, typename Toffset,
typename RegionTy = __ESIMD_NS::region1d_t<Toffset, N, 1>,
typename AccessorTy>
__ESIMD_API
__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> &&
__ESIMD_DNS::get_num_args<Op>() == 1 &&
!std::is_pointer_v<AccessorTy>,
simd<T, N>>
atomic_update(AccessorTy acc, simd_view<Toffset, RegionTy> offsets,
simd<T, N> src0, simd_mask<N> mask) {
__ESIMD_API std::enable_if_t<__ESIMD_DNS::get_num_args<Op>() == 1 &&
!std::is_pointer_v<AccessorTy>,
simd<T, N>>
atomic_update(AccessorTy acc, simd_view<Toffset, RegionTy> offsets,
simd<T, N> src0, simd_mask<N> mask) {
return __ESIMD_ENS::lsc_atomic_update<detail::to_atomic_op<Op>(), T, N>(
acc, offsets, src0, mask);
}
Expand Down Expand Up @@ -3292,8 +3276,7 @@ atomic_update(AccessorTy acc, simd<Toffset, N> offset, simd<T, N> src0,
template <native::lsc::atomic_op Op, typename T, int N, typename Toffset,
typename RegionTy = __ESIMD_NS::region1d_t<Toffset, N, 1>,
typename AccessorTy>
__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> &&
__ESIMD_DNS::get_num_args<Op>() == 2 &&
__ESIMD_API std::enable_if_t<__ESIMD_DNS::get_num_args<Op>() == 2 &&
!std::is_pointer_v<AccessorTy>,
simd<T, N>>
atomic_update(AccessorTy acc, simd_view<Toffset, RegionTy> offsets,
Expand Down
131 changes: 131 additions & 0 deletions sycl/test/esimd/lsc_atomic.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,131 @@
// 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 lsc atomic APIs.

#include <limits>
#include <sycl/ext/intel/esimd.hpp>
#include <sycl/sycl.hpp>
#include <utility>
using namespace sycl::ext::intel::esimd;
using namespace sycl::ext::intel::experimental::esimd;
using namespace sycl;

// --- Postive tests.

void kernel0(uint32_t *ptr) SYCL_ESIMD_FUNCTION {
simd<uint32_t, 32> offsets(0, 1);
lsc_atomic_update<atomic_op::inc, uint32_t, 32>(ptr, offsets, 1);
}
void kernel1(uint32_t *ptr) SYCL_ESIMD_FUNCTION {
simd<uint32_t, 32> offsets(0, 1);
simd<uint32_t, 32> v1(0, 1);
lsc_atomic_update<atomic_op::add, uint32_t, 32>(ptr, offsets, v1, 1);
}
template <class T> void kernel2(T *ptr) SYCL_ESIMD_FUNCTION {
simd<uint32_t, 32> offsets(0, 1);
simd<T, 32> v1(0, 1);
lsc_atomic_update<atomic_op::cmpxchg, T, 32>(ptr, offsets, v1, v1, 1);
}

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);

lsc_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);

lsc_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);

lsc_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);

lsc_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);

lsc_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);

lsc_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: lsc_atomic.cpp:84{{.*}}error: no matching function for call to 'lsc_atomic_update'
lsc_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: lsc_atomic.cpp:94{{.*}}error: no matching function for call to 'lsc_atomic_update'
lsc_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: lsc_atomic.cpp:104{{.*}}error: no matching function for call to 'lsc_atomic_update'
lsc_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: lsc_atomic.cpp:112{{.*}}error: no matching function for call to 'lsc_atomic_update'
lsc_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: lsc_atomic.cpp:121{{.*}}error: no matching function for call to 'lsc_atomic_update'
lsc_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: lsc_atomic.cpp:130{{.*}}error: no matching function for call to 'lsc_atomic_update'
lsc_atomic_update<atomic_op::cmpxchg, uint32_t, 32>(buf, offsets, v1, v1, 1);
}