Skip to content

[SYCL] Update ESIMD tests to use access_mode #22821

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
KornevNikita merged 2 commits into from
Jul 31, 2026
Merged

[SYCL] Update ESIMD tests to use access_mode #22821

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
5182fe9
[SYCL] Update ESIMD tests to use access_mode
KornevNikita Jul 30, 2026
244dc24
update InvokeSimd tests
KornevNikita Jul 30, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 3 additions & 3 deletions sycl/test-e2e/ESIMD/BitonicSortK.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -604,8 +604,8 @@ int BitonicSort::Solve(uint32_t *pInputs, uint32_t *pOutputs, uint32_t size) {
buffer<uint32_t, 1> bufo(pOutputs, range<1>(size));
// enqueue sort265 kernel
auto e = pQueue_->submit([&](handler &cgh) {
auto acci = bufi.get_access<access::mode::read>(cgh);
auto acco = bufo.get_access<access::mode::write>(cgh);
auto acci = bufi.get_access<access_mode::read>(cgh);
auto acco = bufo.get_access<access_mode::write>(cgh);
cgh.parallel_for<class Sort256>(
SortGlobalRange * SortLocalRange, [=](id<1> i) SYCL_ESIMD_KERNEL {
using namespace sycl::ext::intel::esimd;
Expand Down Expand Up @@ -649,7 +649,7 @@ int BitonicSort::Solve(uint32_t *pInputs, uint32_t *pOutputs, uint32_t size) {
for (int j = i; j >= 8; j--) {
buffer<uint32_t, 1> buf(pOutputs, range<1>(size));
mergeEvent[k] = pQueue_->submit([&](handler &cgh) {
auto acc = buf.get_access<access::mode::read_write>(cgh);
auto acc = buf.get_access<access_mode::read_write>(cgh);
cgh.parallel_for<class Merge>(
MergeGlobalRange * MergeLocalRange,
[=](id<1> tid) SYCL_ESIMD_KERNEL {
Expand Down
6 changes: 3 additions & 3 deletions sycl/test-e2e/ESIMD/InlineAsm/asm_glb.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -48,9 +48,9 @@ int main(void) {
<< "\n";

auto e = q.submit([&](handler &cgh) {
auto PA = bufa.get_access<access::mode::read>(cgh);
auto PB = bufb.get_access<access::mode::read>(cgh);
auto PC = bufc.get_access<access::mode::write>(cgh);
auto PA = bufa.get_access<access_mode::read>(cgh);
auto PB = bufb.get_access<access_mode::read>(cgh);
auto PC = bufc.get_access<access_mode::write>(cgh);
cgh.parallel_for<class Test>(
GlobalRange * LocalRange, [=](id<1> i) SYCL_ESIMD_KERNEL {
using namespace sycl::ext::intel::esimd;
Expand Down
6 changes: 3 additions & 3 deletions sycl/test-e2e/ESIMD/InlineAsm/asm_simd_mask.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -44,9 +44,9 @@ int main(void) {
std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";

auto e = q.submit([&](handler &cgh) {
auto PA = bufa.get_access<access::mode::read>(cgh);
auto PB = bufb.get_access<access::mode::read>(cgh);
auto PC = bufc.get_access<access::mode::write>(cgh);
auto PA = bufa.get_access<access_mode::read>(cgh);
auto PB = bufb.get_access<access_mode::read>(cgh);
auto PC = bufc.get_access<access_mode::write>(cgh);
cgh.parallel_for<class Test>(
GlobalRange * LocalRange, [=](id<1> i) SYCL_ESIMD_KERNEL {
using namespace sycl::ext::intel::esimd;
Expand Down
6 changes: 3 additions & 3 deletions sycl/test-e2e/ESIMD/InlineAsm/asm_simd_view.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -44,9 +44,9 @@ int main(void) {
std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";

auto e = q.submit([&](handler &cgh) {
auto PA = bufa.get_access<access::mode::read>(cgh);
auto PB = bufb.get_access<access::mode::read>(cgh);
auto PC = bufc.get_access<access::mode::write>(cgh);
auto PA = bufa.get_access<access_mode::read>(cgh);
auto PB = bufb.get_access<access_mode::read>(cgh);
auto PC = bufc.get_access<access_mode::write>(cgh);
cgh.parallel_for<class Test>(
GlobalRange * LocalRange, [=](id<1> i) SYCL_ESIMD_KERNEL {
using namespace sycl::ext::intel::esimd;
Expand Down
6 changes: 3 additions & 3 deletions sycl/test-e2e/ESIMD/InlineAsm/asm_vadd.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -44,9 +44,9 @@ int main(void) {
std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";

auto e = q.submit([&](handler &cgh) {
auto PA = bufa.get_access<access::mode::read>(cgh);
auto PB = bufb.get_access<access::mode::read>(cgh);
auto PC = bufc.get_access<access::mode::write>(cgh);
auto PA = bufa.get_access<access_mode::read>(cgh);
auto PB = bufb.get_access<access_mode::read>(cgh);
auto PC = bufc.get_access<access_mode::write>(cgh);
cgh.parallel_for<class Test>(
GlobalRange * LocalRange, [=](id<1> i) SYCL_ESIMD_KERNEL {
using namespace sycl::ext::intel::esimd;
Expand Down
4 changes: 2 additions & 2 deletions sycl/test-e2e/ESIMD/acc_gather_scatter_rgba.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -109,8 +109,8 @@ bool test(queue q) {
buffer<T, 1> OutBuf(B, range<1>(size));
range<1> glob_range{numWorkItems};
auto e = q.submit([&](handler &cgh) {
auto InAcc = InBuf.template get_access<access::mode::read_write>(cgh);
auto OutAcc = OutBuf.template get_access<access::mode::read_write>(cgh);
auto InAcc = InBuf.template get_access<access_mode::read_write>(cgh);
auto OutAcc = OutBuf.template get_access<access_mode::read_write>(cgh);
Kernel<T, VL, STRIDE, CH_MASK> kernel(InAcc, OutAcc);
cgh.parallel_for(glob_range, kernel);
});
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/ESIMD/acc_gather_scatter_rgba_stateless_64.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -39,7 +39,7 @@ int main(void) {
<< "\n";
try {
q.submit([&](handler &cgh) {
auto PA = bufa.get_access<access::mode::read_write>(cgh);
auto PA = bufa.get_access<access_mode::read_write>(cgh);
cgh.single_task<class Test>([=]() SYCL_ESIMD_KERNEL {
uint64_t offsetStart = (Size - VL) * sizeof(uint64_t);
simd<uint64_t, VL> offset(offsetStart, sizeof(uint64_t));
Expand Down
4 changes: 2 additions & 2 deletions sycl/test-e2e/ESIMD/accessor.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -36,8 +36,8 @@ int main() {
std::cout << "Running on "
<< q.get_device().get_info<sycl::info::device::name>() << "\n";

auto acc0 = buf0.get_access<access::mode::read_write>(cgh);
auto acc1 = buf1.get_access<access::mode::write>(cgh);
auto acc0 = buf0.get_access<access_mode::read_write>(cgh);
auto acc1 = buf1.get_access<access_mode::write>(cgh);

cgh.parallel_for<class Test>(range<1>(1),
[=](sycl::id<1> i) SYCL_ESIMD_KERNEL {
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/ESIMD/accessor_gather_scatter.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -57,7 +57,7 @@ template <typename T, unsigned VL, unsigned STRIDE> bool test(queue q) {
range<1> glob_range{size / VL};

auto e = q.submit([&](handler &cgh) {
auto acc = buf.template get_access<access::mode::read_write>(cgh);
auto acc = buf.template get_access<access_mode::read_write>(cgh);
Kernel<T, VL, STRIDE> kernel(acc);
cgh.parallel_for(glob_range, kernel);
});
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/ESIMD/accessor_gather_scatter_stateless_64.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -34,7 +34,7 @@ int main(void) {
<< "\n";
try {
q.submit([&](handler &cgh) {
auto PA = bufa.get_access<access::mode::read_write>(cgh);
auto PA = bufa.get_access<access_mode::read_write>(cgh);
cgh.single_task<class Test>([=]() SYCL_ESIMD_KERNEL {
uint64_t offsetStart = (Size - VL) * sizeof(uint64_t);
simd<uint64_t, VL> offset(offsetStart, sizeof(uint64_t));
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/ESIMD/accessor_load_store.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -50,7 +50,7 @@ template <typename T> bool test(queue q, size_t size) {
range<1> glob_range{size};

auto e = q.submit([&](handler &cgh) {
auto acc = buf.template get_access<access::mode::read_write>(cgh);
auto acc = buf.template get_access<access_mode::read_write>(cgh);
Kernel<T> kernel(acc);
cgh.parallel_for(glob_range, kernel);
});
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/ESIMD/accessor_load_store_stateless_64.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -34,7 +34,7 @@ int main(void) {
<< "\n";
try {
q.submit([&](handler &cgh) {
auto PA = bufa.get_access<access::mode::read_write>(cgh);
auto PA = bufa.get_access<access_mode::read_write>(cgh);
cgh.single_task<class Test>([=]() SYCL_ESIMD_KERNEL {
uint64_t offset = (Size - VL) * sizeof(uint64_t);
simd<uint64_t, VL> va = block_load<uint64_t, VL>(PA, 0);
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/ESIMD/accessor_stateless_64.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -36,7 +36,7 @@ int main(void) {
<< "\n";

q.submit([&](handler &cgh) {
auto PA = bufa.get_access<access::mode::read_write>(cgh);
auto PA = bufa.get_access<access_mode::read_write>(cgh);
cgh.single_task<class Test>([=]() SYCL_ESIMD_KERNEL {
uint64_t offset = (Size - VL) * sizeof(uint64_t);
simd<uint64_t, VL> va;
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/ESIMD/accessor_stateless_ctor_64.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -36,7 +36,7 @@ int main(void) {
<< "\n";

q.submit([&](handler &cgh) {
auto PA = bufa.get_access<access::mode::read_write>(cgh);
auto PA = bufa.get_access<access_mode::read_write>(cgh);
cgh.single_task<class Test>([=]() SYCL_ESIMD_KERNEL {
uint64_t offset = (Size - VL) * sizeof(uint64_t);
simd<uint64_t, VL> va(PA, 0);
Expand Down
12 changes: 6 additions & 6 deletions sycl/test-e2e/ESIMD/aot_mixed.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -70,9 +70,9 @@ bool test_esimd(queue q) {
buffer<float, 1> bufc(C, range<1>(Size));

auto e = q.submit([&](handler &cgh) {
auto PA = bufa.get_access<access::mode::read>(cgh);
auto PB = bufb.get_access<access::mode::read>(cgh);
auto PC = bufc.get_access<access::mode::write>(cgh);
auto PA = bufa.get_access<access_mode::read>(cgh);
auto PB = bufb.get_access<access_mode::read>(cgh);
auto PC = bufc.get_access<access_mode::write>(cgh);
cgh.parallel_for<class TestESIMD>(
Size / VL, [=](id<1> i) SYCL_ESIMD_KERNEL {
using namespace sycl::ext::intel::esimd;
Expand Down Expand Up @@ -122,9 +122,9 @@ bool test_sycl(queue q) {
buffer<float, 1> bufc(C, range<1>(Size));

auto e = q.submit([&](handler &cgh) {
auto PA = bufa.get_access<access::mode::read>(cgh);
auto PB = bufb.get_access<access::mode::read>(cgh);
auto PC = bufc.get_access<access::mode::write>(cgh);
auto PA = bufa.get_access<access_mode::read>(cgh);
auto PB = bufb.get_access<access_mode::read>(cgh);
auto PC = bufc.get_access<access_mode::write>(cgh);
cgh.parallel_for<class TestSYCL>(Size,
[=](id<1> i) { PC[i] = PA[i] + PB[i]; });
});
Expand Down
4 changes: 2 additions & 2 deletions sycl/test-e2e/ESIMD/api/ballot.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -38,8 +38,8 @@ template <class T, int N> bool test(queue &Q) {
buffer<decltype(Res), 1> RB(&Res, range<1>(1));

auto E = Q.submit([&](handler &CGH) {
auto In = PB.template get_access<access::mode::read>(CGH);
auto Out = RB.template get_access<access::mode::write>(CGH);
auto In = PB.template get_access<access_mode::read>(CGH);
auto Out = RB.template get_access<access_mode::write>(CGH);

CGH.parallel_for(sycl::range<1>{1}, [=](id<1>) SYCL_ESIMD_KERNEL {
simd<T, N> Mask;
Expand Down
6 changes: 3 additions & 3 deletions sycl/test-e2e/ESIMD/api/functional/operators/operator_assignment_glb.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -50,9 +50,9 @@ int main(void) {
<< "\n";

auto e = q.submit([&](handler &cgh) {
auto PA = bufa.get_access<access::mode::read>(cgh);
auto PB = bufb.get_access<access::mode::read>(cgh);
auto PC = bufc.get_access<access::mode::write>(cgh);
auto PA = bufa.get_access<access_mode::read>(cgh);
auto PB = bufb.get_access<access_mode::read>(cgh);
auto PC = bufc.get_access<access_mode::write>(cgh);
cgh.parallel_for<class Test>(
GlobalRange * LocalRange, [=](id<1> i) SYCL_ESIMD_KERNEL {
using namespace sycl::ext::intel::esimd;
Expand Down
4 changes: 2 additions & 2 deletions sycl/test-e2e/ESIMD/api/functional/operators/operator_assignment_glb_mask.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -48,8 +48,8 @@ int main(void) {
<< "\n";

auto e = q.submit([&](handler &cgh) {
auto PA = bufa.get_access<access::mode::read>(cgh);
auto PB = bufb.get_access<access::mode::write>(cgh);
auto PA = bufa.get_access<access_mode::read>(cgh);
auto PB = bufb.get_access<access_mode::write>(cgh);
cgh.parallel_for<class Test>(
GlobalRange * LocalRange, [=](id<1> i) SYCL_ESIMD_KERNEL {
using namespace sycl::ext::intel::esimd;
Expand Down
4 changes: 2 additions & 2 deletions sycl/test-e2e/ESIMD/api/replicate_smoke.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -98,8 +98,8 @@ bool test_impl(queue q, int offset, T (&&gold)[N]) {
sycl::buffer<T, 1> dst_buf(dm.dst, VL);

q.submit([&](handler &cgh) {
auto src_acc = src_buf.template get_access<access::mode::read>(cgh);
auto dst_acc = dst_buf.template get_access<access::mode::write>(cgh);
auto src_acc = src_buf.template get_access<access_mode::read>(cgh);
auto dst_acc = dst_buf.template get_access<access_mode::write>(cgh);

cgh.single_task([=]() SYCL_ESIMD_KERNEL {
simd<T, VL> src(src_acc, 0);
Expand Down
4 changes: 2 additions & 2 deletions sycl/test-e2e/ESIMD/api/saturation_smoke.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -83,8 +83,8 @@ bool test(queue q) {
sycl::buffer<To, 1> dst_buf(dm.dst, N);

auto e = q.submit([&](handler &cgh) {
auto src_acc = src_buf.template get_access<access::mode::read>(cgh);
auto dst_acc = dst_buf.template get_access<access::mode::write>(cgh);
auto src_acc = src_buf.template get_access<access_mode::read>(cgh);
auto dst_acc = dst_buf.template get_access<access_mode::write>(cgh);

cgh.single_task([=]() SYCL_ESIMD_KERNEL {
simd<From, N> x(src_acc, 0);
Expand Down
6 changes: 3 additions & 3 deletions sycl/test-e2e/ESIMD/api/simd_binop_integer_promotion.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -52,9 +52,9 @@ template <typename T> bool test(queue q) {
range<1> glob_range{1};

auto e = q.submit([&](handler &cgh) {
auto PA = bufA.template get_access<access::mode::read>(cgh);
auto PB = bufB.template get_access<access::mode::read>(cgh);
auto PC = bufC.template get_access<access::mode::write>(cgh);
auto PA = bufA.template get_access<access_mode::read>(cgh);
auto PB = bufB.template get_access<access_mode::read>(cgh);
auto PC = bufC.template get_access<access_mode::write>(cgh);
cgh.parallel_for<KernelName<T>>(
glob_range, [=](id<1> i) SYCL_ESIMD_KERNEL {
using namespace sycl::ext::intel::experimental::esimd;
Expand Down
4 changes: 2 additions & 2 deletions sycl/test-e2e/ESIMD/api/simd_copy_to_from.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -84,8 +84,8 @@ bool testAcc(queue &Q, T *Src, T *Dst, unsigned Off, Flags) {
buffer<T, 1> DstB(Dst, range<1>(Off + N));

Q.submit([&](handler &CGH) {
auto SrcA = SrcB.template get_access<access::mode::read>(CGH);
auto DstA = DstB.template get_access<access::mode::write>(CGH);
auto SrcA = SrcB.template get_access<access_mode::read>(CGH);
auto DstA = DstB.template get_access<access_mode::write>(CGH);

CGH.parallel_for(sycl::range<1>{1}, [=](id<1>) SYCL_ESIMD_KERNEL {
simd<T, N> Vals;
Expand Down
4 changes: 2 additions & 2 deletions sycl/test-e2e/ESIMD/api/simd_copy_to_from_stateful.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -89,8 +89,8 @@ bool testAcc(queue &Q, T *Src, T *Dst, unsigned Off, Flags) {
buffer<T, 1> DstB(Dst, range<1>(Off + N));

Q.submit([&](handler &CGH) {
auto SrcA = SrcB.template get_access<access::mode::read>(CGH);
auto DstA = DstB.template get_access<access::mode::write>(CGH);
auto SrcA = SrcB.template get_access<access_mode::read>(CGH);
auto DstA = DstB.template get_access<access_mode::write>(CGH);

CGH.parallel_for(sycl::range<1>{1}, [=](id<1>) SYCL_ESIMD_KERNEL {
simd<T, N> Vals;
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/ESIMD/api/simd_memory_access.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -87,7 +87,7 @@ template <typename T, int N, bool IsAcc> bool test(queue q, size_t size) {
range<1> glob_range{size / N};

auto e = q.submit([&](handler &cgh) {
auto acc = buf.template get_access<access::mode::read_write>(cgh);
auto acc = buf.template get_access<access_mode::read_write>(cgh);
Kernel<T, N, true> kernel(acc);
cgh.parallel_for(glob_range, kernel);
});
Expand Down
4 changes: 2 additions & 2 deletions sycl/test-e2e/ESIMD/api/simd_negation_operator.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -37,8 +37,8 @@ bool test(queue q, const std::array<T, VL> &input,
range<1> glob_range{1};

auto e = q.submit([&](handler &cgh) {
auto PA = bufA.template get_access<access::mode::read>(cgh);
auto PB = bufB.template get_access<access::mode::write>(cgh);
auto PA = bufA.template get_access<access_mode::read>(cgh);
auto PB = bufB.template get_access<access_mode::write>(cgh);
cgh.parallel_for(glob_range, [=](id<1> i) SYCL_ESIMD_KERNEL {
using namespace sycl::ext::intel::esimd;
simd<T, VL> va;
Expand Down
4 changes: 2 additions & 2 deletions sycl/test-e2e/ESIMD/api/simd_subscript_operator.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -45,8 +45,8 @@ template <class T> bool test(queue &q) {
range<1> glob_range{1};

auto e = q.submit([&](handler &cgh) {
auto PA = bufA.template get_access<access::mode::read>(cgh);
auto PB = bufB.template get_access<access::mode::read_write>(cgh);
auto PA = bufA.template get_access<access_mode::read>(cgh);
auto PB = bufB.template get_access<access_mode::read_write>(cgh);
cgh.parallel_for<T>(glob_range, [=](id<1> i) SYCL_ESIMD_KERNEL {
using namespace sycl::ext::intel::esimd;
unsigned int offset = i * VL * sizeof(T);
Expand Down
4 changes: 2 additions & 2 deletions sycl/test-e2e/ESIMD/api/simd_view_copy_move_assign.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -38,8 +38,8 @@ bool test(queue q, std::string str, F funcUnderTest) {
buffer<T, 1> BufB(B, range<1>(Size));

q.submit([&](handler &cgh) {
auto PA = BufA.template get_access<access::mode::read_write>(cgh);
auto PB = BufB.template get_access<access::mode::read_write>(cgh);
auto PA = BufA.template get_access<access_mode::read_write>(cgh);
auto PB = BufB.template get_access<access_mode::read_write>(cgh);
cgh.parallel_for(range<1>{Size / VL}, [=](id<1> i) SYCL_ESIMD_KERNEL {
using namespace sycl::ext::intel::esimd;
unsigned int offset = i * VL * sizeof(T);
Expand Down
4 changes: 2 additions & 2 deletions sycl/test-e2e/ESIMD/api/simd_view_negation_operator.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -37,8 +37,8 @@ bool test(queue q, const std::array<T, VL> &input,
range<1> glob_range{1};

auto e = q.submit([&](handler &cgh) {
auto PA = bufA.template get_access<access::mode::read>(cgh);
auto PB = bufB.template get_access<access::mode::write>(cgh);
auto PA = bufA.template get_access<access_mode::read>(cgh);
auto PB = bufB.template get_access<access_mode::write>(cgh);
cgh.parallel_for(glob_range, [=](id<1> i) SYCL_ESIMD_KERNEL {
using namespace sycl::ext::intel::esimd;
simd<T, VL> va;
Expand Down
4 changes: 2 additions & 2 deletions sycl/test-e2e/ESIMD/api/simd_view_subscript_operator.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -48,8 +48,8 @@ template <class T> bool test(queue &q) {
range<1> glob_range{1};

auto e = q.submit([&](handler &cgh) {
auto PA = bufA.template get_access<access::mode::read>(cgh);
auto PB = bufB.template get_access<access::mode::read_write>(cgh);
auto PA = bufA.template get_access<access_mode::read>(cgh);
auto PB = bufB.template get_access<access_mode::read_write>(cgh);
cgh.parallel_for<TestID<T>>(glob_range, [=](id<1> i) SYCL_ESIMD_KERNEL {
using namespace sycl::ext::intel::esimd;
simd<T, VL> va;
Expand Down
4 changes: 2 additions & 2 deletions sycl/test-e2e/ESIMD/api/slm_gather_scatter_heavy.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -382,8 +382,8 @@ bool test_impl(queue q) {
range<1> glob_range{size / VL};

auto e = q.submit([&](handler &cgh) {
auto acc_in = buf_in.template get_access<access::mode::read_write>(cgh);
auto acc_out = buf_out.template get_access<access::mode::read_write>(cgh);
auto acc_in = buf_in.template get_access<access_mode::read_write>(cgh);
auto acc_out = buf_out.template get_access<access_mode::read_write>(cgh);
constexpr auto WG_SIZE = STRIDE; // for simplicity of the test
KernelType kernel(acc_in, acc_out, MaskedLane);
cgh.parallel_for(nd_range<1>{glob_range, range<1>(WG_SIZE)}, kernel);
Expand Down
4 changes: 2 additions & 2 deletions sycl/test-e2e/ESIMD/clz_ctz.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -43,8 +43,8 @@ template <typename T, bool CLZ> bool test(queue &q) {
queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());

auto e = q.submit([&](handler &cgh) {
auto PA = bufa.template get_access<access::mode::read>(cgh);
auto PB = bufb.template get_access<access::mode::write>(cgh);
auto PA = bufa.template get_access<access_mode::read>(cgh);
auto PB = bufb.template get_access<access_mode::write>(cgh);
cgh.parallel_for(GlobalRange * LocalRange,
[=](id<1> i) SYCL_ESIMD_KERNEL {
using namespace sycl::ext::intel::esimd;
Expand Down
6 changes: 3 additions & 3 deletions sycl/test-e2e/ESIMD/ext_math.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -346,10 +346,10 @@ bool test(queue &Q, const std::string &Name, InitF Init = InitNarrow<T>{},
sycl::range<1> LocalRange{1};

auto E = Q.submit([&](handler &CGH) {
auto PA = BufA.template get_access<access::mode::read>(CGH);
auto PC = BufC.template get_access<access::mode::write>(CGH);
auto PA = BufA.template get_access<access_mode::read>(CGH);
auto PC = BufC.template get_access<access_mode::write>(CGH);
if constexpr (IsBinOp) {
auto PB = BufB.template get_access<access::mode::read>(CGH);
auto PB = BufB.template get_access<access_mode::read>(CGH);
BinaryDeviceFunc<T, N, Op, Kernel, decltype(PA), decltype(PC)> F(PA, PB,
PC);
CGH.parallel_for(nd_range<1>{GlobalRange, LocalRange}, F);
Expand Down
6 changes: 3 additions & 3 deletions sycl/test-e2e/ESIMD/fma.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -36,9 +36,9 @@ template <typename T> bool test(queue q) {
buffer<T, 1> bufc(C, range<1>(Size));

auto e = q.submit([&](handler &cgh) {
auto PA = bufa.template get_access<access::mode::read>(cgh);
auto PB = bufb.template get_access<access::mode::read>(cgh);
auto PC = bufc.template get_access<access::mode::write>(cgh);
auto PA = bufa.template get_access<access_mode::read>(cgh);
auto PB = bufb.template get_access<access_mode::read>(cgh);
auto PC = bufc.template get_access<access_mode::write>(cgh);
cgh.single_task([=]() SYCL_ESIMD_KERNEL {
using namespace sycl::ext::intel::esimd;
unsigned int offset = 0;
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/ESIMD/fp_call_from_func.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -43,7 +43,7 @@ int main(int argc, char **argv) {
buffer<int, 1> buf(output, range<1>(1));

q.submit([&](handler &cgh) {
auto acc = buf.get_access<access::mode::write>(cgh);
auto acc = buf.get_access<access_mode::write>(cgh);

cgh.parallel_for<KernelID>(
sycl::range<1>{1},
Expand Down
Loading
Loading