@@ -14,8 +14,7 @@ using namespace sycl::intel;
14
14
15
15
// Equivalent to add_test from add.cpp
16
16
// Uses atomic_accessor instead of atomic_ref
17
- template <typename T>
18
- void accessor_test (queue q, size_t N) {
17
+ template <typename T> void accessor_test (queue q, size_t N) {
19
18
T sum = 0 ;
20
19
std::vector<T> output (N, 0 );
21
20
{
@@ -24,15 +23,26 @@ void accessor_test(queue q, size_t N) {
24
23
25
24
q.submit ([&](handler &cgh) {
26
25
#if __cplusplus > 201402L
27
- static_assert (std::is_same<decltype (atomic_accessor (sum_buf, cgh, relaxed_order, device_scope)),
28
- atomic_accessor<T, 1 , intel::memory_order::relaxed, intel::memory_scope::device>>::value,
29
- " atomic_accessor type incorrectly deduced" );
26
+ static_assert (
27
+ std::is_same<decltype (atomic_accessor (sum_buf, cgh, relaxed_order,
28
+ device_scope)),
29
+ atomic_accessor<T, 1 , intel::memory_order::relaxed,
30
+ intel::memory_scope::device>>::value,
31
+ " atomic_accessor type incorrectly deduced" );
30
32
#endif
31
- auto sum = atomic_accessor<T, 1 , intel::memory_order::relaxed, intel::memory_scope::device>(sum_buf, cgh);
32
- auto out = output_buf.template get_access <access::mode::discard_write>(cgh);
33
+ auto sum = atomic_accessor<T, 1 , intel::memory_order::relaxed,
34
+ intel::memory_scope::device>(sum_buf, cgh);
35
+ auto out =
36
+ output_buf.template get_access <access::mode::discard_write>(cgh);
33
37
cgh.parallel_for (range<1 >(N), [=](item<1 > it) {
34
38
int gid = it.get_id (0 );
35
- static_assert (std::is_same<decltype (sum[0 ]), atomic_ref<T, intel::memory_order::relaxed, intel::memory_scope::device, access::address_space::global_space>>::value, " atomic_accessor returns incorrect atomic_ref" );
39
+ static_assert (
40
+ std::is_same<
41
+ decltype (sum[0 ]),
42
+ atomic_ref<T, intel::memory_order::relaxed,
43
+ intel::memory_scope::device,
44
+ access::address_space::global_space>>::value,
45
+ " atomic_accessor returns incorrect atomic_ref" );
36
46
out[gid] = sum[0 ].fetch_add (T (1 ));
37
47
});
38
48
});
@@ -59,13 +69,21 @@ void local_accessor_test(queue q, size_t N, size_t L = 8) {
59
69
{
60
70
buffer<T> output_buf (output.data (), output.size ());
61
71
q.submit ([&](handler &cgh) {
62
- auto sum = atomic_accessor<T, 1 , intel::memory_order::relaxed, intel::memory_scope::device, access::target::local>(1 , cgh);
72
+ auto sum =
73
+ atomic_accessor<T, 1 , intel::memory_order::relaxed,
74
+ intel::memory_scope::device, access::target::local>(
75
+ 1 , cgh);
63
76
auto out = output_buf.template get_access <access::mode::read_write>(cgh);
64
77
cgh.parallel_for (nd_range<1 >(N, L), [=](nd_item<1 > it) {
65
78
int grp = it.get_group (0 );
66
79
sum[0 ].store (0 );
67
80
it.barrier ();
68
- static_assert (std::is_same<decltype (sum[0 ]), atomic_ref<T, intel::memory_order::relaxed, intel::memory_scope::device, access::address_space::local_space>>::value, " local atomic_accessor returns incorrect atomic_ref" );
81
+ static_assert (
82
+ std::is_same<decltype (sum[0 ]),
83
+ atomic_ref<T, intel::memory_order::relaxed,
84
+ intel::memory_scope::device,
85
+ access::address_space::local_space>>::value,
86
+ " local atomic_accessor returns incorrect atomic_ref" );
69
87
T result = sum[0 ].fetch_add (T (1 ));
70
88
if (result == it.get_local_range (0 ) - 1 ) {
71
89
out[grp] = result;
@@ -76,7 +94,8 @@ void local_accessor_test(queue q, size_t N, size_t L = 8) {
76
94
77
95
// All work-items increment by 1, and last in the group writes out old value
78
96
// All values should be L-1
79
- assert (std::all_of (output.begin (), output.end (), [=](T x) { return x == L - 1 ; }));
97
+ assert (std::all_of (output.begin (), output.end (),
98
+ [=](T x) { return x == L - 1 ; }));
80
99
}
81
100
82
101
int main () {
0 commit comments