Skip to content
This repository was archived by the owner on Mar 28, 2023. It is now read-only.

Commit 476f8ce

Browse files
author
Pavel Samolysov
committed
[SYCL][NFC] Fix formatting issues
Signed-off-by: Pavel Samolysov <[email protected]>
1 parent a6bb62d commit 476f8ce

File tree

12 files changed

+63
-70
lines changed

12 files changed

+63
-70
lines changed

‎SYCL/AtomicRef/accessor.cpp

Lines changed: 8 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -38,11 +38,10 @@ template <typename T> void accessor_test(queue q, size_t N) {
3838
cgh.parallel_for(range<1>(N), [=](item<1> it) {
3939
int gid = it.get_id(0);
4040
static_assert(
41-
std::is_same<
42-
decltype(sum[0]),
43-
::sycl::ext::oneapi::atomic_ref<T, memory_order::relaxed,
44-
memory_scope::device,
45-
access::address_space::global_space>>::value,
41+
std::is_same<decltype(sum[0]),
42+
::sycl::ext::oneapi::atomic_ref<
43+
T, memory_order::relaxed, memory_scope::device,
44+
access::address_space::global_space>>::value,
4645
"atomic_accessor returns incorrect atomic_ref");
4746
out[gid] = sum[0].fetch_add(T(1));
4847
});
@@ -79,11 +78,10 @@ void local_accessor_test(queue q, size_t N, size_t L = 8) {
7978
sum[0].store(0);
8079
it.barrier();
8180
static_assert(
82-
std::is_same<
83-
decltype(sum[0]),
84-
::sycl::ext::oneapi::atomic_ref<T, memory_order::relaxed,
85-
memory_scope::device,
86-
access::address_space::local_space>>::value,
81+
std::is_same<decltype(sum[0]),
82+
::sycl::ext::oneapi::atomic_ref<
83+
T, memory_order::relaxed, memory_scope::device,
84+
access::address_space::local_space>>::value,
8785
"local atomic_accessor returns incorrect atomic_ref");
8886
T result = sum[0].fetch_add(T(1));
8987
if (result == it.get_local_range(0) - 1) {

‎SYCL/AtomicRef/add.h

Lines changed: 12 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -24,9 +24,9 @@ void add_fetch_test(queue q, size_t N) {
2424
output_buf.template get_access<access::mode::discard_write>(cgh);
2525
cgh.parallel_for(range<1>(N), [=](item<1> it) {
2626
int gid = it.get_id(0);
27-
auto atm = ::sycl::ext::oneapi::atomic_ref<T, memory_order::relaxed,
28-
memory_scope::device,
29-
access::address_space::global_space>(sum[0]);
27+
auto atm = ::sycl::ext::oneapi::atomic_ref<
28+
T, memory_order::relaxed, memory_scope::device,
29+
access::address_space::global_space>(sum[0]);
3030
out[gid] = atm.fetch_add(Difference(1));
3131
});
3232
});
@@ -60,9 +60,9 @@ void add_plus_equal_test(queue q, size_t N) {
6060
output_buf.template get_access<access::mode::discard_write>(cgh);
6161
cgh.parallel_for(range<1>(N), [=](item<1> it) {
6262
int gid = it.get_id(0);
63-
auto atm = ::sycl::ext::oneapi::atomic_ref<T, memory_order::relaxed,
64-
memory_scope::device,
65-
access::address_space::global_space>(sum[0]);
63+
auto atm = ::sycl::ext::oneapi::atomic_ref<
64+
T, memory_order::relaxed, memory_scope::device,
65+
access::address_space::global_space>(sum[0]);
6666
out[gid] = atm += Difference(1);
6767
});
6868
});
@@ -96,9 +96,9 @@ void add_pre_inc_test(queue q, size_t N) {
9696
output_buf.template get_access<access::mode::discard_write>(cgh);
9797
cgh.parallel_for(range<1>(N), [=](item<1> it) {
9898
int gid = it.get_id(0);
99-
auto atm = ::sycl::ext::oneapi::atomic_ref<T, memory_order::relaxed,
100-
memory_scope::device,
101-
access::address_space::global_space>(sum[0]);
99+
auto atm = ::sycl::ext::oneapi::atomic_ref<
100+
T, memory_order::relaxed, memory_scope::device,
101+
access::address_space::global_space>(sum[0]);
102102
out[gid] = ++atm;
103103
});
104104
});
@@ -132,9 +132,9 @@ void add_post_inc_test(queue q, size_t N) {
132132
output_buf.template get_access<access::mode::discard_write>(cgh);
133133
cgh.parallel_for(range<1>(N), [=](item<1> it) {
134134
int gid = it.get_id(0);
135-
auto atm = ::sycl::ext::oneapi::atomic_ref<T, memory_order::relaxed,
136-
memory_scope::device,
137-
access::address_space::global_space>(sum[0]);
135+
auto atm = ::sycl::ext::oneapi::atomic_ref<
136+
T, memory_order::relaxed, memory_scope::device,
137+
access::address_space::global_space>(sum[0]);
138138
out[gid] = atm++;
139139
});
140140
});

‎SYCL/AtomicRef/assignment.h

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -21,10 +21,10 @@ template <typename T> void assignment_test(queue q, size_t N) {
2121
assignment_buf.template get_access<access::mode::read_write>(cgh);
2222
cgh.parallel_for<assignment_kernel<T>>(range<1>(N), [=](item<1> it) {
2323
size_t gid = it.get_id(0);
24-
auto atm = ::sycl::ext::oneapi::atomic_ref<T,
25-
ext::oneapi::memory_order::relaxed,
26-
ext::oneapi::memory_scope::device,
27-
access::address_space::global_space>(st[0]);
24+
auto atm = ::sycl::ext::oneapi::atomic_ref<
25+
T, ext::oneapi::memory_order::relaxed,
26+
ext::oneapi::memory_scope::device,
27+
access::address_space::global_space>(st[0]);
2828
atm = T(gid);
2929
});
3030
});

‎SYCL/AtomicRef/atomic_memory_order_acq_rel.h

Lines changed: 3 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -19,10 +19,9 @@ template <typename T> void acq_rel_test(queue q, size_t N) {
1919
auto a_acc = a_buf.template get_access<access::mode::read_write>(cgh);
2020
cgh.parallel_for<atomic_memory_order_acq_rel_kernel<T>>(
2121
range<1>(N), [=](item<1> it) {
22-
auto aar =
23-
::sycl::ext::oneapi::atomic_ref<T, memory_order::acq_rel,
24-
memory_scope::device,
25-
access::address_space::global_space>(a_acc[0]);
22+
auto aar = ::sycl::ext::oneapi::atomic_ref<
23+
T, memory_order::acq_rel, memory_scope::device,
24+
access::address_space::global_space>(a_acc[0]);
2625
auto ld = aar.load();
2726
ld += 1;
2827
aar.store(ld);

‎SYCL/AtomicRef/atomic_memory_order_seq_cst.h

Lines changed: 6 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -22,14 +22,12 @@ template <typename T> void seq_cst_test(queue q, size_t N) {
2222
auto b_acc = b_buf.template get_access<access::mode::read_write>(cgh);
2323
cgh.parallel_for<atomic_memory_order_seq_cst_kernel<T>>(
2424
range<1>(N), [=](item<1> it) {
25-
auto aar =
26-
::sycl::ext::oneapi::atomic_ref<T, memory_order::seq_cst,
27-
memory_scope::device,
28-
access::address_space::global_space>(a_acc[0]);
29-
auto bar =
30-
::sycl::ext::oneapi::atomic_ref<T, memory_order::seq_cst,
31-
memory_scope::device,
32-
access::address_space::global_space>(b_acc[0]);
25+
auto aar = ::sycl::ext::oneapi::atomic_ref<
26+
T, memory_order::seq_cst, memory_scope::device,
27+
access::address_space::global_space>(a_acc[0]);
28+
auto bar = ::sycl::ext::oneapi::atomic_ref<
29+
T, memory_order::seq_cst, memory_scope::device,
30+
access::address_space::global_space>(b_acc[0]);
3331
auto ald = aar.load();
3432
auto bld = bar.load();
3533
ald += 1;

‎SYCL/AtomicRef/compare_exchange.h

Lines changed: 3 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -29,10 +29,9 @@ template <typename T> void compare_exchange_test(queue q, size_t N) {
2929
cgh.parallel_for<compare_exchange_kernel<T>>(
3030
range<1>(N), [=](item<1> it) {
3131
size_t gid = it.get_id(0);
32-
auto atm =
33-
::sycl::ext::oneapi::atomic_ref<T, memory_order::relaxed,
34-
memory_scope::device,
35-
access::address_space::global_space>(exc[0]);
32+
auto atm = ::sycl::ext::oneapi::atomic_ref<
33+
T, memory_order::relaxed, memory_scope::device,
34+
access::address_space::global_space>(exc[0]);
3635
T result = T(N); // Avoid copying pointer
3736
bool success = atm.compare_exchange_strong(result, (T)gid);
3837
if (success) {

‎SYCL/AtomicRef/exchange.h

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -27,9 +27,9 @@ template <typename T> void exchange_test(queue q, size_t N) {
2727
output_buf.template get_access<access::mode::discard_write>(cgh);
2828
cgh.parallel_for<exchange_kernel<T>>(range<1>(N), [=](item<1> it) {
2929
size_t gid = it.get_id(0);
30-
auto atm = ::sycl::ext::oneapi::atomic_ref<T, memory_order::relaxed,
31-
memory_scope::device,
32-
access::address_space::global_space>(exc[0]);
30+
auto atm = ::sycl::ext::oneapi::atomic_ref<
31+
T, memory_order::relaxed, memory_scope::device,
32+
access::address_space::global_space>(exc[0]);
3333
out[gid] = atm.exchange(T(gid));
3434
});
3535
});

‎SYCL/AtomicRef/load.h

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -26,9 +26,9 @@ template <typename T> void load_test(queue q, size_t N) {
2626
output_buf.template get_access<access::mode::discard_write>(cgh);
2727
cgh.parallel_for<load_kernel<T>>(range<1>(N), [=](item<1> it) {
2828
size_t gid = it.get_id(0);
29-
auto atm = ::sycl::ext::oneapi::atomic_ref<T, memory_order::relaxed,
30-
memory_scope::device,
31-
access::address_space::global_space>(ld[0]);
29+
auto atm = ::sycl::ext::oneapi::atomic_ref<
30+
T, memory_order::relaxed, memory_scope::device,
31+
access::address_space::global_space>(ld[0]);
3232
out[gid] = atm.load();
3333
});
3434
});

‎SYCL/AtomicRef/max.h

Lines changed: 3 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -24,10 +24,9 @@ template <typename T> void max_test(queue q, size_t N) {
2424
output_buf.template get_access<access::mode::discard_write>(cgh);
2525
cgh.parallel_for(range<1>(N), [=](item<1> it) {
2626
int gid = it.get_id(0);
27-
auto atm = ::sycl::ext::oneapi::atomic_ref<T, memory_order::relaxed,
28-
memory_scope::device,
29-
access::address_space::global_space>(val[0]);
30-
27+
auto atm = ::sycl::ext::oneapi::atomic_ref<
28+
T, memory_order::relaxed, memory_scope::device,
29+
access::address_space::global_space>(val[0]);
3130
// +1 accounts for lowest() returning 0 for unsigned types
3231
out[gid] = atm.fetch_max(T(gid) + 1);
3332
});

‎SYCL/AtomicRef/min.h

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -24,9 +24,9 @@ template <typename T> void min_test(queue q, size_t N) {
2424
output_buf.template get_access<access::mode::discard_write>(cgh);
2525
cgh.parallel_for(range<1>(N), [=](item<1> it) {
2626
int gid = it.get_id(0);
27-
auto atm = ::sycl::ext::oneapi::atomic_ref<T, memory_order::relaxed,
28-
memory_scope::device,
29-
access::address_space::global_space>(val[0]);
27+
auto atm = ::sycl::ext::oneapi::atomic_ref<
28+
T, memory_order::relaxed, memory_scope::device,
29+
access::address_space::global_space>(val[0]);
3030
out[gid] = atm.fetch_min(T(gid));
3131
});
3232
});

‎SYCL/AtomicRef/store.h

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -20,9 +20,9 @@ template <typename T> void store_test(queue q, size_t N) {
2020
auto st = store_buf.template get_access<access::mode::read_write>(cgh);
2121
cgh.parallel_for<store_kernel<T>>(range<1>(N), [=](item<1> it) {
2222
size_t gid = it.get_id(0);
23-
auto atm = ::sycl::ext::oneapi::atomic_ref<T, memory_order::relaxed,
24-
memory_scope::device,
25-
access::address_space::global_space>(st[0]);
23+
auto atm = ::sycl::ext::oneapi::atomic_ref<
24+
T, memory_order::relaxed, memory_scope::device,
25+
access::address_space::global_space>(st[0]);
2626
atm.store(T(gid));
2727
});
2828
});

‎SYCL/AtomicRef/sub.h

Lines changed: 12 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -24,9 +24,9 @@ void sub_fetch_test(queue q, size_t N) {
2424
output_buf.template get_access<access::mode::discard_write>(cgh);
2525
cgh.parallel_for(range<1>(N), [=](item<1> it) {
2626
int gid = it.get_id(0);
27-
auto atm = ::sycl::ext::oneapi::atomic_ref<T, memory_order::relaxed,
28-
memory_scope::device,
29-
access::address_space::global_space>(val[0]);
27+
auto atm = ::sycl::ext::oneapi::atomic_ref<
28+
T, memory_order::relaxed, memory_scope::device,
29+
access::address_space::global_space>(val[0]);
3030
out[gid] = atm.fetch_sub(Difference(1));
3131
});
3232
});
@@ -60,9 +60,9 @@ void sub_plus_equal_test(queue q, size_t N) {
6060
output_buf.template get_access<access::mode::discard_write>(cgh);
6161
cgh.parallel_for(range<1>(N), [=](item<1> it) {
6262
int gid = it.get_id(0);
63-
auto atm = ::sycl::ext::oneapi::atomic_ref<T, memory_order::relaxed,
64-
memory_scope::device,
65-
access::address_space::global_space>(val[0]);
63+
auto atm = ::sycl::ext::oneapi::atomic_ref<
64+
T, memory_order::relaxed, memory_scope::device,
65+
access::address_space::global_space>(val[0]);
6666
out[gid] = atm -= Difference(1);
6767
});
6868
});
@@ -96,9 +96,9 @@ void sub_pre_dec_test(queue q, size_t N) {
9696
output_buf.template get_access<access::mode::discard_write>(cgh);
9797
cgh.parallel_for(range<1>(N), [=](item<1> it) {
9898
int gid = it.get_id(0);
99-
auto atm = ::sycl::ext::oneapi::atomic_ref<T, memory_order::relaxed,
100-
memory_scope::device,
101-
access::address_space::global_space>(val[0]);
99+
auto atm = ::sycl::ext::oneapi::atomic_ref<
100+
T, memory_order::relaxed, memory_scope::device,
101+
access::address_space::global_space>(val[0]);
102102
out[gid] = --atm;
103103
});
104104
});
@@ -132,9 +132,9 @@ void sub_post_dec_test(queue q, size_t N) {
132132
output_buf.template get_access<access::mode::discard_write>(cgh);
133133
cgh.parallel_for(range<1>(N), [=](item<1> it) {
134134
int gid = it.get_id(0);
135-
auto atm = ::sycl::ext::oneapi::atomic_ref<T, memory_order::relaxed,
136-
memory_scope::device,
137-
access::address_space::global_space>(val[0]);
135+
auto atm = ::sycl::ext::oneapi::atomic_ref<
136+
T, memory_order::relaxed, memory_scope::device,
137+
access::address_space::global_space>(val[0]);
138138
out[gid] = atm--;
139139
});
140140
});

0 commit comments

Comments
 (0)