From a12a084acf43f4a4432ed56dd1feb7c369ce5cfb Mon Sep 17 00:00:00 2001 From: "aidan.belton" Date: Tue, 21 Jun 2022 14:26:41 +0100 Subject: [PATCH 1/2] Update tests to use local_accessor --- SYCL/AtomicRef/add.h | 3 +-- SYCL/AtomicRef/and.h | 3 +-- SYCL/AtomicRef/atomic_memory_order_acq_rel.cpp | 6 ++---- SYCL/AtomicRef/atomic_memory_order_seq_cst.cpp | 3 +-- SYCL/AtomicRef/compare_exchange.h | 3 +-- SYCL/AtomicRef/exchange.h | 3 +-- SYCL/AtomicRef/load.h | 3 +-- SYCL/AtomicRef/max.h | 3 +-- SYCL/AtomicRef/min.h | 3 +-- SYCL/AtomicRef/or.h | 3 +-- SYCL/AtomicRef/store.h | 3 +-- SYCL/AtomicRef/sub.h | 3 +-- SYCL/AtomicRef/xor.h | 3 +-- SYCL/Basic/device_event.cpp | 4 +--- SYCL/Basic/group_async_copy.cpp | 3 +-- SYCL/Basic/multi_ptr.cpp | 7 ++----- SYCL/DeviceLib/ITTAnnotations/barrier.cpp | 3 +-- SYCL/DeviceLib/string_test.cpp | 5 +---- SYCL/DiscardEvents/discard_events_accessors.cpp | 5 +---- SYCL/GroupAlgorithm/SYCL2020/sort.cpp | 8 ++------ SYCL/GroupAlgorithm/barrier.cpp | 12 ++++-------- SYCL/Regression/group.cpp | 4 +--- SYCL/Regression/local-arg-align.cpp | 6 ++---- SYCL/Regression/zero_size_local_accessor.cpp | 4 +--- SYCL/SubGroup/load_store.cpp | 6 ++---- SYCL/SubGroup/sub_group_as.cpp | 4 +--- SYCL/SubGroup/sub_group_as_vec.cpp | 4 +--- SYCL/XPTI/buffer/accessors.cpp | 2 +- SYCL/XPTI/kernel/basic.cpp | 2 +- 29 files changed, 37 insertions(+), 84 deletions(-) diff --git a/SYCL/AtomicRef/add.h b/SYCL/AtomicRef/add.h index f5a5330cbc..808b3e5523 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 30bf678a12..9f41ca6e42 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 af95cc48ce..81f113e4a9 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 2d1cdfb83f..39c7c56b53 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 6efa3bca35..7c31f3e283 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 2e7d6e840e..abeb48cf95 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 4a98141d7c..95fec5092d 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 30f94d3e92..a2ef84cc2e 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 bc24ada2af..0cb1781c07 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 84aab0ef8a..4a91428366 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 825b98396c..d0b9e96a0b 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 7fe96e49e0..487a20f9eb 100644 --- a/SYCL/Basic/device_event.cpp +++ b/SYCL/Basic/device_event.cpp @@ -76,9 +76,7 @@ int test_strideN(size_t stride) { myQueue.submit([&](handler &cgh) { auto out_ptr = out_buf.get_access(cgh); - accessor - local_acc(range<1>(16), cgh); + local_accessor local_acc(range<1>(16), cgh); // 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 bc56f4cce0..059aa31370 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 9c993bad2d..03a99aea0c 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 6707b989e5..3842b8eef6 100644 --- a/SYCL/DeviceLib/string_test.cpp +++ b/SYCL/DeviceLib/string_test.cpp @@ -390,10 +390,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 259e05c7ec..d9a5a706e2 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 9cc07785b2..e9cdd8656a 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 035dcb398b..edcba7883f 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 aee4664b62..d76a51a177 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 adbc579d75..9beb6814fe 100644 --- a/SYCL/SubGroup/sub_group_as.cpp +++ b/SYCL/SubGroup/sub_group_as.cpp @@ -36,9 +36,7 @@ int main(int argc, char *argv[]) { queue.submit([&](cl::sycl::handler &cgh) { auto global = buf.get_access(cgh); - sycl::accessor - local(N, cgh); + sycl::local_accessor local(N, cgh); cgh.parallel_for( cl::sycl::nd_range<1>(N, 32), [=](cl::sycl::nd_item<1> it) { diff --git a/SYCL/SubGroup/sub_group_as_vec.cpp b/SYCL/SubGroup/sub_group_as_vec.cpp index d84991b36c..c38f379fa3 100644 --- a/SYCL/SubGroup/sub_group_as_vec.cpp +++ b/SYCL/SubGroup/sub_group_as_vec.cpp @@ -38,9 +38,7 @@ int main(int argc, char *argv[]) { queue.submit([&](cl::sycl::handler &cgh) { auto global = buf.get_access(cgh); - sycl::accessor, 1, sycl::access::mode::read_write, - sycl::access::target::local> - local(N, cgh); + sycl::local_accessor, 1> local(N, cgh); cgh.parallel_for( cl::sycl::nd_range<1>(N, 32), [=](cl::sycl::nd_item<1> it) { cl::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..b80d215a2f 100644 --- a/SYCL/XPTI/buffer/accessors.cpp +++ b/SYCL/XPTI/buffer/accessors.cpp @@ -28,7 +28,7 @@ int main() { // 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); + 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 1f0223c303..902dbdab0b 100644 --- a/SYCL/XPTI/kernel/basic.cpp +++ b/SYCL/XPTI/kernel/basic.cpp @@ -57,7 +57,7 @@ int main() { // 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); + 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( From 52732cc191bf3a1c69baee9e0cd9bf69260d0ce3 Mon Sep 17 00:00:00 2001 From: "aidan.belton" Date: Thu, 1 Sep 2022 17:27:08 +0100 Subject: [PATCH 2/2] Update XPTI tests for local_accessor --- SYCL/XPTI/buffer/accessors.cpp | 2 +- SYCL/XPTI/kernel/basic.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/SYCL/XPTI/buffer/accessors.cpp b/SYCL/XPTI/buffer/accessors.cpp index b80d215a2f..84b9dc7773 100644 --- a/SYCL/XPTI/buffer/accessors.cpp +++ b/SYCL/XPTI/buffer/accessors.cpp @@ -27,7 +27,7 @@ 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 + // 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); diff --git a/SYCL/XPTI/kernel/basic.cpp b/SYCL/XPTI/kernel/basic.cpp index 45d39ca5f1..e2283aee24 100644 --- a/SYCL/XPTI/kernel/basic.cpp +++ b/SYCL/XPTI/kernel/basic.cpp @@ -56,7 +56,7 @@ 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 + // 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