From 6961a1e94b5f4e85b402a3a7160c644f4f0abe58 Mon Sep 17 00:00:00 2001 From: "Kornev, Nikita" Date: Fri, 31 Jul 2026 15:54:13 +0200 Subject: [PATCH 1/2] [SYCL] Update E2E to use access_mode [1/3] access::mode is deprecated in favor of access_mode. PR to add deprecations: https://github.com/intel/llvm/pull/22803 --- .../khr_dynamic_addrspace_cast.cpp | 4 +- .../Assert/Inputs/kernels_in_file2.cpp | 4 +- sycl/test-e2e/Assert/assert_in_one_kernel.hpp | 2 +- sycl/test-e2e/AtomicRef/and.h | 8 +- .../AtomicRef/atomic_memory_order_seq_cst.cpp | 20 ++--- sycl/test-e2e/AtomicRef/load.h | 8 +- sycl/test-e2e/AtomicRef/or.h | 8 +- sycl/test-e2e/AtomicRef/xor.h | 8 +- .../Basic/alloc_pinned_host_memory.cpp | 4 +- .../Basic/buffer/buffer_container.cpp | 2 +- .../buffer/native_buffer_creation_flags.cpp | 2 +- sycl/test-e2e/Basic/built-ins.cpp | 4 +- sycl/test-e2e/Basic/device_event.cpp | 2 +- sycl/test-e2e/Basic/event_release.cpp | 2 +- .../free_function_queries.cpp | 24 +++--- sycl/test-e2e/Basic/group_async_copy.cpp | 4 +- sycl/test-e2e/Basic/half_type.cpp | 26 +++--- .../test-e2e/Basic/handler/handler_mem_op.cpp | 80 +++++++++---------- .../Basic/image/image_accessor_range.cpp | 8 +- .../image/image_accessor_readwrite_half.cpp | 6 +- sycl/test-e2e/Basic/image/image_read.h | 18 ++--- sycl/test-e2e/Basic/image/image_write.h | 24 +++--- .../Basic/kernel_bundle/kernel_bundle_api.cpp | 2 +- sycl/test-e2e/Basic/kernel_info_attr.cpp | 2 +- sycl/test-e2e/Basic/multi_ptr_legacy.hpp | 10 +-- sycl/test-e2e/Basic/multisource_spv_obj.cpp | 8 +- sycl/test-e2e/Basic/parallel_for_indexers.cpp | 8 +- .../Basic/range_offset_fit_in_int.cpp | 24 +++--- sycl/test-e2e/Basic/swizzle_op.cpp | 22 ++--- .../test-e2e/Basic/use_pinned_host_memory.cpp | 2 +- .../enqueueLaunchCustom_check_event_deps.cpp | 6 +- .../discard_events_accessors.cpp | 2 +- .../device_architecture_on_device_aot.cpp | 8 +- .../device_global_operator_passthrough.hpp | 2 +- .../DeviceLib/ITTAnnotations/atomic.cpp | 4 +- .../DeviceLib/bfloat16_conversion_test.hpp | 4 +- sycl/test-e2e/DeviceLib/built-ins/nan.cpp | 4 +- .../DeviceLib/built-ins/scalar_integer.cpp | 66 +++++++-------- .../DeviceLib/built-ins/scalar_relational.cpp | 54 ++++++------- .../DeviceLib/built-ins/vector_integer.cpp | 68 ++++++++-------- .../DeviceLib/c99_complex_math_fp64_test.cpp | 2 +- sycl/test-e2e/DeviceLib/cmath_test.cpp | 4 +- sycl/test-e2e/DeviceLib/imf/fp64_test.cpp | 4 +- .../DeviceLib/math_fp64_windows_test.cpp | 4 +- .../DeviceLib/math_test_marray_vec_common.hpp | 24 +++--- .../DeviceLib/std_complex_math_fp64_test.cpp | 4 +- .../DotProduct/dot_product_int_test.cpp | 32 ++++---- .../custom-command-hip.cpp | 12 +-- .../launch_queries/max_sub_groups.cpp | 4 +- .../cross_queue_dependency.cpp | 8 +- .../reusable_events/event_reuse.cpp | 6 +- .../out_of_order_queue_signal.cpp | 2 +- .../accessor_as_kernel_parameter.cpp | 12 +-- sycl/test-e2e/Functor/functor_inheritance.cpp | 2 +- sycl/test-e2e/GroupAlgorithm/barrier.cpp | 8 +- .../GroupLocalMemory/group_local_memory.cpp | 8 +- sycl/test-e2e/HierPar/hier_par_basic.cpp | 12 +-- .../HostInteropTask/host-task-failure.cpp | 8 +- .../HostInteropTask/interop-task-hip.cpp | 12 +-- .../test-e2e/InlineAsm/asm_16_matrix_mult.cpp | 2 +- sycl/test-e2e/InlineAsm/asm_8_empty.cpp | 2 +- sycl/test-e2e/InlineAsm/asm_decl_in_scope.cpp | 6 +- sycl/test-e2e/InlineAsm/asm_float_neg.cpp | 4 +- sycl/test-e2e/InlineAsm/asm_loop.cpp | 6 +- sycl/test-e2e/InlineAsm/asm_no_output.cpp | 2 +- sycl/test-e2e/InorderQueue/in_order_buffs.cpp | 6 +- sycl/test-e2e/KernelCompiler/opencl.cpp | 4 +- sycl/test-e2e/KernelCompiler/spirv.cpp | 6 +- .../KernelParams/non-standard-layout.cpp | 2 +- ...-for-out-of-registers-on-kernel-launch.cpp | 10 +-- .../Reduction/reduction_nd_N_vars.cpp | 4 +- .../Reduction/reduction_nd_lambda.cpp | 2 +- .../Reduction/reduction_range_N_vars.cpp | 4 +- sycl/test-e2e/Reduction/reduction_utils.hpp | 6 +- sycl/test-e2e/Regression/build_log.cpp | 66 +++++++-------- .../test-e2e/Regression/event_destruction.cpp | 2 +- .../Regression/fsycl-host-compiler.cpp | 2 +- .../Regression/implicit_atomic_conversion.cpp | 2 +- sycl/test-e2e/Regression/isordered.cpp | 2 +- sycl/test-e2e/Regression/kernel_unnamed.cpp | 6 +- sycl/test-e2e/Regression/msvc_crt.cpp | 2 +- .../nontrivial_device_copyable_value.cpp | 2 +- .../Regression/reduction_resource_leak_dw.cpp | 2 +- .../Regression/static-buffer-dtor.cpp | 4 +- .../Sampler/normalized-clamp-linear-float.cpp | 6 +- .../Sampler/normalized-clampedge-nearest.cpp | 6 +- .../Sampler/normalized-none-linear-float.cpp | 6 +- .../Sampler/normalized-repeat-nearest.cpp | 6 +- .../unnormalized-clampedge-linear-float.cpp | 6 +- .../Sampler/unnormalized-none-nearest.cpp | 6 +- sycl/test-e2e/Scheduler/DataMovement.cpp | 2 +- sycl/test-e2e/Scheduler/InOrderQueueDeps.cpp | 2 +- .../Scheduler/ReleaseResourcesTest.cpp | 2 +- .../SpecConstants/2020/kernel-bundle-api.cpp | 12 +-- .../2020/vector-convolution-demo.cpp | 4 +- sycl/test-e2e/SubGroup/common.cpp | 4 +- sycl/test-e2e/SubGroup/reduce.hpp | 4 +- sycl/test-e2e/SubGroupMask/Basic.cpp | 2 +- sycl/test-e2e/Tracing/image_printers.cpp | 2 +- sycl/test-e2e/USM/math.cpp | 6 +- .../test-e2e/WeakObject/weak_object_utils.hpp | 18 ++--- .../dynamic_allocation.cpp | 2 +- ...rk_group_static_memory_namespace_scope.cpp | 4 +- sycl/test-e2e/XPTI/buffer/in_cycle.cpp | 2 +- sycl/test-e2e/XPTI/buffer/recursion.cpp | 2 +- sycl/test-e2e/XPTI/kernel/basic.cpp | 6 +- 106 files changed, 499 insertions(+), 499 deletions(-) diff --git a/sycl/test-e2e/AddressCast/khr_dynamic_addrspace_cast.cpp b/sycl/test-e2e/AddressCast/khr_dynamic_addrspace_cast.cpp index c140a338e3161..269d6180bb83a 100644 --- a/sycl/test-e2e/AddressCast/khr_dynamic_addrspace_cast.cpp +++ b/sycl/test-e2e/AddressCast/khr_dynamic_addrspace_cast.cpp @@ -31,10 +31,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/Inputs/kernels_in_file2.cpp b/sycl/test-e2e/Assert/Inputs/kernels_in_file2.cpp index 3069840a4ad0d..9824244d18d7c 100644 --- a/sycl/test-e2e/Assert/Inputs/kernels_in_file2.cpp +++ b/sycl/test-e2e/Assert/Inputs/kernels_in_file2.cpp @@ -25,7 +25,7 @@ void enqueueKernel_1_fromFile2(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) { @@ -40,7 +40,7 @@ void enqueueKernel_2_fromFile2(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) { diff --git a/sycl/test-e2e/Assert/assert_in_one_kernel.hpp b/sycl/test-e2e/Assert/assert_in_one_kernel.hpp index 29496c1af925e..2cfc263a1b52f 100644 --- a/sycl/test-e2e/Assert/assert_in_one_kernel.hpp +++ b/sycl/test-e2e/Assert/assert_in_one_kernel.hpp @@ -19,7 +19,7 @@ void assertTest() { queue Q; Q.submit([&](handler &CGH) { - auto Acc = Buf.template get_access(CGH); + auto Acc = Buf.template get_access(CGH); CGH.parallel_for( numOfItems, [=](item<1> Item) { kernelFunc(&Acc[0], Item[0]); }); diff --git a/sycl/test-e2e/AtomicRef/and.h b/sycl/test-e2e/AtomicRef/and.h index 54ceb35eab915..abb932d3fda90 100644 --- a/sycl/test-e2e/AtomicRef/and.h +++ b/sycl/test-e2e/AtomicRef/and.h @@ -32,9 +32,9 @@ void and_local_test(queue q) { 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) { @@ -79,9 +79,9 @@ void and_global_test(queue q) { 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); 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/atomic_memory_order_seq_cst.cpp b/sycl/test-e2e/AtomicRef/atomic_memory_order_seq_cst.cpp index 4d9947601e2e2..1b1ee9b56d9e0 100644 --- a/sycl/test-e2e/AtomicRef/atomic_memory_order_seq_cst.cpp +++ b/sycl/test-e2e/AtomicRef/atomic_memory_order_seq_cst.cpp @@ -35,9 +35,9 @@ void check(queue &q, buffer &res_buf, size_t N_iters) { {N_items / 2 * N_iters + 1, N_items / 2 * N_iters + 1}); q.submit([&](handler &cgh) { - auto res = res_buf.template get_access(cgh); + auto res = res_buf.template get_access(cgh); auto checked = - checked_buf.template get_access(cgh); + checked_buf.template get_access(cgh); cgh.parallel_for(nd_range<1>(N_items, 32), [=](nd_item<1> it) { for (int i = it.get_global_id(0); i < N_items / 2 * N_iters + 1; i += N_items) { @@ -48,8 +48,8 @@ void check(queue &q, buffer &res_buf, size_t N_iters) { }); }); q.submit([&](handler &cgh) { - auto res = res_buf.template get_access(cgh); - auto checked = checked_buf.template get_access(cgh); + auto res = res_buf.template get_access(cgh); + auto checked = checked_buf.template get_access(cgh); cgh.parallel_for(nd_range<1>(N_items / 2, 32), [=](nd_item<1> it) { size_t id = it.get_global_id(0); for (int i = 1; i < N_iters; i++) { @@ -65,9 +65,9 @@ void check(queue &q, buffer &res_buf, size_t N_iters) { int err = 0; buffer err_buf(&err, 1); q.submit([&](handler &cgh) { - auto res = res_buf.template get_access(cgh); - auto checked = checked_buf.template get_access(cgh); - auto err = err_buf.template get_access(cgh); + auto res = res_buf.template get_access(cgh); + auto checked = checked_buf.template get_access(cgh); + auto err = err_buf.template get_access(cgh); cgh.parallel_for(nd_range<1>(N_items / 2, 32), [=](nd_item<1> it) { size_t id = it.get_global_id(0); for (int i = 1; i < N_iters; i++) { @@ -92,8 +92,8 @@ template void test_global(size_t N_iters) { buffer val_buf(&val, 1); q.submit([&](handler &cgh) { - auto res = res_buf.template get_access(cgh); - auto val = val_buf.template get_access(cgh); + auto res = res_buf.template get_access(cgh); + auto val = val_buf.template get_access(cgh); // Intentionally using a small work group size. The assumption being that // more sub groups mean more likely failure for the same number of // work-items if sequential consistency does not work @@ -123,7 +123,7 @@ template void test_local(size_t N_iters) { buffer res_buf({N_items / 2, N_iters}); q.submit([&](handler &cgh) { - auto res = res_buf.template get_access(cgh); + auto res = res_buf.template get_access(cgh); local_accessor val(2, cgh); cgh.parallel_for(nd_range<1>(N_items, N_items), [=](nd_item<1> it) { val[0] = 0; diff --git a/sycl/test-e2e/AtomicRef/load.h b/sycl/test-e2e/AtomicRef/load.h index 792e18fce6557..73eceb303ad38 100644 --- a/sycl/test-e2e/AtomicRef/load.h +++ b/sycl/test-e2e/AtomicRef/load.h @@ -33,9 +33,9 @@ void load_local_test(queue q, size_t N) { buffer output_buf(output.data(), output.size()); q.submit([&](handler &cgh) { - auto ld = load_buf.template get_access(cgh); + auto ld = load_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) { int gid = it.get_global_id(0); @@ -69,9 +69,9 @@ void load_global_test(queue q, size_t N) { buffer output_buf(output.data(), output.size()); q.submit([&](handler &cgh) { - auto ld = load_buf.template get_access(cgh); + auto ld = load_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(ld[0]); diff --git a/sycl/test-e2e/AtomicRef/or.h b/sycl/test-e2e/AtomicRef/or.h index cb3f1a51eb3d1..600d5a49ce98e 100644 --- a/sycl/test-e2e/AtomicRef/or.h +++ b/sycl/test-e2e/AtomicRef/or.h @@ -32,9 +32,9 @@ void or_local_test(queue q) { 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) { @@ -79,9 +79,9 @@ void or_global_test(queue q) { 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); 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/xor.h b/sycl/test-e2e/AtomicRef/xor.h index 2f4ce02cf91f3..ffe5d1c739318 100644 --- a/sycl/test-e2e/AtomicRef/xor.h +++ b/sycl/test-e2e/AtomicRef/xor.h @@ -32,9 +32,9 @@ void xor_local_test(queue q) { 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) { @@ -79,9 +79,9 @@ void xor_global_test(queue q) { 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); 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/Basic/alloc_pinned_host_memory.cpp b/sycl/test-e2e/Basic/alloc_pinned_host_memory.cpp index 5fb98ae7d6b59..32309a20e87de 100644 --- a/sycl/test-e2e/Basic/alloc_pinned_host_memory.cpp +++ b/sycl/test-e2e/Basic/alloc_pinned_host_memory.cpp @@ -30,8 +30,8 @@ int main() { {ext::oneapi::property::buffer::use_pinned_host_memory()}); queue myQueue; myQueue.submit([&](handler &cgh) { - auto A = a.get_access(cgh); - auto B = b.get_access(cgh); + auto A = a.get_access(cgh); + auto B = b.get_access(cgh); cgh.parallel_for(range<1>{10}, [=](id<1> index) { B[index] = 0; A[index] = B[index] + 1; diff --git a/sycl/test-e2e/Basic/buffer/buffer_container.cpp b/sycl/test-e2e/Basic/buffer/buffer_container.cpp index 0a263bb199a6c..84e398ac0ec29 100644 --- a/sycl/test-e2e/Basic/buffer/buffer_container.cpp +++ b/sycl/test-e2e/Basic/buffer/buffer_container.cpp @@ -14,7 +14,7 @@ int main() { buffer A(out_data); queue Q; Q.submit([&](handler &h) { - auto out = A.get_access(h); + auto out = A.get_access(h); h.parallel_for(A.get_range(), [out](id<1> i) { out[i] = 1; }); }); diff --git a/sycl/test-e2e/Basic/buffer/native_buffer_creation_flags.cpp b/sycl/test-e2e/Basic/buffer/native_buffer_creation_flags.cpp index 2fa49525a738b..d65ad14b6ba10 100644 --- a/sycl/test-e2e/Basic/buffer/native_buffer_creation_flags.cpp +++ b/sycl/test-e2e/Basic/buffer/native_buffer_creation_flags.cpp @@ -22,7 +22,7 @@ int main() { // buffer is created with the UR_MEM_FLAG_USE_HOST_POINTER flag. // CHECK: <--- urMemBufferCreate // CHECK-SAME: UR_MEM_FLAG_USE_HOST_POINTER - auto BufAcc = Buf.get_access(Cgh); + auto BufAcc = Buf.get_access(Cgh); Cgh.single_task([=]() { int A = BufAcc[0]; }); }); } diff --git a/sycl/test-e2e/Basic/built-ins.cpp b/sycl/test-e2e/Basic/built-ins.cpp index 2984ba302a9d4..e2085d4ee1b7c 100644 --- a/sycl/test-e2e/Basic/built-ins.cpp +++ b/sycl/test-e2e/Basic/built-ins.cpp @@ -41,8 +41,8 @@ int main() { s::buffer BufMin(s::range<1>(1)); s::buffer BufMax(s::range<1>(1)); q.submit([&](s::handler &cgh) { - auto AccMin = BufMin.get_access(cgh); - auto AccMax = BufMax.get_access(cgh); + auto AccMin = BufMin.get_access(cgh); + auto AccMax = BufMax.get_access(cgh); cgh.single_task([=]() { AccMax[0] = s::max(s::float2{0.5f, 2.5f}, s::float2{2.3f, 2.3f}); AccMin[0] = s::min(float{0.5f}, float{2.3f}); diff --git a/sycl/test-e2e/Basic/device_event.cpp b/sycl/test-e2e/Basic/device_event.cpp index c1ee8e38e8587..e840a47ec87c4 100644 --- a/sycl/test-e2e/Basic/device_event.cpp +++ b/sycl/test-e2e/Basic/device_event.cpp @@ -67,7 +67,7 @@ int test_strideN(size_t stride) { buffer out_buf(out_data, range<1>(nElems)); myQueue.submit([&](handler &cgh) { - auto out_ptr = out_buf.get_access(cgh); + auto out_ptr = out_buf.get_access(cgh); local_accessor local_acc(range<1>(16), cgh); // Create work-groups with 16 work items in each group. diff --git a/sycl/test-e2e/Basic/event_release.cpp b/sycl/test-e2e/Basic/event_release.cpp index ce5432c3cf781..a1f39f82c2abd 100644 --- a/sycl/test-e2e/Basic/event_release.cpp +++ b/sycl/test-e2e/Basic/event_release.cpp @@ -21,7 +21,7 @@ int main() { { buffer Buf{&Val, range<1>(1)}; Q.submit([&](handler &Cgh) { - auto Acc = Buf.get_access(Cgh); + auto Acc = Buf.get_access(Cgh); Cgh.single_task([=]() { Acc[0] = Gold; }); }); } diff --git a/sycl/test-e2e/Basic/free_function_queries/free_function_queries.cpp b/sycl/test-e2e/Basic/free_function_queries/free_function_queries.cpp index 9f1120994e3e9..2cb4b7e0c519d 100644 --- a/sycl/test-e2e/Basic/free_function_queries/free_function_queries.cpp +++ b/sycl/test-e2e/Basic/free_function_queries/free_function_queries.cpp @@ -32,12 +32,12 @@ int main() { sycl::buffer results_buf(results, sycl::range<1>(checks_number)); sycl::queue q; 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(n, [=](sycl::id<1> i) { auto that_id = sycl::ext::oneapi::this_work_item::get_nd_item<1>(); results_acc[0] = that_id.get_global_id() == i; @@ -66,12 +66,12 @@ int main() { sycl::buffer results_buf(results, sycl::range<1>(checks_number)); sycl::queue q; 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(n, [=](auto i) { static_assert(std::is_same>::value, "lambda arg type is unexpected"); @@ -104,12 +104,12 @@ int main() { sycl::queue q; sycl::id<1> offset(1); 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( sycl::range<1>{n}, offset, [=](sycl::item<1, true> i) { auto that_id = diff --git a/sycl/test-e2e/Basic/group_async_copy.cpp b/sycl/test-e2e/Basic/group_async_copy.cpp index ccf3de87c58a0..48d2365c03f05 100644 --- a/sycl/test-e2e/Basic/group_async_copy.cpp +++ b/sycl/test-e2e/Basic/group_async_copy.cpp @@ -108,8 +108,8 @@ template int test(size_t Stride) { initOutputBuffer(OutBuf); Q.submit([&](handler &CGH) { - 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/half_type.cpp b/sycl/test-e2e/Basic/half_type.cpp index c24a4651c4bd4..f20d215a7475d 100644 --- a/sycl/test-e2e/Basic/half_type.cpp +++ b/sycl/test-e2e/Basic/half_type.cpp @@ -34,9 +34,9 @@ void verify_add(queue &q, buffer &a, buffer &b, range<1> &r, buffer c{r}; q.submit([&](handler &cgh) { - auto A = a.get_access(cgh); - auto B = b.get_access(cgh); - auto C = c.get_access(cgh); + auto A = a.get_access(cgh); + auto B = b.get_access(cgh); + auto C = c.get_access(cgh); cgh.parallel_for( r, [=](id<1> index) { C[index] = A[index] + B[index]; }); }); @@ -49,9 +49,9 @@ void verify_min(queue &q, buffer &a, buffer &b, range<1> &r, buffer c{r}; q.submit([&](handler &cgh) { - auto A = a.get_access(cgh); - auto B = b.get_access(cgh); - auto C = c.get_access(cgh); + auto A = a.get_access(cgh); + auto B = b.get_access(cgh); + auto C = c.get_access(cgh); cgh.parallel_for( r, [=](id<1> index) { C[index] = A[index] - B[index]; }); }); @@ -64,9 +64,9 @@ void verify_mul(queue &q, buffer &a, buffer &b, range<1> &r, buffer c{r}; q.submit([&](handler &cgh) { - auto A = a.get_access(cgh); - auto B = b.get_access(cgh); - auto C = c.get_access(cgh); + auto A = a.get_access(cgh); + auto B = b.get_access(cgh); + auto C = c.get_access(cgh); cgh.parallel_for( r, [=](id<1> index) { C[index] = A[index] * B[index]; }); }); @@ -79,9 +79,9 @@ void verify_div(queue &q, buffer &a, buffer &b, range<1> &r, buffer c{r}; q.submit([&](handler &cgh) { - auto A = a.get_access(cgh); - auto B = b.get_access(cgh); - auto C = c.get_access(cgh); + auto A = a.get_access(cgh); + auto B = b.get_access(cgh); + auto C = c.get_access(cgh); cgh.parallel_for( r, [=](id<1> index) { C[index] = A[index] / B[index]; }); }); @@ -94,7 +94,7 @@ void verify_vec(queue &q) { int a = 0; buffer e(&a, range<1>(1)); q.submit([&](sycl::handler &cgh) { - auto E = e.get_access(cgh); + auto E = e.get_access(cgh); cgh.single_task([=]() { if (int(hvec.s0()) != 999 || int(hvec.s1()) != 999) E[0] = 1; diff --git a/sycl/test-e2e/Basic/handler/handler_mem_op.cpp b/sycl/test-e2e/Basic/handler/handler_mem_op.cpp index 553e3e8f1c624..a6b87c3e9a912 100644 --- a/sycl/test-e2e/Basic/handler/handler_mem_op.cpp +++ b/sycl/test-e2e/Basic/handler/handler_mem_op.cpp @@ -226,7 +226,7 @@ template void test_fill(T Val) { buffer Buffer(Data, range<1>(Size)); queue Queue; Queue.submit([&](handler &Cgh) { - accessor Accessor( + accessor Accessor( Buffer, Cgh, range<1>(Size)); Cgh.fill(Accessor, Value); }); @@ -247,7 +247,7 @@ template void test_copy_ptr_acc() { buffer Buffer(Data, range<1>(Size)); queue Queue; Queue.submit([&](handler &Cgh) { - accessor Accessor( + accessor Accessor( Buffer, Cgh, range<1>(Size)); Cgh.copy(Values, Accessor); }); @@ -263,8 +263,8 @@ template void test_copy_ptr_acc() { buffer DstBuf(&DstValue, range<1>(1)); queue Queue; Queue.submit([&](handler &Cgh) { - accessor - DstAcc(DstBuf, Cgh); + accessor DstAcc( + DstBuf, Cgh); Cgh.copy(&SrcValue, DstAcc); }); } @@ -283,7 +283,7 @@ template void test_3D_copy_ptr_acc() { buffer Buffer(Data, Range); queue Queue; Queue.submit([&](handler &Cgh) { - accessor Accessor( + accessor Accessor( Buffer, Cgh, Range); Cgh.copy(Values, Accessor); }); @@ -304,7 +304,7 @@ template void test_copy_acc_ptr() { buffer Buffer(Data, range<1>(Size)); queue Queue; Queue.submit([&](handler &Cgh) { - accessor Accessor( + accessor Accessor( Buffer, Cgh, range<1>(Size)); Cgh.copy(Accessor, Values); }); @@ -320,7 +320,7 @@ template void test_copy_acc_ptr() { T Dst[Size] = {0}; queue Queue; Queue.submit([&](handler &Cgh) { - auto Acc = Buffer.template get_access(Cgh); + auto Acc = Buffer.template get_access(Cgh); Cgh.copy(Acc, reinterpret_cast(Dst)); }); Queue.wait(); @@ -336,8 +336,8 @@ template void test_copy_acc_ptr() { buffer SrcBuf(&SrcValue, range<1>(1)); queue Queue; Queue.submit([&](handler &Cgh) { - accessor SrcAcc(SrcBuf, - Cgh); + accessor SrcAcc(SrcBuf, + Cgh); Cgh.copy(SrcAcc, &DstValue); }); } @@ -348,7 +348,7 @@ template void test_copy_acc_ptr() { DstValue = 0; { buffer SrcBuf(&SrcValue, range<1>(1)); - accessor SrcAcc(SrcBuf); { @@ -374,7 +374,7 @@ template void test_3D_copy_acc_ptr() { buffer Buffer(Data, Range); queue Queue; Queue.submit([&](handler &Cgh) { - accessor Accessor( + accessor Accessor( Buffer, Cgh, Range); Cgh.copy(Accessor, Values); }); @@ -395,7 +395,7 @@ template void test_copy_shared_ptr_acc() { buffer Buffer(Data, range<1>(Size)); queue Queue; Queue.submit([&](handler &Cgh) { - accessor Accessor( + accessor Accessor( Buffer, Cgh, range<1>(Size)); Cgh.copy(Values, Accessor); }); @@ -413,7 +413,7 @@ template void test_copy_shared_ptr_const_acc() { buffer Buffer(Data, range<1>(Size)); queue Queue; Queue.submit([&](handler &Cgh) { - accessor Accessor( + accessor Accessor( Buffer, Cgh, range<1>(Size)); Cgh.copy(Values, Accessor); }); @@ -434,7 +434,7 @@ template void test_copy_acc_shared_ptr() { buffer Buffer(Data, range<1>(Size)); queue Queue; Queue.submit([&](handler &Cgh) { - accessor Accessor( + accessor Accessor( Buffer, Cgh, range<1>(Size)); Cgh.copy(Accessor, Values); }); @@ -456,9 +456,9 @@ template void test_copy_acc_acc() { buffer BufferTo(Values, range<1>(Size)); queue Queue; Queue.submit([&](handler &Cgh) { - accessor AccessorFrom( + accessor AccessorFrom( BufferFrom, Cgh, range<1>(Size)); - accessor AccessorTo( + accessor AccessorTo( BufferTo, Cgh, range<1>(Size)); Cgh.copy(AccessorFrom, AccessorTo); }); @@ -481,13 +481,13 @@ template void test_update_host() { Buffer.set_final_data(nullptr); queue Queue; Queue.submit([&](handler &Cgh) { - accessor Accessor( + accessor Accessor( Buffer, Cgh, range<1>(Size)); Cgh.parallel_for>( range<1>{Size}, [=](id<1> Index) { Accessor[Index] = Index.get(0); }); }); Queue.submit([&](handler &Cgh) { - accessor Accessor( + accessor Accessor( Buffer, Cgh, range<1>(Size)); Cgh.update_host(Accessor); }); @@ -511,9 +511,9 @@ template void test_2D_copy_acc_acc() { buffer BufferTo((T *)Values, range<2>(Size, Size)); queue Queue; Queue.submit([&](handler &Cgh) { - accessor AccessorFrom( + accessor AccessorFrom( BufferFrom, Cgh, range<2>(Size, Size)); - accessor AccessorTo( + accessor AccessorTo( BufferTo, Cgh, range<2>(Size, Size)); Cgh.copy(AccessorFrom, AccessorTo); }); @@ -542,9 +542,9 @@ template void test_3D_copy_acc_acc() { buffer BufferTo((T *)Values, range<3>(Size, Size, Size)); queue Queue; Queue.submit([&](handler &Cgh) { - accessor AccessorFrom( + accessor AccessorFrom( BufferFrom, Cgh, range<3>(Size, Size, Size)); - accessor AccessorTo( + accessor AccessorTo( BufferTo, Cgh, range<3>(Size, Size, Size)); Cgh.copy(AccessorFrom, AccessorTo); }); @@ -567,9 +567,9 @@ template void test_0D1D_copy_acc_acc() { buffer BufferTo(&Dst, range<1>(1)); queue Queue; Queue.submit([&](handler &Cgh) { - accessor AccessorFrom( + accessor AccessorFrom( BufferFrom, Cgh); - accessor AccessorTo( + accessor AccessorTo( BufferTo, Cgh); Cgh.copy(AccessorFrom, AccessorTo); }); @@ -584,9 +584,9 @@ template void test_0D1D_copy_acc_acc() { buffer BufferTo(&Dst, range<1>(1)); queue Queue; Queue.submit([&](handler &Cgh) { - accessor AccessorFrom( + accessor AccessorFrom( BufferFrom, Cgh); - accessor AccessorTo( + accessor AccessorTo( BufferTo, Cgh); Cgh.copy(AccessorFrom, AccessorTo); }); @@ -601,9 +601,9 @@ template void test_0D1D_copy_acc_acc() { buffer BufferTo(&Dst, range<1>(1)); queue Queue; Queue.submit([&](handler &Cgh) { - accessor AccessorFrom( + accessor AccessorFrom( BufferFrom, Cgh); - accessor AccessorTo( + accessor AccessorTo( BufferTo, Cgh); Cgh.copy(AccessorFrom, AccessorTo); }); @@ -621,9 +621,9 @@ template void test_1D2D_copy_acc_acc() { buffer BufferTo(&Values[0], range<2>(Size / 2, 2)); queue Queue; Queue.submit([&](handler &Cgh) { - accessor AccessorFrom( + accessor AccessorFrom( BufferFrom, Cgh, range<1>(Size)); - accessor AccessorTo( + accessor AccessorTo( BufferTo, Cgh, range<2>(Size / 2, 2)); Cgh.copy(AccessorFrom, AccessorTo); }); @@ -641,9 +641,9 @@ template void test_1D3D_copy_acc_acc() { buffer BufferTo(&Values[0], range<3>(Size / 4, 2, 2)); queue Queue; Queue.submit([&](handler &Cgh) { - accessor AccessorFrom( + accessor AccessorFrom( BufferFrom, Cgh, range<1>(Size)); - accessor AccessorTo( + accessor AccessorTo( BufferTo, Cgh, range<3>(Size / 4, 2, 2)); Cgh.copy(AccessorFrom, AccessorTo); }); @@ -661,9 +661,9 @@ template void test_2D1D_copy_acc_acc() { buffer BufferTo(&Values[0], range<1>(Size)); queue Queue; Queue.submit([&](handler &Cgh) { - accessor AccessorFrom( + accessor AccessorFrom( BufferFrom, Cgh, range<2>(Size / 2, 2)); - accessor AccessorTo( + accessor AccessorTo( BufferTo, Cgh, range<1>(Size)); Cgh.copy(AccessorFrom, AccessorTo); }); @@ -681,9 +681,9 @@ template void test_2D3D_copy_acc_acc() { buffer BufferTo(&Values[0], range<3>(Size / 4, 2, 2)); queue Queue; Queue.submit([&](handler &Cgh) { - accessor AccessorFrom( + accessor AccessorFrom( BufferFrom, Cgh, range<2>(Size / 2, 2)); - accessor AccessorTo( + accessor AccessorTo( BufferTo, Cgh, range<3>(Size / 4, 2, 2)); Cgh.copy(AccessorFrom, AccessorTo); }); @@ -701,9 +701,9 @@ template void test_3D1D_copy_acc_acc() { buffer BufferTo(&Values[0], range<1>(Size)); queue Queue; Queue.submit([&](handler &Cgh) { - accessor AccessorFrom( + accessor AccessorFrom( BufferFrom, Cgh, range<3>(Size / 4, 2, 2)); - accessor AccessorTo( + accessor AccessorTo( BufferTo, Cgh, range<1>(Size)); Cgh.copy(AccessorFrom, AccessorTo); }); @@ -721,9 +721,9 @@ template void test_3D2D_copy_acc_acc() { buffer BufferTo(&Values[0], range<2>(Size / 2, 2)); queue Queue; Queue.submit([&](handler &Cgh) { - accessor AccessorFrom( + accessor AccessorFrom( BufferFrom, Cgh, range<3>(Size / 4, 2, 2)); - accessor AccessorTo( + accessor AccessorTo( BufferTo, Cgh, range<2>(Size / 2, 2)); Cgh.copy(AccessorFrom, AccessorTo); }); diff --git a/sycl/test-e2e/Basic/image/image_accessor_range.cpp b/sycl/test-e2e/Basic/image/image_accessor_range.cpp index c29cd31d5cfb3..8937b97ab46c3 100644 --- a/sycl/test-e2e/Basic/image/image_accessor_range.cpp +++ b/sycl/test-e2e/Basic/image/image_accessor_range.cpp @@ -30,7 +30,7 @@ void try_1D(queue &Q) { image_channel_type::unsigned_int8, range{M}); Q.submit([&](handler &h) { - accessor acs1(im1, h); + accessor acs1(im1, h); accessor ABX{BX, h}; auto R = acs1.get_range(); std::cout << "Host acs1.get_range()=" << R[0] << "\n"; @@ -60,7 +60,7 @@ void try_2D(queue &Q) { image_channel_type::unsigned_int8, range{M, N}); Q.submit([&](handler &h) { - accessor acs2(im2, h); + accessor acs2(im2, h); accessor ABX{BX, h}; auto R = acs2.get_range(); std::cout << "Host acs2.get_range()=" << R[0] << "," << R[1] << "\n"; @@ -92,8 +92,8 @@ void try_3D(queue &Q) { image_channel_type::unsigned_int8, range{M, N, L}); Q.submit([&](handler &h) { - accessor acs3(im3, - h); + accessor acs3(im3, + h); accessor ABX{BX, h}; auto R = acs3.get_range(); std::cout << "Host acs3.get_range()=" << R[0] << "," << R[1] << "," << R[2] diff --git a/sycl/test-e2e/Basic/image/image_accessor_readwrite_half.cpp b/sycl/test-e2e/Basic/image/image_accessor_readwrite_half.cpp index 52edf12c04f25..418cce5df235c 100644 --- a/sycl/test-e2e/Basic/image/image_accessor_readwrite_half.cpp +++ b/sycl/test-e2e/Basic/image/image_accessor_readwrite_half.cpp @@ -75,7 +75,7 @@ void write_type_order(char *HostPtr, const s::image_channel_order ImgOrder, s::image<1> Img(HostPtr, ImgOrder, ImgType, s::range<1>{10}); s::queue Queue; Queue.submit([&](s::handler &cgh) { - auto WriteAcc = Img.get_access(cgh); + auto WriteAcc = Img.get_access(cgh); cgh.single_task< class kernel_class(ImgType), 0>>( [=]() { WriteAcc.write(Coord, Color); }); @@ -95,8 +95,8 @@ void check_read_type_order(char *HostPtr, const s::image_channel_order ImgOrder, s::queue Queue; s::buffer ReadDataBuf(&ReadData, s::range<1>(1)); Queue.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_read.h b/sycl/test-e2e/Basic/image/image_read.h index f670349df7f63..c13830eb0175a 100644 --- a/sycl/test-e2e/Basic/image/image_read.h +++ b/sycl/test-e2e/Basic/image/image_read.h @@ -33,9 +33,9 @@ bool test1d_coord(s::queue myQueue, dataT *hostPtr, coordT coord, // Do the test by reading a single pixel from the image myQueue.submit([&](s::handler &cgh) { - auto imageAcc = image.get_access(cgh); - s::accessor resultDataAcc(resultDataBuf, - cgh); + auto imageAcc = image.get_access(cgh); + s::accessor resultDataAcc(resultDataBuf, + cgh); cgh.single_task>([=]() { dataT RetColor = imageAcc.read(coord); @@ -74,9 +74,9 @@ bool test2d_coord(s::queue myQueue, dataT *hostPtr, coordT coord, // Do the test by reading a single pixel from the image myQueue.submit([&](s::handler &cgh) { - auto imageAcc = image.get_access(cgh); - s::accessor resultDataAcc(resultDataBuf, - cgh); + auto imageAcc = image.get_access(cgh); + s::accessor resultDataAcc(resultDataBuf, + cgh); cgh.single_task>([=]() { dataT RetColor = imageAcc.read(coord); @@ -114,9 +114,9 @@ bool test3d_coord(s::queue myQueue, dataT *hostPtr, coordT coord, // Do the test by reading a single pixel from the image myQueue.submit([&](s::handler &cgh) { - auto imageAcc = image.get_access(cgh); - s::accessor resultDataAcc(resultDataBuf, - cgh); + auto imageAcc = image.get_access(cgh); + s::accessor resultDataAcc(resultDataBuf, + cgh); cgh.single_task>([=]() { dataT RetColor = imageAcc.read(coord); diff --git a/sycl/test-e2e/Basic/image/image_write.h b/sycl/test-e2e/Basic/image/image_write.h index 4d7e8a21c2e97..35ffc17819c86 100644 --- a/sycl/test-e2e/Basic/image/image_write.h +++ b/sycl/test-e2e/Basic/image/image_write.h @@ -35,15 +35,15 @@ bool test1d_coord(s::queue myQueue, dataT *hostPtr, coordT coord, s::buffer resultDataBuf(&resultData, s::range<1>(1)); myQueue.submit([&](s::handler &cgh) { - auto imageAcc = image.get_access(cgh); + auto imageAcc = image.get_access(cgh); cgh.single_task>( [=]() { imageAcc.write(coord, colour); }); }); myQueue.submit([&](s::handler &cgh) { - auto imageAcc = image.get_access(cgh); - s::accessor resultDataAcc(resultDataBuf, - cgh); + auto imageAcc = image.get_access(cgh); + s::accessor resultDataAcc(resultDataBuf, + cgh); cgh.single_task>([=]() { dataT RetColor = imageAcc.read(coord); @@ -81,16 +81,16 @@ bool test2d_coord(s::queue myQueue, dataT *hostPtr, coordT coord, s::buffer resultDataBuf(&resultData, s::range<1>(1)); myQueue.submit([&](s::handler &cgh) { - auto imageAcc = image.get_access(cgh); + auto imageAcc = image.get_access(cgh); cgh.single_task>( [=]() { imageAcc.write(coord, colour); }); }); myQueue.submit([&](s::handler &cgh) { - auto imageAcc = image.get_access(cgh); - s::accessor resultDataAcc(resultDataBuf, - cgh); + auto imageAcc = image.get_access(cgh); + s::accessor resultDataAcc(resultDataBuf, + cgh); cgh.single_task>([=]() { dataT RetColor = imageAcc.read(coord); @@ -128,15 +128,15 @@ bool test3d_coord(s::queue myQueue, dataT *hostPtr, coordT coord, s::buffer resultDataBuf(&resultData, s::range<1>(1)); myQueue.submit([&](s::handler &cgh) { - auto imageAcc = image.get_access(cgh); + auto imageAcc = image.get_access(cgh); cgh.single_task>( [=]() { imageAcc.write(coord, colour); }); }); myQueue.submit([&](s::handler &cgh) { - auto imageAcc = image.get_access(cgh); - s::accessor resultDataAcc(resultDataBuf, - cgh); + auto imageAcc = image.get_access(cgh); + s::accessor resultDataAcc(resultDataBuf, + cgh); cgh.single_task>([=]() { dataT RetColor = imageAcc.read(coord); diff --git a/sycl/test-e2e/Basic/kernel_bundle/kernel_bundle_api.cpp b/sycl/test-e2e/Basic/kernel_bundle/kernel_bundle_api.cpp index 7d87d2ed13e55..95bd422d32b81 100644 --- a/sycl/test-e2e/Basic/kernel_bundle/kernel_bundle_api.cpp +++ b/sycl/test-e2e/Basic/kernel_bundle/kernel_bundle_api.cpp @@ -240,7 +240,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(KernelBundleExecutable); CGH.single_task([=]() { Acc[0] = 42; }); }); diff --git a/sycl/test-e2e/Basic/kernel_info_attr.cpp b/sycl/test-e2e/Basic/kernel_info_attr.cpp index e185b7b26a5c2..deb1841955af0 100644 --- a/sycl/test-e2e/Basic/kernel_info_attr.cpp +++ b/sycl/test-e2e/Basic/kernel_info_attr.cpp @@ -37,7 +37,7 @@ int main() { kernel krn = KB.get_kernel(KernelID); q.submit([&](handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.single_task([=]() { acc[0] = acc[0] + 1; }); }); diff --git a/sycl/test-e2e/Basic/multi_ptr_legacy.hpp b/sycl/test-e2e/Basic/multi_ptr_legacy.hpp index 45cc9a104c55a..93dcb1f6a05f8 100644 --- a/sycl/test-e2e/Basic/multi_ptr_legacy.hpp +++ b/sycl/test-e2e/Basic/multi_ptr_legacy.hpp @@ -55,10 +55,10 @@ template void testMultPtr() { buffer bufferData_2(data_2, numOfItems); queue myQueue; myQueue.submit([&](handler &cgh) { - accessor accessorData_1(bufferData_1, cgh); - accessor accessorData_2(bufferData_2, cgh); local_accessor localAccessor(numOfItems, cgh); @@ -139,14 +139,14 @@ template void testMultPtrArrowOperator() { buffer result_buf{1}; queue myQueue; myQueue.submit([&](handler &cgh) { - accessor, 1, access::mode::read, access::target::device, + accessor, 1, access_mode::read, access::target::device, access::placeholder::false_t> accessorData_1(bufferData_1, cgh); - accessor, 1, access::mode::read, access::target::constant_buffer, + accessor, 1, access_mode::read, access::target::constant_buffer, access::placeholder::false_t> accessorData_2(bufferData_2, cgh); local_accessor, 1> accessorData_3(1, cgh); - accessor, 1, access::mode::read, access::target::device, + accessor, 1, access_mode::read, access::target::device, access::placeholder::false_t> accessorData_4(bufferData_4, cgh); diff --git a/sycl/test-e2e/Basic/multisource_spv_obj.cpp b/sycl/test-e2e/Basic/multisource_spv_obj.cpp index 7f03c74f51753..03c2788c2022e 100644 --- a/sycl/test-e2e/Basic/multisource_spv_obj.cpp +++ b/sycl/test-e2e/Basic/multisource_spv_obj.cpp @@ -42,7 +42,7 @@ void init_buf(queue &q, buffer &b, range<1> &r, int i); #elif INIT_KERNEL void init_buf(queue &q, buffer &b, range<1> &r, int i) { q.submit([&](handler &cgh) { - auto B = b.get_access(cgh); + auto B = b.get_access(cgh); cgh.parallel_for(r, [=](id<1> index) { B[index] = i; }); }); } @@ -55,9 +55,9 @@ void calc_buf(queue &q, buffer &a, buffer &b, buffer &c, void calc_buf(queue &q, buffer &a, buffer &b, buffer &c, range<1> &r) { q.submit([&](handler &cgh) { - auto A = a.get_access(cgh); - auto B = b.get_access(cgh); - auto C = c.get_access(cgh); + auto A = a.get_access(cgh); + auto B = b.get_access(cgh); + auto C = c.get_access(cgh); cgh.parallel_for( r, [=](id<1> index) { C[index] = A[index] - B[index]; }); }); diff --git a/sycl/test-e2e/Basic/parallel_for_indexers.cpp b/sycl/test-e2e/Basic/parallel_for_indexers.cpp index 59dbf3329e82d..946237d0db7db 100644 --- a/sycl/test-e2e/Basic/parallel_for_indexers.cpp +++ b/sycl/test-e2e/Basic/parallel_for_indexers.cpp @@ -20,7 +20,7 @@ int main() { {property::buffer::use_host_ptr()}); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = b.get_access(cgh); + auto B = b.get_access(cgh); cgh.parallel_for(globalRange, [=](id<1> index) { B[index] = index[0]; }); }); @@ -43,7 +43,7 @@ int main() { {property::buffer::use_host_ptr()}); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = b.get_access(cgh); + auto B = b.get_access(cgh); cgh.parallel_for( globalRange, [=](item<1, false> index) { B[index.get_id()] = int2{index.get_id()[0], index.get_range()[0]}; @@ -72,7 +72,7 @@ int main() { {property::buffer::use_host_ptr()}); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = b.get_access(cgh); + auto B = b.get_access(cgh); cgh.parallel_for( globalRange, globalOffset, [=](item<1> index) { B[index.get_id()] = int3{index.get_id()[0], index.get_range()[0], @@ -107,7 +107,7 @@ int main() { {property::buffer::use_host_ptr()}); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = b.get_access(cgh); + auto B = b.get_access(cgh); cgh.parallel_for(ndRange, [=](nd_item<1> index) { B[index.get_global_id()] = int3{index.get_global_id()[0], index.get_global_range()[0], diff --git a/sycl/test-e2e/Basic/range_offset_fit_in_int.cpp b/sycl/test-e2e/Basic/range_offset_fit_in_int.cpp index 12821742f97cb..e3950d13a8a8f 100644 --- a/sycl/test-e2e/Basic/range_offset_fit_in_int.cpp +++ b/sycl/test-e2e/Basic/range_offset_fit_in_int.cpp @@ -64,7 +64,7 @@ void test() { try { S::buffer Buf{&Data, 1}; Queue.submit([&](S::handler &CGH) { - auto Acc = Buf.get_access(CGH); + auto Acc = Buf.get_access(CGH); CGH.parallel_for(RangeOutOfLimits, [=](S::id<2> Id) { Acc[0] += 1; }); @@ -81,7 +81,7 @@ void test() { try { S::buffer Buf{&Data, 1}; Queue.submit([&](S::handler &CGH) { - auto Acc = Buf.get_access(CGH); + auto Acc = Buf.get_access(CGH); CGH.parallel_for(RangeInLimits, [Acc](S::id<2> Id) { Acc[0] += 1; }); @@ -94,7 +94,7 @@ void test() { try { S::buffer Buf{&Data, 1}; Queue.submit([&](S::handler &CGH) { - auto Acc = Buf.get_access(CGH); + auto Acc = Buf.get_access(CGH); CGH.parallel_for(RangeInLimits_POL, [Acc](S::id<2> Id) { Acc[0] += 1; }); @@ -111,7 +111,7 @@ void test() { try { S::buffer Buf{&Data, 1}; Queue.submit([&](S::handler &CGH) { - auto Acc = Buf.get_access(CGH); + auto Acc = Buf.get_access(CGH); CGH.parallel_for(RangeOutOfLimits, OffsetInLimits, [Acc](S::id<2> Id) { Acc[0] += 1; }); @@ -128,7 +128,7 @@ void test() { try { S::buffer Buf{&Data, 1}; Queue.submit([&](S::handler &CGH) { - auto Acc = Buf.get_access(CGH); + auto Acc = Buf.get_access(CGH); CGH.parallel_for( RangeInLimits_Large, OffsetInLimits_Large, @@ -146,7 +146,7 @@ void test() { try { S::buffer Buf{&Data, 1}; Queue.submit([&](S::handler &CGH) { - auto Acc = Buf.get_access(CGH); + auto Acc = Buf.get_access(CGH); CGH.parallel_for(RangeInLimits, OffsetOutOfLimits, [Acc](S::id<2> Id) { Acc[0] += 1; }); @@ -163,7 +163,7 @@ void test() { try { S::buffer Buf{&Data, 1}; Queue.submit([&](S::handler &CGH) { - auto Acc = Buf.get_access(CGH); + auto Acc = Buf.get_access(CGH); CGH.parallel_for(RangeInLimits, OffsetInLimits, [Acc](S::id<2> Id) { Acc[0] += 1; }); @@ -176,7 +176,7 @@ void test() { try { S::buffer Buf{&Data, 1}; Queue.submit([&](S::handler &CGH) { - auto Acc = Buf.get_access(CGH); + auto Acc = Buf.get_access(CGH); CGH.parallel_for( NDRange_ROL_LIL_OIL, [Acc](S::nd_item<2> Id) { Acc[0] += 1; }); @@ -193,7 +193,7 @@ void test() { try { S::buffer Buf{&Data, 1}; Queue.submit([&](S::handler &CGH) { - auto Acc = Buf.get_access(CGH); + auto Acc = Buf.get_access(CGH); CGH.parallel_for( NDRange_RIL_LIL_OOL, [Acc](S::nd_item<2> Id) { Acc[0] += 1; }); @@ -210,7 +210,7 @@ void test() { try { S::buffer Buf{&Data, 1}; Queue.submit([&](S::handler &CGH) { - auto Acc = Buf.get_access(CGH); + auto Acc = Buf.get_access(CGH); CGH.parallel_for( NDRange_RIL_LIL_OIL, [Acc](S::nd_item<2> Id) { Acc[0] += 1; }); @@ -223,7 +223,7 @@ void test() { try { S::buffer Buf{&Data, 1}; Queue.submit([&](S::handler &CGH) { - auto Acc = Buf.get_access(CGH); + auto Acc = Buf.get_access(CGH); CGH.parallel_for( NDRange_RIL_LIL_OIL_POL, [Acc](S::nd_item<2> Id) { Acc[0] += 1; }); @@ -240,7 +240,7 @@ void test() { try { S::buffer Buf{&Data, 1}; Queue.submit([&](S::handler &CGH) { - auto Acc = Buf.get_access(CGH); + auto Acc = Buf.get_access(CGH); CGH.parallel_for( NDRange_RIL_LIL_OIL_POL, [Acc](S::nd_item<2> Id) { Acc[0] += 1; }); diff --git a/sycl/test-e2e/Basic/swizzle_op.cpp b/sycl/test-e2e/Basic/swizzle_op.cpp index 0d800faa5769b..42186afb2511b 100644 --- a/sycl/test-e2e/Basic/swizzle_op.cpp +++ b/sycl/test-e2e/Basic/swizzle_op.cpp @@ -22,7 +22,7 @@ int main() { buffer b(results, range<1>(3)); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = b.get_access(cgh); + auto B = b.get_access(cgh); cgh.single_task([=]() { float2 ab = {4, 2}; float c = ab.x() * ab.y(); @@ -43,7 +43,7 @@ int main() { buffer b(results, range<1>(3)); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = b.get_access(cgh); + auto B = b.get_access(cgh); cgh.single_task([=]() { float2 ab = {4, 2}; float c = ab.x() * 2; @@ -64,7 +64,7 @@ int main() { buffer b(results, range<1>(3)); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = b.get_access(cgh); + auto B = b.get_access(cgh); cgh.single_task([=]() { float2 ab = {4, 2}; float c = 4 * ab.y(); @@ -85,7 +85,7 @@ int main() { buffer b(results, range<1>(4)); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = b.get_access(cgh); + auto B = b.get_access(cgh); cgh.single_task([=]() { float2 ab = {4, 2}; float2 c = {0, 0}; @@ -109,7 +109,7 @@ int main() { buffer b(results, range<1>(4)); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = b.get_access(cgh); + auto B = b.get_access(cgh); cgh.single_task([=]() { float2 ab = {4, 2}; float2 c = {0, 0}; @@ -133,7 +133,7 @@ int main() { buffer b(results, range<1>(4)); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = b.get_access(cgh); + auto B = b.get_access(cgh); cgh.single_task([=]() { float2 ab = {4, 2}; float2 c = {0, 0}; @@ -157,7 +157,7 @@ int main() { buffer b(results, range<1>(6)); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = b.get_access(cgh); + auto B = b.get_access(cgh); cgh.single_task([=]() { uchar4 abc = {4, 2, 1, 0}; @@ -192,7 +192,7 @@ int main() { buffer b(results, range<1>(4)); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = b.get_access(cgh); + auto B = b.get_access(cgh); cgh.single_task([=]() { uchar4 cba; unsigned char x = 1; @@ -232,7 +232,7 @@ int main() { buffer b(results, range<1>(4)); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = b.get_access(cgh); + auto B = b.get_access(cgh); cgh.single_task([=]() { uchar4 vec; unsigned int add = 254; @@ -261,7 +261,7 @@ int main() { buffer b((int3 *)FF, range<1>(2)); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = b.get_access(cgh); + auto B = b.get_access(cgh); cgh.parallel_for( range<1>{2}, [=](id<1> ID) { B[ID] = int3{(int3)ID[0]} / B[ID]; }); }); @@ -279,7 +279,7 @@ int main() { buffer b(&result, range<1>(1)); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = b.get_access(cgh); + auto B = b.get_access(cgh); cgh.single_task([=]() { int3 testVec1 = {2, 2, 2}; int3 testVec2 = {1, 1, 1}; diff --git a/sycl/test-e2e/Basic/use_pinned_host_memory.cpp b/sycl/test-e2e/Basic/use_pinned_host_memory.cpp index cba9c575e8497..b4c557ab50ad6 100644 --- a/sycl/test-e2e/Basic/use_pinned_host_memory.cpp +++ b/sycl/test-e2e/Basic/use_pinned_host_memory.cpp @@ -20,7 +20,7 @@ int main() { sycl::queue Q; Q.submit([&](sycl::handler &CGH) { - auto Acc = Buf.get_access(CGH); + auto Acc = Buf.get_access(CGH); CGH.single_task([=]() {}); }); diff --git a/sycl/test-e2e/ClusterLaunch/enqueueLaunchCustom_check_event_deps.cpp b/sycl/test-e2e/ClusterLaunch/enqueueLaunchCustom_check_event_deps.cpp index 50a488d861874..3a69d39fdda36 100644 --- a/sycl/test-e2e/ClusterLaunch/enqueueLaunchCustom_check_event_deps.cpp +++ b/sycl/test-e2e/ClusterLaunch/enqueueLaunchCustom_check_event_deps.cpp @@ -53,7 +53,7 @@ int main() { { sycl::buffer Buff(HostArray.data(), 4096); Queue.submit([&](sycl::handler &CGH) { - auto Acc = Buff.template get_access(CGH); + auto Acc = Buff.template get_access(CGH); CGH.parallel_for(4096, [=](auto i) { Acc[i] = 1; }); }); Queue.submit([&](sycl::handler &CGH) { @@ -61,12 +61,12 @@ int main() { cuda::cluster_size ClusterDims(sycl::range{2}); properties ClusterLaunchProperty{ClusterDims}; - auto Acc = Buff.template get_access(CGH); + auto Acc = Buff.template get_access(CGH); CGH.parallel_for(sycl::nd_range({4096}, {32}), KernelFunctor(ClusterLaunchProperty, Acc)); }); Queue.submit([&](sycl::handler &CGH) { - auto Acc = Buff.template get_access(CGH); + auto Acc = Buff.template get_access(CGH); CGH.parallel_for(4096, [=](auto i) { Acc[i] *= 5; }); }); } diff --git a/sycl/test-e2e/DeprecatedFeatures/DiscardEvents/discard_events_accessors.cpp b/sycl/test-e2e/DeprecatedFeatures/DiscardEvents/discard_events_accessors.cpp index fffd5d22625f7..b0c7afd42170e 100644 --- a/sycl/test-e2e/DeprecatedFeatures/DiscardEvents/discard_events_accessors.cpp +++ b/sycl/test-e2e/DeprecatedFeatures/DiscardEvents/discard_events_accessors.cpp @@ -69,7 +69,7 @@ int main(int Argc, const char *Argv[]) { RunKernelHelper(Q, [&](int *Harray) { sycl::buffer Buf(Range); Q.submit([&](sycl::handler &CGH) { - auto Acc = Buf.get_access(CGH); + auto Acc = Buf.get_access(CGH); CGH.parallel_for( Range, [=](sycl::item<1> itemID) { size_t i = itemID.get_id(0); diff --git a/sycl/test-e2e/DeviceArchitecture/device_architecture_on_device_aot.cpp b/sycl/test-e2e/DeviceArchitecture/device_architecture_on_device_aot.cpp index 0050f08c752a5..c728a8cbdae6c 100644 --- a/sycl/test-e2e/DeviceArchitecture/device_architecture_on_device_aot.cpp +++ b/sycl/test-e2e/DeviceArchitecture/device_architecture_on_device_aot.cpp @@ -22,7 +22,7 @@ int main() { // test if_architecture_is q.submit([&](handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.single_task([=]() { if_architecture_is([&]() { acc[0] = 1; @@ -32,7 +32,7 @@ int main() { // test else_if_architecture_is q.submit([&](handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.single_task([=]() { if_architecture_is([&]() { acc[1] = 0; @@ -44,7 +44,7 @@ int main() { // test otherwise q.submit([&](handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.single_task([=]() { if_architecture_is([&]() { acc[2] = 0; @@ -55,7 +55,7 @@ int main() { // test more than one architecture template parameter is passed to // if_architecture_is q.submit([&](handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.single_task([=]() { if_architecture_is( [&]() { acc[3] = 4; }) diff --git a/sycl/test-e2e/DeviceGlobal/device_global_operator_passthrough.hpp b/sycl/test-e2e/DeviceGlobal/device_global_operator_passthrough.hpp index 514176b53b413..719e3cfe7e762 100644 --- a/sycl/test-e2e/DeviceGlobal/device_global_operator_passthrough.hpp +++ b/sycl/test-e2e/DeviceGlobal/device_global_operator_passthrough.hpp @@ -19,7 +19,7 @@ int test() { { buffer OutBuf{&Out, 1}; Q.submit([&](handler &CGH) { - auto OutAcc = OutBuf.get_access(CGH); + auto OutAcc = OutBuf.get_access(CGH); CGH.single_task([=]() { OutAcc[0] = DeviceGlobalVar; }); }); } diff --git a/sycl/test-e2e/DeviceLib/ITTAnnotations/atomic.cpp b/sycl/test-e2e/DeviceLib/ITTAnnotations/atomic.cpp index 272c21885365f..0af04f8247b41 100644 --- a/sycl/test-e2e/DeviceLib/ITTAnnotations/atomic.cpp +++ b/sycl/test-e2e/DeviceLib/ITTAnnotations/atomic.cpp @@ -19,9 +19,9 @@ int main() { // ITT start/finish annotations and ITT atomic start/finish annotations. q.submit([&](handler &cgh) { auto source_acc = - source_buf.template get_access(cgh); + source_buf.template get_access(cgh); auto target_acc = - target_buf.template get_access(cgh); + target_buf.template get_access(cgh); cgh.single_task([=]() { auto source_atomic = atomic_ref void check_scalar_nan(s::queue &Queue) { { s::buffer Buf(&Data, s::range<1>(1)); Queue.submit([&](s::handler &Cgh) { - auto Acc = Buf.template get_access(Cgh); + auto Acc = Buf.template get_access(Cgh); Cgh.single_task>([=]() { Acc[0] = s::nan(T{0}); }); }); Queue.wait_and_throw(); @@ -40,7 +40,7 @@ template void check_vec_nan(s::queue &Queue) { { s::buffer, 1> VBuf(&VData, s::range<1>(1)); Queue.submit([&](s::handler &Cgh) { - auto VAcc = VBuf.template get_access(Cgh); + auto VAcc = VBuf.template get_access(Cgh); Cgh.single_task>( [=]() { VAcc[0] = s::nan(s::vec{0}); }); }); diff --git a/sycl/test-e2e/DeviceLib/built-ins/scalar_integer.cpp b/sycl/test-e2e/DeviceLib/built-ins/scalar_integer.cpp index 0c359e10fc8d2..b879ce8c3ba25 100644 --- a/sycl/test-e2e/DeviceLib/built-ins/scalar_integer.cpp +++ b/sycl/test-e2e/DeviceLib/built-ins/scalar_integer.cpp @@ -18,7 +18,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(5, 2); }); }); } @@ -32,7 +32,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(5u, 2u); }); }); } @@ -46,7 +46,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::min(5, 2); }); }); } @@ -60,7 +60,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::min(5ll, 2ll); }); }); @@ -75,7 +75,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::min(5u, 2u); }); }); } @@ -89,7 +89,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::min(5ull, 2ull); }); }); @@ -104,7 +104,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::abs(-5); }); }); } @@ -118,7 +118,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::abs_diff(-5, -1); }); }); @@ -133,7 +133,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::abs_diff((unsigned char)3, (unsigned char)250); }); @@ -149,7 +149,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::add_sat(0x7FFFFFFF, 100); }); }); @@ -164,7 +164,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::hadd(0x0000007F, 0x00000020); }); }); @@ -179,7 +179,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::rhadd(0x0000007F, 0x00000020); }); }); @@ -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::clamp(5, 10, 30); }); }); @@ -209,7 +209,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::clz(0x0FFFFFFF); }); }); } @@ -223,7 +223,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::ctz(0x7FFFFFF0); }); }); @@ -238,7 +238,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::mad_hi(0x10000000, 0x00000100, 0x00000001); }); // 2^28 * 2^8 = 2^36 -> 0x10 00000000. @@ -254,7 +254,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::mad_sat(0x10000000, 0x00000100, 0x00000001); }); // 2^31 * 2^8 = 2^39 -> 0x80 00000000 -> reuslt is saturated in the @@ -272,7 +272,7 @@ int main() { sycl::buffer buf(&r, sycl::range<1>(1)); sycl::queue q; q.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.single_task([=]() { signed char inputData_0(-17); signed char inputData_1(-10); @@ -292,7 +292,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::mul_hi(0x10000000, 0x00000100); }); // 2^28 * 2^8 = 2^36 -> 0x10 00000000. @@ -308,7 +308,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::mul_hi(-0x10000000, 0x00000100); }); // -2^28 * 2^8 = -2^36 -> -0x10 (FFFFFFF0) 00000000. @@ -324,7 +324,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::mul_hi(-0x10000000, 0x00000101); }); // -2^28 * (2^8 + 1) = -2^36 - 2^28 -> -0x11 (FFFFFFEF) -0x10000000 @@ -341,7 +341,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::rotate(0x11100000, 12); }); }); @@ -356,7 +356,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::rotate((char)0xe0, (char)50); }); }); @@ -371,7 +371,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::sub_sat(x, y); }); }); @@ -396,7 +396,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::upsample((unsigned char)0x10, (unsigned char)0x10); }); @@ -412,7 +412,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::upsample((int8_t)0x10, (uint8_t)0x10); }); }); @@ -427,7 +427,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::upsample((unsigned short)0x0010, (unsigned short)0x0010); }); @@ -443,7 +443,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::upsample((short)0x0010, (unsigned short)0x0010); }); @@ -459,7 +459,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::upsample(0x00000010u, 0x00000010u); }); }); @@ -474,7 +474,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::upsample(0x00000010, 0x00000010u); }); }); @@ -489,7 +489,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::popcount(0x000000FF); }); }); @@ -504,7 +504,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::mad24((int)0xFFFFFFFF, (int)20, (int)20); }); }); @@ -519,7 +519,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::mul24((int)0xFFFFFFFF, (int)20); }); }); diff --git a/sycl/test-e2e/DeviceLib/built-ins/scalar_relational.cpp b/sycl/test-e2e/DeviceLib/built-ins/scalar_relational.cpp index d67cc5d79e5df..f493a8108f3e0 100644 --- a/sycl/test-e2e/DeviceLib/built-ins/scalar_relational.cpp +++ b/sycl/test-e2e/DeviceLib/built-ins/scalar_relational.cpp @@ -20,7 +20,7 @@ int main() { { s::buffer BufR(&r, s::range<1>(1)); myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::isequal(float{10.5f}, float{10.5f}); }); }); @@ -34,7 +34,7 @@ int main() { { s::buffer BufR(&r, s::range<1>(1)); myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::isnotequal(float{0.4f}, float{0.5f}); }); }); @@ -48,7 +48,7 @@ int main() { { s::buffer BufR(&r, s::range<1>(1)); myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::isgreater(float{0.6f}, float{0.5f}); }); }); @@ -62,7 +62,7 @@ int main() { { s::buffer BufR(&r, s::range<1>(1)); myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::isgreaterequal(float{0.5f}, float{0.5f}); }); }); @@ -76,7 +76,7 @@ int main() { { s::buffer BufR(&r, s::range<1>(1)); myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::isless(float{0.4f}, float{0.5f}); }); }); @@ -90,7 +90,7 @@ int main() { { s::buffer BufR(&r, s::range<1>(1)); myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::islessequal(float{0.5f}, float{0.5f}); }); }); @@ -104,7 +104,7 @@ int main() { { s::buffer BufR(&r, s::range<1>(1)); myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::islessgreater(float{0.5f}, float{0.5f}); }); }); @@ -118,7 +118,7 @@ int main() { { s::buffer BufR(&r, s::range<1>(1)); myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::isfinite(float{NAN}); }); }); @@ -132,7 +132,7 @@ int main() { { s::buffer BufR(&r, s::range<1>(1)); myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::isinf(float{INFINITY}); }); }); @@ -146,7 +146,7 @@ int main() { { s::buffer BufR(&r, s::range<1>(1)); myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::isnan(float{NAN}); }); }); @@ -160,7 +160,7 @@ int main() { { s::buffer BufR(&r, s::range<1>(1)); myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::isnormal(float{INFINITY}); }); }); @@ -174,7 +174,7 @@ int main() { { s::buffer BufR(&r, s::range<1>(1)); myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::isnormal(double{INFINITY}); }); }); @@ -188,7 +188,7 @@ int main() { { s::buffer BufR(&r, s::range<1>(1)); myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::isordered(float{4.0f}, float{NAN}); }); }); @@ -202,7 +202,7 @@ int main() { { s::buffer BufR(&r, s::range<1>(1)); myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::isunordered(float{4.0f}, float{NAN}); }); }); @@ -216,7 +216,7 @@ int main() { { s::buffer BufR(&r, s::range<1>(1)); myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::signbit(float{-12.0f}); }); }); @@ -230,7 +230,7 @@ int main() { { s::buffer BufR(&r, s::range<1>(1)); myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::any(int{12}); }); }); @@ -243,7 +243,7 @@ int main() { { s::buffer BufR(&r, s::range<1>(1)); myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task([=]() { AccR[0] = s::any(int{0}); }); }); } @@ -256,7 +256,7 @@ int main() { { s::buffer BufR(&r, s::range<1>(1)); myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::any(int{-12}); }); }); @@ -270,7 +270,7 @@ int main() { { s::buffer BufR(&r, s::range<1>(1)); myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::all(int{12}); }); }); @@ -284,7 +284,7 @@ int main() { { s::buffer BufR(&r, s::range<1>(1)); myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task([=]() { AccR[0] = s::all(int{0}); }); }); } @@ -297,7 +297,7 @@ int main() { { s::buffer BufR(&r, s::range<1>(1)); myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::all(int{-12}); }); }); @@ -311,7 +311,7 @@ int main() { { s::buffer BufR(&r, s::range<1>(1)); myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task([=]() { AccR[0] = s::bitselect(float{112.112}, float{34.34}, float{3.3}); }); @@ -326,7 +326,7 @@ int main() { { s::buffer BufR(&r, s::range<1>(1)); myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task([=]() { AccR[0] = s::select(float{34.34}, float{123.123}, int{1}); }); @@ -341,7 +341,7 @@ int main() { { s::buffer BufR(&r, s::range<1>(1)); myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task([=]() { AccR[0] = s::select(float{34.34}, float{123.123}, int{0}); }); @@ -356,7 +356,7 @@ int main() { { s::buffer BufR(&r, s::range<1>(1)); myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task([=]() { AccR[0] = s::select(float{34.34}, float{123.123}, int{-1}); }); @@ -372,7 +372,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::select(34.34f, 123.123f, true); }); }); @@ -387,7 +387,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::select(34.34f, 123.123f, false); }); }); diff --git a/sycl/test-e2e/DeviceLib/built-ins/vector_integer.cpp b/sycl/test-e2e/DeviceLib/built-ins/vector_integer.cpp index 876f800743769..90003329ec9fb 100644 --- a/sycl/test-e2e/DeviceLib/built-ins/vector_integer.cpp +++ b/sycl/test-e2e/DeviceLib/built-ins/vector_integer.cpp @@ -18,7 +18,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::int2{5, 3}, s::int2{2, 7}); }); @@ -37,7 +37,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::uint2{5, 3}, s::uint2{2, 7}); }); @@ -56,7 +56,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::int2{5, 3}, int{2}); }); @@ -76,7 +76,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(T{5, 3}, 2ll); }); }); @@ -94,7 +94,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::uint2{5, 3}, 2u); }); @@ -114,7 +114,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(T{5, 3}, 2ull); }); }); @@ -132,7 +132,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::min(s::int2{5, 3}, s::int2{2, 7}); }); @@ -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::min(s::uint2{5, 3}, s::uint2{2, 7}); }); @@ -170,7 +170,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::min(s::int2{5, 3}, int{2}); }); @@ -189,7 +189,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::min(s::uint2{5, 3}, 2u); }); @@ -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::abs(s::int2{-5, -2}); }); @@ -228,7 +228,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::abs(T{-5, -2}); }); }); } @@ -245,7 +245,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::abs_diff(s::int2{-5, -2}, s::int2{-1, -1}); }); @@ -264,7 +264,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::add_sat(s::int2{0x7FFFFFFF, 0x7FFFFFFF}, s::int2{100, 90}); @@ -284,7 +284,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::hadd(s::int2{0x0000007F, 0x0000007F}, s::int2{0x00000020, 0x00000020}); @@ -304,7 +304,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::rhadd(s::int2{0x0000007F, 0x0000007F}, s::int2{0x00000020, 0x00000020}); @@ -324,7 +324,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::clamp(s::int2{5, 5}, s::int2{10, 10}, s::int2{30, 30}); }); @@ -343,7 +343,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::clamp(s::int2{5, 5}, int{10}, int{30}); }); @@ -362,7 +362,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::clz(s::int2{0x0FFFFFFF, 0x0FFFFFFF}); }); @@ -381,7 +381,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::ctz(s::int2{0x7FFFFFF0, 0x7FFFFFF0}); }); @@ -400,7 +400,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::mad_hi(s::int2{0x10000000, 0x10000000}, s::int2{0x00000100, 0x00000100}, s::int2{1, 1}); @@ -420,7 +420,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::mad_sat(s::int2{0x10000000, 0x10000000}, s::int2{0x00000100, 0x00000100}, s::int2{1, 1}); @@ -440,7 +440,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::mul_hi(s::int2{0x10000000, 0x10000000}, s::int2{0x00000100, 0x00000100}); @@ -460,7 +460,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::rotate(s::int2{0x11100000, 0x11100000}, s::int2{12, 12}); }); @@ -480,7 +480,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::sub_sat(x, y); }); }); @@ -512,7 +512,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::upsample(s::uchar2{0x10, 0x10}, s::uchar2{0x10, 0x10}); }); @@ -531,7 +531,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::upsample(s::char2{0x10, 0x10}, s::uchar2{0x10, 0x10}); }); @@ -550,7 +550,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::upsample(s::ushort2{0x0010, 0x0010}, s::ushort2{0x0010, 0x0010}); @@ -570,7 +570,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::upsample(s::short2{0x0010, 0x0010}, s::ushort2{0x0010, 0x0010}); @@ -590,7 +590,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::upsample(s::uint2{0x00000010, 0x00000010}, s::uint2{0x00000010, 0x00000010}); @@ -610,7 +610,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::upsample(s::int2{0x00000010, 0x00000010}, s::uint2{0x00000010, 0x00000010}); @@ -630,7 +630,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::popcount(s::int2{0x000000FF, 0x000000FF}); }); @@ -649,7 +649,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::mad24(s::int2{0xFFFFFFFF, 0xFFFFFFFF}, s::int2{20, 20}, s::int2{20, 20}); @@ -669,7 +669,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::mul24(s::int2{0xFFFFFFFF, 0xFFFFFFFF}, s::int2{20, 20}); }); diff --git a/sycl/test-e2e/DeviceLib/c99_complex_math_fp64_test.cpp b/sycl/test-e2e/DeviceLib/c99_complex_math_fp64_test.cpp index 617cd67a559cc..30e40e8c2a747 100644 --- a/sycl/test-e2e/DeviceLib/c99_complex_math_fp64_test.cpp +++ b/sycl/test-e2e/DeviceLib/c99_complex_math_fp64_test.cpp @@ -58,7 +58,7 @@ double _Complex catanh(double _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_d(double _Complex x, double _Complex y) { return approx_equal_fp(__real__ x, __real__ y) && diff --git a/sycl/test-e2e/DeviceLib/cmath_test.cpp b/sycl/test-e2e/DeviceLib/cmath_test.cpp index c774e53ad3b8e..d3aa2cc266daa 100644 --- a/sycl/test-e2e/DeviceLib/cmath_test.cpp +++ b/sycl/test-e2e/DeviceLib/cmath_test.cpp @@ -18,8 +18,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 71 diff --git a/sycl/test-e2e/DeviceLib/imf/fp64_test.cpp b/sycl/test-e2e/DeviceLib/imf/fp64_test.cpp index 89c94e176e470..59351289eccaa 100644 --- a/sycl/test-e2e/DeviceLib/imf/fp64_test.cpp +++ b/sycl/test-e2e/DeviceLib/imf/fp64_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; int main(int, char **) { s::queue device_queue(s::default_selector_v); diff --git a/sycl/test-e2e/DeviceLib/math_fp64_windows_test.cpp b/sycl/test-e2e/DeviceLib/math_fp64_windows_test.cpp index a228324c43400..bc74864605c6b 100644 --- a/sycl/test-e2e/DeviceLib/math_fp64_windows_test.cpp +++ b/sycl/test-e2e/DeviceLib/math_fp64_windows_test.cpp @@ -11,8 +11,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 38 diff --git a/sycl/test-e2e/DeviceLib/math_test_marray_vec_common.hpp b/sycl/test-e2e/DeviceLib/math_test_marray_vec_common.hpp index ad5dcd766ab65..989174c2e056c 100644 --- a/sycl/test-e2e/DeviceLib/math_test_marray_vec_common.hpp +++ b/sycl/test-e2e/DeviceLib/math_test_marray_vec_common.hpp @@ -33,9 +33,9 @@ template bool checkEqual(marray A, size_t B) { buffer buffer1(&result, 1); \ buffer buffer2(&input, 1); \ deviceQueue.submit([&](handler &cgh) { \ - accessor res_access( \ - buffer1, cgh); \ - accessor input_access( \ + accessor res_access(buffer1, \ + cgh); \ + accessor input_access( \ buffer2, cgh); \ cgh.single_task>( \ [=]() { res_access[0] = NAME(input_access[0]); }); \ @@ -94,11 +94,11 @@ OPERATOR(trunc) buffer buffer2(&input1, 1); \ buffer buffer3(&input2, 1); \ deviceQueue.submit([&](handler &cgh) { \ - accessor res_access( \ - buffer1, cgh); \ - accessor input1_access( \ + accessor res_access(buffer1, \ + cgh); \ + accessor input1_access( \ buffer2, cgh); \ - accessor input2_access( \ + accessor input2_access( \ buffer3, cgh); \ cgh.single_task>([=]() { \ res_access[0] = NAME(input1_access[0], input2_access[0]); \ @@ -135,13 +135,13 @@ OPERATOR_2(remainder) buffer buffer3(&input2, 1); \ buffer buffer4(&input3, 1); \ deviceQueue.submit([&](handler &cgh) { \ - accessor res_access( \ - buffer1, cgh); \ - accessor input1_access( \ + accessor res_access(buffer1, \ + cgh); \ + accessor input1_access( \ buffer2, cgh); \ - accessor input2_access( \ + accessor input2_access( \ buffer3, cgh); \ - accessor input3_access( \ + accessor input3_access( \ buffer4, cgh); \ cgh.single_task>([=]() { \ res_access[0] = \ diff --git a/sycl/test-e2e/DeviceLib/std_complex_math_fp64_test.cpp b/sycl/test-e2e/DeviceLib/std_complex_math_fp64_test.cpp index c3fa5003e764d..1b206c6c09f1d 100644 --- a/sycl/test-e2e/DeviceLib/std_complex_math_fp64_test.cpp +++ b/sycl/test-e2e/DeviceLib/std_complex_math_fp64_test.cpp @@ -11,8 +11,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_int_test.cpp b/sycl/test-e2e/DotProduct/dot_product_int_test.cpp index 0a0542149d866..ecdeae9bfa854 100644 --- a/sycl/test-e2e/DotProduct/dot_product_int_test.cpp +++ b/sycl/test-e2e/DotProduct/dot_product_int_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/EnqueueNativeCommand/custom-command-hip.cpp b/sycl/test-e2e/EnqueueNativeCommand/custom-command-hip.cpp index b93c66a467a8f..7ff494cc94e5e 100644 --- a/sycl/test-e2e/EnqueueNativeCommand/custom-command-hip.cpp +++ b/sycl/test-e2e/EnqueueNativeCommand/custom-command-hip.cpp @@ -35,8 +35,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 HipStream = IH.get_native_queue(); @@ -61,7 +61,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; }; @@ -72,8 +72,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] = B1Init; @@ -87,7 +87,7 @@ void test_ht_buffer(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.ext_codeplay_enqueue_native_command(Func); }); diff --git a/sycl/test-e2e/Experimental/launch_queries/max_sub_groups.cpp b/sycl/test-e2e/Experimental/launch_queries/max_sub_groups.cpp index 9b0e71864dfca..3985e786db09c 100644 --- a/sycl/test-e2e/Experimental/launch_queries/max_sub_groups.cpp +++ b/sycl/test-e2e/Experimental/launch_queries/max_sub_groups.cpp @@ -21,7 +21,7 @@ namespace kernels { template using sycl_global_accessor = - sycl::accessor; class TestKernel { @@ -50,7 +50,7 @@ int main() { sycl::buffer buf{sycl::range<1>{1}}; auto launchRange = sycl::nd_range<1>{sycl::range<1>{1}, sycl::range<1>{1}}; q.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.parallel_for(launchRange, kernels::TestKernel{acc}); }).wait(); diff --git a/sycl/test-e2e/Experimental/reusable_events/cross_queue_dependency.cpp b/sycl/test-e2e/Experimental/reusable_events/cross_queue_dependency.cpp index 9533249683cbf..b97384106481f 100644 --- a/sycl/test-e2e/Experimental/reusable_events/cross_queue_dependency.cpp +++ b/sycl/test-e2e/Experimental/reusable_events/cross_queue_dependency.cpp @@ -32,7 +32,7 @@ int main() { // Launch kernel on q1 q1.submit([&](sycl::handler &cgh) { - auto acc1 = buf1.get_access(cgh); + auto acc1 = buf1.get_access(cgh); cgh.parallel_for( sycl::range<1>(N), [=](sycl::id<1> idx) { acc1[idx] = acc1[idx] * 2; }); }); @@ -45,9 +45,9 @@ int main() { // Launch kernel on q2 q2.submit([&](sycl::handler &cgh) { - auto acc1 = buf1.get_access(cgh); - auto acc2 = buf2.get_access(cgh); - auto acc_result = buf_result.get_access(cgh); + auto acc1 = buf1.get_access(cgh); + auto acc2 = buf2.get_access(cgh); + auto acc_result = buf_result.get_access(cgh); cgh.parallel_for(sycl::range<1>(N), [=](sycl::id<1> idx) { acc_result[idx] = acc1[idx] + acc2[idx]; }); diff --git a/sycl/test-e2e/Experimental/reusable_events/event_reuse.cpp b/sycl/test-e2e/Experimental/reusable_events/event_reuse.cpp index f1cf3440223f4..b9030a28d3812 100644 --- a/sycl/test-e2e/Experimental/reusable_events/event_reuse.cpp +++ b/sycl/test-e2e/Experimental/reusable_events/event_reuse.cpp @@ -23,7 +23,7 @@ int main() { // First operation 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; }); }); @@ -32,7 +32,7 @@ int main() { // Reuse event for second operation 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] + 10; }); }); @@ -41,7 +41,7 @@ int main() { // Reuse event for third operation 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/out_of_order_queue_signal.cpp b/sycl/test-e2e/Experimental/reusable_events/out_of_order_queue_signal.cpp index c9fe4341d8f4b..971c674f3801e 100644 --- a/sycl/test-e2e/Experimental/reusable_events/out_of_order_queue_signal.cpp +++ b/sycl/test-e2e/Experimental/reusable_events/out_of_order_queue_signal.cpp @@ -21,7 +21,7 @@ int main() { sycl::buffer buf(data.data(), sycl::range<1>(N)); 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] = 5; }); }); diff --git a/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp b/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp index cd4bcfa532451..ebd6ba2fdd96b 100644 --- a/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp +++ b/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp @@ -18,7 +18,7 @@ namespace ns { template SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel)) void nsNdRangeFreeFunc( - sycl::accessor Accessor, + sycl::accessor Accessor, int Value) { auto Item = syclext::this_work_item::get_nd_item().get_global_id(); Accessor[Item] = Value; @@ -28,11 +28,11 @@ void nsNdRangeFreeFunc( template SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel)) void ndRangeFreeFuncMultipleParameters( - sycl::accessor InputAAcc, sycl::accessor InputBAcc, - sycl::accessor ResultAcc) { + sycl::accessor ResultAcc) { auto Item = syclext::this_work_item::get_nd_item().get_global_id(); ResultAcc[Item] = InputAAcc[Item] + InputBAcc[Item]; } @@ -97,13 +97,13 @@ int runNdRangeTestMultipleParameters(sycl::queue &Queue, sycl::context &Context, sycl::buffer ResultBuffer(ResultData.data(), NdRange.get_global_range()); Queue.submit([&](sycl::handler &Handler) { - sycl::accessor InputAAcc{InputABuffer, Handler}; - sycl::accessor InputBAcc{InputBBuffer, Handler}; - sycl::accessor ResultAcc{ + sycl::accessor ResultAcc{ ResultBuffer, Handler}; Handler.set_args(InputAAcc, InputBAcc, ResultAcc); Handler.parallel_for(NdRange, UsedKernel); diff --git a/sycl/test-e2e/Functor/functor_inheritance.cpp b/sycl/test-e2e/Functor/functor_inheritance.cpp index bcb8f2d9394cc..4b83b288917e1 100644 --- a/sycl/test-e2e/Functor/functor_inheritance.cpp +++ b/sycl/test-e2e/Functor/functor_inheritance.cpp @@ -4,7 +4,7 @@ #include #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; struct SecondBase { diff --git a/sycl/test-e2e/GroupAlgorithm/barrier.cpp b/sycl/test-e2e/GroupAlgorithm/barrier.cpp index d6c85823fe430..0d521de41d7b6 100644 --- a/sycl/test-e2e/GroupAlgorithm/barrier.cpp +++ b/sycl/test-e2e/GroupAlgorithm/barrier.cpp @@ -23,7 +23,7 @@ void basic() { buffer buf(data.data(), N); q.submit([&](handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); local_accessor loc(N, cgh); local_accessor loc_barrier(2, cgh); cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> item) { @@ -66,9 +66,9 @@ void interface() { buffer test2_buf(test2.data(), N); q.submit([&](handler &cgh) { - auto data_acc = data_buf.get_access(cgh); - auto test1_acc = test1_buf.get_access(cgh); - auto test2_acc = test2_buf.get_access(cgh); + auto data_acc = data_buf.get_access(cgh); + auto test1_acc = test1_buf.get_access(cgh); + auto test2_acc = test2_buf.get_access(cgh); local_accessor loc(N, cgh); local_accessor loc_barrier(2, cgh); cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> item) { diff --git a/sycl/test-e2e/GroupLocalMemory/group_local_memory.cpp b/sycl/test-e2e/GroupLocalMemory/group_local_memory.cpp index ca19a74d59a9a..7253cddb0f9ea 100644 --- a/sycl/test-e2e/GroupLocalMemory/group_local_memory.cpp +++ b/sycl/test-e2e/GroupLocalMemory/group_local_memory.cpp @@ -41,8 +41,8 @@ int main() { buffer CounterBuf{CounterVec.data(), range<1>(WgCount)}; Q.submit([&](handler &Cgh) { - auto Acc = Buf.get_access(Cgh); - auto CounterAcc = CounterBuf.get_access(Cgh); + auto Acc = Buf.get_access(Cgh); + auto CounterAcc = CounterBuf.get_access(Cgh); Cgh.parallel_for( nd_range<1>(range<1>(Size), range<1>(WgSize)), [=](nd_item<1> Item) { // Some alternative (and functionally equivalent) ways to use this @@ -80,7 +80,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) { multi_ptr 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) { multi_ptr; + accessor; struct PFWIFunctor { PFWIFunctor(size_t wg_chunk, size_t wg_size, size_t wg_offset, @@ -134,7 +134,7 @@ int main() { const int addend = 10; myQueue.submit([&](handler &cgh) { - auto dev_ptr = buf.get_access(cgh); + auto dev_ptr = buf.get_access(cgh); // number of 'buf' elements per work group: size_t wg_chunk = (range_length + N_WG - 1) / N_WG; PFWGFunctor PFWG(wg_chunk, range_length, addend, N_ITER, dev_ptr); @@ -154,7 +154,7 @@ int main() { std::memset(ptr, 0, range_length * sizeof(ptr[0])); buffer buf(ptr, range<1>(range_length)); myQueue.submit([&](handler &cgh) { - auto dev_ptr = buf.get_access(cgh); + auto dev_ptr = buf.get_access(cgh); cgh.parallel_for_work_group( range<1>(N_WG), range<1>(WG_SIZE_PHYSICAL), [=](group<1> g) { @@ -194,7 +194,7 @@ int main() { std::memset(ptr, 0, range_length * sizeof(ptr[0])); buffer buf(ptr, range<1>(range_length)); myQueue.submit([&](handler &cgh) { - auto dev_ptr = buf.get_access(cgh); + auto dev_ptr = buf.get_access(cgh); cgh.parallel_for_work_group( range<1>(N_WG), range<1>(WG_SIZE_PHYSICAL), [=](group<1> g) { @@ -237,7 +237,7 @@ int main() { std::memset(ptr, 0, range_length * sizeof(ptr[0])); buffer buf(ptr, range<1>(range_length)); myQueue.submit([&](handler &cgh) { - auto dev_ptr = buf.get_access(cgh); + auto dev_ptr = buf.get_access(cgh); cgh.parallel_for_work_group( range<2>(N_WG, 1), range<2>(WG_X_SIZE, WG_Y_SIZE), [=](group<2> g) { @@ -278,7 +278,7 @@ int main() { std::memset(ptr, 0, range_length * sizeof(ptr[0])); buffer buf(ptr, range<1>(range_length)); myQueue.submit([&](handler &cgh) { - auto dev_ptr = buf.get_access(cgh); + auto dev_ptr = buf.get_access(cgh); cgh.parallel_for_work_group( range<1>(N_WG), range<1>(WG_SIZE_PHYSICAL), [=](group<1> g) { diff --git a/sycl/test-e2e/HostInteropTask/host-task-failure.cpp b/sycl/test-e2e/HostInteropTask/host-task-failure.cpp index 88855a03e78c9..a24170b3564a3 100644 --- a/sycl/test-e2e/HostInteropTask/host-task-failure.cpp +++ b/sycl/test-e2e/HostInteropTask/host-task-failure.cpp @@ -18,8 +18,8 @@ template class Init; 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); CGH.host_task([=]() { for (size_t Idx = 0; Idx < SrcA.size(); ++Idx) @@ -31,8 +31,8 @@ void copy(buffer &Src, buffer &Dst, 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; diff --git a/sycl/test-e2e/HostInteropTask/interop-task-hip.cpp b/sycl/test-e2e/HostInteropTask/interop-task-hip.cpp index 3d854590ad43c..43bb76ebadd57 100644 --- a/sycl/test-e2e/HostInteropTask/interop-task-hip.cpp +++ b/sycl/test-e2e/HostInteropTask/interop-task-hip.cpp @@ -36,8 +36,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 HipStream = IH.get_native_queue(); @@ -62,7 +62,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; }; @@ -73,8 +73,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] = B1Init; @@ -88,7 +88,7 @@ void test_ht_buffer(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); }); diff --git a/sycl/test-e2e/InlineAsm/asm_16_matrix_mult.cpp b/sycl/test-e2e/InlineAsm/asm_16_matrix_mult.cpp index e73772aa00885..d4bca3eb82966 100644 --- a/sycl/test-e2e/InlineAsm/asm_16_matrix_mult.cpp +++ b/sycl/test-e2e/InlineAsm/asm_16_matrix_mult.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_empty.cpp b/sycl/test-e2e/InlineAsm/asm_8_empty.cpp index f784205bc6029..acbab149fae62 100644 --- a/sycl/test-e2e/InlineAsm/asm_8_empty.cpp +++ b/sycl/test-e2e/InlineAsm/asm_8_empty.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_decl_in_scope.cpp b/sycl/test-e2e/InlineAsm/asm_decl_in_scope.cpp index 3e642d09edc67..9884db6234db1 100644 --- a/sycl/test-e2e/InlineAsm/asm_decl_in_scope.cpp +++ b/sycl/test-e2e/InlineAsm/asm_decl_in_scope.cpp @@ -18,13 +18,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_float_neg.cpp b/sycl/test-e2e/InlineAsm/asm_float_neg.cpp index 89a6b2f249f0e..27ff61a407260 100644 --- a/sycl/test-e2e/InlineAsm/asm_float_neg.cpp +++ b/sycl/test-e2e/InlineAsm/asm_float_neg.cpp @@ -17,10 +17,10 @@ struct KernelFunctor : WithInputBuffers, WithOutputBuffer { void operator()(sycl::handler &cgh) { auto A = - this->getInputBuffer().template get_access( + this->getInputBuffer().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/InlineAsm/asm_loop.cpp b/sycl/test-e2e/InlineAsm/asm_loop.cpp index c9c74b835e2a0..dbf96ca05c21b 100644 --- a/sycl/test-e2e/InlineAsm/asm_loop.cpp +++ b/sycl/test-e2e/InlineAsm/asm_loop.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>( sycl::range<1>{this->getOutputBufferSize()}, diff --git a/sycl/test-e2e/InlineAsm/asm_no_output.cpp b/sycl/test-e2e/InlineAsm/asm_no_output.cpp index 2445b56cf7cf3..e2143a00333f6 100644 --- a/sycl/test-e2e/InlineAsm/asm_no_output.cpp +++ b/sycl/test-e2e/InlineAsm/asm_no_output.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/InorderQueue/in_order_buffs.cpp b/sycl/test-e2e/InorderQueue/in_order_buffs.cpp index 52618ac084027..864bf0705c092 100644 --- a/sycl/test-e2e/InorderQueue/in_order_buffs.cpp +++ b/sycl/test-e2e/InorderQueue/in_order_buffs.cpp @@ -32,7 +32,7 @@ int main() { buffer bufA(dataA, range<1>(dataSize)); buffer bufB(dataB, range<1>(dataSize)); Queue.submit([&](handler &cgh) { - auto writeBuffer = bufA.get_access(cgh); + auto writeBuffer = bufA.get_access(cgh); // Create a range. auto myRange = range<1>(dataSize); @@ -44,8 +44,8 @@ int main() { }); Queue.submit([&](handler &cgh) { - auto writeBuffer = bufB.get_access(cgh); - auto readBuffer = bufA.get_access(cgh); + auto writeBuffer = bufB.get_access(cgh); + auto readBuffer = bufA.get_access(cgh); // Create a range. auto myRange = range<1>(dataSize); diff --git a/sycl/test-e2e/KernelCompiler/opencl.cpp b/sycl/test-e2e/KernelCompiler/opencl.cpp index 5fbcfa994b507..de709fcff425b 100644 --- a/sycl/test-e2e/KernelCompiler/opencl.cpp +++ b/sycl/test-e2e/KernelCompiler/opencl.cpp @@ -62,8 +62,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/KernelCompiler/spirv.cpp b/sycl/test-e2e/KernelCompiler/spirv.cpp index 3f39146643948..ef4ffd3a95ea6 100644 --- a/sycl/test-e2e/KernelCompiler/spirv.cpp +++ b/sycl/test-e2e/KernelCompiler/spirv.cpp @@ -61,8 +61,8 @@ void testSimpleKernel(sycl::queue &q, const sycl::kernel &kernel, sycl::buffer output_buffer(output_array.data(), sycl::range<1>(N)); q.submit([&](sycl::handler &cgh) { - cgh.set_arg(0, input_buffer.get_access(cgh)); - cgh.set_arg(1, output_buffer.get_access(cgh)); + cgh.set_arg(0, input_buffer.get_access(cgh)); + cgh.set_arg(1, output_buffer.get_access(cgh)); cgh.parallel_for(sycl::range<1>{N}, kernel); }); @@ -98,7 +98,7 @@ void testParam(sycl::queue &q, const sycl::kernel &kernel) { cgh.set_arg(1, b_ptr); // Pass sycl::accessor for OpTypePointer(CrossWorkgroup) parameter. cgh.set_arg( - 2, output_buffer.template get_access(cgh)); + 2, output_buffer.template get_access(cgh)); // Pass sycl::local_accessor for OpTypePointer(Workgroup) parameter. cgh.set_arg(3, local); cgh.parallel_for(sycl::range<1>{1}, kernel); diff --git a/sycl/test-e2e/KernelParams/non-standard-layout.cpp b/sycl/test-e2e/KernelParams/non-standard-layout.cpp index 39d6640b97c4f..16271ea54962c 100644 --- a/sycl/test-e2e/KernelParams/non-standard-layout.cpp +++ b/sycl/test-e2e/KernelParams/non-standard-layout.cpp @@ -21,7 +21,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; }); }); } diff --git a/sycl/test-e2e/OptionalKernelFeatures/throw-exception-for-out-of-registers-on-kernel-launch.cpp b/sycl/test-e2e/OptionalKernelFeatures/throw-exception-for-out-of-registers-on-kernel-launch.cpp index 7915459199c93..9844edd108488 100644 --- a/sycl/test-e2e/OptionalKernelFeatures/throw-exception-for-out-of-registers-on-kernel-launch.cpp +++ b/sycl/test-e2e/OptionalKernelFeatures/throw-exception-for-out-of-registers-on-kernel-launch.cpp @@ -65,11 +65,11 @@ int main() { try { q.submit([&](sycl::handler &h) { - auto input1 = valuesBuf1.get_access(h); - auto input2 = valuesBuf2.get_access(h); - auto input3 = valuesBuf3.get_access(h); - auto input4 = valuesBuf4.get_access(h); - auto output = outputBuf.get_access(h); + auto input1 = valuesBuf1.get_access(h); + auto input2 = valuesBuf2.get_access(h); + auto input3 = valuesBuf3.get_access(h); + auto input4 = valuesBuf4.get_access(h); + auto output = outputBuf.get_access(h); h.parallel_for( sycl::nd_range<1>{{GLOBAL_WORK_SIZE}, {local_size}}, [=](sycl::id<1> i) { diff --git a/sycl/test-e2e/Reduction/reduction_nd_N_vars.cpp b/sycl/test-e2e/Reduction/reduction_nd_N_vars.cpp index 28813bcc09ff1..a9bb32e874d75 100644 --- a/sycl/test-e2e/Reduction/reduction_nd_N_vars.cpp +++ b/sycl/test-e2e/Reduction/reduction_nd_N_vars.cpp @@ -12,8 +12,8 @@ using namespace sycl; template class KNameGroup; template class KName; -constexpr access::mode RW = access::mode::read_write; -constexpr access::mode DW = access::mode::discard_write; +constexpr access_mode RW = access_mode::read_write; +constexpr access_mode DW = access_mode::discard_write; template void printNVarsTestLabel(const RangeT &Range, bool ToCERR = false) { diff --git a/sycl/test-e2e/Reduction/reduction_nd_lambda.cpp b/sycl/test-e2e/Reduction/reduction_nd_lambda.cpp index aedd72bd89649..44ec77de22ed4 100644 --- a/sycl/test-e2e/Reduction/reduction_nd_lambda.cpp +++ b/sycl/test-e2e/Reduction/reduction_nd_lambda.cpp @@ -8,7 +8,7 @@ using namespace sycl; -template class MName; +template class MName; int NumErrors = 0; diff --git a/sycl/test-e2e/Reduction/reduction_range_N_vars.cpp b/sycl/test-e2e/Reduction/reduction_range_N_vars.cpp index e3aec17126342..af5faa26db2dc 100644 --- a/sycl/test-e2e/Reduction/reduction_range_N_vars.cpp +++ b/sycl/test-e2e/Reduction/reduction_range_N_vars.cpp @@ -12,8 +12,8 @@ using namespace sycl; template class KNameGroup; template class KName; -constexpr access::mode RW = access::mode::read_write; -constexpr access::mode DW = access::mode::discard_write; +constexpr access_mode RW = access_mode::read_write; +constexpr access_mode DW = access_mode::discard_write; template void printNVarsTestLabel(const RangeT &Range, bool ToCERR = false) { diff --git a/sycl/test-e2e/Reduction/reduction_utils.hpp b/sycl/test-e2e/Reduction/reduction_utils.hpp index eeff47dd77c38..439b7a733f1a4 100644 --- a/sycl/test-e2e/Reduction/reduction_utils.hpp +++ b/sycl/test-e2e/Reduction/reduction_utils.hpp @@ -373,7 +373,7 @@ int testInner(queue &Q, OptionalIdentity Identity, T Init, } }; - auto In = InBuf.template get_access(CGH); + auto In = InBuf.template get_access(CGH); auto Redu = CreateReduction(); if constexpr (IsRange) CGH.parallel_for(Range, Redu, [=](item Id, auto &Sum) { @@ -484,7 +484,7 @@ int testUSMInner(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>( Range, Redu, [=](item Id, auto &Sum) { @@ -498,7 +498,7 @@ int testUSMInner(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/build_log.cpp b/sycl/test-e2e/Regression/build_log.cpp index f395eef0e4729..eb9d01c075da4 100644 --- a/sycl/test-e2e/Regression/build_log.cpp +++ b/sycl/test-e2e/Regression/build_log.cpp @@ -87,39 +87,39 @@ void test() { buffer ah{&app_var_37, range(1)}; buffer ai{&app_var_38, range(1)}; b.submit([&](handler &aj) { - auto ak = c.get_access(aj); - auto al = d.get_access(aj); - auto am = e.get_access(aj); - auto an = f.get_access(aj); - auto ao = g.get_access(aj); - auto ap = h.get_access(aj); - auto aq = i.get_access(aj); - auto ar = j.get_access(aj); - auto as = k.get_access(aj); - auto at = l.get_access(aj); - auto au = m.get_access(aj); - auto av = n.get_access(aj); - auto aw = o.get_access(aj); - auto ax = p.get_access(aj); - auto ay = q.get_access(aj); - auto az = r.get_access(aj); - auto ba = s.get_access(aj); - auto bb = t.get_access(aj); - auto bc = u.get_access(aj); - auto bd = v.get_access(aj); - auto be = w.get_access(aj); - auto bf = x.get_access(aj); - auto bg = y.get_access(aj); - auto bh = z.get_access(aj); - auto bi = aa.get_access(aj); - auto bj = ab.get_access(aj); - auto bk = ac.get_access(aj); - auto bl = ad.get_access(aj); - auto bm = ae.get_access(aj); - auto bn = af.get_access(aj); - auto bo = ag.get_access(aj); - auto bp = ah.get_access(aj); - auto bq = ai.get_access(aj); + auto ak = c.get_access(aj); + auto al = d.get_access(aj); + auto am = e.get_access(aj); + auto an = f.get_access(aj); + auto ao = g.get_access(aj); + auto ap = h.get_access(aj); + auto aq = i.get_access(aj); + auto ar = j.get_access(aj); + auto as = k.get_access(aj); + auto at = l.get_access(aj); + auto au = m.get_access(aj); + auto av = n.get_access(aj); + auto aw = o.get_access(aj); + auto ax = p.get_access(aj); + auto ay = q.get_access(aj); + auto az = r.get_access(aj); + auto ba = s.get_access(aj); + auto bb = t.get_access(aj); + auto bc = u.get_access(aj); + auto bd = v.get_access(aj); + auto be = w.get_access(aj); + auto bf = x.get_access(aj); + auto bg = y.get_access(aj); + auto bh = z.get_access(aj); + auto bi = aa.get_access(aj); + auto bj = ab.get_access(aj); + auto bk = ac.get_access(aj); + auto bl = ad.get_access(aj); + auto bm = ae.get_access(aj); + auto bn = af.get_access(aj); + auto bo = ag.get_access(aj); + auto bp = ah.get_access(aj); + auto bq = ai.get_access(aj); aj.single_task([=] { au[0] = am[0]; if (ao[0] ? as[0] ? ar[0] : 9 : ap[0]) diff --git a/sycl/test-e2e/Regression/event_destruction.cpp b/sycl/test-e2e/Regression/event_destruction.cpp index 2d048824456c4..63e11c398c568 100644 --- a/sycl/test-e2e/Regression/event_destruction.cpp +++ b/sycl/test-e2e/Regression/event_destruction.cpp @@ -21,7 +21,7 @@ int main() { sycl::buffer Buf(1); for (size_t Idx = 0; Idx < ITERS; ++Idx) { auto Event = Q.submit([&](sycl::handler &cgh) { - auto Acc = Buf.get_access(cgh); + auto Acc = Buf.get_access(cgh); cgh.single_task([=]() { Acc[0] = 1; }); }); Event.wait(); diff --git a/sycl/test-e2e/Regression/fsycl-host-compiler.cpp b/sycl/test-e2e/Regression/fsycl-host-compiler.cpp index bb08f7955da1f..cc3ba7ae21fbc 100644 --- a/sycl/test-e2e/Regression/fsycl-host-compiler.cpp +++ b/sycl/test-e2e/Regression/fsycl-host-compiler.cpp @@ -26,7 +26,7 @@ int main() { buffer b(data, range<1>(3), {property::buffer::use_host_ptr()}); queue q; q.submit([&](handler &cgh) { - auto B = b.get_access(cgh); + auto B = b.get_access(cgh); cgh.parallel_for(range<1>(3), [=](id<1> idx) { B[idx] = 1; }); }); } diff --git a/sycl/test-e2e/Regression/implicit_atomic_conversion.cpp b/sycl/test-e2e/Regression/implicit_atomic_conversion.cpp index 00ea968a5b53f..c15d66965411f 100644 --- a/sycl/test-e2e/Regression/implicit_atomic_conversion.cpp +++ b/sycl/test-e2e/Regression/implicit_atomic_conversion.cpp @@ -12,7 +12,7 @@ void test_conversion(queue q) { buffer in_buf(&init, 1); q.submit([&](handler &cgh) { - auto in = in_buf.template get_access(cgh); + auto in = in_buf.template get_access(cgh); cgh.single_task([=]() { sycl::atomic atm = in[0]; atm.store(42); diff --git a/sycl/test-e2e/Regression/isordered.cpp b/sycl/test-e2e/Regression/isordered.cpp index 05ec9c0b73574..d53b5ae49b6ef 100644 --- a/sycl/test-e2e/Regression/isordered.cpp +++ b/sycl/test-e2e/Regression/isordered.cpp @@ -13,7 +13,7 @@ int main() { { sycl::buffer buffer(&kernelResult[0], ndRng); testQueue.submit([&](sycl::handler &h) { - auto resultPtr = buffer.template get_access(h); + auto resultPtr = buffer.template get_access(h); h.single_task([=]() { float inputData_0F(0.1); float inputData_1F(0.5); diff --git a/sycl/test-e2e/Regression/kernel_unnamed.cpp b/sycl/test-e2e/Regression/kernel_unnamed.cpp index bff1a8b82cc70..dec04158e73ef 100644 --- a/sycl/test-e2e/Regression/kernel_unnamed.cpp +++ b/sycl/test-e2e/Regression/kernel_unnamed.cpp @@ -18,7 +18,7 @@ static int NumTestCases = 0; template void foo(sycl::queue &deviceQueue, sycl::buffer &buf, F f) { deviceQueue.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.single_task([=]() { acc[0] = f(acc[0], GOLD); }); }); } @@ -33,7 +33,7 @@ struct Wrapper { sycl::queue deviceQueue; sycl::buffer buf(arr, 1); deviceQueue.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.single_task([=]() { acc[0] += GOLD; }); }); ++NumTestCases; @@ -43,7 +43,7 @@ struct Wrapper { [] {}(); #endif deviceQueue.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.single_task([=]() { acc[0] += GOLD; }); }); ++NumTestCases; diff --git a/sycl/test-e2e/Regression/msvc_crt.cpp b/sycl/test-e2e/Regression/msvc_crt.cpp index 9d59547e50d7b..fd299b06e0b9f 100644 --- a/sycl/test-e2e/Regression/msvc_crt.cpp +++ b/sycl/test-e2e/Regression/msvc_crt.cpp @@ -27,7 +27,7 @@ int main() { buffer b(data, range<1>(3), {property::buffer::use_host_ptr()}); queue q; q.submit([&](handler &cgh) { - auto B = b.get_access(cgh); + auto B = b.get_access(cgh); cgh.parallel_for(range<1>(3), [=](id<1> idx) { B[idx] = 1; }); }); } diff --git a/sycl/test-e2e/Regression/nontrivial_device_copyable_value.cpp b/sycl/test-e2e/Regression/nontrivial_device_copyable_value.cpp index 6c8c41292c2b8..ac6e1b0ab776b 100644 --- a/sycl/test-e2e/Regression/nontrivial_device_copyable_value.cpp +++ b/sycl/test-e2e/Regression/nontrivial_device_copyable_value.cpp @@ -28,7 +28,7 @@ int main() { buffer Buff{&Res, range<1>{1}}; Q.submit([&](sycl::handler &cgh) { - auto Acc = Buff.get_access(cgh); + auto Acc = Buff.get_access(cgh); cgh.single_task([=] { Acc[0] = C.i; }); }); } diff --git a/sycl/test-e2e/Regression/reduction_resource_leak_dw.cpp b/sycl/test-e2e/Regression/reduction_resource_leak_dw.cpp index 46bbf02ee6881..2e85d201e021a 100644 --- a/sycl/test-e2e/Regression/reduction_resource_leak_dw.cpp +++ b/sycl/test-e2e/Regression/reduction_resource_leak_dw.cpp @@ -26,7 +26,7 @@ int main() { buffer OutBuf(1); buffer InBuf(49 * 5); Q.submit([&](handler &CGH) { - auto In = InBuf.get_access(CGH); + auto In = InBuf.get_access(CGH); auto Redu = reduction(OutBuf, CGH, 0, BOp, {property::reduction::initialize_to_identity{}}); CGH.parallel_for( diff --git a/sycl/test-e2e/Regression/static-buffer-dtor.cpp b/sycl/test-e2e/Regression/static-buffer-dtor.cpp index 82633f4cfb047..39e3c911c162c 100644 --- a/sycl/test-e2e/Regression/static-buffer-dtor.cpp +++ b/sycl/test-e2e/Regression/static-buffer-dtor.cpp @@ -24,7 +24,7 @@ int main() { sycl::range<1>(256)}; sycl::queue q; q.submit([&](sycl::handler &cgh) { - auto acc = bufs[0].get_access(cgh); + auto acc = bufs[0].get_access(cgh); cgh.single_task([=]() { for (int i = 0; i < 256; i++) { acc[i] = 24; @@ -33,7 +33,7 @@ int main() { }); q.submit([&](sycl::handler &cgh) { - auto acc = bufs[1].get_access(cgh); + auto acc = bufs[1].get_access(cgh); cgh.single_task([=]() { for (int i = 0; i < 256; i++) { acc[i] = 25; diff --git a/sycl/test-e2e/Sampler/normalized-clamp-linear-float.cpp b/sycl/test-e2e/Sampler/normalized-clamp-linear-float.cpp index ebd04c2b11946..b7ed3593c9c79 100644 --- a/sycl/test-e2e/Sampler/normalized-clamp-linear-float.cpp +++ b/sycl/test-e2e/Sampler/normalized-clamp-linear-float.cpp @@ -62,7 +62,7 @@ void test_normalized_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); @@ -80,8 +80,8 @@ void test_normalized_clamp_linear_sampler(image_channel_order ChanOrder, sampler(normalized, 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/normalized-clampedge-nearest.cpp b/sycl/test-e2e/Sampler/normalized-clampedge-nearest.cpp index 0a15fde4db2fb..575065d3636ad 100644 --- a/sycl/test-e2e/Sampler/normalized-clampedge-nearest.cpp +++ b/sycl/test-e2e/Sampler/normalized-clampedge-nearest.cpp @@ -51,7 +51,7 @@ void test_normalized_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); @@ -69,8 +69,8 @@ void test_normalized_clampedge_nearest_sampler(image_channel_order ChanOrder, sampler(normalized, 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/Sampler/normalized-none-linear-float.cpp b/sycl/test-e2e/Sampler/normalized-none-linear-float.cpp index 5348e08a3257c..3ea8728d83335 100644 --- a/sycl/test-e2e/Sampler/normalized-none-linear-float.cpp +++ b/sycl/test-e2e/Sampler/normalized-none-linear-float.cpp @@ -59,7 +59,7 @@ void test_normalized_none_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); @@ -76,8 +76,8 @@ void test_normalized_none_linear_sampler(image_channel_order ChanOrder, auto Norm_None_Linear_sampler = sampler(normalized, none, 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-repeat-nearest.cpp b/sycl/test-e2e/Sampler/normalized-repeat-nearest.cpp index ecd9a8f0b333e..486e2e9b1bff2 100644 --- a/sycl/test-e2e/Sampler/normalized-repeat-nearest.cpp +++ b/sycl/test-e2e/Sampler/normalized-repeat-nearest.cpp @@ -57,7 +57,7 @@ void test_normalized_repeat_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); @@ -74,8 +74,8 @@ void test_normalized_repeat_nearest_sampler(image_channel_order ChanOrder, auto Norm_Repeat_Nearest_sampler = sampler(normalized, repeat, 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-clampedge-linear-float.cpp b/sycl/test-e2e/Sampler/unnormalized-clampedge-linear-float.cpp index 63d3c060269ba..8e8d738320b85 100644 --- a/sycl/test-e2e/Sampler/unnormalized-clampedge-linear-float.cpp +++ b/sycl/test-e2e/Sampler/unnormalized-clampedge-linear-float.cpp @@ -55,7 +55,7 @@ void test_unnormalized_clampedge_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); @@ -73,8 +73,8 @@ void test_unnormalized_clampedge_linear_sampler(image_channel_order ChanOrder, sampler(unnormalized, clamp_edge, 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-none-nearest.cpp b/sycl/test-e2e/Sampler/unnormalized-none-nearest.cpp index 1efb39ce40259..85d22a4b7b582 100644 --- a/sycl/test-e2e/Sampler/unnormalized-none-nearest.cpp +++ b/sycl/test-e2e/Sampler/unnormalized-none-nearest.cpp @@ -47,7 +47,7 @@ void test_unnormalized_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); @@ -64,8 +64,8 @@ void test_unnormalized_none_nearest_sampler(image_channel_order ChanOrder, auto UnNorm_None_Nearest_sampler = sampler(unnormalized, 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/Scheduler/DataMovement.cpp b/sycl/test-e2e/Scheduler/DataMovement.cpp index c06f11fef7c59..fb4a1fc35fba9 100644 --- a/sycl/test-e2e/Scheduler/DataMovement.cpp +++ b/sycl/test-e2e/Scheduler/DataMovement.cpp @@ -18,7 +18,7 @@ #include "../helpers.hpp" -using sycl_access_mode = sycl::access::mode; +using sycl_access_mode = sycl::access_mode; template class CustomAllocator { public: diff --git a/sycl/test-e2e/Scheduler/InOrderQueueDeps.cpp b/sycl/test-e2e/Scheduler/InOrderQueueDeps.cpp index 6f5d4cdbe602e..fbecc4b7ad45b 100644 --- a/sycl/test-e2e/Scheduler/InOrderQueueDeps.cpp +++ b/sycl/test-e2e/Scheduler/InOrderQueueDeps.cpp @@ -16,7 +16,7 @@ class KernelA; void submitKernel(sycl::queue &Queue, sycl::buffer &Buf) { Queue.submit([&](sycl::handler &Cgh) { - auto BufAcc = Buf.get_access(Cgh); + auto BufAcc = Buf.get_access(Cgh); Cgh.single_task([=]() { (void)BufAcc[0]; }); }); } diff --git a/sycl/test-e2e/Scheduler/ReleaseResourcesTest.cpp b/sycl/test-e2e/Scheduler/ReleaseResourcesTest.cpp index 355f1a19e3c48..42323cf034069 100644 --- a/sycl/test-e2e/Scheduler/ReleaseResourcesTest.cpp +++ b/sycl/test-e2e/Scheduler/ReleaseResourcesTest.cpp @@ -14,7 +14,7 @@ #include "../helpers.hpp" -using sycl_access_mode = sycl::access::mode; +using sycl_access_mode = sycl::access_mode; int main() { bool Failed = false; diff --git a/sycl/test-e2e/SpecConstants/2020/kernel-bundle-api.cpp b/sycl/test-e2e/SpecConstants/2020/kernel-bundle-api.cpp index cb64134e2c03e..bea397c59aad0 100644 --- a/sycl/test-e2e/SpecConstants/2020/kernel-bundle-api.cpp +++ b/sycl/test-e2e/SpecConstants/2020/kernel-bundle-api.cpp @@ -95,10 +95,10 @@ bool test_default_values(sycl::queue q) { q.submit([&](sycl::handler &cgh) { cgh.use_kernel_bundle(exec_bundle); - auto int_acc = int_buffer.get_access(cgh); - auto float_acc = float_buffer.get_access(cgh); + auto int_acc = int_buffer.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(); float_acc[0] = kh.get_specialization_constant(); @@ -236,10 +236,10 @@ bool test_set_and_get_on_device(sycl::queue q) { q.submit([&](sycl::handler &cgh) { cgh.use_kernel_bundle(exec_bundle); - auto int_acc = int_buffer.get_access(cgh); - auto float_acc = float_buffer.get_access(cgh); + auto int_acc = int_buffer.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(); diff --git a/sycl/test-e2e/SpecConstants/2020/vector-convolution-demo.cpp b/sycl/test-e2e/SpecConstants/2020/vector-convolution-demo.cpp index e3b67f85cc1e5..9d2546560c8b3 100644 --- a/sycl/test-e2e/SpecConstants/2020/vector-convolution-demo.cpp +++ b/sycl/test-e2e/SpecConstants/2020/vector-convolution-demo.cpp @@ -85,8 +85,8 @@ void do_conv(buffer in, buffer out, CP coeff_provider) { queue myQueue; myQueue.submit([&](handler &cgh) { - auto in_acc = in.template get_access(cgh); - auto out_acc = out.template get_access(cgh); + auto in_acc = in.template get_access(cgh); + auto out_acc = out.template get_access(cgh); // Set the coefficient of the convolution as constant. // This will build a specific kernel the coefficient available as literals. diff --git a/sycl/test-e2e/SubGroup/common.cpp b/sycl/test-e2e/SubGroup/common.cpp index dfa5808b46ff2..0fb0df8fd4535 100644 --- a/sycl/test-e2e/SubGroup/common.cpp +++ b/sycl/test-e2e/SubGroup/common.cpp @@ -27,8 +27,8 @@ void check(queue &Queue, unsigned int G, unsigned int L) { buffer syclbuf(G); buffer sgsizebuf(1); Queue.submit([&](handler &cgh) { - auto sgsizeacc = sgsizebuf.get_access(cgh); - auto syclacc = syclbuf.get_access(cgh); + auto sgsizeacc = sgsizebuf.get_access(cgh); + auto syclacc = syclbuf.get_access(cgh); cgh.parallel_for(NdRange, [=](nd_item<1> NdItem) { sycl::sub_group SG = NdItem.get_sub_group(); syclacc[NdItem.get_global_id()].local_id = SG.get_local_id().get(0); diff --git a/sycl/test-e2e/SubGroup/reduce.hpp b/sycl/test-e2e/SubGroup/reduce.hpp index 514f07cd0f54a..c122618c9413d 100644 --- a/sycl/test-e2e/SubGroup/reduce.hpp +++ b/sycl/test-e2e/SubGroup/reduce.hpp @@ -20,8 +20,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) { sycl::sub_group sg = NdItem.get_sub_group(); diff --git a/sycl/test-e2e/SubGroupMask/Basic.cpp b/sycl/test-e2e/SubGroupMask/Basic.cpp index 040403d2f4afd..9de9a6d496b9d 100644 --- a/sycl/test-e2e/SubGroupMask/Basic.cpp +++ b/sycl/test-e2e/SubGroupMask/Basic.cpp @@ -36,7 +36,7 @@ int main() { buffer resbuf(&Res, range<1>(1)); Queue.submit([&](handler &cgh) { - auto resacc = resbuf.get_access(cgh); + auto resacc = resbuf.get_access(cgh); cgh.parallel_for( NdRange, [=](nd_item<1> NdItem) [[sycl::reqd_sub_group_size(32)]] { diff --git a/sycl/test-e2e/Tracing/image_printers.cpp b/sycl/test-e2e/Tracing/image_printers.cpp index 6d8a5be459e6a..87d68c0cb8892 100644 --- a/sycl/test-e2e/Tracing/image_printers.cpp +++ b/sycl/test-e2e/Tracing/image_printers.cpp @@ -21,7 +21,7 @@ int main() { const sycl::image_channel_order ChanOrder = sycl::image_channel_order::rgba; const sycl::image_channel_type ChanType = sycl::image_channel_type::fp32; - constexpr auto SYCLWrite = sycl::access::mode::write; + constexpr auto SYCLWrite = sycl::access_mode::write; const sycl::range<2> ImgSize(4, 4); diff --git a/sycl/test-e2e/USM/math.cpp b/sycl/test-e2e/USM/math.cpp index 5c7b9265e26b2..ea9d4ff2a6c30 100644 --- a/sycl/test-e2e/USM/math.cpp +++ b/sycl/test-e2e/USM/math.cpp @@ -78,7 +78,7 @@ int main() { int *BufI = (int *)s::malloc_shared( sizeof(int) * 2, myQueue.get_device(), myQueue.get_context()); myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task([=]() { AccR[0] = s::lgamma_r( float{10.f}, @@ -103,7 +103,7 @@ int main() { int *BufI = (int *)s::malloc_shared( sizeof(int) * 2, myQueue.get_device(), myQueue.get_context()); myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task([=]() { AccR[0] = s::lgamma_r( float{-2.4f}, @@ -128,7 +128,7 @@ int main() { s::int2 *BufI = (s::int2 *)s::malloc_shared( sizeof(s::int2) * 2, myQueue.get_device(), myQueue.get_context()); myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task([=]() { AccR[0] = s::lgamma_r( s::float2{10.f, -2.4f}, diff --git a/sycl/test-e2e/WeakObject/weak_object_utils.hpp b/sycl/test-e2e/WeakObject/weak_object_utils.hpp index aa6a9a280bfa4..3f222260e4a32 100644 --- a/sycl/test-e2e/WeakObject/weak_object_utils.hpp +++ b/sycl/test-e2e/WeakObject/weak_object_utils.hpp @@ -53,9 +53,9 @@ template