[SYCL][Reduction] Group counter should use at least memory_order::acq…

…_rel

From https://en.cppreference.com/w/cpp/atomic/memory_order:

> Atomic operations tagged memory_order_relaxed are not synchronization
> operations; they do not impose an order among concurrent memory
> accesses. They only guarantee atomicity and modification order
> consistency.

Yet we want to use that counter exactly for the synchronization purposes
- to decide which WG finished last and should perform the final step.

The same issue had been fixed in DPC++ at intel/llvm#8058.

core/src/SYCL/Kokkos_SYCL_ParallelReduce_MDRange.hpp

@@ -215,7 +215,7 @@ class Kokkos::Impl::ParallelReduce<CombinedFunctorReducerType,
                     value_count, reducer, false, wgroup_size);
 
                 if (local_id == 0) {
-                  sycl::atomic_ref<unsigned, sycl::memory_order::relaxed,
+                  sycl::atomic_ref<unsigned, sycl::memory_order::acq_rel,
                                    sycl::memory_scope::device,
                                    sycl::access::address_space::global_space>
                       scratch_flags_ref(*scratch_flags);
@@ -260,7 +260,7 @@ class Kokkos::Impl::ParallelReduce<CombinedFunctorReducerType,
                     device_accessible_result_ptr, reducer, false, wgroup_size);
 
                 if (local_id == 0) {
-                  sycl::atomic_ref<unsigned, sycl::memory_order::relaxed,
+                  sycl::atomic_ref<unsigned, sycl::memory_order::acq_rel,
                                    sycl::memory_scope::device,
                                    sycl::access::address_space::global_space>
                       scratch_flags_ref(*scratch_flags);

core/src/SYCL/Kokkos_SYCL_ParallelReduce_Range.hpp

@@ -160,7 +160,7 @@ class Kokkos::Impl::ParallelReduce<CombinedFunctorReducerType,
                     value_count, reducer, false, std::min(size, wgroup_size));
 
                 if (local_id == 0) {
-                  sycl::atomic_ref<unsigned, sycl::memory_order::relaxed,
+                  sycl::atomic_ref<unsigned, sycl::memory_order::acq_rel,
                                    sycl::memory_scope::device,
                                    sycl::access::address_space::global_space>
                       scratch_flags_ref(*scratch_flags);
@@ -202,7 +202,7 @@ class Kokkos::Impl::ParallelReduce<CombinedFunctorReducerType,
                     std::min(size, wgroup_size));
 
                 if (local_id == 0) {
-                  sycl::atomic_ref<unsigned, sycl::memory_order::relaxed,
+                  sycl::atomic_ref<unsigned, sycl::memory_order::acq_rel,
                                    sycl::memory_scope::device,
                                    sycl::access::address_space::global_space>
                       scratch_flags_ref(*scratch_flags);

core/src/SYCL/Kokkos_SYCL_ParallelReduce_Team.hpp

@@ -208,7 +208,7 @@ class Kokkos::Impl::ParallelReduce<CombinedFunctorReducerType,
                                                 item.get_local_range()[1]));
 
                   if (local_id == 0) {
-                    sycl::atomic_ref<unsigned, sycl::memory_order::relaxed,
+                    sycl::atomic_ref<unsigned, sycl::memory_order::acq_rel,
                                      sycl::memory_scope::device,
                                      sycl::access::address_space::global_space>
                         scratch_flags_ref(*scratch_flags);
@@ -260,7 +260,7 @@ class Kokkos::Impl::ParallelReduce<CombinedFunctorReducerType,
                                                 item.get_local_range()[1]));
 
                   if (local_id == 0) {
-                    sycl::atomic_ref<unsigned, sycl::memory_order::relaxed,
+                    sycl::atomic_ref<unsigned, sycl::memory_order::acq_rel,
                                      sycl::memory_scope::device,
                                      sycl::access::address_space::global_space>
                         scratch_flags_ref(*scratch_flags);