[SYCL] Added tests for atomics with various memory orders and scopes

SYCL/AtomicRef/add.cpp

@@ -9,12 +9,6 @@
 #include <iostream>
 using namespace sycl;
 
-// Floating-point types do not support pre- or post-increment
-template <> void add_test<float>(queue q, size_t N) {
-  add_fetch_test<float>(q, N);
-  add_plus_equal_test<float>(q, N);
-}
-
 int main() {
   queue q;
 

SYCL/AtomicRef/add.h

@@ -4,31 +4,84 @@
 #include <algorithm>
 #include <cassert>
 #include <numeric>
+#include <type_traits>
 #include <vector>
 
 using namespace sycl;
 using namespace sycl::ext::oneapi;
 
-template <typename T, typename Difference = T>
+template <typename T, typename Difference = T,
+          memory_order order = memory_order::relaxed,
+          memory_scope scope = memory_scope::device>
+void add_fetch_local_test(queue q, size_t N) {
+  T sum = 0;
+  std::vector<T> output(N);
+  std::fill(output.begin(), output.end(), T(123456));
+  {
+    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 out =
+           output_buf.template get_access<access::mode::discard_write>(cgh);
+       accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
+                                                                           cgh);
+
+       cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
+         int gid = it.get_global_id(0);
+         if (gid == 0)
+           loc[0] = 0;
+         it.barrier(access::fence_space::local_space);
+         auto atm = atomic_ref < T,
+              (order == memory_order::acquire || order == memory_order::release)
+                  ? memory_order::relaxed
+                  : order,
+              scope, access::address_space::local_space > (loc[0]);
+         out[gid] = atm.fetch_add(Difference(1), order);
+         it.barrier(access::fence_space::local_space);
+         if (gid == 0)
+           sum[0] = loc[0];
+       });
+     }).wait_and_throw();
+  }
+
+  // All work-items increment by 1, so final value should be equal to N
+  assert(sum == T(N));
+
+  // Fetch returns original value: will be in [0, N-1]
+  auto min_e = std::min_element(output.begin(), output.end());
+  auto max_e = std::max_element(output.begin(), output.end());
+  assert(*min_e == 0 && *max_e == T(N - 1));
+
+  // Intermediate values should be unique
+  std::sort(output.begin(), output.end());
+  assert(std::unique(output.begin(), output.end()) == output.end());
+}
+
+template <typename T, typename Difference = T,
+          memory_order order = memory_order::relaxed,
+          memory_scope scope = memory_scope::device>
 void add_fetch_test(queue q, size_t N) {
   T sum = 0;
   std::vector<T> output(N);
   std::fill(output.begin(), output.end(), T(0));
   {
     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 out =
-          output_buf.template get_access<access::mode::discard_write>(cgh);
-      cgh.parallel_for(range<1>(N), [=](item<1> it) {
-        int gid = it.get_id(0);
-        auto atm = atomic_ref<T, memory_order::relaxed, memory_scope::device,
-                              access::address_space::global_space>(sum[0]);
-        out[gid] = atm.fetch_add(Difference(1));
-      });
-    });
+       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) {
+         int gid = it.get_id(0);
+         auto atm = atomic_ref < T,
+              (order == memory_order::acquire || order == memory_order::release)
+                  ? memory_order::relaxed
+                  : order,
+              scope, access::address_space::global_space > (sum[0]);
+         out[gid] = atm.fetch_add(Difference(1), order);
+       });
+     }).wait_and_throw();
   }
 
   // All work-items increment by 1, so final value should be equal to N
@@ -37,14 +90,16 @@ void add_fetch_test(queue q, size_t N) {
   // Fetch returns original value: will be in [0, N-1]
   auto min_e = std::min_element(output.begin(), output.end());
   auto max_e = std::max_element(output.begin(), output.end());
-  assert(*min_e == T(0) && *max_e == T(N - 1));
+  assert(*min_e == 0 && *max_e == T(N - 1));
 
   // Intermediate values should be unique
   std::sort(output.begin(), output.end());
   assert(std::unique(output.begin(), output.end()) == output.end());
 }
 
-template <typename T, typename Difference = T>
+template <typename T, typename Difference = T,
+          memory_order order = memory_order::relaxed,
+          memory_scope scope = memory_scope::device>
 void add_plus_equal_test(queue q, size_t N) {
   T sum = 0;
   std::vector<T> output(N);
@@ -59,8 +114,11 @@ void add_plus_equal_test(queue q, size_t N) {
           output_buf.template get_access<access::mode::discard_write>(cgh);
       cgh.parallel_for(range<1>(N), [=](item<1> it) {
         int gid = it.get_id(0);
-        auto atm = atomic_ref<T, memory_order::relaxed, memory_scope::device,
-                              access::address_space::global_space>(sum[0]);
+        auto atm = atomic_ref < T,
+             (order == memory_order::acquire || order == memory_order::release)
+                 ? memory_order::relaxed
+                 : order,
+             scope, access::address_space::global_space > (sum[0]);
         out[gid] = atm += Difference(1);
       });
     });
@@ -79,7 +137,9 @@ void add_plus_equal_test(queue q, size_t N) {
   assert(std::unique(output.begin(), output.end()) == output.end());
 }
 
-template <typename T, typename Difference = T>
+template <typename T, typename Difference = T,
+          memory_order order = memory_order::relaxed,
+          memory_scope scope = memory_scope::device>
 void add_pre_inc_test(queue q, size_t N) {
   T sum = 0;
   std::vector<T> output(N);
@@ -94,8 +154,11 @@ void add_pre_inc_test(queue q, size_t N) {
           output_buf.template get_access<access::mode::discard_write>(cgh);
       cgh.parallel_for(range<1>(N), [=](item<1> it) {
         int gid = it.get_id(0);
-        auto atm = atomic_ref<T, memory_order::relaxed, memory_scope::device,
-                              access::address_space::global_space>(sum[0]);
+        auto atm = atomic_ref < T,
+             (order == memory_order::acquire || order == memory_order::release)
+                 ? memory_order::relaxed
+                 : order,
+             scope, access::address_space::global_space > (sum[0]);
         out[gid] = ++atm;
       });
     });
@@ -114,7 +177,9 @@ void add_pre_inc_test(queue q, size_t N) {
   assert(std::unique(output.begin(), output.end()) == output.end());
 }
 
-template <typename T, typename Difference = T>
+template <typename T, typename Difference = T,
+          memory_order order = memory_order::relaxed,
+          memory_scope scope = memory_scope::device>
 void add_post_inc_test(queue q, size_t N) {
   T sum = 0;
   std::vector<T> output(N);
@@ -129,8 +194,11 @@ void add_post_inc_test(queue q, size_t N) {
           output_buf.template get_access<access::mode::discard_write>(cgh);
       cgh.parallel_for(range<1>(N), [=](item<1> it) {
         int gid = it.get_id(0);
-        auto atm = atomic_ref<T, memory_order::relaxed, memory_scope::device,
-                              access::address_space::global_space>(sum[0]);
+        auto atm = atomic_ref < T,
+             (order == memory_order::acquire || order == memory_order::release)
+                 ? memory_order::relaxed
+                 : order,
+             scope, access::address_space::global_space > (sum[0]);
         out[gid] = atm++;
       });
     });
@@ -149,10 +217,15 @@ void add_post_inc_test(queue q, size_t N) {
   assert(std::unique(output.begin(), output.end()) == output.end());
 }
 
-template <typename T, typename Difference = T>
+template <typename T, typename Difference = T,
+          memory_order order = memory_order::relaxed,
+          memory_scope scope = memory_scope::device>
 void add_test(queue q, size_t N) {
-  add_fetch_test<T, Difference>(q, N);
-  add_plus_equal_test<T, Difference>(q, N);
-  add_pre_inc_test<T, Difference>(q, N);
-  add_post_inc_test<T, Difference>(q, N);
+  add_fetch_local_test<T, Difference, order, scope>(q, N);
+  add_fetch_test<T, Difference, order, scope>(q, N);
+  add_plus_equal_test<T, Difference, order, scope>(q, N);
+  if constexpr (!std::is_floating_point_v<T>) {
+    add_pre_inc_test<T, Difference, order, scope>(q, N);
+    add_post_inc_test<T, Difference, order, scope>(q, N);
+  }
 }

SYCL/AtomicRef/add_atomic64.cpp

@@ -9,12 +9,6 @@
 #include <iostream>
 using namespace sycl;
 
-// Floating-point types do not support pre- or post-increment
-template <> void add_test<double>(queue q, size_t N) {
-  add_fetch_test<double>(q, N);
-  add_plus_equal_test<double>(q, N);
-}
-
 int main() {
   queue q;
 

SYCL/AtomicRef/add_orders_scopes.cpp

@@ -0,0 +1,49 @@
+// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-targets=%sycl_triple %s -o %t.out \
+// RUN: -Xsycl-target-backend --cuda-gpu-arch=sm_70
+// RUN: %HOST_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+#define SYCL_USE_NATIVE_FP_ATOMICS
+
+#include "add.h"
+#include <iostream>
+using namespace sycl;
+
+template <typename T, typename Difference = T,
+          memory_order order = memory_order::relaxed>
+void add_test_scopes(queue q, size_t N) {
+  add_test<T, Difference, order, memory_scope::system>(q, N);
+  add_test<T, Difference, order, memory_scope::device>(q, N);
+  add_test<T, Difference, order, memory_scope::work_group>(q, N);
+  add_test<T, Difference, order, memory_scope::sub_group>(q, N);
+}
+
+template <typename T, typename Difference = T>
+void add_test_orders_scopes(queue q, size_t N) {
+  add_test_scopes<T, Difference, memory_order::relaxed>(q, N);
+  add_test_scopes<T, Difference, memory_order::acquire>(q, N);
+  add_test_scopes<T, Difference, memory_order::release>(q, N);
+  add_test_scopes<T, Difference, memory_order::acq_rel>(q, N);
+}
+
+int main() {
+  queue q;
+
+  constexpr int N = 32;
+  add_test_orders_scopes<int>(q, N);
+  add_test_orders_scopes<float>(q, N);
+  add_test_orders_scopes<unsigned int>(q, N);
+  add_test_orders_scopes<double>(q, N);
+  add_test_orders_scopes<long>(q, N);
+  add_test_orders_scopes<unsigned long>(q, N);
+
+  // Include long long tests if they are 64 bits wide
+  if constexpr (sizeof(long long) == 8) {
+    add_test_orders_scopes<long long>(q, N);
+    add_test_orders_scopes<unsigned long long>(q, N);
+  }
+
+  std::cout << "Test passed." << std::endl;
+}

SYCL/AtomicRef/and.h

@@ -0,0 +1,96 @@
+#pragma once
+
+#include <CL/sycl.hpp>
+#include <algorithm>
+#include <cassert>
+#include <numeric>
+#include <vector>
+
+using namespace sycl;
+using namespace sycl::ext::oneapi;
+
+template <typename T, memory_order order = memory_order::relaxed,
+          memory_scope scope = memory_scope::device>
+void and_local_test(queue q) {
+  const size_t N = 32;
+  T cum = 0;
+  std::vector<T> output(N);
+  std::fill(output.begin(), output.end(), T(123456));
+  {
+    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 out =
+           output_buf.template get_access<access::mode::discard_write>(cgh);
+       accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
+                                                                           cgh);
+
+       cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
+         int gid = it.get_global_id(0);
+         if (gid == 0)
+           loc[0] = T((1ll << N) - 1);
+         it.barrier(access::fence_space::local_space);
+         auto atm = atomic_ref < T,
+              (order == memory_order::acquire || order == memory_order::release)
+                  ? memory_order::relaxed
+                  : order,
+              scope, access::address_space::local_space > (loc[0]);
+         out[gid] = atm.fetch_and(~T(1ll << gid), order);
+         it.barrier(access::fence_space::local_space);
+         if (gid == 0)
+           cum[0] = loc[0];
+       });
+     }).wait_and_throw();
+  }
+
+  // Final value should be equal to 0
+  assert(cum == 0);
+
+  // All other values should be unique; each work-item sets one bit to 0
+  std::sort(output.begin(), output.end());
+  assert(std::unique(output.begin(), output.end()) == output.end());
+}
+
+template <typename T, memory_order order = memory_order::relaxed,
+          memory_scope scope = memory_scope::device>
+void and_global_test(queue q) {
+  const size_t N = 32;
+  const T initial = T((1ll << N) - 1);
+  T cum = initial;
+  std::vector<T> output(N);
+  std::fill(output.begin(), output.end(), T(0));
+  {
+    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 out =
+          output_buf.template get_access<access::mode::discard_write>(cgh);
+      cgh.parallel_for(range<1>(N), [=](item<1> it) {
+        size_t gid = it.get_id(0);
+        auto atm = atomic_ref < T,
+             (order == memory_order::acquire || order == memory_order::release)
+                 ? memory_order::relaxed
+                 : order,
+             scope, access::address_space::global_space > (cum[0]);
+        out[gid] = atm.fetch_and(~T(1ll << gid), order);
+      });
+    });
+  }
+
+  // Final value should be equal to 0
+  assert(cum == 0);
+
+  // All other values should be unique; each work-item sets one bit to 0
+  std::sort(output.begin(), output.end());
+  assert(std::unique(output.begin(), output.end()) == output.end());
+}
+
+template <typename T, memory_order order = memory_order::relaxed,
+          memory_scope scope = memory_scope::device>
+void and_test(queue q) {
+  and_local_test<T, order, scope>(q);
+  and_global_test<T, order, scope>(q);
+}