Skip to content

Commit 1a40033

Browse files
fineg74fineg74
authored
[SYCL][ESIMD] Add compile time checks for lsc_block_load/store, lsc_gather/scatter and lsc_prefetch API (#11744)
1 parent e49a0c3 commit 1a40033

File tree

3 files changed

+202
-53
lines changed

3 files changed

+202
-53
lines changed

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

Lines changed: 72 additions & 53 deletions
Original file line numberDiff line numberDiff line change
@@ -818,13 +818,13 @@ template <typename T, int NElts = 1,
818818
lsc_data_size DS = lsc_data_size::default_size,
819819
cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
820820
int N, typename AccessorTy>
821-
__ESIMD_API std::enable_if_t<
822-
!std::is_pointer_v<AccessorTy> &&
823-
!sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>,
824-
__ESIMD_NS::simd<T, N * NElts>>
825-
lsc_gather(AccessorTy acc,
826-
__ESIMD_NS::simd<__ESIMD_DNS::DeviceAccessorOffsetT, N> offsets,
827-
__ESIMD_NS::simd_mask<N> pred = 1) {
821+
__ESIMD_API
822+
std::enable_if_t<__ESIMD_DNS::is_device_accessor_with_v<
823+
AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read>,
824+
__ESIMD_NS::simd<T, N * NElts>>
825+
lsc_gather(AccessorTy acc,
826+
__ESIMD_NS::simd<__ESIMD_DNS::DeviceAccessorOffsetT, N> offsets,
827+
__ESIMD_NS::simd_mask<N> pred = 1) {
828828
#ifdef __ESIMD_FORCE_STATELESS_MEM
829829
return lsc_gather<T, NElts, DS, L1H, L3H>(
830830
reinterpret_cast<T *>(acc.get_pointer().get()), offsets, pred);
@@ -854,8 +854,8 @@ template <typename T, int NElts = 1,
854854
cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
855855
int N, typename AccessorTy, typename Toffset>
856856
__ESIMD_API std::enable_if_t<
857-
!std::is_pointer_v<AccessorTy> &&
858-
!sycl::detail::acc_properties::is_local_accessor_v<AccessorTy> &&
857+
__ESIMD_DNS::is_device_accessor_with_v<
858+
AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read> &&
859859
std::is_integral_v<Toffset> && !std::is_same_v<Toffset, uint64_t>,
860860
__ESIMD_NS::simd<T, N * NElts>>
861861
lsc_gather(AccessorTy acc, __ESIMD_NS::simd<Toffset, N> offsets,
@@ -869,11 +869,12 @@ template <typename T, int NElts = 1,
869869
lsc_data_size DS = lsc_data_size::default_size,
870870
cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
871871
int N, typename AccessorTy>
872-
__ESIMD_API std::enable_if_t<
873-
sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>,
874-
__ESIMD_NS::simd<T, N * NElts>>
875-
lsc_gather(AccessorTy acc, __ESIMD_NS::simd<uint32_t, N> offsets,
876-
__ESIMD_NS::simd_mask<N> pred = 1) {
872+
__ESIMD_API
873+
std::enable_if_t<__ESIMD_DNS::is_local_accessor_with_v<
874+
AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read>,
875+
__ESIMD_NS::simd<T, N * NElts>>
876+
lsc_gather(AccessorTy acc, __ESIMD_NS::simd<uint32_t, N> offsets,
877+
__ESIMD_NS::simd_mask<N> pred = 1) {
877878
return lsc_slm_gather<T, NElts, DS>(
878879
offsets + __ESIMD_DNS::localAccessorToOffset(acc), pred);
879880
}
@@ -903,14 +904,14 @@ template <typename T, int NElts = 1,
903904
lsc_data_size DS = lsc_data_size::default_size,
904905
cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
905906
int N, typename AccessorTy>
906-
__ESIMD_API std::enable_if_t<
907-
!std::is_pointer_v<AccessorTy> &&
908-
!sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>,
909-
__ESIMD_NS::simd<T, N * NElts>>
910-
lsc_gather(AccessorTy acc,
911-
__ESIMD_NS::simd<__ESIMD_DNS::DeviceAccessorOffsetT, N> offsets,
912-
__ESIMD_NS::simd_mask<N> pred,
913-
__ESIMD_NS::simd<T, N * NElts> pass_thru) {
907+
__ESIMD_API
908+
std::enable_if_t<__ESIMD_DNS::is_device_accessor_with_v<
909+
AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read>,
910+
__ESIMD_NS::simd<T, N * NElts>>
911+
lsc_gather(AccessorTy acc,
912+
__ESIMD_NS::simd<__ESIMD_DNS::DeviceAccessorOffsetT, N> offsets,
913+
__ESIMD_NS::simd_mask<N> pred,
914+
__ESIMD_NS::simd<T, N * NElts> pass_thru) {
914915
#ifdef __ESIMD_FORCE_STATELESS_MEM
915916
return lsc_gather<T, NElts, DS, L1H, L3H>(
916917
reinterpret_cast<T *>(acc.get_pointer().get()), offsets, pred, pass_thru);
@@ -943,8 +944,8 @@ template <typename T, int NElts = 1,
943944
cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
944945
int N, typename AccessorTy, typename Toffset>
945946
__ESIMD_API std::enable_if_t<
946-
!std::is_pointer_v<AccessorTy> &&
947-
!sycl::detail::acc_properties::is_local_accessor_v<AccessorTy> &&
947+
__ESIMD_DNS::is_device_accessor_with_v<
948+
AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read> &&
948949
std::is_integral_v<Toffset> && !std::is_same_v<Toffset, uint64_t>,
949950
__ESIMD_NS::simd<T, N * NElts>>
950951
lsc_gather(AccessorTy acc, __ESIMD_NS::simd<Toffset, N> offsets,
@@ -1130,8 +1131,8 @@ template <typename T, int NElts, lsc_data_size DS = lsc_data_size::default_size,
11301131
typename AccessorTy,
11311132
typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
11321133
__ESIMD_API std::enable_if_t<
1133-
!std::is_pointer_v<AccessorTy> &&
1134-
!sycl::detail::acc_properties::is_local_accessor_v<AccessorTy> &&
1134+
__ESIMD_DNS::is_device_accessor_with_v<
1135+
AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read> &&
11351136
__ESIMD_NS::is_simd_flag_type_v<FlagsT>,
11361137
__ESIMD_NS::simd<T, NElts>>
11371138
lsc_block_load(AccessorTy acc, __ESIMD_DNS::DeviceAccessorOffsetT offset,
@@ -1145,7 +1146,8 @@ template <typename T, int NElts, lsc_data_size DS = lsc_data_size::default_size,
11451146
typename AccessorTy,
11461147
typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
11471148
__ESIMD_API std::enable_if_t<
1148-
sycl::detail::acc_properties::is_local_accessor_v<AccessorTy> &&
1149+
__ESIMD_DNS::is_local_accessor_with_v<
1150+
AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read> &&
11491151
__ESIMD_NS::is_simd_flag_type_v<FlagsT>,
11501152
__ESIMD_NS::simd<T, NElts>>
11511153
lsc_block_load(AccessorTy acc, uint32_t offset,
@@ -1185,15 +1187,31 @@ template <typename T, int NElts, lsc_data_size DS = lsc_data_size::default_size,
11851187
cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
11861188
typename AccessorTy,
11871189
typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
1188-
__ESIMD_API std::enable_if_t<!std::is_pointer_v<AccessorTy> &&
1189-
__ESIMD_NS::is_simd_flag_type_v<FlagsT>,
1190-
__ESIMD_NS::simd<T, NElts>>
1190+
__ESIMD_API std::enable_if_t<
1191+
__ESIMD_DNS::is_device_accessor_with_v<
1192+
AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read> &&
1193+
__ESIMD_NS::is_simd_flag_type_v<FlagsT>,
1194+
__ESIMD_NS::simd<T, NElts>>
11911195
lsc_block_load(AccessorTy acc, __ESIMD_DNS::DeviceAccessorOffsetT offset,
11921196
FlagsT flags) {
11931197
return lsc_block_load<T, NElts, DS, L1H, L3H>(
11941198
acc, offset, __ESIMD_NS::simd_mask<1>(1), flags);
11951199
}
11961200

1201+
template <typename T, int NElts, lsc_data_size DS = lsc_data_size::default_size,
1202+
cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
1203+
typename AccessorTy,
1204+
typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
1205+
__ESIMD_API std::enable_if_t<
1206+
__ESIMD_DNS::is_local_accessor_with_v<
1207+
AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read> &&
1208+
__ESIMD_NS::is_simd_flag_type_v<FlagsT>,
1209+
__ESIMD_NS::simd<T, NElts>>
1210+
lsc_block_load(AccessorTy acc, uint32_t offset, FlagsT flags) {
1211+
return lsc_block_load<T, NElts, DS, L1H, L3H>(
1212+
acc, offset, __ESIMD_NS::simd_mask<1>(1), flags);
1213+
}
1214+
11971215
/// Accessor-based transposed gather with 1 channel.
11981216
/// Supported platforms: DG2, PVC
11991217
/// VISA instruction: lsc_load.ugm
@@ -1231,8 +1249,8 @@ template <typename T, int NElts, lsc_data_size DS = lsc_data_size::default_size,
12311249
typename AccessorTy,
12321250
typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
12331251
__ESIMD_API std::enable_if_t<
1234-
!std::is_pointer_v<AccessorTy> &&
1235-
!sycl::detail::acc_properties::is_local_accessor_v<AccessorTy> &&
1252+
__ESIMD_DNS::is_device_accessor_with_v<
1253+
AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read> &&
12361254
__ESIMD_NS::is_simd_flag_type_v<FlagsT>,
12371255
__ESIMD_NS::simd<T, NElts>>
12381256
lsc_block_load(AccessorTy acc, __ESIMD_DNS::DeviceAccessorOffsetT offset,
@@ -1247,7 +1265,8 @@ template <typename T, int NElts, lsc_data_size DS = lsc_data_size::default_size,
12471265
typename AccessorTy,
12481266
typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
12491267
__ESIMD_API std::enable_if_t<
1250-
sycl::detail::acc_properties::is_local_accessor_v<AccessorTy> &&
1268+
__ESIMD_DNS::is_local_accessor_with_v<
1269+
AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read> &&
12511270
__ESIMD_NS::is_simd_flag_type_v<FlagsT>,
12521271
__ESIMD_NS::simd<T, NElts>>
12531272
lsc_block_load(AccessorTy acc, uint32_t offset, __ESIMD_NS::simd_mask<1> pred,
@@ -1377,9 +1396,8 @@ template <typename T, int NElts = 1,
13771396
lsc_data_size DS = lsc_data_size::default_size,
13781397
cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
13791398
int N, typename AccessorTy>
1380-
__ESIMD_API std::enable_if_t<
1381-
!std::is_pointer_v<AccessorTy> &&
1382-
!sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>>
1399+
__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_device_accessor_with_v<
1400+
AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read>>
13831401
lsc_prefetch(AccessorTy acc,
13841402
#ifdef __ESIMD_FORCE_STATELESS_MEM
13851403
__ESIMD_NS::simd<uint64_t, N> offsets,
@@ -1414,8 +1432,8 @@ template <typename T, int NElts = 1,
14141432
cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
14151433
int N, typename AccessorTy, typename Toffset>
14161434
__ESIMD_API std::enable_if_t<
1417-
!std::is_pointer_v<AccessorTy> &&
1418-
!sycl::detail::acc_properties::is_local_accessor_v<AccessorTy> &&
1435+
__ESIMD_DNS::is_device_accessor_with_v<
1436+
AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read> &&
14191437
std::is_integral_v<Toffset> && !std::is_same_v<Toffset, uint64_t>>
14201438
lsc_prefetch(AccessorTy acc, __ESIMD_NS::simd<Toffset, N> offsets,
14211439
__ESIMD_NS::simd_mask<N> pred = 1) {
@@ -1443,9 +1461,8 @@ template <typename T, int NElts = 1,
14431461
lsc_data_size DS = lsc_data_size::default_size,
14441462
cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
14451463
typename AccessorTy>
1446-
__ESIMD_API std::enable_if_t<
1447-
!std::is_pointer_v<AccessorTy> &&
1448-
!sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>>
1464+
__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_device_accessor_with_v<
1465+
AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read>>
14491466
lsc_prefetch(AccessorTy acc, __ESIMD_DNS::DeviceAccessorOffsetT offset) {
14501467
#ifdef __ESIMD_FORCE_STATELESS_MEM
14511468
lsc_prefetch<T, NElts, DS, L1H, L3H>(
@@ -1619,9 +1636,8 @@ template <typename T, int NElts = 1,
16191636
lsc_data_size DS = lsc_data_size::default_size,
16201637
cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
16211638
int N, typename AccessorTy>
1622-
__ESIMD_API std::enable_if_t<
1623-
!std::is_pointer_v<AccessorTy> &&
1624-
!sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>>
1639+
__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_device_accessor_with_v<
1640+
AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_write>>
16251641
lsc_scatter(AccessorTy acc,
16261642
__ESIMD_NS::simd<__ESIMD_DNS::DeviceAccessorOffsetT, N> offsets,
16271643
__ESIMD_NS::simd<T, N * NElts> vals,
@@ -1656,8 +1672,8 @@ template <typename T, int NElts = 1,
16561672
cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
16571673
int N, typename AccessorTy, typename Toffset>
16581674
__ESIMD_API std::enable_if_t<
1659-
!std::is_pointer_v<AccessorTy> &&
1660-
!sycl::detail::acc_properties::is_local_accessor_v<AccessorTy> &&
1675+
__ESIMD_DNS::is_device_accessor_with_v<
1676+
AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_write> &&
16611677
std::is_integral_v<Toffset> && !std::is_same_v<Toffset, uint64_t>>
16621678
lsc_scatter(AccessorTy acc, __ESIMD_NS::simd<Toffset, N> offsets,
16631679
__ESIMD_NS::simd<T, N * NElts> vals,
@@ -1671,8 +1687,8 @@ template <typename T, int NElts = 1,
16711687
lsc_data_size DS = lsc_data_size::default_size,
16721688
cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
16731689
int N, typename AccessorTy>
1674-
__ESIMD_API std::enable_if_t<
1675-
sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>>
1690+
__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_local_accessor_with_v<
1691+
AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_write>>
16761692
lsc_scatter(AccessorTy acc, __ESIMD_NS::simd<uint32_t, N> offsets,
16771693
__ESIMD_NS::simd<T, N * NElts> vals,
16781694
__ESIMD_NS::simd_mask<N> pred = 1) {
@@ -1798,8 +1814,8 @@ template <typename T, int NElts, lsc_data_size DS = lsc_data_size::default_size,
17981814
typename AccessorTy,
17991815
typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
18001816
__ESIMD_API std::enable_if_t<
1801-
!std::is_pointer_v<AccessorTy> &&
1802-
!sycl::detail::acc_properties::is_local_accessor_v<AccessorTy> &&
1817+
__ESIMD_DNS::is_device_accessor_with_v<
1818+
AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_write> &&
18031819
__ESIMD_NS::is_simd_flag_type_v<FlagsT>>
18041820
lsc_block_store(AccessorTy acc, __ESIMD_DNS::DeviceAccessorOffsetT offset,
18051821
__ESIMD_NS::simd<T, NElts> vals,
@@ -1813,7 +1829,8 @@ template <typename T, int NElts, lsc_data_size DS = lsc_data_size::default_size,
18131829
typename AccessorTy,
18141830
typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
18151831
__ESIMD_API std::enable_if_t<
1816-
sycl::detail::acc_properties::is_local_accessor_v<AccessorTy> &&
1832+
__ESIMD_DNS::is_local_accessor_with_v<
1833+
AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_write> &&
18171834
__ESIMD_NS::is_simd_flag_type_v<FlagsT>>
18181835
lsc_block_store(AccessorTy acc, uint32_t offset,
18191836
__ESIMD_NS::simd<T, NElts> vals, FlagsT flags = FlagsT{}) {
@@ -1855,8 +1872,10 @@ template <typename T, int NElts, lsc_data_size DS = lsc_data_size::default_size,
18551872
cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
18561873
typename AccessorTy,
18571874
typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
1858-
__ESIMD_API std::enable_if_t<!std::is_pointer_v<AccessorTy> &&
1859-
__ESIMD_NS::is_simd_flag_type_v<FlagsT>>
1875+
__ESIMD_API std::enable_if_t<
1876+
__ESIMD_DNS::is_accessor_with_v<
1877+
AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_write> &&
1878+
__ESIMD_NS::is_simd_flag_type_v<FlagsT>>
18601879
lsc_block_store(AccessorTy acc, __ESIMD_DNS::DeviceAccessorOffsetT offset,
18611880
__ESIMD_NS::simd<T, NElts> vals, FlagsT flags) {
18621881
lsc_block_store<T, NElts, DS, L1H, L3H>(acc, offset, vals,

‎sycl/test/esimd/lsc_block_load_store.cpp‎

Lines changed: 59 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,59 @@
1+
// RUN: not %clangxx -fsycl -fsyntax-only -Wno-unused-command-line-argument %s 2>&1 | FileCheck %s --implicit-check-not="warning:" --implicit-check-not="error:"
2+
3+
// This test checks that lsc_block_load/store API gets successfully compiled.
4+
5+
#include <limits>
6+
#include <sycl/ext/intel/esimd.hpp>
7+
#include <sycl/sycl.hpp>
8+
#include <utility>
9+
10+
using namespace sycl::ext::intel::esimd;
11+
using namespace sycl::ext::intel::experimental::esimd;
12+
using namespace sycl;
13+
14+
SYCL_EXTERNAL void
15+
kernel1(accessor<int, 1, access::mode::read_write, access::target::device> &buf)
16+
SYCL_ESIMD_FUNCTION {
17+
simd<int, 32> v1(0, 1);
18+
auto v0 = lsc_block_load<int, 32>(buf, 0);
19+
v0 = v0 + v1;
20+
lsc_block_store<int, 32>(buf, 0, v0);
21+
}
22+
23+
SYCL_EXTERNAL void kernel2(int *ptr) SYCL_ESIMD_FUNCTION {
24+
simd<int, 32> v1(0, 1);
25+
auto v0 = lsc_block_load<int, 32>(ptr);
26+
v0 = v0 + v1;
27+
lsc_block_store<int, 32>(ptr, v0);
28+
}
29+
30+
// --- Negative tests.
31+
32+
// Incompatible mode (write).
33+
SYCL_EXTERNAL void
34+
kernel4(accessor<int, 1, access::mode::write, access::target::device> &buf)
35+
SYCL_ESIMD_FUNCTION {
36+
simd<int, 32> v;
37+
// CHECK: lsc_block_load_store.cpp:39{{.*}}error: no matching function
38+
// function for call to 'lsc_block_load'
39+
v = lsc_block_load<int, 32>(buf, 0);
40+
}
41+
42+
// Incompatible mode (read).
43+
SYCL_EXTERNAL void
44+
kernel5(accessor<int, 1, access::mode::read, access::target::device> &buf)
45+
SYCL_ESIMD_FUNCTION {
46+
simd<int, 32> v(0, 1);
47+
// CHECK: lsc_block_load_store.cpp:49{{.*}}error: no matching function
48+
// function for call to 'lsc_block_store'
49+
lsc_block_store<int, 32>(buf, 0, v);
50+
}
51+
52+
// Incompatible mode (read).
53+
SYCL_EXTERNAL void
54+
kernel6(local_accessor<const int, 1> &buf) SYCL_ESIMD_FUNCTION {
55+
simd<int, 32> v(0, 1);
56+
// CHECK: lsc_block_load_store.cpp:58{{.*}}error: no matching function
57+
// function for call to 'lsc_block_store'
58+
lsc_block_store<int, 32>(buf, 0, v);
59+
}

‎sycl/test/esimd/lsc_gather_scatter.cpp‎

Lines changed: 71 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,71 @@
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 -fsycl-device-only -fsyntax-only -fsycl-esimd-force-stateless-mem -Wno-unused-command-line-argument %s 2>&1 | FileCheck %s --implicit-check-not="warning:" --implicit-check-not="error:"
3+
4+
#include <limits>
5+
#include <sycl/ext/intel/esimd.hpp>
6+
#include <sycl/sycl.hpp>
7+
#include <utility>
8+
9+
using namespace sycl::ext::intel::esimd;
10+
using namespace sycl::ext::intel::experimental::esimd;
11+
using namespace sycl;
12+
13+
void kernel(accessor<int, 1, access::mode::read_write, access::target::device>
14+
&buf) SYCL_ESIMD_FUNCTION {
15+
simd<uint32_t, 32> offsets(0, sizeof(int));
16+
simd<int, 32> v1(0, 1);
17+
18+
lsc_prefetch<uint32_t, 1, lsc_data_size::default_size, cache_hint::cached,
19+
cache_hint::cached>(buf, offsets);
20+
21+
auto v0 = lsc_gather<int>(buf, offsets);
22+
23+
v0 = v0 + v1;
24+
25+
lsc_scatter<int>(buf, offsets, v0);
26+
}
27+
28+
// --- Negative tests.
29+
30+
// Incompatible mode (write).
31+
SYCL_EXTERNAL void
32+
kernel2(accessor<int, 1, access::mode::write, access::target::device> &buf)
33+
SYCL_ESIMD_FUNCTION {
34+
simd<int, 32> v;
35+
simd<uint32_t, 32> offset(0, 1);
36+
// CHECK: lsc_gather_scatter.cpp:38{{.*}}error: no matching function
37+
// function for call to 'lsc_gather'
38+
v = lsc_gather<int>(buf, offset);
39+
}
40+
41+
// Incompatible mode (write).
42+
SYCL_EXTERNAL void
43+
kernel3(accessor<int, 1, access::mode::write, access::target::device> &buf)
44+
SYCL_ESIMD_FUNCTION {
45+
simd<uint32_t, 32> offset(0, 1);
46+
// CHECK: lsc_gather_scatter.cpp:48{{.*}}error: no matching function
47+
// function for call to 'lsc_prefetch'
48+
lsc_prefetch<int, 1, lsc_data_size::default_size, cache_hint::cached,
49+
cache_hint::cached>(buf, offset);
50+
}
51+
52+
// Incompatible mode (read).
53+
SYCL_EXTERNAL void
54+
kernel4(accessor<int, 1, access::mode::read, access::target::device> &buf)
55+
SYCL_ESIMD_FUNCTION {
56+
simd<int, 32> v(0, 1);
57+
simd<uint32_t, 32> offset(0, 1);
58+
// CHECK: lsc_gather_scatter.cpp:60{{.*}}error: no matching function
59+
// function for call to 'lsc_scatter'
60+
lsc_scatter<int>(buf, offset, v);
61+
}
62+
63+
// Incompatible mode (read).
64+
SYCL_EXTERNAL void
65+
kernel5(local_accessor<const int, 1> &buf) SYCL_ESIMD_FUNCTION {
66+
simd<int, 32> v(0, 1);
67+
simd<uint32_t, 32> offset(0, 1);
68+
// CHECK: lsc_gather_scatter.cpp:70{{.*}}error: no matching function
69+
// function for call to 'lsc_scatter'
70+
lsc_scatter<int>(buf, offset, v);
71+
}

0 commit comments

Comments
 (0)