Skip to content
This repository was archived by the owner on Mar 28, 2023. It is now read-only.

Commit ccef0d4

Browse files
AidanBeltonSAidanBeltonS
authored
[SYCL] Update tests to use local_accessor (#1063)
This PR updates tests that use `target::local` to local_accessor. In all cases the change should not functionally change the test. The goal is to move from the deprecated `target::local` accessor to `local_accessor`. Depends on: intel/llvm#6341
1 parent 483fb75 commit ccef0d4

29 files changed

+68
-77
lines changed

‎SYCL/AtomicRef/add.h

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -30,8 +30,7 @@ void add_fetch_local_test(queue q, size_t N) {
3030
auto sum = sum_buf.template get_access<access::mode::read_write>(cgh);
3131
auto out =
3232
output_buf.template get_access<access::mode::discard_write>(cgh);
33-
accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
34-
cgh);
33+
local_accessor<T, 1> loc(1, cgh);
3534

3635
cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
3736
int gid = it.get_global_id(0);

‎SYCL/AtomicRef/and.h

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -30,8 +30,7 @@ void and_local_test(queue q) {
3030
auto cum = cum_buf.template get_access<access::mode::read_write>(cgh);
3131
auto out =
3232
output_buf.template get_access<access::mode::discard_write>(cgh);
33-
accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
34-
cgh);
33+
local_accessor<T, 1> loc(1, cgh);
3534

3635
cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
3736
int gid = it.get_global_id(0);

‎SYCL/AtomicRef/atomic_memory_order_acq_rel.cpp

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -76,8 +76,7 @@ template <memory_order order> void test_acquire_local() {
7676
q.submit([&](handler &cgh) {
7777
auto error =
7878
error_buf.template get_access<access::mode::read_write>(cgh);
79-
accessor<int, 1, access::mode::read_write, access::target::local> val(
80-
2, cgh);
79+
local_accessor<int, 1> val(2, cgh);
8180
cgh.parallel_for(
8281
nd_range<1>(global_size, local_size), [=](nd_item<1> it) {
8382
size_t lid = it.get_local_id(0);
@@ -168,8 +167,7 @@ template <memory_order order> void test_release_local() {
168167
q.submit([&](handler &cgh) {
169168
auto error =
170169
error_buf.template get_access<access::mode::read_write>(cgh);
171-
accessor<int, 1, access::mode::read_write, access::target::local> val(
172-
2, cgh);
170+
local_accessor<int, 1> val(2, cgh);
173171
cgh.parallel_for(
174172
nd_range<1>(global_size, local_size), [=](nd_item<1> it) {
175173
size_t lid = it.get_local_id(0);

‎SYCL/AtomicRef/atomic_memory_order_seq_cst.cpp

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -120,8 +120,7 @@ template <memory_order order> void test_local() {
120120

121121
q.submit([&](handler &cgh) {
122122
auto res = res_buf.template get_access<access::mode::discard_write>(cgh);
123-
accessor<int, 1, access::mode::read_write, access::target::local> val(2,
124-
cgh);
123+
local_accessor<int, 1> val(2, cgh);
125124
cgh.parallel_for(nd_range<1>(N_items, N_items), [=](nd_item<1> it) {
126125
val[0] = 0;
127126
it.barrier(access::fence_space::local_space);

‎SYCL/AtomicRef/compare_exchange.h

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -32,8 +32,7 @@ void compare_exchange_local_test(queue q, size_t N) {
3232
cgh);
3333
auto out =
3434
output_buf.template get_access<access::mode::discard_write>(cgh);
35-
accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
36-
cgh);
35+
local_accessor<T, 1> loc(1, cgh);
3736

3837
cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
3938
int gid = it.get_global_id(0);

‎SYCL/AtomicRef/exchange.h

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -30,8 +30,7 @@ void exchange_local_test(queue q, size_t N) {
3030
auto cum = cum_buf.template get_access<access::mode::read_write>(cgh);
3131
auto out =
3232
output_buf.template get_access<access::mode::discard_write>(cgh);
33-
accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
34-
cgh);
33+
local_accessor<T, 1> loc(1, cgh);
3534

3635
cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
3736
int gid = it.get_global_id(0);

‎SYCL/AtomicRef/load.h

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -31,8 +31,7 @@ void load_local_test(queue q, size_t N) {
3131
auto ld = load_buf.template get_access<access::mode::read_write>(cgh);
3232
auto out =
3333
output_buf.template get_access<access::mode::discard_write>(cgh);
34-
accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
35-
cgh);
34+
local_accessor<T, 1> loc(1, cgh);
3635
cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
3736
int gid = it.get_global_id(0);
3837
if (gid == 0)

‎SYCL/AtomicRef/max.h

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -30,8 +30,7 @@ void max_local_test(queue q, size_t N) {
3030
auto cum = cum_buf.template get_access<access::mode::read_write>(cgh);
3131
auto out =
3232
output_buf.template get_access<access::mode::discard_write>(cgh);
33-
accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
34-
cgh);
33+
local_accessor<T, 1> loc(1, cgh);
3534

3635
cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
3736
int gid = it.get_global_id(0);

‎SYCL/AtomicRef/min.h

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -30,8 +30,7 @@ void min_local_test(queue q, size_t N) {
3030
auto cum = cum_buf.template get_access<access::mode::read_write>(cgh);
3131
auto out =
3232
output_buf.template get_access<access::mode::discard_write>(cgh);
33-
accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
34-
cgh);
33+
local_accessor<T, 1> loc(1, cgh);
3534

3635
cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
3736
int gid = it.get_global_id(0);

‎SYCL/AtomicRef/or.h

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -30,8 +30,7 @@ void or_local_test(queue q) {
3030
auto cum = cum_buf.template get_access<access::mode::read_write>(cgh);
3131
auto out =
3232
output_buf.template get_access<access::mode::discard_write>(cgh);
33-
accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
34-
cgh);
33+
local_accessor<T, 1> loc(1, cgh);
3534

3635
cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
3736
int gid = it.get_global_id(0);

‎SYCL/AtomicRef/store.h

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -51,8 +51,7 @@ void store_local_test(queue q, size_t N) {
5151
buffer<T> store_buf(&store, 1);
5252
q.submit([&](handler &cgh) {
5353
auto st = store_buf.template get_access<access::mode::read_write>(cgh);
54-
accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
55-
cgh);
54+
local_accessor<T, 1> loc(1, cgh);
5655
cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
5756
size_t gid = it.get_global_id(0);
5857
auto atm = AtomicRef<T, memory_order::relaxed, scope, space>(loc[0]);

‎SYCL/AtomicRef/sub.h

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -30,8 +30,7 @@ void sub_fetch_local_test(queue q, size_t N) {
3030
auto sum = sum_buf.template get_access<access::mode::read_write>(cgh);
3131
auto out =
3232
output_buf.template get_access<access::mode::discard_write>(cgh);
33-
accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
34-
cgh);
33+
local_accessor<T, 1> loc(1, cgh);
3534

3635
cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
3736
int gid = it.get_global_id(0);

‎SYCL/AtomicRef/xor.h

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -30,8 +30,7 @@ void xor_local_test(queue q) {
3030
auto cum = cum_buf.template get_access<access::mode::read_write>(cgh);
3131
auto out =
3232
output_buf.template get_access<access::mode::discard_write>(cgh);
33-
accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
34-
cgh);
33+
local_accessor<T, 1> loc(1, cgh);
3534

3635
cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
3736
int gid = it.get_global_id(0);

‎SYCL/Basic/device_event.cpp

Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -5,6 +5,13 @@
55
// TODO: nd_item::barrier() is not implemented on HOST
66
// RUNx: %HOST_RUN_PLACEHOLDER %t.run
77
//
8+
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DUSE_DEPRECATED_LOCAL_ACC %s -o %t.run
9+
// RUN: %GPU_RUN_PLACEHOLDER %t.run
10+
// RUN: %CPU_RUN_PLACEHOLDER %t.run
11+
// RUN: %ACC_RUN_PLACEHOLDER %t.run
12+
// TODO: nd_item::barrier() is not implemented on HOST
13+
// RUNx: %HOST_RUN_PLACEHOLDER %t.run
14+
//
815
// Returns error "Barrier is not supported on the host device
916
// yet." with Nvidia.
1017
// XFAIL: hip_nvidia
@@ -76,8 +83,12 @@ int test_strideN(size_t stride) {
7683

7784
myQueue.submit([&](handler &cgh) {
7885
auto out_ptr = out_buf.get_access<access::mode::write>(cgh);
86+
#ifdef USE_DEPRECATED_LOCAL_ACC
7987
accessor<sycl::cl_int, 1, access::mode::read_write, access::target::local>
8088
local_acc(range<1>(16), cgh);
89+
#else
90+
local_accessor<sycl::cl_int, 1> local_acc(range<1>(16), cgh);
91+
#endif
8192

8293
// Create work-groups with 16 work items in each group.
8394
auto myRange = nd_range<1>(range<1>(nElems), range<1>(workGroupSize));

‎SYCL/Basic/group_async_copy.cpp

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -112,8 +112,7 @@ template <typename T> int test(size_t Stride) {
112112
Q.submit([&](handler &CGH) {
113113
auto In = InBuf.template get_access<access::mode::read>(CGH);
114114
auto Out = OutBuf.template get_access<access::mode::write>(CGH);
115-
accessor<T, 1, access::mode::read_write, access::target::local> Local(
116-
range<1>{WorkGroupSize}, CGH);
115+
local_accessor<T, 1> Local(range<1>{WorkGroupSize}, CGH);
117116

118117
nd_range<1> NDR{range<1>(NElems), range<1>(WorkGroupSize)};
119118
CGH.parallel_for<KernelName<T>>(NDR, [=](nd_item<1> NDId) {

‎SYCL/Basic/multi_ptr.cpp

Lines changed: 2 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -71,8 +71,7 @@ template <typename T> void testMultPtr() {
7171
accessor<T, 1, access::mode::read_write, access::target::device,
7272
access::placeholder::false_t>
7373
accessorData_2(bufferData_2, cgh);
74-
accessor<T, 1, access::mode::read_write, access::target::local>
75-
localAccessor(numOfItems, cgh);
74+
local_accessor<T, 1> localAccessor(numOfItems, cgh);
7675

7776
cgh.parallel_for<class testMultPtrKernel<T>>(range<1>{10}, [=](id<1> wiID) {
7877
auto ptr_1 = make_ptr<T, access::address_space::global_space>(
@@ -136,9 +135,7 @@ template <typename T> void testMultPtrArrowOperator() {
136135
accessor<point<T>, 1, access::mode::read, access::target::constant_buffer,
137136
access::placeholder::false_t>
138137
accessorData_2(bufferData_2, cgh);
139-
accessor<point<T>, 1, access::mode::read_write, access::target::local,
140-
access::placeholder::false_t>
141-
accessorData_3(1, cgh);
138+
local_accessor<point<T>, 1> accessorData_3(1, cgh);
142139
accessor<point<T>, 1, access::mode::read, access::target::device,
143140
access::placeholder::false_t>
144141
accessorData_4(bufferData_4, cgh);

‎SYCL/DeviceLib/ITTAnnotations/barrier.cpp

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -24,8 +24,7 @@ int main() {
2424
// ITT start/finish annotations and ITT wg_barrier/wi_resume annotations.
2525
q.submit([&](handler &cgh) {
2626
auto acc = buf.get_access<access::mode::read_write>(cgh);
27-
accessor<int, 1, access::mode::read_write, access::target::local>
28-
local_acc(local_range, cgh);
27+
local_accessor<int, 1> local_acc(local_range, cgh);
2928
cgh.parallel_for<class simple_barrier_kernel>(
3029
nd_range<1>(num_items, local_range), [=](nd_item<1> item) {
3130
size_t idx = item.get_global_linear_id();

‎SYCL/DeviceLib/string_test.cpp

Lines changed: 1 addition & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -395,10 +395,7 @@ bool kernel_test_memcpy_addr_space(sycl::queue &deviceQueue) {
395395
sycl::access::placeholder::false_t>
396396
src_acc(buffer1, cgh);
397397

398-
sycl::accessor<char, 1, sycl::access::mode::read_write,
399-
sycl::access::target::local,
400-
sycl::access::placeholder::false_t>
401-
local_acc(sycl::range<1>(16), cgh);
398+
sycl::local_accessor<char, 1> local_acc(sycl::range<1>(16), cgh);
402399

403400
sycl::accessor<char, 1, sycl::access::mode::write,
404401
sycl::access::target::device,

‎SYCL/DiscardEvents/discard_events_accessors.cpp

Lines changed: 1 addition & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -60,10 +60,7 @@ int main(int Argc, const char *Argv[]) {
6060
RunKernelHelper(Q, [&](int *Harray) {
6161
Q.submit([&](sycl::handler &CGH) {
6262
const size_t LocalMemSize = BUFFER_SIZE;
63-
using LocalAccessor =
64-
sycl::accessor<int, 1, sycl::access::mode::read_write,
65-
sycl::access::target::local>;
66-
LocalAccessor LocalAcc(LocalMemSize, CGH);
63+
sycl::local_accessor<int, 1> LocalAcc(LocalMemSize, CGH);
6764

6865
CGH.parallel_for<class kernel_using_local_memory>(
6966
Range, [=](sycl::item<1> itemID) {

‎SYCL/GroupAlgorithm/SYCL2020/sort.cpp

Lines changed: 2 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -112,9 +112,7 @@ int test_sort_over_group(sycl::queue &q, std::size_t local,
112112
<< std::endl;
113113
q.submit([&](sycl::handler &h) {
114114
auto aI1 = sycl::accessor(bufI1, h);
115-
sycl::accessor<std::byte, 1, sycl::access_mode::read_write,
116-
sycl::access::target::local>
117-
scratch({local_memory_size}, h);
115+
sycl::local_accessor<std::byte, 1> scratch({local_memory_size}, h);
118116

119117
h.parallel_for<sort_over_group_kernel_name<int_wrapper<dim>, T, Compare>>(
120118
sycl::nd_range<dim>(local_range, local_range),
@@ -167,9 +165,7 @@ int test_joint_sort(sycl::queue &q, std::size_t n_items, std::size_t local,
167165
<< std::endl;
168166
q.submit([&](sycl::handler &h) {
169167
auto aI1 = sycl::accessor(bufI1, h);
170-
sycl::accessor<std::byte, 1, sycl::access_mode::read_write,
171-
sycl::access::target::local>
172-
scratch({local_memory_size}, h);
168+
sycl::local_accessor<std::byte, 1> scratch({local_memory_size}, h);
173169

174170
h.parallel_for<joint_sort_kernel_name<T, Compare>>(
175171
sycl::nd_range<1>{{n_groups * local}, {local}},

‎SYCL/GroupAlgorithm/barrier.cpp

Lines changed: 4 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -22,10 +22,8 @@ void basic() {
2222

2323
q.submit([&](handler &cgh) {
2424
auto acc = buf.get_access<access::mode::read_write>(cgh);
25-
accessor<int, 1, access::mode::read_write, access::target::local> loc(
26-
N, cgh);
27-
accessor<barrier, 1, access::mode::read_write, access::target::local>
28-
loc_barrier(2, cgh);
25+
local_accessor<int, 1> loc(N, cgh);
26+
local_accessor<barrier, 1> loc_barrier(2, cgh);
2927
cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> item) {
3028
size_t idx = item.get_local_linear_id();
3129
loc[idx] = acc[idx];
@@ -69,10 +67,8 @@ void interface() {
6967
auto data_acc = data_buf.get_access<access::mode::read_write>(cgh);
7068
auto test1_acc = test1_buf.get_access<access::mode::read_write>(cgh);
7169
auto test2_acc = test2_buf.get_access<access::mode::read_write>(cgh);
72-
accessor<int, 1, access::mode::read_write, access::target::local> loc(
73-
N, cgh);
74-
accessor<barrier, 1, access::mode::read_write, access::target::local>
75-
loc_barrier(2, cgh);
70+
local_accessor<int, 1> loc(N, cgh);
71+
local_accessor<barrier, 1> loc_barrier(2, cgh);
7672
cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> item) {
7773
size_t idx = item.get_local_linear_id();
7874
if (idx == 0) {

‎SYCL/Regression/group.cpp

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -187,9 +187,7 @@ bool group__async_work_group_copy() {
187187

188188
Q.submit([&](handler &cgh) {
189189
auto AccGlobal = Buf.get_access<access::mode::read_write>(cgh);
190-
accessor<DataType, DIMS, access::mode::read_write,
191-
access::target::local>
192-
AccLocal(LocalRange, cgh);
190+
local_accessor<DataType, DIMS> AccLocal(LocalRange, cgh);
193191

194192
cgh.parallel_for<class group__async_work_group_copy>(
195193
nd_range<2>{GlobalRange, LocalRange}, [=](nd_item<DIMS> I) {

‎SYCL/Regression/local-arg-align.cpp

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -26,10 +26,8 @@ int main(int argc, char *argv[]) {
2626

2727
q.submit([&](sycl::handler &h) {
2828
// Use two local buffers, one with an int and one with a double4
29-
accessor<cl_int, 1, access::mode::read_write, access::target::local> a(1,
30-
h);
31-
accessor<double4, 1, access::mode::read_write, access::target::local> b(1,
32-
h);
29+
local_accessor<cl_int, 1> a(1, h);
30+
local_accessor<double4, 1> b(1, h);
3331

3432
auto ares = res.get_access<access::mode::read_write>(h);
3533

‎SYCL/Regression/zero_size_local_accessor.cpp

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -17,9 +17,7 @@
1717
int main() {
1818
sycl::queue Q;
1919
Q.submit([&](sycl::handler &CGH) {
20-
sycl::accessor<uint8_t, 1, sycl::access_mode::read_write,
21-
sycl::access::target::local>
22-
ZeroSizeLocalAcc(sycl::range<1>(0), CGH);
20+
sycl::local_accessor<uint8_t, 1> ZeroSizeLocalAcc(sycl::range<1>(0), CGH);
2321
CGH.single_task([=]() {
2422
if (ZeroSizeLocalAcc.get_range()[0])
2523
ZeroSizeLocalAcc[0] = 1;

‎SYCL/SubGroup/load_store.cpp

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -44,8 +44,7 @@ template <typename T, int N> void check(queue &Queue) {
4444
Queue.submit([&](handler &cgh) {
4545
auto acc = syclbuf.template get_access<access::mode::read_write>(cgh);
4646
auto sgsizeacc = sgsizebuf.get_access<access::mode::read_write>(cgh);
47-
accessor<T, 1, access::mode::read_write, access::target::local> LocalMem(
48-
{L + max_sg_size * N}, cgh);
47+
local_accessor<T, 1> LocalMem({L + max_sg_size * N}, cgh);
4948
cgh.parallel_for<sycl_subgr<T, N>>(NdRange, [=](nd_item<1> NdItem) {
5049
ext::oneapi::sub_group SG = NdItem.get_sub_group();
5150
auto SGid = SG.get_group_id().get(0);
@@ -132,8 +131,7 @@ template <typename T> void check(queue &Queue) {
132131
Queue.submit([&](handler &cgh) {
133132
auto acc = syclbuf.template get_access<access::mode::read_write>(cgh);
134133
auto sgsizeacc = sgsizebuf.get_access<access::mode::read_write>(cgh);
135-
accessor<T, 1, access::mode::read_write, access::target::local> LocalMem(
136-
{L}, cgh);
134+
local_accessor<T, 1> LocalMem({L}, cgh);
137135
cgh.parallel_for<sycl_subgr<T, 0>>(NdRange, [=](nd_item<1> NdItem) {
138136
ext::oneapi::sub_group SG = NdItem.get_sub_group();
139137
if (NdItem.get_global_id(0) == 0)

‎SYCL/SubGroup/sub_group_as.cpp

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4,6 +4,12 @@
44
// RUN: %CPU_RUN_PLACEHOLDER %t.out
55
// RUN: %ACC_RUN_PLACEHOLDER %t.out
66
//
7+
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DUSE_DEPRECATED_LOCAL_ACC %s -o %t.out
8+
// Sub-groups are not suported on Host
9+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
10+
// RUN: %CPU_RUN_PLACEHOLDER %t.out
11+
// RUN: %ACC_RUN_PLACEHOLDER %t.out
12+
//
713
// Missing __spirv_GenericCastToPtrExplicit_ToLocal,
814
// __spirv_SubgroupInvocationId, __spirv_GenericCastToPtrExplicit_ToGlobal,
915
// __spirv_SubgroupBlockReadINTEL, __assert_fail,
@@ -36,9 +42,13 @@ int main(int argc, char *argv[]) {
3642
queue.submit([&](sycl::handler &cgh) {
3743
auto global = buf.get_access<sycl::access::mode::read_write,
3844
sycl::access::target::device>(cgh);
45+
#ifdef USE_DEPRECATED_LOCAL_ACC
3946
sycl::accessor<int, 1, sycl::access::mode::read_write,
4047
sycl::access::target::local>
4148
local(N, cgh);
49+
#else
50+
sycl::local_accessor<int, 1> local(N, cgh);
51+
#endif
4252

4353
cgh.parallel_for<class test>(
4454
sycl::nd_range<1>(N, 32), [=](sycl::nd_item<1> it) {

‎SYCL/SubGroup/sub_group_as_vec.cpp

Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4,6 +4,13 @@
44
// RUN: %CPU_RUN_PLACEHOLDER %t.out
55
// RUN: %ACC_RUN_PLACEHOLDER %t.out
66
//
7+
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DUSE_DEPRECATED_LOCAL_ACC %s -o %t.out
8+
// Sub-groups are not suported on Host
9+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
10+
// RUN: %CPU_RUN_PLACEHOLDER %t.out
11+
// RUN: %ACC_RUN_PLACEHOLDER %t.out
12+
//
13+
//
714
// Missing __spirv_GenericCastToPtrExplicit_ToLocal,
815
// __spirv_SubgroupLocalInvocationId, __spirv_GenericCastToPtrExplicit_ToGlobal,
916
// __spirv_SubgroupBlockReadINTEL, __assert_fail,
@@ -38,9 +45,13 @@ int main(int argc, char *argv[]) {
3845
queue.submit([&](sycl::handler &cgh) {
3946
auto global = buf.get_access<sycl::access::mode::read_write,
4047
sycl::access::target::device>(cgh);
48+
#ifdef DUSE_DEPRECATED_LOCAL_ACC
4149
sycl::accessor<sycl::vec<int, 2>, 1, sycl::access::mode::read_write,
4250
sycl::access::target::local>
4351
local(N, cgh);
52+
#else
53+
sycl::local_accessor<sycl::vec<int, 2>, 1> local(N, cgh);
54+
#endif
4455
cgh.parallel_for<class test>(
4556
sycl::nd_range<1>(N, 32), [=](sycl::nd_item<1> it) {
4657
sycl::ext::oneapi::sub_group sg = it.get_sub_group();

0 commit comments

Comments
 (0)