From eb527d43610f42d0abb25358aa812b12d3eea1b4 Mon Sep 17 00:00:00 2001 From: "Kornev, Nikita" Date: Fri, 31 Jul 2026 15:56:55 +0200 Subject: [PATCH 1/4] [SYCL] Update E2E to use access_mode [3/3] access::mode is deprecated in favor of access_mode. PR to add deprecations: https://github.com/intel/llvm/pull/22803 1st - https://github.com/intel/llvm/pull/22842 2nd - https://github.com/intel/llvm/pull/22843 --- .../AddressCast/khr_static_addrspace_cast.cpp | 4 +- sycl/test-e2e/Assert/assert_in_kernels.hpp | 6 +-- .../assert_in_simultaneously_multiple_tus.cpp | 2 +- sycl/test-e2e/AtomicRef/assignment.h | 2 +- sycl/test-e2e/AtomicRef/compare_exchange.h | 8 ++-- sycl/test-e2e/AtomicRef/max.h | 8 ++-- sycl/test-e2e/AtomicRef/store.h | 4 +- sycl/test-e2e/Basic/access_to_subset.cpp | 2 +- sycl/test-e2e/Basic/bit_cast/bit_cast.cpp | 2 +- .../Basic/buffer/buffer_full_copy.cpp | 16 +++---- sycl/test-e2e/Basic/buffer/reinterpret.cpp | 8 ++-- sycl/test-e2e/Basic/char_builtins.cpp | 8 ++-- sycl/test-e2e/Basic/event.cpp | 2 +- sycl/test-e2e/Basic/fill_accessor.cpp | 2 +- .../free_function_queries_sub_group.cpp | 8 ++-- .../Basic/group_async_copy_legacy.cpp | 4 +- .../handler/handler_copy_with_offset.cpp | 4 +- sycl/test-e2e/Basic/host-task-dependency.cpp | 20 ++++---- .../image/image_accessor_readsampler.cpp | 4 +- sycl/test-e2e/Basic/image/image_array.cpp | 6 +-- sycl/test-e2e/Basic/image/image_sample.cpp | 12 ++--- sycl/test-e2e/Basic/image/srgba-read.cpp | 4 +- .../kernel_bundle/kernel_bundle_api_hip.cpp | 2 +- sycl/test-e2e/Basic/linear-sub_group.cpp | 2 +- .../multi_ptr_null_relational_operators.cpp | 2 +- sycl/test-e2e/Basic/nested_queue_submit.cpp | 2 +- .../Basic/parallel_for_user_types.cpp | 4 +- sycl/test-e2e/Basic/span.cpp | 4 +- sycl/test-e2e/Basic/sycl-namespace.cpp | 2 +- sycl/test-e2e/Basic/vector/byte.cpp | 2 +- sycl/test-e2e/Config/env_vars.cpp | 2 +- .../DeprecatedFeatures/set_arg_interop.cpp | 14 +++--- .../DeviceGlobal/device_global_arrow.hpp | 2 +- .../DeviceGlobal/device_global_static.cpp | 2 +- .../DeviceLib/ITTAnnotations/barrier.cpp | 2 +- .../built-ins/ext_native_math_common.hpp | 4 +- .../DeviceLib/built-ins/scalar_common.cpp | 2 +- .../DeviceLib/built-ins/scalar_math.cpp | 48 +++++++++---------- .../DeviceLib/built-ins/vector_common.cpp | 4 +- .../DeviceLib/built-ins/vector_math.cpp | 22 ++++----- .../DeviceLib/c99_complex_math_test.cpp | 2 +- sycl/test-e2e/DeviceLib/complex_utils.hpp | 8 ++-- .../DeviceLib/imf/simd_emulate_test.cpp | 4 +- .../test-e2e/DeviceLib/math_override_test.cpp | 4 +- sycl/test-e2e/DeviceLib/math_windows_test.cpp | 4 +- .../DeviceLib/std_complex_math_test.cpp | 4 +- .../DotProduct/dot_product_vec_test.cpp | 32 ++++++------- .../launch_queries/max_num_work_groups.cpp | 8 ++-- .../launch_queries/max_work_group_size.cpp | 4 +- .../reusable_events/event_as_dependency.cpp | 4 +- .../reusable_events/in_order_queue_signal.cpp | 4 +- .../reusable_events/queue_returned_events.cpp | 2 +- .../FreeFunctionKernels/enum_parameter.cpp | 18 +++---- sycl/test-e2e/Functor/kernel_functor.cpp | 2 +- sycl/test-e2e/GroupAlgorithm/leader.cpp | 2 +- .../GroupLocalMemory/no_early_opt.cpp | 4 +- .../HostInteropTask/host-task-dependency3.cpp | 4 +- sycl/test-e2e/HostInteropTask/host-task.cpp | 8 ++-- .../test-e2e/HostInteropTask/interop-task.cpp | 16 +++---- .../InlineAsm/asm_16_no_input_int.cpp | 2 +- .../test-e2e/InlineAsm/asm_8_no_input_int.cpp | 2 +- sycl/test-e2e/InlineAsm/asm_float_add.cpp | 6 +-- sycl/test-e2e/InlineAsm/asm_if.cpp | 2 +- sycl/test-e2e/InlineAsm/asm_mul.cpp | 6 +-- sycl/test-e2e/InlineAsm/asm_plus_mod.cpp | 4 +- .../Inputs/incrementing_lib.cpp | 2 +- sycl/test-e2e/KernelCompiler/opencl_cache.cpp | 4 +- .../array-kernel-param-nested-run.cpp | 12 ++--- .../KernelParams/struct_kernel_param.cpp | 4 +- .../test-e2e/Reduction/reduction_big_data.cpp | 2 +- .../Reduction/reduction_nd_conditional.cpp | 2 +- .../reduction_range_1d_reducer_skip.cpp | 2 +- sycl/test-e2e/Reduction/reduction_usm.cpp | 4 +- .../2020-spec-constants-debug-info.cpp | 2 +- .../commandlist/Inputs/FindPrimesSYCL.cpp | 4 +- .../Regression/fp16-with-unnamed-lambda.cpp | 2 +- sycl/test-e2e/Regression/group.cpp | 6 +-- ...implicit_kernel_bundle_image_filtering.cpp | 2 +- .../kernel_bundle_ignore_sycl_external.cpp | 2 +- sycl/test-e2e/Regression/local-arg-align.cpp | 2 +- sycl/test-e2e/Regression/multiple-targets.cpp | 10 ++-- .../Regression/private_array_init_test.cpp | 2 +- .../Regression/same_unnamed_kernels.cpp | 6 +-- sycl/test-e2e/Sampler/basic-rw-float.cpp | 6 +-- .../Sampler/normalized-clamp-nearest.cpp | 6 +-- .../normalized-mirror-linear-float.cpp | 6 +-- .../Sampler/normalized-none-nearest.cpp | 6 +-- .../unnormalized-clamp-linear-float.cpp | 6 +-- .../unnormalized-clampedge-nearest.cpp | 6 +-- .../Scheduler/BasicSchedulerTests.cpp | 2 +- .../test-e2e/Scheduler/DeleteCmdException.cpp | 2 +- sycl/test-e2e/Scheduler/MemObjRemapping.cpp | 8 ++-- .../SpecConstants/2020/handler-api.cpp | 16 +++---- .../2020/non_native/Inputs/common.cpp | 4 +- sycl/test-e2e/SubGroup/barrier.cpp | 4 +- sycl/test-e2e/SubGroup/generic_reduce.cpp | 4 +- sycl/test-e2e/SubGroup/scan.hpp | 6 +-- sycl/test-e2e/SubGroupMask/GroupSize.cpp | 2 +- sycl/test-e2e/USM/fill.cpp | 4 +- sycl/test-e2e/USM/memcpy.cpp | 2 +- .../dynamic_alloc_local_accessor.cpp | 2 +- .../work_group_static_memory_block_scope.cpp | 4 +- sycl/test-e2e/XPTI/buffer/accessors.cpp | 16 +++---- .../test-e2e/XPTI/buffer/multiple_buffers.cpp | 4 +- sycl/test-e2e/XPTI/buffer/sub_buffer.cpp | 2 +- sycl/test-e2e/XPTI/kernel/content.cpp | 4 +- 106 files changed, 303 insertions(+), 303 deletions(-) diff --git a/sycl/test-e2e/AddressCast/khr_static_addrspace_cast.cpp b/sycl/test-e2e/AddressCast/khr_static_addrspace_cast.cpp index 9742ef4babfdc..313c262e7fbd7 100644 --- a/sycl/test-e2e/AddressCast/khr_static_addrspace_cast.cpp +++ b/sycl/test-e2e/AddressCast/khr_static_addrspace_cast.cpp @@ -24,10 +24,10 @@ int main() { Queue .submit([&](sycl::handler &cgh) { auto GlobalAccessor = - GlobalBuffer.get_access(cgh); + GlobalBuffer.get_access(cgh); auto LocalAccessor = sycl::local_accessor(1, cgh); auto ResultAccessor = - ResultBuffer.get_access(cgh); + ResultBuffer.get_access(cgh); cgh.parallel_for( sycl::nd_range<1>(NItems, 1), [=](sycl::nd_item<1> Item) { bool Success = true; diff --git a/sycl/test-e2e/Assert/assert_in_kernels.hpp b/sycl/test-e2e/Assert/assert_in_kernels.hpp index 93740c8a15dd1..75813f9d55d6d 100644 --- a/sycl/test-e2e/Assert/assert_in_kernels.hpp +++ b/sycl/test-e2e/Assert/assert_in_kernels.hpp @@ -14,7 +14,7 @@ void kernelFunc1(int *Buf, int wiID) { void assertTest1(queue &Q, buffer &Buf) { Q.submit([&](handler &CGH) { - auto Acc = Buf.template get_access(CGH); + auto Acc = Buf.template get_access(CGH); CGH.parallel_for( Buf.get_range(), [=](sycl::id<1> wiID) { kernelFunc1(&Acc[0], wiID); }); @@ -29,7 +29,7 @@ void kernelFunc2(int *Buf, int wiID) { void assertTest2(queue &Q, buffer &Buf) { Q.submit([&](handler &CGH) { - auto Acc = Buf.template get_access(CGH); + auto Acc = Buf.template get_access(CGH); CGH.parallel_for( Buf.get_range(), [=](sycl::id<1> wiID) { kernelFunc2(&Acc[0], wiID); }); @@ -44,7 +44,7 @@ void kernelFunc3(int *Buf, int wiID) { void assertTest3(queue &Q, buffer &Buf) { Q.submit([&](handler &CGH) { - auto Acc = Buf.template get_access(CGH); + auto Acc = Buf.template get_access(CGH); CGH.parallel_for( Buf.get_range(), [=](sycl::id<1> wiID) { kernelFunc3(&Acc[0], wiID); }); diff --git a/sycl/test-e2e/Assert/assert_in_simultaneously_multiple_tus.cpp b/sycl/test-e2e/Assert/assert_in_simultaneously_multiple_tus.cpp index 143d0b225c066..3f267230f96ae 100644 --- a/sycl/test-e2e/Assert/assert_in_simultaneously_multiple_tus.cpp +++ b/sycl/test-e2e/Assert/assert_in_simultaneously_multiple_tus.cpp @@ -52,7 +52,7 @@ template void enqueueKernel(queue *Q) { sycl::buffer Buf(numOfItems); Q->submit([&](handler &CGH) { - auto Acc = Buf.template get_access(CGH); + auto Acc = Buf.template get_access(CGH); CGH.parallel_for(numOfItems, [=](sycl::id<1> wiID) { Acc[wiID] = 0; diff --git a/sycl/test-e2e/AtomicRef/assignment.h b/sycl/test-e2e/AtomicRef/assignment.h index 660cc7bfbb327..8f7904451a1f8 100644 --- a/sycl/test-e2e/AtomicRef/assignment.h +++ b/sycl/test-e2e/AtomicRef/assignment.h @@ -27,7 +27,7 @@ void assignment_test(queue q, size_t N) { buffer assignment_buf(&assignment, 1); q.submit([&](handler &cgh) { auto st = - assignment_buf.template get_access(cgh); + assignment_buf.template get_access(cgh); cgh.parallel_for>( range<1>(N), [=](item<1> it) { size_t gid = it.get_id(0); diff --git a/sycl/test-e2e/AtomicRef/compare_exchange.h b/sycl/test-e2e/AtomicRef/compare_exchange.h index 7da21730e4793..d54355db4c99c 100644 --- a/sycl/test-e2e/AtomicRef/compare_exchange.h +++ b/sycl/test-e2e/AtomicRef/compare_exchange.h @@ -33,10 +33,10 @@ void compare_exchange_local_test(queue q, size_t N) { buffer output_buf(output.data(), output.size()); q.submit([&](handler &cgh) { auto compare_exchange = - compare_exchange_buf.template get_access( + compare_exchange_buf.template get_access( cgh); auto out = - output_buf.template get_access(cgh); + output_buf.template get_access(cgh); local_accessor loc(1, cgh); cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) { @@ -88,10 +88,10 @@ void compare_exchange_global_test(queue q, size_t N) { q.submit([&](handler &cgh) { auto exc = - compare_exchange_buf.template get_access( + compare_exchange_buf.template get_access( cgh); auto out = - output_buf.template get_access(cgh); + output_buf.template get_access(cgh); cgh.parallel_for(range<1>(N), [=](item<1> it) { size_t gid = it.get_id(0); auto atm = AtomicRef < T, diff --git a/sycl/test-e2e/AtomicRef/max.h b/sycl/test-e2e/AtomicRef/max.h index c9b6da5bc1688..c99c3ef7a3e6e 100644 --- a/sycl/test-e2e/AtomicRef/max.h +++ b/sycl/test-e2e/AtomicRef/max.h @@ -32,9 +32,9 @@ void max_local_test(queue q, size_t N) { buffer cum_buf(&cum, 1); buffer output_buf(output.data(), output.size()); q.submit([&](handler &cgh) { - auto cum = cum_buf.template get_access(cgh); + auto cum = cum_buf.template get_access(cgh); auto out = - output_buf.template get_access(cgh); + output_buf.template get_access(cgh); local_accessor loc(1, cgh); cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) { @@ -83,9 +83,9 @@ void max_global_test(queue q, size_t N) { buffer output_buf(output.data(), output.size()); q.submit([&](handler &cgh) { - auto val = val_buf.template get_access(cgh); + auto val = val_buf.template get_access(cgh); auto out = - output_buf.template get_access(cgh); + output_buf.template get_access(cgh); cgh.parallel_for(range<1>(N), [=](item<1> it) { int gid = it.get_id(0); auto atm = AtomicRef < T, diff --git a/sycl/test-e2e/AtomicRef/store.h b/sycl/test-e2e/AtomicRef/store.h index 8773c80f113d3..fa6a3db28fff0 100644 --- a/sycl/test-e2e/AtomicRef/store.h +++ b/sycl/test-e2e/AtomicRef/store.h @@ -29,7 +29,7 @@ void store_global_test(queue q, size_t N) { { buffer store_buf(&store, 1); q.submit([&](handler &cgh) { - auto st = store_buf.template get_access(cgh); + auto st = store_buf.template get_access(cgh); cgh.parallel_for(range<1>(N), [=](item<1> it) { size_t gid = it.get_id(0); auto atm = AtomicRef(st[0]); @@ -82,7 +82,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); + auto st = store_buf.template get_access(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); diff --git a/sycl/test-e2e/Basic/access_to_subset.cpp b/sycl/test-e2e/Basic/access_to_subset.cpp index 3fc81f477c0de..d927696d7a39c 100644 --- a/sycl/test-e2e/Basic/access_to_subset.cpp +++ b/sycl/test-e2e/Basic/access_to_subset.cpp @@ -12,7 +12,7 @@ #include using namespace sycl; -using acc_w = accessor; +using acc_w = accessor; int main() { diff --git a/sycl/test-e2e/Basic/bit_cast/bit_cast.cpp b/sycl/test-e2e/Basic/bit_cast/bit_cast.cpp index fc3c700462e22..2c61124daf571 100644 --- a/sycl/test-e2e/Basic/bit_cast/bit_cast.cpp +++ b/sycl/test-e2e/Basic/bit_cast/bit_cast.cpp @@ -10,7 +10,7 @@ #include #include -constexpr sycl::access::mode sycl_write = sycl::access::mode::write; +constexpr sycl::access_mode sycl_write = sycl::access_mode::write; template class BitCastKernel; diff --git a/sycl/test-e2e/Basic/buffer/buffer_full_copy.cpp b/sycl/test-e2e/Basic/buffer/buffer_full_copy.cpp index 5106e86bb50a1..2333a8865909a 100644 --- a/sycl/test-e2e/Basic/buffer/buffer_full_copy.cpp +++ b/sycl/test-e2e/Basic/buffer/buffer_full_copy.cpp @@ -22,7 +22,7 @@ void check_copy_device_to_host(sycl::queue &Queue) { // Submit kernel where you'll fill the buffer Queue.submit([&](sycl::handler &cgh) { - auto acc = simple_buffer.get_access(cgh); + auto acc = simple_buffer.get_access(cgh); cgh.fill(acc, 13); }); @@ -69,7 +69,7 @@ void check_fill(sycl::queue &Queue) { } auto e = Queue.submit([&](sycl::handler &cgh) { - auto a = buf_1.template get_access(cgh, size / 2, + auto a = buf_1.template get_access(cgh, size / 2, offset); cgh.fill(a, (float)1337.0); }); @@ -102,8 +102,8 @@ void check_copy_host_to_device(sycl::queue &Queue) { expected_res_2[i] = expected_res_1[offset + i]; auto e = Queue.submit([&](sycl::handler &cgh) { - auto a = buf_1.get_access(cgh, size / 2, offset); - auto b = buf_2.get_access(cgh, size / 2); + auto a = buf_1.get_access(cgh, size / 2, offset); + auto b = buf_2.get_access(cgh, size / 2); cgh.copy(a, b); }); e.wait(); @@ -141,10 +141,10 @@ void check_copy_host_to_device(sycl::queue &Queue) { expected_res_4[i][j] = expected_res_3[i + offset][j + offset]; e = Queue.submit([&](sycl::handler &cgh) { - auto a = buf_3.get_access( + auto a = buf_3.get_access( cgh, {size / 2, size / 2}, {offset, offset}); auto b = - buf_4.get_access(cgh, {size / 2, size / 2}); + buf_4.get_access(cgh, {size / 2, size / 2}); cgh.copy(a, b); }); e.wait(); @@ -190,9 +190,9 @@ void check_copy_host_to_device(sycl::queue &Queue) { expected_res_5[i + offset][j + offset][k + offset]; e = Queue.submit([&](sycl::handler &cgh) { - auto a = buf_5.get_access( + auto a = buf_5.get_access( cgh, {size / 2, size / 2, size / 2}, {offset, offset, offset}); - auto b = buf_6.get_access( + auto b = buf_6.get_access( cgh, {size / 2, size / 2, size / 2}); cgh.copy(a, b); }); diff --git a/sycl/test-e2e/Basic/buffer/reinterpret.cpp b/sycl/test-e2e/Basic/buffer/reinterpret.cpp index 8dba511eb60de..1502e4938268d 100644 --- a/sycl/test-e2e/Basic/buffer/reinterpret.cpp +++ b/sycl/test-e2e/Basic/buffer/reinterpret.cpp @@ -39,7 +39,7 @@ void execute_kernel(sycl::queue &cmd_queue, std::vector &data, cmd_queue.submit([&](sycl::handler &cgh) { auto rb_acc = - reinterpret_subbuf.template get_access(cgh); + reinterpret_subbuf.template get_access(cgh); cgh.parallel_for( reinterpret_subbuf.get_range(), [=](sycl::id index) { rb_acc[index] = val; }); @@ -61,7 +61,7 @@ int main() { sycl::buffer buf_i(r1); auto buf_char = buf_i.reinterpret(r2); cmd_queue.submit([&](sycl::handler &cgh) { - auto acc = buf_char.get_access(cgh); + auto acc = buf_char.get_access(cgh); cgh.parallel_for(r2, [=](sycl::id<1> i) { acc[i] = UCHAR_MAX; }); }); @@ -81,7 +81,7 @@ int main() { sycl::buffer buf_1d(r1d); auto buf_2d = buf_1d.reinterpret(r2d); cmd_queue.submit([&](sycl::handler &cgh) { - auto acc2d = buf_2d.get_access(cgh); + auto acc2d = buf_2d.get_access(cgh); cgh.parallel_for(r2d, [=](sycl::item<2> itemID) { size_t i = itemID.get_id(0); size_t j = itemID.get_id(1); @@ -194,7 +194,7 @@ int main() { cmd_queue.submit([&](sycl::handler &cgh) { auto rb_acc = - reinterpret_subbuf.get_access(cgh); + reinterpret_subbuf.get_access(cgh); cgh.parallel_for( reinterpret_subbuf.get_range(), [=](sycl::id<1> index) { rb_acc[index] = 13; }); diff --git a/sycl/test-e2e/Basic/char_builtins.cpp b/sycl/test-e2e/Basic/char_builtins.cpp index f31c69c51a5e1..3ca6df038a080 100644 --- a/sycl/test-e2e/Basic/char_builtins.cpp +++ b/sycl/test-e2e/Basic/char_builtins.cpp @@ -46,9 +46,9 @@ template int doCharTest(const T *A, const T *B, T *C) { buffer BBuf(B, BufferSize); buffer CBuf(C, BufferSize); Q.submit([&](handler &CGH) { - auto AAcc = ABuf.template get_access(CGH); - auto BAcc = BBuf.template get_access(CGH); - auto CAcc = CBuf.template get_access(CGH); + auto AAcc = ABuf.template get_access(CGH); + auto BAcc = BBuf.template get_access(CGH); + auto CAcc = CBuf.template get_access(CGH); CGH.single_task<>([=]() { CAcc[0] = clz(AAcc[0]); CAcc[1] = ctz(AAcc[1]); @@ -69,7 +69,7 @@ template int doCharTest(const T *A, const T *B, T *C) { { buffer ABuf(A, BufferSize); Q.submit([&](handler &CGH) { - auto AAcc = ABuf.template get_access(CGH); + auto AAcc = ABuf.template get_access(CGH); local_accessor Local(range<1>{WorkGroupSize}, CGH); nd_range<1> NDR{range<1>(NElems), range<1>(WorkGroupSize)}; diff --git a/sycl/test-e2e/Basic/event.cpp b/sycl/test-e2e/Basic/event.cpp index 513e1aa8d6304..aca2fe33185e7 100644 --- a/sycl/test-e2e/Basic/event.cpp +++ b/sycl/test-e2e/Basic/event.cpp @@ -70,7 +70,7 @@ int main() { float Scalar = 2.0; sycl::buffer Buf(&Data, sycl::range<1>(1)); auto Event = Queue.submit([&](sycl::handler &cgh) { - auto Acc = Buf.get_access(cgh); + auto Acc = Buf.get_access(cgh); cgh.single_task([=]() { Acc[0] = Scalar; }); }); diff --git a/sycl/test-e2e/Basic/fill_accessor.cpp b/sycl/test-e2e/Basic/fill_accessor.cpp index 7327e4b5859f0..cd0d7f58852ef 100644 --- a/sycl/test-e2e/Basic/fill_accessor.cpp +++ b/sycl/test-e2e/Basic/fill_accessor.cpp @@ -52,7 +52,7 @@ void CheckFillZeroDimAccessor(queue &Q, T Init, T Expected) { { buffer Buffer(Data.data(), Range); Q.submit([&](handler &CGH) { - accessor Accessor(Buffer, CGH); + accessor Accessor(Buffer, CGH); CGH.fill(Accessor, Expected); }).wait_and_throw(); } diff --git a/sycl/test-e2e/Basic/free_function_queries/free_function_queries_sub_group.cpp b/sycl/test-e2e/Basic/free_function_queries/free_function_queries_sub_group.cpp index 7d8e94bbfb132..f6639ce3527a8 100644 --- a/sycl/test-e2e/Basic/free_function_queries/free_function_queries_sub_group.cpp +++ b/sycl/test-e2e/Basic/free_function_queries/free_function_queries_sub_group.cpp @@ -30,12 +30,12 @@ int main() { sycl::queue q; sycl::nd_range<1> NDR(sycl::range<1>{n}, sycl::range<1>{2}); q.submit([&](sycl::handler &cgh) { - sycl::accessor - acc(buf.get_access(cgh)); - sycl::accessor(cgh)); + sycl::accessor - results_acc(results_buf.get_access(cgh)); + results_acc(results_buf.get_access(cgh)); cgh.parallel_for(NDR, [=](auto nd_i) { static_assert(std::is_same>::value, "lambda arg type is unexpected"); diff --git a/sycl/test-e2e/Basic/group_async_copy_legacy.cpp b/sycl/test-e2e/Basic/group_async_copy_legacy.cpp index 9b2279fdcbb83..2cba047661d9c 100644 --- a/sycl/test-e2e/Basic/group_async_copy_legacy.cpp +++ b/sycl/test-e2e/Basic/group_async_copy_legacy.cpp @@ -114,8 +114,8 @@ template int test(size_t Stride) { // In must be read_write for legacy multi_ptr as it would otherwise produce // a const multi_ptr which was not a valid source argument to // async_work_group_copy in SYCL 1.2.1. - auto In = InBuf.template get_access(CGH); - auto Out = OutBuf.template get_access(CGH); + auto In = InBuf.template get_access(CGH); + auto Out = OutBuf.template get_access(CGH); local_accessor Local(range<1>{WorkGroupSize}, CGH); nd_range<1> NDR{range<1>(NElems), range<1>(WorkGroupSize)}; diff --git a/sycl/test-e2e/Basic/handler/handler_copy_with_offset.cpp b/sycl/test-e2e/Basic/handler/handler_copy_with_offset.cpp index c54047f49b8d2..1f2dfaf10aa0e 100644 --- a/sycl/test-e2e/Basic/handler/handler_copy_with_offset.cpp +++ b/sycl/test-e2e/Basic/handler/handler_copy_with_offset.cpp @@ -19,8 +19,8 @@ #include using namespace sycl; -constexpr access::mode read = access::mode::read; -constexpr access::mode write = access::mode::write; +constexpr access_mode read = access_mode::read; +constexpr access_mode write = access_mode::write; constexpr access::target ondevice = access::target::device; int main() { diff --git a/sycl/test-e2e/Basic/host-task-dependency.cpp b/sycl/test-e2e/Basic/host-task-dependency.cpp index be6ea2c94a29f..50839dfb62818 100644 --- a/sycl/test-e2e/Basic/host-task-dependency.cpp +++ b/sycl/test-e2e/Basic/host-task-dependency.cpp @@ -27,9 +27,9 @@ struct Context { S::event HostTask_CopyBuf1ToBuf2(Context *Ctx) { S::event Event = Ctx->Queue.submit([&](S::handler &CGH) { - S::host_accessor CopierSrcAcc(Ctx->Buf1, + S::host_accessor CopierSrcAcc(Ctx->Buf1, CGH); - S::host_accessor CopierDstAcc(Ctx->Buf2, + S::host_accessor CopierDstAcc(Ctx->Buf2, CGH); auto CopierHostTask = [=] { @@ -54,21 +54,21 @@ S::event HostTask_CopyBuf1ToBuf2(Context *Ctx) { void Thread1Fn(Context *Ctx) { // 0. initialize resulting buffer with apriori wrong result { - S::host_accessor Acc(Ctx->Buf1); + S::host_accessor Acc(Ctx->Buf1); for (size_t Idx = 0; Idx < Acc.size(); ++Idx) Acc[Idx] = -1; } { - S::host_accessor Acc(Ctx->Buf2); + S::host_accessor Acc(Ctx->Buf2); for (size_t Idx = 0; Idx < Acc.size(); ++Idx) Acc[Idx] = -2; } { - S::host_accessor Acc(Ctx->Buf3); + S::host_accessor Acc(Ctx->Buf3); for (size_t Idx = 0; Idx < Acc.size(); ++Idx) Acc[Idx] = -3; @@ -76,7 +76,7 @@ void Thread1Fn(Context *Ctx) { // 1. submit task writing to buffer 1 Ctx->Queue.submit([&](S::handler &CGH) { - S::accessor + S::accessor GeneratorAcc(Ctx->Buf1, CGH); auto GeneratorKernel = [GeneratorAcc] { @@ -92,9 +92,9 @@ void Thread1Fn(Context *Ctx) { // 3. submit simple task to move data between two buffers Ctx->Queue.submit([&](S::handler &CGH) { - S::accessor + S::accessor SrcAcc(Ctx->Buf2, CGH); - S::accessor + S::accessor DstAcc(Ctx->Buf3, CGH); CGH.depends_on(HostTaskEvent); @@ -109,7 +109,7 @@ void Thread1Fn(Context *Ctx) { // 4. check data in buffer #3 { - S::host_accessor Acc(Ctx->Buf3); + S::host_accessor Acc(Ctx->Buf3); bool Failure = false; @@ -154,7 +154,7 @@ void test() { // 3. check via host accessor that buf 2 contains valid data { - S::host_accessor ResultAcc(Ctx.Buf2); + S::host_accessor ResultAcc(Ctx.Buf2); bool Failure = false; for (size_t Idx = 0; Idx < ResultAcc.size(); ++Idx) { diff --git a/sycl/test-e2e/Basic/image/image_accessor_readsampler.cpp b/sycl/test-e2e/Basic/image/image_accessor_readsampler.cpp index faa3a54967efc..1c22e0086f9fe 100644 --- a/sycl/test-e2e/Basic/image/image_accessor_readsampler.cpp +++ b/sycl/test-e2e/Basic/image/image_accessor_readsampler.cpp @@ -67,8 +67,8 @@ void checkReadSampler(char *host_ptr, s::sampler Sampler, s::float4 Coord, s::queue myQueue; s::buffer ReadDataBuf(&ReadData, s::range<1>(1)); myQueue.submit([&](s::handler &cgh) { - auto ReadAcc = Img.get_access(cgh); - s::accessor ReadDataBufAcc( + auto ReadAcc = Img.get_access(cgh); + s::accessor ReadDataBufAcc( ReadDataBuf, cgh); cgh.single_task>([=]() { diff --git a/sycl/test-e2e/Basic/image/image_array.cpp b/sycl/test-e2e/Basic/image/image_array.cpp index 2b5b826d7ba75..81c9f2741dec3 100644 --- a/sycl/test-e2e/Basic/image/image_array.cpp +++ b/sycl/test-e2e/Basic/image/image_array.cpp @@ -57,9 +57,9 @@ int main() { TestQueue Q{sycl::default_selector_v}; - constexpr auto SYCLRead = sycl::access::mode::read; - constexpr auto SYCLWrite = sycl::access::mode::write; - constexpr auto SYCLReadWrite = sycl::access::mode::read_write; + constexpr auto SYCLRead = sycl::access_mode::read; + constexpr auto SYCLWrite = sycl::access_mode::write; + constexpr auto SYCLReadWrite = sycl::access_mode::read_write; Q.submit([&](sycl::handler &CGH) { auto ImgAcc = Img.get_access(CGH); diff --git a/sycl/test-e2e/Basic/image/image_sample.cpp b/sycl/test-e2e/Basic/image/image_sample.cpp index 22d59061b7a90..1bfa52c5152f4 100644 --- a/sycl/test-e2e/Basic/image/image_sample.cpp +++ b/sycl/test-e2e/Basic/image/image_sample.cpp @@ -51,8 +51,8 @@ bool test1d_coord(dataT *hostPtr, coordT coord, dataT expectedColour) { // Do the test by reading a single pixel from the image s::queue myQueue(s::default_selector_v); myQueue.submit([&](s::handler &cgh) { - auto imageAcc = image.get_access(cgh); - s::accessor resultDataAcc(resultDataBuf, + auto imageAcc = image.get_access(cgh); + s::accessor resultDataAcc(resultDataBuf, cgh); cgh.single_task([=]() { @@ -86,8 +86,8 @@ bool test2d_coord(dataT *hostPtr, coordT coord, dataT expectedColour) { // Do the test by reading a single pixel from the image s::queue myQueue(s::default_selector_v); myQueue.submit([&](s::handler &cgh) { - auto imageAcc = image.get_access(cgh); - s::accessor resultDataAcc(resultDataBuf, + auto imageAcc = image.get_access(cgh); + s::accessor resultDataAcc(resultDataBuf, cgh); cgh.single_task([=]() { @@ -121,8 +121,8 @@ bool test3d_coord(dataT *hostPtr, coordT coord, dataT expectedColour) { // Do the test by reading a single pixel from the image s::queue myQueue(s::default_selector_v); myQueue.submit([&](s::handler &cgh) { - auto imageAcc = image.get_access(cgh); - s::accessor resultDataAcc(resultDataBuf, + auto imageAcc = image.get_access(cgh); + s::accessor resultDataAcc(resultDataBuf, cgh); cgh.single_task([=]() { diff --git a/sycl/test-e2e/Basic/image/srgba-read.cpp b/sycl/test-e2e/Basic/image/srgba-read.cpp index 91f0d4052bca8..e999daf6b5777 100644 --- a/sycl/test-e2e/Basic/image/srgba-read.cpp +++ b/sycl/test-e2e/Basic/image/srgba-read.cpp @@ -53,8 +53,8 @@ void test_rd(image_channel_order ChanOrder, image_channel_type ChanType) { Q.submit([&](handler &cgh) { auto image_acc = - image_2D.get_access(cgh); - auto test_acc = testResults.get_access(cgh); + image_2D.get_access(cgh); + auto test_acc = testResults.get_access(cgh); cgh.single_task([=]() { int i = 0; // the index for writing into the testResult buffer. diff --git a/sycl/test-e2e/Basic/kernel_bundle/kernel_bundle_api_hip.cpp b/sycl/test-e2e/Basic/kernel_bundle/kernel_bundle_api_hip.cpp index db558009a6f03..ec437a364ba4a 100644 --- a/sycl/test-e2e/Basic/kernel_bundle/kernel_bundle_api_hip.cpp +++ b/sycl/test-e2e/Basic/kernel_bundle/kernel_bundle_api_hip.cpp @@ -147,7 +147,7 @@ int main() { // The error handling tests rely on Kernel3Name binary existing. sycl::buffer Buf(sycl::range<1>{1}); Q.submit([&](sycl::handler &CGH) { - auto Acc = Buf.get_access(CGH); + auto Acc = Buf.get_access(CGH); CGH.single_task([=]() { Acc[0] = 42; }); }); diff --git a/sycl/test-e2e/Basic/linear-sub_group.cpp b/sycl/test-e2e/Basic/linear-sub_group.cpp index 684ff7d7fbdcc..a08876c8ec83f 100644 --- a/sycl/test-e2e/Basic/linear-sub_group.cpp +++ b/sycl/test-e2e/Basic/linear-sub_group.cpp @@ -28,7 +28,7 @@ int main(int argc, char *argv[]) { { buffer output_buf(output.data(), range<2>(outer, inner)); q.submit([&](handler &cgh) { - auto output = output_buf.get_access(cgh); + auto output = output_buf.get_access(cgh); cgh.parallel_for( nd_range<2>(range<2>(outer, inner), range<2>(outer, inner)), [=](nd_item<2> it) { diff --git a/sycl/test-e2e/Basic/multi_ptr_null_relational_operators.cpp b/sycl/test-e2e/Basic/multi_ptr_null_relational_operators.cpp index 6aaca3eab4ad8..16897929bae8a 100644 --- a/sycl/test-e2e/Basic/multi_ptr_null_relational_operators.cpp +++ b/sycl/test-e2e/Basic/multi_ptr_null_relational_operators.cpp @@ -33,7 +33,7 @@ void nullptrRelationalOperatorTest() { sycl::buffer buf(1); queue .submit([&](sycl::handler &cgh) { - auto dev_acc = buf.get_access(cgh); + auto dev_acc = buf.get_access(cgh); if constexpr (address_space == sycl::access::address_space::local_space) { sycl::local_accessor locAcc(1, cgh); diff --git a/sycl/test-e2e/Basic/nested_queue_submit.cpp b/sycl/test-e2e/Basic/nested_queue_submit.cpp index 8e06abb7ef095..4b50ce612198b 100644 --- a/sycl/test-e2e/Basic/nested_queue_submit.cpp +++ b/sycl/test-e2e/Basic/nested_queue_submit.cpp @@ -20,7 +20,7 @@ void nestedSubmitParallelFor(sycl::queue &q) { { sycl::buffer buf(array.data(), sycl::range<1>{n}); q.submit([&](sycl::handler &h) { - auto acc = buf.get_access(h); + auto acc = buf.get_access(h); q.parallel_for(sycl::range<1>{n}, [=](sycl::id<1> i) { acc[i] = float(0.0); }); }); diff --git a/sycl/test-e2e/Basic/parallel_for_user_types.cpp b/sycl/test-e2e/Basic/parallel_for_user_types.cpp index c77d2eb4183a1..344a24add3ed6 100644 --- a/sycl/test-e2e/Basic/parallel_for_user_types.cpp +++ b/sycl/test-e2e/Basic/parallel_for_user_types.cpp @@ -40,7 +40,7 @@ int main() { // Check user defined sycl::item wrapper sycl::buffer data_buf(data, sz); q.submit([&](sycl::handler &h) { - auto buf_acc = data_buf.get_access(h); + auto buf_acc = data_buf.get_access(h); h.parallel_for(sycl::range<1>{sz}, [=](item_wrapper<1> item) { buf_acc[item.get()] += 1; }); }); @@ -60,7 +60,7 @@ int main() { // Check user defined sycl::nd_item wrapper q.submit([&](sycl::handler &h) { - auto buf_acc = data_buf.get_access(h); + auto buf_acc = data_buf.get_access(h); h.parallel_for(sycl::nd_range<1>{sz, 2}, [=](nd_item_wrapper<1> item) { buf_acc[item.get()] += 1; }); }); diff --git a/sycl/test-e2e/Basic/span.cpp b/sycl/test-e2e/Basic/span.cpp index 2a112ba4740d5..730c4e5cf209b 100644 --- a/sycl/test-e2e/Basic/span.cpp +++ b/sycl/test-e2e/Basic/span.cpp @@ -41,7 +41,7 @@ void testSpanCapture() { usm_span[0] += 100; // 101 modify first value using span affordance. event E = Q.submit([&](handler &cgh) { - auto can_read_from_span_acc = SpanRead.get_access(cgh); + auto can_read_from_span_acc = SpanRead.get_access(cgh); cgh.single_task([=] { // read from the spans. can_read_from_span_acc[0] = vecUSM_span[0]; @@ -84,7 +84,7 @@ void testSpanOnDevice() { buffer SpanRead(NumberOfReadTestsRange); event E = Q.submit([&](handler &cgh) { - auto can_read_from_span_acc = SpanRead.get_access(cgh); + auto can_read_from_span_acc = SpanRead.get_access(cgh); cgh.single_task([=] { // create a span on device, pass it to function that modifies it // read values back out. diff --git a/sycl/test-e2e/Basic/sycl-namespace.cpp b/sycl/test-e2e/Basic/sycl-namespace.cpp index e4c8a64e611cf..af06eb1b48f17 100644 --- a/sycl/test-e2e/Basic/sycl-namespace.cpp +++ b/sycl/test-e2e/Basic/sycl-namespace.cpp @@ -9,7 +9,7 @@ int main() { { sycl::buffer b(&r, 1); q.submit([&](sycl::handler &h) { - auto a = b.get_access(h); + auto a = b.get_access(h); h.single_task([=]() { a[0] = 42; }); }); } diff --git a/sycl/test-e2e/Basic/vector/byte.cpp b/sycl/test-e2e/Basic/vector/byte.cpp index d74ef5f97f7c4..212b1f549d83a 100644 --- a/sycl/test-e2e/Basic/vector/byte.cpp +++ b/sycl/test-e2e/Basic/vector/byte.cpp @@ -72,7 +72,7 @@ int main() { sycl::queue Queue; Queue .submit([&](sycl::handler &cgh) { - auto Acc = Buf.get_access(cgh); + auto Acc = Buf.get_access(cgh); cgh.single_task([=]() { // load sycl::multi_ptr(cgh); + auto B = buf.get_access(cgh); cgh.single_task([=]() { B[0] = 0; }); }); assert(!shouldCrash); diff --git a/sycl/test-e2e/DeprecatedFeatures/set_arg_interop.cpp b/sycl/test-e2e/DeprecatedFeatures/set_arg_interop.cpp index 84381d070d0d7..cbe74b2a2f222 100644 --- a/sycl/test-e2e/DeprecatedFeatures/set_arg_interop.cpp +++ b/sycl/test-e2e/DeprecatedFeatures/set_arg_interop.cpp @@ -58,8 +58,8 @@ int main() { buffer FirstBuffer(Array, range<1>(1)); buffer SecondBuffer(&Value, range<1>(1)); Queue.submit([&](handler &CGH) { - CGH.set_arg(0, FirstBuffer.get_access(CGH)); - CGH.set_arg(1, SecondBuffer.get_access(CGH)); + CGH.set_arg(0, FirstBuffer.get_access(CGH)); + CGH.set_arg(1, SecondBuffer.get_access(CGH)); CGH.single_task(FirstKernel); }); } @@ -71,8 +71,8 @@ int main() { { buffer FirstBuffer(Array, range<1>(Count)); Queue.submit([&](handler &CGH) { - auto Acc = FirstBuffer.get_access(CGH); - CGH.set_arg(0, FirstBuffer.get_access(CGH)); + auto Acc = FirstBuffer.get_access(CGH); + CGH.set_arg(0, FirstBuffer.get_access(CGH)); CGH.parallel_for(range<1>{Count}, SecondKernel); }); } @@ -105,9 +105,9 @@ int main() { { buffer FirstBuffer(Array, range<1>(Count)); Queue.submit([&](handler &CGH) { - auto Acc = FirstBuffer.get_access(CGH); - CGH.set_arg(0, FirstBuffer.get_access(CGH)); - CGH.set_arg(1, sycl::accessor(CGH); + CGH.set_arg(0, FirstBuffer.get_access(CGH)); + CGH.set_arg(1, sycl::accessor( sycl::range<1>(Count), CGH)); CGH.parallel_for(range<1>{Count}, ThirdKernel); diff --git a/sycl/test-e2e/DeviceGlobal/device_global_arrow.hpp b/sycl/test-e2e/DeviceGlobal/device_global_arrow.hpp index ba869734243f9..4807adc34f05f 100644 --- a/sycl/test-e2e/DeviceGlobal/device_global_arrow.hpp +++ b/sycl/test-e2e/DeviceGlobal/device_global_arrow.hpp @@ -32,7 +32,7 @@ int test() { { buffer OutBuf{Out, 2}; Q.submit([&](handler &CGH) { - auto OutAcc = OutBuf.get_access(CGH); + auto OutAcc = OutBuf.get_access(CGH); CGH.single_task([=]() { OutAcc[0] = DeviceGlobalVar1->getX(); OutAcc[1] = DeviceGlobalVar2->getX(); diff --git a/sycl/test-e2e/DeviceGlobal/device_global_static.cpp b/sycl/test-e2e/DeviceGlobal/device_global_static.cpp index 0202c10bbfe44..a6a054766b619 100644 --- a/sycl/test-e2e/DeviceGlobal/device_global_static.cpp +++ b/sycl/test-e2e/DeviceGlobal/device_global_static.cpp @@ -27,7 +27,7 @@ int main() { { buffer OutBuf(&OutVal, 1); Q.submit([&](handler &CGH) { - auto OutAcc = OutBuf.get_access(CGH); + auto OutAcc = OutBuf.get_access(CGH); CGH.single_task([=]() { OutAcc[0] = DeviceGlobalVar.get()[0]; }); }); } diff --git a/sycl/test-e2e/DeviceLib/ITTAnnotations/barrier.cpp b/sycl/test-e2e/DeviceLib/ITTAnnotations/barrier.cpp index e748e1253dde4..281f22783f7c8 100644 --- a/sycl/test-e2e/DeviceLib/ITTAnnotations/barrier.cpp +++ b/sycl/test-e2e/DeviceLib/ITTAnnotations/barrier.cpp @@ -19,7 +19,7 @@ int main() { // Ensure that a simple kernel gets run when instrumented with // ITT start/finish annotations and ITT wg_barrier/wi_resume annotations. q.submit([&](handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); local_accessor local_acc(local_range, cgh); cgh.parallel_for( nd_range<1>(num_items, local_range), [=](nd_item<1> item) { diff --git a/sycl/test-e2e/DeviceLib/built-ins/ext_native_math_common.hpp b/sycl/test-e2e/DeviceLib/built-ins/ext_native_math_common.hpp index a67320dcb6798..7c84f7f763a32 100644 --- a/sycl/test-e2e/DeviceLib/built-ins/ext_native_math_common.hpp +++ b/sycl/test-e2e/DeviceLib/built-ins/ext_native_math_common.hpp @@ -33,7 +33,7 @@ void native_tanh_tester(sycl::queue q, T val, T up, T lo) { { sycl::buffer BufR(&r, sycl::range<1>(1)); q.submit([&](sycl::handler &cgh) { - auto AccR = BufR.template get_access(cgh); + auto AccR = BufR.template get_access(cgh); cgh.single_task([=]() { AccR[0] = sycl::ext::oneapi::experimental::native::tanh(AccR[0]); }); @@ -54,7 +54,7 @@ void native_exp2_tester(sycl::queue q, T val, T up, T lo) { { sycl::buffer BufR(&r, sycl::range<1>(1)); q.submit([&](sycl::handler &cgh) { - auto AccR = BufR.template get_access(cgh); + auto AccR = BufR.template get_access(cgh); cgh.single_task([=]() { AccR[0] = sycl::ext::oneapi::experimental::native::exp2(AccR[0]); }); diff --git a/sycl/test-e2e/DeviceLib/built-ins/scalar_common.cpp b/sycl/test-e2e/DeviceLib/built-ins/scalar_common.cpp index 5e12037006e91..83812a0e3da42 100644 --- a/sycl/test-e2e/DeviceLib/built-ins/scalar_common.cpp +++ b/sycl/test-e2e/DeviceLib/built-ins/scalar_common.cpp @@ -17,7 +17,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::max(float{0.5f}, float{2.3f}); }); }); diff --git a/sycl/test-e2e/DeviceLib/built-ins/scalar_math.cpp b/sycl/test-e2e/DeviceLib/built-ins/scalar_math.cpp index ffd405ef5b21b..1e0bd105455ea 100644 --- a/sycl/test-e2e/DeviceLib/built-ins/scalar_math.cpp +++ b/sycl/test-e2e/DeviceLib/built-ins/scalar_math.cpp @@ -21,7 +21,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task([=]() { AccR[0] = s::acos(float{0.5}); }); }); } @@ -35,7 +35,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::acosh(float{2.4}); }); }); @@ -50,7 +50,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task([=]() { AccR[0] = s::asin(float{0.5}); }); }); } @@ -64,7 +64,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::asinh(float{0.5}); }); }); @@ -79,7 +79,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task([=]() { AccR[0] = s::atan(float{0.5}); }); }); } @@ -93,7 +93,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::atanh(float{0.5}); }); }); @@ -108,7 +108,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::cbrt(float{27.0}); }); }); @@ -123,7 +123,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task([=]() { AccR[0] = s::ceil(float{0.5}); }); }); } @@ -137,7 +137,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task([=]() { AccR[0] = s::cos(float{0.5}); }); }); } @@ -151,7 +151,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task([=]() { AccR[0] = s::cosh(float{0.5}); }); }); } @@ -165,7 +165,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::cospi(float{0.1}); }); }); @@ -180,7 +180,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task([=]() { AccR[0] = s::erfc(float{0.5}); }); }); } @@ -194,7 +194,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task([=]() { AccR[0] = s::erf(float{0.5}); }); }); } @@ -208,7 +208,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task([=]() { AccR[0] = s::exp(float{0.5}); }); }); } @@ -222,7 +222,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task([=]() { AccR[0] = s::exp2(float{8.0}); }); }); } @@ -236,7 +236,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task([=]() { AccR[0] = s::exp10(float{2}); }); }); } @@ -250,7 +250,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::expm1(float{0.5}); }); }); @@ -265,7 +265,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::fabs(float{-0.5}); }); }); @@ -280,7 +280,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::floor(float{0.5}); }); }); @@ -295,7 +295,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::fmax(float{0.5}, float{0.8}); }); }); @@ -310,7 +310,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::fmin(float{0.5}, float{0.8}); }); }); @@ -325,7 +325,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::fmod(float{5.1}, float{3.0}); }); }); @@ -340,7 +340,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::lgamma(float{10.f}); }); }); @@ -355,7 +355,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::lgamma(float{-2.4f}); }); }); diff --git a/sycl/test-e2e/DeviceLib/built-ins/vector_common.cpp b/sycl/test-e2e/DeviceLib/built-ins/vector_common.cpp index d817eacf61079..d12e009972cce 100644 --- a/sycl/test-e2e/DeviceLib/built-ins/vector_common.cpp +++ b/sycl/test-e2e/DeviceLib/built-ins/vector_common.cpp @@ -17,7 +17,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task([=]() { AccR[0] = s::max(s::float2{0.5f, 3.4f}, s::float2{2.3f, 0.4f}); }); @@ -36,7 +36,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task([=]() { AccR[0] = s::max(s::float2{0.5f, 3.4f}, float{3.0f}); }); diff --git a/sycl/test-e2e/DeviceLib/built-ins/vector_math.cpp b/sycl/test-e2e/DeviceLib/built-ins/vector_math.cpp index a47101562556d..dba593e17ea9d 100644 --- a/sycl/test-e2e/DeviceLib/built-ins/vector_math.cpp +++ b/sycl/test-e2e/DeviceLib/built-ins/vector_math.cpp @@ -20,7 +20,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task([=]() { AccR[0] = s::fmin(s::float2{0.5f, 3.4f}, s::float2{2.3f, 0.4f}); }); @@ -39,7 +39,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task([=]() { AccR[0] = s::fabs(s::float2{-1.0f, 2.0f}); }); @@ -58,7 +58,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task([=]() { AccR[0] = s::floor(s::float2{1.4f, 2.8f}); }); @@ -77,7 +77,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task([=]() { AccR[0] = s::ceil(s::float2{1.4f, 2.8f}); }); @@ -99,8 +99,8 @@ int main() { s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); - auto AccI = BufI.get_access(cgh); + auto AccR = BufR.get_access(cgh); + auto AccI = BufI.get_access(cgh); cgh.single_task([=]() { s::global_ptr Iptr(AccI); AccR[0] = s::fract(s::float2{1.5f, 2.5f}, Iptr); @@ -128,8 +128,8 @@ int main() { s::buffer BufI(&i, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); - auto AccI = BufI.get_access(cgh); + auto AccR = BufR.get_access(cgh); + auto AccI = BufI.get_access(cgh); cgh.single_task([=]() { s::float2 temp(0.0); s::private_ptr Iptr(&temp); @@ -157,7 +157,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task([=]() { AccR[0] = s::lgamma(s::float2{10.f, -2.4f}); }); @@ -180,8 +180,8 @@ int main() { s::buffer BufI(&i, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); - auto AccI = BufI.get_access(cgh); + auto AccR = BufR.get_access(cgh); + auto AccI = BufI.get_access(cgh); cgh.single_task([=]() { s::int2 temp(0.0); s::private_ptr Iptr(&temp); diff --git a/sycl/test-e2e/DeviceLib/c99_complex_math_test.cpp b/sycl/test-e2e/DeviceLib/c99_complex_math_test.cpp index 8f9e1de2bd023..160b19ea412e8 100644 --- a/sycl/test-e2e/DeviceLib/c99_complex_math_test.cpp +++ b/sycl/test-e2e/DeviceLib/c99_complex_math_test.cpp @@ -58,7 +58,7 @@ float _Complex catanhf(float _Complex z); } namespace s = sycl; -constexpr s::access::mode sycl_write = s::access::mode::write; +constexpr s::access_mode sycl_write = s::access_mode::write; bool approx_equal_cmplx_f(float _Complex x, float _Complex y) { return approx_equal_fp(__real__ x, __real__ y) && diff --git a/sycl/test-e2e/DeviceLib/complex_utils.hpp b/sycl/test-e2e/DeviceLib/complex_utils.hpp index 26bd41fa489c1..d75342f6ac38c 100644 --- a/sycl/test-e2e/DeviceLib/complex_utils.hpp +++ b/sycl/test-e2e/DeviceLib/complex_utils.hpp @@ -309,9 +309,9 @@ int device_complex_test_mul(sycl::queue &deviceQueue, complex_mul_result.data(), numOfMulOutput); deviceQueue.submit([&](sycl::handler &cgh) { auto complex_mul_access = - buffer_complex_mul.template get_access(cgh); + buffer_complex_mul.template get_access(cgh); auto complex_mul_res_access = - buffer_complex_mul_res.template get_access( + buffer_complex_mul_res.template get_access( cgh); cgh.single_task([=]() { size_t i, j; @@ -349,9 +349,9 @@ int device_complex_test_div(sycl::queue &deviceQueue, complex_div_result.data(), numOfDivOutput); deviceQueue.submit([&](sycl::handler &cgh) { auto complex_div_access = - buffer_complex_div.template get_access(cgh); + buffer_complex_div.template get_access(cgh); auto complex_div_res_access = - buffer_complex_div_res.template get_access( + buffer_complex_div_res.template get_access( cgh); cgh.single_task([=]() { size_t i, j; diff --git a/sycl/test-e2e/DeviceLib/imf/simd_emulate_test.cpp b/sycl/test-e2e/DeviceLib/imf/simd_emulate_test.cpp index 72ff03f6568e4..69a4548dabad5 100644 --- a/sycl/test-e2e/DeviceLib/imf/simd_emulate_test.cpp +++ b/sycl/test-e2e/DeviceLib/imf/simd_emulate_test.cpp @@ -7,8 +7,8 @@ #include namespace s = sycl; -constexpr s::access::mode sycl_read = s::access::mode::read; -constexpr s::access::mode sycl_write = s::access::mode::write; +constexpr s::access_mode sycl_read = s::access_mode::read; +constexpr s::access_mode sycl_write = s::access_mode::write; void run_vabs2_4_test(s::queue &queue) { bool pass = true; diff --git a/sycl/test-e2e/DeviceLib/math_override_test.cpp b/sycl/test-e2e/DeviceLib/math_override_test.cpp index 0a1f4e6cb4220..c0f83d31796ea 100644 --- a/sycl/test-e2e/DeviceLib/math_override_test.cpp +++ b/sycl/test-e2e/DeviceLib/math_override_test.cpp @@ -5,8 +5,8 @@ #include "math_utils.hpp" namespace s = sycl; -constexpr s::access::mode sycl_read = s::access::mode::read; -constexpr s::access::mode sycl_write = s::access::mode::write; +constexpr s::access_mode sycl_read = s::access_mode::read; +constexpr s::access_mode sycl_write = s::access_mode::write; // Dummy function provided by user to override device library // version. diff --git a/sycl/test-e2e/DeviceLib/math_windows_test.cpp b/sycl/test-e2e/DeviceLib/math_windows_test.cpp index e03a9498aae7f..63b15de65b2dc 100644 --- a/sycl/test-e2e/DeviceLib/math_windows_test.cpp +++ b/sycl/test-e2e/DeviceLib/math_windows_test.cpp @@ -13,8 +13,8 @@ #include namespace s = sycl; -constexpr s::access::mode sycl_read = s::access::mode::read; -constexpr s::access::mode sycl_write = s::access::mode::write; +constexpr s::access_mode sycl_read = s::access_mode::read; +constexpr s::access_mode sycl_write = s::access_mode::write; #define TEST_NUM 35 diff --git a/sycl/test-e2e/DeviceLib/std_complex_math_test.cpp b/sycl/test-e2e/DeviceLib/std_complex_math_test.cpp index 45950309fa90f..06714e5762da5 100644 --- a/sycl/test-e2e/DeviceLib/std_complex_math_test.cpp +++ b/sycl/test-e2e/DeviceLib/std_complex_math_test.cpp @@ -10,8 +10,8 @@ using std::complex; namespace s = sycl; -constexpr s::access::mode sycl_read = s::access::mode::read; -constexpr s::access::mode sycl_write = s::access::mode::write; +constexpr s::access_mode sycl_read = s::access_mode::read; +constexpr s::access_mode sycl_write = s::access_mode::write; template bool approx_equal_cmplx(complex x, complex y) { return approx_equal_fp(x.real(), y.real()) && diff --git a/sycl/test-e2e/DotProduct/dot_product_vec_test.cpp b/sycl/test-e2e/DotProduct/dot_product_vec_test.cpp index 28ac47cc66993..0ecf18f1f4487 100644 --- a/sycl/test-e2e/DotProduct/dot_product_vec_test.cpp +++ b/sycl/test-e2e/DotProduct/dot_product_vec_test.cpp @@ -66,10 +66,10 @@ static bool testss(queue &Q) { buffer Dbuf(D, range<1>(RangeLength)); Q.submit([&](handler &cgh) { - auto Ap = Abuf.get_access(cgh); - auto Bp = Bbuf.get_access(cgh); - auto Cp = Cbuf.get_access(cgh); - auto Dp = Dbuf.get_access(cgh); + auto Ap = Abuf.get_access(cgh); + auto Bp = Bbuf.get_access(cgh); + auto Cp = Cbuf.get_access(cgh); + auto Dp = Dbuf.get_access(cgh); cgh.parallel_for(range<1>(RangeLength), [=](id<1> I) { Dp[I] = dot_acc(Ap[I], Bp[I], Cp[I]); @@ -105,10 +105,10 @@ static bool testuu(queue &Q) { buffer Dbuf(D, range<1>(RangeLength)); Q.submit([&](handler &cgh) { - auto Ap = Abuf.get_access(cgh); - auto Bp = Bbuf.get_access(cgh); - auto Cp = Cbuf.get_access(cgh); - auto Dp = Dbuf.get_access(cgh); + auto Ap = Abuf.get_access(cgh); + auto Bp = Bbuf.get_access(cgh); + auto Cp = Cbuf.get_access(cgh); + auto Dp = Dbuf.get_access(cgh); cgh.parallel_for(range<1>(RangeLength), [=](id<1> I) { Dp[I] = dot_acc(Ap[I], Bp[I], Cp[I]); @@ -144,10 +144,10 @@ static bool testsu(queue &Q) { buffer Dbuf(D, range<1>(RangeLength)); Q.submit([&](handler &cgh) { - auto Ap = Abuf.get_access(cgh); - auto Bp = Bbuf.get_access(cgh); - auto Cp = Cbuf.get_access(cgh); - auto Dp = Dbuf.get_access(cgh); + auto Ap = Abuf.get_access(cgh); + auto Bp = Bbuf.get_access(cgh); + auto Cp = Cbuf.get_access(cgh); + auto Dp = Dbuf.get_access(cgh); cgh.parallel_for(range<1>(RangeLength), [=](id<1> I) { Dp[I] = dot_acc(Ap[I], Bp[I], Cp[I]); @@ -183,10 +183,10 @@ static bool testus(queue &Q) { buffer Dbuf(D, range<1>(RangeLength)); Q.submit([&](handler &cgh) { - auto Ap = Abuf.get_access(cgh); - auto Bp = Bbuf.get_access(cgh); - auto Cp = Cbuf.get_access(cgh); - auto Dp = Dbuf.get_access(cgh); + auto Ap = Abuf.get_access(cgh); + auto Bp = Bbuf.get_access(cgh); + auto Cp = Cbuf.get_access(cgh); + auto Dp = Dbuf.get_access(cgh); cgh.parallel_for(range<1>(RangeLength), [=](id<1> I) { Dp[I] = dot_acc(Ap[I], Bp[I], Cp[I]); diff --git a/sycl/test-e2e/Experimental/launch_queries/max_num_work_groups.cpp b/sycl/test-e2e/Experimental/launch_queries/max_num_work_groups.cpp index 133176ef36430..bd6d638dc3ad5 100644 --- a/sycl/test-e2e/Experimental/launch_queries/max_num_work_groups.cpp +++ b/sycl/test-e2e/Experimental/launch_queries/max_num_work_groups.cpp @@ -20,7 +20,7 @@ namespace kernels { template using sycl_global_accessor = - sycl::accessor; class TestKernel { @@ -110,7 +110,7 @@ int test_max_num_work_groups(sycl::queue &q, const sycl::device &dev) { auto launch_range = sycl::nd_range<1>{sycl::range<1>{NumWorkItems}, sycl::range<1>{workGroupSize}}; q.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); if constexpr (KernelName::HasLocalMemory) { sycl::local_accessor loc_acc{ sycl::range<1>{workGroupSize}, cgh}; @@ -142,7 +142,7 @@ int test_max_num_work_groups(sycl::queue &q, const sycl::device &dev) { sycl::range<1>{workGroupSize}}; q.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); if constexpr (KernelName::HasLocalMemory) { sycl::local_accessor loc_acc{sycl::range<1>{localSize}, cgh}; @@ -179,7 +179,7 @@ int test_max_num_work_groups(sycl::queue &q, const sycl::device &dev) { sycl::range<1>{workGroupSize}}; q.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); if constexpr (KernelName::HasLocalMemory) { sycl::local_accessor loc_acc{sycl::range<1>{localSize}, cgh}; diff --git a/sycl/test-e2e/Experimental/launch_queries/max_work_group_size.cpp b/sycl/test-e2e/Experimental/launch_queries/max_work_group_size.cpp index 0de747e8a18f3..eddaaba1251e2 100644 --- a/sycl/test-e2e/Experimental/launch_queries/max_work_group_size.cpp +++ b/sycl/test-e2e/Experimental/launch_queries/max_work_group_size.cpp @@ -15,7 +15,7 @@ namespace kernels { template using sycl_global_accessor = - sycl::accessor; class TestKernel { public: @@ -54,7 +54,7 @@ int main() { auto launch_range = sycl::nd_range<1>{sycl::range<1>{numWorkItems}, sycl::range<1>{1}}; q.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.parallel_for(launch_range, kernels::TestKernel{acc}); }).wait(); diff --git a/sycl/test-e2e/Experimental/reusable_events/event_as_dependency.cpp b/sycl/test-e2e/Experimental/reusable_events/event_as_dependency.cpp index 4eb97915e4a46..c7131f54dca7d 100644 --- a/sycl/test-e2e/Experimental/reusable_events/event_as_dependency.cpp +++ b/sycl/test-e2e/Experimental/reusable_events/event_as_dependency.cpp @@ -23,7 +23,7 @@ int main() { // First kernel q.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.parallel_for(sycl::range<1>(N), [=](sycl::id<1> idx) { acc[idx] = 10; }); }); @@ -32,7 +32,7 @@ int main() { // Second kernel depends on event q.submit([&](sycl::handler &cgh) { cgh.depends_on(event); - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.parallel_for( sycl::range<1>(N), [=](sycl::id<1> idx) { acc[idx] = acc[idx] + 7; }); }); diff --git a/sycl/test-e2e/Experimental/reusable_events/in_order_queue_signal.cpp b/sycl/test-e2e/Experimental/reusable_events/in_order_queue_signal.cpp index 31a10fdafda5d..8c84aeb905245 100644 --- a/sycl/test-e2e/Experimental/reusable_events/in_order_queue_signal.cpp +++ b/sycl/test-e2e/Experimental/reusable_events/in_order_queue_signal.cpp @@ -23,13 +23,13 @@ int main() { // Submit multiple operations q.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.parallel_for(sycl::range<1>(N), [=](sycl::id<1> idx) { acc[idx] = 1; }); }); q.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.parallel_for( sycl::range<1>(N), [=](sycl::id<1> idx) { acc[idx] = acc[idx] + 2; }); }); diff --git a/sycl/test-e2e/Experimental/reusable_events/queue_returned_events.cpp b/sycl/test-e2e/Experimental/reusable_events/queue_returned_events.cpp index e3e6216bf4240..4d4ea7097e9f4 100644 --- a/sycl/test-e2e/Experimental/reusable_events/queue_returned_events.cpp +++ b/sycl/test-e2e/Experimental/reusable_events/queue_returned_events.cpp @@ -19,7 +19,7 @@ int main() { sycl::buffer buf(data.data(), sycl::range<1>(N)); auto event = q.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.parallel_for(sycl::range<1>(N), [=](sycl::id<1> idx) { acc[idx] = 42; }); }); diff --git a/sycl/test-e2e/FreeFunctionKernels/enum_parameter.cpp b/sycl/test-e2e/FreeFunctionKernels/enum_parameter.cpp index 76d7c62ae54a2..718740a4f2e53 100644 --- a/sycl/test-e2e/FreeFunctionKernels/enum_parameter.cpp +++ b/sycl/test-e2e/FreeFunctionKernels/enum_parameter.cpp @@ -9,20 +9,20 @@ #include using namespace sycl; -using rAccType = sycl::accessor; -using wAccType = sycl::accessor; -using rwAccType = sycl::accessor; +using rAccType = sycl::accessor; +using wAccType = sycl::accessor; +using rwAccType = sycl::accessor; using flagType = sycl::accessor; -template -bool hasAccessorMode(sycl::accessor acc, access::mode mode) { +template +bool hasAccessorMode(sycl::accessor acc, access_mode mode) { return modeT == mode; } template SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel)) -void verifyAccessorMode(AccType acc, access::mode accMode, flagType flagAcc) { +void verifyAccessorMode(AccType acc, access_mode accMode, flagType flagAcc) { flagAcc[0] = hasAccessorMode(acc, accMode); } @@ -33,9 +33,9 @@ int main() { sycl::kernel Kernel2 = getKernel>(Context); sycl::kernel Kernel3 = getKernel>(Context); - const auto rMode = sycl::access::mode::read; - const auto wMode = sycl::access::mode::write; - const auto rwMode = sycl::access::mode::read_write; + const auto rMode = sycl::access_mode::read; + const auto wMode = sycl::access_mode::write; + const auto rwMode = sycl::access_mode::read_write; bool flag1, flag2, flag3; flag1 = (flag2 = (flag3 = false)); diff --git a/sycl/test-e2e/Functor/kernel_functor.cpp b/sycl/test-e2e/Functor/kernel_functor.cpp index a4b8a7579f4cb..dc271b8ed6b1a 100644 --- a/sycl/test-e2e/Functor/kernel_functor.cpp +++ b/sycl/test-e2e/Functor/kernel_functor.cpp @@ -14,7 +14,7 @@ #include -constexpr auto sycl_read_write = sycl::access::mode::read_write; +constexpr auto sycl_read_write = sycl::access_mode::read_write; constexpr auto sycl_device = sycl::access::target::device; // Case 1: diff --git a/sycl/test-e2e/GroupAlgorithm/leader.cpp b/sycl/test-e2e/GroupAlgorithm/leader.cpp index 7b23f27c1c7ea..e4ad383b48d35 100644 --- a/sycl/test-e2e/GroupAlgorithm/leader.cpp +++ b/sycl/test-e2e/GroupAlgorithm/leader.cpp @@ -17,7 +17,7 @@ void test(queue q) { buffer out_buf(&out, 1); q.submit([&](handler &cgh) { - auto out = out_buf.template get_access(cgh); + auto out = out_buf.template get_access(cgh); cgh.parallel_for(nd_range<2>(R, R), [=](nd_item<2> it) { group<2> g = it.get_group(); if (g.leader()) { diff --git a/sycl/test-e2e/GroupLocalMemory/no_early_opt.cpp b/sycl/test-e2e/GroupLocalMemory/no_early_opt.cpp index b0b3a0e8a8309..405e58727faeb 100644 --- a/sycl/test-e2e/GroupLocalMemory/no_early_opt.cpp +++ b/sycl/test-e2e/GroupLocalMemory/no_early_opt.cpp @@ -30,8 +30,8 @@ int main() { buffer BufB{VecB.data(), range<1>(Size)}; Q.submit([&](handler &Cgh) { - auto AccA = BufA.get_access(Cgh); - auto AccB = BufB.get_access(Cgh); + auto AccA = BufA.get_access(Cgh); + auto AccB = BufB.get_access(Cgh); Cgh.parallel_for( nd_range<1>(range<1>(Size), range<1>(WgSize)), [=](nd_item<1> Item) { multi_ptr(CGH); + auto Acc0 = B0.get_access(CGH); CGH.single_task([=] { Acc0[1] = 1 * Idx; }); }); @@ -65,7 +65,7 @@ void test(size_t Count) { Q.submit([&](handler &CGH) { std::cout << "Submit Kernel-2" << std::endl; - auto Acc1 = B1.get_access(CGH); + auto Acc1 = B1.get_access(CGH); CGH.single_task([=] { Acc1[2] = 1 * Idx; }); }); diff --git a/sycl/test-e2e/HostInteropTask/host-task.cpp b/sycl/test-e2e/HostInteropTask/host-task.cpp index 3d6423aa6e0f8..779aaf189810a 100644 --- a/sycl/test-e2e/HostInteropTask/host-task.cpp +++ b/sycl/test-e2e/HostInteropTask/host-task.cpp @@ -28,7 +28,7 @@ void test1(queue &Q) { buffer Buffer{BUFFER_SIZE}; Q.submit([&](handler &CGH) { - auto Acc = Buffer.get_access(CGH); + auto Acc = Buffer.get_access(CGH); CGH.host_task([=] { /* A no-op */ }); }); @@ -41,15 +41,15 @@ void test2(queue &Q) { buffer Buffer2{BUFFER_SIZE}; Q.submit([&](handler &CGH) { - auto Acc = Buffer1.template get_access(CGH); + auto Acc = Buffer1.template get_access(CGH); auto Kernel = [=](item<1> Id) { Acc[Id] = 123; }; CGH.parallel_for>(Acc.size(), Kernel); }); Q.submit([&](handler &CGH) { - auto AccSrc = Buffer1.template get_access(CGH); - auto AccDst = Buffer2.template get_access(CGH); + auto AccSrc = Buffer1.template get_access(CGH); + auto AccDst = Buffer2.template get_access(CGH); auto Func = [=] { for (size_t Idx = 0; Idx < AccDst.size(); ++Idx) diff --git a/sycl/test-e2e/HostInteropTask/interop-task.cpp b/sycl/test-e2e/HostInteropTask/interop-task.cpp index d0ca0d2c0659f..203baea4ca5ae 100644 --- a/sycl/test-e2e/HostInteropTask/interop-task.cpp +++ b/sycl/test-e2e/HostInteropTask/interop-task.cpp @@ -34,8 +34,8 @@ void checkBufferValues(BufferT Buffer, ValueT Value) { template void copy(buffer &Src, buffer &Dst, queue &Q) { Q.submit([&](handler &CGH) { - auto SrcA = Src.template get_access(CGH); - auto DstA = Dst.template get_access(CGH); + auto SrcA = Src.template get_access(CGH); + auto DstA = Dst.template get_access(CGH); auto Func = [=](interop_handle IH) { auto NativeQ = IH.get_native_queue(); @@ -60,7 +60,7 @@ void copy(buffer &Src, buffer &Dst, queue &Q) { template void modify(buffer &B, queue &Q) { Q.submit([&](handler &CGH) { - auto Acc = B.template get_access(CGH); + auto Acc = B.template get_access(CGH); auto Kernel = [=](item<1> Id) { Acc[Id] += 1; }; @@ -71,8 +71,8 @@ template void modify(buffer &B, queue &Q) { template void init(buffer &B1, buffer &B2, queue &Q) { Q.submit([&](handler &CGH) { - auto Acc1 = B1.template get_access(CGH); - auto Acc2 = B2.template get_access(CGH); + auto Acc1 = B1.template get_access(CGH); + auto Acc2 = B2.template get_access(CGH); CGH.parallel_for>(BUFFER_SIZE, [=](item<1> Id) { Acc1[Id] = -1; @@ -156,7 +156,7 @@ void test3(queue &Q) { buffer Buffer{BUFFER_SIZE * 128}; event Event = Q.submit([&](handler &CGH) { - auto Acc1 = Buffer.get_access(CGH); + auto Acc1 = Buffer.get_access(CGH); CGH.parallel_for(BUFFER_SIZE, [=](item<1> Id) { Acc1[Id] = 123; }); @@ -178,7 +178,7 @@ void test4(queue &Q) { buffer Buffer{BUFFER_SIZE}; Q.submit([&](handler &CGH) { - auto Acc = Buffer.get_access(CGH); + auto Acc = Buffer.get_access(CGH); auto Func = [=](interop_handle IH) { /*A no-op */ }; CGH.host_task(Func); }); @@ -189,7 +189,7 @@ void test5(queue &Q) { buffer Buffer2{BUFFER_SIZE}; Q.submit([&](handler &CGH) { - auto Acc = Buffer1.template get_access(CGH); + auto Acc = Buffer1.template get_access(CGH); auto Kernel = [=](item<1> Id) { Acc[Id] = 123; }; CGH.parallel_for(Acc.get_count(), Kernel); diff --git a/sycl/test-e2e/InlineAsm/asm_16_no_input_int.cpp b/sycl/test-e2e/InlineAsm/asm_16_no_input_int.cpp index e73772aa00885..d4bca3eb82966 100644 --- a/sycl/test-e2e/InlineAsm/asm_16_no_input_int.cpp +++ b/sycl/test-e2e/InlineAsm/asm_16_no_input_int.cpp @@ -15,7 +15,7 @@ template struct KernelFunctor : WithOutputBuffer { void operator()(sycl::handler &cgh) { auto C = - this->getOutputBuffer().template get_access( + this->getOutputBuffer().template get_access( cgh); cgh.parallel_for>( sycl::range<1>{this->getOutputBufferSize()}, diff --git a/sycl/test-e2e/InlineAsm/asm_8_no_input_int.cpp b/sycl/test-e2e/InlineAsm/asm_8_no_input_int.cpp index 7fba71b894bdb..cf7405532bf6f 100644 --- a/sycl/test-e2e/InlineAsm/asm_8_no_input_int.cpp +++ b/sycl/test-e2e/InlineAsm/asm_8_no_input_int.cpp @@ -13,7 +13,7 @@ template struct KernelFunctor : WithOutputBuffer { void operator()(sycl::handler &cgh) { auto C = - this->getOutputBuffer().template get_access( + this->getOutputBuffer().template get_access( cgh); cgh.parallel_for>( sycl::range<1>{this->getOutputBufferSize()}, diff --git a/sycl/test-e2e/InlineAsm/asm_float_add.cpp b/sycl/test-e2e/InlineAsm/asm_float_add.cpp index 0466540dfd494..9f3cbc60dc4bb 100644 --- a/sycl/test-e2e/InlineAsm/asm_float_add.cpp +++ b/sycl/test-e2e/InlineAsm/asm_float_add.cpp @@ -19,13 +19,13 @@ struct KernelFunctor : WithInputBuffers, WithOutputBuffer { void operator()(sycl::handler &cgh) { auto A = - this->getInputBuffer(0).template get_access( + this->getInputBuffer(0).template get_access( cgh); auto B = - this->getInputBuffer(1).template get_access( + this->getInputBuffer(1).template get_access( cgh); auto C = - this->getOutputBuffer().template get_access( + this->getOutputBuffer().template get_access( cgh); cgh.parallel_for>( diff --git a/sycl/test-e2e/InlineAsm/asm_if.cpp b/sycl/test-e2e/InlineAsm/asm_if.cpp index 83b9a921fcc48..de6072c3b2798 100644 --- a/sycl/test-e2e/InlineAsm/asm_if.cpp +++ b/sycl/test-e2e/InlineAsm/asm_if.cpp @@ -13,7 +13,7 @@ template struct KernelFunctor : WithOutputBuffer { void operator()(sycl::handler &CGH) { auto C = - this->getOutputBuffer().template get_access( + this->getOutputBuffer().template get_access( CGH); bool switchField = false; CGH.parallel_for>( diff --git a/sycl/test-e2e/InlineAsm/asm_mul.cpp b/sycl/test-e2e/InlineAsm/asm_mul.cpp index 6030bb5b144b9..ed37ebfc85439 100644 --- a/sycl/test-e2e/InlineAsm/asm_mul.cpp +++ b/sycl/test-e2e/InlineAsm/asm_mul.cpp @@ -17,13 +17,13 @@ struct KernelFunctor : WithInputBuffers, WithOutputBuffer { input1.size()) {} void operator()(sycl::handler &cgh) { auto A = - this->getInputBuffer(0).template get_access( + this->getInputBuffer(0).template get_access( cgh); auto B = - this->getInputBuffer(1).template get_access( + this->getInputBuffer(1).template get_access( cgh); auto C = - this->getOutputBuffer().template get_access( + this->getOutputBuffer().template get_access( cgh); cgh.parallel_for>( diff --git a/sycl/test-e2e/InlineAsm/asm_plus_mod.cpp b/sycl/test-e2e/InlineAsm/asm_plus_mod.cpp index 7ac0c3af60150..96b305016eba9 100644 --- a/sycl/test-e2e/InlineAsm/asm_plus_mod.cpp +++ b/sycl/test-e2e/InlineAsm/asm_plus_mod.cpp @@ -17,10 +17,10 @@ struct KernelFunctor : WithInputBuffers, WithOutputBuffer { void operator()(sycl::handler &cgh) { auto A = - this->getInputBuffer(0).template get_access( + this->getInputBuffer(0).template get_access( cgh); auto B = - this->getOutputBuffer().template get_access( + this->getOutputBuffer().template get_access( cgh); cgh.parallel_for>( diff --git a/sycl/test-e2e/IntermediateLib/Inputs/incrementing_lib.cpp b/sycl/test-e2e/IntermediateLib/Inputs/incrementing_lib.cpp index 9c8faf38700cc..ea3dbe0c4565c 100644 --- a/sycl/test-e2e/IntermediateLib/Inputs/incrementing_lib.cpp +++ b/sycl/test-e2e/IntermediateLib/Inputs/incrementing_lib.cpp @@ -31,7 +31,7 @@ extern "C" API_EXPORT void performIncrementation(sycl::queue &q, sycl::buffer &buf) { sycl::range<1> r = buf.get_range(); q.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.parallel_for( r, [=](sycl::id<1> idx) { acc[idx] += INC; }); }); diff --git a/sycl/test-e2e/KernelCompiler/opencl_cache.cpp b/sycl/test-e2e/KernelCompiler/opencl_cache.cpp index e154bfc7cdb7f..7b13ab079372a 100644 --- a/sycl/test-e2e/KernelCompiler/opencl_cache.cpp +++ b/sycl/test-e2e/KernelCompiler/opencl_cache.cpp @@ -53,8 +53,8 @@ void testSyclKernel(sycl::queue &Q, sycl::kernel Kernel, int multiplier, sycl::buffer OutputBuf(OutputArray, sycl::range<1>(N)); Q.submit([&](sycl::handler &CGH) { - CGH.set_arg(0, InputBuf.get_access(CGH)); - CGH.set_arg(1, OutputBuf.get_access(CGH)); + CGH.set_arg(0, InputBuf.get_access(CGH)); + CGH.set_arg(1, OutputBuf.get_access(CGH)); CGH.parallel_for(sycl::range<1>{N}, Kernel); }); diff --git a/sycl/test-e2e/KernelParams/array-kernel-param-nested-run.cpp b/sycl/test-e2e/KernelParams/array-kernel-param-nested-run.cpp index 3ca90a0f07fe8..aad4a8e9d5e5b 100644 --- a/sycl/test-e2e/KernelParams/array-kernel-param-nested-run.cpp +++ b/sycl/test-e2e/KernelParams/array-kernel-param-nested-run.cpp @@ -58,7 +58,7 @@ bool test_accessor_array_in_struct(queue &myQueue) { myQueue.submit([&](handler &cgh) { using Accessor = - accessor; + accessor; struct S { int w; @@ -68,11 +68,11 @@ bool test_accessor_array_in_struct(queue &myQueue) { int z; } S = {3, 3, - {in_buffer1.get_access(cgh), - in_buffer2.get_access(cgh)}, + {in_buffer1.get_access(cgh), + in_buffer2.get_access(cgh)}, 7, 7}; - auto output_accessor = out_buffer.get_access(cgh); + auto output_accessor = out_buffer.get_access(cgh); cgh.parallel_for( num_items, [=](sycl::id<1> index) { @@ -103,8 +103,8 @@ bool test_templated_array_in_struct(queue &myQueue) { myQueue.submit([&](handler &cgh) { using Accessor = - accessor; - auto output_accessor = out_buffer.get_access(cgh); + accessor; + auto output_accessor = out_buffer.get_access(cgh); cgh.parallel_for( num_items, [=](sycl::id<1> index) { diff --git a/sycl/test-e2e/KernelParams/struct_kernel_param.cpp b/sycl/test-e2e/KernelParams/struct_kernel_param.cpp index bb9206b295762..a55a6986d6963 100644 --- a/sycl/test-e2e/KernelParams/struct_kernel_param.cpp +++ b/sycl/test-e2e/KernelParams/struct_kernel_param.cpp @@ -60,7 +60,7 @@ bool test0() { buffer Buf(&S0, range<1>(1)); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = Buf.get_access(cgh); + auto B = Buf.get_access(cgh); cgh.single_task([=] { B[0] = S; }); }); } @@ -87,7 +87,7 @@ bool test1() { buffer Buffer((unsigned int *)result, range<1>(4)); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = Buffer.get_access(cgh); + auto B = Buffer.get_access(cgh); cgh.parallel_for(range<1>{4}, [=](id<1> index) { B[index.get(0)] = index.get(0) > 2 ? ice2 : ice.get(index.get(0)); }); diff --git a/sycl/test-e2e/Reduction/reduction_big_data.cpp b/sycl/test-e2e/Reduction/reduction_big_data.cpp index 6fc5468aaf92a..f317bea555d6c 100644 --- a/sycl/test-e2e/Reduction/reduction_big_data.cpp +++ b/sycl/test-e2e/Reduction/reduction_big_data.cpp @@ -52,7 +52,7 @@ int test(queue &Q, T Identity) { // Compute. Q.submit([&](handler &CGH) { - auto In = InBuf.template get_access(CGH); + auto In = InBuf.template get_access(CGH); CGH.parallel_for(NDRange, sycl::reduction(OutBuf, CGH, Identity, BOp), [=](nd_item<1> NDIt, auto &Sum) { if (NDIt.get_global_linear_id() < NWorkItems) diff --git a/sycl/test-e2e/Reduction/reduction_nd_conditional.cpp b/sycl/test-e2e/Reduction/reduction_nd_conditional.cpp index f8a6d0fb8027a..795567d5a889c 100644 --- a/sycl/test-e2e/Reduction/reduction_nd_conditional.cpp +++ b/sycl/test-e2e/Reduction/reduction_nd_conditional.cpp @@ -41,7 +41,7 @@ int test(queue &Q, T Identity, size_t WGSize, size_t NWItems) { // Compute. Q.submit([&](handler &CGH) { - auto In = InBuf.template get_access(CGH); + auto In = InBuf.template get_access(CGH); auto Redu = sycl::reduction(OutBuf, CGH, Identity, BOp); CGH.parallel_for(NDRange, Redu, [=](nd_item<1> NDIt, auto &Sum) { size_t I = NDIt.get_global_linear_id(); diff --git a/sycl/test-e2e/Reduction/reduction_range_1d_reducer_skip.cpp b/sycl/test-e2e/Reduction/reduction_range_1d_reducer_skip.cpp index 2f5077c41f03f..ef5d50b24cd5a 100644 --- a/sycl/test-e2e/Reduction/reduction_range_1d_reducer_skip.cpp +++ b/sycl/test-e2e/Reduction/reduction_range_1d_reducer_skip.cpp @@ -31,7 +31,7 @@ int main() { size_t MaxWGSize = Q.get_device().get_info(); - constexpr access::mode RW = access::mode::read_write; + constexpr access_mode RW = access_mode::read_write; // Fast-reduce and Fast-atomics. Try various range types/sizes. tests(Q, 0, 99, std::plus{}, range<1>(1)); tests(Q, 0, 99, std::plus<>{}, range<1>(2)); diff --git a/sycl/test-e2e/Reduction/reduction_usm.cpp b/sycl/test-e2e/Reduction/reduction_usm.cpp index 5f86545e878f3..ecc8866f13068 100644 --- a/sycl/test-e2e/Reduction/reduction_usm.cpp +++ b/sycl/test-e2e/Reduction/reduction_usm.cpp @@ -52,7 +52,7 @@ int test(queue &Q, OptionalIdentity Identity, T Init, } }; - auto In = InBuf.template get_access(CGH); + auto In = InBuf.template get_access(CGH); auto Redu = CreateReduction(); CGH.parallel_for>( NDRange, Redu, [=](nd_item<1> NDIt, auto &Sum) { @@ -65,7 +65,7 @@ int test(queue &Q, OptionalIdentity Identity, T Init, if (AllocType == usm::alloc::device) { buffer Buf(&ComputedOut, range<1>(1)); Q.submit([&](handler &CGH) { - auto OutAcc = Buf.template get_access(CGH); + auto OutAcc = Buf.template get_access(CGH); CGH.single_task>( [=]() { OutAcc[0] = *ReduVarPtr; }); }).wait(); diff --git a/sycl/test-e2e/Regression/2020-spec-constants-debug-info.cpp b/sycl/test-e2e/Regression/2020-spec-constants-debug-info.cpp index 9fbb08b0deda3..945885295449a 100644 --- a/sycl/test-e2e/Regression/2020-spec-constants-debug-info.cpp +++ b/sycl/test-e2e/Regression/2020-spec-constants-debug-info.cpp @@ -19,7 +19,7 @@ int main() { { sycl::buffer Buf{sycl::range{1}}; q.submit([&](sycl::handler &cgh) { - auto Acc = Buf.get_access(cgh); + auto Acc = Buf.get_access(cgh); cgh.set_specialization_constant(1); cgh.single_task([=](sycl::kernel_handler kh) { Acc[0] = kh.get_specialization_constant(); diff --git a/sycl/test-e2e/Regression/commandlist/Inputs/FindPrimesSYCL.cpp b/sycl/test-e2e/Regression/commandlist/Inputs/FindPrimesSYCL.cpp index 7e460f4cb73be..e2da59e61cfda 100644 --- a/sycl/test-e2e/Regression/commandlist/Inputs/FindPrimesSYCL.cpp +++ b/sycl/test-e2e/Regression/commandlist/Inputs/FindPrimesSYCL.cpp @@ -9,8 +9,8 @@ #include #include -constexpr sycl::access::mode sycl_read = sycl::access::mode::read; -constexpr sycl::access::mode sycl_write = sycl::access::mode::write; +constexpr sycl::access_mode sycl_read = sycl::access_mode::read; +constexpr sycl::access_mode sycl_write = sycl::access_mode::write; using namespace std; diff --git a/sycl/test-e2e/Regression/fp16-with-unnamed-lambda.cpp b/sycl/test-e2e/Regression/fp16-with-unnamed-lambda.cpp index 4686ccc76efc6..8f15a0fa7ba16 100644 --- a/sycl/test-e2e/Regression/fp16-with-unnamed-lambda.cpp +++ b/sycl/test-e2e/Regression/fp16-with-unnamed-lambda.cpp @@ -24,7 +24,7 @@ int main() { sycl::buffer Buf(1); Q.submit([&](sycl::handler &CGH) { - auto Acc = Buf.get_access(CGH); + auto Acc = Buf.get_access(CGH); CGH.single_task([=]() { Acc[0] = 1; }); }); diff --git a/sycl/test-e2e/Regression/group.cpp b/sycl/test-e2e/Regression/group.cpp index 5c1d256550c0c..837878fd1e9fe 100644 --- a/sycl/test-e2e/Regression/group.cpp +++ b/sycl/test-e2e/Regression/group.cpp @@ -40,7 +40,7 @@ bool group__get_group_range() { queue Q(AsyncHandler{}); Q.submit([&](handler &cgh) { - auto Ptr = Buf.get_access(cgh); + auto Ptr = Buf.get_access(cgh); cgh.parallel_for( nd_range<3>{GlobalRange, LocalRange}, [=](nd_item I) { @@ -106,7 +106,7 @@ bool group__get_linear_id() { queue Q(AsyncHandler{}); Q.submit([&](handler &cgh) { - auto Ptr = Buf.get_access(cgh); + auto Ptr = Buf.get_access(cgh); cgh.parallel_for( nd_range<3>{GlobalRange, LocalRange}, [=](nd_item I) { @@ -184,7 +184,7 @@ bool group__async_work_group_copy() { queue Q(AsyncHandler{}); Q.submit([&](handler &cgh) { - auto AccGlobal = Buf.get_access(cgh); + auto AccGlobal = Buf.get_access(cgh); local_accessor AccLocal(LocalRange, cgh); cgh.parallel_for( diff --git a/sycl/test-e2e/Regression/implicit_kernel_bundle_image_filtering.cpp b/sycl/test-e2e/Regression/implicit_kernel_bundle_image_filtering.cpp index ee47211d8e232..f47b54f3f211c 100644 --- a/sycl/test-e2e/Regression/implicit_kernel_bundle_image_filtering.cpp +++ b/sycl/test-e2e/Regression/implicit_kernel_bundle_image_filtering.cpp @@ -23,7 +23,7 @@ int main() { { sycl::buffer Buf(&Ret, 1); Q.submit([&](sycl::handler &CGH) { - auto Acc = Buf.template get_access(CGH); + auto Acc = Buf.template get_access(CGH); CGH.set_specialization_constant(42); CGH.single_task([=](sycl::kernel_handler H) { Acc[0] = H.get_specialization_constant(); diff --git a/sycl/test-e2e/Regression/kernel_bundle_ignore_sycl_external.cpp b/sycl/test-e2e/Regression/kernel_bundle_ignore_sycl_external.cpp index 61b2ca057fa04..f6f2df0b2886f 100644 --- a/sycl/test-e2e/Regression/kernel_bundle_ignore_sycl_external.cpp +++ b/sycl/test-e2e/Regression/kernel_bundle_ignore_sycl_external.cpp @@ -35,7 +35,7 @@ int main() { sycl::buffer Buf(sycl::range<1>{1}); Q.submit([&](sycl::handler &CGH) { - auto Acc = Buf.get_access(CGH); + auto Acc = Buf.get_access(CGH); CGH.use_kernel_bundle(KernelBundle); CGH.single_task([=]() { Acc[0] = 42; }); }); diff --git a/sycl/test-e2e/Regression/local-arg-align.cpp b/sycl/test-e2e/Regression/local-arg-align.cpp index 718aa69f9bff6..71ddb2c9166e3 100644 --- a/sycl/test-e2e/Regression/local-arg-align.cpp +++ b/sycl/test-e2e/Regression/local-arg-align.cpp @@ -32,7 +32,7 @@ int main(int argc, char *argv[]) { local_accessor a(1, h); local_accessor b(1, h); - auto ares = res.get_access(h); + auto ares = res.get_access(h); // Manually capture kernel arguments to ensure an order with the int // argument first and the float4 argument second. If the two arguments are diff --git a/sycl/test-e2e/Regression/multiple-targets.cpp b/sycl/test-e2e/Regression/multiple-targets.cpp index 7ccea010b8e4c..bef32697dba06 100644 --- a/sycl/test-e2e/Regression/multiple-targets.cpp +++ b/sycl/test-e2e/Regression/multiple-targets.cpp @@ -27,7 +27,7 @@ int main() { sycl::buffer A_Buf(A_Data, sycl::range<1>(10)); Q.submit([&](sycl::handler &Cgh) { - auto A_Acc = A_Buf.get_access(Cgh); + auto A_Acc = A_Buf.get_access(Cgh); Cgh.parallel_for(sycl::range<1>{5}, [=](sycl::id<1> index) { A_Acc[index] = 5; }); }); @@ -40,8 +40,8 @@ int main() { sycl::buffer C_Buf(C_Data, sycl::range<1>(10)); Q.submit([&](sycl::handler &Cgh) { - auto B_Acc = B_Buf.get_access(Cgh); - auto C_Acc = C_Buf.get_access(Cgh); + auto B_Acc = B_Buf.get_access(Cgh); + auto C_Acc = C_Buf.get_access(Cgh); Cgh.parallel_for(sycl::range<1>{5}, [=](sycl::id<1> index) { B_Acc[index] += C_Acc[index]; }); @@ -55,8 +55,8 @@ int main() { sycl::buffer C_Buf(C_Data, sycl::range<1>(10)); Q.submit([&](sycl::handler &Cgh) { - auto B_Acc = B_Buf.get_access(Cgh); - auto C_Acc = C_Buf.get_access(Cgh); + auto B_Acc = B_Buf.get_access(Cgh); + auto C_Acc = C_Buf.get_access(Cgh); Cgh.parallel_for(sycl::range<1>{5}, [=](sycl::id<1> index) { C_Acc[index] = B_Acc[index]; }); }); diff --git a/sycl/test-e2e/Regression/private_array_init_test.cpp b/sycl/test-e2e/Regression/private_array_init_test.cpp index 286204bf8e12f..7b54c0531f35b 100644 --- a/sycl/test-e2e/Regression/private_array_init_test.cpp +++ b/sycl/test-e2e/Regression/private_array_init_test.cpp @@ -30,7 +30,7 @@ int main() { { s::buffer buf(&data, s::range<1>(1)); q.submit([&](s::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.single_task([=]() { // Test that SYCL program is not crashed if it contains a class/struct // that has private array member which is initialized by zeroes. diff --git a/sycl/test-e2e/Regression/same_unnamed_kernels.cpp b/sycl/test-e2e/Regression/same_unnamed_kernels.cpp index 13460332777ce..1c3bb044fa5ce 100644 --- a/sycl/test-e2e/Regression/same_unnamed_kernels.cpp +++ b/sycl/test-e2e/Regression/same_unnamed_kernels.cpp @@ -14,7 +14,7 @@ template void run(sycl::queue &q, B &buf, const F &func) { auto e = q.submit([&](sycl::handler &cgh) { - auto acc = buf.template get_access(cgh); + auto acc = buf.template get_access(cgh); cgh.single_task([=]() { func(acc); }); }); e.wait(); @@ -29,11 +29,11 @@ int main() { sycl::buffer bufB(B, 1); run(q, bufA, - [&](const sycl::accessor &acc) { + [&](const sycl::accessor &acc) { acc[0] = 0; }); run(q, bufB, - [&](const sycl::accessor &acc) { + [&](const sycl::accessor &acc) { acc[0] *= 2; }); diff --git a/sycl/test-e2e/Sampler/basic-rw-float.cpp b/sycl/test-e2e/Sampler/basic-rw-float.cpp index 43ec0a2fdade3..c4c3ed36b07a0 100644 --- a/sycl/test-e2e/Sampler/basic-rw-float.cpp +++ b/sycl/test-e2e/Sampler/basic-rw-float.cpp @@ -37,7 +37,7 @@ void test_rw(image_channel_order ChanOrder, image_channel_type ChanType) { // - create an image image<1> image_1D(ChanOrder, ChanType, ImgRange_1D); event E_Setup = Q.submit([&](handler &cgh) { - auto image_acc = image_1D.get_access(cgh); + auto image_acc = image_1D.get_access(cgh); cgh.single_task([=]() { image_acc.write(0, leftEdge); image_acc.write(1, body); @@ -51,8 +51,8 @@ void test_rw(image_channel_order ChanOrder, image_channel_type ChanType) { buffer testResults((range<1>(numTests))); event E_Test = Q.submit([&](handler &cgh) { - auto image_acc = image_1D.get_access(cgh); - auto test_acc = testResults.get_access(cgh); + auto image_acc = image_1D.get_access(cgh); + auto test_acc = testResults.get_access(cgh); cgh.single_task([=]() { int i = 0; // the index for writing into the testResult buffer. diff --git a/sycl/test-e2e/Sampler/normalized-clamp-nearest.cpp b/sycl/test-e2e/Sampler/normalized-clamp-nearest.cpp index 03d6fb2dcbcbb..159d1c2760248 100644 --- a/sycl/test-e2e/Sampler/normalized-clamp-nearest.cpp +++ b/sycl/test-e2e/Sampler/normalized-clamp-nearest.cpp @@ -48,7 +48,7 @@ void test_normalized_clamp_nearest_sampler(image_channel_order ChanOrder, // - create an image image<1> image_1D(ChanOrder, ChanType, ImgRange_1D); event E_Setup = Q.submit([&](handler &cgh) { - auto image_acc = image_1D.get_access(cgh); + auto image_acc = image_1D.get_access(cgh); cgh.single_task([=]() { image_acc.write(0, leftEdge); image_acc.write(1, body); @@ -66,8 +66,8 @@ void test_normalized_clamp_nearest_sampler(image_channel_order ChanOrder, sampler(normalized, addressing_mode::clamp, nearest); event E_Test = Q.submit([&](handler &cgh) { - auto image_acc = image_1D.get_access(cgh); - auto test_acc = testResults.get_access(cgh); + auto image_acc = image_1D.get_access(cgh); + auto test_acc = testResults.get_access(cgh); cgh.single_task([=]() { int i = 0; // the index for writing into the testResult buffer. diff --git a/sycl/test-e2e/Sampler/normalized-mirror-linear-float.cpp b/sycl/test-e2e/Sampler/normalized-mirror-linear-float.cpp index 50b204671f3c8..9387a1216b855 100644 --- a/sycl/test-e2e/Sampler/normalized-mirror-linear-float.cpp +++ b/sycl/test-e2e/Sampler/normalized-mirror-linear-float.cpp @@ -65,7 +65,7 @@ void test_normalized_mirrored_linear_sampler(image_channel_order ChanOrder, // - create an image image<1> image_1D(ChanOrder, ChanType, ImgRange_1D); event E_Setup = Q.submit([&](handler &cgh) { - auto image_acc = image_1D.get_access(cgh); + auto image_acc = image_1D.get_access(cgh); cgh.single_task([=]() { image_acc.write(0, leftEdge); image_acc.write(1, body); @@ -82,8 +82,8 @@ void test_normalized_mirrored_linear_sampler(image_channel_order ChanOrder, auto Norm_Mirror_Linear_sampler = sampler(normalized, mirrored, linear); event E_Test = Q.submit([&](handler &cgh) { - auto image_acc = image_1D.get_access(cgh); - auto test_acc = testResults.get_access(cgh); + auto image_acc = image_1D.get_access(cgh); + auto test_acc = testResults.get_access(cgh); cgh.single_task([=]() { int i = 0; // the index for writing into the testResult buffer. diff --git a/sycl/test-e2e/Sampler/normalized-none-nearest.cpp b/sycl/test-e2e/Sampler/normalized-none-nearest.cpp index ff2ccf987d4df..b08b252fcf928 100644 --- a/sycl/test-e2e/Sampler/normalized-none-nearest.cpp +++ b/sycl/test-e2e/Sampler/normalized-none-nearest.cpp @@ -52,7 +52,7 @@ void test_normalized_none_nearest_sampler(image_channel_order ChanOrder, // - create an image image<1> image_1D(ChanOrder, ChanType, ImgRange_1D); event E_Setup = Q.submit([&](handler &cgh) { - auto image_acc = image_1D.get_access(cgh); + auto image_acc = image_1D.get_access(cgh); cgh.single_task([=]() { image_acc.write(0, leftEdge); image_acc.write(1, body); @@ -69,8 +69,8 @@ void test_normalized_none_nearest_sampler(image_channel_order ChanOrder, auto Norm_None_Nearest_sampler = sampler(normalized, none, nearest); event E_Test = Q.submit([&](handler &cgh) { - auto image_acc = image_1D.get_access(cgh); - auto test_acc = testResults.get_access(cgh); + auto image_acc = image_1D.get_access(cgh); + auto test_acc = testResults.get_access(cgh); cgh.single_task([=]() { int i = 0; // the index for writing into the testResult buffer. diff --git a/sycl/test-e2e/Sampler/unnormalized-clamp-linear-float.cpp b/sycl/test-e2e/Sampler/unnormalized-clamp-linear-float.cpp index 3fd60031f632a..1c725eb35d4e1 100644 --- a/sycl/test-e2e/Sampler/unnormalized-clamp-linear-float.cpp +++ b/sycl/test-e2e/Sampler/unnormalized-clamp-linear-float.cpp @@ -56,7 +56,7 @@ void test_unnormalized_clamp_linear_sampler(image_channel_order ChanOrder, // - create an image image<1> image_1D(ChanOrder, ChanType, ImgRange_1D); event E_Setup = Q.submit([&](handler &cgh) { - auto image_acc = image_1D.get_access(cgh); + auto image_acc = image_1D.get_access(cgh); cgh.single_task([=]() { image_acc.write(0, leftEdge); image_acc.write(1, body); @@ -74,8 +74,8 @@ void test_unnormalized_clamp_linear_sampler(image_channel_order ChanOrder, sampler(unnormalized, addressing_mode::clamp, linear); event E_Test = Q.submit([&](handler &cgh) { - auto image_acc = image_1D.get_access(cgh); - auto test_acc = testResults.get_access(cgh); + auto image_acc = image_1D.get_access(cgh); + auto test_acc = testResults.get_access(cgh); cgh.single_task([=]() { int i = 0; // the index for writing into the testResult buffer. diff --git a/sycl/test-e2e/Sampler/unnormalized-clampedge-nearest.cpp b/sycl/test-e2e/Sampler/unnormalized-clampedge-nearest.cpp index c1e6d0da6c6f8..3fdcc7f372965 100644 --- a/sycl/test-e2e/Sampler/unnormalized-clampedge-nearest.cpp +++ b/sycl/test-e2e/Sampler/unnormalized-clampedge-nearest.cpp @@ -48,7 +48,7 @@ void test_unnormalized_clampedge_nearest_sampler(image_channel_order ChanOrder, // - create an image image<1> image_1D(ChanOrder, ChanType, ImgRange_1D); event E_Setup = Q.submit([&](handler &cgh) { - auto image_acc = image_1D.get_access(cgh); + auto image_acc = image_1D.get_access(cgh); cgh.single_task([=]() { image_acc.write(0, leftEdge); image_acc.write(1, body); @@ -66,8 +66,8 @@ void test_unnormalized_clampedge_nearest_sampler(image_channel_order ChanOrder, sampler(unnormalized, clamp_edge, nearest); event E_Test = Q.submit([&](handler &cgh) { - auto image_acc = image_1D.get_access(cgh); - auto test_acc = testResults.get_access(cgh); + auto image_acc = image_1D.get_access(cgh); + auto test_acc = testResults.get_access(cgh); cgh.single_task([=]() { int i = 0; // the index for writing into the testResult buffer. diff --git a/sycl/test-e2e/Scheduler/BasicSchedulerTests.cpp b/sycl/test-e2e/Scheduler/BasicSchedulerTests.cpp index a644c9c9a95ae..ae4b094732671 100644 --- a/sycl/test-e2e/Scheduler/BasicSchedulerTests.cpp +++ b/sycl/test-e2e/Scheduler/BasicSchedulerTests.cpp @@ -12,7 +12,7 @@ #include -using sycl_access_mode = sycl::access::mode; +using sycl_access_mode = sycl::access_mode; // Execute functor provided passing a queue created with default device selector // and async handler then waits for the tasks submitted to queue to finish. If diff --git a/sycl/test-e2e/Scheduler/DeleteCmdException.cpp b/sycl/test-e2e/Scheduler/DeleteCmdException.cpp index 6ff7fc9c9e89b..82a9441f77fe8 100644 --- a/sycl/test-e2e/Scheduler/DeleteCmdException.cpp +++ b/sycl/test-e2e/Scheduler/DeleteCmdException.cpp @@ -27,7 +27,7 @@ void test_exception(sycl::queue &q, sycl::buffer &buf, // Will throw when submitted q.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.parallel_for(illegal_range, [=](sycl::nd_item<1> nd_item) { acc[nd_item.get_global_linear_id()] = 42; // will not be reached }); diff --git a/sycl/test-e2e/Scheduler/MemObjRemapping.cpp b/sycl/test-e2e/Scheduler/MemObjRemapping.cpp index 55181dfbcc565..d72c7470ba492 100644 --- a/sycl/test-e2e/Scheduler/MemObjRemapping.cpp +++ b/sycl/test-e2e/Scheduler/MemObjRemapping.cpp @@ -21,7 +21,7 @@ int main() { buffer BufA{Range}; Q.submit([&](handler &Cgh) { - auto AccA = BufA.get_access(Cgh); + auto AccA = BufA.get_access(Cgh); Cgh.parallel_for(Range, [=](id<1> Idx) { AccA[Idx] = Idx[0]; }); }); @@ -29,7 +29,7 @@ int main() { // Check access mode flags // CHECK: <--- urEnqueueMemBufferMap // CHECK: .mapFlags = UR_MAP_FLAG_READ - auto AccA = BufA.get_access(); + auto AccA = BufA.get_access(); for (std::size_t I = 0; I < Size; ++I) { assert(AccA[I] == I); } @@ -38,13 +38,13 @@ int main() { // CHECK: <--- urEnqueueMemUnmap // CHECK: <--- urEnqueueMemBufferMap // CHECK: .mapFlags = UR_MAP_FLAG_READ | UR_MAP_FLAG_WRITE - auto AccA = BufA.get_access(); + auto AccA = BufA.get_access(); for (std::size_t I = 0; I < Size; ++I) AccA[I] = 2 * I; } // CHECK-NOT: <--- urEnqueueMemBufferMap - auto AccA = BufA.get_access(); + auto AccA = BufA.get_access(); for (std::size_t I = 0; I < Size; ++I) { assert(AccA[I] == 2 * I); } diff --git a/sycl/test-e2e/SpecConstants/2020/handler-api.cpp b/sycl/test-e2e/SpecConstants/2020/handler-api.cpp index 7090d595fc12f..3149a61b8cd9d 100644 --- a/sycl/test-e2e/SpecConstants/2020/handler-api.cpp +++ b/sycl/test-e2e/SpecConstants/2020/handler-api.cpp @@ -73,11 +73,11 @@ bool test_default_values(sycl::queue q) { sycl::buffer custom_type_buffer(1); q.submit([&](sycl::handler &cgh) { - auto int_acc = int_buffer.get_access(cgh); - auto int_acc2 = int_buffer2.get_access(cgh); - auto float_acc = float_buffer.get_access(cgh); + auto int_acc = int_buffer.get_access(cgh); + auto int_acc2 = int_buffer2.get_access(cgh); + auto float_acc = float_buffer.get_access(cgh); auto custom_type_acc = - custom_type_buffer.get_access(cgh); + custom_type_buffer.get_access(cgh); cgh.single_task([=](sycl::kernel_handler kh) { int_acc[0] = kh.get_specialization_constant(); int_acc2[0] = kh.get_specialization_constant(); @@ -168,11 +168,11 @@ bool test_set_and_get_on_device(sycl::queue q) { custom_type new_custom_type_value('b', 1.0f, 12); q.submit([&](sycl::handler &cgh) { - auto int_acc = int_buffer.get_access(cgh); - auto int_acc2 = int_buffer2.get_access(cgh); - auto float_acc = float_buffer.get_access(cgh); + auto int_acc = int_buffer.get_access(cgh); + auto int_acc2 = int_buffer2.get_access(cgh); + auto float_acc = float_buffer.get_access(cgh); auto custom_type_acc = - custom_type_buffer.get_access(cgh); + custom_type_buffer.get_access(cgh); cgh.set_specialization_constant(new_int_value); cgh.set_specialization_constant(new_int_value2); diff --git a/sycl/test-e2e/SpecConstants/2020/non_native/Inputs/common.cpp b/sycl/test-e2e/SpecConstants/2020/non_native/Inputs/common.cpp index c5d4a3a746c09..5812844f84f65 100644 --- a/sycl/test-e2e/SpecConstants/2020/non_native/Inputs/common.cpp +++ b/sycl/test-e2e/SpecConstants/2020/non_native/Inputs/common.cpp @@ -25,7 +25,7 @@ int main() { sycl::buffer Buf{sycl::range{1}}; Q.submit([&](sycl::handler &CGH) { CGH.set_specialization_constant(1); - auto Acc = Buf.get_access(CGH); + auto Acc = Buf.get_access(CGH); CGH.single_task([=](sycl::kernel_handler KH) { Acc[0] = KH.get_specialization_constant(); }); @@ -37,7 +37,7 @@ int main() { { sycl::buffer Buf{sycl::range{1}}; Q.submit([&](sycl::handler &CGH) { - auto Acc = Buf.get_access(CGH); + auto Acc = Buf.get_access(CGH); CGH.set_specialization_constant(TestStruct{1, 2}); const auto SC = CGH.get_specialization_constant(); assert(SC == 42); diff --git a/sycl/test-e2e/SubGroup/barrier.cpp b/sycl/test-e2e/SubGroup/barrier.cpp index f009b98432b87..3af81618e519b 100644 --- a/sycl/test-e2e/SubGroup/barrier.cpp +++ b/sycl/test-e2e/SubGroup/barrier.cpp @@ -28,8 +28,8 @@ template void check(queue &Queue, size_t G = 240, size_t L = 60) { buffer addbuf(data.data(), range<1>(G)); buffer sgsizebuf(1); Queue.submit([&](handler &cgh) { - auto addacc = addbuf.template get_access(cgh); - auto sgsizeacc = sgsizebuf.get_access(cgh); + auto addacc = addbuf.template get_access(cgh); + auto sgsizeacc = sgsizebuf.get_access(cgh); cgh.parallel_for>(NdRange, [=](nd_item<1> NdItem) { sycl::sub_group SG = NdItem.get_sub_group(); diff --git a/sycl/test-e2e/SubGroup/generic_reduce.cpp b/sycl/test-e2e/SubGroup/generic_reduce.cpp index bca6342a2b0b1..f6d77198257ba 100644 --- a/sycl/test-e2e/SubGroup/generic_reduce.cpp +++ b/sycl/test-e2e/SubGroup/generic_reduce.cpp @@ -27,8 +27,8 @@ void check_op(queue &Queue, T init, BinaryOperation op, bool skip_init = false, buffer buf(G); buffer sgsizebuf(1); Queue.submit([&](handler &cgh) { - auto sgsizeacc = sgsizebuf.get_access(cgh); - auto acc = buf.template get_access(cgh); + auto sgsizeacc = sgsizebuf.get_access(cgh); + auto acc = buf.template get_access(cgh); cgh.parallel_for(NdRange, [=](nd_item<1> NdItem) { auto sg = NdItem.get_sub_group(); if (skip_init) { diff --git a/sycl/test-e2e/SubGroup/scan.hpp b/sycl/test-e2e/SubGroup/scan.hpp index 241b5d8a711e0..804bb8205533a 100644 --- a/sycl/test-e2e/SubGroup/scan.hpp +++ b/sycl/test-e2e/SubGroup/scan.hpp @@ -22,9 +22,9 @@ void check_op(queue &Queue, T init, BinaryOperation op, bool skip_init = false, buffer exbuf(G), inbuf(G); buffer sgsizebuf(1); Queue.submit([&](handler &cgh) { - auto sgsizeacc = sgsizebuf.get_access(cgh); - auto exacc = exbuf.template get_access(cgh); - auto inacc = inbuf.template get_access(cgh); + auto sgsizeacc = sgsizebuf.get_access(cgh); + auto exacc = exbuf.template get_access(cgh); + auto inacc = inbuf.template get_access(cgh); cgh.parallel_for( NdRange, [=](nd_item<1> NdItem) { sycl::sub_group sg = NdItem.get_sub_group(); diff --git a/sycl/test-e2e/SubGroupMask/GroupSize.cpp b/sycl/test-e2e/SubGroupMask/GroupSize.cpp index ff8f7e7fec634..aa7760df51f66 100644 --- a/sycl/test-e2e/SubGroupMask/GroupSize.cpp +++ b/sycl/test-e2e/SubGroupMask/GroupSize.cpp @@ -38,7 +38,7 @@ template void test(queue Queue) { buffer resbuf(Res, range<1>(32 / SGSize)); Queue.submit([&](handler &cgh) { - auto resacc = resbuf.template get_access(cgh); + auto resacc = resbuf.template get_access(cgh); cgh.parallel_for>( NdRange, diff --git a/sycl/test-e2e/USM/fill.cpp b/sycl/test-e2e/USM/fill.cpp index c4033e4672591..1afcdd4d0c5d4 100644 --- a/sycl/test-e2e/USM/fill.cpp +++ b/sycl/test-e2e/USM/fill.cpp @@ -113,7 +113,7 @@ void runDeviceTests(device dev, context ctxt, queue q, T val) { { buffer buf{&out[0], range<1>{N}}; q.submit([&](handler &h) { - auto acc = buf.template get_access(h); + auto acc = buf.template get_access(h); h.parallel_for>( range<1>(N), [=](id<1> item) { acc[item] = array[item]; }); }).wait(); @@ -134,7 +134,7 @@ void runDeviceTests(device dev, context ctxt, queue q, T val) { { buffer buf{&out[0], range<1>{N}}; q.submit([&](handler &h) { - auto acc = buf.template get_access(h); + auto acc = buf.template get_access(h); h.parallel_for>( range<1>(N), [=](id<1> item) { acc[item] = array[item]; }); }).wait(); diff --git a/sycl/test-e2e/USM/memcpy.cpp b/sycl/test-e2e/USM/memcpy.cpp index c12dcde5fcb66..65a8094cf921d 100644 --- a/sycl/test-e2e/USM/memcpy.cpp +++ b/sycl/test-e2e/USM/memcpy.cpp @@ -43,7 +43,7 @@ void check_on_device(queue q, int *arr) { { buffer buf{&out[0], range<1>{N}}; q.submit([&](handler &h) { - auto acc = buf.template get_access(h); + auto acc = buf.template get_access(h); h.parallel_for( range<1>(N), [=](id<1> item) { acc[item] = arr[item]; }); }).wait(); diff --git a/sycl/test-e2e/WorkGroupScratchMemory/dynamic_alloc_local_accessor.cpp b/sycl/test-e2e/WorkGroupScratchMemory/dynamic_alloc_local_accessor.cpp index f84f09ba75b36..7d70ff592f90d 100644 --- a/sycl/test-e2e/WorkGroupScratchMemory/dynamic_alloc_local_accessor.cpp +++ b/sycl/test-e2e/WorkGroupScratchMemory/dynamic_alloc_local_accessor.cpp @@ -69,7 +69,7 @@ int main() { buffer Buf{Vec.data(), range<1>(Size)}; Q.submit([&](handler &Cgh) { - auto Acc = Buf.get_access(Cgh); + auto Acc = Buf.get_access(Cgh); sycl_ext::work_group_scratch_size static_size(WgSize * RepeatWG * sizeof(int)); sycl_ext::properties properties{static_size}; diff --git a/sycl/test-e2e/WorkGroupStatic/work_group_static_memory_block_scope.cpp b/sycl/test-e2e/WorkGroupStatic/work_group_static_memory_block_scope.cpp index b09327ff7f3e0..a6db7a906952b 100644 --- a/sycl/test-e2e/WorkGroupStatic/work_group_static_memory_block_scope.cpp +++ b/sycl/test-e2e/WorkGroupStatic/work_group_static_memory_block_scope.cpp @@ -42,7 +42,7 @@ int main() { buffer Buf{Vec.data(), range<1>(Size)}; Q.submit([&](handler &Cgh) { - auto Acc = Buf.get_access(Cgh); + auto Acc = Buf.get_access(Cgh); Cgh.parallel_for( nd_range<1>(range<1>(Size), range<1>(WgSize)), [=](nd_item<1> Item) { sycl::ext::oneapi::experimental::work_group_static @@ -68,7 +68,7 @@ int main() { buffer Buf{Vec.data(), range<1>(Size)}; Q.submit([&](handler &Cgh) { - auto Acc = Buf.get_access(Cgh); + auto Acc = Buf.get_access(Cgh); Cgh.parallel_for( nd_range<1>(range<1>(Size), range<1>(WgSize)), [=](nd_item<1> Item) { sycl::ext::oneapi::experimental::work_group_static localIDBuff; diff --git a/sycl/test-e2e/XPTI/buffer/accessors.cpp b/sycl/test-e2e/XPTI/buffer/accessors.cpp index 2bac9a54b2403..7adbc6d5ba1d9 100644 --- a/sycl/test-e2e/XPTI/buffer/accessors.cpp +++ b/sycl/test-e2e/XPTI/buffer/accessors.cpp @@ -26,17 +26,17 @@ int main() { Queue.submit([&](sycl::handler &cgh) { // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID1:.+]]|2015|1024|{{.*}}accessors.cpp:[[# @LINE + 1]]:19 - auto A1 = Buf.get_access(cgh); + auto A1 = Buf.get_access(cgh); // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID2:.*]]|2014|1025|{{.*}}accessors.cpp:[[# @LINE + 1]]:19 - auto A2 = Buf.get_access(cgh); + auto A2 = Buf.get_access(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]]:19 - auto A4 = Buf.get_access(cgh); + auto A4 = Buf.get_access(cgh); // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID5:.*]]|2014|1028|{{.*}}accessors.cpp:[[# @LINE + 1]]:19 - auto A5 = Buf.get_access(cgh); + auto A5 = Buf.get_access(cgh); // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID6:.*]]|2014|1029|{{.*}}accessors.cpp:[[# @LINE + 1]]:19 - auto A6 = Buf.get_access(cgh); + auto A6 = Buf.get_access(cgh); cgh.parallel_for(NDRange, [=](sycl::nd_item<1>) { (void)A1; (void)A2; @@ -53,11 +53,11 @@ int main() { // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID9:.*]]|2018|1026|{{.*}}accessors.cpp:[[# @LINE + 1]]:25 { sycl::host_accessor HA(Buf, sycl::read_write); } // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID10:.*]]|2018|1027|{{.*}}accessors.cpp:[[# @LINE + 1]]:19 - { auto HA = Buf.get_access(); } + { auto HA = Buf.get_access(); } // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID11:.*]]|2018|1028|{{.*}}accessors.cpp:[[# @LINE + 1]]:19 - { auto HA = Buf.get_access(); } + { auto HA = Buf.get_access(); } // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID12:.*]]|2018|1029|{{.*}}accessors.cpp:[[# @LINE + 1]]:19 - { auto HA = Buf.get_access(); } + { auto HA = Buf.get_access(); } return 0; } diff --git a/sycl/test-e2e/XPTI/buffer/multiple_buffers.cpp b/sycl/test-e2e/XPTI/buffer/multiple_buffers.cpp index a7b5358007698..bd9db0ff85953 100644 --- a/sycl/test-e2e/XPTI/buffer/multiple_buffers.cpp +++ b/sycl/test-e2e/XPTI/buffer/multiple_buffers.cpp @@ -26,8 +26,8 @@ int main() { // CHECK:{{[0-9]+}}|Associate buffer|[[USERID2]]|[[BEID2:.*]] Queue.submit([&](sycl::handler &cgh) { // Get write only access to the buffer on a device. - auto Accessor1 = Buffer1.get_access(cgh); - auto Accessor2 = Buffer2.get_access(cgh); + auto Accessor1 = Buffer1.get_access(cgh); + auto Accessor2 = Buffer2.get_access(cgh); // Execute kernel. cgh.parallel_for(NumOfWorkItems, [=](sycl::id<1> WIid) { Accessor1[WIid] = static_cast(WIid.get(0)); diff --git a/sycl/test-e2e/XPTI/buffer/sub_buffer.cpp b/sycl/test-e2e/XPTI/buffer/sub_buffer.cpp index 0a225a138981c..157645b8d951a 100644 --- a/sycl/test-e2e/XPTI/buffer/sub_buffer.cpp +++ b/sycl/test-e2e/XPTI/buffer/sub_buffer.cpp @@ -25,7 +25,7 @@ int main() { Queue.submit([&](sycl::handler &cgh) { // CHECK: {{[0-9]+}}|Construct accessor|[[USERID1]]|[[ACCID1:.*]]|2014|1025|{{.*}}sub_buffer.cpp:[[# @LINE + 1]]:34 - auto Accessor1 = SubBuffer.get_access(cgh); + auto Accessor1 = SubBuffer.get_access(cgh); // CHECK:{{[0-9]+}}|Associate buffer|[[USERID1]]|[[BEID1:.*]] // CHECK:{{[0-9]+}}|Associate buffer|[[USERID1]]|[[BEID2:.*]] cgh.parallel_for( diff --git a/sycl/test-e2e/XPTI/kernel/content.cpp b/sycl/test-e2e/XPTI/kernel/content.cpp index f17e2defa5893..3df834749b5e7 100644 --- a/sycl/test-e2e/XPTI/kernel/content.cpp +++ b/sycl/test-e2e/XPTI/kernel/content.cpp @@ -46,8 +46,8 @@ int main() { buffer in_buf(input.data(), input.size()); buffer out_buf(input.size()); myQueue.submit([&](handler &cgh) { - auto in = in_buf.template get_access(cgh); - auto out = out_buf.template get_access(cgh); + auto in = in_buf.template get_access(cgh); + auto out = out_buf.template get_access(cgh); // CHECK-OPT:Node create|{{.*}}test2{{.*}}|{{.*}}.cpp:[[# @LINE - 3 ]]:13|{128, 4, 2}, {32, 2, 1}, {16, 1, 0}, 2 // CHECK-NOOPT:Node create|{{.*}}test2{{.*}}|{{.*}}.cpp:[[# @LINE - 4 ]]:13|{128, 4, 2}, {32, 2, 1}, {16, 1, 0}, 8 cgh.parallel_for( From 33411b85dd29b5656c23aee797777fa1a76a2c05 Mon Sep 17 00:00:00 2001 From: "Kornev, Nikita" Date: Fri, 31 Jul 2026 16:10:26 +0200 Subject: [PATCH 2/4] format --- .../Basic/buffer/buffer_full_copy.cpp | 2 +- sycl/test-e2e/Basic/host-task-dependency.cpp | 9 ++++----- sycl/test-e2e/Basic/image/image_sample.cpp | 6 +++--- sycl/test-e2e/XPTI/buffer/accessors.cpp | 19 ++++++++++++++----- 4 files changed, 22 insertions(+), 14 deletions(-) diff --git a/sycl/test-e2e/Basic/buffer/buffer_full_copy.cpp b/sycl/test-e2e/Basic/buffer/buffer_full_copy.cpp index 2333a8865909a..2041955eb7b9d 100644 --- a/sycl/test-e2e/Basic/buffer/buffer_full_copy.cpp +++ b/sycl/test-e2e/Basic/buffer/buffer_full_copy.cpp @@ -70,7 +70,7 @@ void check_fill(sycl::queue &Queue) { auto e = Queue.submit([&](sycl::handler &cgh) { auto a = buf_1.template get_access(cgh, size / 2, - offset); + offset); cgh.fill(a, (float)1337.0); }); e.wait(); diff --git a/sycl/test-e2e/Basic/host-task-dependency.cpp b/sycl/test-e2e/Basic/host-task-dependency.cpp index 50839dfb62818..218db85007b06 100644 --- a/sycl/test-e2e/Basic/host-task-dependency.cpp +++ b/sycl/test-e2e/Basic/host-task-dependency.cpp @@ -27,10 +27,9 @@ struct Context { S::event HostTask_CopyBuf1ToBuf2(Context *Ctx) { S::event Event = Ctx->Queue.submit([&](S::handler &CGH) { - S::host_accessor CopierSrcAcc(Ctx->Buf1, - CGH); + S::host_accessor CopierSrcAcc(Ctx->Buf1, CGH); S::host_accessor CopierDstAcc(Ctx->Buf2, - CGH); + CGH); auto CopierHostTask = [=] { for (size_t Idx = 0; Idx < CopierDstAcc.size(); ++Idx) @@ -92,8 +91,8 @@ void Thread1Fn(Context *Ctx) { // 3. submit simple task to move data between two buffers Ctx->Queue.submit([&](S::handler &CGH) { - S::accessor - SrcAcc(Ctx->Buf2, CGH); + S::accessor SrcAcc( + Ctx->Buf2, CGH); S::accessor DstAcc(Ctx->Buf3, CGH); diff --git a/sycl/test-e2e/Basic/image/image_sample.cpp b/sycl/test-e2e/Basic/image/image_sample.cpp index 1bfa52c5152f4..78dc940303642 100644 --- a/sycl/test-e2e/Basic/image/image_sample.cpp +++ b/sycl/test-e2e/Basic/image/image_sample.cpp @@ -53,7 +53,7 @@ bool test1d_coord(dataT *hostPtr, coordT coord, dataT expectedColour) { myQueue.submit([&](s::handler &cgh) { auto imageAcc = image.get_access(cgh); s::accessor resultDataAcc(resultDataBuf, - cgh); + cgh); cgh.single_task([=]() { dataT RetColor = imageAcc.read(coord, testSampler); @@ -88,7 +88,7 @@ bool test2d_coord(dataT *hostPtr, coordT coord, dataT expectedColour) { myQueue.submit([&](s::handler &cgh) { auto imageAcc = image.get_access(cgh); s::accessor resultDataAcc(resultDataBuf, - cgh); + cgh); cgh.single_task([=]() { dataT RetColor = imageAcc.read(coord, testSampler); @@ -123,7 +123,7 @@ bool test3d_coord(dataT *hostPtr, coordT coord, dataT expectedColour) { myQueue.submit([&](s::handler &cgh) { auto imageAcc = image.get_access(cgh); s::accessor resultDataAcc(resultDataBuf, - cgh); + cgh); cgh.single_task([=]() { dataT RetColor = imageAcc.read(coord, testSampler); diff --git a/sycl/test-e2e/XPTI/buffer/accessors.cpp b/sycl/test-e2e/XPTI/buffer/accessors.cpp index 7adbc6d5ba1d9..993710b12c27c 100644 --- a/sycl/test-e2e/XPTI/buffer/accessors.cpp +++ b/sycl/test-e2e/XPTI/buffer/accessors.cpp @@ -26,7 +26,8 @@ int main() { Queue.submit([&](sycl::handler &cgh) { // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID1:.+]]|2015|1024|{{.*}}accessors.cpp:[[# @LINE + 1]]:19 - auto A1 = Buf.get_access(cgh); + auto A1 = + Buf.get_access(cgh); // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID2:.*]]|2014|1025|{{.*}}accessors.cpp:[[# @LINE + 1]]:19 auto A2 = Buf.get_access(cgh); // CHECK: {{[0-9]+}}|Construct accessor|0x0|[[ACCID3:.*]]|2016|1026|{{.*}}accessors.cpp:[[# @LINE + 1]]:34 @@ -34,7 +35,9 @@ int main() { // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID4:.*]]|2014|1027|{{.*}}accessors.cpp:[[# @LINE + 1]]:19 auto A4 = Buf.get_access(cgh); // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID5:.*]]|2014|1028|{{.*}}accessors.cpp:[[# @LINE + 1]]:19 - auto A5 = Buf.get_access(cgh); + auto A5 = + Buf.get_access( + cgh); // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID6:.*]]|2014|1029|{{.*}}accessors.cpp:[[# @LINE + 1]]:19 auto A6 = Buf.get_access(cgh); cgh.parallel_for(NDRange, [=](sycl::nd_item<1>) { @@ -53,11 +56,17 @@ int main() { // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID9:.*]]|2018|1026|{{.*}}accessors.cpp:[[# @LINE + 1]]:25 { sycl::host_accessor HA(Buf, sycl::read_write); } // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID10:.*]]|2018|1027|{{.*}}accessors.cpp:[[# @LINE + 1]]:19 - { auto HA = Buf.get_access(); } + { + auto HA = Buf.get_access(); + } // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID11:.*]]|2018|1028|{{.*}}accessors.cpp:[[# @LINE + 1]]:19 - { auto HA = Buf.get_access(); } + { + auto HA = Buf.get_access(); + } // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID12:.*]]|2018|1029|{{.*}}accessors.cpp:[[# @LINE + 1]]:19 - { auto HA = Buf.get_access(); } + { + auto HA = Buf.get_access(); + } return 0; } From e13a11d1364f43721c66e3197a410d4446e27c5b Mon Sep 17 00:00:00 2001 From: "Kornev, Nikita" Date: Fri, 31 Jul 2026 16:21:44 +0200 Subject: [PATCH 3/4] fix XPTI/buffer/accessors.cpp after format --- sycl/test-e2e/XPTI/buffer/accessors.cpp | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/sycl/test-e2e/XPTI/buffer/accessors.cpp b/sycl/test-e2e/XPTI/buffer/accessors.cpp index 993710b12c27c..f1a66de4185b6 100644 --- a/sycl/test-e2e/XPTI/buffer/accessors.cpp +++ b/sycl/test-e2e/XPTI/buffer/accessors.cpp @@ -25,7 +25,7 @@ int main() { sycl::nd_range<1> NDRange(Buf.size(), Buf.size()); Queue.submit([&](sycl::handler &cgh) { - // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID1:.+]]|2015|1024|{{.*}}accessors.cpp:[[# @LINE + 1]]:19 + // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID1:.+]]|2015|1024|{{.*}}accessors.cpp:[[# @LINE + 2]]:13 auto A1 = Buf.get_access(cgh); // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID2:.*]]|2014|1025|{{.*}}accessors.cpp:[[# @LINE + 1]]:19 @@ -34,7 +34,7 @@ int main() { sycl::local_accessor A3(Range, cgh); // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID4:.*]]|2014|1027|{{.*}}accessors.cpp:[[# @LINE + 1]]:19 auto A4 = Buf.get_access(cgh); - // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID5:.*]]|2014|1028|{{.*}}accessors.cpp:[[# @LINE + 1]]:19 + // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID5:.*]]|2014|1028|{{.*}}accessors.cpp:[[# @LINE + 2]]:13 auto A5 = Buf.get_access( cgh); @@ -55,15 +55,15 @@ int main() { { sycl::host_accessor HA(Buf, sycl::write_only); } // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID9:.*]]|2018|1026|{{.*}}accessors.cpp:[[# @LINE + 1]]:25 { sycl::host_accessor HA(Buf, sycl::read_write); } - // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID10:.*]]|2018|1027|{{.*}}accessors.cpp:[[# @LINE + 1]]:19 + // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID10:.*]]|2018|1027|{{.*}}accessors.cpp:[[# @LINE + 2]]:19 { auto HA = Buf.get_access(); } - // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID11:.*]]|2018|1028|{{.*}}accessors.cpp:[[# @LINE + 1]]:19 + // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID11:.*]]|2018|1028|{{.*}}accessors.cpp:[[# @LINE + 2]]:19 { auto HA = Buf.get_access(); } - // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID12:.*]]|2018|1029|{{.*}}accessors.cpp:[[# @LINE + 1]]:19 + // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID12:.*]]|2018|1029|{{.*}}accessors.cpp:[[# @LINE + 2]]:19 { auto HA = Buf.get_access(); } From 5d598d89768b04c1895eedfe27e5a8bb2713518e Mon Sep 17 00:00:00 2001 From: "Kornev, Nikita" Date: Fri, 31 Jul 2026 17:49:12 +0200 Subject: [PATCH 4/4] drop sycl:: where possible --- sycl/test-e2e/Assert/assert_in_kernels.hpp | 6 +++--- .../assert_in_simultaneously_multiple_tus.cpp | 2 +- .../HostInteropTask/host-task-dependency3.cpp | 4 ++-- sycl/test-e2e/HostInteropTask/host-task.cpp | 8 ++++---- sycl/test-e2e/HostInteropTask/interop-task.cpp | 16 ++++++++-------- 5 files changed, 18 insertions(+), 18 deletions(-) diff --git a/sycl/test-e2e/Assert/assert_in_kernels.hpp b/sycl/test-e2e/Assert/assert_in_kernels.hpp index 75813f9d55d6d..2181f710c382d 100644 --- a/sycl/test-e2e/Assert/assert_in_kernels.hpp +++ b/sycl/test-e2e/Assert/assert_in_kernels.hpp @@ -14,7 +14,7 @@ void kernelFunc1(int *Buf, int wiID) { void assertTest1(queue &Q, buffer &Buf) { Q.submit([&](handler &CGH) { - auto Acc = Buf.template get_access(CGH); + auto Acc = Buf.template get_access(CGH); CGH.parallel_for( Buf.get_range(), [=](sycl::id<1> wiID) { kernelFunc1(&Acc[0], wiID); }); @@ -29,7 +29,7 @@ void kernelFunc2(int *Buf, int wiID) { void assertTest2(queue &Q, buffer &Buf) { Q.submit([&](handler &CGH) { - auto Acc = Buf.template get_access(CGH); + auto Acc = Buf.template get_access(CGH); CGH.parallel_for( Buf.get_range(), [=](sycl::id<1> wiID) { kernelFunc2(&Acc[0], wiID); }); @@ -44,7 +44,7 @@ void kernelFunc3(int *Buf, int wiID) { void assertTest3(queue &Q, buffer &Buf) { Q.submit([&](handler &CGH) { - auto Acc = Buf.template get_access(CGH); + auto Acc = Buf.template get_access(CGH); CGH.parallel_for( Buf.get_range(), [=](sycl::id<1> wiID) { kernelFunc3(&Acc[0], wiID); }); diff --git a/sycl/test-e2e/Assert/assert_in_simultaneously_multiple_tus.cpp b/sycl/test-e2e/Assert/assert_in_simultaneously_multiple_tus.cpp index 3f267230f96ae..a4e45aad405f0 100644 --- a/sycl/test-e2e/Assert/assert_in_simultaneously_multiple_tus.cpp +++ b/sycl/test-e2e/Assert/assert_in_simultaneously_multiple_tus.cpp @@ -52,7 +52,7 @@ template void enqueueKernel(queue *Q) { sycl::buffer Buf(numOfItems); Q->submit([&](handler &CGH) { - auto Acc = Buf.template get_access(CGH); + auto Acc = Buf.template get_access(CGH); CGH.parallel_for(numOfItems, [=](sycl::id<1> wiID) { Acc[wiID] = 0; diff --git a/sycl/test-e2e/HostInteropTask/host-task-dependency3.cpp b/sycl/test-e2e/HostInteropTask/host-task-dependency3.cpp index 17b2c6ecd8053..8a31542b8d46b 100644 --- a/sycl/test-e2e/HostInteropTask/host-task-dependency3.cpp +++ b/sycl/test-e2e/HostInteropTask/host-task-dependency3.cpp @@ -57,7 +57,7 @@ void test(size_t Count) { Q.submit([&](handler &CGH) { std::cout << "Submit Kernel-1" << std::endl; - auto Acc0 = B0.get_access(CGH); + auto Acc0 = B0.get_access(CGH); CGH.single_task([=] { Acc0[1] = 1 * Idx; }); }); @@ -65,7 +65,7 @@ void test(size_t Count) { Q.submit([&](handler &CGH) { std::cout << "Submit Kernel-2" << std::endl; - auto Acc1 = B1.get_access(CGH); + auto Acc1 = B1.get_access(CGH); CGH.single_task([=] { Acc1[2] = 1 * Idx; }); }); diff --git a/sycl/test-e2e/HostInteropTask/host-task.cpp b/sycl/test-e2e/HostInteropTask/host-task.cpp index 779aaf189810a..4fae53e8d61ac 100644 --- a/sycl/test-e2e/HostInteropTask/host-task.cpp +++ b/sycl/test-e2e/HostInteropTask/host-task.cpp @@ -28,7 +28,7 @@ void test1(queue &Q) { buffer Buffer{BUFFER_SIZE}; Q.submit([&](handler &CGH) { - auto Acc = Buffer.get_access(CGH); + auto Acc = Buffer.get_access(CGH); CGH.host_task([=] { /* A no-op */ }); }); @@ -41,15 +41,15 @@ void test2(queue &Q) { buffer Buffer2{BUFFER_SIZE}; Q.submit([&](handler &CGH) { - auto Acc = Buffer1.template get_access(CGH); + auto Acc = Buffer1.template get_access(CGH); auto Kernel = [=](item<1> Id) { Acc[Id] = 123; }; CGH.parallel_for>(Acc.size(), Kernel); }); Q.submit([&](handler &CGH) { - auto AccSrc = Buffer1.template get_access(CGH); - auto AccDst = Buffer2.template get_access(CGH); + auto AccSrc = Buffer1.template get_access(CGH); + auto AccDst = Buffer2.template get_access(CGH); auto Func = [=] { for (size_t Idx = 0; Idx < AccDst.size(); ++Idx) diff --git a/sycl/test-e2e/HostInteropTask/interop-task.cpp b/sycl/test-e2e/HostInteropTask/interop-task.cpp index 203baea4ca5ae..f4c0e32a6ef1b 100644 --- a/sycl/test-e2e/HostInteropTask/interop-task.cpp +++ b/sycl/test-e2e/HostInteropTask/interop-task.cpp @@ -34,8 +34,8 @@ void checkBufferValues(BufferT Buffer, ValueT Value) { template void copy(buffer &Src, buffer &Dst, queue &Q) { Q.submit([&](handler &CGH) { - auto SrcA = Src.template get_access(CGH); - auto DstA = Dst.template get_access(CGH); + auto SrcA = Src.template get_access(CGH); + auto DstA = Dst.template get_access(CGH); auto Func = [=](interop_handle IH) { auto NativeQ = IH.get_native_queue(); @@ -60,7 +60,7 @@ void copy(buffer &Src, buffer &Dst, queue &Q) { template void modify(buffer &B, queue &Q) { Q.submit([&](handler &CGH) { - auto Acc = B.template get_access(CGH); + auto Acc = B.template get_access(CGH); auto Kernel = [=](item<1> Id) { Acc[Id] += 1; }; @@ -71,8 +71,8 @@ template void modify(buffer &B, queue &Q) { template void init(buffer &B1, buffer &B2, queue &Q) { Q.submit([&](handler &CGH) { - auto Acc1 = B1.template get_access(CGH); - auto Acc2 = B2.template get_access(CGH); + auto Acc1 = B1.template get_access(CGH); + auto Acc2 = B2.template get_access(CGH); CGH.parallel_for>(BUFFER_SIZE, [=](item<1> Id) { Acc1[Id] = -1; @@ -156,7 +156,7 @@ void test3(queue &Q) { buffer Buffer{BUFFER_SIZE * 128}; event Event = Q.submit([&](handler &CGH) { - auto Acc1 = Buffer.get_access(CGH); + auto Acc1 = Buffer.get_access(CGH); CGH.parallel_for(BUFFER_SIZE, [=](item<1> Id) { Acc1[Id] = 123; }); @@ -178,7 +178,7 @@ void test4(queue &Q) { buffer Buffer{BUFFER_SIZE}; Q.submit([&](handler &CGH) { - auto Acc = Buffer.get_access(CGH); + auto Acc = Buffer.get_access(CGH); auto Func = [=](interop_handle IH) { /*A no-op */ }; CGH.host_task(Func); }); @@ -189,7 +189,7 @@ void test5(queue &Q) { buffer Buffer2{BUFFER_SIZE}; Q.submit([&](handler &CGH) { - auto Acc = Buffer1.template get_access(CGH); + auto Acc = Buffer1.template get_access(CGH); auto Kernel = [=](item<1> Id) { Acc[Id] = 123; }; CGH.parallel_for(Acc.get_count(), Kernel);