Skip to content

Commit 8ceefe5

Browse files
fineg74fineg74
authored
[SYCL][ESIMD] Add compile time checks for lsc_atomic_update accessor based API (#11849)
1 parent b1adf03 commit 8ceefe5

File tree

4 files changed

+190
-59
lines changed

4 files changed

+190
-59
lines changed

‎sycl/include/sycl/ext/intel/esimd/detail/sycl_util.hpp‎

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -105,6 +105,16 @@ inline constexpr bool is_accessor_with_v =
105105
is_device_accessor_with_v<T, Capability> ||
106106
is_local_accessor_with_v<T, Capability>;
107107

108+
template <typename T>
109+
inline constexpr bool is_rw_device_accessor_v =
110+
is_device_accessor_with_v<T, accessor_mode_cap::can_read> &&
111+
is_device_accessor_with_v<T, accessor_mode_cap::can_write>;
112+
113+
template <typename T>
114+
inline constexpr bool is_rw_local_accessor_v =
115+
is_local_accessor_with_v<T, accessor_mode_cap::can_read> &&
116+
is_local_accessor_with_v<T, accessor_mode_cap::can_write>;
117+
108118
template <typename T, accessor_mode_cap_val_t Capability, typename RetT>
109119
using EnableIfAccessor =
110120
std::enable_if_t<detail::is_device_accessor_with_v<T, Capability>, RetT>;

‎sycl/include/sycl/ext/intel/esimd/memory.hpp‎

Lines changed: 11 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -3755,6 +3755,7 @@ template <atomic_op Op, typename T, int N, lsc_data_size DS, cache_hint L1H,
37553755
__ESIMD_API std::enable_if_t<get_num_args<Op>() == 0, simd<T, N>>
37563756
atomic_update_impl(T *p, simd<Toffset, N> offsets, simd_mask<N> pred) {
37573757
static_assert(sizeof(T) > 1, "Unsupported data type");
3758+
static_assert(std::is_integral_v<Toffset>, "Unsupported offset type");
37583759
check_atomic<Op, T, N, 0, /*IsLSC*/ true>();
37593760
check_lsc_data_size<T, DS>();
37603761
check_cache_hint<cache_action::atomic, L1H, L2H>();
@@ -3795,6 +3796,7 @@ __ESIMD_API std::enable_if_t<get_num_args<Op>() == 1, simd<T, N>>
37953796
atomic_update_impl(T *p, simd<Toffset, N> offsets, simd<T, N> src0,
37963797
simd_mask<N> pred) {
37973798
static_assert(sizeof(T) > 1, "Unsupported data type");
3799+
static_assert(std::is_integral_v<Toffset>, "Unsupported offset type");
37983800
check_lsc_data_size<T, DS>();
37993801
check_atomic<Op, T, N, 1, /*IsLSC*/ true>();
38003802
check_cache_hint<cache_action::atomic, L1H, L2H>();
@@ -3837,6 +3839,7 @@ __ESIMD_API std::enable_if_t<get_num_args<Op>() == 2, simd<T, N>>
38373839
atomic_update_impl(T *p, simd<Toffset, N> offsets, simd<T, N> src0,
38383840
simd<T, N> src1, simd_mask<N> pred) {
38393841
static_assert(sizeof(T) > 1, "Unsupported data type");
3842+
static_assert(std::is_integral_v<Toffset>, "Unsupported offset type");
38403843
check_lsc_data_size<T, DS>();
38413844
check_atomic<Op, T, N, 2, /*IsLSC*/ true>();
38423845
check_cache_hint<cache_action::atomic, L1H, L2H>();
@@ -3880,8 +3883,10 @@ template <atomic_op Op, typename T, int N,
38803883
typename AccessorTy, typename Toffset>
38813884
__ESIMD_API std::enable_if_t<
38823885
get_num_args<Op>() == 0 &&
3883-
sycl::detail::acc_properties::is_accessor_v<AccessorTy> &&
3884-
!sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>,
3886+
__ESIMD_DNS::is_device_accessor_with_v<
3887+
AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read> &&
3888+
__ESIMD_DNS::is_device_accessor_with_v<
3889+
AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_write>,
38853890
simd<T, N>>
38863891
atomic_update_impl(AccessorTy acc, simd<Toffset, N> byte_offsets,
38873892
simd_mask<N> pred) {
@@ -3933,8 +3938,10 @@ template <atomic_op Op, typename T, int N, lsc_data_size DS, cache_hint L1H,
39333938
cache_hint L2H, typename AccessorTy, typename Toffset>
39343939
__ESIMD_API std::enable_if_t<
39353940
get_num_args<Op>() == 1 &&
3936-
sycl::detail::acc_properties::is_accessor_v<AccessorTy> &&
3937-
!sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>,
3941+
__ESIMD_DNS::is_device_accessor_with_v<
3942+
AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read> &&
3943+
__ESIMD_DNS::is_device_accessor_with_v<
3944+
AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_write>,
39383945
simd<T, N>>
39393946
atomic_update_impl(AccessorTy acc, simd<Toffset, N> byte_offset,
39403947
simd<T, N> src0, simd_mask<N> pred) {

‎sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp‎

Lines changed: 38 additions & 55 deletions
Original file line numberDiff line numberDiff line change
@@ -2873,8 +2873,11 @@ template <__ESIMD_NS::atomic_op Op, typename T, int N,
28732873
cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
28742874
typename AccessorTy, typename Toffset>
28752875
__ESIMD_API std::enable_if_t<
2876-
sycl::detail::acc_properties::is_accessor_v<AccessorTy> &&
2877-
!sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>,
2876+
__ESIMD_DNS::is_device_accessor_with_v<
2877+
AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read> &&
2878+
(Op == __ESIMD_NS::atomic_op::load ||
2879+
__ESIMD_DNS::is_device_accessor_with_v<
2880+
AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_write>),
28782881
__ESIMD_NS::simd<T, N>>
28792882
lsc_atomic_update(AccessorTy acc, __ESIMD_NS::simd<Toffset, N> offsets,
28802883
__ESIMD_NS::simd_mask<N> pred) {
@@ -2901,9 +2904,8 @@ template <__ESIMD_NS::atomic_op Op, typename T, int N,
29012904
lsc_data_size DS = lsc_data_size::default_size,
29022905
cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
29032906
typename AccessorTy>
2904-
__ESIMD_API std::enable_if_t<
2905-
sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>,
2906-
__ESIMD_NS::simd<T, N>>
2907+
__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_rw_local_accessor_v<AccessorTy>,
2908+
__ESIMD_NS::simd<T, N>>
29072909
lsc_atomic_update(AccessorTy acc, __ESIMD_NS::simd<uint32_t, N> offsets,
29082910
__ESIMD_NS::simd_mask<N> pred) {
29092911
return lsc_slm_atomic_update<Op, T, N, DS>(
@@ -2932,10 +2934,8 @@ template <__ESIMD_NS::atomic_op Op, typename T, int N,
29322934
lsc_data_size DS = lsc_data_size::default_size,
29332935
cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
29342936
typename AccessorTy, typename Toffset>
2935-
__ESIMD_API std::enable_if_t<
2936-
sycl::detail::acc_properties::is_accessor_v<AccessorTy> &&
2937-
!sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>,
2938-
__ESIMD_NS::simd<T, N>>
2937+
__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_rw_device_accessor_v<AccessorTy>,
2938+
__ESIMD_NS::simd<T, N>>
29392939
lsc_atomic_update(AccessorTy acc, __ESIMD_NS::simd<Toffset, N> offsets,
29402940
__ESIMD_NS::simd<T, N> src0, __ESIMD_NS::simd_mask<N> pred) {
29412941
return __ESIMD_DNS::atomic_update_impl<Op, T, N, DS, L1H, L3H>(acc, offsets,
@@ -2962,9 +2962,8 @@ template <__ESIMD_NS::atomic_op Op, typename T, int N,
29622962
lsc_data_size DS = lsc_data_size::default_size,
29632963
cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
29642964
typename AccessorTy>
2965-
__ESIMD_API std::enable_if_t<
2966-
sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>,
2967-
__ESIMD_NS::simd<T, N>>
2965+
__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_rw_local_accessor_v<AccessorTy>,
2966+
__ESIMD_NS::simd<T, N>>
29682967
lsc_atomic_update(AccessorTy acc, __ESIMD_NS::simd<uint32_t, N> offsets,
29692968
__ESIMD_NS::simd<T, N> src0, __ESIMD_NS::simd_mask<N> pred) {
29702969
return lsc_slm_atomic_update<Op, T, N, DS>(
@@ -2994,10 +2993,8 @@ template <__ESIMD_NS::atomic_op Op, typename T, int N,
29942993
lsc_data_size DS = lsc_data_size::default_size,
29952994
cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
29962995
typename AccessorTy, typename Toffset>
2997-
__ESIMD_API std::enable_if_t<
2998-
sycl::detail::acc_properties::is_accessor_v<AccessorTy> &&
2999-
!sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>,
3000-
__ESIMD_NS::simd<T, N>>
2996+
__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_rw_device_accessor_v<AccessorTy>,
2997+
__ESIMD_NS::simd<T, N>>
30012998
lsc_atomic_update(AccessorTy acc, __ESIMD_NS::simd<Toffset, N> offsets,
30022999
__ESIMD_NS::simd<T, N> src0, __ESIMD_NS::simd<T, N> src1,
30033000
__ESIMD_NS::simd_mask<N> pred) {
@@ -3052,9 +3049,8 @@ template <__ESIMD_NS::atomic_op Op, typename T, int N,
30523049
lsc_data_size DS = lsc_data_size::default_size,
30533050
cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
30543051
typename AccessorTy>
3055-
__ESIMD_API std::enable_if_t<
3056-
sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>,
3057-
__ESIMD_NS::simd<T, N>>
3052+
__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_rw_local_accessor_v<AccessorTy>,
3053+
__ESIMD_NS::simd<T, N>>
30583054
lsc_atomic_update(AccessorTy acc, __ESIMD_NS::simd<uint32_t, N> offsets,
30593055
__ESIMD_NS::simd<T, N> src0, __ESIMD_NS::simd<T, N> src1,
30603056
__ESIMD_NS::simd_mask<N> pred) {
@@ -3131,9 +3127,7 @@ atomic_update(T *p, simd<Toffset, N> offset, simd_mask<N> mask) {
31313127

31323128
template <native::lsc::atomic_op Op, typename T, int N, typename Toffset,
31333129
typename RegionTy = __ESIMD_NS::region1d_t<Toffset, N, 1>>
3134-
__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> &&
3135-
__ESIMD_DNS::get_num_args<Op>() == 0,
3136-
simd<T, N>>
3130+
__ESIMD_API std::enable_if_t<__ESIMD_DNS::get_num_args<Op>() == 0, simd<T, N>>
31373131
atomic_update(T *p, simd_view<Toffset, RegionTy> offsets,
31383132
simd_mask<N> mask = 1) {
31393133
return __ESIMD_ENS::lsc_atomic_update<detail::to_atomic_op<Op>(), T, N>(
@@ -3151,24 +3145,20 @@ atomic_update(T *p, Toffset offset, simd_mask<N> mask = 1) {
31513145

31523146
/// LSC version of the single-argument atomic update.
31533147
template <native::lsc::atomic_op Op, typename T, int N, typename Toffset>
3154-
__ESIMD_API
3155-
__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> &&
3156-
__ESIMD_DNS::get_num_args<Op>() == 1,
3157-
simd<T, N>>
3158-
atomic_update(T *p, simd<Toffset, N> offset, simd<T, N> src0,
3159-
simd_mask<N> mask) {
3148+
__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> &&
3149+
__ESIMD_DNS::get_num_args<Op>() == 1,
3150+
simd<T, N>>
3151+
atomic_update(T *p, simd<Toffset, N> offset, simd<T, N> src0,
3152+
simd_mask<N> mask) {
31603153
return __ESIMD_ENS::lsc_atomic_update<detail::to_atomic_op<Op>(), T, N>(
31613154
p, offset, src0, mask);
31623155
}
31633156

31643157
template <native::lsc::atomic_op Op, typename T, int N, typename Toffset,
31653158
typename RegionTy = __ESIMD_NS::region1d_t<Toffset, N, 1>>
3166-
__ESIMD_API
3167-
__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> &&
3168-
__ESIMD_DNS::get_num_args<Op>() == 1,
3169-
simd<T, N>>
3170-
atomic_update(T *p, simd_view<Toffset, RegionTy> offsets, simd<T, N> src0,
3171-
simd_mask<N> mask = 1) {
3159+
__ESIMD_API std::enable_if_t<__ESIMD_DNS::get_num_args<Op>() == 1, simd<T, N>>
3160+
atomic_update(T *p, simd_view<Toffset, RegionTy> offsets, simd<T, N> src0,
3161+
simd_mask<N> mask = 1) {
31723162
return __ESIMD_ENS::lsc_atomic_update<detail::to_atomic_op<Op>(), T, N>(
31733163
p, offsets, src0, mask);
31743164
}
@@ -3198,9 +3188,7 @@ atomic_update(T *p, simd<Toffset, N> offset, simd<T, N> src0, simd<T, N> src1,
31983188

31993189
template <native::lsc::atomic_op Op, typename T, int N, typename Toffset,
32003190
typename RegionTy = __ESIMD_NS::region1d_t<Toffset, N, 1>>
3201-
__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> &&
3202-
__ESIMD_DNS::get_num_args<Op>() == 2,
3203-
simd<T, N>>
3191+
__ESIMD_API std::enable_if_t<__ESIMD_DNS::get_num_args<Op>() == 2, simd<T, N>>
32043192
atomic_update(T *p, simd_view<Toffset, RegionTy> offsets, simd<T, N> src0,
32053193
simd<T, N> src1, simd_mask<N> mask = 1) {
32063194
return __ESIMD_ENS::lsc_atomic_update<detail::to_atomic_op<Op>(), T, N>(
@@ -3231,8 +3219,7 @@ atomic_update(AccessorTy acc, simd<Toffset, N> offset, simd_mask<N> mask) {
32313219
template <native::lsc::atomic_op Op, typename T, int N, typename Toffset,
32323220
typename RegionTy = __ESIMD_NS::region1d_t<Toffset, N, 1>,
32333221
typename AccessorTy>
3234-
__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> &&
3235-
__ESIMD_DNS::get_num_args<Op>() == 0 &&
3222+
__ESIMD_API std::enable_if_t<__ESIMD_DNS::get_num_args<Op>() == 0 &&
32363223
!std::is_pointer_v<AccessorTy>,
32373224
simd<T, N>>
32383225
atomic_update(AccessorTy acc, simd_view<Toffset, RegionTy> offsets,
@@ -3255,27 +3242,24 @@ atomic_update(AccessorTy acc, Toffset offset, simd_mask<N> mask) {
32553242
/// LSC version of the single-argument atomic update.
32563243
template <native::lsc::atomic_op Op, typename T, int N, typename Toffset,
32573244
typename AccessorTy>
3258-
__ESIMD_API
3259-
__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> &&
3260-
__ESIMD_DNS::get_num_args<Op>() == 1 &&
3261-
!std::is_pointer_v<AccessorTy>,
3262-
simd<T, N>>
3263-
atomic_update(AccessorTy acc, simd<Toffset, N> offset, simd<T, N> src0,
3264-
simd_mask<N> mask) {
3245+
__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> &&
3246+
__ESIMD_DNS::get_num_args<Op>() == 1 &&
3247+
!std::is_pointer_v<AccessorTy>,
3248+
simd<T, N>>
3249+
atomic_update(AccessorTy acc, simd<Toffset, N> offset, simd<T, N> src0,
3250+
simd_mask<N> mask) {
32653251
return __ESIMD_ENS::lsc_atomic_update<detail::to_atomic_op<Op>(), T, N>(
32663252
acc, offset, src0, mask);
32673253
}
32683254

32693255
template <native::lsc::atomic_op Op, typename T, int N, typename Toffset,
32703256
typename RegionTy = __ESIMD_NS::region1d_t<Toffset, N, 1>,
32713257
typename AccessorTy>
3272-
__ESIMD_API
3273-
__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> &&
3274-
__ESIMD_DNS::get_num_args<Op>() == 1 &&
3275-
!std::is_pointer_v<AccessorTy>,
3276-
simd<T, N>>
3277-
atomic_update(AccessorTy acc, simd_view<Toffset, RegionTy> offsets,
3278-
simd<T, N> src0, simd_mask<N> mask) {
3258+
__ESIMD_API std::enable_if_t<__ESIMD_DNS::get_num_args<Op>() == 1 &&
3259+
!std::is_pointer_v<AccessorTy>,
3260+
simd<T, N>>
3261+
atomic_update(AccessorTy acc, simd_view<Toffset, RegionTy> offsets,
3262+
simd<T, N> src0, simd_mask<N> mask) {
32793263
return __ESIMD_ENS::lsc_atomic_update<detail::to_atomic_op<Op>(), T, N>(
32803264
acc, offsets, src0, mask);
32813265
}
@@ -3311,8 +3295,7 @@ atomic_update(AccessorTy acc, simd<Toffset, N> offset, simd<T, N> src0,
33113295
template <native::lsc::atomic_op Op, typename T, int N, typename Toffset,
33123296
typename RegionTy = __ESIMD_NS::region1d_t<Toffset, N, 1>,
33133297
typename AccessorTy>
3314-
__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> &&
3315-
__ESIMD_DNS::get_num_args<Op>() == 2 &&
3298+
__ESIMD_API std::enable_if_t<__ESIMD_DNS::get_num_args<Op>() == 2 &&
33163299
!std::is_pointer_v<AccessorTy>,
33173300
simd<T, N>>
33183301
atomic_update(AccessorTy acc, simd_view<Toffset, RegionTy> offsets,

‎sycl/test/esimd/lsc_atomic.cpp‎

Lines changed: 131 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,131 @@
1+
// 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:"
2+
// 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:"
3+
4+
// This test checks compilation of ESIMD lsc atomic APIs.
5+
6+
#include <limits>
7+
#include <sycl/ext/intel/esimd.hpp>
8+
#include <sycl/sycl.hpp>
9+
#include <utility>
10+
using namespace sycl::ext::intel::esimd;
11+
using namespace sycl::ext::intel::experimental::esimd;
12+
using namespace sycl;
13+
14+
// --- Postive tests.
15+
16+
void kernel0(uint32_t *ptr) SYCL_ESIMD_FUNCTION {
17+
simd<uint32_t, 32> offsets(0, 1);
18+
lsc_atomic_update<atomic_op::inc, uint32_t, 32>(ptr, offsets, 1);
19+
}
20+
void kernel1(uint32_t *ptr) SYCL_ESIMD_FUNCTION {
21+
simd<uint32_t, 32> offsets(0, 1);
22+
simd<uint32_t, 32> v1(0, 1);
23+
lsc_atomic_update<atomic_op::add, uint32_t, 32>(ptr, offsets, v1, 1);
24+
}
25+
template <class T> void kernel2(T *ptr) SYCL_ESIMD_FUNCTION {
26+
simd<uint32_t, 32> offsets(0, 1);
27+
simd<T, 32> v1(0, 1);
28+
lsc_atomic_update<atomic_op::cmpxchg, T, 32>(ptr, offsets, v1, v1, 1);
29+
}
30+
31+
template void kernel2<uint32_t>(uint32_t *) SYCL_ESIMD_FUNCTION;
32+
33+
void kernel3(accessor<uint32_t, 1, access::mode::read_write,
34+
access::target::device> &buf) SYCL_ESIMD_FUNCTION {
35+
simd<uint32_t, 32> offsets(0, 1);
36+
37+
lsc_atomic_update<atomic_op::inc, uint32_t, 32>(buf, offsets, 1);
38+
}
39+
40+
void kernel4(accessor<uint32_t, 1, access::mode::read_write,
41+
access::target::device> &buf) SYCL_ESIMD_FUNCTION {
42+
simd<uint32_t, 32> offsets(0, 1);
43+
simd<uint32_t, 32> v1(0, 1);
44+
45+
lsc_atomic_update<atomic_op::add, uint32_t, 32>(buf, offsets, v1, 1);
46+
}
47+
48+
void kernel5(accessor<uint32_t, 1, access::mode::read_write,
49+
access::target::device> &buf) SYCL_ESIMD_FUNCTION {
50+
simd<uint32_t, 32> offsets(0, 1);
51+
simd<uint32_t, 32> v1(0, 1);
52+
53+
lsc_atomic_update<atomic_op::cmpxchg, uint32_t, 32>(buf, offsets, v1, v1, 1);
54+
}
55+
56+
void kernel6(local_accessor<uint32_t, 1> &buf) SYCL_ESIMD_FUNCTION {
57+
simd<uint32_t, 32> offsets(0, 1);
58+
59+
lsc_atomic_update<atomic_op::inc, uint32_t, 32>(buf, offsets, 1);
60+
}
61+
62+
void kernel7(local_accessor<uint32_t, 1> &buf) SYCL_ESIMD_FUNCTION {
63+
simd<uint32_t, 32> offsets(0, 1);
64+
simd<uint32_t, 32> v1(0, 1);
65+
66+
lsc_atomic_update<atomic_op::add, uint32_t, 32>(buf, offsets, v1, 1);
67+
}
68+
69+
void kernel8(local_accessor<uint32_t, 1> &buf) SYCL_ESIMD_FUNCTION {
70+
simd<uint32_t, 32> offsets(0, 1);
71+
simd<uint32_t, 32> v1(0, 1);
72+
73+
lsc_atomic_update<atomic_op::cmpxchg, uint32_t, 32>(buf, offsets, v1, v1, 1);
74+
}
75+
76+
// --- Negative tests.
77+
78+
// Incompatible mode (read).
79+
void kernel9(accessor<uint32_t, 1, access::mode::read, access::target::device>
80+
&buf) SYCL_ESIMD_FUNCTION {
81+
simd<uint32_t, 32> offsets(0, 1);
82+
83+
// CHECK: lsc_atomic.cpp:84{{.*}}error: no matching function for call to 'lsc_atomic_update'
84+
lsc_atomic_update<atomic_op::inc, uint32_t, 32>(buf, offsets, 1);
85+
}
86+
87+
// Incompatible mode (read).
88+
void kernel10(accessor<uint32_t, 1, access::mode::read, access::target::device>
89+
&buf) SYCL_ESIMD_FUNCTION {
90+
simd<uint32_t, 32> offsets(0, 1);
91+
simd<uint32_t, 32> v1(0, 1);
92+
93+
// CHECK: lsc_atomic.cpp:94{{.*}}error: no matching function for call to 'lsc_atomic_update'
94+
lsc_atomic_update<atomic_op::add, uint32_t, 32>(buf, offsets, v1, 1);
95+
}
96+
97+
// Incompatible mode (read).
98+
void kernel11(accessor<uint32_t, 1, access::mode::read, access::target::device>
99+
&buf) SYCL_ESIMD_FUNCTION {
100+
simd<uint32_t, 32> offsets(0, 1);
101+
simd<uint32_t, 32> v1(0, 1);
102+
103+
// CHECK: lsc_atomic.cpp:104{{.*}}error: no matching function for call to 'lsc_atomic_update'
104+
lsc_atomic_update<atomic_op::cmpxchg, uint32_t, 32>(buf, offsets, v1, v1, 1);
105+
}
106+
107+
// Incompatible mode (read).
108+
void kernel12(local_accessor<const uint32_t, 1> &buf) SYCL_ESIMD_FUNCTION {
109+
simd<uint32_t, 32> offsets(0, 1);
110+
111+
// CHECK: lsc_atomic.cpp:112{{.*}}error: no matching function for call to 'lsc_atomic_update'
112+
lsc_atomic_update<atomic_op::inc, uint32_t, 32>(buf, offsets, 1);
113+
}
114+
115+
// Incompatible mode (read).
116+
void kernel13(local_accessor<const uint32_t, 1> &buf) SYCL_ESIMD_FUNCTION {
117+
simd<uint32_t, 32> offsets(0, 1);
118+
simd<uint32_t, 32> v1(0, 1);
119+
120+
// CHECK: lsc_atomic.cpp:121{{.*}}error: no matching function for call to 'lsc_atomic_update'
121+
lsc_atomic_update<atomic_op::add, uint32_t, 32>(buf, offsets, v1, 1);
122+
}
123+
124+
// Incompatible mode (read).
125+
void kernel8(const local_accessor<const uint32_t, 1> &buf) SYCL_ESIMD_FUNCTION {
126+
simd<uint32_t, 32> offsets(0, 1);
127+
simd<uint32_t, 32> v1(0, 1);
128+
129+
// CHECK: lsc_atomic.cpp:130{{.*}}error: no matching function for call to 'lsc_atomic_update'
130+
lsc_atomic_update<atomic_op::cmpxchg, uint32_t, 32>(buf, offsets, v1, v1, 1);
131+
}

0 commit comments

Comments
 (0)