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

Commit 080cbe0

Browse files
author
Pavel Samolysov
authored
Add tests for newly introduced sycl::atomic_ref class (#548)
Signed-off-by: Pavel Samolysov <[email protected]>
1 parent 6bb47ee commit 080cbe0

15 files changed

+220
-111
lines changed

‎SYCL/AtomicRef/add.cpp

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -11,8 +11,10 @@ using namespace sycl;
1111

1212
// Floating-point types do not support pre- or post-increment
1313
template <> void add_test<float>(queue q, size_t N) {
14-
add_fetch_test<float>(q, N);
15-
add_plus_equal_test<float>(q, N);
14+
add_fetch_test<::sycl::ext::oneapi::atomic_ref, float>(q, N);
15+
add_fetch_test<::sycl::atomic_ref, float>(q, N);
16+
add_plus_equal_test<::sycl::ext::oneapi::atomic_ref, float>(q, N);
17+
add_plus_equal_test<::sycl::atomic_ref, float>(q, N);
1618
}
1719

1820
int main() {

‎SYCL/AtomicRef/add.h

Lines changed: 28 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -9,7 +9,9 @@
99
using namespace sycl;
1010
using namespace sycl::ext::oneapi;
1111

12-
template <typename T, typename Difference = T>
12+
template <template <typename, memory_order, memory_scope, access::address_space>
13+
class AtomicRef,
14+
typename T, typename Difference = T>
1315
void add_fetch_test(queue q, size_t N) {
1416
T sum = 0;
1517
std::vector<T> output(N);
@@ -24,9 +26,8 @@ void add_fetch_test(queue q, size_t N) {
2426
output_buf.template get_access<access::mode::discard_write>(cgh);
2527
cgh.parallel_for(range<1>(N), [=](item<1> it) {
2628
int gid = it.get_id(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]);
29+
auto atm = AtomicRef<T, memory_order::relaxed, memory_scope::device,
30+
access::address_space::global_space>(sum[0]);
3031
out[gid] = atm.fetch_add(Difference(1));
3132
});
3233
});
@@ -45,7 +46,9 @@ void add_fetch_test(queue q, size_t N) {
4546
assert(std::unique(output.begin(), output.end()) == output.end());
4647
}
4748

48-
template <typename T, typename Difference = T>
49+
template <template <typename, memory_order, memory_scope, access::address_space>
50+
class AtomicRef,
51+
typename T, typename Difference = T>
4952
void add_plus_equal_test(queue q, size_t N) {
5053
T sum = 0;
5154
std::vector<T> output(N);
@@ -60,9 +63,8 @@ void add_plus_equal_test(queue q, size_t N) {
6063
output_buf.template get_access<access::mode::discard_write>(cgh);
6164
cgh.parallel_for(range<1>(N), [=](item<1> it) {
6265
int gid = it.get_id(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]);
66+
auto atm = AtomicRef<T, memory_order::relaxed, memory_scope::device,
67+
access::address_space::global_space>(sum[0]);
6668
out[gid] = atm += Difference(1);
6769
});
6870
});
@@ -81,7 +83,9 @@ void add_plus_equal_test(queue q, size_t N) {
8183
assert(std::unique(output.begin(), output.end()) == output.end());
8284
}
8385

84-
template <typename T, typename Difference = T>
86+
template <template <typename, memory_order, memory_scope, access::address_space>
87+
class AtomicRef,
88+
typename T, typename Difference = T>
8589
void add_pre_inc_test(queue q, size_t N) {
8690
T sum = 0;
8791
std::vector<T> output(N);
@@ -96,9 +100,8 @@ void add_pre_inc_test(queue q, size_t N) {
96100
output_buf.template get_access<access::mode::discard_write>(cgh);
97101
cgh.parallel_for(range<1>(N), [=](item<1> it) {
98102
int gid = it.get_id(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]);
103+
auto atm = AtomicRef<T, memory_order::relaxed, memory_scope::device,
104+
access::address_space::global_space>(sum[0]);
102105
out[gid] = ++atm;
103106
});
104107
});
@@ -117,7 +120,9 @@ void add_pre_inc_test(queue q, size_t N) {
117120
assert(std::unique(output.begin(), output.end()) == output.end());
118121
}
119122

120-
template <typename T, typename Difference = T>
123+
template <template <typename, memory_order, memory_scope, access::address_space>
124+
class AtomicRef,
125+
typename T, typename Difference = T>
121126
void add_post_inc_test(queue q, size_t N) {
122127
T sum = 0;
123128
std::vector<T> output(N);
@@ -132,9 +137,8 @@ void add_post_inc_test(queue q, size_t N) {
132137
output_buf.template get_access<access::mode::discard_write>(cgh);
133138
cgh.parallel_for(range<1>(N), [=](item<1> it) {
134139
int gid = it.get_id(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]);
140+
auto atm = AtomicRef<T, memory_order::relaxed, memory_scope::device,
141+
access::address_space::global_space>(sum[0]);
138142
out[gid] = atm++;
139143
});
140144
});
@@ -155,8 +159,12 @@ void add_post_inc_test(queue q, size_t N) {
155159

156160
template <typename T, typename Difference = T>
157161
void add_test(queue q, size_t N) {
158-
add_fetch_test<T, Difference>(q, N);
159-
add_plus_equal_test<T, Difference>(q, N);
160-
add_pre_inc_test<T, Difference>(q, N);
161-
add_post_inc_test<T, Difference>(q, N);
162+
add_fetch_test<::sycl::ext::oneapi::atomic_ref, T, Difference>(q, N);
163+
add_fetch_test<::sycl::atomic_ref, T, Difference>(q, N);
164+
add_plus_equal_test<::sycl::ext::oneapi::atomic_ref, T, Difference>(q, N);
165+
add_plus_equal_test<::sycl::atomic_ref, T, Difference>(q, N);
166+
add_pre_inc_test<::sycl::ext::oneapi::atomic_ref, T, Difference>(q, N);
167+
add_pre_inc_test<::sycl::atomic_ref, T, Difference>(q, N);
168+
add_post_inc_test<::sycl::ext::oneapi::atomic_ref, T, Difference>(q, N);
169+
add_post_inc_test<::sycl::atomic_ref, T, Difference>(q, N);
162170
}

‎SYCL/AtomicRef/add_atomic64.cpp

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -11,8 +11,10 @@ using namespace sycl;
1111

1212
// Floating-point types do not support pre- or post-increment
1313
template <> void add_test<double>(queue q, size_t N) {
14-
add_fetch_test<double>(q, N);
15-
add_plus_equal_test<double>(q, N);
14+
add_fetch_test<::sycl::ext::oneapi::atomic_ref, double>(q, N);
15+
add_fetch_test<::sycl::atomic_ref, double>(q, N);
16+
add_plus_equal_test<::sycl::ext::oneapi::atomic_ref, double>(q, N);
17+
add_plus_equal_test<::sycl::atomic_ref, double>(q, N);
1618
}
1719

1820
int main() {

‎SYCL/AtomicRef/assignment.h

Lines changed: 20 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -9,24 +9,29 @@
99
using namespace sycl;
1010
using namespace sycl::ext::oneapi;
1111

12-
template <typename T> class assignment_kernel;
12+
template <template <typename, memory_order, memory_scope, access::address_space>
13+
class AtomicRef,
14+
typename T>
15+
class assignment_kernel;
1316

14-
template <typename T> void assignment_test(queue q, size_t N) {
17+
template <template <typename, memory_order, memory_scope, access::address_space>
18+
class AtomicRef,
19+
typename T>
20+
void assignment_test(queue q, size_t N) {
1521
T initial = T(N);
1622
T assignment = initial;
1723
{
1824
buffer<T> assignment_buf(&assignment, 1);
1925
q.submit([&](handler &cgh) {
2026
auto st =
2127
assignment_buf.template get_access<access::mode::read_write>(cgh);
22-
cgh.parallel_for<assignment_kernel<T>>(range<1>(N), [=](item<1> it) {
23-
size_t gid = it.get_id(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]);
28-
atm = T(gid);
29-
});
28+
cgh.parallel_for<assignment_kernel<AtomicRef, T>>(
29+
range<1>(N), [=](item<1> it) {
30+
size_t gid = it.get_id(0);
31+
auto atm = AtomicRef<T, memory_order::relaxed, memory_scope::device,
32+
access::address_space::global_space>(st[0]);
33+
atm = T(gid);
34+
});
3035
});
3136
}
3237

@@ -35,3 +40,8 @@ template <typename T> void assignment_test(queue q, size_t N) {
3540
assert(assignment != initial);
3641
assert(assignment >= T(0) && assignment <= T(N - 1));
3742
}
43+
44+
template <typename T> void assignment_test(queue q, size_t N) {
45+
assignment_test<::sycl::ext::oneapi::atomic_ref, T>(q, N);
46+
assignment_test<::sycl::atomic_ref, T>(q, N);
47+
}

‎SYCL/AtomicRef/atomic_memory_order_acq_rel.h

Lines changed: 16 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -8,20 +8,25 @@
88
using namespace sycl;
99
using namespace sycl::ext::oneapi;
1010

11-
template <typename T> class atomic_memory_order_acq_rel_kernel;
11+
template <template <typename, memory_order, memory_scope, access::address_space>
12+
class AtomicRef,
13+
typename T>
14+
class atomic_memory_order_acq_rel_kernel;
1215

13-
template <typename T> void acq_rel_test(queue q, size_t N) {
16+
template <template <typename, memory_order, memory_scope, access::address_space>
17+
class AtomicRef,
18+
typename T>
19+
void acq_rel_test(queue q, size_t N) {
1420
T a = 0;
1521
{
1622
buffer<T> a_buf(&a, 1);
1723

1824
q.submit([&](handler &cgh) {
1925
auto a_acc = a_buf.template get_access<access::mode::read_write>(cgh);
20-
cgh.parallel_for<atomic_memory_order_acq_rel_kernel<T>>(
26+
cgh.parallel_for<atomic_memory_order_acq_rel_kernel<AtomicRef, T>>(
2127
range<1>(N), [=](item<1> it) {
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]);
28+
auto aar = AtomicRef<T, memory_order::acq_rel, memory_scope::device,
29+
access::address_space::global_space>(a_acc[0]);
2530
auto ld = aar.load();
2631
ld += 1;
2732
aar.store(ld);
@@ -32,3 +37,8 @@ template <typename T> void acq_rel_test(queue q, size_t N) {
3237
// All work-items increment by 1, so final value should be equal to N
3338
assert(a == T(N));
3439
}
40+
41+
template <typename T> void acq_rel_test(queue q, size_t N) {
42+
acq_rel_test<::sycl::ext::oneapi::atomic_ref, T>(q, N);
43+
acq_rel_test<::sycl::atomic_ref, T>(q, N);
44+
}

‎SYCL/AtomicRef/atomic_memory_order_seq_cst.h

Lines changed: 18 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -8,9 +8,15 @@
88
using namespace sycl;
99
using namespace sycl::ext::oneapi;
1010

11-
template <typename T> class atomic_memory_order_seq_cst_kernel;
11+
template <template <typename, memory_order, memory_scope, access::address_space>
12+
class AtomicRef,
13+
typename T>
14+
class atomic_memory_order_seq_cst_kernel;
1215

13-
template <typename T> void seq_cst_test(queue q, size_t N) {
16+
template <template <typename, memory_order, memory_scope, access::address_space>
17+
class AtomicRef,
18+
typename T>
19+
void seq_cst_test(queue q, size_t N) {
1420
T a = 0;
1521
T b = 0;
1622
{
@@ -20,14 +26,12 @@ template <typename T> void seq_cst_test(queue q, size_t N) {
2026
q.submit([&](handler &cgh) {
2127
auto a_acc = a_buf.template get_access<access::mode::read_write>(cgh);
2228
auto b_acc = b_buf.template get_access<access::mode::read_write>(cgh);
23-
cgh.parallel_for<atomic_memory_order_seq_cst_kernel<T>>(
29+
cgh.parallel_for<atomic_memory_order_seq_cst_kernel<AtomicRef, T>>(
2430
range<1>(N), [=](item<1> it) {
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]);
31+
auto aar = AtomicRef<T, memory_order::seq_cst, memory_scope::device,
32+
access::address_space::global_space>(a_acc[0]);
33+
auto bar = AtomicRef<T, memory_order::seq_cst, memory_scope::device,
34+
access::address_space::global_space>(b_acc[0]);
3135
auto ald = aar.load();
3236
auto bld = bar.load();
3337
ald += 1;
@@ -46,3 +50,8 @@ template <typename T> void seq_cst_test(queue q, size_t N) {
4650
rsum += i;
4751
assert(b == T(rsum));
4852
}
53+
54+
template <typename T> void seq_cst_test(queue q, size_t N) {
55+
seq_cst_test<::sycl::ext::oneapi::atomic_ref, T>(q, N);
56+
seq_cst_test<::sycl::atomic_ref, T>(q, N);
57+
}

‎SYCL/AtomicRef/compare_exchange.h

Lines changed: 16 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -9,9 +9,15 @@
99
using namespace sycl;
1010
using namespace sycl::ext::oneapi;
1111

12-
template <typename T> class compare_exchange_kernel;
12+
template <template <typename, memory_order, memory_scope, access::address_space>
13+
class AtomicRef,
14+
typename T>
15+
class compare_exchange_kernel;
1316

14-
template <typename T> void compare_exchange_test(queue q, size_t N) {
17+
template <template <typename, memory_order, memory_scope, access::address_space>
18+
class AtomicRef,
19+
typename T>
20+
void compare_exchange_test(queue q, size_t N) {
1521
const T initial = T(N);
1622
T compare_exchange = initial;
1723
std::vector<T> output(N);
@@ -26,12 +32,11 @@ template <typename T> void compare_exchange_test(queue q, size_t N) {
2632
cgh);
2733
auto out =
2834
output_buf.template get_access<access::mode::discard_write>(cgh);
29-
cgh.parallel_for<compare_exchange_kernel<T>>(
35+
cgh.parallel_for<compare_exchange_kernel<AtomicRef, T>>(
3036
range<1>(N), [=](item<1> it) {
3137
size_t gid = it.get_id(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]);
38+
auto atm = AtomicRef<T, memory_order::relaxed, memory_scope::device,
39+
access::address_space::global_space>(exc[0]);
3540
T result = T(N); // Avoid copying pointer
3641
bool success = atm.compare_exchange_strong(result, (T)gid);
3742
if (success) {
@@ -51,3 +56,8 @@ template <typename T> void compare_exchange_test(queue q, size_t N) {
5156
assert(output[i] == T(i) || output[i] == initial);
5257
}
5358
}
59+
60+
template <typename T> void compare_exchange_test(queue q, size_t N) {
61+
compare_exchange_test<::sycl::ext::oneapi::atomic_ref, T>(q, N);
62+
compare_exchange_test<::sycl::atomic_ref, T>(q, N);
63+
}

‎SYCL/AtomicRef/exchange.h

Lines changed: 20 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -9,9 +9,15 @@
99
using namespace sycl;
1010
using namespace sycl::ext::oneapi;
1111

12-
template <typename T> class exchange_kernel;
12+
template <template <typename, memory_order, memory_scope, access::address_space>
13+
class AtomicRef,
14+
typename T>
15+
class exchange_kernel;
1316

14-
template <typename T> void exchange_test(queue q, size_t N) {
17+
template <template <typename, memory_order, memory_scope, access::address_space>
18+
class AtomicRef,
19+
typename T>
20+
void exchange_test(queue q, size_t N) {
1521
const T initial = T(N);
1622
T exchange = initial;
1723
std::vector<T> output(N);
@@ -25,13 +31,13 @@ template <typename T> void exchange_test(queue q, size_t N) {
2531
exchange_buf.template get_access<access::mode::read_write>(cgh);
2632
auto out =
2733
output_buf.template get_access<access::mode::discard_write>(cgh);
28-
cgh.parallel_for<exchange_kernel<T>>(range<1>(N), [=](item<1> it) {
29-
size_t gid = it.get_id(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]);
33-
out[gid] = atm.exchange(T(gid));
34-
});
34+
cgh.parallel_for<exchange_kernel<AtomicRef, T>>(
35+
range<1>(N), [=](item<1> it) {
36+
size_t gid = it.get_id(0);
37+
auto atm = AtomicRef<T, memory_order::relaxed, memory_scope::device,
38+
access::address_space::global_space>(exc[0]);
39+
out[gid] = atm.exchange(T(gid));
40+
});
3541
});
3642
}
3743

@@ -43,3 +49,8 @@ template <typename T> void exchange_test(queue q, size_t N) {
4349
std::sort(output.begin(), output.end());
4450
assert(std::unique(output.begin(), output.end()) == output.end());
4551
}
52+
53+
template <typename T> void exchange_test(queue q, size_t N) {
54+
exchange_test<::sycl::ext::oneapi::atomic_ref, T>(q, N);
55+
exchange_test<::sycl::atomic_ref, T>(q, N);
56+
}

0 commit comments

Comments
 (0)