diff --git a/sycl/test-e2e/ESIMD/BitonicSortK.hpp b/sycl/test-e2e/ESIMD/BitonicSortK.hpp index 58cdd434b0d5f..25a28fceced1a 100644 --- a/sycl/test-e2e/ESIMD/BitonicSortK.hpp +++ b/sycl/test-e2e/ESIMD/BitonicSortK.hpp @@ -604,8 +604,8 @@ int BitonicSort::Solve(uint32_t *pInputs, uint32_t *pOutputs, uint32_t size) { buffer bufo(pOutputs, range<1>(size)); // enqueue sort265 kernel auto e = pQueue_->submit([&](handler &cgh) { - auto acci = bufi.get_access(cgh); - auto acco = bufo.get_access(cgh); + auto acci = bufi.get_access(cgh); + auto acco = bufo.get_access(cgh); cgh.parallel_for( SortGlobalRange * SortLocalRange, [=](id<1> i) SYCL_ESIMD_KERNEL { using namespace sycl::ext::intel::esimd; @@ -649,7 +649,7 @@ int BitonicSort::Solve(uint32_t *pInputs, uint32_t *pOutputs, uint32_t size) { for (int j = i; j >= 8; j--) { buffer buf(pOutputs, range<1>(size)); mergeEvent[k] = pQueue_->submit([&](handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.parallel_for( MergeGlobalRange * MergeLocalRange, [=](id<1> tid) SYCL_ESIMD_KERNEL { diff --git a/sycl/test-e2e/ESIMD/InlineAsm/asm_glb.cpp b/sycl/test-e2e/ESIMD/InlineAsm/asm_glb.cpp index ea401219d50f7..18c78ad95f6e1 100644 --- a/sycl/test-e2e/ESIMD/InlineAsm/asm_glb.cpp +++ b/sycl/test-e2e/ESIMD/InlineAsm/asm_glb.cpp @@ -48,9 +48,9 @@ int main(void) { << "\n"; auto e = q.submit([&](handler &cgh) { - auto PA = bufa.get_access(cgh); - auto PB = bufb.get_access(cgh); - auto PC = bufc.get_access(cgh); + auto PA = bufa.get_access(cgh); + auto PB = bufb.get_access(cgh); + auto PC = bufc.get_access(cgh); cgh.parallel_for( GlobalRange * LocalRange, [=](id<1> i) SYCL_ESIMD_KERNEL { using namespace sycl::ext::intel::esimd; diff --git a/sycl/test-e2e/ESIMD/InlineAsm/asm_simd_mask.cpp b/sycl/test-e2e/ESIMD/InlineAsm/asm_simd_mask.cpp index 8a93bc0aa6076..4603c99a0410b 100644 --- a/sycl/test-e2e/ESIMD/InlineAsm/asm_simd_mask.cpp +++ b/sycl/test-e2e/ESIMD/InlineAsm/asm_simd_mask.cpp @@ -44,9 +44,9 @@ int main(void) { std::cout << "Running on " << dev.get_info() << "\n"; auto e = q.submit([&](handler &cgh) { - auto PA = bufa.get_access(cgh); - auto PB = bufb.get_access(cgh); - auto PC = bufc.get_access(cgh); + auto PA = bufa.get_access(cgh); + auto PB = bufb.get_access(cgh); + auto PC = bufc.get_access(cgh); cgh.parallel_for( GlobalRange * LocalRange, [=](id<1> i) SYCL_ESIMD_KERNEL { using namespace sycl::ext::intel::esimd; diff --git a/sycl/test-e2e/ESIMD/InlineAsm/asm_simd_view.cpp b/sycl/test-e2e/ESIMD/InlineAsm/asm_simd_view.cpp index fcbed19e4ab5f..df2a3ccc322d4 100644 --- a/sycl/test-e2e/ESIMD/InlineAsm/asm_simd_view.cpp +++ b/sycl/test-e2e/ESIMD/InlineAsm/asm_simd_view.cpp @@ -44,9 +44,9 @@ int main(void) { std::cout << "Running on " << dev.get_info() << "\n"; auto e = q.submit([&](handler &cgh) { - auto PA = bufa.get_access(cgh); - auto PB = bufb.get_access(cgh); - auto PC = bufc.get_access(cgh); + auto PA = bufa.get_access(cgh); + auto PB = bufb.get_access(cgh); + auto PC = bufc.get_access(cgh); cgh.parallel_for( GlobalRange * LocalRange, [=](id<1> i) SYCL_ESIMD_KERNEL { using namespace sycl::ext::intel::esimd; diff --git a/sycl/test-e2e/ESIMD/InlineAsm/asm_vadd.cpp b/sycl/test-e2e/ESIMD/InlineAsm/asm_vadd.cpp index e80b0c1c7592d..caa131d676991 100644 --- a/sycl/test-e2e/ESIMD/InlineAsm/asm_vadd.cpp +++ b/sycl/test-e2e/ESIMD/InlineAsm/asm_vadd.cpp @@ -44,9 +44,9 @@ int main(void) { std::cout << "Running on " << dev.get_info() << "\n"; auto e = q.submit([&](handler &cgh) { - auto PA = bufa.get_access(cgh); - auto PB = bufb.get_access(cgh); - auto PC = bufc.get_access(cgh); + auto PA = bufa.get_access(cgh); + auto PB = bufb.get_access(cgh); + auto PC = bufc.get_access(cgh); cgh.parallel_for( GlobalRange * LocalRange, [=](id<1> i) SYCL_ESIMD_KERNEL { using namespace sycl::ext::intel::esimd; diff --git a/sycl/test-e2e/ESIMD/acc_gather_scatter_rgba.hpp b/sycl/test-e2e/ESIMD/acc_gather_scatter_rgba.hpp index fdce70bc61512..95b905a4c2bc0 100644 --- a/sycl/test-e2e/ESIMD/acc_gather_scatter_rgba.hpp +++ b/sycl/test-e2e/ESIMD/acc_gather_scatter_rgba.hpp @@ -109,8 +109,8 @@ bool test(queue q) { buffer OutBuf(B, range<1>(size)); range<1> glob_range{numWorkItems}; auto e = q.submit([&](handler &cgh) { - auto InAcc = InBuf.template get_access(cgh); - auto OutAcc = OutBuf.template get_access(cgh); + auto InAcc = InBuf.template get_access(cgh); + auto OutAcc = OutBuf.template get_access(cgh); Kernel kernel(InAcc, OutAcc); cgh.parallel_for(glob_range, kernel); }); diff --git a/sycl/test-e2e/ESIMD/acc_gather_scatter_rgba_stateless_64.cpp b/sycl/test-e2e/ESIMD/acc_gather_scatter_rgba_stateless_64.cpp index 049a762499add..375487b5cfb98 100644 --- a/sycl/test-e2e/ESIMD/acc_gather_scatter_rgba_stateless_64.cpp +++ b/sycl/test-e2e/ESIMD/acc_gather_scatter_rgba_stateless_64.cpp @@ -39,7 +39,7 @@ int main(void) { << "\n"; try { q.submit([&](handler &cgh) { - auto PA = bufa.get_access(cgh); + auto PA = bufa.get_access(cgh); cgh.single_task([=]() SYCL_ESIMD_KERNEL { uint64_t offsetStart = (Size - VL) * sizeof(uint64_t); simd offset(offsetStart, sizeof(uint64_t)); diff --git a/sycl/test-e2e/ESIMD/accessor.hpp b/sycl/test-e2e/ESIMD/accessor.hpp index 0456a5e1b5f49..e18aeaf87e60a 100644 --- a/sycl/test-e2e/ESIMD/accessor.hpp +++ b/sycl/test-e2e/ESIMD/accessor.hpp @@ -36,8 +36,8 @@ int main() { std::cout << "Running on " << q.get_device().get_info() << "\n"; - auto acc0 = buf0.get_access(cgh); - auto acc1 = buf1.get_access(cgh); + auto acc0 = buf0.get_access(cgh); + auto acc1 = buf1.get_access(cgh); cgh.parallel_for(range<1>(1), [=](sycl::id<1> i) SYCL_ESIMD_KERNEL { diff --git a/sycl/test-e2e/ESIMD/accessor_gather_scatter.hpp b/sycl/test-e2e/ESIMD/accessor_gather_scatter.hpp index 9fc706b4fa133..de37569242357 100644 --- a/sycl/test-e2e/ESIMD/accessor_gather_scatter.hpp +++ b/sycl/test-e2e/ESIMD/accessor_gather_scatter.hpp @@ -57,7 +57,7 @@ template bool test(queue q) { range<1> glob_range{size / VL}; auto e = q.submit([&](handler &cgh) { - auto acc = buf.template get_access(cgh); + auto acc = buf.template get_access(cgh); Kernel kernel(acc); cgh.parallel_for(glob_range, kernel); }); diff --git a/sycl/test-e2e/ESIMD/accessor_gather_scatter_stateless_64.cpp b/sycl/test-e2e/ESIMD/accessor_gather_scatter_stateless_64.cpp index e5037b42c934f..bc87ccd801946 100644 --- a/sycl/test-e2e/ESIMD/accessor_gather_scatter_stateless_64.cpp +++ b/sycl/test-e2e/ESIMD/accessor_gather_scatter_stateless_64.cpp @@ -34,7 +34,7 @@ int main(void) { << "\n"; try { q.submit([&](handler &cgh) { - auto PA = bufa.get_access(cgh); + auto PA = bufa.get_access(cgh); cgh.single_task([=]() SYCL_ESIMD_KERNEL { uint64_t offsetStart = (Size - VL) * sizeof(uint64_t); simd offset(offsetStart, sizeof(uint64_t)); diff --git a/sycl/test-e2e/ESIMD/accessor_load_store.hpp b/sycl/test-e2e/ESIMD/accessor_load_store.hpp index a345cbdbc29e3..aae2e397273f9 100644 --- a/sycl/test-e2e/ESIMD/accessor_load_store.hpp +++ b/sycl/test-e2e/ESIMD/accessor_load_store.hpp @@ -50,7 +50,7 @@ template bool test(queue q, size_t size) { range<1> glob_range{size}; auto e = q.submit([&](handler &cgh) { - auto acc = buf.template get_access(cgh); + auto acc = buf.template get_access(cgh); Kernel kernel(acc); cgh.parallel_for(glob_range, kernel); }); diff --git a/sycl/test-e2e/ESIMD/accessor_load_store_stateless_64.cpp b/sycl/test-e2e/ESIMD/accessor_load_store_stateless_64.cpp index 5ba14f718f359..4a76b8c6f7342 100644 --- a/sycl/test-e2e/ESIMD/accessor_load_store_stateless_64.cpp +++ b/sycl/test-e2e/ESIMD/accessor_load_store_stateless_64.cpp @@ -34,7 +34,7 @@ int main(void) { << "\n"; try { q.submit([&](handler &cgh) { - auto PA = bufa.get_access(cgh); + auto PA = bufa.get_access(cgh); cgh.single_task([=]() SYCL_ESIMD_KERNEL { uint64_t offset = (Size - VL) * sizeof(uint64_t); simd va = block_load(PA, 0); diff --git a/sycl/test-e2e/ESIMD/accessor_stateless_64.cpp b/sycl/test-e2e/ESIMD/accessor_stateless_64.cpp index 1345d5bdb72ee..523487f7890e9 100644 --- a/sycl/test-e2e/ESIMD/accessor_stateless_64.cpp +++ b/sycl/test-e2e/ESIMD/accessor_stateless_64.cpp @@ -36,7 +36,7 @@ int main(void) { << "\n"; q.submit([&](handler &cgh) { - auto PA = bufa.get_access(cgh); + auto PA = bufa.get_access(cgh); cgh.single_task([=]() SYCL_ESIMD_KERNEL { uint64_t offset = (Size - VL) * sizeof(uint64_t); simd va; diff --git a/sycl/test-e2e/ESIMD/accessor_stateless_ctor_64.cpp b/sycl/test-e2e/ESIMD/accessor_stateless_ctor_64.cpp index 5598318e9ae9e..c84602a32fee2 100644 --- a/sycl/test-e2e/ESIMD/accessor_stateless_ctor_64.cpp +++ b/sycl/test-e2e/ESIMD/accessor_stateless_ctor_64.cpp @@ -36,7 +36,7 @@ int main(void) { << "\n"; q.submit([&](handler &cgh) { - auto PA = bufa.get_access(cgh); + auto PA = bufa.get_access(cgh); cgh.single_task([=]() SYCL_ESIMD_KERNEL { uint64_t offset = (Size - VL) * sizeof(uint64_t); simd va(PA, 0); diff --git a/sycl/test-e2e/ESIMD/aot_mixed.cpp b/sycl/test-e2e/ESIMD/aot_mixed.cpp index 9a4294354b214..0b678400af5cd 100644 --- a/sycl/test-e2e/ESIMD/aot_mixed.cpp +++ b/sycl/test-e2e/ESIMD/aot_mixed.cpp @@ -70,9 +70,9 @@ bool test_esimd(queue q) { buffer bufc(C, range<1>(Size)); auto e = q.submit([&](handler &cgh) { - auto PA = bufa.get_access(cgh); - auto PB = bufb.get_access(cgh); - auto PC = bufc.get_access(cgh); + auto PA = bufa.get_access(cgh); + auto PB = bufb.get_access(cgh); + auto PC = bufc.get_access(cgh); cgh.parallel_for( Size / VL, [=](id<1> i) SYCL_ESIMD_KERNEL { using namespace sycl::ext::intel::esimd; @@ -122,9 +122,9 @@ bool test_sycl(queue q) { buffer bufc(C, range<1>(Size)); auto e = q.submit([&](handler &cgh) { - auto PA = bufa.get_access(cgh); - auto PB = bufb.get_access(cgh); - auto PC = bufc.get_access(cgh); + auto PA = bufa.get_access(cgh); + auto PB = bufb.get_access(cgh); + auto PC = bufc.get_access(cgh); cgh.parallel_for(Size, [=](id<1> i) { PC[i] = PA[i] + PB[i]; }); }); diff --git a/sycl/test-e2e/ESIMD/api/ballot.cpp b/sycl/test-e2e/ESIMD/api/ballot.cpp index 917a2d3052dda..7ec62ebc3cea9 100644 --- a/sycl/test-e2e/ESIMD/api/ballot.cpp +++ b/sycl/test-e2e/ESIMD/api/ballot.cpp @@ -38,8 +38,8 @@ template bool test(queue &Q) { buffer RB(&Res, range<1>(1)); auto E = Q.submit([&](handler &CGH) { - auto In = PB.template get_access(CGH); - auto Out = RB.template get_access(CGH); + auto In = PB.template get_access(CGH); + auto Out = RB.template get_access(CGH); CGH.parallel_for(sycl::range<1>{1}, [=](id<1>) SYCL_ESIMD_KERNEL { simd Mask; diff --git a/sycl/test-e2e/ESIMD/api/functional/operators/operator_assignment_glb.cpp b/sycl/test-e2e/ESIMD/api/functional/operators/operator_assignment_glb.cpp index 11a3a02aede09..383896272a10e 100644 --- a/sycl/test-e2e/ESIMD/api/functional/operators/operator_assignment_glb.cpp +++ b/sycl/test-e2e/ESIMD/api/functional/operators/operator_assignment_glb.cpp @@ -50,9 +50,9 @@ int main(void) { << "\n"; auto e = q.submit([&](handler &cgh) { - auto PA = bufa.get_access(cgh); - auto PB = bufb.get_access(cgh); - auto PC = bufc.get_access(cgh); + auto PA = bufa.get_access(cgh); + auto PB = bufb.get_access(cgh); + auto PC = bufc.get_access(cgh); cgh.parallel_for( GlobalRange * LocalRange, [=](id<1> i) SYCL_ESIMD_KERNEL { using namespace sycl::ext::intel::esimd; diff --git a/sycl/test-e2e/ESIMD/api/functional/operators/operator_assignment_glb_mask.cpp b/sycl/test-e2e/ESIMD/api/functional/operators/operator_assignment_glb_mask.cpp index 4537ddd210cb9..e2426bb5ee9b9 100644 --- a/sycl/test-e2e/ESIMD/api/functional/operators/operator_assignment_glb_mask.cpp +++ b/sycl/test-e2e/ESIMD/api/functional/operators/operator_assignment_glb_mask.cpp @@ -48,8 +48,8 @@ int main(void) { << "\n"; auto e = q.submit([&](handler &cgh) { - auto PA = bufa.get_access(cgh); - auto PB = bufb.get_access(cgh); + auto PA = bufa.get_access(cgh); + auto PB = bufb.get_access(cgh); cgh.parallel_for( GlobalRange * LocalRange, [=](id<1> i) SYCL_ESIMD_KERNEL { using namespace sycl::ext::intel::esimd; diff --git a/sycl/test-e2e/ESIMD/api/replicate_smoke.cpp b/sycl/test-e2e/ESIMD/api/replicate_smoke.cpp index 6ae830f3a940d..a897181729d03 100644 --- a/sycl/test-e2e/ESIMD/api/replicate_smoke.cpp +++ b/sycl/test-e2e/ESIMD/api/replicate_smoke.cpp @@ -98,8 +98,8 @@ bool test_impl(queue q, int offset, T (&&gold)[N]) { sycl::buffer dst_buf(dm.dst, VL); q.submit([&](handler &cgh) { - auto src_acc = src_buf.template get_access(cgh); - auto dst_acc = dst_buf.template get_access(cgh); + auto src_acc = src_buf.template get_access(cgh); + auto dst_acc = dst_buf.template get_access(cgh); cgh.single_task([=]() SYCL_ESIMD_KERNEL { simd src(src_acc, 0); diff --git a/sycl/test-e2e/ESIMD/api/saturation_smoke.cpp b/sycl/test-e2e/ESIMD/api/saturation_smoke.cpp index ca1ba1ecc58ef..5c1f6373ca857 100644 --- a/sycl/test-e2e/ESIMD/api/saturation_smoke.cpp +++ b/sycl/test-e2e/ESIMD/api/saturation_smoke.cpp @@ -83,8 +83,8 @@ bool test(queue q) { sycl::buffer dst_buf(dm.dst, N); auto e = q.submit([&](handler &cgh) { - auto src_acc = src_buf.template get_access(cgh); - auto dst_acc = dst_buf.template get_access(cgh); + auto src_acc = src_buf.template get_access(cgh); + auto dst_acc = dst_buf.template get_access(cgh); cgh.single_task([=]() SYCL_ESIMD_KERNEL { simd x(src_acc, 0); diff --git a/sycl/test-e2e/ESIMD/api/simd_binop_integer_promotion.cpp b/sycl/test-e2e/ESIMD/api/simd_binop_integer_promotion.cpp index 861e8a78155bd..ce12d1d7b4e68 100644 --- a/sycl/test-e2e/ESIMD/api/simd_binop_integer_promotion.cpp +++ b/sycl/test-e2e/ESIMD/api/simd_binop_integer_promotion.cpp @@ -52,9 +52,9 @@ template bool test(queue q) { range<1> glob_range{1}; auto e = q.submit([&](handler &cgh) { - auto PA = bufA.template get_access(cgh); - auto PB = bufB.template get_access(cgh); - auto PC = bufC.template get_access(cgh); + auto PA = bufA.template get_access(cgh); + auto PB = bufB.template get_access(cgh); + auto PC = bufC.template get_access(cgh); cgh.parallel_for>( glob_range, [=](id<1> i) SYCL_ESIMD_KERNEL { using namespace sycl::ext::intel::experimental::esimd; diff --git a/sycl/test-e2e/ESIMD/api/simd_copy_to_from.cpp b/sycl/test-e2e/ESIMD/api/simd_copy_to_from.cpp index 297d94a9410a5..f381e86429557 100644 --- a/sycl/test-e2e/ESIMD/api/simd_copy_to_from.cpp +++ b/sycl/test-e2e/ESIMD/api/simd_copy_to_from.cpp @@ -84,8 +84,8 @@ bool testAcc(queue &Q, T *Src, T *Dst, unsigned Off, Flags) { buffer DstB(Dst, range<1>(Off + N)); Q.submit([&](handler &CGH) { - auto SrcA = SrcB.template get_access(CGH); - auto DstA = DstB.template get_access(CGH); + auto SrcA = SrcB.template get_access(CGH); + auto DstA = DstB.template get_access(CGH); CGH.parallel_for(sycl::range<1>{1}, [=](id<1>) SYCL_ESIMD_KERNEL { simd Vals; diff --git a/sycl/test-e2e/ESIMD/api/simd_copy_to_from_stateful.cpp b/sycl/test-e2e/ESIMD/api/simd_copy_to_from_stateful.cpp index db5a547178355..6e49c300afd24 100644 --- a/sycl/test-e2e/ESIMD/api/simd_copy_to_from_stateful.cpp +++ b/sycl/test-e2e/ESIMD/api/simd_copy_to_from_stateful.cpp @@ -89,8 +89,8 @@ bool testAcc(queue &Q, T *Src, T *Dst, unsigned Off, Flags) { buffer DstB(Dst, range<1>(Off + N)); Q.submit([&](handler &CGH) { - auto SrcA = SrcB.template get_access(CGH); - auto DstA = DstB.template get_access(CGH); + auto SrcA = SrcB.template get_access(CGH); + auto DstA = DstB.template get_access(CGH); CGH.parallel_for(sycl::range<1>{1}, [=](id<1>) SYCL_ESIMD_KERNEL { simd Vals; diff --git a/sycl/test-e2e/ESIMD/api/simd_memory_access.cpp b/sycl/test-e2e/ESIMD/api/simd_memory_access.cpp index b07e042309422..e4f6ea6aa062e 100644 --- a/sycl/test-e2e/ESIMD/api/simd_memory_access.cpp +++ b/sycl/test-e2e/ESIMD/api/simd_memory_access.cpp @@ -87,7 +87,7 @@ template bool test(queue q, size_t size) { range<1> glob_range{size / N}; auto e = q.submit([&](handler &cgh) { - auto acc = buf.template get_access(cgh); + auto acc = buf.template get_access(cgh); Kernel kernel(acc); cgh.parallel_for(glob_range, kernel); }); diff --git a/sycl/test-e2e/ESIMD/api/simd_negation_operator.cpp b/sycl/test-e2e/ESIMD/api/simd_negation_operator.cpp index 685ba01f4346d..f2c11be215625 100644 --- a/sycl/test-e2e/ESIMD/api/simd_negation_operator.cpp +++ b/sycl/test-e2e/ESIMD/api/simd_negation_operator.cpp @@ -37,8 +37,8 @@ bool test(queue q, const std::array &input, range<1> glob_range{1}; auto e = q.submit([&](handler &cgh) { - auto PA = bufA.template get_access(cgh); - auto PB = bufB.template get_access(cgh); + auto PA = bufA.template get_access(cgh); + auto PB = bufB.template get_access(cgh); cgh.parallel_for(glob_range, [=](id<1> i) SYCL_ESIMD_KERNEL { using namespace sycl::ext::intel::esimd; simd va; diff --git a/sycl/test-e2e/ESIMD/api/simd_subscript_operator.cpp b/sycl/test-e2e/ESIMD/api/simd_subscript_operator.cpp index 86f6b7fe998bd..1dcefb430a1d2 100644 --- a/sycl/test-e2e/ESIMD/api/simd_subscript_operator.cpp +++ b/sycl/test-e2e/ESIMD/api/simd_subscript_operator.cpp @@ -45,8 +45,8 @@ template bool test(queue &q) { range<1> glob_range{1}; auto e = q.submit([&](handler &cgh) { - auto PA = bufA.template get_access(cgh); - auto PB = bufB.template get_access(cgh); + auto PA = bufA.template get_access(cgh); + auto PB = bufB.template get_access(cgh); cgh.parallel_for(glob_range, [=](id<1> i) SYCL_ESIMD_KERNEL { using namespace sycl::ext::intel::esimd; unsigned int offset = i * VL * sizeof(T); diff --git a/sycl/test-e2e/ESIMD/api/simd_view_copy_move_assign.cpp b/sycl/test-e2e/ESIMD/api/simd_view_copy_move_assign.cpp index 9619a2937d406..b2fe194789028 100644 --- a/sycl/test-e2e/ESIMD/api/simd_view_copy_move_assign.cpp +++ b/sycl/test-e2e/ESIMD/api/simd_view_copy_move_assign.cpp @@ -38,8 +38,8 @@ bool test(queue q, std::string str, F funcUnderTest) { buffer BufB(B, range<1>(Size)); q.submit([&](handler &cgh) { - auto PA = BufA.template get_access(cgh); - auto PB = BufB.template get_access(cgh); + auto PA = BufA.template get_access(cgh); + auto PB = BufB.template get_access(cgh); cgh.parallel_for(range<1>{Size / VL}, [=](id<1> i) SYCL_ESIMD_KERNEL { using namespace sycl::ext::intel::esimd; unsigned int offset = i * VL * sizeof(T); diff --git a/sycl/test-e2e/ESIMD/api/simd_view_negation_operator.cpp b/sycl/test-e2e/ESIMD/api/simd_view_negation_operator.cpp index eeacce508b783..a3d9077ff6d11 100644 --- a/sycl/test-e2e/ESIMD/api/simd_view_negation_operator.cpp +++ b/sycl/test-e2e/ESIMD/api/simd_view_negation_operator.cpp @@ -37,8 +37,8 @@ bool test(queue q, const std::array &input, range<1> glob_range{1}; auto e = q.submit([&](handler &cgh) { - auto PA = bufA.template get_access(cgh); - auto PB = bufB.template get_access(cgh); + auto PA = bufA.template get_access(cgh); + auto PB = bufB.template get_access(cgh); cgh.parallel_for(glob_range, [=](id<1> i) SYCL_ESIMD_KERNEL { using namespace sycl::ext::intel::esimd; simd va; diff --git a/sycl/test-e2e/ESIMD/api/simd_view_subscript_operator.cpp b/sycl/test-e2e/ESIMD/api/simd_view_subscript_operator.cpp index 6a8732d8a001c..188722b3a4659 100644 --- a/sycl/test-e2e/ESIMD/api/simd_view_subscript_operator.cpp +++ b/sycl/test-e2e/ESIMD/api/simd_view_subscript_operator.cpp @@ -48,8 +48,8 @@ template bool test(queue &q) { range<1> glob_range{1}; auto e = q.submit([&](handler &cgh) { - auto PA = bufA.template get_access(cgh); - auto PB = bufB.template get_access(cgh); + auto PA = bufA.template get_access(cgh); + auto PB = bufB.template get_access(cgh); cgh.parallel_for>(glob_range, [=](id<1> i) SYCL_ESIMD_KERNEL { using namespace sycl::ext::intel::esimd; simd va; diff --git a/sycl/test-e2e/ESIMD/api/slm_gather_scatter_heavy.cpp b/sycl/test-e2e/ESIMD/api/slm_gather_scatter_heavy.cpp index 59b498027d6b8..0c54c571ec17e 100644 --- a/sycl/test-e2e/ESIMD/api/slm_gather_scatter_heavy.cpp +++ b/sycl/test-e2e/ESIMD/api/slm_gather_scatter_heavy.cpp @@ -382,8 +382,8 @@ bool test_impl(queue q) { range<1> glob_range{size / VL}; auto e = q.submit([&](handler &cgh) { - auto acc_in = buf_in.template get_access(cgh); - auto acc_out = buf_out.template get_access(cgh); + auto acc_in = buf_in.template get_access(cgh); + auto acc_out = buf_out.template get_access(cgh); constexpr auto WG_SIZE = STRIDE; // for simplicity of the test KernelType kernel(acc_in, acc_out, MaskedLane); cgh.parallel_for(nd_range<1>{glob_range, range<1>(WG_SIZE)}, kernel); diff --git a/sycl/test-e2e/ESIMD/clz_ctz.cpp b/sycl/test-e2e/ESIMD/clz_ctz.cpp index 3e8352072ee52..694bed582ee55 100644 --- a/sycl/test-e2e/ESIMD/clz_ctz.cpp +++ b/sycl/test-e2e/ESIMD/clz_ctz.cpp @@ -43,8 +43,8 @@ template bool test(queue &q) { queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler()); auto e = q.submit([&](handler &cgh) { - auto PA = bufa.template get_access(cgh); - auto PB = bufb.template get_access(cgh); + auto PA = bufa.template get_access(cgh); + auto PB = bufb.template get_access(cgh); cgh.parallel_for(GlobalRange * LocalRange, [=](id<1> i) SYCL_ESIMD_KERNEL { using namespace sycl::ext::intel::esimd; diff --git a/sycl/test-e2e/ESIMD/ext_math.cpp b/sycl/test-e2e/ESIMD/ext_math.cpp index 0261047b0213a..c7938fd22b019 100644 --- a/sycl/test-e2e/ESIMD/ext_math.cpp +++ b/sycl/test-e2e/ESIMD/ext_math.cpp @@ -346,10 +346,10 @@ bool test(queue &Q, const std::string &Name, InitF Init = InitNarrow{}, sycl::range<1> LocalRange{1}; auto E = Q.submit([&](handler &CGH) { - auto PA = BufA.template get_access(CGH); - auto PC = BufC.template get_access(CGH); + auto PA = BufA.template get_access(CGH); + auto PC = BufC.template get_access(CGH); if constexpr (IsBinOp) { - auto PB = BufB.template get_access(CGH); + auto PB = BufB.template get_access(CGH); BinaryDeviceFunc F(PA, PB, PC); CGH.parallel_for(nd_range<1>{GlobalRange, LocalRange}, F); diff --git a/sycl/test-e2e/ESIMD/fma.cpp b/sycl/test-e2e/ESIMD/fma.cpp index 6133b1059fd59..e74e12bd4e853 100644 --- a/sycl/test-e2e/ESIMD/fma.cpp +++ b/sycl/test-e2e/ESIMD/fma.cpp @@ -36,9 +36,9 @@ template bool test(queue q) { buffer bufc(C, range<1>(Size)); auto e = q.submit([&](handler &cgh) { - auto PA = bufa.template get_access(cgh); - auto PB = bufb.template get_access(cgh); - auto PC = bufc.template get_access(cgh); + auto PA = bufa.template get_access(cgh); + auto PB = bufb.template get_access(cgh); + auto PC = bufc.template get_access(cgh); cgh.single_task([=]() SYCL_ESIMD_KERNEL { using namespace sycl::ext::intel::esimd; unsigned int offset = 0; diff --git a/sycl/test-e2e/ESIMD/fp_call_from_func.cpp b/sycl/test-e2e/ESIMD/fp_call_from_func.cpp index 110079c725772..2be2511b4077e 100644 --- a/sycl/test-e2e/ESIMD/fp_call_from_func.cpp +++ b/sycl/test-e2e/ESIMD/fp_call_from_func.cpp @@ -43,7 +43,7 @@ int main(int argc, char **argv) { buffer buf(output, range<1>(1)); q.submit([&](handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.parallel_for( sycl::range<1>{1}, diff --git a/sycl/test-e2e/ESIMD/fp_in_phi.cpp b/sycl/test-e2e/ESIMD/fp_in_phi.cpp index c8763f23598aa..d0991044fb6ae 100644 --- a/sycl/test-e2e/ESIMD/fp_in_phi.cpp +++ b/sycl/test-e2e/ESIMD/fp_in_phi.cpp @@ -38,8 +38,8 @@ bool test(queue q, bool flag) { buffer y_buf(Y.data(), Y.size()); q.submit([&](handler &cgh) { - auto o_acc = o_buf.get_access(cgh); - auto y_acc = y_buf.get_access(cgh); + auto o_acc = o_buf.get_access(cgh); + auto y_acc = y_buf.get_access(cgh); cgh.parallel_for(sycl::range<1>{1}, [=](id<1> i) SYCL_ESIMD_KERNEL { diff --git a/sycl/test-e2e/ESIMD/fp_in_select.cpp b/sycl/test-e2e/ESIMD/fp_in_select.cpp index 83007653a556a..a0edaeb8b971f 100644 --- a/sycl/test-e2e/ESIMD/fp_in_select.cpp +++ b/sycl/test-e2e/ESIMD/fp_in_select.cpp @@ -31,7 +31,7 @@ bool test(queue q, bool flag) { buffer buf(output, range<1>(1)); q.submit([&](handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.parallel_for(sycl::range<1>{1}, [=](id<1> i) SYCL_ESIMD_KERNEL { diff --git a/sycl/test-e2e/ESIMD/frem.cpp b/sycl/test-e2e/ESIMD/frem.cpp index 7b5e3c50ea87c..89f2b0ea6b8f1 100644 --- a/sycl/test-e2e/ESIMD/frem.cpp +++ b/sycl/test-e2e/ESIMD/frem.cpp @@ -38,9 +38,9 @@ template bool test(queue q) { buffer bufc(C, range<1>(Size)); auto e = q.submit([&](handler &cgh) { - auto PA = bufa.template get_access(cgh); - auto PB = bufb.template get_access(cgh); - auto PC = bufc.template get_access(cgh); + auto PA = bufa.template get_access(cgh); + auto PB = bufb.template get_access(cgh); + auto PC = bufc.template get_access(cgh); cgh.single_task([=]() SYCL_ESIMD_KERNEL { using namespace sycl::ext::intel::esimd; unsigned int offset = 0; diff --git a/sycl/test-e2e/ESIMD/grf.cpp b/sycl/test-e2e/ESIMD/grf.cpp index fd73ac9056ad7..75b98c3e3d6a4 100644 --- a/sycl/test-e2e/ESIMD/grf.cpp +++ b/sycl/test-e2e/ESIMD/grf.cpp @@ -66,7 +66,7 @@ int main(void) { try { buffer bufa(A.data(), range<1>(Size)); auto e = q.submit([&](handler &cgh) { - auto PA = bufa.get_access(cgh); + auto PA = bufa.get_access(cgh); cgh.parallel_for(Size, [=](id<1> i) { PA[i] = PA[i] + 1; }); }); @@ -86,7 +86,7 @@ int main(void) { try { buffer bufa(A.data(), range<1>(Size)); auto e = q.submit([&](handler &cgh) { - auto PA = bufa.get_access(cgh); + auto PA = bufa.get_access(cgh); cgh.parallel_for(Size, [=](id<1> i) SYCL_ESIMD_KERNEL { unsigned int offset = i * VL * sizeof(float); simd va; @@ -116,7 +116,7 @@ int main(void) { sycl::ext::oneapi::experimental::properties prop{grf_size<256>}; #endif auto e = q.submit([&](handler &cgh) { - auto PA = bufa.get_access(cgh); + auto PA = bufa.get_access(cgh); cgh.parallel_for( Size, prop, [=](id<1> i) SYCL_ESIMD_KERNEL { unsigned int offset = i * VL * sizeof(float); diff --git a/sycl/test-e2e/ESIMD/histogram.cpp b/sycl/test-e2e/ESIMD/histogram.cpp index a77a666eefd74..f9605248a9aec 100644 --- a/sycl/test-e2e/ESIMD/histogram.cpp +++ b/sycl/test-e2e/ESIMD/histogram.cpp @@ -155,7 +155,7 @@ int main(int argc, char *argv[]) { nd_range<1> Range(GlobalRange, LocalRange); auto e = q.submit([&](handler &cgh) { - auto readAcc = Img.get_access(cgh); + auto readAcc = Img.get_access(cgh); cgh.parallel_for( Range, [=](nd_item<1> ndi) SYCL_ESIMD_KERNEL { diff --git a/sycl/test-e2e/ESIMD/histogram_2d.cpp b/sycl/test-e2e/ESIMD/histogram_2d.cpp index a7971b94e28cd..4728f383ac901 100644 --- a/sycl/test-e2e/ESIMD/histogram_2d.cpp +++ b/sycl/test-e2e/ESIMD/histogram_2d.cpp @@ -157,7 +157,7 @@ int main(int argc, char *argv[]) { nd_range<2> Range(GlobalRange, LocalRange); auto e = q.submit([&](handler &cgh) { - auto readAcc = Img.get_access(cgh); + auto readAcc = Img.get_access(cgh); cgh.parallel_for( Range, [=](nd_item<2> ndi) SYCL_ESIMD_KERNEL { diff --git a/sycl/test-e2e/ESIMD/linear/linear.cpp b/sycl/test-e2e/ESIMD/linear/linear.cpp index e395ec88ea01d..d6fb2fdd5d73e 100644 --- a/sycl/test-e2e/ESIMD/linear/linear.cpp +++ b/sycl/test-e2e/ESIMD/linear/linear.cpp @@ -86,8 +86,8 @@ int main(int argc, char *argv[]) { for (int iter = 0; iter <= num_iters; ++iter) { auto e = q.submit([&](handler &cgh) { - auto accInput = imgInput.get_access(cgh); - auto accOutput = imgOutput.get_access(cgh); + auto accInput = imgInput.get_access(cgh); + auto accOutput = imgOutput.get_access(cgh); cgh.parallel_for( GlobalRange * LocalRange, [=](item<2> it) SYCL_ESIMD_KERNEL { diff --git a/sycl/test-e2e/ESIMD/local_accessor_copy_to_from.cpp b/sycl/test-e2e/ESIMD/local_accessor_copy_to_from.cpp index c88f5d7017b98..2395de92443da 100644 --- a/sycl/test-e2e/ESIMD/local_accessor_copy_to_from.cpp +++ b/sycl/test-e2e/ESIMD/local_accessor_copy_to_from.cpp @@ -48,7 +48,7 @@ template bool test(queue q) { nd_range<1> NDRange{range<1>{GlobalRange}, range<1>{LocalRange}}; q.submit([&](handler &CGH) { auto LocalAcc = local_accessor(Size, CGH); - auto Acc = buf.template get_access(CGH); + auto Acc = buf.template get_access(CGH); CGH.parallel_for(NDRange, [=](nd_item<1> Item) SYCL_ESIMD_KERNEL { uint32_t GID = Item.get_global_id(0); uint32_t LID = Item.get_local_id(0); diff --git a/sycl/test-e2e/ESIMD/local_accessor_gather_scatter.cpp b/sycl/test-e2e/ESIMD/local_accessor_gather_scatter.cpp index 1484d97f1fe41..3d49ae7973ef6 100644 --- a/sycl/test-e2e/ESIMD/local_accessor_gather_scatter.cpp +++ b/sycl/test-e2e/ESIMD/local_accessor_gather_scatter.cpp @@ -45,7 +45,7 @@ template bool test(queue q) { nd_range<1> glob_range{range<1>{1}, range<1>{1}}; q.submit([&](handler &cgh) { - auto acc = buf.template get_access(cgh); + auto acc = buf.template get_access(cgh); auto LocalAcc = local_accessor(size * STRIDE, cgh); cgh.parallel_for(glob_range, [=](nd_item<1> ndi) SYCL_ESIMD_KERNEL { using namespace sycl::ext::intel::esimd; diff --git a/sycl/test-e2e/ESIMD/local_accessor_gather_scatter_rgba.cpp b/sycl/test-e2e/ESIMD/local_accessor_gather_scatter_rgba.cpp index 4498927793fd7..1a558b726b72a 100644 --- a/sycl/test-e2e/ESIMD/local_accessor_gather_scatter_rgba.cpp +++ b/sycl/test-e2e/ESIMD/local_accessor_gather_scatter_rgba.cpp @@ -83,7 +83,7 @@ template bool test(queue q) { buffer OutBuf(A, range<1>(size)); q.submit([&](handler &cgh) { - auto OutAcc = OutBuf.template get_access(cgh); + auto OutAcc = OutBuf.template get_access(cgh); auto LocalAcc = local_accessor(VL * NUM_RGBA_CHANNELS, cgh); cgh.parallel_for(Range, [=](nd_item<1> ndi) SYCL_ESIMD_KERNEL { diff --git a/sycl/test-e2e/ESIMD/lsc/Inputs/lsc_surf.hpp b/sycl/test-e2e/ESIMD/lsc/Inputs/lsc_surf.hpp index 58996e80505cc..8673b398f1cf0 100644 --- a/sycl/test-e2e/ESIMD/lsc/Inputs/lsc_surf.hpp +++ b/sycl/test-e2e/ESIMD/lsc/Inputs/lsc_surf.hpp @@ -44,11 +44,11 @@ int main() { auto buf_3 = buffer{vec_3}; auto buf_4 = buffer{vec_4}; q.submit([&](handler &h) { - auto access_0 = buf_0.template get_access(h); - auto access_1 = buf_1.template get_access(h); - auto access_2 = buf_2.template get_access(h); - auto access_3 = buf_3.template get_access(h); - auto access_4 = buf_4.template get_access(h); + auto access_0 = buf_0.template get_access(h); + auto access_1 = buf_1.template get_access(h); + auto access_2 = buf_2.template get_access(h); + auto access_3 = buf_3.template get_access(h); + auto access_4 = buf_4.template get_access(h); h.parallel_for( range<1>{size / SIMDSize}, [=](id<1> id) SYCL_ESIMD_KERNEL { auto offset = id[0] * SIMDSize * sizeof(int); diff --git a/sycl/test-e2e/ESIMD/lsc/Inputs/lsc_surf_load.hpp b/sycl/test-e2e/ESIMD/lsc/Inputs/lsc_surf_load.hpp index 54be6b640fc5f..2a7a40467dff4 100644 --- a/sycl/test-e2e/ESIMD/lsc/Inputs/lsc_surf_load.hpp +++ b/sycl/test-e2e/ESIMD/lsc/Inputs/lsc_surf_load.hpp @@ -77,8 +77,8 @@ bool test(uint32_t pmask = 0xffffffff) { buffer bufi(in.data(), in.size()); auto e = q.submit([&](handler &cgh) { - auto acco = bufo.template get_access(cgh); - auto acci = bufi.template get_access(cgh); + auto acco = bufo.template get_access(cgh); + auto acci = bufi.template get_access(cgh); cgh.parallel_for>( Range, [=](sycl::nd_item<1> ndi) SYCL_ESIMD_KERNEL { uint16_t globalID = ndi.get_global_id(0); diff --git a/sycl/test-e2e/ESIMD/lsc/Inputs/lsc_surf_store.hpp b/sycl/test-e2e/ESIMD/lsc/Inputs/lsc_surf_store.hpp index 99e2293f3f475..0bd93accb0c03 100644 --- a/sycl/test-e2e/ESIMD/lsc/Inputs/lsc_surf_store.hpp +++ b/sycl/test-e2e/ESIMD/lsc/Inputs/lsc_surf_store.hpp @@ -73,7 +73,7 @@ bool test(uint32_t pmask = 0xffffffff) { buffer bufo(out.data(), out.size()); auto e = q.submit([&](handler &cgh) { - auto acco = bufo.template get_access(cgh); + auto acco = bufo.template get_access(cgh); cgh.parallel_for>( Range, [=](sycl::nd_item<1> ndi) SYCL_ESIMD_KERNEL { uint16_t globalID = ndi.get_global_id(0); diff --git a/sycl/test-e2e/ESIMD/lsc/atomic_smoke.cpp b/sycl/test-e2e/ESIMD/lsc/atomic_smoke.cpp index abe6983c8ae11..0fee7c85c0eb0 100644 --- a/sycl/test-e2e/ESIMD/lsc/atomic_smoke.cpp +++ b/sycl/test-e2e/ESIMD/lsc/atomic_smoke.cpp @@ -325,7 +325,7 @@ bool test(queue q, const Config &cfg) { try { buffer buf(arr, range<1>(size)); auto e = q.submit([&](handler &cgh) { - auto accessor = buf.template get_access(cgh); + auto accessor = buf.template get_access(cgh); cgh.parallel_for>( rng, [=](nd_item<1> gid) SYCL_ESIMD_KERNEL { int i = gid.get_global_id(0); diff --git a/sycl/test-e2e/ESIMD/lsc/lsc_argument_type_deduction.cpp b/sycl/test-e2e/ESIMD/lsc/lsc_argument_type_deduction.cpp index 2de56d3dfcedf..9a83113b5a84f 100644 --- a/sycl/test-e2e/ESIMD/lsc/lsc_argument_type_deduction.cpp +++ b/sycl/test-e2e/ESIMD/lsc/lsc_argument_type_deduction.cpp @@ -31,7 +31,7 @@ template int testAccessor(queue q) { try { auto buf_0 = buffer{vec_0}; q.submit([&](handler &h) { - auto access_0 = buf_0.template get_access(h); + auto access_0 = buf_0.template get_access(h); h.parallel_for( range<1>{size / SIMDSize}, [=](id<1> id) SYCL_ESIMD_KERNEL { diff --git a/sycl/test-e2e/ESIMD/lsc/lsc_block_load_store_stateless_64.cpp b/sycl/test-e2e/ESIMD/lsc/lsc_block_load_store_stateless_64.cpp index b1f590fcfa0b3..a739ef3b755ce 100644 --- a/sycl/test-e2e/ESIMD/lsc/lsc_block_load_store_stateless_64.cpp +++ b/sycl/test-e2e/ESIMD/lsc/lsc_block_load_store_stateless_64.cpp @@ -35,7 +35,7 @@ int main(void) { << "\n"; try { q.submit([&](handler &cgh) { - auto PA = bufa.get_access(cgh); + auto PA = bufa.get_access(cgh); cgh.single_task([=]() SYCL_ESIMD_KERNEL { uint64_t offset = (Size - VL) * sizeof(uint64_t); #ifdef TEST_FLAG diff --git a/sycl/test-e2e/ESIMD/lsc/lsc_gather_scatter_stateless_64.cpp b/sycl/test-e2e/ESIMD/lsc/lsc_gather_scatter_stateless_64.cpp index 1bd973a278fe4..a49e949e269f3 100644 --- a/sycl/test-e2e/ESIMD/lsc/lsc_gather_scatter_stateless_64.cpp +++ b/sycl/test-e2e/ESIMD/lsc/lsc_gather_scatter_stateless_64.cpp @@ -35,7 +35,7 @@ int main() { << "\n"; try { q.submit([&](handler &cgh) { - auto PA = bufa.get_access(cgh); + auto PA = bufa.get_access(cgh); cgh.single_task([=]() SYCL_ESIMD_KERNEL { uint64_t offsetStart = (Size - VL) * sizeof(uint64_t); simd offset(offsetStart, sizeof(uint64_t)); diff --git a/sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_gather_scatter.cpp b/sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_gather_scatter.cpp index 3486390d2a74f..6b1701bfeec98 100644 --- a/sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_gather_scatter.cpp +++ b/sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_gather_scatter.cpp @@ -47,7 +47,7 @@ template bool test(queue q) { nd_range<1> glob_range{range<1>{1}, range<1>{1}}; q.submit([&](handler &cgh) { - auto acc = buf.template get_access(cgh); + auto acc = buf.template get_access(cgh); auto LocalAcc = local_accessor(size * STRIDE, cgh); cgh.parallel_for(glob_range, [=](nd_item<1> ndi) SYCL_ESIMD_KERNEL { using namespace sycl::ext::intel::esimd; diff --git a/sycl/test-e2e/ESIMD/lsc/lsc_predicate.cpp b/sycl/test-e2e/ESIMD/lsc/lsc_predicate.cpp index cc4b23e6d512b..638b99ee2b589 100644 --- a/sycl/test-e2e/ESIMD/lsc/lsc_predicate.cpp +++ b/sycl/test-e2e/ESIMD/lsc/lsc_predicate.cpp @@ -34,8 +34,8 @@ template int testAccessor(queue q) { auto buf_0 = buffer{vec_0}; auto buf_2 = buffer{vec_2}; q.submit([&](handler &h) { - auto access_0 = buf_0.template get_access(h); - auto access_2 = buf_2.template get_access(h); + auto access_0 = buf_0.template get_access(h); + auto access_2 = buf_2.template get_access(h); h.parallel_for( range<1>{size / SIMDSize}, [=](id<1> id) SYCL_ESIMD_KERNEL { diff --git a/sycl/test-e2e/ESIMD/lsc/lsc_predicate_stateless.cpp b/sycl/test-e2e/ESIMD/lsc/lsc_predicate_stateless.cpp index 6267f0d20d690..ae38caffe13f2 100644 --- a/sycl/test-e2e/ESIMD/lsc/lsc_predicate_stateless.cpp +++ b/sycl/test-e2e/ESIMD/lsc/lsc_predicate_stateless.cpp @@ -36,8 +36,8 @@ template int testAccessor(queue q) { auto buf_0 = buffer{vec_0}; auto buf_2 = buffer{vec_2}; q.submit([&](handler &h) { - auto access_0 = buf_0.template get_access(h); - auto access_2 = buf_2.template get_access(h); + auto access_0 = buf_0.template get_access(h); + auto access_2 = buf_2.template get_access(h); h.parallel_for( range<1>{size / SIMDSize}, [=](id<1> id) SYCL_ESIMD_KERNEL { diff --git a/sycl/test-e2e/ESIMD/lsc/lsc_slm.cpp b/sycl/test-e2e/ESIMD/lsc/lsc_slm.cpp index 51517031b1830..c7d88f8a54c80 100644 --- a/sycl/test-e2e/ESIMD/lsc/lsc_slm.cpp +++ b/sycl/test-e2e/ESIMD/lsc/lsc_slm.cpp @@ -42,11 +42,11 @@ int main() { auto buf_3 = buffer{vec_3}; auto buf_4 = buffer{vec_4}; q.submit([&](handler &h) { - auto access_0 = buf_0.template get_access(h); - auto access_1 = buf_1.template get_access(h); - auto access_2 = buf_2.template get_access(h); - auto access_3 = buf_3.template get_access(h); - auto access_4 = buf_4.template get_access(h); + auto access_0 = buf_0.template get_access(h); + auto access_1 = buf_1.template get_access(h); + auto access_2 = buf_2.template get_access(h); + auto access_3 = buf_3.template get_access(h); + auto access_4 = buf_4.template get_access(h); h.parallel_for( range<1>{size / SIMDSize}, [=](id<1> id) SYCL_ESIMD_KERNEL { auto offset = id * SIMDSize * sizeof(int); diff --git a/sycl/test-e2e/ESIMD/mandelbrot/mandelbrot.cpp b/sycl/test-e2e/ESIMD/mandelbrot/mandelbrot.cpp index 3fc8516a4d315..ede66f38c9907 100644 --- a/sycl/test-e2e/ESIMD/mandelbrot/mandelbrot.cpp +++ b/sycl/test-e2e/ESIMD/mandelbrot/mandelbrot.cpp @@ -118,7 +118,7 @@ int main(int argc, char *argv[]) { for (int iter = 0; iter <= num_iters; ++iter) { auto e = q.submit([&](sycl::handler &cgh) { auto accOutput = - imgOutput.get_access(cgh); + imgOutput.get_access(cgh); cgh.parallel_for( GlobalRange * LocalRange, [=](item<2> it) SYCL_ESIMD_KERNEL { diff --git a/sycl/test-e2e/ESIMD/mandelbrot/mandelbrot_spec.cpp b/sycl/test-e2e/ESIMD/mandelbrot/mandelbrot_spec.cpp index 8a1c96304e7a8..85584157642cb 100644 --- a/sycl/test-e2e/ESIMD/mandelbrot/mandelbrot_spec.cpp +++ b/sycl/test-e2e/ESIMD/mandelbrot/mandelbrot_spec.cpp @@ -144,7 +144,7 @@ int main(int argc, char *argv[]) { for (int iter = 0; iter <= num_iters; ++iter) { auto e = q.submit([&](sycl::handler &cgh) { auto accOutput = - imgOutput.get_access(cgh); + imgOutput.get_access(cgh); cgh.set_specialization_constant(crunch); cgh.set_specialization_constant(xoff); diff --git a/sycl/test-e2e/ESIMD/mask_compress_store_acc.cpp b/sycl/test-e2e/ESIMD/mask_compress_store_acc.cpp index db4669aa444ed..fb36495c0a54e 100644 --- a/sycl/test-e2e/ESIMD/mask_compress_store_acc.cpp +++ b/sycl/test-e2e/ESIMD/mask_compress_store_acc.cpp @@ -35,7 +35,7 @@ template bool test() { auto e = Queue.submit([&](sycl::handler &cgh) { auto OutputAcc_out = - OutputAcc_buffer.get_access(cgh); + OutputAcc_buffer.get_access(cgh); auto kernel = ([=]() [[intel::sycl_explicit_simd]] { simd Input(1, 1); diff --git a/sycl/test-e2e/ESIMD/mask_expand_load.cpp b/sycl/test-e2e/ESIMD/mask_expand_load.cpp index f586c3aeb075f..9ee815573a2d5 100644 --- a/sycl/test-e2e/ESIMD/mask_expand_load.cpp +++ b/sycl/test-e2e/ESIMD/mask_expand_load.cpp @@ -44,7 +44,7 @@ template bool test(sycl::queue &Queue) { buffer InputAcc_buffer(Input.data(), Input.size()); Queue.submit([&](sycl::handler &cgh) { - auto InputAcc = InputAcc_buffer.get_access(cgh); + auto InputAcc = InputAcc_buffer.get_access(cgh); auto Kernel = ([=]() SYCL_ESIMD_KERNEL { simd_mask Mask; for (int i = 0; i < N; i++) diff --git a/sycl/test-e2e/ESIMD/matrix_transpose.cpp b/sycl/test-e2e/ESIMD/matrix_transpose.cpp index 928e3ac7212f8..18b55c5beb6b9 100644 --- a/sycl/test-e2e/ESIMD/matrix_transpose.cpp +++ b/sycl/test-e2e/ESIMD/matrix_transpose.cpp @@ -303,7 +303,7 @@ bool runTest(unsigned MZ, unsigned block_size, unsigned num_iters, double etime = 0; if (block_size == 16 && MZ >= 16) { auto e = q.submit([&](handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.parallel_for( Range, [=](nd_item<2> ndi) SYCL_ESIMD_KERNEL { transpose16(acc, MZ, ndi.get_global_id(0), @@ -315,7 +315,7 @@ bool runTest(unsigned MZ, unsigned block_size, unsigned num_iters, etime = esimd_test::report_time("kernel time", e, e); } else if (block_size == 8) { auto e = q.submit([&](handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.parallel_for( Range, [=](nd_item<2> ndi) SYCL_ESIMD_KERNEL { transpose8(acc, MZ, ndi.get_global_id(0), ndi.get_global_id(1)); diff --git a/sycl/test-e2e/ESIMD/matrix_transpose2.cpp b/sycl/test-e2e/ESIMD/matrix_transpose2.cpp index bbf26b2797e74..57724d35904dc 100644 --- a/sycl/test-e2e/ESIMD/matrix_transpose2.cpp +++ b/sycl/test-e2e/ESIMD/matrix_transpose2.cpp @@ -303,8 +303,8 @@ bool runTest(unsigned MZ, unsigned block_size, unsigned num_iters, double etime = 0; if (block_size == 16 && MZ >= 16) { auto e = q.submit([&](handler &cgh) { - auto accInput = imgM.get_access(cgh); - auto accOutput = imgM.get_access(cgh); + auto accInput = imgM.get_access(cgh); + auto accOutput = imgM.get_access(cgh); cgh.parallel_for( Range, [=](nd_item<2> ndi) SYCL_ESIMD_KERNEL { transpose16(accInput, accOutput, MZ, ndi.get_global_id(0), @@ -315,8 +315,8 @@ bool runTest(unsigned MZ, unsigned block_size, unsigned num_iters, etime = esimd_test::report_time("kernel time", e, e); } else if (block_size == 8) { auto e = q.submit([&](handler &cgh) { - auto accInput = imgM.get_access(cgh); - auto accOutput = imgM.get_access(cgh); + auto accInput = imgM.get_access(cgh); + auto accOutput = imgM.get_access(cgh); cgh.parallel_for( Range, [=](nd_item<2> ndi) SYCL_ESIMD_KERNEL { transpose8(accInput, accOutput, MZ, ndi.get_global_id(0), diff --git a/sycl/test-e2e/ESIMD/named_barriers/allocate_barrier.cpp b/sycl/test-e2e/ESIMD/named_barriers/allocate_barrier.cpp index aeae189271d8a..04789a991e7c3 100644 --- a/sycl/test-e2e/ESIMD/named_barriers/allocate_barrier.cpp +++ b/sycl/test-e2e/ESIMD/named_barriers/allocate_barrier.cpp @@ -95,7 +95,7 @@ bool test(QueueTY q) { sycl::nd_range<1> Range{GlobalRange * LocalRange, LocalRange}; auto e = q.submit([&](handler &cgh) { - auto acc = buf.template get_access(cgh); + auto acc = buf.template get_access(cgh); cgh.parallel_for>( Range, [=](sycl::nd_item<1> ndi) SYCL_ESIMD_KERNEL { // Threads - 1 named barriers required diff --git a/sycl/test-e2e/ESIMD/named_barriers/exec_in_order.cpp b/sycl/test-e2e/ESIMD/named_barriers/exec_in_order.cpp index 5b601d9f39e7e..00d2c6c09b807 100644 --- a/sycl/test-e2e/ESIMD/named_barriers/exec_in_order.cpp +++ b/sycl/test-e2e/ESIMD/named_barriers/exec_in_order.cpp @@ -59,7 +59,7 @@ bool test(QueueTY q) { sycl::nd_range<1> Range{GlobalRange * LocalRange, LocalRange}; auto e = q.submit([&](handler &cgh) { - auto acc = buf.template get_access(cgh); + auto acc = buf.template get_access(cgh); cgh.parallel_for>( Range, [=](sycl::nd_item<1> ndi) SYCL_ESIMD_KERNEL { // Threads - 1 named barriers required diff --git a/sycl/test-e2e/ESIMD/named_barriers/exec_in_order_branched.cpp b/sycl/test-e2e/ESIMD/named_barriers/exec_in_order_branched.cpp index 0aa5fc6dd98bc..8b897cdba8c1a 100644 --- a/sycl/test-e2e/ESIMD/named_barriers/exec_in_order_branched.cpp +++ b/sycl/test-e2e/ESIMD/named_barriers/exec_in_order_branched.cpp @@ -58,7 +58,7 @@ bool test(QueueTY q) { sycl::nd_range<1> Range{GlobalRange * LocalRange, LocalRange}; auto e = q.submit([&](handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.parallel_for>( Range, [=](sycl::nd_item<1> ndi) SYCL_ESIMD_KERNEL { // Threads - 1 named barriers required diff --git a/sycl/test-e2e/ESIMD/named_barriers/loop.cpp b/sycl/test-e2e/ESIMD/named_barriers/loop.cpp index 8c186cbe4d50f..42a1657afe2fe 100644 --- a/sycl/test-e2e/ESIMD/named_barriers/loop.cpp +++ b/sycl/test-e2e/ESIMD/named_barriers/loop.cpp @@ -47,7 +47,7 @@ bool test(QueueTY q) { sycl::nd_range<1> Range{GlobalRange * LocalRange, LocalRange}; auto e = q.submit([&](handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.parallel_for>( Range, [=](sycl::nd_item<1> ndi) SYCL_ESIMD_KERNEL { // 2 named barriers, id 0 reserved for unnamed diff --git a/sycl/test-e2e/ESIMD/named_barriers/loop_extended.cpp b/sycl/test-e2e/ESIMD/named_barriers/loop_extended.cpp index 474f3bdece79c..3d890561d42b7 100644 --- a/sycl/test-e2e/ESIMD/named_barriers/loop_extended.cpp +++ b/sycl/test-e2e/ESIMD/named_barriers/loop_extended.cpp @@ -48,7 +48,7 @@ bool test(QueueTY q) { sycl::nd_range<1> Range{GlobalRange * LocalRange, LocalRange}; auto e = q.submit([&](handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.parallel_for>( Range, [=](sycl::nd_item<1> ndi) SYCL_ESIMD_KERNEL { // 2 named barriers, id 0 reserved for unnamed diff --git a/sycl/test-e2e/ESIMD/named_barriers/multiple_wg.cpp b/sycl/test-e2e/ESIMD/named_barriers/multiple_wg.cpp index 25b6b53f9d0de..a089719fce05f 100644 --- a/sycl/test-e2e/ESIMD/named_barriers/multiple_wg.cpp +++ b/sycl/test-e2e/ESIMD/named_barriers/multiple_wg.cpp @@ -49,7 +49,7 @@ bool test(QueueTY q) { sycl::nd_range<1> Range{GlobalRange * LocalRange, LocalRange}; auto e = q.submit([&](handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.parallel_for>( Range, [=](sycl::nd_item<1> ndi) SYCL_ESIMD_KERNEL { // 1 named barrier, id 0 reserved for unnamed diff --git a/sycl/test-e2e/ESIMD/named_barriers/single_wg.cpp b/sycl/test-e2e/ESIMD/named_barriers/single_wg.cpp index 694737455e6b1..4be50fbe0b1bf 100644 --- a/sycl/test-e2e/ESIMD/named_barriers/single_wg.cpp +++ b/sycl/test-e2e/ESIMD/named_barriers/single_wg.cpp @@ -51,7 +51,7 @@ bool test(QueueTY q) { sycl::nd_range<1> Range{GlobalRange * LocalRange, LocalRange}; auto e = q.submit([&](handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.parallel_for>( Range, [=](sycl::nd_item<1> ndi) SYCL_ESIMD_KERNEL { // 1 named barrier, id 0 reserved for unnamed diff --git a/sycl/test-e2e/ESIMD/noinline_call_from_func.cpp b/sycl/test-e2e/ESIMD/noinline_call_from_func.cpp index 9a4ac7ec55e16..65ab47a525569 100644 --- a/sycl/test-e2e/ESIMD/noinline_call_from_func.cpp +++ b/sycl/test-e2e/ESIMD/noinline_call_from_func.cpp @@ -40,7 +40,7 @@ int main(int argc, char **argv) { buffer buf(output, range<1>(1)); q.submit([&](handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.parallel_for( sycl::range<1>{1}, diff --git a/sycl/test-e2e/ESIMD/popcount.cpp b/sycl/test-e2e/ESIMD/popcount.cpp index 1b66793b46435..c54064900509f 100644 --- a/sycl/test-e2e/ESIMD/popcount.cpp +++ b/sycl/test-e2e/ESIMD/popcount.cpp @@ -43,8 +43,8 @@ template bool test(queue &q) { queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler()); auto e = q.submit([&](handler &cgh) { - auto PA = bufa.template get_access(cgh); - auto PB = bufb.template get_access(cgh); + auto PA = bufa.template get_access(cgh); + auto PB = bufb.template get_access(cgh); cgh.parallel_for(GlobalRange * LocalRange, [=](id<1> i) SYCL_ESIMD_KERNEL { using namespace sycl::ext::intel::esimd; diff --git a/sycl/test-e2e/ESIMD/regression/big_const_initializer.cpp b/sycl/test-e2e/ESIMD/regression/big_const_initializer.cpp index 220436c9f85f9..f64e0031f07de 100644 --- a/sycl/test-e2e/ESIMD/regression/big_const_initializer.cpp +++ b/sycl/test-e2e/ESIMD/regression/big_const_initializer.cpp @@ -45,7 +45,7 @@ int main(int argc, char **argv) { try { queue.submit([&](sycl::handler &cgh) { - auto r_acc = r.template get_access(cgh); + auto r_acc = r.template get_access(cgh); cgh.parallel_for( sycl::range<1>{nsamples / SIMD_WIDTH}, [=](sycl::item<1> item) SYCL_ESIMD_KERNEL { diff --git a/sycl/test-e2e/ESIMD/regression/copyto_char_test.cpp b/sycl/test-e2e/ESIMD/regression/copyto_char_test.cpp index 9eaeb21c3c861..6e42fd7e8af7f 100644 --- a/sycl/test-e2e/ESIMD/regression/copyto_char_test.cpp +++ b/sycl/test-e2e/ESIMD/regression/copyto_char_test.cpp @@ -56,7 +56,7 @@ template int test_to_copy() { DataT *init_ref_ptr = initial_ref_data.data(); DataT *ref_data_for_fill_ptr = ref_data_for_fill.data(); auto acc = - output_buf.template get_access(cgh); + output_buf.template get_access(cgh); cgh.single_task([=]() SYCL_ESIMD_KERNEL { simd src_simd_obj; diff --git a/sycl/test-e2e/ESIMD/regression/fmod_compatibility_test.cpp b/sycl/test-e2e/ESIMD/regression/fmod_compatibility_test.cpp index d449c11858f59..526be3b8d9f96 100644 --- a/sycl/test-e2e/ESIMD/regression/fmod_compatibility_test.cpp +++ b/sycl/test-e2e/ESIMD/regression/fmod_compatibility_test.cpp @@ -16,7 +16,7 @@ #include #include -constexpr auto sycl_write = sycl::access::mode::write; +constexpr auto sycl_write = sycl::access_mode::write; #define SIMD 16 int test_fmod(float x, float y) { diff --git a/sycl/test-e2e/ESIMD/regression/half_conversion_test.cpp b/sycl/test-e2e/ESIMD/regression/half_conversion_test.cpp index 063592c31696b..66b8789b37cb8 100644 --- a/sycl/test-e2e/ESIMD/regression/half_conversion_test.cpp +++ b/sycl/test-e2e/ESIMD/regression/half_conversion_test.cpp @@ -42,7 +42,7 @@ template bool test(queue q, int inc) { std::cout << "Running on " << q.get_device().get_info<::sycl::info::device::name>() << "\n"; - auto acc = buf.template get_access(cgh); + auto acc = buf.template get_access(cgh); cgh.single_task([=]() SYCL_ESIMD_KERNEL { simd offsets(0); simd vec = gather(acc, offsets); diff --git a/sycl/test-e2e/ESIMD/regression/operator_decrement.cpp b/sycl/test-e2e/ESIMD/regression/operator_decrement.cpp index 88efb3c32ae2e..44e0fa4132226 100644 --- a/sycl/test-e2e/ESIMD/regression/operator_decrement.cpp +++ b/sycl/test-e2e/ESIMD/regression/operator_decrement.cpp @@ -40,7 +40,7 @@ int main(void) { std::cout << "Running on " << dev.get_info() << "\n"; auto e = q.submit([&](handler &cgh) { - auto PA = bufa.get_access(cgh); + auto PA = bufa.get_access(cgh); cgh.parallel_for( GlobalRange * LocalRange, [=](id<1> i) SYCL_ESIMD_KERNEL { using namespace sycl::ext::intel::esimd; diff --git a/sycl/test-e2e/ESIMD/regression/sycl_esimd_mixed_unnamed.cpp b/sycl/test-e2e/ESIMD/regression/sycl_esimd_mixed_unnamed.cpp index bff75087186a6..a69324b793a3e 100644 --- a/sycl/test-e2e/ESIMD/regression/sycl_esimd_mixed_unnamed.cpp +++ b/sycl/test-e2e/ESIMD/regression/sycl_esimd_mixed_unnamed.cpp @@ -61,7 +61,7 @@ int main(void) { std::cout << "Running on " << dev.get_info() << "\n"; auto e = q.submit([&](handler &cgh) { - auto PA = bufa.get_access(cgh); + auto PA = bufa.get_access(cgh); cgh.parallel_for(GlobalRange * LocalRange, [=](id<1> i) { PA[i] = PA[i] + 1; }); }); @@ -92,7 +92,7 @@ int main(void) { std::cout << "Running on " << dev.get_info() << "\n"; auto e = q.submit([&](handler &cgh) { - auto PA = bufa.get_access(cgh); + auto PA = bufa.get_access(cgh); cgh.parallel_for(GlobalRange * LocalRange, [=](id<1> i) SYCL_ESIMD_KERNEL { using namespace sycl::ext::intel::esimd; diff --git a/sycl/test-e2e/ESIMD/regression/tanh_fix_test.cpp b/sycl/test-e2e/ESIMD/regression/tanh_fix_test.cpp index 1dec9ecd151c3..0408dda7df133 100644 --- a/sycl/test-e2e/ESIMD/regression/tanh_fix_test.cpp +++ b/sycl/test-e2e/ESIMD/regression/tanh_fix_test.cpp @@ -16,7 +16,7 @@ #include #include -constexpr auto sycl_write = sycl::access::mode::write; +constexpr auto sycl_write = sycl::access_mode::write; #define SIMD 16 int test_tanh(float x) { diff --git a/sycl/test-e2e/ESIMD/regression/unused_load.cpp b/sycl/test-e2e/ESIMD/regression/unused_load.cpp index 1400cf5d31744..865cd42632fc7 100644 --- a/sycl/test-e2e/ESIMD/regression/unused_load.cpp +++ b/sycl/test-e2e/ESIMD/regression/unused_load.cpp @@ -34,7 +34,7 @@ int main() { std::cout << "Running on " << q.get_device().get_info() << "\n"; - auto acc0 = buf0.get_access(cgh); + auto acc0 = buf0.get_access(cgh); cgh.parallel_for(range<1>(1), [=](sycl::id<1> i) SYCL_ESIMD_KERNEL { diff --git a/sycl/test-e2e/ESIMD/regression/variable_gather_mask.cpp b/sycl/test-e2e/ESIMD/regression/variable_gather_mask.cpp index 99ec0d7a100cb..aa397d6c84d59 100644 --- a/sycl/test-e2e/ESIMD/regression/variable_gather_mask.cpp +++ b/sycl/test-e2e/ESIMD/regression/variable_gather_mask.cpp @@ -139,7 +139,7 @@ int main(int argc, char **argv) { sycl::buffer Bbuf(B, range<1>(VL)); return q.submit([&](handler &cgh) { - auto b = Bbuf.template get_access(cgh); + auto b = Bbuf.template get_access(cgh); KernelAcc kernel(b); cgh.single_task(kernel); }); diff --git a/sycl/test-e2e/ESIMD/spec_const/Inputs/spec-const-2020-common.hpp b/sycl/test-e2e/ESIMD/spec_const/Inputs/spec-const-2020-common.hpp index e573a6cd8c73c..827080ec98f33 100644 --- a/sycl/test-e2e/ESIMD/spec_const/Inputs/spec-const-2020-common.hpp +++ b/sycl/test-e2e/ESIMD/spec_const/Inputs/spec-const-2020-common.hpp @@ -49,7 +49,7 @@ int main(int argc, char **argv) { sycl::buffer buf(output.data(), output.size()); q.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); if (i % 2 != 0) cgh.set_specialization_constant(REDEF_VAL); cgh.single_task([=](kernel_handler kh) SYCL_ESIMD_KERNEL { diff --git a/sycl/test-e2e/ESIMD/spec_const/spec_const_redefine.cpp b/sycl/test-e2e/ESIMD/spec_const/spec_const_redefine.cpp index 7d9910b3a24f4..beab9b40139d5 100644 --- a/sycl/test-e2e/ESIMD/spec_const/spec_const_redefine.cpp +++ b/sycl/test-e2e/ESIMD/spec_const/spec_const_redefine.cpp @@ -64,7 +64,7 @@ int main(int argc, char **argv) { cgh.set_specialization_constant(9999); cgh.set_specialization_constant(sc_set[0]); cgh.set_specialization_constant(sc_set[1]); - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.single_task( [=](kernel_handler kh) SYCL_ESIMD_KERNEL { auto SC0Val = kh.get_specialization_constant(); diff --git a/sycl/test-e2e/ESIMD/sycl_esimd_mix.cpp b/sycl/test-e2e/ESIMD/sycl_esimd_mix.cpp index 60cd44faa88b7..7f8681b20efab 100644 --- a/sycl/test-e2e/ESIMD/sycl_esimd_mix.cpp +++ b/sycl/test-e2e/ESIMD/sycl_esimd_mix.cpp @@ -61,7 +61,7 @@ int main(void) { std::cout << "Running on " << dev.get_info() << "\n"; auto e = q.submit([&](handler &cgh) { - auto PA = bufa.get_access(cgh); + auto PA = bufa.get_access(cgh); cgh.parallel_for(GlobalRange * LocalRange, [=](id<1> i) { PA[i] = PA[i] + 1; }); }); @@ -92,7 +92,7 @@ int main(void) { std::cout << "Running on " << dev.get_info() << "\n"; auto e = q.submit([&](handler &cgh) { - auto PA = bufa.get_access(cgh); + auto PA = bufa.get_access(cgh); cgh.parallel_for( GlobalRange * LocalRange, [=](id<1> i) SYCL_ESIMD_KERNEL { using namespace sycl::ext::intel::esimd; diff --git a/sycl/test-e2e/ESIMD/template.cpp b/sycl/test-e2e/ESIMD/template.cpp index 2b916643f46da..0b86ed4c67511 100644 --- a/sycl/test-e2e/ESIMD/template.cpp +++ b/sycl/test-e2e/ESIMD/template.cpp @@ -25,9 +25,9 @@ sycl::event createKernel(sycl::queue &q, buffer &bufa, buffer &bufb, // We need that many threads in each group range<1> LocalRange{1}; return q.submit([&](handler &cgh) { - auto PA = bufa.template get_access(cgh); - auto PB = bufb.template get_access(cgh); - auto PC = bufc.template get_access(cgh); + auto PA = bufa.template get_access(cgh); + auto PB = bufb.template get_access(cgh); + auto PC = bufc.template get_access(cgh); cgh.parallel_for(GlobalRange * LocalRange, [=](id<1> i) SYCL_ESIMD_KERNEL { using namespace sycl::ext::intel::esimd; diff --git a/sycl/test-e2e/ESIMD/unified_memory_api/Inputs/atomic_update.hpp b/sycl/test-e2e/ESIMD/unified_memory_api/Inputs/atomic_update.hpp index df7d32bd7c246..5a5615cecbce1 100644 --- a/sycl/test-e2e/ESIMD/unified_memory_api/Inputs/atomic_update.hpp +++ b/sycl/test-e2e/ESIMD/unified_memory_api/Inputs/atomic_update.hpp @@ -281,7 +281,7 @@ bool test_acc(queue &q, const Config &cfg) { buffer arr_buf(arr, range<1>(size)); auto e = q.submit([&](handler &cgh) { auto arr_acc = - arr_buf.template get_access(cgh); + arr_buf.template get_access(cgh); cgh.parallel_for(rng, [=](nd_item<1> ndi) SYCL_ESIMD_KERNEL { int i = ndi.get_global_id(0); simd offsets(cfg.start_ind * sizeof(T), diff --git a/sycl/test-e2e/ESIMD/vadd_1d.cpp b/sycl/test-e2e/ESIMD/vadd_1d.cpp index 6375017602ef6..9a72a27c5a1ec 100644 --- a/sycl/test-e2e/ESIMD/vadd_1d.cpp +++ b/sycl/test-e2e/ESIMD/vadd_1d.cpp @@ -43,9 +43,9 @@ int main(void) { std::cout << "Running on " << dev.get_info() << "\n"; auto e = q.submit([&](handler &cgh) { - auto PA = bufa.get_access(cgh); - auto PB = bufb.get_access(cgh); - auto PC = bufc.get_access(cgh); + auto PA = bufa.get_access(cgh); + auto PB = bufb.get_access(cgh); + auto PC = bufc.get_access(cgh); cgh.parallel_for( GlobalRange * LocalRange, [=](id<1> i) SYCL_ESIMD_KERNEL { using namespace sycl::ext::intel::esimd; diff --git a/sycl/test-e2e/ESIMD/vadd_2d.cpp b/sycl/test-e2e/ESIMD/vadd_2d.cpp index 96d1e6f81a08d..8068672f1022b 100644 --- a/sycl/test-e2e/ESIMD/vadd_2d.cpp +++ b/sycl/test-e2e/ESIMD/vadd_2d.cpp @@ -49,9 +49,9 @@ int main(void) { std::cout << "Running on " << dev.get_info() << "\n"; auto e = q.submit([&](handler &cgh) { - auto accA = imgA.get_access(cgh); - auto accB = imgB.get_access(cgh); - auto accC = imgC.get_access(cgh); + auto accA = imgA.get_access(cgh); + auto accB = imgB.get_access(cgh); + auto accC = imgC.get_access(cgh); cgh.parallel_for( GlobalRange * LocalRange, [=](id<1> i) SYCL_ESIMD_KERNEL { diff --git a/sycl/test-e2e/ESIMD/vadd_2d_acc.cpp b/sycl/test-e2e/ESIMD/vadd_2d_acc.cpp index 83ed144893fb8..1218f4aa5b3f5 100644 --- a/sycl/test-e2e/ESIMD/vadd_2d_acc.cpp +++ b/sycl/test-e2e/ESIMD/vadd_2d_acc.cpp @@ -49,9 +49,9 @@ int main(void) { std::cout << "Running on " << dev.get_info() << "\n"; auto e = q.submit([&](handler &cgh) { - auto PA = bufa.get_access(cgh); - auto PB = bufb.get_access(cgh); - auto PC = bufc.get_access(cgh); + auto PA = bufa.get_access(cgh); + auto PB = bufb.get_access(cgh); + auto PC = bufc.get_access(cgh); cgh.parallel_for( Range, [=](nd_item<2> ndi) SYCL_ESIMD_KERNEL { using namespace sycl::ext::intel::esimd; diff --git a/sycl/test-e2e/InvokeSimd/Feature/popcnt_emu.cpp b/sycl/test-e2e/InvokeSimd/Feature/popcnt_emu.cpp index 85ab93d7d43a9..d5d1b8bc35e14 100644 --- a/sycl/test-e2e/InvokeSimd/Feature/popcnt_emu.cpp +++ b/sycl/test-e2e/InvokeSimd/Feature/popcnt_emu.cpp @@ -108,7 +108,7 @@ int main(void) { << "\n"; auto e = q.submit([&](handler &cgh) { - auto out_accessor = buf.get_access(cgh); + auto out_accessor = buf.get_access(cgh); cgh.parallel_for( nd_range<1>(GlobalRange, LocalRange), diff --git a/sycl/test-e2e/InvokeSimd/Feature/scale.cpp b/sycl/test-e2e/InvokeSimd/Feature/scale.cpp index 5e4af520b43c3..b62cdc0776c76 100644 --- a/sycl/test-e2e/InvokeSimd/Feature/scale.cpp +++ b/sycl/test-e2e/InvokeSimd/Feature/scale.cpp @@ -90,8 +90,8 @@ template bool test(QueueTY q) { sycl::range<1> LocalRange{VL}; auto e = q.submit([&](handler &cgh) { - auto acca = bufa.template get_access(cgh); - auto accc = bufc.template get_access(cgh); + auto acca = bufa.template get_access(cgh); + auto accc = bufc.template get_access(cgh); cgh.parallel_for>( nd_range<1>(GlobalRange, LocalRange), diff --git a/sycl/test-e2e/InvokeSimd/Feature/void_retval.cpp b/sycl/test-e2e/InvokeSimd/Feature/void_retval.cpp index 5989e49dfe41d..ceaedf17d3164 100644 --- a/sycl/test-e2e/InvokeSimd/Feature/void_retval.cpp +++ b/sycl/test-e2e/InvokeSimd/Feature/void_retval.cpp @@ -91,9 +91,9 @@ int main(void) { << "\n"; auto e = q.submit([&](handler &cgh) { - auto PA = bufa.get_access(cgh); - auto PB = bufb.get_access(cgh); - auto PC = bufc.get_access(cgh); + auto PA = bufa.get_access(cgh); + auto PB = bufb.get_access(cgh); + auto PC = bufc.get_access(cgh); cgh.parallel_for( nd_range<1>(GlobalRange, LocalRange), diff --git a/sycl/test-e2e/InvokeSimd/Regression/call_vadd_1d_spill.cpp b/sycl/test-e2e/InvokeSimd/Regression/call_vadd_1d_spill.cpp index 538b7e370823c..f56acde724af0 100644 --- a/sycl/test-e2e/InvokeSimd/Regression/call_vadd_1d_spill.cpp +++ b/sycl/test-e2e/InvokeSimd/Regression/call_vadd_1d_spill.cpp @@ -93,16 +93,16 @@ bool test(QueueTY q, float *A, float *B, float *C, float *P, float *Q, float *R, sycl::range<1> LocalRange{VL}; auto e = q.submit([&](handler &cgh) { - auto PA = bufa.get_access(cgh); - auto PB = bufb.get_access(cgh); - auto PP = bufp.get_access(cgh); - auto PQ = bufq.get_access(cgh); - auto PR = bufr.get_access(cgh); - auto PX = bufx.get_access(cgh); - auto PY = bufy.get_access(cgh); - auto PZ = bufz.get_access(cgh); - - auto PC = bufc.get_access(cgh); + auto PA = bufa.get_access(cgh); + auto PB = bufb.get_access(cgh); + auto PP = bufp.get_access(cgh); + auto PQ = bufq.get_access(cgh); + auto PR = bufr.get_access(cgh); + auto PX = bufx.get_access(cgh); + auto PY = bufy.get_access(cgh); + auto PZ = bufz.get_access(cgh); + + auto PC = bufc.get_access(cgh); cgh.parallel_for>( nd_range<1>(GlobalRange, LocalRange), diff --git a/sycl/test-e2e/InvokeSimd/Regression/debug_symbols.cpp b/sycl/test-e2e/InvokeSimd/Regression/debug_symbols.cpp index 11c118dad7615..d46c59842c76c 100644 --- a/sycl/test-e2e/InvokeSimd/Regression/debug_symbols.cpp +++ b/sycl/test-e2e/InvokeSimd/Regression/debug_symbols.cpp @@ -73,9 +73,9 @@ int main(void) { << "\n"; auto e = q.submit([&](handler &cgh) { - auto PA = bufa.get_access(cgh); - auto PB = bufb.get_access(cgh); - auto PC = bufc.get_access(cgh); + auto PA = bufa.get_access(cgh); + auto PB = bufb.get_access(cgh); + auto PC = bufc.get_access(cgh); cgh.parallel_for( nd_range<1>(GlobalRange, LocalRange), diff --git a/sycl/test-e2e/InvokeSimd/Spec/ESIMD_to_unmarked_function.cpp b/sycl/test-e2e/InvokeSimd/Spec/ESIMD_to_unmarked_function.cpp index 2e24ad430ffa1..f3cdb91e8c46a 100644 --- a/sycl/test-e2e/InvokeSimd/Spec/ESIMD_to_unmarked_function.cpp +++ b/sycl/test-e2e/InvokeSimd/Spec/ESIMD_to_unmarked_function.cpp @@ -107,9 +107,9 @@ int main(void) { << "\n"; auto e = q.submit([&](handler &cgh) { - auto PA = bufa.get_access(cgh); - auto PB = bufb.get_access(cgh); - auto PC = bufc.get_access(cgh); + auto PA = bufa.get_access(cgh); + auto PB = bufb.get_access(cgh); + auto PC = bufc.get_access(cgh); cgh.parallel_for( nd_range<1>(GlobalRange, LocalRange), diff --git a/sycl/test-e2e/InvokeSimd/Spec/function_overloads.cpp b/sycl/test-e2e/InvokeSimd/Spec/function_overloads.cpp index 71be83ded52d1..5d7d05ef1c3aa 100644 --- a/sycl/test-e2e/InvokeSimd/Spec/function_overloads.cpp +++ b/sycl/test-e2e/InvokeSimd/Spec/function_overloads.cpp @@ -105,8 +105,8 @@ int main(void) { << "\n"; auto e = q.submit([&](handler &cgh) { - auto PA = bufa.get_access(cgh); - auto PC = bufc.get_access(cgh); + auto PA = bufa.get_access(cgh); + auto PC = bufc.get_access(cgh); cgh.parallel_for( nd_range<1>(GlobalRange, LocalRange), diff --git a/sycl/test-e2e/InvokeSimd/Spec/nested_ESIMD_to_ESIMD.cpp b/sycl/test-e2e/InvokeSimd/Spec/nested_ESIMD_to_ESIMD.cpp index 762f98480bcd8..52dddcdea084a 100644 --- a/sycl/test-e2e/InvokeSimd/Spec/nested_ESIMD_to_ESIMD.cpp +++ b/sycl/test-e2e/InvokeSimd/Spec/nested_ESIMD_to_ESIMD.cpp @@ -91,9 +91,9 @@ int main(void) { << "\n"; auto e = q.submit([&](handler &cgh) { - auto PA = bufa.get_access(cgh); - auto PB = bufb.get_access(cgh); - auto PC = bufc.get_access(cgh); + auto PA = bufa.get_access(cgh); + auto PB = bufb.get_access(cgh); + auto PC = bufc.get_access(cgh); cgh.parallel_for( nd_range<1>(GlobalRange, LocalRange), diff --git a/sycl/test-e2e/InvokeSimd/Spec/nested_SPMD_to_ESIMD.cpp b/sycl/test-e2e/InvokeSimd/Spec/nested_SPMD_to_ESIMD.cpp index b0c062a2d72fe..87537b4987223 100644 --- a/sycl/test-e2e/InvokeSimd/Spec/nested_SPMD_to_ESIMD.cpp +++ b/sycl/test-e2e/InvokeSimd/Spec/nested_SPMD_to_ESIMD.cpp @@ -96,9 +96,9 @@ int main(void) { << "\n"; auto e = q.submit([&](handler &cgh) { - auto PA = bufa.get_access(cgh); - auto PB = bufb.get_access(cgh); - auto PC = bufc.get_access(cgh); + auto PA = bufa.get_access(cgh); + auto PB = bufb.get_access(cgh); + auto PC = bufc.get_access(cgh); cgh.parallel_for( nd_range<1>(GlobalRange, LocalRange), diff --git a/sycl/test-e2e/InvokeSimd/Spec/uniform_retval.cpp b/sycl/test-e2e/InvokeSimd/Spec/uniform_retval.cpp index bd2458ff278b8..9656ac15ed8e4 100644 --- a/sycl/test-e2e/InvokeSimd/Spec/uniform_retval.cpp +++ b/sycl/test-e2e/InvokeSimd/Spec/uniform_retval.cpp @@ -160,8 +160,8 @@ template bool test(QueueTY q) { sycl::range<1> LocalRange{VL}; auto e = q.submit([&](handler &cgh) { - auto acca = bufa.template get_access(cgh); - auto accc = bufc.template get_access(cgh); + auto acca = bufa.template get_access(cgh); + auto accc = bufc.template get_access(cgh); cgh.parallel_for>( nd_range<1>(GlobalRange, LocalRange),