diff --git a/sycl/test-e2e/AddressCast/dynamic_address_cast.cpp b/sycl/test-e2e/AddressCast/dynamic_address_cast.cpp index bf77bc1aed306..728a01b680426 100644 --- a/sycl/test-e2e/AddressCast/dynamic_address_cast.cpp +++ b/sycl/test-e2e/AddressCast/dynamic_address_cast.cpp @@ -33,10 +33,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/AddressCast/static_address_cast.cpp b/sycl/test-e2e/AddressCast/static_address_cast.cpp index e7630f1715e5a..ce5834b47eebe 100644 --- a/sycl/test-e2e/AddressCast/static_address_cast.cpp +++ b/sycl/test-e2e/AddressCast/static_address_cast.cpp @@ -25,10 +25,10 @@ int main() { Queue .submit([&](sycl::handler &cgh) { auto GlobalAccessor = - GlobalBuffer.get_access(cgh); + GlobalBuffer.get_access(cgh); auto LocalAccessor = sycl::local_accessor(1, cgh); auto ResultAccessor = - ResultBuffer.get_access(cgh); + ResultBuffer.get_access(cgh); cgh.parallel_for( sycl::nd_range<1>(NItems, 1), [=](sycl::nd_item<1> Item) { bool Success = true; diff --git a/sycl/test-e2e/Assert/assert_in_multiple_tus.hpp b/sycl/test-e2e/Assert/assert_in_multiple_tus.hpp index 39d77db87edfe..d9e884d8fbd6d 100644 --- a/sycl/test-e2e/Assert/assert_in_multiple_tus.hpp +++ b/sycl/test-e2e/Assert/assert_in_multiple_tus.hpp @@ -29,7 +29,7 @@ void enqueueKernel_1_fromFile1(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( sycl::nd_range(Buf.get_range(), sycl::range<1>(4)), diff --git a/sycl/test-e2e/AtomicRef/add.h b/sycl/test-e2e/AtomicRef/add.h index 69b0eaa16b2ee..1513179672d44 100644 --- a/sycl/test-e2e/AtomicRef/add.h +++ b/sycl/test-e2e/AtomicRef/add.h @@ -32,9 +32,9 @@ void add_fetch_local_test(queue q, size_t N) { buffer sum_buf(&sum, 1); buffer output_buf(output.data(), output.size()); q.submit([&](handler &cgh) { - auto sum = sum_buf.template get_access(cgh); + auto sum = sum_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) { @@ -81,9 +81,9 @@ void add_fetch_test(queue q, size_t N) { buffer sum_buf(&sum, 1); buffer output_buf(output.data(), output.size()); q.submit([&](handler &cgh) { - auto sum = sum_buf.template get_access(cgh); + auto sum = sum_buf.template get_access(cgh); auto out = - output_buf.template get_access(cgh); + output_buf.template get_access(cgh); cgh.parallel_for(range<1>(N), [=](item<1> it) { int gid = it.get_id(0); auto atm = AtomicRef < T, @@ -164,9 +164,9 @@ void add_plus_equal_test(queue q, size_t N) { buffer output_buf(output.data(), output.size()); q.submit([&](handler &cgh) { - auto sum = sum_buf.template get_access(cgh); + auto sum = sum_buf.template get_access(cgh); auto out = - output_buf.template get_access(cgh); + output_buf.template get_access(cgh); cgh.parallel_for(range<1>(N), [=](item<1> it) { int gid = it.get_id(0); auto atm = AtomicRef < T, @@ -247,9 +247,9 @@ void add_pre_inc_test(queue q, size_t N) { buffer output_buf(output.data(), output.size()); q.submit([&](handler &cgh) { - auto sum = sum_buf.template get_access(cgh); + auto sum = sum_buf.template get_access(cgh); auto out = - output_buf.template get_access(cgh); + output_buf.template get_access(cgh); cgh.parallel_for(range<1>(N), [=](item<1> it) { int gid = it.get_id(0); auto atm = AtomicRef < T, @@ -329,9 +329,9 @@ void add_post_inc_test(queue q, size_t N) { buffer output_buf(output.data(), output.size()); q.submit([&](handler &cgh) { - auto sum = sum_buf.template get_access(cgh); + auto sum = sum_buf.template get_access(cgh); auto out = - output_buf.template get_access(cgh); + output_buf.template get_access(cgh); cgh.parallel_for(range<1>(N), [=](item<1> it) { int gid = it.get_id(0); auto atm = AtomicRef < T, diff --git a/sycl/test-e2e/AtomicRef/atomic_memory_order_acq_rel.cpp b/sycl/test-e2e/AtomicRef/atomic_memory_order_acq_rel.cpp index 495348aa7c5e4..70814a9227966 100644 --- a/sycl/test-e2e/AtomicRef/atomic_memory_order_acq_rel.cpp +++ b/sycl/test-e2e/AtomicRef/atomic_memory_order_acq_rel.cpp @@ -22,9 +22,8 @@ template void test_acquire_global() { buffer val_buf(val, 2); q.submit([&](handler &cgh) { - auto error = - error_buf.template get_access(cgh); - auto val = val_buf.template get_access(cgh); + auto error = error_buf.template get_access(cgh); + auto val = val_buf.template get_access(cgh); cgh.parallel_for(range<1>(N_items), [=](item<1> it) { volatile int *val_p = val.get_multi_ptr().get(); auto atm0 = @@ -66,8 +65,7 @@ template void test_acquire_local() { buffer val_buf(val, 2); q.submit([&](handler &cgh) { - auto error = - error_buf.template get_access(cgh); + auto error = error_buf.template get_access(cgh); local_accessor val(2, cgh); cgh.parallel_for( nd_range<1>(global_size, local_size), [=](nd_item<1> it) { @@ -114,9 +112,8 @@ template void test_release_global() { buffer val_buf(val, 2); q.submit([&](handler &cgh) { - auto error = - error_buf.template get_access(cgh); - auto val = val_buf.template get_access(cgh); + auto error = error_buf.template get_access(cgh); + auto val = val_buf.template get_access(cgh); cgh.parallel_for(range<1>(N_items), [=](item<1> it) { volatile int *val_p = val.get_multi_ptr().get(); auto atm0 = @@ -158,8 +155,7 @@ template void test_release_local() { buffer val_buf(val, 2); q.submit([&](handler &cgh) { - auto error = - error_buf.template get_access(cgh); + auto error = error_buf.template get_access(cgh); local_accessor val(2, cgh); cgh.parallel_for( nd_range<1>(global_size, local_size), [=](nd_item<1> it) { diff --git a/sycl/test-e2e/AtomicRef/exchange.h b/sycl/test-e2e/AtomicRef/exchange.h index fd2ec48433fcf..1d4400b6538fb 100644 --- a/sycl/test-e2e/AtomicRef/exchange.h +++ b/sycl/test-e2e/AtomicRef/exchange.h @@ -32,9 +32,9 @@ void exchange_local_test(queue q, size_t N) { buffer cum_buf(&cum, 1); buffer output_buf(output.data(), output.size()); q.submit([&](handler &cgh) { - auto cum = cum_buf.template get_access(cgh); + auto cum = cum_buf.template get_access(cgh); auto out = - output_buf.template get_access(cgh); + output_buf.template get_access(cgh); local_accessor loc(1, cgh); cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) { @@ -79,10 +79,9 @@ void exchange_global_test(queue q, size_t N) { buffer output_buf(output.data(), output.size()); q.submit([&](handler &cgh) { - auto exc = - exchange_buf.template get_access(cgh); + auto exc = exchange_buf.template get_access(cgh); auto out = - output_buf.template get_access(cgh); + output_buf.template get_access(cgh); cgh.parallel_for(range<1>(N), [=](item<1> it) { size_t gid = it.get_id(0); auto atm = AtomicRef < T, diff --git a/sycl/test-e2e/AtomicRef/min.h b/sycl/test-e2e/AtomicRef/min.h index fd623d4a3d155..f82d0bf33d4d0 100644 --- a/sycl/test-e2e/AtomicRef/min.h +++ b/sycl/test-e2e/AtomicRef/min.h @@ -32,9 +32,9 @@ void min_local_test(queue q, size_t N) { buffer cum_buf(&cum, 1); buffer output_buf(output.data(), output.size()); q.submit([&](handler &cgh) { - auto cum = cum_buf.template get_access(cgh); + auto cum = cum_buf.template get_access(cgh); auto out = - output_buf.template get_access(cgh); + output_buf.template get_access(cgh); local_accessor loc(1, cgh); cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) { @@ -83,9 +83,9 @@ void min_global_test(queue q, size_t N) { buffer output_buf(output.data(), output.size()); q.submit([&](handler &cgh) { - auto val = val_buf.template get_access(cgh); + auto val = val_buf.template get_access(cgh); auto out = - output_buf.template get_access(cgh); + output_buf.template get_access(cgh); cgh.parallel_for(range<1>(N), [=](item<1> it) { int gid = it.get_id(0); auto atm = AtomicRef < T, diff --git a/sycl/test-e2e/AtomicRef/sub.h b/sycl/test-e2e/AtomicRef/sub.h index ab5ef47efe931..3fb6fc8731066 100644 --- a/sycl/test-e2e/AtomicRef/sub.h +++ b/sycl/test-e2e/AtomicRef/sub.h @@ -32,9 +32,9 @@ void sub_fetch_local_test(queue q, size_t N) { buffer sum_buf(&sum, 1); buffer output_buf(output.data(), output.size()); q.submit([&](handler &cgh) { - auto sum = sum_buf.template get_access(cgh); + auto sum = sum_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) { @@ -82,9 +82,9 @@ void sub_fetch_test(queue q, size_t N) { buffer output_buf(output.data(), output.size()); q.submit([&](handler &cgh) { - auto val = val_buf.template get_access(cgh); + auto val = val_buf.template get_access(cgh); auto out = - output_buf.template get_access(cgh); + output_buf.template get_access(cgh); cgh.parallel_for(range<1>(N), [=](item<1> it) { int gid = it.get_id(0); auto atm = AtomicRef < T, @@ -165,9 +165,9 @@ void sub_minus_equal_test(queue q, size_t N) { buffer output_buf(output.data(), output.size()); q.submit([&](handler &cgh) { - auto val = val_buf.template get_access(cgh); + auto val = val_buf.template get_access(cgh); auto out = - output_buf.template get_access(cgh); + output_buf.template get_access(cgh); cgh.parallel_for(range<1>(N), [=](item<1> it) { int gid = it.get_id(0); auto atm = AtomicRef < T, @@ -248,9 +248,9 @@ void sub_pre_dec_test(queue q, size_t N) { buffer output_buf(output.data(), output.size()); q.submit([&](handler &cgh) { - auto val = val_buf.template get_access(cgh); + auto val = val_buf.template get_access(cgh); auto out = - output_buf.template get_access(cgh); + output_buf.template get_access(cgh); cgh.parallel_for(range<1>(N), [=](item<1> it) { int gid = it.get_id(0); auto atm = AtomicRef < T, @@ -331,9 +331,9 @@ void sub_post_dec_test(queue q, size_t N) { buffer output_buf(output.data(), output.size()); q.submit([&](handler &cgh) { - auto val = val_buf.template get_access(cgh); + auto val = val_buf.template get_access(cgh); auto out = - output_buf.template get_access(cgh); + output_buf.template get_access(cgh); cgh.parallel_for(range<1>(N), [=](item<1> it) { int gid = it.get_id(0); auto atm = AtomicRef < T, diff --git a/sycl/test-e2e/Basic/accessor/accessor.cpp b/sycl/test-e2e/Basic/accessor/accessor.cpp index b6b7036f1a558..3de1a02daae10 100644 --- a/sycl/test-e2e/Basic/accessor/accessor.cpp +++ b/sycl/test-e2e/Basic/accessor/accessor.cpp @@ -35,7 +35,7 @@ struct IdxID3 { template using AccAlias = - sycl::accessor; + sycl::accessor; template struct InheritedAccessor : public AccAlias { @@ -67,9 +67,9 @@ template struct Wrapper3 { Wrapper2 w2; }; -using ResAccT = sycl::accessor; -using AccT = sycl::accessor; -using AccCT = sycl::accessor; +using ResAccT = sycl::accessor; +using AccT = sycl::accessor; +using AccCT = sycl::accessor; void implicit_conversion(const AccCT &acc, const ResAccT &res_acc) { auto v = acc[0]; @@ -188,7 +188,7 @@ void testLocalAccIters(std::vector &vec, bool testConstIter = false, sycl::queue queue; sycl::buffer buf(vec.data(), vec.size()); queue.submit([&](sycl::handler &cgh) { - auto globAcc = buf.get_access(cgh); + auto globAcc = buf.get_access(cgh); if (test2D) { sycl::local_accessor locAcc(sycl::range<2>{2, 16}, cgh); testLocalAccItersImpl(cgh, globAcc, locAcc, testConstIter); @@ -268,7 +268,7 @@ int main() { {sycl::property::buffer::use_host_ptr()}); Queue.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); assert(!acc.is_placeholder()); assert(acc.byte_size() == sizeof(int)); assert(acc.size() == 1); @@ -289,7 +289,7 @@ int main() { {sycl::property::buffer::use_host_ptr()}); Queue.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.parallel_for(Range, [=](sycl::item<2> itemID) { acc[itemID.get_id(0)][itemID.get_id(1)] += itemID.get_linear_id(); }); @@ -316,7 +316,7 @@ int main() { { sycl::buffer Buf(Arr, 1); queue.submit([&](sycl::handler &cgh) { - auto acc = sycl::accessor(range, cgh); cgh.parallel_for( sycl::nd_range<2>{range, range}, [=](sycl::nd_item<2>) { @@ -337,7 +337,7 @@ int main() { {sycl::property::buffer::use_host_ptr()}); Queue.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.parallel_for(Range, [=](sycl::item<3> itemID) { acc[itemID.get_id(0)][itemID.get_id(1)][itemID.get_id(2)] += itemID.get_linear_id(); @@ -366,7 +366,7 @@ int main() { constexpr int dims = 1; using data_loc = int; - constexpr auto mode_loc = sycl::access::mode::read_write; + constexpr auto mode_loc = sycl::access_mode::read_write; constexpr auto target_loc = sycl::target::local; const auto range_loc = sycl::range<1>(1); @@ -391,7 +391,7 @@ int main() { sycl::buffer buf(sycl::range<1>(3)); Queue.submit([&](sycl::handler &cgh) { - auto dev_acc = buf.get_access(cgh); + auto dev_acc = buf.get_access(cgh); cgh.parallel_for( sycl::range<1>{3}, [=](sycl::id<1> index) { dev_acc[index] = 42; }); @@ -414,7 +414,7 @@ int main() { sycl::buffer buf(sycl::range<1>(3)); Queue.submit([&](sycl::handler &cgh) { - auto dev_acc = buf.get_access(cgh); + auto dev_acc = buf.get_access(cgh); cgh.parallel_for( sycl::range<1>{3}, [=](sycl::id<1> index) { dev_acc[index] = 42; }); @@ -435,7 +435,7 @@ int main() { sycl::buffer buf((int *)array, sycl::range<1>(10), {sycl::property::buffer::use_host_ptr()}); queue.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); auto acc_wrapped = AccWrapper{acc}; cgh.parallel_for( sycl::range<1>(buf.size()), [=](sycl::item<1> it) { @@ -466,8 +466,8 @@ int main() { sycl::buffer buf2((int *)array2, sycl::range<1>(10), {sycl::property::buffer::use_host_ptr()}); queue.submit([&](sycl::handler &cgh) { - auto acc1 = buf1.get_access(cgh); - auto acc2 = buf2.get_access(cgh); + auto acc1 = buf1.get_access(cgh); + auto acc2 = buf2.get_access(cgh); auto acc_wrapped = AccsWrapper{10, acc1, 5, acc2}; cgh.parallel_for( @@ -503,7 +503,7 @@ int main() { sycl::buffer buf((int *)array, sycl::range<1>(10), {sycl::property::buffer::use_host_ptr()}); queue.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); auto acc_wrapped = AccWrapper{acc}; Wrapper1 wr1; auto wr2 = Wrapper2{wr1, acc_wrapped}; @@ -533,8 +533,8 @@ int main() { sycl::buffer buf(array, sycl::range<1>(3)); queue.submit([&](sycl::handler &cgh) { - auto acc1 = buf.get_access(cgh); - auto acc2 = buf.get_access(cgh); + auto acc1 = buf.get_access(cgh); + auto acc2 = buf.get_access(cgh); cgh.parallel_for( sycl::range<1>{3}, @@ -559,7 +559,7 @@ int main() { sycl::buffer b(&data, sycl::range<1>(1)); sycl::queue queue; queue.submit([&](sycl::handler &cgh) { - sycl::accessor B(b, cgh); cgh.single_task([=]() { @@ -591,13 +591,13 @@ int main() { sycl::queue queue; queue.submit([&](sycl::handler &cgh) { - sycl::accessor acc1(buf1, cgh); - sycl::accessor acc2(buf2, cgh); - sycl::accessor acc3(buf3, cgh, sycl::range<1>(1)); cgh.single_task([=]() { @@ -607,11 +607,11 @@ int main() { }); }); - sycl::accessor + sycl::accessor acc4(buf1); - sycl::accessor + sycl::accessor acc5(buf2); - sycl::accessor + sycl::accessor acc6(buf3, sycl::range<1>(1)); assert(acc4 == 2); @@ -631,10 +631,9 @@ int main() { sycl::queue queue; queue.submit([&](sycl::handler &cgh) { - sycl::accessor + sycl::accessor D(d, cgh); - sycl::accessor C(c, cgh); @@ -661,10 +660,10 @@ int main() { sycl::buffer d(&data, sycl::range<1>(1)); sycl::buffer c(&cnst, sycl::range<1>(1)); - sycl::accessor D(d); - sycl::accessor C(c); @@ -695,7 +694,7 @@ int main() { sycl::range<1> r(4); sycl::buffer b(r); try { - sycl::accessor acc(b); @@ -726,7 +725,7 @@ int main() { sycl::range<1> r(4); sycl::buffer b(r); try { - using AccT = sycl::accessor; AccT acc(b); @@ -760,7 +759,7 @@ int main() { sycl::range<1> r(4); sycl::buffer b(r); try { - using AccT = sycl::accessor; AccT acc(b); @@ -797,7 +796,7 @@ int main() { sycl::range<1> r(4); sycl::buffer b(r); try { - using AccT = sycl::accessor; AccT acc(b); @@ -852,10 +851,9 @@ int main() { sycl::queue queue; queue.submit([&](sycl::handler &cgh) { - sycl::accessor + sycl::accessor AccA(A, cgh); - sycl::accessor AccB(B, cgh); InheritedAccessor AccC(C, cgh); @@ -888,7 +886,7 @@ int main() { { sycl::buffer buf(array, sycl::range<1>(3)); q.submit([&](sycl::handler &h) { - auto acc = buf.get_access(h); + auto acc = buf.get_access(h); h.parallel_for(3, [=](sycl::id<1> i) { for (int j = 0; j < 3; ++j) { acc[i][j] = j + i * 10; @@ -925,7 +923,7 @@ int main() { // illegal ranges try { - auto acc = b.get_access(illegalR, offset); + auto acc = b.get_access(illegalR, offset); assert(false && "operation should not have succeeded"); } catch (sycl::exception &e) { assert(e.code() == sycl::errc::invalid && "errc should be errc::invalid"); @@ -933,7 +931,7 @@ int main() { try { sycl::queue q; q.submit([&](sycl::handler &cgh) { - auto acc = b.get_access(cgh, illegalR); + auto acc = b.get_access(cgh, illegalR); }); q.wait_and_throw(); assert(false && @@ -1174,7 +1172,7 @@ int main() { sycl::queue Queue; int Data[] = {32, 32}; - using HostTaskAcc = sycl::accessor; // Explicit block to prompt copy-back to Data @@ -1455,7 +1453,7 @@ int main() { queue .submit([&](sycl::handler &cgh) { auto B = - IBuf.template get_access(cgh); + IBuf.template get_access(cgh); cgh.single_task([=]() { for (size_t I = 0; I < B.size(); ++I) @@ -1530,7 +1528,7 @@ int main() { // default constructed accessor can be passed to a kernel (2). { - using AccT = sycl::accessor; + using AccT = sycl::accessor; AccT acc; assert(acc.empty()); sycl::queue q; diff --git a/sycl/test-e2e/Basic/buffer/buffer.cpp b/sycl/test-e2e/Basic/buffer/buffer.cpp index 6444d6f9867f6..f263c5c98ea31 100644 --- a/sycl/test-e2e/Basic/buffer/buffer.cpp +++ b/sycl/test-e2e/Basic/buffer/buffer.cpp @@ -33,7 +33,7 @@ int main() { buffer b(data1, range<1>(10), {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(range<1>{10}, [=](id<1> index) { B[index] = 0; }); }); @@ -49,7 +49,7 @@ int main() { buffer b(data1.data(), range<1>(10)); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = b.get_access(cgh); + auto B = b.get_access(cgh); cgh.parallel_for(range<1>{10}, [=](id<1> index) { B[index] = 0; }); }); @@ -72,9 +72,9 @@ int main() { buffer e(res.data(), range<1>(bufsSize)); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = b.get_access(cgh); - auto D = d.get_access(cgh); - auto E = e.get_access(cgh); + auto B = b.get_access(cgh); + auto D = d.get_access(cgh); + auto E = e.get_access(cgh); cgh.parallel_for(range<1>{bufsSize}, [=](id<1> index) { B[index]++; D[index] = B[index] + B[index] + 1; @@ -162,7 +162,7 @@ int main() { Buffer.set_final_data(nullptr); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = Buffer.get_access(cgh); + auto B = Buffer.get_access(cgh); cgh.parallel_for(range<1>{10}, [=](id<1> index) { B[index] = 1; }); }); @@ -177,7 +177,7 @@ int main() { Buffer.set_final_data(result); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = Buffer.get_access(cgh); + auto B = Buffer.get_access(cgh); cgh.parallel_for(range<1>{10}, [=](id<1> index) { B[index] = 1; }); }); @@ -193,7 +193,7 @@ int main() { Buffer.set_final_data(resultWeak); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = Buffer.get_access(cgh); + auto B = Buffer.get_access(cgh); cgh.parallel_for(range<1>{10}, [=](id<1> index) { B[index] = 1; }); }); @@ -209,7 +209,7 @@ int main() { Buffer.set_final_data(result); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = Buffer.get_access(cgh); + auto B = Buffer.get_access(cgh); cgh.parallel_for( range<1>{10}, [=](id<1> index) { B[index] = 1; }); }); @@ -230,7 +230,7 @@ int main() { queue myQueue; resultShared.reset(); myQueue.submit([&](handler &cgh) { - auto B = Buffer.get_access(cgh); + auto B = Buffer.get_access(cgh); cgh.parallel_for( range<1>{10}, [=](id<1> index) { B[index] = 1; }); }); @@ -246,7 +246,7 @@ int main() { Buffer.set_final_data(result.begin()); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = Buffer.get_access(cgh); + auto B = Buffer.get_access(cgh); cgh.parallel_for( range<1>{10}, [=](id<1> index) { B[index] = 1; }); }); @@ -262,7 +262,7 @@ int main() { Buffer.set_final_data(nullptr); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = Buffer.get_access(cgh); + auto B = Buffer.get_access(cgh); cgh.parallel_for( range<1>{10}, [=](id<1> index) { B[index] = 1; }); }); @@ -278,7 +278,7 @@ int main() { Buffer.set_final_data(result); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = Buffer.get_access(cgh); + auto B = Buffer.get_access(cgh); cgh.parallel_for( range<1>{10}, [=](id<1> index) { B[index] = 1; }); }); @@ -295,7 +295,7 @@ int main() { Buffer.set_final_data(resultWeak); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = Buffer.get_access(cgh); + auto B = Buffer.get_access(cgh); cgh.parallel_for( range<1>{10}, [=](id<1> index) { B[index] = 1; }); }); @@ -316,7 +316,7 @@ int main() { queue myQueue; resultShared.reset(); myQueue.submit([&](handler &cgh) { - auto B = Buffer.get_access(cgh); + auto B = Buffer.get_access(cgh); cgh.parallel_for( range<1>{10}, [=](id<1> index) { B[index] = 1; }); }); @@ -333,7 +333,7 @@ int main() { Buffer.set_final_data(result.begin()); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = Buffer.get_access(cgh); + auto B = Buffer.get_access(cgh); cgh.parallel_for( range<1>{10}, [=](id<1> index) { B[index] = 1; }); }); @@ -348,13 +348,13 @@ int main() { Buffer.set_final_data((int *)result); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = Buffer.get_access(cgh); + auto B = Buffer.get_access(cgh); cgh.parallel_for( range<2>{20, 20}, [=](id<2> index) { B[index] = 0; }); }); myQueue.submit([&](handler &cgh) { - auto B = Buffer.get_access(cgh); + auto B = Buffer.get_access(cgh); cgh.parallel_for( range<2>{10, 10}, [=](id<2> index) { B[index] = 1; }); }); @@ -388,13 +388,13 @@ int main() { Buffer.set_final_data((int *)result); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = Buffer.get_access(cgh); + auto B = Buffer.get_access(cgh); cgh.parallel_for( range<2>{20, 20}, [=](id<2> index) { B[index] = 0; }); }); myQueue.submit([&](handler &cgh) { - accessor B(Buffer, cgh, range<2>(10, 10), id<2>(10, 10)); cgh.parallel_for( @@ -429,7 +429,7 @@ int main() { buffer b(data1.begin() + 2, data1.begin() + 5); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = b.get_access(cgh); + auto B = b.get_access(cgh); cgh.parallel_for( range<1>{3}, [=](id<1> index) { B[index] = 20; }); }); @@ -448,7 +448,7 @@ int main() { buffer b(data1.cbegin() + 2, data1.cbegin() + 5); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = b.get_access(cgh); + auto B = b.get_access(cgh); cgh.parallel_for( range<1>{3}, [=](id<1> index) { B[index] = 20; }); }); @@ -467,7 +467,7 @@ int main() { b.set_final_data(data2.begin() + 2); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = b.get_access(cgh); + auto B = b.get_access(cgh); cgh.parallel_for( range<1>{3}, [=](id<1> index) { B[index] = 20; }); }); @@ -490,7 +490,7 @@ int main() { b.set_write_back(true); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = b.get_access(cgh); + auto B = b.get_access(cgh); cgh.parallel_for(range<1>{10}, [=](id<1> index) { B[index] = 0; }); }); @@ -552,11 +552,11 @@ int main() { Queue.submit([&](sycl::handler &cgh) { auto Accessor_bool = - buf_bool_shrd.get_access(cgh); + buf_bool_shrd.get_access(cgh); auto Accessor_int = - buf_int_shrd.get_access(cgh); + buf_int_shrd.get_access(cgh); auto Accessor_float = - buf_float_shrd.get_access(cgh); + buf_float_shrd.get_access(cgh); cgh.parallel_for(r, [=](sycl::id<1> WIid) { Accessor_bool[WIid] = true; Accessor_int[WIid] = 3; @@ -584,7 +584,7 @@ int main() { b.set_write_back(false); queue myQueue; myQueue.submit([&](handler &cgh) { - auto B = b.get_access(cgh); + auto B = b.get_access(cgh); cgh.parallel_for(range<1>{10}, [=](id<1> index) { B[index] = 0; }); }); @@ -605,8 +605,8 @@ int main() { buffer b(data2.data(), range<1>(10)); 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) { A[index] = 0; }); }); @@ -623,10 +623,10 @@ int main() { { buffer a(data1.data(), range<1>(10)); buffer b(data2); - accessor A(a); - accessor B(b); queue myQueue; @@ -670,8 +670,8 @@ int main() { assert(AccA.size() == AccB.size()); } - auto AH0 = host_accessor(Buf_1); - auto BH0 = host_accessor(Buf_2); + auto AH0 = host_accessor(Buf_1); + auto BH0 = host_accessor(Buf_2); assert(AH0.byte_size() == sizeof(char)); assert(BH0.byte_size() == sizeof(char)); assert(AH0.size() == 1); @@ -680,10 +680,10 @@ int main() { queue Queue; Queue.submit([&](handler &CGH) { auto AK0 = - accessor( + accessor( Buf_1, CGH); auto BK0 = - accessor( + accessor( Buf_2, CGH); assert(AK0.byte_size() == sizeof(char)); assert(BK0.byte_size() == sizeof(char)); diff --git a/sycl/test-e2e/Basic/buffer/buffer_release.cpp b/sycl/test-e2e/Basic/buffer/buffer_release.cpp index 92f48f2027cd3..3e0cef483d5d1 100644 --- a/sycl/test-e2e/Basic/buffer/buffer_release.cpp +++ b/sycl/test-e2e/Basic/buffer/buffer_release.cpp @@ -18,7 +18,7 @@ void test() { HostAcc.reset(new host_accessor(Buffer)); // Host accessor should block kernel execution Q.submit([&](handler &CGH) { - auto Acc = Buffer.template get_access(CGH); + auto Acc = Buffer.template get_access(CGH); CGH.parallel_for(BUFFER_SIZE, [=](item<1> Id) { Acc[Id] = 0; }); }); diff --git a/sycl/test-e2e/Basic/buffer/subbuffer.cpp b/sycl/test-e2e/Basic/buffer/subbuffer.cpp index 0ea777a23ecad..d588ddc39c837 100644 --- a/sycl/test-e2e/Basic/buffer/subbuffer.cpp +++ b/sycl/test-e2e/Basic/buffer/subbuffer.cpp @@ -41,7 +41,7 @@ void checkHostAccessor(sycl::queue &q) { } q.submit([&](sycl::handler &cgh) { - auto acc = subbuf.get_access(cgh); + auto acc = subbuf.get_access(cgh); cgh.parallel_for(sycl::range<1>(10), [=](sycl::id<1> i) { acc[i] *= -10; }); }); @@ -85,7 +85,7 @@ void check1DSubBuffer(sycl::queue &q) { q.submit([&](sycl::handler &cgh) { // accessor starts at the third element of the subbuffer // and can read for 7 more (ie to the end of the subbuffer) - auto acc = subbuf.get_access( + auto acc = subbuf.get_access( cgh, sycl::range<1>(subbuffer_access_range), sycl::id<1>(offset_inside_subbuf)); // subrange is made negative. ( 32 33 34 -35 -36 -37 -38 -39 -40 -41) @@ -95,8 +95,8 @@ void check1DSubBuffer(sycl::queue &q) { // copy results of last operation back to buf2/vec2 q.submit([&](sycl::handler &cgh) { - auto acc_sub = subbuf.get_access(cgh); - auto acc_buf = buf2.get_access(cgh); + auto acc_sub = subbuf.get_access(cgh); + auto acc_buf = buf2.get_access(cgh); cgh.parallel_for( subbuf.get_range(), [=](sycl::id<1> i) { acc_buf[i] = acc_sub[i]; }); }); @@ -106,7 +106,7 @@ void check1DSubBuffer(sycl::queue &q) { // (..29 30 31 | 320 330 340 -350 -360 -370 -380 -390 -400 -410 | 42 43 44 // ...) q.submit([&](sycl::handler &cgh) { - auto acc_sub = subbuf.get_access( + auto acc_sub = subbuf.get_access( cgh, sycl::range<1>(subbuf_size)); cgh.parallel_for( sycl::range<1>(subbuf_size), @@ -248,7 +248,6 @@ void checkExceptions() { } void copyBlock() { - using typename sycl::access::mode; using buffer = sycl::buffer; auto CopyF = [](buffer &Buffer, buffer &Block, size_t Idx, size_t BlockSize) { @@ -312,13 +311,13 @@ void checkMultipleContexts() { sycl::buffer subbuf2(buf, sycl::id<1>(N / 2), sycl::range<1>(N / 2)); queue1.submit([&](sycl::handler &cgh) { - auto bufacc = subbuf1.get_access(cgh); + auto bufacc = subbuf1.get_access(cgh); cgh.parallel_for( sycl::range<1>(N / 2), [=](sycl::id<1> idx) { bufacc[idx[0]] = 1; }); }); queue2.submit([&](sycl::handler &cgh) { - auto bufacc = subbuf2.get_access(cgh); + auto bufacc = subbuf2.get_access(cgh); cgh.parallel_for( sycl::range<1>(N / 2), [=](sycl::id<1> idx) { bufacc[idx[0]] = 2; }); }); @@ -331,7 +330,7 @@ void checkMultipleContexts() { sycl::buffer subbuf1(buf, sycl::id<1>(N / 2), sycl::range<1>(N / 2)); queue1.submit([&](sycl::handler &cgh) { - auto bufacc = subbuf1.get_access(cgh); + auto bufacc = subbuf1.get_access(cgh); cgh.parallel_for( sycl::range<1>(N / 2), [=](sycl::id<1> idx) { bufacc[idx[0]] = -1; }); }); diff --git a/sycl/test-e2e/Basic/compare_exchange_strong.cpp b/sycl/test-e2e/Basic/compare_exchange_strong.cpp index 92e0a30f04d98..fca6d96079cff 100644 --- a/sycl/test-e2e/Basic/compare_exchange_strong.cpp +++ b/sycl/test-e2e/Basic/compare_exchange_strong.cpp @@ -16,8 +16,8 @@ int main() { buffer buf(&data, range<1>(1)); testQueue.submit([&](handler &cgh) { - auto globAcc = buf.template get_access(cgh); - auto resultAcc = resultBuf.template get_access(cgh); + auto globAcc = buf.template get_access(cgh); + auto resultAcc = resultBuf.template get_access(cgh); cgh.single_task([=]() { auto a = globAcc[0]; char result = 0; diff --git a/sycl/test-e2e/Basic/event_profiling_info.cpp b/sycl/test-e2e/Basic/event_profiling_info.cpp index 8339491e93ad7..c4c4c3d894c47 100644 --- a/sycl/test-e2e/Basic/event_profiling_info.cpp +++ b/sycl/test-e2e/Basic/event_profiling_info.cpp @@ -53,9 +53,9 @@ int main() { // buffer copy queue copyQueue{Dev, sycl::property::queue::enable_profiling()}; event copyEvent = copyQueue.submit([&](sycl::handler &Cgh) { - accessor AccessorFrom( + accessor AccessorFrom( BufferFrom, Cgh, range<1>(Size)); - accessor AccessorTo( + accessor AccessorTo( BufferTo, Cgh, range<1>(Size)); Cgh.copy(AccessorFrom, AccessorTo); }); diff --git a/sycl/test-e2e/Basic/fill_accessor_ur.cpp b/sycl/test-e2e/Basic/fill_accessor_ur.cpp index 8f938aa77d815..247a041d46d4b 100644 --- a/sycl/test-e2e/Basic/fill_accessor_ur.cpp +++ b/sycl/test-e2e/Basic/fill_accessor_ur.cpp @@ -20,7 +20,7 @@ void testFill_Buffer1D() { sycl::queue q; std::cout << "start testFill_Buffer1D" << std::endl; q.submit([&](sycl::handler &cgh) { - auto acc1D = buffer_1D.get_access(cgh); + auto acc1D = buffer_1D.get_access(cgh); // should stage urEnqueueMemBufferFill cgh.fill(acc1D, float{1}); }); @@ -29,7 +29,7 @@ void testFill_Buffer1D() { std::cout << "start testFill_Buffer1D -- OFFSET" << std::endl; q.submit([&](sycl::handler &cgh) { auto acc1DOffset = - buffer_1D.get_access(cgh, {4}, {2}); + buffer_1D.get_access(cgh, {4}, {2}); // despite being offset, should stage urEnqueueMemBufferFill cgh.fill(acc1DOffset, float{2}); }); @@ -50,7 +50,7 @@ void testFill_Buffer2D() { sycl::queue q; std::cout << "start testFill_Buffer2D" << std::endl; q.submit([&](sycl::handler &cgh) { - auto acc2D = buffer_2D.get_access(cgh); + auto acc2D = buffer_2D.get_access(cgh); // should stage urEnqueueMemBufferFill cgh.fill(acc2D, float{3}); }); @@ -59,7 +59,7 @@ void testFill_Buffer2D() { std::cout << "start testFill_Buffer2D -- OFFSET" << std::endl; q.submit([&](sycl::handler &cgh) { auto acc2D = - buffer_2D.get_access(cgh, {8, 12}, {2, 2}); + buffer_2D.get_access(cgh, {8, 12}, {2, 2}); // "ranged accessor" will have to be handled by custom kernel: // urEnqueueKernelLaunchWithArgsExp cgh.fill(acc2D, float{4}); @@ -81,7 +81,7 @@ void testFill_Buffer3D() { sycl::queue q; std::cout << "start testFill_Buffer3D" << std::endl; q.submit([&](sycl::handler &cgh) { - auto acc3D = buffer_3D.get_access(cgh); + auto acc3D = buffer_3D.get_access(cgh); // should stage urEnqueueMemBufferFill cgh.fill(acc3D, float{5}); }); @@ -89,7 +89,7 @@ void testFill_Buffer3D() { std::cout << "start testFill_Buffer3D -- OFFSET" << std::endl; q.submit([&](sycl::handler &cgh) { - auto acc3D = buffer_3D.get_access( + auto acc3D = buffer_3D.get_access( cgh, {4, 8, 12}, {3, 3, 3}); // "ranged accessor" will have to be handled by custom kernel: // urEnqueueKernelLaunchWithArgsExp @@ -113,7 +113,7 @@ void testFill_ZeroDim() { sycl::queue q; std::cout << "start testFill_ZeroDim" << std::endl; q.submit([&](sycl::handler &cgh) { - sycl::accessor Acc0(Buffer, cgh); + sycl::accessor Acc0(Buffer, cgh); cgh.fill(Acc0, float{1}); }); q.wait(); diff --git a/sycl/test-e2e/Basic/get_backend.cpp b/sycl/test-e2e/Basic/get_backend.cpp index fbb3090ba7f4f..81d521c5286d6 100644 --- a/sycl/test-e2e/Basic/get_backend.cpp +++ b/sycl/test-e2e/Basic/get_backend.cpp @@ -57,7 +57,7 @@ int main() { buffer buf{range<1>(1)}; event e = q.submit([&](handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.fill(acc, 0); }); if (e.get_backend() != plt.get_backend()) { diff --git a/sycl/test-e2e/Basic/half_builtins.cpp b/sycl/test-e2e/Basic/half_builtins.cpp index 2fcdd1c4b4b68..4ed22ec856251 100644 --- a/sycl/test-e2e/Basic/half_builtins.cpp +++ b/sycl/test-e2e/Basic/half_builtins.cpp @@ -178,10 +178,10 @@ int main() { buffer c_buf(&c, 1); buffer err_buf(&err, 1); q.submit([&](handler &cgh) { - auto A = a_buf.get_access(cgh); - auto B = b_buf.get_access(cgh); - auto C = c_buf.get_access(cgh); - auto err = err_buf.get_access(cgh); + auto A = a_buf.get_access(cgh); + auto B = b_buf.get_access(cgh); + auto C = c_buf.get_access(cgh); + auto err = err_buf.get_access(cgh); cgh.parallel_for(SZ_max, [=](item<1> index) { size_t i = index.get_id(0); TEST_BUILTIN_1(sycl::fabs); diff --git a/sycl/test-e2e/Basic/handler/handler_generic_integral_lambda.cpp b/sycl/test-e2e/Basic/handler/handler_generic_integral_lambda.cpp index b68022ff45141..e91bb4873bd61 100644 --- a/sycl/test-e2e/Basic/handler/handler_generic_integral_lambda.cpp +++ b/sycl/test-e2e/Basic/handler/handler_generic_integral_lambda.cpp @@ -23,9 +23,9 @@ int main() { sycl::buffer buf(data, sycl::range<1>(length)); sycl::queue q; q.submit([&](sycl::handler &cgh) { - sycl::accessor - acc(buf.get_access(cgh)); + acc(buf.get_access(cgh)); cgh.parallel_for( length, [=](auto item) { acc[item.get_id()] = item; }); }); @@ -41,9 +41,9 @@ int main() { sycl::buffer buf(data, sycl::range<1>(length)); sycl::queue q; q.submit([&](sycl::handler &cgh) { - sycl::accessor - acc(buf.get_access(cgh)); + acc(buf.get_access(cgh)); cgh.parallel_for( length, [=](std::size_t item) { acc[item] = item; }); }); @@ -59,9 +59,9 @@ int main() { sycl::buffer buf(data, sycl::range<1>(length)); sycl::queue q; q.submit([&](sycl::handler &cgh) { - sycl::accessor - acc(buf.get_access(cgh)); + acc(buf.get_access(cgh)); cgh.parallel_for(length, [=](int item) { acc[item] = item; }); }); diff --git a/sycl/test-e2e/Basic/image/image.cpp b/sycl/test-e2e/Basic/image/image.cpp index 79734d985df3b..b9a66a42e3400 100644 --- a/sycl/test-e2e/Basic/image/image.cpp +++ b/sycl/test-e2e/Basic/image/image.cpp @@ -29,8 +29,8 @@ 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 SYCLRead = sycl::access::mode::read; - constexpr auto SYCLWrite = sycl::access::mode::write; + constexpr auto SYCLRead = sycl::access_mode::read; + constexpr auto SYCLWrite = sycl::access_mode::write; const sycl::range<2> Img1Size(4, 4); const sycl::range<2> Img2Size(4, 4); @@ -95,7 +95,7 @@ int main() { constexpr int dims = 1; using data_img = sycl::float4; - constexpr auto mode_img = sycl::access::mode::read; + constexpr auto mode_img = sycl::access_mode::read; constexpr auto target_img = sycl::target::image; const auto range_img = sycl::range(3); auto image = sycl::image(sycl::image_channel_order::rgba, @@ -126,7 +126,7 @@ int main() { sycl::queue Q; Q.submit([&](sycl::handler &CGH) { auto ImgAcc = - Img.get_access(CGH); + Img.get_access(CGH); sycl::nd_range<2> Rng(sycl::range<2>(NX, NY), sycl::range<2>(16, 16)); diff --git a/sycl/test-e2e/Basic/image/image_accessor_readwrite.cpp b/sycl/test-e2e/Basic/image/image_accessor_readwrite.cpp index 220d5e31b3f0f..a13e493385656 100644 --- a/sycl/test-e2e/Basic/image/image_accessor_readwrite.cpp +++ b/sycl/test-e2e/Basic/image/image_accessor_readwrite.cpp @@ -108,7 +108,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); }); @@ -128,8 +128,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_max_size.cpp b/sycl/test-e2e/Basic/image/image_max_size.cpp index a2c0e22a283f0..24e4c23ad50a2 100644 --- a/sycl/test-e2e/Basic/image/image_max_size.cpp +++ b/sycl/test-e2e/Basic/image/image_max_size.cpp @@ -56,8 +56,8 @@ bool testND(queue &Q, size_t XSize, size_t YSize, size_t ZSize = 1) { image_channel_type::unsigned_int32, ImgRange); Q.submit([&](handler &CGH) { - auto AAcc = ImgA.template get_access(CGH); - auto BAcc = ImgB.template get_access(CGH); + auto AAcc = ImgA.template get_access(CGH); + auto BAcc = ImgB.template get_access(CGH); CGH.parallel_for>( ImgRange, [=](id Id) { // Use int2 for 2D and int4 for 3D images. diff --git a/sycl/test-e2e/Basic/image/image_trace.cpp b/sycl/test-e2e/Basic/image/image_trace.cpp index 9437468d61a2f..7c76e8ed31ed7 100644 --- a/sycl/test-e2e/Basic/image/image_trace.cpp +++ b/sycl/test-e2e/Basic/image/image_trace.cpp @@ -33,8 +33,8 @@ 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 SYCLRead = sycl::access::mode::read; - constexpr auto SYCLWrite = sycl::access::mode::write; + constexpr auto SYCLRead = sycl::access_mode::read; + constexpr auto SYCLWrite = sycl::access_mode::write; const sycl::range<2> Img1Size(4, 4); const sycl::range<2> Img2Size(4, 4); diff --git a/sycl/test-e2e/Basic/interop/make_kernel_subdevice_l0.cpp b/sycl/test-e2e/Basic/interop/make_kernel_subdevice_l0.cpp index 587f2664d6d14..3cb393659fe01 100644 --- a/sycl/test-e2e/Basic/interop/make_kernel_subdevice_l0.cpp +++ b/sycl/test-e2e/Basic/interop/make_kernel_subdevice_l0.cpp @@ -133,7 +133,7 @@ int main() { } Q.submit([&](handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.set_args(acc); cgh.parallel_for(range<1>(1), K); }); diff --git a/sycl/test-e2e/Basic/kernel_info.cpp b/sycl/test-e2e/Basic/kernel_info.cpp index 73d5a6fd7514d..edc2c19b6bd89 100644 --- a/sycl/test-e2e/Basic/kernel_info.cpp +++ b/sycl/test-e2e/Basic/kernel_info.cpp @@ -49,7 +49,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.hpp b/sycl/test-e2e/Basic/multi_ptr.hpp index 812dab1652ae3..744d820a99d94 100644 --- a/sycl/test-e2e/Basic/multi_ptr.hpp +++ b/sycl/test-e2e/Basic/multi_ptr.hpp @@ -86,19 +86,19 @@ template void testMultPtr() { buffer bufferData_5(data_5, numOfItems); queue myQueue; myQueue.submit([&](handler &cgh) { - accessor accessorData_1(bufferData_1, cgh); - accessor accessorData_2(bufferData_2, cgh); - accessor accessorData_3(bufferData_3, cgh); - accessor accessorData_4(bufferData_4, cgh); - accessor accessorData_5(bufferData_5, cgh); local_accessor localAccessor(numOfItems, cgh); @@ -172,11 +172,11 @@ void testMultPtrArrowOperator() { buffer, 1> bufferData_3(data_3, numOfItems); 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); local_accessor, 1> accessorData_2(1, cgh); - accessor, 1, access::mode::read, access::target::device, + accessor, 1, access_mode::read, access::target::device, access::placeholder::false_t> accessorData_3(bufferData_3, cgh); diff --git a/sycl/test-e2e/Basic/multisource.cpp b/sycl/test-e2e/Basic/multisource.cpp index db7b4dbf39e74..6d6fb4815a6c9 100644 --- a/sycl/test-e2e/Basic/multisource.cpp +++ b/sycl/test-e2e/Basic/multisource.cpp @@ -30,7 +30,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; }); }); } @@ -43,9 +43,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/offset-accessor-get_pointer.cpp b/sycl/test-e2e/Basic/offset-accessor-get_pointer.cpp index 76a5135ff6563..01475ff22fb12 100644 --- a/sycl/test-e2e/Basic/offset-accessor-get_pointer.cpp +++ b/sycl/test-e2e/Basic/offset-accessor-get_pointer.cpp @@ -17,7 +17,7 @@ using namespace sycl; void test_across_ranges() { - constexpr auto r_w = access::mode::read_write; + constexpr auto r_w = access_mode::read_write; constexpr unsigned long width = 4; constexpr unsigned long count = width * width; constexpr unsigned long count3D = width * width * width; // 64 diff --git a/sycl/test-e2e/Basic/queue/queue_shortcut_functions.cpp b/sycl/test-e2e/Basic/queue/queue_shortcut_functions.cpp index 26eba4b57e28e..77dafe85907ec 100644 --- a/sycl/test-e2e/Basic/queue/queue_shortcut_functions.cpp +++ b/sycl/test-e2e/Basic/queue/queue_shortcut_functions.cpp @@ -28,7 +28,7 @@ int main() { { sycl::buffer DataBuffer(Data.data(), sycl::range<1>(Size)); - sycl::accessor> @@ -57,7 +57,7 @@ int main() { { sycl::buffer Buf(Data.data(), sycl::range<1>(Size)); - sycl::accessor> @@ -96,7 +96,7 @@ int main() { { sycl::buffer Buf(DataPtr, sycl::range<1>(Size)); - sycl::accessor> @@ -135,12 +135,12 @@ int main() { sycl::buffer RefBuf(ReferenceData.data(), sycl::range<1>(Size)); sycl::buffer DataBuf(Data.data(), sycl::range<1>(Size)); - sycl::accessor> RefAcc(RefBuf); - sycl::accessor> diff --git a/sycl/test-e2e/Basic/subdevice_pi.cpp b/sycl/test-e2e/Basic/subdevice_pi.cpp index 43595063abc1e..56fca44e0ab50 100644 --- a/sycl/test-e2e/Basic/subdevice_pi.cpp +++ b/sycl/test-e2e/Basic/subdevice_pi.cpp @@ -22,7 +22,7 @@ static void log_pi(const char *msg) { std::cout << msg << std::endl; } static void use_mem(buffer buf, queue q) { q.submit([&](handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.parallel_for(range<1>(buf.size()), [=](item<1> itemID) { acc[itemID] += 1; }); }); diff --git a/sycl/test-e2e/Basic/unused_pointer.cpp b/sycl/test-e2e/Basic/unused_pointer.cpp index 89e24720ff738..7e0f6f6549bde 100644 --- a/sycl/test-e2e/Basic/unused_pointer.cpp +++ b/sycl/test-e2e/Basic/unused_pointer.cpp @@ -30,7 +30,7 @@ int main(int argc, char **argv) { { sycl::buffer buf(&data, 1); queue.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.single_task([=]() { acc[0] = obj.data; }); }); } diff --git a/sycl/test-e2e/Basic/vector/operators.cpp b/sycl/test-e2e/Basic/vector/operators.cpp index d56cc6cdc8cb5..17a3789b2d3c0 100644 --- a/sycl/test-e2e/Basic/vector/operators.cpp +++ b/sycl/test-e2e/Basic/vector/operators.cpp @@ -49,7 +49,7 @@ int main() { { s::buffer Buf(&res, s::range<1>(1)); Queue.submit([&](s::handler &cgh) { - auto Acc = Buf.get_access(cgh); + auto Acc = Buf.get_access(cgh); cgh.single_task([=]() { vec_type vec1(2); vec_type vec2(2); @@ -69,7 +69,7 @@ int main() { { s::buffer Buf(&res, s::range<1>(1)); Queue.submit([&](s::handler &cgh) { - auto Acc = Buf.get_access(cgh); + auto Acc = Buf.get_access(cgh); cgh.single_task([=]() { vec_type vec(2); char rhs_scalar = 2; @@ -95,7 +95,7 @@ int main() { { s::buffer Buf(&res, s::range<1>(1)); Queue.submit([&](s::handler &cgh) { - auto Acc = Buf.get_access(cgh); + auto Acc = Buf.get_access(cgh); cgh.single_task([=]() { s::uint4 vec1(42, 42, 42, 0); s::uint4 vec2(0, 42, 42, 0); @@ -114,7 +114,7 @@ int main() { { s::buffer Buf(&res, s::range<1>(1)); Queue.submit([&](s::handler &cgh) { - auto Acc = Buf.get_access(cgh); + auto Acc = Buf.get_access(cgh); cgh.single_task([=]() { s::double4 vec1(0.5, 10.1, 10.2, 10.3); s::double4 vec2(10.5, 0.1, 0.2, 0.3); @@ -133,7 +133,7 @@ int main() { { s::buffer Buf(&res, s::range<1>(1)); Queue.submit([&](s::handler &cgh) { - auto Acc = Buf.get_access(cgh); + auto Acc = Buf.get_access(cgh); cgh.single_task([=]() { s::char4 vec1(0, 0, 42, 42); s::char4 vec2(42, 0, 0, -42); @@ -152,7 +152,7 @@ int main() { { s::buffer Buf(&res, s::range<1>(1)); Queue.submit([&](s::handler &cgh) { - auto Acc = Buf.get_access(cgh); + auto Acc = Buf.get_access(cgh); cgh.single_task([=]() { s::half4 vec1(0, 0, 42, 42); s::half4 vec2(42, 0, 0, -42); @@ -173,7 +173,7 @@ int main() { { s::buffer Buf(&res, s::range<1>(1)); Queue.submit([&](s::handler &cgh) { - auto Acc = Buf.get_access(cgh); + auto Acc = Buf.get_access(cgh); cgh.single_task([=]() { s::ulong4 vec1(0, 0, 42, 42); s::ulong4 vec2{42, 0, 0, 0}; @@ -192,7 +192,7 @@ int main() { { s::buffer Buf(&res, s::range<1>(1)); Queue.submit([&](s::handler &cgh) { - auto Acc = Buf.get_access(cgh); + auto Acc = Buf.get_access(cgh); cgh.single_task([=]() { s::ushort4 vec1(0, 0, 42, 42); s::ushort4 vec2{42, 0, 0, 42}; @@ -211,7 +211,7 @@ int main() { { s::buffer Buf(&res, s::range<1>(1)); Queue.submit([&](s::handler &cgh) { - auto Acc = Buf.get_access(cgh); + auto Acc = Buf.get_access(cgh); cgh.single_task([=]() { s::int4 vec1(0, 0, 42, 42); s::int4 vec2{42, 0, 0, 42}; @@ -230,7 +230,7 @@ int main() { { s::buffer Buf(&res, s::range<1>(1)); Queue.submit([&](s::handler &cgh) { - auto Acc = Buf.get_access(cgh); + auto Acc = Buf.get_access(cgh); cgh.single_task([=]() { s::int4 vec1(0, 0, 42, 42); s::int4 vec2{42, 0, 0, 42}; @@ -250,7 +250,7 @@ int main() { { s::buffer Buf(&res, s::range<1>(1)); Queue.submit([&](s::handler &cgh) { - auto Acc = Buf.get_access(cgh); + auto Acc = Buf.get_access(cgh); cgh.single_task([=]() { s::float4 vec1(4.5f, 0, 3.5f, -10.0f); Acc[0] = vec1.template as(); diff --git a/sycl/test-e2e/Config/kernel_from_file.cpp b/sycl/test-e2e/Config/kernel_from_file.cpp index c8b880e6c9ff5..d9182d72ecce2 100644 --- a/sycl/test-e2e/Config/kernel_from_file.cpp +++ b/sycl/test-e2e/Config/kernel_from_file.cpp @@ -26,7 +26,7 @@ int main(int argc, char **argv) { buffer buf(&data, range<1>(1)); event e = myQueue.submit([&](handler &cgh) { - auto ptr = buf.get_access(cgh); + auto ptr = buf.get_access(cgh); cgh.single_task([=]() { #ifdef USED_KERNEL diff --git a/sycl/test-e2e/DeprecatedFeatures/subbuffer_interop.cpp b/sycl/test-e2e/DeprecatedFeatures/subbuffer_interop.cpp index bb41c3e0ebe95..37d6a41ebcbbb 100644 --- a/sycl/test-e2e/DeprecatedFeatures/subbuffer_interop.cpp +++ b/sycl/test-e2e/DeprecatedFeatures/subbuffer_interop.cpp @@ -93,8 +93,8 @@ int main() { kernel TestKernel(clKernel, TestQueue.get_context()); TestQueue.submit([&](handler &cgh) { - auto a_acc = BufA.get_access(cgh); - auto b_acc = BufB.get_access(cgh); + auto a_acc = BufA.get_access(cgh); + auto b_acc = BufB.get_access(cgh); cgh.set_arg(0, a_acc); cgh.set_arg(1, b_acc); @@ -159,9 +159,9 @@ int main() { kernel TestKernel(clKernel, TestQueue.get_context()); TestQueue.submit([&](handler &cgh) { - auto a_acc = BufA.get_access(cgh); - auto b_acc = BufB.get_access(cgh); - auto c_acc = BufC.get_access(cgh); + auto a_acc = BufA.get_access(cgh); + auto b_acc = BufB.get_access(cgh); + auto c_acc = BufC.get_access(cgh); cgh.set_arg(0, a_acc); cgh.set_arg(1, b_acc); cgh.set_arg(2, c_acc); @@ -241,8 +241,8 @@ int main() { kernel TestKernel2(clKernel2, TestQueue.get_context()); TestQueue.submit([&](handler &cgh) { - auto a_acc = BufA.get_access(cgh); - auto b_acc = BufB.get_access(cgh); + auto a_acc = BufA.get_access(cgh); + auto b_acc = BufB.get_access(cgh); cgh.set_arg(0, a_acc); cgh.set_arg(1, b_acc); @@ -250,8 +250,8 @@ int main() { }); TestQueue.submit([&](handler &cgh) { - auto a_acc = BufA.get_access(cgh); - auto c_acc = BufC.get_access(cgh); + auto a_acc = BufA.get_access(cgh); + auto c_acc = BufC.get_access(cgh); cgh.set_arg(0, a_acc); cgh.set_arg(1, c_acc); @@ -314,13 +314,13 @@ int main() { // Test offsets { - auto host_acc = subbuf_copy->get_access(); + auto host_acc = subbuf_copy->get_access(); std::cout << "On host: offset = " << host_acc[0] << std::endl; assert(host_acc[0] == 256 && "Invalid subbuffer origin"); } Q.submit([&](sycl::handler &cgh) { - auto acc = subbuf_copy->get_access(cgh); + auto acc = subbuf_copy->get_access(cgh); cgh.set_args(acc); cgh.single_task(kernel_sycl); }); @@ -328,7 +328,7 @@ int main() { Q.wait_and_throw(); { - auto host_acc = subbuf_copy->get_access(); + auto host_acc = subbuf_copy->get_access(); std::cout << "On host: offset = " << host_acc[0] << std::endl; assert(host_acc[0] == 256 * 3 && "Invalid subbuffer origin"); } diff --git a/sycl/test-e2e/DeviceGlobal/device_global_device_only.hpp b/sycl/test-e2e/DeviceGlobal/device_global_device_only.hpp index e5406bb0a3679..58bb3d206d244 100644 --- a/sycl/test-e2e/DeviceGlobal/device_global_device_only.hpp +++ b/sycl/test-e2e/DeviceGlobal/device_global_device_only.hpp @@ -15,7 +15,7 @@ int test() { { buffer OutBuf(&OutVal, 1); Q.submit([&](handler &CGH) { - auto OutAcc = OutBuf.get_access(CGH); + auto OutAcc = OutBuf.get_access(CGH); CGH.single_task([=]() { OutAcc[0] = DeviceGlobalVar.get()[0]; }); }); } diff --git a/sycl/test-e2e/DeviceGlobal/device_global_subscript.hpp b/sycl/test-e2e/DeviceGlobal/device_global_subscript.hpp index 475a1f10dd17d..b8753d2890e9e 100644 --- a/sycl/test-e2e/DeviceGlobal/device_global_subscript.hpp +++ b/sycl/test-e2e/DeviceGlobal/device_global_subscript.hpp @@ -22,7 +22,7 @@ int test() { { buffer OutBuf{Out, 2}; Q.submit([&](handler &CGH) { - auto OutAcc = OutBuf.get_access(CGH); + auto OutAcc = OutBuf.get_access(CGH); CGH.single_task([=]() { OutAcc[0] = DeviceGlobalVar1[2]; OutAcc[1] = DeviceGlobalVar2[1]; diff --git a/sycl/test-e2e/DeviceLib/bfloat16_conversion_dlopen_test.hpp b/sycl/test-e2e/DeviceLib/bfloat16_conversion_dlopen_test.hpp index c879cdc2b3180..6260912b20d16 100644 --- a/sycl/test-e2e/DeviceLib/bfloat16_conversion_dlopen_test.hpp +++ b/sycl/test-e2e/DeviceLib/bfloat16_conversion_dlopen_test.hpp @@ -7,8 +7,8 @@ using namespace sycl; -constexpr access::mode sycl_read = access::mode::read; -constexpr access::mode sycl_write = access::mode::write; +constexpr access_mode sycl_read = access_mode::read; +constexpr access_mode sycl_write = access_mode::write; using BFP = sycl::ext::oneapi::bfloat16; diff --git a/sycl/test-e2e/DeviceLib/built-ins/fast-math-flag.cpp b/sycl/test-e2e/DeviceLib/built-ins/fast-math-flag.cpp index 0b5eae2289cc7..b9c5b840e4f91 100644 --- a/sycl/test-e2e/DeviceLib/built-ins/fast-math-flag.cpp +++ b/sycl/test-e2e/DeviceLib/built-ins/fast-math-flag.cpp @@ -32,9 +32,9 @@ bool checkEqual(marray A, marray B) { buffer input_buff(&input, 1); \ buffer res_buff(&res[0], sycl::range<1>(2)); \ deviceQueue.submit([&](handler &cgh) { \ - accessor res_acc(res_buff, \ - cgh); \ - accessor input_acc( \ + accessor res_acc(res_buff, \ + cgh); \ + accessor input_acc( \ input_buff, cgh); \ cgh.single_task([=]() { \ res_acc[0] = sycl::native::func(input_acc[0]); \ @@ -54,9 +54,9 @@ bool checkEqual(marray A, marray B) { buffer input_buff(&input[0], range<1>(2)); \ buffer res_buff(&res[0], range<1>(2)); \ deviceQueue.submit([&](handler &cgh) { \ - accessor res_acc(res_buff, \ - cgh); \ - accessor input_acc( \ + accessor res_acc(res_buff, \ + cgh); \ + accessor input_acc( \ input_buff, cgh); \ cgh.single_task([=]() { \ res_acc[0] = sycl::native::func(input_acc[0], input_acc[1]); \ diff --git a/sycl/test-e2e/DeviceLib/built-ins/scalar_geometric.cpp b/sycl/test-e2e/DeviceLib/built-ins/scalar_geometric.cpp index 634706c578ccf..8223729ad12aa 100644 --- a/sycl/test-e2e/DeviceLib/built-ins/scalar_geometric.cpp +++ b/sycl/test-e2e/DeviceLib/built-ins/scalar_geometric.cpp @@ -17,7 +17,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::dot(float{0.5}, float{1.6}); }); }); @@ -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::distance(float{1.f}, float{3.f}); }); }); @@ -47,7 +47,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::length(float{1.f}); }); }); @@ -62,7 +62,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::normalize(float{2.f}); }); }); @@ -77,7 +77,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::fast_distance(float{1.f}, float{3.f}); }); }); @@ -92,7 +92,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::fast_length(float{2.f}); }); }); @@ -107,7 +107,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::fast_normalize(float{2.f}); }); }); diff --git a/sycl/test-e2e/DeviceLib/built-ins/scalar_math_2.cpp b/sycl/test-e2e/DeviceLib/built-ins/scalar_math_2.cpp index 919e6c8d13dae..aa92a8d8685e8 100644 --- a/sycl/test-e2e/DeviceLib/built-ins/scalar_math_2.cpp +++ b/sycl/test-e2e/DeviceLib/built-ins/scalar_math_2.cpp @@ -20,7 +20,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::acospi(float{0.5}); }); }); @@ -35,7 +35,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::asinpi(float{0.5}); }); }); @@ -50,7 +50,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::atan2(float{0.5}, float{0.5}); }); }); @@ -65,7 +65,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::atanpi(float{0.5}); }); }); @@ -80,7 +80,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::atan2pi(float{0.5}, float{0.5}); }); }); @@ -95,7 +95,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::copysign(float{1}, float{-0.5}); }); }); @@ -110,7 +110,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::fdim(float{1.6}, float{0.6}); }); }); @@ -125,7 +125,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::fma(float{0.5}, float{10.0}, float{3.0}); }); }); @@ -143,8 +143,8 @@ int main() { {s::property::buffer::use_host_ptr()}); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); - auto AccI = BufI.get_access(cgh); + auto AccR = BufR.get_access(cgh); + auto AccI = BufI.get_access(cgh); cgh.single_task([=]() { s::global_ptr Iptr(AccI); AccR[0] = s::fract(float{1.5}, Iptr); @@ -165,8 +165,8 @@ int main() { {s::property::buffer::use_host_ptr()}); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); - auto AccI = BufI.get_access(cgh); + auto AccR = BufR.get_access(cgh); + auto AccI = BufI.get_access(cgh); cgh.single_task([=]() { float temp(0.0); s::private_ptr Iptr(&temp); @@ -189,8 +189,8 @@ int main() { {s::property::buffer::use_host_ptr()}); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); - auto AccI = BufI.get_access(cgh); + auto AccR = BufR.get_access(cgh); + auto AccI = BufI.get_access(cgh); cgh.single_task([=]() { int temp(0.0); s::private_ptr Iptr(&temp); @@ -213,8 +213,8 @@ int main() { {s::property::buffer::use_host_ptr()}); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); - auto AccI = BufI.get_access(cgh); + auto AccR = BufR.get_access(cgh); + auto AccI = BufI.get_access(cgh); cgh.single_task([=]() { int temp(0.0); s::private_ptr Iptr(&temp); diff --git a/sycl/test-e2e/DeviceLib/built-ins/vector_geometric.cpp b/sycl/test-e2e/DeviceLib/built-ins/vector_geometric.cpp index 6fa9948e77c5c..1fa4108a4917f 100644 --- a/sycl/test-e2e/DeviceLib/built-ins/vector_geometric.cpp +++ b/sycl/test-e2e/DeviceLib/built-ins/vector_geometric.cpp @@ -23,7 +23,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::dot( s::float2{ @@ -43,7 +43,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::cross( s::float4{ @@ -79,7 +79,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::cross( s::double4{ @@ -115,7 +115,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::distance( s::float2{ @@ -138,7 +138,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::length(s::float2{ 1.f, @@ -156,7 +156,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::normalize(s::float2{ 1.f, @@ -178,7 +178,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::fast_distance( s::float2{ @@ -201,7 +201,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::fast_length(s::float2{ 1.f, @@ -219,7 +219,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::fast_normalize(s::float2{ 1.f, diff --git a/sycl/test-e2e/DeviceLib/built-ins/vector_relational.cpp b/sycl/test-e2e/DeviceLib/built-ins/vector_relational.cpp index 77a80886d8756..ad7bccbacea5d 100644 --- a/sycl/test-e2e/DeviceLib/built-ins/vector_relational.cpp +++ b/sycl/test-e2e/DeviceLib/built-ins/vector_relational.cpp @@ -20,7 +20,7 @@ int main() { s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccR = BufR.get_access(cgh); + auto AccR = BufR.get_access(cgh); cgh.single_task([=]() { AccR[0] = s::isequal(s::float4{0.5f, 0.6f, NAN, INFINITY}, s::float4{0.5f, 0.5f, 0.5f, 0.5f}); @@ -45,7 +45,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::isnotequal(s::float4{0.5f, 0.6f, NAN, INFINITY}, s::float4{0.5f, 0.5f, 0.5f, 0.5f}); @@ -70,7 +70,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::isgreater(s::float4{0.5f, 0.6f, NAN, INFINITY}, s::float4{0.5f, 0.5f, 0.5f, 0.5f}); @@ -95,7 +95,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::isgreaterequal(s::float4{0.5f, 0.6f, NAN, INFINITY}, s::float4{0.5f, 0.5f, 0.5f, 0.5f}); @@ -120,7 +120,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::isless(s::float4{0.5f, 0.4f, NAN, INFINITY}, s::float4{0.5f, 0.5f, 0.5f, 0.5f}); @@ -145,7 +145,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::islessequal(s::float4{0.5f, 0.4f, NAN, INFINITY}, s::float4{0.5f, 0.5f, 0.5f, 0.5f}); @@ -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::islessgreater(s::float4{0.5f, 0.4f, NAN, INFINITY}, s::float4{0.5f, 0.5f, 0.5f, INFINITY}); @@ -196,7 +196,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::isfinite(s::float4{0.5f, 0.4f, NAN, INFINITY}); }); @@ -220,7 +220,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::isinf(s::float4{0.5f, 0.4f, NAN, INFINITY}); }); @@ -244,7 +244,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::isnan(s::float4{0.5f, 0.4f, NAN, INFINITY}); }); @@ -268,7 +268,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::isnormal(s::float4{0.5f, 0.4f, NAN, INFINITY}); }); @@ -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::isordered(s::float4{0.5f, 0.6f, NAN, INFINITY}, s::float4{0.5f, 0.5f, 0.5f, 0.5f}); @@ -317,7 +317,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::isunordered(s::float4{0.5f, 0.6f, NAN, INFINITY}, s::float4{0.5f, 0.5f, 0.5f, 0.5f}); @@ -342,7 +342,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::signbit(s::float3{0.5f, -12.0f, INFINITY}); }); @@ -365,7 +365,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::any(s::int4{-12, -12, 0, 1}); }); @@ -384,7 +384,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::any(s::int4{-12, -12, -12, -12}); }); @@ -403,7 +403,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::any(s::int4{0, 0, 0, 0}); }); @@ -422,7 +422,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::any(s::int4{12, 12, 12, 12}); }); @@ -441,7 +441,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::all(s::int4{-12, -12, -12, -12}); // Infinity (positive or negative) or Nan are not integers. @@ -463,7 +463,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::all(s::int4{-12, -12, -12, -12}); }); @@ -482,7 +482,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::all(s::int4{0, 0, 0, 0}); }); @@ -501,7 +501,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::all(s::int4{12, 12, 12, 12}); }); @@ -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::bitselect(s::float4{112.112f, 12.12f, 0, 0.0f}, s::float4{34.34f, 23.23f, 1, 0.0f}, @@ -546,7 +546,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(s::float4{112.112f, 34.34f, 112.112f, 34.34f}, s::float4{34.34f, 112.112f, 34.34f, 112.112f}, @@ -579,10 +579,10 @@ int main() { s::buffer> R(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { - auto AccA = A.get_access(cgh); - auto AccB = B.get_access(cgh); - auto AccM = M.get_access(cgh); - auto AccR = R.get_access(cgh); + auto AccA = A.get_access(cgh); + auto AccB = B.get_access(cgh); + auto AccM = M.get_access(cgh); + auto AccR = R.get_access(cgh); cgh.single_task( [=]() { AccR[0] = s::select(AccA[0], AccB[0], AccM[0]); }); }); diff --git a/sycl/test-e2e/DeviceLib/cmath_fp64_test.cpp b/sycl/test-e2e/DeviceLib/cmath_fp64_test.cpp index 5a8ba80a547d9..423e0d984571c 100644 --- a/sycl/test-e2e/DeviceLib/cmath_fp64_test.cpp +++ b/sycl/test-e2e/DeviceLib/cmath_fp64_test.cpp @@ -14,8 +14,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 75 diff --git a/sycl/test-e2e/DeviceLib/imf/fp32_test.cpp b/sycl/test-e2e/DeviceLib/imf/fp32_test.cpp index 2db3f16209bfa..e7d48256a70f7 100644 --- a/sycl/test-e2e/DeviceLib/imf/fp32_test.cpp +++ b/sycl/test-e2e/DeviceLib/imf/fp32_test.cpp @@ -5,8 +5,8 @@ #include "imf_utils.hpp" #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_test.cpp b/sycl/test-e2e/DeviceLib/math_fp64_test.cpp index 2e3eb85c6f7da..7384c704d58f5 100644 --- a/sycl/test-e2e/DeviceLib/math_fp64_test.cpp +++ b/sycl/test-e2e/DeviceLib/math_fp64_test.cpp @@ -12,8 +12,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 62 diff --git a/sycl/test-e2e/DeviceLib/math_test.cpp b/sycl/test-e2e/DeviceLib/math_test.cpp index f96f7b8be97f8..76eb8513182d1 100644 --- a/sycl/test-e2e/DeviceLib/math_test.cpp +++ b/sycl/test-e2e/DeviceLib/math_test.cpp @@ -10,8 +10,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 60 diff --git a/sycl/test-e2e/DeviceLib/rand_test.cpp b/sycl/test-e2e/DeviceLib/rand_test.cpp index 6a29935dfcf36..150145d0dbec8 100644 --- a/sycl/test-e2e/DeviceLib/rand_test.cpp +++ b/sycl/test-e2e/DeviceLib/rand_test.cpp @@ -40,7 +40,7 @@ int main() { // Submit a command group to the queue Q.submit([&](sycl::handler &cgh) { // Get access to the buffers - auto accessC = bufferC.get_access(cgh); + auto accessC = bufferC.get_access(cgh); // Execute the kernel cgh.parallel_for( diff --git a/sycl/test-e2e/DeviceLib/string_test.cpp b/sycl/test-e2e/DeviceLib/string_test.cpp index 66dca441ddeed..35d39a87d7ad7 100644 --- a/sycl/test-e2e/DeviceLib/string_test.cpp +++ b/sycl/test-e2e/DeviceLib/string_test.cpp @@ -31,8 +31,8 @@ bool kernel_test_memcpy(sycl::queue &deviceQueue) { sycl::buffer buffer1(src, sycl::range<1>(20)); sycl::buffer buffer2(dst, sycl::range<1>(20)); deviceQueue.submit([&](sycl::handler &cgh) { - auto dst_acc = buffer2.get_access(cgh); - auto src_acc = buffer1.get_access(cgh); + auto dst_acc = buffer2.get_access(cgh); + auto src_acc = buffer1.get_access(cgh); cgh.single_task([=]() { memcpy(dst_acc.get_multi_ptr().get(), src_acc.get_multi_ptr().get(), 20); @@ -59,7 +59,7 @@ bool kernel_test_strcpy(sycl::queue &deviceQueue) { { sycl::buffer dst_buffer(dst, sycl::range<1>(5)); deviceQueue.submit([&](sycl::handler &cgh) { - auto dst_acc = dst_buffer.get_access(cgh); + auto dst_acc = dst_buffer.get_access(cgh); cgh.single_task([=]() { char str[20] = "abcdefg012345xyzvvv"; strcpy(dst_acc[0], str); @@ -101,7 +101,7 @@ bool kernel_test_strncpy(sycl::queue &deviceQueue) { { sycl::buffer dst_buffer(dst, sycl::range<1>(3)); deviceQueue.submit([&](sycl::handler &cgh) { - auto dst_acc = dst_buffer.get_access(cgh); + auto dst_acc = dst_buffer.get_access(cgh); cgh.single_task([=]() { char str[20] = "abcdefg012345xyzvvv"; strncpy(dst_acc[0], str, 19); @@ -148,7 +148,7 @@ bool kernel_test_strcmp(sycl::queue &deviceQueue) { { sycl::buffer res_buffer(res, sycl::range<1>(5)); deviceQueue.submit([&](sycl::handler &cgh) { - auto res_acc = res_buffer.get_access(cgh); + auto res_acc = res_buffer.get_access(cgh); cgh.single_task([=]() { char str1[20] = "abcdefg012"; char str2[20] = "abcd"; @@ -175,7 +175,7 @@ bool kernel_test_strncmp(sycl::queue &deviceQueue) { { sycl::buffer res_buffer(res, sycl::range<1>(10)); deviceQueue.submit([&](sycl::handler &cgh) { - auto res_acc = res_buffer.get_access(cgh); + auto res_acc = res_buffer.get_access(cgh); cgh.single_task([=]() { char str1[20] = "abcdefg012"; char str2[20] = "abcd"; @@ -216,8 +216,8 @@ bool kernel_test_strlen(sycl::queue &deviceQueue) { sycl::buffer buffer1(src, sycl::range<1>(20)); sycl::buffer buffer2(len, sycl::range<1>(5)); deviceQueue.submit([&](sycl::handler &cgh) { - auto len_acc = buffer2.get_access(cgh); - auto src_acc = buffer1.get_access(cgh); + auto len_acc = buffer2.get_access(cgh); + auto src_acc = buffer1.get_access(cgh); cgh.single_task([=]() { len_acc[0] = strlen(src_acc.get_multi_ptr().get()); @@ -348,7 +348,7 @@ bool kernel_test_memset(sycl::queue &deviceQueue) { { sycl::buffer buffer1(dst, sycl::range<1>(20)); deviceQueue.submit([&](sycl::handler &cgh) { - auto dst_acc = buffer1.get_access(cgh); + auto dst_acc = buffer1.get_access(cgh); cgh.single_task([=]() { memset(dst_acc.get_multi_ptr().get(), 'P', 18); @@ -459,15 +459,15 @@ bool kernel_test_memcmp(sycl::queue &deviceQueue) { sycl::buffer buffer_str7(str7, sycl::range<1>(sizeof(str7))); sycl::buffer buffer_str8(str8, sycl::range<1>(sizeof(str8))); deviceQueue.submit([&](sycl::handler &cgh) { - auto results_acc = buffer1.get_access(cgh); - auto str1_acc = buffer_str1.get_access(cgh); - auto str2_acc = buffer_str2.get_access(cgh); - auto str3_acc = buffer_str3.get_access(cgh); - auto str4_acc = buffer_str4.get_access(cgh); - auto str5_acc = buffer_str5.get_access(cgh); - auto str6_acc = buffer_str6.get_access(cgh); - auto str7_acc = buffer_str7.get_access(cgh); - auto str8_acc = buffer_str8.get_access(cgh); + auto results_acc = buffer1.get_access(cgh); + auto str1_acc = buffer_str1.get_access(cgh); + auto str2_acc = buffer_str2.get_access(cgh); + auto str3_acc = buffer_str3.get_access(cgh); + auto str4_acc = buffer_str4.get_access(cgh); + auto str5_acc = buffer_str5.get_access(cgh); + auto str6_acc = buffer_str6.get_access(cgh); + auto str7_acc = buffer_str7.get_access(cgh); + auto str8_acc = buffer_str8.get_access(cgh); cgh.single_task([=]() { results_acc[0] = memcmp( str1_acc.get_multi_ptr().get(), @@ -521,7 +521,7 @@ bool kernel_test_memcmp_align(sycl::queue &deviceQueue) { { sycl::buffer cmp_buf(cmps, sycl::range<1>{16}); deviceQueue.submit([&](sycl::handler &cgh) { - auto cmp_acc = cmp_buf.get_access(cgh); + auto cmp_acc = cmp_buf.get_access(cgh); cgh.single_task([=]() { uint8_t s1_buf[32], s2_buf[32]; uint8_t *s1_ptr = &s1_buf[0]; @@ -592,19 +592,19 @@ bool kernel_test_memcpy_addr_space(sycl::queue &deviceQueue) { sycl::buffer buffer2(dst, sycl::range<1>(16)); sycl::buffer buffer3(dst1, sycl::range<1>(16)); deviceQueue.submit([&](sycl::handler &cgh) { - sycl::accessor src_acc(buffer1, cgh); sycl::local_accessor local_acc(sycl::range<1>(16), cgh); - sycl::accessor dst_acc(buffer2, cgh); - sycl::accessor dst1_acc(buffer3, cgh); diff --git a/sycl/test-e2e/EnqueueNativeCommand/custom-command-cuda.cpp b/sycl/test-e2e/EnqueueNativeCommand/custom-command-cuda.cpp index 686b44bde7013..076e290635cda 100644 --- a/sycl/test-e2e/EnqueueNativeCommand/custom-command-cuda.cpp +++ b/sycl/test-e2e/EnqueueNativeCommand/custom-command-cuda.cpp @@ -33,8 +33,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 Stream = IH.get_native_queue(); @@ -54,7 +54,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; }; @@ -65,8 +65,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; @@ -80,7 +80,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_group_size.cpp b/sycl/test-e2e/Experimental/launch_queries/max_sub_group_size.cpp index 406ba3b3afbd3..f826c0de45ca0 100644 --- a/sycl/test-e2e/Experimental/launch_queries/max_sub_group_size.cpp +++ b/sycl/test-e2e/Experimental/launch_queries/max_sub_group_size.cpp @@ -25,7 +25,7 @@ namespace kernels { template using sycl_global_accessor = - sycl::accessor; class TestKernel { @@ -54,7 +54,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/launch_queries/max_work_item_sizes.cpp b/sycl/test-e2e/Experimental/launch_queries/max_work_item_sizes.cpp index 75278f40f6aa3..7872ec35eff11 100644 --- a/sycl/test-e2e/Experimental/launch_queries/max_work_item_sizes.cpp +++ b/sycl/test-e2e/Experimental/launch_queries/max_work_item_sizes.cpp @@ -17,7 +17,7 @@ namespace kernels { template using sycl_global_accessor = - sycl::accessor; class TestKernel { @@ -74,7 +74,7 @@ int main() { auto LaunchRange = sycl::nd_range<1>{sycl::range<1>{MaxWorkGroupSizeActual}, sycl::range<1>{MaxWorkGroupSizeActual}}; 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/event_reassociation_dependency.cpp b/sycl/test-e2e/Experimental/reusable_events/event_reassociation_dependency.cpp index 7dc45633512c0..4920fa261e772 100644 --- a/sycl/test-e2e/Experimental/reusable_events/event_reassociation_dependency.cpp +++ b/sycl/test-e2e/Experimental/reusable_events/event_reassociation_dependency.cpp @@ -25,7 +25,7 @@ int main() { // First operation q.submit([&](sycl::handler &cgh) { - auto acc = buf1.get_access(cgh); + auto acc = buf1.get_access(cgh); cgh.parallel_for(sycl::range<1>(N), [=](sycl::id<1> idx) { acc[idx] = 10; }); }); @@ -34,7 +34,7 @@ int main() { // Second operation depends on event auto evt2 = q.submit([&](sycl::handler &cgh) { cgh.depends_on(event); - auto acc = buf2.get_access(cgh); + auto acc = buf2.get_access(cgh); cgh.parallel_for(sycl::range<1>(N), [=](sycl::id<1> idx) { acc[idx] = 20; }); }); diff --git a/sycl/test-e2e/Experimental/reusable_events/kernel_timing.cpp b/sycl/test-e2e/Experimental/reusable_events/kernel_timing.cpp index fccae8afa3ab7..5a7f9e41f0c8c 100644 --- a/sycl/test-e2e/Experimental/reusable_events/kernel_timing.cpp +++ b/sycl/test-e2e/Experimental/reusable_events/kernel_timing.cpp @@ -32,7 +32,7 @@ int main() { syclex::enqueue_signal_event(q, start_event); q.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.parallel_for(sycl::range<1>(N), [=](sycl::id<1> idx) { for (int i = 0; i < 100; ++i) { @@ -42,7 +42,7 @@ int main() { }); 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) { for (int i = 0; i < 100; ++i) { diff --git a/sycl/test-e2e/Experimental/reusable_events/wait_multiple_events.cpp b/sycl/test-e2e/Experimental/reusable_events/wait_multiple_events.cpp index f5ee50c6c0572..7e76ee5c78681 100644 --- a/sycl/test-e2e/Experimental/reusable_events/wait_multiple_events.cpp +++ b/sycl/test-e2e/Experimental/reusable_events/wait_multiple_events.cpp @@ -34,7 +34,7 @@ int main() { // Launch on q1 q1.submit([&](sycl::handler &cgh) { - auto acc = buf1.get_access(cgh); + auto acc = buf1.get_access(cgh); cgh.parallel_for( sycl::range<1>(N), [=](sycl::id<1> idx) { acc[idx] = acc[idx] * 3; }); }); @@ -42,7 +42,7 @@ int main() { // Launch on q2 q2.submit([&](sycl::handler &cgh) { - auto acc = buf2.get_access(cgh); + auto acc = buf2.get_access(cgh); cgh.parallel_for( sycl::range<1>(N), [=](sycl::id<1> idx) { acc[idx] = acc[idx] * 5; }); }); @@ -54,9 +54,9 @@ int main() { // Combine results on q3 q3.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/FreeFunctionKernels/template_specialization.cpp b/sycl/test-e2e/FreeFunctionKernels/template_specialization.cpp index 3d0a0b9b6be7d..d7129e6707e54 100644 --- a/sycl/test-e2e/FreeFunctionKernels/template_specialization.cpp +++ b/sycl/test-e2e/FreeFunctionKernels/template_specialization.cpp @@ -140,7 +140,7 @@ void test_accessor() { kernel Kernel = bundle.get_kernel(id); sycl::buffer buf(data, sycl::range<1>(N)); Q.submit([&](handler &h) { - auto acc = buf.get_access(h); + auto acc = buf.get_access(h); h.set_args(acc); h.parallel_for(nd_range{{1}, {1}}, Kernel); }); diff --git a/sycl/test-e2e/GroupAlgorithm/back_to_back_collectives.cpp b/sycl/test-e2e/GroupAlgorithm/back_to_back_collectives.cpp index 5678ae9412de5..d4f574a653ec2 100644 --- a/sycl/test-e2e/GroupAlgorithm/back_to_back_collectives.cpp +++ b/sycl/test-e2e/GroupAlgorithm/back_to_back_collectives.cpp @@ -35,10 +35,10 @@ int main() { buffer EScanBuf(EScan.data(), N); buffer IScanBuf(IScan.data(), N); q.submit([&](handler &h) { - auto Input = InputBuf.get_access(h); - auto Sum = SumBuf.get_access(h); - auto EScan = EScanBuf.get_access(h); - auto IScan = IScanBuf.get_access(h); + auto Input = InputBuf.get_access(h); + auto Sum = SumBuf.get_access(h); + auto EScan = EScanBuf.get_access(h); + auto IScan = IScanBuf.get_access(h); h.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) { size_t i = it.get_global_id(0); auto g = it.get_group(); diff --git a/sycl/test-e2e/GroupInterface/leader_of.cpp b/sycl/test-e2e/GroupInterface/leader_of.cpp index 7931eae21b775..76d1ac5cbd361 100644 --- a/sycl/test-e2e/GroupInterface/leader_of.cpp +++ b/sycl/test-e2e/GroupInterface/leader_of.cpp @@ -17,7 +17,7 @@ void test(queue q) { buffer out_buf(&out, 1); q.submit([&](handler &cgh) { - auto out = out_buf.template get_access(cgh); + auto out = out_buf.template get_access(cgh); cgh.parallel_for(nd_range<2>(R, R), [=](nd_item<2> it) { khr::work_group<2> g = it.get_group(); if (khr::leader_of(g)) { diff --git a/sycl/test-e2e/HierPar/Inputs/hier_par_wgscope_impl.hpp b/sycl/test-e2e/HierPar/Inputs/hier_par_wgscope_impl.hpp index 987af01a12361..2a131a557e953 100644 --- a/sycl/test-e2e/HierPar/Inputs/hier_par_wgscope_impl.hpp +++ b/sycl/test-e2e/HierPar/Inputs/hier_par_wgscope_impl.hpp @@ -42,7 +42,7 @@ static bool testWgScope(queue &Q) { int N = N_INNER_ITER; Q.submit([&](handler &cgh) { - auto DevPtr = Buf.get_access(cgh); + auto DevPtr = Buf.get_access(cgh); cgh.parallel_for_work_group( range<1>(N_WG), range<1>(PHYS_WG_SIZE), [=](group<1> G) { @@ -202,7 +202,7 @@ bool testPrivateMemory(queue &Q) { std::memset(Ptr, 0, RangeLength * sizeof(Ptr[0])); buffer Buf(Ptr, range<1>(RangeLength)); Q.submit([&](handler &cgh) { - auto DevPtr = Buf.get_access(cgh); + auto DevPtr = 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) { diff --git a/sycl/test-e2e/HostInteropTask/host-task-dependency4.cpp b/sycl/test-e2e/HostInteropTask/host-task-dependency4.cpp index bb315329195bf..ca0e58b827b2f 100644 --- a/sycl/test-e2e/HostInteropTask/host-task-dependency4.cpp +++ b/sycl/test-e2e/HostInteropTask/host-task-dependency4.cpp @@ -6,7 +6,7 @@ sycl::event submit(sycl::queue &Q, sycl::buffer &B) { return Q.submit([&](sycl::handler &CGH) { - auto A = B.template get_access(CGH); + auto A = B.template get_access(CGH); CGH.host_task([=]() { (void)A; }); }); } diff --git a/sycl/test-e2e/HostInteropTask/interop-task-cuda.cpp b/sycl/test-e2e/HostInteropTask/interop-task-cuda.cpp index adde6ed4d8d03..8d735c63b4253 100644 --- a/sycl/test-e2e/HostInteropTask/interop-task-cuda.cpp +++ b/sycl/test-e2e/HostInteropTask/interop-task-cuda.cpp @@ -33,8 +33,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 Stream = IH.get_native_queue(); @@ -59,7 +59,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; }; @@ -70,8 +70,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; @@ -85,7 +85,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_empty.cpp b/sycl/test-e2e/InlineAsm/asm_16_empty.cpp index 3579ef55b4eee..426e05d734c2b 100644 --- a/sycl/test-e2e/InlineAsm/asm_16_empty.cpp +++ b/sycl/test-e2e/InlineAsm/asm_16_empty.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_16_no_opts.cpp b/sycl/test-e2e/InlineAsm/asm_16_no_opts.cpp index 289969b5e8c5a..55b80447eb090 100644 --- a/sycl/test-e2e/InlineAsm/asm_16_no_opts.cpp +++ b/sycl/test-e2e/InlineAsm/asm_16_no_opts.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_arbitrary_ops_order.cpp b/sycl/test-e2e/InlineAsm/asm_arbitrary_ops_order.cpp index 86d3098648d83..7b418565e746b 100644 --- a/sycl/test-e2e/InlineAsm/asm_arbitrary_ops_order.cpp +++ b/sycl/test-e2e/InlineAsm/asm_arbitrary_ops_order.cpp @@ -19,16 +19,16 @@ 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->getInputBuffer(2).template get_access( + this->getInputBuffer(2).template get_access( cgh); auto D = - this->getOutputBuffer().template get_access( + this->getOutputBuffer().template get_access( cgh); cgh.parallel_for>( diff --git a/sycl/test-e2e/InlineAsm/asm_float_imm_arg.cpp b/sycl/test-e2e/InlineAsm/asm_float_imm_arg.cpp index dddeba3764ce6..242150d475412 100644 --- a/sycl/test-e2e/InlineAsm/asm_float_imm_arg.cpp +++ b/sycl/test-e2e/InlineAsm/asm_float_imm_arg.cpp @@ -19,10 +19,10 @@ struct KernelFunctor : WithInputBuffers, WithOutputBuffer { void operator()(sycl::handler &cgh) { auto A = - this->getInputBuffer(0).template get_access( + this->getInputBuffer(0).template get_access( cgh); auto B = - this->getOutputBuffer().template get_access( + this->getOutputBuffer().template get_access( cgh); cgh.parallel_for>( diff --git a/sycl/test-e2e/InlineAsm/asm_imm_arg.cpp b/sycl/test-e2e/InlineAsm/asm_imm_arg.cpp index 0cc620f806c63..553fdd53d36d1 100644 --- a/sycl/test-e2e/InlineAsm/asm_imm_arg.cpp +++ b/sycl/test-e2e/InlineAsm/asm_imm_arg.cpp @@ -18,10 +18,10 @@ struct KernelFunctor : WithInputBuffers, WithOutputBuffer { void operator()(sycl::handler &cgh) { auto A = - this->getInputBuffer(0).template get_access( + this->getInputBuffer(0).template get_access( cgh); auto B = - this->getOutputBuffer().template get_access( + this->getOutputBuffer().template get_access( cgh); cgh.parallel_for>( diff --git a/sycl/test-e2e/InlineAsm/asm_multiple_instructions.cpp b/sycl/test-e2e/InlineAsm/asm_multiple_instructions.cpp index 0904725a47684..d67b341da5613 100644 --- a/sycl/test-e2e/InlineAsm/asm_multiple_instructions.cpp +++ b/sycl/test-e2e/InlineAsm/asm_multiple_instructions.cpp @@ -22,15 +22,15 @@ struct KernelFunctor : WithInputBuffers, WithOutputBuffer { void operator()(sycl::handler &cgh) { auto A = this->getInputBuffer(0) - .template get_access(cgh); + .template get_access(cgh); auto B = - this->getInputBuffer(1).template get_access( + this->getInputBuffer(1).template get_access( cgh); auto C = - this->getInputBuffer(2).template get_access( + this->getInputBuffer(2).template get_access( cgh); auto D = - this->getOutputBuffer().template get_access( + this->getOutputBuffer().template get_access( cgh); cgh.parallel_for>( diff --git a/sycl/test-e2e/InlineAsm/asm_switch.cpp b/sycl/test-e2e/InlineAsm/asm_switch.cpp index 4d9c18eab108f..8dbb37ce6c810 100644 --- a/sycl/test-e2e/InlineAsm/asm_switch.cpp +++ b/sycl/test-e2e/InlineAsm/asm_switch.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); int switchField = 2; CGH.parallel_for>( diff --git a/sycl/test-e2e/KernelAndProgram/cache_env_vars.hpp b/sycl/test-e2e/KernelAndProgram/cache_env_vars.hpp index a8a81fc270133..b4ceb2c4e73fa 100644 --- a/sycl/test-e2e/KernelAndProgram/cache_env_vars.hpp +++ b/sycl/test-e2e/KernelAndProgram/cache_env_vars.hpp @@ -90,7 +90,7 @@ int main(int argc, char **argv) { sycl::range<1> NumOfWorkItems{buf.size()}; q.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.parallel_for( NumOfWorkItems, [=](sycl::id<1> WIid) { TARGET_IMAGE(acc[0]) }); }); diff --git a/sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp b/sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp index d661c6cfe66ea..d4324a0967a58 100644 --- a/sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp +++ b/sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp @@ -69,7 +69,7 @@ void test_local_accessor() { sycl::buffer b(init, N_wg); q.submit([&](handler &cgh) { - auto acc_global = b.get_access(cgh); + auto acc_global = b.get_access(cgh); local_accessor acc_local(N_slm, cgh); cgh.set_arg(0, acc_global); diff --git a/sycl/test-e2e/KernelParams/array-kernel-param-run.cpp b/sycl/test-e2e/KernelParams/array-kernel-param-run.cpp index 87e93b1a60a9a..9301c1a546477 100644 --- a/sycl/test-e2e/KernelParams/array-kernel-param-run.cpp +++ b/sycl/test-e2e/KernelParams/array-kernel-param-run.cpp @@ -67,7 +67,7 @@ bool test_one_array(queue &myQueue) { auto out_buffer = buffer(&output[0], num_items); myQueue.submit([&](handler &cgh) { - auto output_accessor = out_buffer.get_access(cgh); + auto output_accessor = out_buffer.get_access(cgh); cgh.parallel_for(num_items, [=](sycl::id<1> index) { output_accessor[index] = input1[0][index] + input2[2][1][index] + 1; @@ -90,7 +90,7 @@ bool test_two_arrays(queue &myQueue) { auto out_buffer = buffer(&output[0], num_items); myQueue.submit([&](handler &cgh) { - auto output_accessor = out_buffer.get_access(cgh); + auto output_accessor = out_buffer.get_access(cgh); cgh.parallel_for(num_items, [=](sycl::id<1> index) { output_accessor[index] = input1[index] + input2[index]; @@ -118,10 +118,10 @@ bool test_accessor_arrays_1(queue &myQueue) { myQueue.submit([&](handler &cgh) { using Accessor = - accessor; + accessor; Accessor a[2] = { - in_buffer1.get_access(cgh), - in_buffer2.get_access(cgh), + in_buffer1.get_access(cgh), + in_buffer2.get_access(cgh), }; cgh.parallel_for( @@ -150,12 +150,12 @@ bool test_accessor_arrays_2(queue &myQueue) { myQueue.submit([&](handler &cgh) { using Accessor = - accessor; - Accessor a[4] = {in_buffer1.get_access(cgh), - in_buffer2.get_access(cgh), - in_buffer1.get_access(cgh), - in_buffer2.get_access(cgh)}; - auto output_accessor = out_buffer.get_access(cgh); + accessor; + Accessor a[4] = {in_buffer1.get_access(cgh), + in_buffer2.get_access(cgh), + in_buffer1.get_access(cgh), + in_buffer2.get_access(cgh)}; + auto output_accessor = out_buffer.get_access(cgh); cgh.parallel_for( num_items, [=](sycl::id<1> index) { diff --git a/sycl/test-e2e/KernelParams/union_kernel_param.cpp b/sycl/test-e2e/KernelParams/union_kernel_param.cpp index b81fe04cba0aa..cf1e30d5f5b1b 100644 --- a/sycl/test-e2e/KernelParams/union_kernel_param.cpp +++ b/sycl/test-e2e/KernelParams/union_kernel_param.cpp @@ -25,7 +25,7 @@ int main(int argc, char **argv) { { sycl::buffer buf(&myfloat, 1); queue.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.single_task([=]() { acc[0] = x.myfloat; }); }); } diff --git a/sycl/test-e2e/Reduction/reduction_ctor.cpp b/sycl/test-e2e/Reduction/reduction_ctor.cpp index bf59f3b91a15b..7f28b5db8365f 100644 --- a/sycl/test-e2e/Reduction/reduction_ctor.cpp +++ b/sycl/test-e2e/Reduction/reduction_ctor.cpp @@ -59,9 +59,9 @@ void testKnown(T Identity, BinaryOperation BOp, T A, T B) { Q.submit([&](handler &CGH) { // Reduction needs a device accessor as a parameter. // This accessor is not really used in this test. - accessor - ReduRWAcc(ReduBuf, CGH); - accessor + accessor ReduRWAcc( + ReduBuf, CGH); + accessor ReduDWAcc(ReduBuf, CGH); auto Redu = sycl::reduction(ReduBuf, CGH, BOp); auto ReduUSM = sycl::reduction(ReduUSMPtr, BOp); @@ -91,9 +91,9 @@ void testUnknown(T Identity, BinaryOperation BOp, T A, T B) { Q.submit([&](handler &CGH) { // Reduction needs a device accessor as a parameter. // This accessor is not really used in this test. - accessor - ReduRWAcc(ReduBuf, CGH); - accessor + accessor ReduRWAcc( + ReduBuf, CGH); + accessor ReduDWAcc(ReduBuf, CGH); auto Redu = sycl::reduction(ReduBuf, CGH, Identity, BOp); auto ReduUSM = sycl::reduction(ReduUSMPtr, Identity, BOp); diff --git a/sycl/test-e2e/Reduction/reduction_nd_ext_type.hpp b/sycl/test-e2e/Reduction/reduction_nd_ext_type.hpp index 54491ca236529..d8824049854e4 100644 --- a/sycl/test-e2e/Reduction/reduction_nd_ext_type.hpp +++ b/sycl/test-e2e/Reduction/reduction_nd_ext_type.hpp @@ -26,8 +26,8 @@ template int runTests(sycl::aspect ExtAspect) { return 0; } - 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; tests(Q, 1, 77, std::multiplies{}, 4, 4, init_to_identity()); tests(Q, 1, 77, std::multiplies{}, 4, 8); diff --git a/sycl/test-e2e/Reduction/reduction_range_1d_rw.cpp b/sycl/test-e2e/Reduction/reduction_range_1d_rw.cpp index 57d6f652aa8b3..c53e56d08687b 100644 --- a/sycl/test-e2e/Reduction/reduction_range_1d_rw.cpp +++ b/sycl/test-e2e/Reduction/reduction_range_1d_rw.cpp @@ -21,7 +21,7 @@ int main() { size_t MaxWGSize = Q.get_device().get_info(); - constexpr access::mode RW = access::mode::read_write; + constexpr access_mode RW = access_mode::read_write; // Fast-reduce and Fast-atomics. Try various range types/sizes. tests(Q, 0, 99, std::plus{}, range<1>(1)); tests(Q, 0, 99, std::plus<>{}, range<1>(2)); diff --git a/sycl/test-e2e/Reduction/reduction_usm_dw.cpp b/sycl/test-e2e/Reduction/reduction_usm_dw.cpp index a3ca89baeea56..6fccc6b8a28c9 100644 --- a/sycl/test-e2e/Reduction/reduction_usm_dw.cpp +++ b/sycl/test-e2e/Reduction/reduction_usm_dw.cpp @@ -43,7 +43,7 @@ int test(queue &Q, T Identity, T Init, size_t WGSize, size_t NWItems, // Compute. Q.submit([&](handler &CGH) { - auto In = InBuf.template get_access(CGH); + auto In = InBuf.template get_access(CGH); auto Redu = reduction(ReduVarPtr, Identity, BOp, {property::reduction::initialize_to_identity{}}); @@ -58,7 +58,7 @@ int test(queue &Q, T Identity, T Init, size_t WGSize, size_t NWItems, 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/atomic_load.cpp b/sycl/test-e2e/Regression/atomic_load.cpp index 3f5c31b9ad0f6..32c183de208c8 100644 --- a/sycl/test-e2e/Regression/atomic_load.cpp +++ b/sycl/test-e2e/Regression/atomic_load.cpp @@ -14,7 +14,7 @@ template void kernel_func(T val) { buffer buf(&data, range<1>(1)); testQueue.submit([&](handler &cgh) { - auto GlobAcc = buf.template get_access(cgh); + auto GlobAcc = buf.template get_access(cgh); cgh.single_task>([=]() { auto a = GlobAcc[0]; T var = a.load(); diff --git a/sycl/test-e2e/Regression/complex_global_object.cpp b/sycl/test-e2e/Regression/complex_global_object.cpp index 4a34cdf968a1a..2fe355143ef04 100644 --- a/sycl/test-e2e/Regression/complex_global_object.cpp +++ b/sycl/test-e2e/Regression/complex_global_object.cpp @@ -18,7 +18,7 @@ class ComplexClass { sycl::queue Queue{Device}; sycl::buffer Buf{sycl::range<1>{16}}; Queue.submit([&](sycl::handler &CGH) { - auto Acc = Buf.get_access(CGH); + auto Acc = Buf.get_access(CGH); CGH.single_task([=]() { Acc[0] = 42; }); }); } diff --git a/sycl/test-e2e/Regression/fsycl-host-compiler-win.cpp b/sycl/test-e2e/Regression/fsycl-host-compiler-win.cpp index 4ffe42c45c52f..37fc4d5c074b6 100644 --- a/sycl/test-e2e/Regression/fsycl-host-compiler-win.cpp +++ b/sycl/test-e2e/Regression/fsycl-host-compiler-win.cpp @@ -29,7 +29,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/image_access.cpp b/sycl/test-e2e/Regression/image_access.cpp index 901c78a86a939..eeb95d70a7af7 100644 --- a/sycl/test-e2e/Regression/image_access.cpp +++ b/sycl/test-e2e/Regression/image_access.cpp @@ -24,7 +24,7 @@ int main() { sycl::queue Queue; Queue.submit([&](sycl::handler &CGH) { - sycl::accessor A(Image, CGH); @@ -32,7 +32,7 @@ int main() { }); Queue.wait_and_throw(); - sycl::accessor A(Image); diff --git a/sycl/test-e2e/Regression/implicit_offset_debug_info.cpp b/sycl/test-e2e/Regression/implicit_offset_debug_info.cpp index d54b81a60b097..15cbce2504e74 100644 --- a/sycl/test-e2e/Regression/implicit_offset_debug_info.cpp +++ b/sycl/test-e2e/Regression/implicit_offset_debug_info.cpp @@ -12,7 +12,7 @@ int main() { queue q; buffer t1(10); q.submit([&](handler &cgh) { - auto table = t1.get_access(cgh); + auto table = t1.get_access(cgh); cgh.parallel_for(10, [=](id<1> gtid) { table[gtid] = gtid[0]; }); }); diff --git a/sycl/test-e2e/Regression/kernel_name_class.cpp b/sycl/test-e2e/Regression/kernel_name_class.cpp index c4478d075bd89..80d431f346bad 100644 --- a/sycl/test-e2e/Regression/kernel_name_class.cpp +++ b/sycl/test-e2e/Regression/kernel_name_class.cpp @@ -68,7 +68,7 @@ struct Wrapper { // PD-- // bool as kernel name deviceQueue.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.single_task([=]() { acc[0] += GOLD; }); }); ++NumTestCases; @@ -76,7 +76,7 @@ struct Wrapper { // PI-- // traditional in-place incomplete type deviceQueue.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.single_task([=]() { acc[0] += GOLD; }); }); ++NumTestCases; @@ -85,7 +85,7 @@ struct Wrapper { // a class completely defined within a namespace at // translation unit scope deviceQueue.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.single_task([=]() { acc[0] += GOLD; }); }); ++NumTestCases; @@ -96,7 +96,7 @@ struct Wrapper { // local scope class KernelName2; deviceQueue.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.single_task([=]() { acc[0] += GOLD; }); }); ++NumTestCases; @@ -107,7 +107,7 @@ struct Wrapper { // kernel name is a class defined in local scope class KernelName2a {}; deviceQueue.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.single_task([=]() { acc[0] += GOLD; }); }); ++NumTestCases; @@ -117,7 +117,7 @@ struct Wrapper { // an incomplete class forward-declared in a namespace at // translation unit scope deviceQueue.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.single_task([=]() { acc[0] += GOLD; }); }); ++NumTestCases; @@ -126,7 +126,7 @@ struct Wrapper { // an incomplete template specialization class with defined class as // argument declared in a namespace at translation unit scope deviceQueue.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.single_task>( [=]() { acc[0] += GOLD; }); }); @@ -136,7 +136,7 @@ struct Wrapper { // an incomplete template specialization class with incomplete class as // argument forward-declared in a namespace at translation unit scope deviceQueue.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.single_task>( [=]() { acc[0] += GOLD; }); }); @@ -146,7 +146,7 @@ struct Wrapper { // a defined template specialization class with defined class as argument // declared in a namespace at translation unit scope deviceQueue.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.single_task>( [=]() { acc[0] += GOLD; }); }); @@ -156,7 +156,7 @@ struct Wrapper { // a defined template specialization class with incomplete class as // argument forward-declared in a namespace at translation unit scope deviceQueue.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.single_task>( [=]() { acc[0] += GOLD; }); }); @@ -166,7 +166,7 @@ struct Wrapper { // an incomplete template specialization class with incomplete class as // argument forward-declared "in-place" deviceQueue.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.single_task>( [=]() { acc[0] += GOLD; }); }); @@ -179,7 +179,7 @@ struct Wrapper { // argument forward-declared locally class KernelName6; deviceQueue.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.single_task>( [=]() { acc[0] += GOLD; }); }); @@ -190,7 +190,7 @@ struct Wrapper { // a defined template specialization class with incomplete class as // argument forward-declared "in-place" deviceQueue.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.single_task>( [=]() { acc[0] += GOLD; }); }); @@ -202,7 +202,7 @@ struct Wrapper { // incomplete class as argument forward-declared locally class KernelName6a; deviceQueue.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.single_task>( [=]() { acc[0] += GOLD; }); }); @@ -216,7 +216,7 @@ struct Wrapper { // defined locally class KernelName9 {}; deviceQueue.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.single_task>( [=]() { acc[0] += GOLD; }); }); @@ -229,7 +229,7 @@ struct Wrapper { // an incomplete template specialization class with a defined class as // argument declared in the containing class deviceQueue.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.single_task>([=]() { acc[0] += GOLD; }); }); ++NumTestCases; @@ -241,7 +241,7 @@ struct Wrapper { // an incomplete template specialization class with an incomplete class as // argument declared in the containing class deviceQueue.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.single_task>([=]() { acc[0] += GOLD; }); }); ++NumTestCases; @@ -251,7 +251,7 @@ struct Wrapper { // translation unit scope with a defined class as argument declared in // a namespace at translation unit scope deviceQueue.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.single_task>>( [=]() { acc[0] += GOLD; }); }); @@ -262,7 +262,7 @@ struct Wrapper { // namespace at translation unit scope with a defined class as argument // declared in a namespace at translation unit scope deviceQueue.submit([&](sycl::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/Regression/mad_sat.cpp b/sycl/test-e2e/Regression/mad_sat.cpp index 964997171672f..c7e683c07285c 100644 --- a/sycl/test-e2e/Regression/mad_sat.cpp +++ b/sycl/test-e2e/Regression/mad_sat.cpp @@ -27,7 +27,7 @@ int main() { sycl::buffer buffer(1); testQueue.submit([&](sycl::handler &h) { - auto resultPtr = buffer.template get_access(h); + auto resultPtr = buffer.template get_access(h); h.single_task([=]() { resultPtr[0] = sycl::mad_sat(inputData_0, inputData_1, inputData_2); }); @@ -51,7 +51,7 @@ int main() { sycl::buffer buffer(1); testQueue.submit([&](sycl::handler &h) { - auto resultPtr = buffer.template get_access(h); + auto resultPtr = buffer.template get_access(h); h.single_task([=]() { resultPtr[0] = sycl::mad_sat(inputData_0, inputData_1, inputData_2); }); @@ -82,7 +82,7 @@ int main() { sycl::buffer buffer(1); testQueue.submit([&](sycl::handler &h) { - auto resultPtr = buffer.template get_access(h); + auto resultPtr = buffer.template get_access(h); h.single_task([=]() { resultPtr[0] = sycl::mad_sat(inputData_0, inputData_1, inputData_2); }); @@ -125,7 +125,7 @@ int main() { sycl::buffer buffer(1); testQueue.submit([&](sycl::handler &h) { - auto resultPtr = buffer.template get_access(h); + auto resultPtr = buffer.template get_access(h); h.single_task([=]() { resultPtr[0] = sycl::mad_sat(inputData_0, inputData_1, inputData_2); diff --git a/sycl/test-e2e/Regression/multithread_write_accessor.cpp b/sycl/test-e2e/Regression/multithread_write_accessor.cpp index 4c854eb861b4b..8cb4e155cef68 100644 --- a/sycl/test-e2e/Regression/multithread_write_accessor.cpp +++ b/sycl/test-e2e/Regression/multithread_write_accessor.cpp @@ -21,7 +21,7 @@ class KernelA; void threadFunction(sycl::buffer &Buf) { sycl::queue Q; Q.submit([&](sycl::handler &Cgh) { - auto Acc = Buf.get_access(Cgh); + auto Acc = Buf.get_access(Cgh); Cgh.single_task([=]() { Acc[0] += 1; }); }); } diff --git a/sycl/test-e2e/Regression/queue_submit.cpp b/sycl/test-e2e/Regression/queue_submit.cpp index 3fca0f9fa2063..33d90586904ad 100755 --- a/sycl/test-e2e/Regression/queue_submit.cpp +++ b/sycl/test-e2e/Regression/queue_submit.cpp @@ -24,7 +24,7 @@ void submit(sycl::queue &queue, sycl::kernel &kernel) { sycl::buffer result_buf{data, sycl::range<1>{DIMS}}; queue.submit([&](sycl::handler &cgh) { auto result_acc = - result_buf.get_access(cgh); + result_buf.get_access(cgh); cgh.set_arg(0, result_acc); cgh.parallel_for(sycl::range<1>{DIMS}, kernel); }); @@ -51,7 +51,7 @@ void run_test(size_t numThreads) { queue .submit([&](sycl::handler &cgh) { auto result_acc = - result_buf.get_access(cgh); + result_buf.get_access(cgh); cgh.parallel_for( sycl::range<1>{DIMS}, [=](sycl::id<1> idx) { result_acc[idx] = idx[0]; }); diff --git a/sycl/test-e2e/Regression/set-arg-local-accessor.cpp b/sycl/test-e2e/Regression/set-arg-local-accessor.cpp index bd630da4343e3..09762d5f4e701 100644 --- a/sycl/test-e2e/Regression/set-arg-local-accessor.cpp +++ b/sycl/test-e2e/Regression/set-arg-local-accessor.cpp @@ -53,7 +53,7 @@ int main() { buffer b(init, N_wg); Q.submit([&](handler &cgh) { - auto acc_global = b.get_access(cgh); + auto acc_global = b.get_access(cgh); local_accessor acc_local(N_slm, cgh); cgh.set_arg(0, acc_global); diff --git a/sycl/test-e2e/Sampler/basic-rw.cpp b/sycl/test-e2e/Sampler/basic-rw.cpp index 66789e0d560a8..71f3b913fd293 100644 --- a/sycl/test-e2e/Sampler/basic-rw.cpp +++ b/sycl/test-e2e/Sampler/basic-rw.cpp @@ -47,7 +47,7 @@ void test_rw(image_channel_order ChanOrder, image_channel_type ChanType) { // - create an image image<1> image_1D(ChanOrder, ChanType, ImgRange_1D); event E_Setup = Q.submit([&](handler &cgh) { - auto image_acc = image_1D.get_access(cgh); + auto image_acc = image_1D.get_access(cgh); cgh.single_task([=]() { image_acc.write(0, leftEdge); image_acc.write(1, body); @@ -61,8 +61,8 @@ void test_rw(image_channel_order ChanOrder, image_channel_type ChanType) { buffer testResults((range<1>(numTests))); event E_Test = Q.submit([&](handler &cgh) { - auto image_acc = image_1D.get_access(cgh); - auto test_acc = testResults.get_access(cgh); + auto image_acc = image_1D.get_access(cgh); + auto test_acc = testResults.get_access(cgh); cgh.single_task([=]() { int i = 0; // the index for writing into the testResult buffer. diff --git a/sycl/test-e2e/Sampler/normalized-clampedge-linear-float.cpp b/sycl/test-e2e/Sampler/normalized-clampedge-linear-float.cpp index 8c9605f52a90f..1a7d2d9312be4 100644 --- a/sycl/test-e2e/Sampler/normalized-clampedge-linear-float.cpp +++ b/sycl/test-e2e/Sampler/normalized-clampedge-linear-float.cpp @@ -61,7 +61,7 @@ void test_normalized_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); @@ -79,8 +79,8 @@ void test_normalized_clampedge_linear_sampler(image_channel_order ChanOrder, sampler(normalized, 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/normalized-mirror-nearest.cpp b/sycl/test-e2e/Sampler/normalized-mirror-nearest.cpp index 94eb84064ca8b..530bcbcee47e3 100644 --- a/sycl/test-e2e/Sampler/normalized-mirror-nearest.cpp +++ b/sycl/test-e2e/Sampler/normalized-mirror-nearest.cpp @@ -60,7 +60,7 @@ void test_normalized_mirrored_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); @@ -77,8 +77,8 @@ void test_normalized_mirrored_nearest_sampler(image_channel_order ChanOrder, auto Norm_Mirror_Nearest_sampler = sampler(normalized, mirrored, 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-repeat-linear-float.cpp b/sycl/test-e2e/Sampler/normalized-repeat-linear-float.cpp index c4155118c68da..55fb376516c5a 100644 --- a/sycl/test-e2e/Sampler/normalized-repeat-linear-float.cpp +++ b/sycl/test-e2e/Sampler/normalized-repeat-linear-float.cpp @@ -64,7 +64,7 @@ void test_normalized_repeat_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); @@ -81,8 +81,8 @@ void test_normalized_repeat_linear_sampler(image_channel_order ChanOrder, auto Norm_Repeat_Linear_sampler = sampler(normalized, repeat, 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-clamp-nearest.cpp b/sycl/test-e2e/Sampler/unnormalized-clamp-nearest.cpp index 21b79344bca9d..6eb32879e72fe 100644 --- a/sycl/test-e2e/Sampler/unnormalized-clamp-nearest.cpp +++ b/sycl/test-e2e/Sampler/unnormalized-clamp-nearest.cpp @@ -46,7 +46,7 @@ void test_unnormalized_clamp_nearest_sampler(image_channel_order ChanOrder, // - create an image image<1> image_1D(ChanOrder, ChanType, ImgRange_1D); event E_Setup = Q.submit([&](handler &cgh) { - auto image_acc = image_1D.get_access(cgh); + auto image_acc = image_1D.get_access(cgh); cgh.single_task([=]() { image_acc.write(0, leftEdge); image_acc.write(1, body); @@ -64,8 +64,8 @@ void test_unnormalized_clamp_nearest_sampler(image_channel_order ChanOrder, sampler(unnormalized, addressing_mode::clamp, nearest); event E_Test = Q.submit([&](handler &cgh) { - auto image_acc = image_1D.get_access(cgh); - auto test_acc = testResults.get_access(cgh); + auto image_acc = image_1D.get_access(cgh); + auto test_acc = testResults.get_access(cgh); cgh.single_task([=]() { int i = 0; // the index for writing into the testResult buffer. diff --git a/sycl/test-e2e/Sampler/unnormalized-none-linear-float.cpp b/sycl/test-e2e/Sampler/unnormalized-none-linear-float.cpp index e51bc4ba1b8ad..dcb604fed9b0b 100644 --- a/sycl/test-e2e/Sampler/unnormalized-none-linear-float.cpp +++ b/sycl/test-e2e/Sampler/unnormalized-none-linear-float.cpp @@ -70,7 +70,7 @@ void test_unnormalized_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); @@ -87,8 +87,8 @@ void test_unnormalized_none_linear_sampler(image_channel_order ChanOrder, auto UnNorm_None_Linear_sampler = sampler(unnormalized, 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/Scheduler/CommandCleanupThreadSafety.cpp b/sycl/test-e2e/Scheduler/CommandCleanupThreadSafety.cpp index e112428e2ddb6..24e01f7f727a0 100644 --- a/sycl/test-e2e/Scheduler/CommandCleanupThreadSafety.cpp +++ b/sycl/test-e2e/Scheduler/CommandCleanupThreadSafety.cpp @@ -16,7 +16,7 @@ class Foo; event submitTask(queue &Q, buffer &Buf) { return Q.submit([&](handler &Cgh) { - auto Acc = Buf.get_access(Cgh); + auto Acc = Buf.get_access(Cgh); Cgh.single_task([=]() { Acc[0] = 42; }); }); } diff --git a/sycl/test-e2e/Scheduler/HostAccDestruction.cpp b/sycl/test-e2e/Scheduler/HostAccDestruction.cpp index 30e98974b5f1d..84a12be35bdd9 100644 --- a/sycl/test-e2e/Scheduler/HostAccDestruction.cpp +++ b/sycl/test-e2e/Scheduler/HostAccDestruction.cpp @@ -20,7 +20,7 @@ int main() { sycl::queue q; auto host_acc = buf.get_host_access(); q.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); cgh.parallel_for( sycl::range<1>{size}, [=](sycl::id<1> id) { (void)acc[id]; }); }); diff --git a/sycl/test-e2e/Scheduler/MultipleDevices.cpp b/sycl/test-e2e/Scheduler/MultipleDevices.cpp index f848ce79ab9a4..2974c200c818b 100644 --- a/sycl/test-e2e/Scheduler/MultipleDevices.cpp +++ b/sycl/test-e2e/Scheduler/MultipleDevices.cpp @@ -24,21 +24,21 @@ int multidevice_test(queue MyQueue1, queue MyQueue2) { buffer BufD(range<1>{N}); MyQueue1.submit([&](handler &cgh) { - auto A = BufA.get_access(cgh); + auto A = BufA.get_access(cgh); cgh.parallel_for( range<1>{N}, [=](id<1> index) { A[index[0]] = index[0]; }); }); MyQueue2.submit([&](handler &cgh) { - auto B = BufB.get_access(cgh); + auto B = BufB.get_access(cgh); cgh.parallel_for( range<1>{N}, [=](id<1> index) { B[index[0]] = N - index[0]; }); }); MyQueue2.submit([&](handler &cgh) { - auto A = BufA.get_access(cgh); - auto B = BufB.get_access(cgh); - auto C = BufC.get_access(cgh); + auto A = BufA.get_access(cgh); + auto B = BufB.get_access(cgh); + auto C = BufC.get_access(cgh); cgh.parallel_for(range<1>{N}, [=](id<1> index) { B[index[0]] = B[index[0]] + A[index[0]]; C[index[0]] = B[index[0]] - index[0]; @@ -46,15 +46,15 @@ int multidevice_test(queue MyQueue1, queue MyQueue2) { }); MyQueue2.submit([&](handler &cgh) { - auto D = BufD.get_access(cgh); + auto D = BufD.get_access(cgh); cgh.parallel_for(range<1>{N}, [=](id<1> index) { D[index[0]] = 1; }); }); MyQueue1.submit([&](handler &cgh) { - auto B = BufB.get_access(cgh); - auto C = BufC.get_access(cgh); - auto D = BufD.get_access(cgh); + auto B = BufB.get_access(cgh); + auto C = BufC.get_access(cgh); + auto D = BufD.get_access(cgh); cgh.parallel_for(range<1>{N}, [=](id<1> index) { D[index[0]] = D[index[0]] + B[index[0]] - C[index[0]]; }); diff --git a/sycl/test-e2e/SpecConstants/2020/host_apis.cpp b/sycl/test-e2e/SpecConstants/2020/host_apis.cpp index 5a99b61c4941a..7b072f425ec09 100644 --- a/sycl/test-e2e/SpecConstants/2020/host_apis.cpp +++ b/sycl/test-e2e/SpecConstants/2020/host_apis.cpp @@ -60,7 +60,7 @@ int main() { sycl::buffer Buf{sycl::range{1}}; Q.submit([&](sycl::handler &CGH) { CGH.use_kernel_bundle(ExecBundle); - auto Acc = Buf.get_access(CGH); + auto Acc = Buf.get_access(CGH); CGH.single_task([=](sycl::kernel_handler KH) { Acc[0] = KH.get_specialization_constant(); }); @@ -72,7 +72,7 @@ int main() { { sycl::buffer Buf{sycl::range{1}}; Q.submit([&](sycl::handler &CGH) { - auto Acc = Buf.get_access(CGH); + auto Acc = Buf.get_access(CGH); CGH.set_specialization_constant(TestStruct{1, 2}); const auto SC = CGH.get_specialization_constant(); assert(SC == 42); diff --git a/sycl/test-e2e/SpecConstants/2020/non_native/spec_const_alignment_test.cpp b/sycl/test-e2e/SpecConstants/2020/non_native/spec_const_alignment_test.cpp index b10a5ced01975..6d007586498f7 100644 --- a/sycl/test-e2e/SpecConstants/2020/non_native/spec_const_alignment_test.cpp +++ b/sycl/test-e2e/SpecConstants/2020/non_native/spec_const_alignment_test.cpp @@ -95,10 +95,10 @@ int main() { } } - auto acc0 = b0.template get_access(cgh); - auto acc1 = b1.template get_access(cgh); - auto acc2 = b2.template get_access(cgh); - auto accc = bc.template get_access(cgh); + auto acc0 = b0.template get_access(cgh); + auto acc1 = b1.template get_access(cgh); + auto acc2 = b2.template get_access(cgh); + auto accc = bc.template get_access(cgh); cgh.single_task([=](sycl::kernel_handler kh) { acc0[0] = kh.get_specialization_constant(); diff --git a/sycl/test-e2e/SubGroup/broadcast.hpp b/sycl/test-e2e/SubGroup/broadcast.hpp index a4a9736aa2eac..b2ca9d644a2ac 100644 --- a/sycl/test-e2e/SubGroup/broadcast.hpp +++ b/sycl/test-e2e/SubGroup/broadcast.hpp @@ -17,8 +17,8 @@ template void check(queue &Queue) { buffer syclbuf(G); buffer sgsizebuf(1); Queue.submit([&](handler &cgh) { - auto syclacc = syclbuf.template get_access(cgh); - auto sgsizeacc = sgsizebuf.get_access(cgh); + auto syclacc = syclbuf.template get_access(cgh); + auto sgsizeacc = sgsizebuf.get_access(cgh); cgh.parallel_for>(NdRange, [=](nd_item<1> NdItem) { sycl::sub_group SG = NdItem.get_sub_group(); /*Broadcast GID of element with SGLID == SGID % SGMLR*/ diff --git a/sycl/test-e2e/SubGroup/info.cpp b/sycl/test-e2e/SubGroup/info.cpp index bc97ad8959fc4..cd80970096f7a 100644 --- a/sycl/test-e2e/SubGroup/info.cpp +++ b/sycl/test-e2e/SubGroup/info.cpp @@ -34,9 +34,9 @@ int main() { buffer ABuf{GlobalRange}, BBuf{GlobalRange}, CBuf{GlobalRange}; Queue.submit([&](sycl::handler &cgh) { - auto A = ABuf.get_access(cgh); - auto B = BBuf.get_access(cgh); - auto C = CBuf.get_access(cgh); + auto A = ABuf.get_access(cgh); + auto B = BBuf.get_access(cgh); + auto C = CBuf.get_access(cgh); cgh.parallel_for( nd_range<2>(GlobalRange, range<2>(10, 20)), [=](nd_item<2> index) { const id<2> GlobalID = index.get_global_id(); diff --git a/sycl/test-e2e/SubGroup/vote.cpp b/sycl/test-e2e/SubGroup/vote.cpp index 5fb9ac0066194..37dd2d025baaa 100644 --- a/sycl/test-e2e/SubGroup/vote.cpp +++ b/sycl/test-e2e/SubGroup/vote.cpp @@ -22,8 +22,8 @@ void check(queue Queue, const int G, const int L, const int D, const int R) { // Initialise buffer with zeros Queue.submit([&](handler &cgh) { - auto sganyacc = sganybuf.get_access(cgh); - auto sgallacc = sgallbuf.get_access(cgh); + auto sganyacc = sganybuf.get_access(cgh); + auto sgallacc = sgallbuf.get_access(cgh); cgh.parallel_for(range<1>{(unsigned)G}, [=](id<1> index) { sganyacc[index] = 0; sgallacc[index] = 0; @@ -31,8 +31,8 @@ void check(queue Queue, const int G, const int L, const int D, const int R) { }); Queue.submit([&](handler &cgh) { - auto sganyacc = sganybuf.get_access(cgh); - auto sgallacc = sgallbuf.get_access(cgh); + auto sganyacc = sganybuf.get_access(cgh); + auto sgallacc = sgallbuf.get_access(cgh); cgh.parallel_for(NdRange, [=](nd_item<1> NdItem) { sganyacc[NdItem.get_global_id()] = 0; sgallacc[NdItem.get_global_id()] = 0; @@ -40,8 +40,8 @@ void check(queue Queue, const int G, const int L, const int D, const int R) { }); Queue.submit([&](handler &cgh) { - auto sganyacc = sganybuf.get_access(cgh); - auto sgallacc = sgallbuf.get_access(cgh); + auto sganyacc = sganybuf.get_access(cgh); + auto sgallacc = sgallbuf.get_access(cgh); cgh.parallel_for(NdRange, [=](nd_item<1> NdItem) { sycl::sub_group SG = NdItem.get_sub_group(); /* Set to 1 if any local ID in subgroup devided by D has remainder R */ diff --git a/sycl/test-e2e/Tracing/buffer_printers.cpp b/sycl/test-e2e/Tracing/buffer_printers.cpp index d712b0009c1ae..4905526b90631 100644 --- a/sycl/test-e2e/Tracing/buffer_printers.cpp +++ b/sycl/test-e2e/Tracing/buffer_printers.cpp @@ -27,9 +27,9 @@ int main() { buffer DstBuffer(DstData.data(), range<2>(Height, Width)); Queue.submit([&](handler &CGH) { - auto Read = SrcBuffer.get_access( + auto Read = SrcBuffer.get_access( CGH, range<2>(Height, Width), id(Height, Width)); - auto Write = DstBuffer.get_access(CGH); + auto Write = DstBuffer.get_access(CGH); CGH.copy(Read, Write); }); } @@ -44,7 +44,7 @@ int main() { sycl::buffer SubBuf(Buf, Size / 2, 8); Queue.submit([&](sycl::handler &CGH) { - auto Acc = SubBuf.get_access(CGH); + auto Acc = SubBuf.get_access(CGH); CGH.single_task([=]() {}); }); } diff --git a/sycl/test-e2e/USM/free_during_kernel_execution.cpp b/sycl/test-e2e/USM/free_during_kernel_execution.cpp index 8fd28516f9036..d0ce5715da3b7 100644 --- a/sycl/test-e2e/USM/free_during_kernel_execution.cpp +++ b/sycl/test-e2e/USM/free_during_kernel_execution.cpp @@ -23,7 +23,7 @@ int main() { auto *USM = sycl::malloc_host(1, Queue.get_context()); Queue.submit([&](sycl::handler &cgh) { - auto Accessor = Buffer.get_access(cgh); + auto Accessor = Buffer.get_access(cgh); cgh.single_task([=]() { for (int I = 0; I < N; ++I) { Accessor[I] = I; diff --git a/sycl/test-e2e/USM/memset.cpp b/sycl/test-e2e/USM/memset.cpp index 3f80088dc12b5..e073a979ac347 100644 --- a/sycl/test-e2e/USM/memset.cpp +++ b/sycl/test-e2e/USM/memset.cpp @@ -84,7 +84,7 @@ int main() { { buffer buf{&out[0], range<1>{N}}; q.submit([&](handler &h) { - auto acc = buf.template get_access(h); + auto acc = buf.template get_access(h); h.parallel_for( range<1>(N), [=](id<1> item) { acc[item] = array[item]; }); }); @@ -109,7 +109,7 @@ int main() { { buffer buf{&out[0], range<1>{N}}; q.submit([&](handler &h) { - auto acc = buf.template get_access(h); + auto acc = buf.template get_access(h); h.parallel_for( range<1>(N), [=](id<1> item) { acc[item] = array[item]; }); }); diff --git a/sycl/test-e2e/WorkGroupScratchMemory/dynamic_alloc_ptr_alias.cpp b/sycl/test-e2e/WorkGroupScratchMemory/dynamic_alloc_ptr_alias.cpp index 0572e699e145e..3bdd1d72a99d7 100644 --- a/sycl/test-e2e/WorkGroupScratchMemory/dynamic_alloc_ptr_alias.cpp +++ b/sycl/test-e2e/WorkGroupScratchMemory/dynamic_alloc_ptr_alias.cpp @@ -63,7 +63,7 @@ int main() { buffer Buf{Vec.data(), range<1>(Size)}; Q.submit([&](handler &Cgh) { - auto Acc = Buf.get_access(Cgh); + auto Acc = Buf.get_access(Cgh); sycl_ext::work_group_scratch_size static_size(WgSize * RepeatWG * sizeof(int)); sycl_ext::properties properties{static_size}; diff --git a/sycl/test-e2e/WorkGroupStatic/work_group_static_memory_class_scope.cpp b/sycl/test-e2e/WorkGroupStatic/work_group_static_memory_class_scope.cpp index 64920b8a78719..f341885a789f1 100644 --- a/sycl/test-e2e/WorkGroupStatic/work_group_static_memory_class_scope.cpp +++ b/sycl/test-e2e/WorkGroupStatic/work_group_static_memory_class_scope.cpp @@ -50,7 +50,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) { LocalMem::localIDBuff[Item.get_local_linear_id()] = @@ -74,7 +74,7 @@ int main() { buffer Buf{Vec.data(), range<1>(Size)}; Q.submit([&](handler &Cgh) { - auto Acc = Buf.get_access(Cgh); + auto Acc = Buf.get_access(Cgh); Cgh.parallel_for( nd_range<1>(range<1>(Size), range<1>(WgSize)), [=](nd_item<1> Item) { sycl::ext::oneapi::experimental::work_group_static localIDBuff; diff --git a/sycl/test-e2e/XPTI/buffer/host_array.cpp b/sycl/test-e2e/XPTI/buffer/host_array.cpp index b2c56dd14528d..2ec80a6764588 100644 --- a/sycl/test-e2e/XPTI/buffer/host_array.cpp +++ b/sycl/test-e2e/XPTI/buffer/host_array.cpp @@ -24,7 +24,7 @@ int main() { // CHECK:{{[0-9]+}}|Associate buffer|[[USERID1]]|[[BEID1:.*]] Queue.submit([&](sycl::handler &cgh) { // Get write only access to the buffer on a device. - auto Accessor1 = Buffer1.get_access(cgh); + auto Accessor1 = Buffer1.get_access(cgh); // Execute kernel. cgh.parallel_for(NumOfWorkItems, [=](sycl::id<1> WIid) { Accessor1[WIid] = static_cast(WIid.get(0)); diff --git a/sycl/test-e2e/XPTI/buffer/multiple_queues.cpp b/sycl/test-e2e/XPTI/buffer/multiple_queues.cpp index 434fed2206019..b110a932a272e 100644 --- a/sycl/test-e2e/XPTI/buffer/multiple_queues.cpp +++ b/sycl/test-e2e/XPTI/buffer/multiple_queues.cpp @@ -29,7 +29,7 @@ int main() { Queue1.submit([&](sycl::handler &cgh) { // CHECK: {{[0-9]+}}|Construct accessor|[[USERID1]]|[[ACCID1:.*]]|2014|1025|{{.*}}multiple_queues.cpp:[[# @LINE + 1]]:32 - auto Accessor1 = Buffer1.get_access(cgh); + auto Accessor1 = Buffer1.get_access(cgh); // CHECK:{{[0-9]+}}|Associate buffer|[[USERID1]]|[[BEID1:.*]] cgh.parallel_for(NumOfWorkItems, [=](sycl::id<1> WIid) { Accessor1[WIid] = static_cast(WIid.get(0)); @@ -39,7 +39,7 @@ int main() { Queue2.submit([&](sycl::handler &cgh) { // CHECK: {{[0-9]+}}|Construct accessor|[[USERID1]]|[[ACCID2:.*]]|2014|1025|{{.*}}multiple_queues.cpp:[[# @LINE + 1]]:32 - auto Accessor1 = Buffer1.get_access(cgh); + auto Accessor1 = Buffer1.get_access(cgh); // CHECK:{{[0-9]+}}|Associate buffer|[[USERID1]]|[[BEID2:.*]] cgh.parallel_for(NumOfWorkItems, [=](sycl::id<1> WIid) { Accessor1[WIid] *= static_cast(WIid.get(0)); diff --git a/sycl/test-e2e/XPTI/buffer/use_host_ptr.cpp b/sycl/test-e2e/XPTI/buffer/use_host_ptr.cpp index a06e34fe878b1..2b4ecae0a373f 100644 --- a/sycl/test-e2e/XPTI/buffer/use_host_ptr.cpp +++ b/sycl/test-e2e/XPTI/buffer/use_host_ptr.cpp @@ -25,7 +25,7 @@ int main() { // CHECK:{{[0-9]+}}|Associate buffer|[[USERID1]]|[[BEID1:.*]] Queue.submit([&](sycl::handler &cgh) { // Get write only access to the buffer on a device. - auto Accessor1 = Buffer1.get_access(cgh); + auto Accessor1 = Buffer1.get_access(cgh); // Execute kernel. cgh.parallel_for(NumOfWorkItems, [=](sycl::id<1> WIid) { Accessor1[WIid] = static_cast(WIid.get(0));