Skip to content

[SYCL] Update E2E to use access_mode [2/3] #22843

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
Aug 3, 2026
Merged

[SYCL] Update E2E to use access_mode [2/3] #22843

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
6961a1e
[SYCL] Update E2E to use access_mode [1/3]
KornevNikita Jul 31, 2026
4523f69
drop sycl:: where possible
KornevNikita Jul 31, 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
4 changes: 2 additions & 2 deletions sycl/test-e2e/AddressCast/khr_dynamic_addrspace_cast.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -31,10 +31,10 @@ int main() {
Queue
.submit([&](sycl::handler &cgh) {
auto GlobalAccessor =
GlobalBuffer.get_access<sycl::access::mode::read_write>(cgh);
GlobalBuffer.get_access<sycl::access_mode::read_write>(cgh);
auto LocalAccessor = sycl::local_accessor<int>(1, cgh);
auto ResultAccessor =
ResultBuffer.get_access<sycl::access::mode::write>(cgh);
ResultBuffer.get_access<sycl::access_mode::write>(cgh);
cgh.parallel_for<class Kernel>(
sycl::nd_range<1>(NItems, 1), [=](sycl::nd_item<1> Item) {
bool Success = true;
Expand Down
4 changes: 2 additions & 2 deletions sycl/test-e2e/Assert/Inputs/kernels_in_file2.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -25,7 +25,7 @@ void enqueueKernel_1_fromFile2(queue *Q) {
sycl::buffer<int, 1> Buf(numOfItems);

Q->submit([&](handler &CGH) {
auto Acc = Buf.template get_access<mode::read_write>(CGH);
auto Acc = Buf.template get_access<access_mode::read_write>(CGH);

CGH.parallel_for<class kernel1_from_separate_file>(numOfItems,
[=](sycl::id<1> wiID) {
Expand All @@ -40,7 +40,7 @@ void enqueueKernel_2_fromFile2(queue *Q) {
sycl::buffer<int, 1> Buf(numOfItems);

Q->submit([&](handler &CGH) {
auto Acc = Buf.template get_access<mode::read_write>(CGH);
auto Acc = Buf.template get_access<access_mode::read_write>(CGH);

CGH.parallel_for<class kernel2_from_separate_file>(numOfItems,
[=](sycl::id<1> wiID) {
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/Assert/assert_in_one_kernel.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -19,7 +19,7 @@ void assertTest() {

queue Q;
Q.submit([&](handler &CGH) {
auto Acc = Buf.template get_access<mode::read_write>(CGH);
auto Acc = Buf.template get_access<access_mode::read_write>(CGH);

CGH.parallel_for<class TheKernel>(
numOfItems, [=](item<1> Item) { kernelFunc(&Acc[0], Item[0]); });
Expand Down
8 changes: 4 additions & 4 deletions sycl/test-e2e/AtomicRef/and.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -32,9 +32,9 @@ void and_local_test(queue q) {
buffer<T> cum_buf(&cum, 1);
buffer<T> output_buf(output.data(), output.size());
q.submit([&](handler &cgh) {
auto cum = cum_buf.template get_access<access::mode::read_write>(cgh);
auto cum = cum_buf.template get_access<access_mode::read_write>(cgh);
auto out =
output_buf.template get_access<access::mode::discard_write>(cgh);
output_buf.template get_access<access_mode::discard_write>(cgh);
Comment thread
KornevNikita marked this conversation as resolved.
local_accessor<T, 1> loc(1, cgh);

cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
Expand Down Expand Up @@ -79,9 +79,9 @@ void and_global_test(queue q) {
buffer<T> output_buf(output.data(), output.size());

q.submit([&](handler &cgh) {
auto cum = cum_buf.template get_access<access::mode::read_write>(cgh);
auto cum = cum_buf.template get_access<access_mode::read_write>(cgh);
auto out =
output_buf.template get_access<access::mode::discard_write>(cgh);
output_buf.template get_access<access_mode::discard_write>(cgh);
cgh.parallel_for(range<1>(N), [=](item<1> it) {
size_t gid = it.get_id(0);
auto atm = AtomicRef < T,
Expand Down
20 changes: 10 additions & 10 deletions sycl/test-e2e/AtomicRef/atomic_memory_order_seq_cst.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -35,9 +35,9 @@ void check(queue &q, buffer<int, 2> &res_buf, size_t N_iters) {
{N_items / 2 * N_iters + 1, N_items / 2 * N_iters + 1});

q.submit([&](handler &cgh) {
auto res = res_buf.template get_access<access::mode::read>(cgh);
auto res = res_buf.template get_access<access_mode::read>(cgh);
auto checked =
checked_buf.template get_access<access::mode::discard_write>(cgh);
checked_buf.template get_access<access_mode::discard_write>(cgh);
cgh.parallel_for(nd_range<1>(N_items, 32), [=](nd_item<1> it) {
for (int i = it.get_global_id(0); i < N_items / 2 * N_iters + 1;
i += N_items) {
Expand All @@ -48,8 +48,8 @@ void check(queue &q, buffer<int, 2> &res_buf, size_t N_iters) {
});
});
q.submit([&](handler &cgh) {
auto res = res_buf.template get_access<access::mode::read>(cgh);
auto checked = checked_buf.template get_access<access::mode::write>(cgh);
auto res = res_buf.template get_access<access_mode::read>(cgh);
auto checked = checked_buf.template get_access<access_mode::write>(cgh);
cgh.parallel_for(nd_range<1>(N_items / 2, 32), [=](nd_item<1> it) {
size_t id = it.get_global_id(0);
for (int i = 1; i < N_iters; i++) {
Expand All @@ -65,9 +65,9 @@ void check(queue &q, buffer<int, 2> &res_buf, size_t N_iters) {
int err = 0;
buffer<int> err_buf(&err, 1);
q.submit([&](handler &cgh) {
auto res = res_buf.template get_access<access::mode::read>(cgh);
auto checked = checked_buf.template get_access<access::mode::read>(cgh);
auto err = err_buf.template get_access<access::mode::write>(cgh);
auto res = res_buf.template get_access<access_mode::read>(cgh);
auto checked = checked_buf.template get_access<access_mode::read>(cgh);
auto err = err_buf.template get_access<access_mode::write>(cgh);
cgh.parallel_for(nd_range<1>(N_items / 2, 32), [=](nd_item<1> it) {
size_t id = it.get_global_id(0);
for (int i = 1; i < N_iters; i++) {
Expand All @@ -92,8 +92,8 @@ template <memory_order order> void test_global(size_t N_iters) {
buffer<int> val_buf(&val, 1);

q.submit([&](handler &cgh) {
auto res = res_buf.template get_access<access::mode::discard_write>(cgh);
auto val = val_buf.template get_access<access::mode::read_write>(cgh);
auto res = res_buf.template get_access<access_mode::discard_write>(cgh);
auto val = val_buf.template get_access<access_mode::read_write>(cgh);
// Intentionally using a small work group size. The assumption being that
// more sub groups mean more likely failure for the same number of
// work-items if sequential consistency does not work
Expand Down Expand Up @@ -123,7 +123,7 @@ template <memory_order order> void test_local(size_t N_iters) {
buffer<int, 2> res_buf({N_items / 2, N_iters});

q.submit([&](handler &cgh) {
auto res = res_buf.template get_access<access::mode::discard_write>(cgh);
auto res = res_buf.template get_access<access_mode::discard_write>(cgh);
local_accessor<int, 1> val(2, cgh);
cgh.parallel_for(nd_range<1>(N_items, N_items), [=](nd_item<1> it) {
val[0] = 0;
Expand Down
8 changes: 4 additions & 4 deletions sycl/test-e2e/AtomicRef/load.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -33,9 +33,9 @@ void load_local_test(queue q, size_t N) {
buffer<T> output_buf(output.data(), output.size());

q.submit([&](handler &cgh) {
auto ld = load_buf.template get_access<access::mode::read_write>(cgh);
auto ld = load_buf.template get_access<access_mode::read_write>(cgh);
auto out =
output_buf.template get_access<access::mode::discard_write>(cgh);
output_buf.template get_access<access_mode::discard_write>(cgh);
local_accessor<T, 1> loc(1, cgh);
cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
int gid = it.get_global_id(0);
Expand Down Expand Up @@ -69,9 +69,9 @@ void load_global_test(queue q, size_t N) {
buffer<T> output_buf(output.data(), output.size());

q.submit([&](handler &cgh) {
auto ld = load_buf.template get_access<access::mode::read_write>(cgh);
auto ld = load_buf.template get_access<access_mode::read_write>(cgh);
auto out =
output_buf.template get_access<access::mode::discard_write>(cgh);
output_buf.template get_access<access_mode::discard_write>(cgh);
cgh.parallel_for(range<1>(N), [=](item<1> it) {
size_t gid = it.get_id(0);
auto atm = AtomicRef<T, memory_order::relaxed, scope, space>(ld[0]);
Expand Down
8 changes: 4 additions & 4 deletions sycl/test-e2e/AtomicRef/or.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -32,9 +32,9 @@ void or_local_test(queue q) {
buffer<T> cum_buf(&cum, 1);
buffer<T> output_buf(output.data(), output.size());
q.submit([&](handler &cgh) {
auto cum = cum_buf.template get_access<access::mode::read_write>(cgh);
auto cum = cum_buf.template get_access<access_mode::read_write>(cgh);
auto out =
output_buf.template get_access<access::mode::discard_write>(cgh);
output_buf.template get_access<access_mode::discard_write>(cgh);
local_accessor<T, 1> loc(1, cgh);

cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
Expand Down Expand Up @@ -79,9 +79,9 @@ void or_global_test(queue q) {
buffer<T> output_buf(output.data(), output.size());

q.submit([&](handler &cgh) {
auto cum = cum_buf.template get_access<access::mode::read_write>(cgh);
auto cum = cum_buf.template get_access<access_mode::read_write>(cgh);
auto out =
output_buf.template get_access<access::mode::discard_write>(cgh);
output_buf.template get_access<access_mode::discard_write>(cgh);
cgh.parallel_for(range<1>(N), [=](item<1> it) {
size_t gid = it.get_id(0);
auto atm = AtomicRef < T,
Expand Down
8 changes: 4 additions & 4 deletions sycl/test-e2e/AtomicRef/xor.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -32,9 +32,9 @@ void xor_local_test(queue q) {
buffer<T> cum_buf(&cum, 1);
buffer<T> output_buf(output.data(), output.size());
q.submit([&](handler &cgh) {
auto cum = cum_buf.template get_access<access::mode::read_write>(cgh);
auto cum = cum_buf.template get_access<access_mode::read_write>(cgh);
auto out =
output_buf.template get_access<access::mode::discard_write>(cgh);
output_buf.template get_access<access_mode::discard_write>(cgh);
local_accessor<T, 1> loc(1, cgh);

cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
Expand Down Expand Up @@ -79,9 +79,9 @@ void xor_global_test(queue q) {
buffer<T> output_buf(output.data(), output.size());

q.submit([&](handler &cgh) {
auto cum = cum_buf.template get_access<access::mode::read_write>(cgh);
auto cum = cum_buf.template get_access<access_mode::read_write>(cgh);
auto out =
output_buf.template get_access<access::mode::discard_write>(cgh);
output_buf.template get_access<access_mode::discard_write>(cgh);
cgh.parallel_for(range<1>(N), [=](item<1> it) {
size_t gid = it.get_id(0);
auto atm = AtomicRef < T,
Expand Down
4 changes: 2 additions & 2 deletions sycl/test-e2e/Basic/alloc_pinned_host_memory.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -30,8 +30,8 @@ int main() {
{ext::oneapi::property::buffer::use_pinned_host_memory()});
queue myQueue;
myQueue.submit([&](handler &cgh) {
auto A = a.get_access<access::mode::read_write>(cgh);
auto B = b.get_access<access::mode::read_write>(cgh);
auto A = a.get_access<access_mode::read_write>(cgh);
auto B = b.get_access<access_mode::read_write>(cgh);
cgh.parallel_for<class init_b>(range<1>{10}, [=](id<1> index) {
B[index] = 0;
A[index] = B[index] + 1;
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/Basic/buffer/buffer_container.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -14,7 +14,7 @@ int main() {
buffer A(out_data);
queue Q;
Q.submit([&](handler &h) {
auto out = A.get_access<access::mode::write>(h);
auto out = A.get_access<access_mode::write>(h);
h.parallel_for<class containerBuffer>(A.get_range(),
[out](id<1> i) { out[i] = 1; });
});
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/Basic/buffer/native_buffer_creation_flags.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -22,7 +22,7 @@ int main() {
// buffer is created with the UR_MEM_FLAG_USE_HOST_POINTER flag.
// CHECK: <--- urMemBufferCreate
// CHECK-SAME: UR_MEM_FLAG_USE_HOST_POINTER
auto BufAcc = Buf.get_access<access::mode::read>(Cgh);
auto BufAcc = Buf.get_access<access_mode::read>(Cgh);
Cgh.single_task<Foo>([=]() { int A = BufAcc[0]; });
});
}
4 changes: 2 additions & 2 deletions sycl/test-e2e/Basic/built-ins.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -41,8 +41,8 @@ int main() {
s::buffer<float, 1> BufMin(s::range<1>(1));
s::buffer<s::float2, 1> BufMax(s::range<1>(1));
q.submit([&](s::handler &cgh) {
auto AccMin = BufMin.get_access<s::access::mode::write>(cgh);
auto AccMax = BufMax.get_access<s::access::mode::write>(cgh);
auto AccMin = BufMin.get_access<s::access_mode::write>(cgh);
auto AccMax = BufMax.get_access<s::access_mode::write>(cgh);
cgh.single_task<class common>([=]() {
AccMax[0] = s::max(s::float2{0.5f, 2.5f}, s::float2{2.3f, 2.3f});
AccMin[0] = s::min(float{0.5f}, float{2.3f});
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/Basic/device_event.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -67,7 +67,7 @@ int test_strideN(size_t stride) {
buffer<int, 1> out_buf(out_data, range<1>(nElems));

myQueue.submit([&](handler &cgh) {
auto out_ptr = out_buf.get_access<access::mode::write>(cgh);
auto out_ptr = out_buf.get_access<access_mode::write>(cgh);
local_accessor<int, 1> local_acc(range<1>(16), cgh);

// Create work-groups with 16 work items in each group.
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/Basic/event_release.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,7 +21,7 @@ int main() {
{
buffer<int, 1> Buf{&Val, range<1>(1)};
Q.submit([&](handler &Cgh) {
auto Acc = Buf.get_access<access::mode::discard_write>(Cgh);
auto Acc = Buf.get_access<access_mode::discard_write>(Cgh);
Cgh.single_task<Foo>([=]() { Acc[0] = Gold; });
});
}
Expand Down
24 changes: 12 additions & 12 deletions sycl/test-e2e/Basic/free_function_queries/free_function_queries.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -32,12 +32,12 @@ int main() {
sycl::buffer<int> results_buf(results, sycl::range<1>(checks_number));
sycl::queue q;
q.submit([&](sycl::handler &cgh) {
sycl::accessor<int, 1, sycl::access::mode::write,
sycl::accessor<int, 1, sycl::access_mode::write,
sycl::access::target::device>
acc(buf.get_access<sycl::access::mode::write>(cgh));
sycl::accessor<int, 1, sycl::access::mode::write,
acc(buf.get_access<sycl::access_mode::write>(cgh));
sycl::accessor<int, 1, sycl::access_mode::write,
sycl::access::target::device>
results_acc(results_buf.get_access<sycl::access::mode::write>(cgh));
results_acc(results_buf.get_access<sycl::access_mode::write>(cgh));
cgh.parallel_for<class IdTest>(n, [=](sycl::id<1> i) {
auto that_id = sycl::ext::oneapi::this_work_item::get_nd_item<1>();
results_acc[0] = that_id.get_global_id() == i;
Expand Down Expand Up @@ -66,12 +66,12 @@ int main() {
sycl::buffer<int> results_buf(results, sycl::range<1>(checks_number));
sycl::queue q;
q.submit([&](sycl::handler &cgh) {
sycl::accessor<int, 1, sycl::access::mode::write,
sycl::accessor<int, 1, sycl::access_mode::write,
sycl::access::target::device>
acc(buf.get_access<sycl::access::mode::write>(cgh));
sycl::accessor<int, 1, sycl::access::mode::write,
acc(buf.get_access<sycl::access_mode::write>(cgh));
sycl::accessor<int, 1, sycl::access_mode::write,
sycl::access::target::device>
results_acc(results_buf.get_access<sycl::access::mode::write>(cgh));
results_acc(results_buf.get_access<sycl::access_mode::write>(cgh));
cgh.parallel_for<class ItemTest>(n, [=](auto i) {
static_assert(std::is_same<decltype(i), sycl::item<1>>::value,
"lambda arg type is unexpected");
Expand Down Expand Up @@ -104,12 +104,12 @@ int main() {
sycl::queue q;
sycl::id<1> offset(1);
q.submit([&](sycl::handler &cgh) {
sycl::accessor<int, 1, sycl::access::mode::write,
sycl::accessor<int, 1, sycl::access_mode::write,
sycl::access::target::device>
acc(buf.get_access<sycl::access::mode::write>(cgh));
sycl::accessor<int, 1, sycl::access::mode::write,
acc(buf.get_access<sycl::access_mode::write>(cgh));
sycl::accessor<int, 1, sycl::access_mode::write,
sycl::access::target::device>
results_acc(results_buf.get_access<sycl::access::mode::write>(cgh));
results_acc(results_buf.get_access<sycl::access_mode::write>(cgh));
cgh.parallel_for<class ItemOffsetTest>(
sycl::range<1>{n}, offset, [=](sycl::item<1, true> i) {
auto that_id =
Expand Down
4 changes: 2 additions & 2 deletions sycl/test-e2e/Basic/group_async_copy.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -108,8 +108,8 @@ template <typename T> int test(size_t Stride) {
initOutputBuffer(OutBuf);

Q.submit([&](handler &CGH) {
auto In = InBuf.template get_access<access::mode::read>(CGH);
auto Out = OutBuf.template get_access<access::mode::write>(CGH);
auto In = InBuf.template get_access<access_mode::read>(CGH);
auto Out = OutBuf.template get_access<access_mode::write>(CGH);
local_accessor<T, 1> Local(range<1>{WorkGroupSize}, CGH);

nd_range<1> NDR{range<1>(NElems), range<1>(WorkGroupSize)};
Expand Down
26 changes: 13 additions & 13 deletions sycl/test-e2e/Basic/half_type.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -34,9 +34,9 @@ void verify_add(queue &q, buffer<half, 1> &a, buffer<half, 1> &b, range<1> &r,
buffer<half, 1> c{r};

q.submit([&](handler &cgh) {
auto A = a.get_access<access::mode::read>(cgh);
auto B = b.get_access<access::mode::read>(cgh);
auto C = c.get_access<access::mode::write>(cgh);
auto A = a.get_access<access_mode::read>(cgh);
auto B = b.get_access<access_mode::read>(cgh);
auto C = c.get_access<access_mode::write>(cgh);
cgh.parallel_for<class calc_add>(
r, [=](id<1> index) { C[index] = A[index] + B[index]; });
});
Expand All @@ -49,9 +49,9 @@ void verify_min(queue &q, buffer<half, 1> &a, buffer<half, 1> &b, range<1> &r,
buffer<half, 1> c{r};

q.submit([&](handler &cgh) {
auto A = a.get_access<access::mode::read>(cgh);
auto B = b.get_access<access::mode::read>(cgh);
auto C = c.get_access<access::mode::write>(cgh);
auto A = a.get_access<access_mode::read>(cgh);
auto B = b.get_access<access_mode::read>(cgh);
auto C = c.get_access<access_mode::write>(cgh);
cgh.parallel_for<class calc_min>(
r, [=](id<1> index) { C[index] = A[index] - B[index]; });
});
Expand All @@ -64,9 +64,9 @@ void verify_mul(queue &q, buffer<half, 1> &a, buffer<half, 1> &b, range<1> &r,
buffer<half, 1> c{r};

q.submit([&](handler &cgh) {
auto A = a.get_access<access::mode::read>(cgh);
auto B = b.get_access<access::mode::read>(cgh);
auto C = c.get_access<access::mode::write>(cgh);
auto A = a.get_access<access_mode::read>(cgh);
auto B = b.get_access<access_mode::read>(cgh);
auto C = c.get_access<access_mode::write>(cgh);
cgh.parallel_for<class calc_mul>(
r, [=](id<1> index) { C[index] = A[index] * B[index]; });
});
Expand All @@ -79,9 +79,9 @@ void verify_div(queue &q, buffer<half, 1> &a, buffer<half, 1> &b, range<1> &r,
buffer<half, 1> c{r};

q.submit([&](handler &cgh) {
auto A = a.get_access<access::mode::read>(cgh);
auto B = b.get_access<access::mode::read>(cgh);
auto C = c.get_access<access::mode::write>(cgh);
auto A = a.get_access<access_mode::read>(cgh);
auto B = b.get_access<access_mode::read>(cgh);
auto C = c.get_access<access_mode::write>(cgh);
cgh.parallel_for<class calc_div>(
r, [=](id<1> index) { C[index] = A[index] / B[index]; });
});
Expand All @@ -94,7 +94,7 @@ void verify_vec(queue &q) {
int a = 0;
buffer<int, 1> e(&a, range<1>(1));
q.submit([&](sycl::handler &cgh) {
auto E = e.get_access<access::mode::write>(cgh);
auto E = e.get_access<access_mode::write>(cgh);
cgh.single_task<class vec_of_half>([=]() {
if (int(hvec.s0()) != 999 || int(hvec.s1()) != 999)
E[0] = 1;
Expand Down
Loading
Loading