Skip to content

[SYCL] Update SYCL tests to use access_mode #22825

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

Closed
KornevNikita wants to merge 6 commits into from
Closed

[SYCL] Update SYCL tests to use access_mode #22825

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
c555be9
[SYCL] Update SYCL tests to use access_mode
KornevNikita Jul 30, 2026
a54e525
restore SpecConstants tests
KornevNikita Jul 30, 2026
6db9c90
restore tools tests
KornevNikita Jul 30, 2026
fd1c4b0
format
KornevNikita Jul 30, 2026
0f86892
restore-discard-write
KornevNikita Jul 31, 2026
1d379c8
restore
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/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/AddressCast/khr_static_addrspace_cast.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -24,10 +24,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
10 changes: 5 additions & 5 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,7 +32,7 @@ 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);
local_accessor<T, 1> loc(1, cgh);
Expand Down Expand Up @@ -81,7 +81,7 @@ 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);
cgh.parallel_for(range<1>(N), [=](item<1> it) {
Expand Down Expand Up @@ -164,7 +164,7 @@ 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);
cgh.parallel_for(range<1>(N), [=](item<1> it) {
Expand Down Expand Up @@ -247,7 +247,7 @@ 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);
cgh.parallel_for(range<1>(N), [=](item<1> it) {
Expand Down Expand Up @@ -329,7 +329,7 @@ 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);
cgh.parallel_for(range<1>(N), [=](item<1> it) {
Expand Down
4 changes: 2 additions & 2 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,7 +32,7 @@ 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);
local_accessor<T, 1> loc(1, cgh);
Expand Down Expand Up @@ -79,7 +79,7 @@ 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);
cgh.parallel_for(range<1>(N), [=](item<1> it) {
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/AtomicRef/assignment.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -27,7 +27,7 @@ void assignment_test(queue q, size_t N) {
buffer<T> assignment_buf(&assignment, 1);
q.submit([&](handler &cgh) {
auto st =
assignment_buf.template get_access<access::mode::read_write>(cgh);
assignment_buf.template get_access<access_mode::read_write>(cgh);
cgh.parallel_for<assignment_kernel<AtomicRef, address_space, T>>(
range<1>(N), [=](item<1> it) {
size_t gid = it.get_id(0);
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
14 changes: 7 additions & 7 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,7 +35,7 @@ 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);
cgh.parallel_for(nd_range<1>(N_items, 32), [=](nd_item<1> it) {
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 @@ -93,7 +93,7 @@ template <memory_order order> void test_global(size_t N_iters) {

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 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
4 changes: 2 additions & 2 deletions sycl/test-e2e/AtomicRef/compare_exchange.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -33,7 +33,7 @@ void compare_exchange_local_test(queue q, size_t N) {
buffer<T> output_buf(output.data(), output.size());
q.submit([&](handler &cgh) {
auto compare_exchange =
compare_exchange_buf.template get_access<access::mode::read_write>(
compare_exchange_buf.template get_access<access_mode::read_write>(
cgh);
auto out =
output_buf.template get_access<access::mode::discard_write>(cgh);
Expand Down Expand Up @@ -88,7 +88,7 @@ void compare_exchange_global_test(queue q, size_t N) {

q.submit([&](handler &cgh) {
auto exc =
compare_exchange_buf.template get_access<access::mode::read_write>(
compare_exchange_buf.template get_access<access_mode::read_write>(
cgh);
auto out =
output_buf.template get_access<access::mode::discard_write>(cgh);
Expand Down
5 changes: 2 additions & 3 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,7 +32,7 @@ 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);
local_accessor<T, 1> loc(1, cgh);
Expand Down Expand Up @@ -79,8 +79,7 @@ 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);
cgh.parallel_for(range<1>(N), [=](item<1> it) {
Expand Down
4 changes: 2 additions & 2 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,7 +33,7 @@ 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);
local_accessor<T, 1> loc(1, cgh);
Expand Down Expand Up @@ -69,7 +69,7 @@ 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);
cgh.parallel_for(range<1>(N), [=](item<1> it) {
Expand Down
4 changes: 2 additions & 2 deletions sycl/test-e2e/AtomicRef/max.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -32,7 +32,7 @@ void max_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);
local_accessor<T, 1> loc(1, cgh);
Expand Down Expand Up @@ -83,7 +83,7 @@ void max_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);
cgh.parallel_for(range<1>(N), [=](item<1> it) {
Expand Down
4 changes: 2 additions & 2 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,7 +32,7 @@ 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);
local_accessor<T, 1> loc(1, cgh);
Expand Down Expand Up @@ -83,7 +83,7 @@ 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);
cgh.parallel_for(range<1>(N), [=](item<1> it) {
Expand Down
4 changes: 2 additions & 2 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,7 +32,7 @@ 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);
local_accessor<T, 1> loc(1, cgh);
Expand Down Expand Up @@ -79,7 +79,7 @@ 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);
cgh.parallel_for(range<1>(N), [=](item<1> it) {
Expand Down
4 changes: 2 additions & 2 deletions sycl/test-e2e/AtomicRef/store.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -29,7 +29,7 @@ void store_global_test(queue q, size_t N) {
{
buffer<T> store_buf(&store, 1);
q.submit([&](handler &cgh) {
auto st = store_buf.template get_access<access::mode::read_write>(cgh);
auto st = store_buf.template get_access<access_mode::read_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>(st[0]);
Expand Down Expand Up @@ -82,7 +82,7 @@ void store_local_test(queue q, size_t N) {
{
buffer<T> store_buf(&store, 1);
q.submit([&](handler &cgh) {
auto st = store_buf.template get_access<access::mode::read_write>(cgh);
auto st = store_buf.template get_access<access_mode::read_write>(cgh);
local_accessor<T, 1> loc(1, cgh);
cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
size_t gid = it.get_global_id(0);
Expand Down
10 changes: 5 additions & 5 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,7 +32,7 @@ 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);
local_accessor<T, 1> loc(1, cgh);
Expand Down Expand Up @@ -82,7 +82,7 @@ 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);
cgh.parallel_for(range<1>(N), [=](item<1> it) {
Expand Down Expand Up @@ -165,7 +165,7 @@ 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);
cgh.parallel_for(range<1>(N), [=](item<1> it) {
Expand Down Expand Up @@ -248,7 +248,7 @@ 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);
cgh.parallel_for(range<1>(N), [=](item<1> it) {
Expand Down Expand Up @@ -331,7 +331,7 @@ 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);
cgh.parallel_for(range<1>(N), [=](item<1> it) {
Expand Down
Loading
Loading