diff --git a/SYCL/AtomicRef/add.h b/SYCL/AtomicRef/add.h index fe5deef656..48f2434220 100644 --- a/SYCL/AtomicRef/add.h +++ b/SYCL/AtomicRef/add.h @@ -30,8 +30,7 @@ void add_fetch_local_test(queue q, size_t N) { auto sum = sum_buf.template get_access(cgh); auto out = output_buf.template get_access(cgh); - accessor loc(1, - cgh); + local_accessor loc(1, cgh); cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) { int gid = it.get_global_id(0); diff --git a/SYCL/AtomicRef/and.h b/SYCL/AtomicRef/and.h index 888257c78c..5762125a40 100755 --- a/SYCL/AtomicRef/and.h +++ b/SYCL/AtomicRef/and.h @@ -30,8 +30,7 @@ void and_local_test(queue q) { auto cum = cum_buf.template get_access(cgh); auto out = output_buf.template get_access(cgh); - accessor loc(1, - cgh); + local_accessor loc(1, cgh); cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) { int gid = it.get_global_id(0); diff --git a/SYCL/AtomicRef/atomic_memory_order_acq_rel.cpp b/SYCL/AtomicRef/atomic_memory_order_acq_rel.cpp index 36e8cd8925..a02d7dd174 100644 --- a/SYCL/AtomicRef/atomic_memory_order_acq_rel.cpp +++ b/SYCL/AtomicRef/atomic_memory_order_acq_rel.cpp @@ -76,8 +76,7 @@ template void test_acquire_local() { q.submit([&](handler &cgh) { auto error = error_buf.template get_access(cgh); - accessor val( - 2, cgh); + local_accessor val(2, cgh); cgh.parallel_for( nd_range<1>(global_size, local_size), [=](nd_item<1> it) { size_t lid = it.get_local_id(0); @@ -168,8 +167,7 @@ template void test_release_local() { q.submit([&](handler &cgh) { auto error = error_buf.template get_access(cgh); - accessor val( - 2, cgh); + local_accessor val(2, cgh); cgh.parallel_for( nd_range<1>(global_size, local_size), [=](nd_item<1> it) { size_t lid = it.get_local_id(0); diff --git a/SYCL/AtomicRef/atomic_memory_order_seq_cst.cpp b/SYCL/AtomicRef/atomic_memory_order_seq_cst.cpp index 91ae2f8027..97ed6200dd 100755 --- a/SYCL/AtomicRef/atomic_memory_order_seq_cst.cpp +++ b/SYCL/AtomicRef/atomic_memory_order_seq_cst.cpp @@ -120,8 +120,7 @@ template void test_local() { q.submit([&](handler &cgh) { auto res = res_buf.template get_access(cgh); - accessor val(2, - cgh); + local_accessor val(2, cgh); cgh.parallel_for(nd_range<1>(N_items, N_items), [=](nd_item<1> it) { val[0] = 0; it.barrier(access::fence_space::local_space); diff --git a/SYCL/AtomicRef/compare_exchange.h b/SYCL/AtomicRef/compare_exchange.h index 4cded0f877..01bef9d567 100644 --- a/SYCL/AtomicRef/compare_exchange.h +++ b/SYCL/AtomicRef/compare_exchange.h @@ -32,8 +32,7 @@ void compare_exchange_local_test(queue q, size_t N) { cgh); auto out = output_buf.template get_access(cgh); - accessor loc(1, - cgh); + local_accessor loc(1, cgh); cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) { int gid = it.get_global_id(0); diff --git a/SYCL/AtomicRef/exchange.h b/SYCL/AtomicRef/exchange.h index c8f73b42ba..ae7d182c86 100644 --- a/SYCL/AtomicRef/exchange.h +++ b/SYCL/AtomicRef/exchange.h @@ -30,8 +30,7 @@ void exchange_local_test(queue q, size_t N) { auto cum = cum_buf.template get_access(cgh); auto out = output_buf.template get_access(cgh); - accessor loc(1, - cgh); + local_accessor loc(1, cgh); cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) { int gid = it.get_global_id(0); diff --git a/SYCL/AtomicRef/load.h b/SYCL/AtomicRef/load.h index cccc8138a2..f40cd60fdd 100644 --- a/SYCL/AtomicRef/load.h +++ b/SYCL/AtomicRef/load.h @@ -31,8 +31,7 @@ void load_local_test(queue q, size_t N) { auto ld = load_buf.template get_access(cgh); auto out = output_buf.template get_access(cgh); - accessor loc(1, - cgh); + local_accessor loc(1, cgh); cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) { int gid = it.get_global_id(0); if (gid == 0) diff --git a/SYCL/AtomicRef/max.h b/SYCL/AtomicRef/max.h index 235356276e..0eed88cac6 100644 --- a/SYCL/AtomicRef/max.h +++ b/SYCL/AtomicRef/max.h @@ -30,8 +30,7 @@ void max_local_test(queue q, size_t N) { auto cum = cum_buf.template get_access(cgh); auto out = output_buf.template get_access(cgh); - accessor loc(1, - cgh); + local_accessor loc(1, cgh); cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) { int gid = it.get_global_id(0); diff --git a/SYCL/AtomicRef/min.h b/SYCL/AtomicRef/min.h index c4f2c04fd2..5b1ff078ba 100644 --- a/SYCL/AtomicRef/min.h +++ b/SYCL/AtomicRef/min.h @@ -30,8 +30,7 @@ void min_local_test(queue q, size_t N) { auto cum = cum_buf.template get_access(cgh); auto out = output_buf.template get_access(cgh); - accessor loc(1, - cgh); + local_accessor loc(1, cgh); cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) { int gid = it.get_global_id(0); diff --git a/SYCL/AtomicRef/or.h b/SYCL/AtomicRef/or.h index ff887e0144..5395dba28d 100755 --- a/SYCL/AtomicRef/or.h +++ b/SYCL/AtomicRef/or.h @@ -30,8 +30,7 @@ void or_local_test(queue q) { auto cum = cum_buf.template get_access(cgh); auto out = output_buf.template get_access(cgh); - accessor loc(1, - cgh); + local_accessor loc(1, cgh); cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) { int gid = it.get_global_id(0); diff --git a/SYCL/AtomicRef/store.h b/SYCL/AtomicRef/store.h index 241baf9aa9..a2ea45aaab 100644 --- a/SYCL/AtomicRef/store.h +++ b/SYCL/AtomicRef/store.h @@ -51,8 +51,7 @@ void store_local_test(queue q, size_t N) { buffer store_buf(&store, 1); q.submit([&](handler &cgh) { auto st = store_buf.template get_access(cgh); - accessor loc(1, - cgh); + local_accessor loc(1, cgh); cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) { size_t gid = it.get_global_id(0); auto atm = AtomicRef(loc[0]); diff --git a/SYCL/AtomicRef/sub.h b/SYCL/AtomicRef/sub.h index 860fd4c326..2f0cab5f43 100644 --- a/SYCL/AtomicRef/sub.h +++ b/SYCL/AtomicRef/sub.h @@ -30,8 +30,7 @@ void sub_fetch_local_test(queue q, size_t N) { auto sum = sum_buf.template get_access(cgh); auto out = output_buf.template get_access(cgh); - accessor loc(1, - cgh); + local_accessor loc(1, cgh); cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) { int gid = it.get_global_id(0); diff --git a/SYCL/AtomicRef/xor.h b/SYCL/AtomicRef/xor.h index 556e269523..0d64dab878 100755 --- a/SYCL/AtomicRef/xor.h +++ b/SYCL/AtomicRef/xor.h @@ -30,8 +30,7 @@ void xor_local_test(queue q) { auto cum = cum_buf.template get_access(cgh); auto out = output_buf.template get_access(cgh); - accessor loc(1, - cgh); + local_accessor loc(1, cgh); cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) { int gid = it.get_global_id(0); diff --git a/SYCL/Basic/device_event.cpp b/SYCL/Basic/device_event.cpp index 6be9d45e00..c2bfcd5c3f 100644 --- a/SYCL/Basic/device_event.cpp +++ b/SYCL/Basic/device_event.cpp @@ -5,6 +5,13 @@ // TODO: nd_item::barrier() is not implemented on HOST // RUNx: %HOST_RUN_PLACEHOLDER %t.run // +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DUSE_DEPRECATED_LOCAL_ACC %s -o %t.run +// RUN: %GPU_RUN_PLACEHOLDER %t.run +// RUN: %CPU_RUN_PLACEHOLDER %t.run +// RUN: %ACC_RUN_PLACEHOLDER %t.run +// TODO: nd_item::barrier() is not implemented on HOST +// RUNx: %HOST_RUN_PLACEHOLDER %t.run +// // Returns error "Barrier is not supported on the host device // yet." with Nvidia. // XFAIL: hip_nvidia @@ -76,8 +83,12 @@ int test_strideN(size_t stride) { myQueue.submit([&](handler &cgh) { auto out_ptr = out_buf.get_access(cgh); +#ifdef USE_DEPRECATED_LOCAL_ACC accessor local_acc(range<1>(16), cgh); +#else + local_accessor local_acc(range<1>(16), cgh); +#endif // Create work-groups with 16 work items in each group. auto myRange = nd_range<1>(range<1>(nElems), range<1>(workGroupSize)); diff --git a/SYCL/Basic/group_async_copy.cpp b/SYCL/Basic/group_async_copy.cpp index 5b52d568a1..056774d60d 100644 --- a/SYCL/Basic/group_async_copy.cpp +++ b/SYCL/Basic/group_async_copy.cpp @@ -112,8 +112,7 @@ template int test(size_t Stride) { Q.submit([&](handler &CGH) { auto In = InBuf.template get_access(CGH); auto Out = OutBuf.template get_access(CGH); - accessor Local( - range<1>{WorkGroupSize}, CGH); + local_accessor Local(range<1>{WorkGroupSize}, CGH); nd_range<1> NDR{range<1>(NElems), range<1>(WorkGroupSize)}; CGH.parallel_for>(NDR, [=](nd_item<1> NDId) { diff --git a/SYCL/Basic/multi_ptr.cpp b/SYCL/Basic/multi_ptr.cpp index e1b05aefa0..adf34d20b4 100644 --- a/SYCL/Basic/multi_ptr.cpp +++ b/SYCL/Basic/multi_ptr.cpp @@ -71,8 +71,7 @@ template void testMultPtr() { accessor accessorData_2(bufferData_2, cgh); - accessor - localAccessor(numOfItems, cgh); + local_accessor localAccessor(numOfItems, cgh); cgh.parallel_for>(range<1>{10}, [=](id<1> wiID) { auto ptr_1 = make_ptr( @@ -136,9 +135,7 @@ template void testMultPtrArrowOperator() { accessor, 1, access::mode::read, access::target::constant_buffer, access::placeholder::false_t> accessorData_2(bufferData_2, cgh); - accessor, 1, access::mode::read_write, access::target::local, - access::placeholder::false_t> - accessorData_3(1, cgh); + local_accessor, 1> accessorData_3(1, cgh); accessor, 1, access::mode::read, access::target::device, access::placeholder::false_t> accessorData_4(bufferData_4, cgh); diff --git a/SYCL/DeviceLib/ITTAnnotations/barrier.cpp b/SYCL/DeviceLib/ITTAnnotations/barrier.cpp index c84fb74474..1f65b70290 100644 --- a/SYCL/DeviceLib/ITTAnnotations/barrier.cpp +++ b/SYCL/DeviceLib/ITTAnnotations/barrier.cpp @@ -24,8 +24,7 @@ int main() { // ITT start/finish annotations and ITT wg_barrier/wi_resume annotations. q.submit([&](handler &cgh) { auto acc = buf.get_access(cgh); - accessor - local_acc(local_range, cgh); + local_accessor local_acc(local_range, cgh); cgh.parallel_for( nd_range<1>(num_items, local_range), [=](nd_item<1> item) { size_t idx = item.get_global_linear_id(); diff --git a/SYCL/DeviceLib/string_test.cpp b/SYCL/DeviceLib/string_test.cpp index b35e21d6e5..d667dd1e33 100644 --- a/SYCL/DeviceLib/string_test.cpp +++ b/SYCL/DeviceLib/string_test.cpp @@ -395,10 +395,7 @@ bool kernel_test_memcpy_addr_space(sycl::queue &deviceQueue) { sycl::access::placeholder::false_t> src_acc(buffer1, cgh); - sycl::accessor - local_acc(sycl::range<1>(16), cgh); + sycl::local_accessor local_acc(sycl::range<1>(16), cgh); sycl::accessor; - LocalAccessor LocalAcc(LocalMemSize, CGH); + sycl::local_accessor LocalAcc(LocalMemSize, CGH); CGH.parallel_for( Range, [=](sycl::item<1> itemID) { diff --git a/SYCL/GroupAlgorithm/SYCL2020/sort.cpp b/SYCL/GroupAlgorithm/SYCL2020/sort.cpp index 85f31e78b0..b3e3875b2c 100644 --- a/SYCL/GroupAlgorithm/SYCL2020/sort.cpp +++ b/SYCL/GroupAlgorithm/SYCL2020/sort.cpp @@ -112,9 +112,7 @@ int test_sort_over_group(sycl::queue &q, std::size_t local, << std::endl; q.submit([&](sycl::handler &h) { auto aI1 = sycl::accessor(bufI1, h); - sycl::accessor - scratch({local_memory_size}, h); + sycl::local_accessor scratch({local_memory_size}, h); h.parallel_for, T, Compare>>( sycl::nd_range(local_range, local_range), @@ -167,9 +165,7 @@ int test_joint_sort(sycl::queue &q, std::size_t n_items, std::size_t local, << std::endl; q.submit([&](sycl::handler &h) { auto aI1 = sycl::accessor(bufI1, h); - sycl::accessor - scratch({local_memory_size}, h); + sycl::local_accessor scratch({local_memory_size}, h); h.parallel_for>( sycl::nd_range<1>{{n_groups * local}, {local}}, diff --git a/SYCL/GroupAlgorithm/barrier.cpp b/SYCL/GroupAlgorithm/barrier.cpp index c4d9238acd..76b7d92543 100644 --- a/SYCL/GroupAlgorithm/barrier.cpp +++ b/SYCL/GroupAlgorithm/barrier.cpp @@ -22,10 +22,8 @@ void basic() { q.submit([&](handler &cgh) { auto acc = buf.get_access(cgh); - accessor loc( - N, cgh); - accessor - loc_barrier(2, cgh); + local_accessor loc(N, cgh); + local_accessor loc_barrier(2, cgh); cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> item) { size_t idx = item.get_local_linear_id(); loc[idx] = acc[idx]; @@ -69,10 +67,8 @@ void interface() { auto data_acc = data_buf.get_access(cgh); auto test1_acc = test1_buf.get_access(cgh); auto test2_acc = test2_buf.get_access(cgh); - accessor loc( - N, cgh); - accessor - loc_barrier(2, cgh); + local_accessor loc(N, cgh); + local_accessor loc_barrier(2, cgh); cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> item) { size_t idx = item.get_local_linear_id(); if (idx == 0) { diff --git a/SYCL/Regression/group.cpp b/SYCL/Regression/group.cpp index 605bdfc0e3..d0be836278 100644 --- a/SYCL/Regression/group.cpp +++ b/SYCL/Regression/group.cpp @@ -187,9 +187,7 @@ bool group__async_work_group_copy() { Q.submit([&](handler &cgh) { auto AccGlobal = Buf.get_access(cgh); - accessor - AccLocal(LocalRange, cgh); + local_accessor AccLocal(LocalRange, cgh); cgh.parallel_for( nd_range<2>{GlobalRange, LocalRange}, [=](nd_item I) { diff --git a/SYCL/Regression/local-arg-align.cpp b/SYCL/Regression/local-arg-align.cpp index 72c676a6c4..69f07b908a 100644 --- a/SYCL/Regression/local-arg-align.cpp +++ b/SYCL/Regression/local-arg-align.cpp @@ -26,10 +26,8 @@ int main(int argc, char *argv[]) { q.submit([&](sycl::handler &h) { // Use two local buffers, one with an int and one with a double4 - accessor a(1, - h); - accessor b(1, - h); + local_accessor a(1, h); + local_accessor b(1, h); auto ares = res.get_access(h); diff --git a/SYCL/Regression/zero_size_local_accessor.cpp b/SYCL/Regression/zero_size_local_accessor.cpp index be7e608e49..587644e368 100644 --- a/SYCL/Regression/zero_size_local_accessor.cpp +++ b/SYCL/Regression/zero_size_local_accessor.cpp @@ -17,9 +17,7 @@ int main() { sycl::queue Q; Q.submit([&](sycl::handler &CGH) { - sycl::accessor - ZeroSizeLocalAcc(sycl::range<1>(0), CGH); + sycl::local_accessor ZeroSizeLocalAcc(sycl::range<1>(0), CGH); CGH.single_task([=]() { if (ZeroSizeLocalAcc.get_range()[0]) ZeroSizeLocalAcc[0] = 1; diff --git a/SYCL/SubGroup/load_store.cpp b/SYCL/SubGroup/load_store.cpp index 854608dce7..2f0fb9b7ea 100644 --- a/SYCL/SubGroup/load_store.cpp +++ b/SYCL/SubGroup/load_store.cpp @@ -44,8 +44,7 @@ template void check(queue &Queue) { Queue.submit([&](handler &cgh) { auto acc = syclbuf.template get_access(cgh); auto sgsizeacc = sgsizebuf.get_access(cgh); - accessor LocalMem( - {L + max_sg_size * N}, cgh); + local_accessor LocalMem({L + max_sg_size * N}, cgh); cgh.parallel_for>(NdRange, [=](nd_item<1> NdItem) { ext::oneapi::sub_group SG = NdItem.get_sub_group(); auto SGid = SG.get_group_id().get(0); @@ -132,8 +131,7 @@ template void check(queue &Queue) { Queue.submit([&](handler &cgh) { auto acc = syclbuf.template get_access(cgh); auto sgsizeacc = sgsizebuf.get_access(cgh); - accessor LocalMem( - {L}, cgh); + local_accessor LocalMem({L}, cgh); cgh.parallel_for>(NdRange, [=](nd_item<1> NdItem) { ext::oneapi::sub_group SG = NdItem.get_sub_group(); if (NdItem.get_global_id(0) == 0) diff --git a/SYCL/SubGroup/sub_group_as.cpp b/SYCL/SubGroup/sub_group_as.cpp index 643ee8e44e..14dc1d8bcc 100644 --- a/SYCL/SubGroup/sub_group_as.cpp +++ b/SYCL/SubGroup/sub_group_as.cpp @@ -4,6 +4,12 @@ // RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %ACC_RUN_PLACEHOLDER %t.out // +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DUSE_DEPRECATED_LOCAL_ACC %s -o %t.out +// Sub-groups are not suported on Host +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %ACC_RUN_PLACEHOLDER %t.out +// // Missing __spirv_GenericCastToPtrExplicit_ToLocal, // __spirv_SubgroupInvocationId, __spirv_GenericCastToPtrExplicit_ToGlobal, // __spirv_SubgroupBlockReadINTEL, __assert_fail, @@ -36,9 +42,13 @@ int main(int argc, char *argv[]) { queue.submit([&](sycl::handler &cgh) { auto global = buf.get_access(cgh); +#ifdef USE_DEPRECATED_LOCAL_ACC sycl::accessor local(N, cgh); +#else + sycl::local_accessor local(N, cgh); +#endif cgh.parallel_for( sycl::nd_range<1>(N, 32), [=](sycl::nd_item<1> it) { diff --git a/SYCL/SubGroup/sub_group_as_vec.cpp b/SYCL/SubGroup/sub_group_as_vec.cpp index 93c8467089..3f6cd3c901 100644 --- a/SYCL/SubGroup/sub_group_as_vec.cpp +++ b/SYCL/SubGroup/sub_group_as_vec.cpp @@ -4,6 +4,13 @@ // RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %ACC_RUN_PLACEHOLDER %t.out // +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DUSE_DEPRECATED_LOCAL_ACC %s -o %t.out +// Sub-groups are not suported on Host +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %ACC_RUN_PLACEHOLDER %t.out +// +// // Missing __spirv_GenericCastToPtrExplicit_ToLocal, // __spirv_SubgroupLocalInvocationId, __spirv_GenericCastToPtrExplicit_ToGlobal, // __spirv_SubgroupBlockReadINTEL, __assert_fail, @@ -38,9 +45,13 @@ int main(int argc, char *argv[]) { queue.submit([&](sycl::handler &cgh) { auto global = buf.get_access(cgh); +#ifdef DUSE_DEPRECATED_LOCAL_ACC sycl::accessor, 1, sycl::access::mode::read_write, sycl::access::target::local> local(N, cgh); +#else + sycl::local_accessor, 1> local(N, cgh); +#endif cgh.parallel_for( sycl::nd_range<1>(N, 32), [=](sycl::nd_item<1> it) { sycl::ext::oneapi::sub_group sg = it.get_sub_group(); diff --git a/SYCL/XPTI/buffer/accessors.cpp b/SYCL/XPTI/buffer/accessors.cpp index c57f425e64..84b9dc7773 100644 --- a/SYCL/XPTI/buffer/accessors.cpp +++ b/SYCL/XPTI/buffer/accessors.cpp @@ -27,8 +27,8 @@ int main() { auto A1 = Buf.get_access(cgh); // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID2:.*]]|2014|1025|{{.*}}accessors.cpp:[[# @LINE + 1]]:15 auto A2 = Buf.get_access(cgh); - // CHECK: {{[0-9]+}}|Construct accessor|0x0|[[ACCID3:.*]]|2016|1026|{{.*}}accessors.cpp:[[# @LINE + 1]]:61 - sycl::accessor A3(Range, cgh); + // CHECK: {{[0-9]+}}|Construct accessor|0x0|[[ACCID3:.*]]|2016|1026|{{.*}}accessors.cpp:[[# @LINE + 1]]:34 + sycl::local_accessor A3(Range, cgh); // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID4:.*]]|2014|1027|{{.*}}accessors.cpp:[[# @LINE + 1]]:15 auto A4 = Buf.get_access(cgh); // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID5:.*]]|2014|1028|{{.*}}accessors.cpp:[[# @LINE + 1]]:15 diff --git a/SYCL/XPTI/kernel/basic.cpp b/SYCL/XPTI/kernel/basic.cpp index e9ce40bcde..e2283aee24 100644 --- a/SYCL/XPTI/kernel/basic.cpp +++ b/SYCL/XPTI/kernel/basic.cpp @@ -56,8 +56,8 @@ int main() { .submit([&](sycl::handler &cgh) { // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID1:.+]]|2014|1026|{{.*}}.cpp:[[# @LINE + 1]]:19 auto A1 = Buf.get_access(cgh); - // CHECK: {{[0-9]+}}|Construct accessor|0x0|[[ACCID2:.*]]|2016|1026|{{.*}}.cpp:[[# @LINE + 1]]:65 - sycl::accessor A2(Range, cgh); + // CHECK: {{[0-9]+}}|Construct accessor|0x0|[[ACCID2:.*]]|2016|1026|{{.*}}.cpp:[[# @LINE + 1]]:38 + sycl::local_accessor A2(Range, cgh); // CHECK-OPT:Node create|{{.*}}FillBuffer{{.*}}|{{.*}}.cpp:[[# @LINE - 6 ]]:3|{5, 1, 1}, {0, 0, 0}, {0, 0, 0}, 6 // CHECK-NOOPT:Node create|{{.*}}FillBuffer{{.*}}|{{.*}}.cpp:[[# @LINE - 7 ]]:3|{5, 1, 1}, {0, 0, 0}, {0, 0, 0}, 12 cgh.parallel_for(