Skip to content

[SYCL] Update E2E to use access_mode [1/3] #22842

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
againull merged 2 commits into from
Jul 31, 2026
Merged

[SYCL] Update E2E to use access_mode [1/3] #22842

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
9b69625
[SYCL] Update E2E to use access_mode
KornevNikita Jul 31, 2026
89dabdb
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/dynamic_address_cast.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -33,10 +33,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/AddressCast/static_address_cast.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -25,10 +25,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
2 changes: 1 addition & 1 deletion sycl/test-e2e/Assert/assert_in_multiple_tus.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -29,7 +29,7 @@ void enqueueKernel_1_fromFile1(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 Kernel_1>(
sycl::nd_range(Buf.get_range(), sycl::range<1>(4)),
Expand Down
20 changes: 10 additions & 10 deletions sycl/test-e2e/AtomicRef/add.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -32,9 +32,9 @@ void add_fetch_local_test(queue q, size_t N) {
buffer<T> sum_buf(&sum, 1);
buffer<T> output_buf(output.data(), output.size());
q.submit([&](handler &cgh) {
auto sum = sum_buf.template get_access<access::mode::read_write>(cgh);
auto sum = sum_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 @@ -81,9 +81,9 @@ void add_fetch_test(queue q, size_t N) {
buffer<T> sum_buf(&sum, 1);
buffer<T> output_buf(output.data(), output.size());
q.submit([&](handler &cgh) {
auto sum = sum_buf.template get_access<access::mode::read_write>(cgh);
auto sum = sum_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) {
int gid = it.get_id(0);
auto atm = AtomicRef < T,
Expand Down Expand Up @@ -164,9 +164,9 @@ void add_plus_equal_test(queue q, size_t N) {
buffer<T> output_buf(output.data(), output.size());

q.submit([&](handler &cgh) {
auto sum = sum_buf.template get_access<access::mode::read_write>(cgh);
auto sum = sum_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) {
int gid = it.get_id(0);
auto atm = AtomicRef < T,
Expand Down Expand Up @@ -247,9 +247,9 @@ void add_pre_inc_test(queue q, size_t N) {
buffer<T> output_buf(output.data(), output.size());

q.submit([&](handler &cgh) {
auto sum = sum_buf.template get_access<access::mode::read_write>(cgh);
auto sum = sum_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) {
int gid = it.get_id(0);
auto atm = AtomicRef < T,
Expand Down Expand Up @@ -329,9 +329,9 @@ void add_post_inc_test(queue q, size_t N) {
buffer<T> output_buf(output.data(), output.size());

q.submit([&](handler &cgh) {
auto sum = sum_buf.template get_access<access::mode::read_write>(cgh);
auto sum = sum_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) {
int gid = it.get_id(0);
auto atm = AtomicRef < T,
Expand Down
16 changes: 6 additions & 10 deletions sycl/test-e2e/AtomicRef/atomic_memory_order_acq_rel.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -22,9 +22,8 @@ template <memory_order order> void test_acquire_global() {
buffer<int> val_buf(val, 2);

q.submit([&](handler &cgh) {
auto error =
error_buf.template get_access<access::mode::read_write>(cgh);
auto val = val_buf.template get_access<access::mode::read_write>(cgh);
auto error = error_buf.template get_access<access_mode::read_write>(cgh);
auto val = val_buf.template get_access<access_mode::read_write>(cgh);
cgh.parallel_for(range<1>(N_items), [=](item<1> it) {
volatile int *val_p = val.get_multi_ptr<access::decorated::no>().get();
auto atm0 =
Expand Down Expand Up @@ -66,8 +65,7 @@ template <memory_order order> void test_acquire_local() {
buffer<int> val_buf(val, 2);

q.submit([&](handler &cgh) {
auto error =
error_buf.template get_access<access::mode::read_write>(cgh);
auto error = error_buf.template get_access<access_mode::read_write>(cgh);
local_accessor<int, 1> val(2, cgh);
cgh.parallel_for(
nd_range<1>(global_size, local_size), [=](nd_item<1> it) {
Expand Down Expand Up @@ -114,9 +112,8 @@ template <memory_order order> void test_release_global() {
buffer<int> val_buf(val, 2);

q.submit([&](handler &cgh) {
auto error =
error_buf.template get_access<access::mode::read_write>(cgh);
auto val = val_buf.template get_access<access::mode::read_write>(cgh);
auto error = error_buf.template get_access<access_mode::read_write>(cgh);
auto val = val_buf.template get_access<access_mode::read_write>(cgh);
cgh.parallel_for(range<1>(N_items), [=](item<1> it) {
volatile int *val_p = val.get_multi_ptr<access::decorated::no>().get();
auto atm0 =
Expand Down Expand Up @@ -158,8 +155,7 @@ template <memory_order order> void test_release_local() {
buffer<int> val_buf(val, 2);

q.submit([&](handler &cgh) {
auto error =
error_buf.template get_access<access::mode::read_write>(cgh);
auto error = error_buf.template get_access<access_mode::read_write>(cgh);
local_accessor<int, 1> val(2, cgh);
cgh.parallel_for(
nd_range<1>(global_size, local_size), [=](nd_item<1> it) {
Expand Down
9 changes: 4 additions & 5 deletions sycl/test-e2e/AtomicRef/exchange.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -32,9 +32,9 @@ void exchange_local_test(queue q, size_t N) {
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,10 +79,9 @@ void exchange_global_test(queue q, size_t N) {
buffer<T> output_buf(output.data(), output.size());

q.submit([&](handler &cgh) {
auto exc =
exchange_buf.template get_access<access::mode::read_write>(cgh);
auto exc = exchange_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/min.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -32,9 +32,9 @@ void min_local_test(queue q, size_t N) {
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 @@ -83,9 +83,9 @@ void min_global_test(queue q, size_t N) {
buffer<T> output_buf(output.data(), output.size());

q.submit([&](handler &cgh) {
auto val = val_buf.template get_access<access::mode::read_write>(cgh);
auto val = val_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) {
int gid = it.get_id(0);
auto atm = AtomicRef < T,
Expand Down
20 changes: 10 additions & 10 deletions sycl/test-e2e/AtomicRef/sub.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -32,9 +32,9 @@ void sub_fetch_local_test(queue q, size_t N) {
buffer<T> sum_buf(&sum, 1);
buffer<T> output_buf(output.data(), output.size());
q.submit([&](handler &cgh) {
auto sum = sum_buf.template get_access<access::mode::read_write>(cgh);
auto sum = sum_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 @@ -82,9 +82,9 @@ void sub_fetch_test(queue q, size_t N) {
buffer<T> output_buf(output.data(), output.size());

q.submit([&](handler &cgh) {
auto val = val_buf.template get_access<access::mode::read_write>(cgh);
auto val = val_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) {
int gid = it.get_id(0);
auto atm = AtomicRef < T,
Expand Down Expand Up @@ -165,9 +165,9 @@ void sub_minus_equal_test(queue q, size_t N) {
buffer<T> output_buf(output.data(), output.size());

q.submit([&](handler &cgh) {
auto val = val_buf.template get_access<access::mode::read_write>(cgh);
auto val = val_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) {
int gid = it.get_id(0);
auto atm = AtomicRef < T,
Expand Down Expand Up @@ -248,9 +248,9 @@ void sub_pre_dec_test(queue q, size_t N) {
buffer<T> output_buf(output.data(), output.size());

q.submit([&](handler &cgh) {
auto val = val_buf.template get_access<access::mode::read_write>(cgh);
auto val = val_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) {
int gid = it.get_id(0);
auto atm = AtomicRef < T,
Expand Down Expand Up @@ -331,9 +331,9 @@ void sub_post_dec_test(queue q, size_t N) {
buffer<T> output_buf(output.data(), output.size());

q.submit([&](handler &cgh) {
auto val = val_buf.template get_access<access::mode::read_write>(cgh);
auto val = val_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) {
int gid = it.get_id(0);
auto atm = AtomicRef < T,
Expand Down
Loading
Loading