[SYCL] Disambiguate atomic_ref references

SYCL/AtomicRef/accessor.cpp

@@ -38,10 +38,10 @@ template <typename T> void accessor_test(queue q, size_t N) {
       cgh.parallel_for(range<1>(N), [=](item<1> it) {
         int gid = it.get_id(0);
         static_assert(
-            std::is_same<
-                decltype(sum[0]),
-                atomic_ref<T, memory_order::relaxed, memory_scope::device,
-                           access::address_space::global_space>>::value,
+            std::is_same<decltype(sum[0]),
+                         ::sycl::ext::oneapi::atomic_ref<
+                             T, memory_order::relaxed, memory_scope::device,
+                             access::address_space::global_space>>::value,
             "atomic_accessor returns incorrect atomic_ref");
         out[gid] = sum[0].fetch_add(T(1));
       });
@@ -78,10 +78,10 @@ void local_accessor_test(queue q, size_t N, size_t L = 8) {
         sum[0].store(0);
         it.barrier();
         static_assert(
-            std::is_same<
-                decltype(sum[0]),
-                atomic_ref<T, memory_order::relaxed, memory_scope::device,
-                           access::address_space::local_space>>::value,
+            std::is_same<decltype(sum[0]),
+                         ::sycl::ext::oneapi::atomic_ref<
+                             T, memory_order::relaxed, memory_scope::device,
+                             access::address_space::local_space>>::value,
             "local atomic_accessor returns incorrect atomic_ref");
         T result = sum[0].fetch_add(T(1));
         if (result == it.get_local_range(0) - 1) {

SYCL/AtomicRef/add.h

@@ -24,8 +24,9 @@ void add_fetch_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 = ::sycl::ext::oneapi::atomic_ref<
+            T, memory_order::relaxed, memory_scope::device,
+            access::address_space::global_space>(sum[0]);
         out[gid] = atm.fetch_add(Difference(1));
       });
     });
@@ -59,8 +60,9 @@ 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 = ::sycl::ext::oneapi::atomic_ref<
+            T, memory_order::relaxed, memory_scope::device,
+            access::address_space::global_space>(sum[0]);
         out[gid] = atm += Difference(1);
       });
     });
@@ -94,8 +96,9 @@ 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 = ::sycl::ext::oneapi::atomic_ref<
+            T, memory_order::relaxed, memory_scope::device,
+            access::address_space::global_space>(sum[0]);
         out[gid] = ++atm;
       });
     });
@@ -129,8 +132,9 @@ 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 = ::sycl::ext::oneapi::atomic_ref<
+            T, memory_order::relaxed, memory_scope::device,
+            access::address_space::global_space>(sum[0]);
         out[gid] = atm++;
       });
     });

SYCL/AtomicRef/assignment.h

@@ -21,9 +21,10 @@ template <typename T> void assignment_test(queue q, size_t N) {
           assignment_buf.template get_access<access::mode::read_write>(cgh);
       cgh.parallel_for<assignment_kernel<T>>(range<1>(N), [=](item<1> it) {
         size_t gid = it.get_id(0);
-        auto atm = atomic_ref<T, ext::oneapi::memory_order::relaxed,
-                              ext::oneapi::memory_scope::device,
-                              access::address_space::global_space>(st[0]);
+        auto atm = ::sycl::ext::oneapi::atomic_ref<
+            T, ext::oneapi::memory_order::relaxed,
+            ext::oneapi::memory_scope::device,
+            access::address_space::global_space>(st[0]);
         atm = T(gid);
       });
     });

SYCL/AtomicRef/atomic_memory_order_acq_rel.h

@@ -19,9 +19,9 @@ template <typename T> void acq_rel_test(queue q, size_t N) {
       auto a_acc = a_buf.template get_access<access::mode::read_write>(cgh);
       cgh.parallel_for<atomic_memory_order_acq_rel_kernel<T>>(
           range<1>(N), [=](item<1> it) {
-            auto aar =
-                atomic_ref<T, memory_order::acq_rel, memory_scope::device,
-                           access::address_space::global_space>(a_acc[0]);
+            auto aar = ::sycl::ext::oneapi::atomic_ref<
+                T, memory_order::acq_rel, memory_scope::device,
+                access::address_space::global_space>(a_acc[0]);
             auto ld = aar.load();
             ld += 1;
             aar.store(ld);

SYCL/AtomicRef/atomic_memory_order_seq_cst.h

@@ -22,12 +22,12 @@ template <typename T> void seq_cst_test(queue q, size_t N) {
       auto b_acc = b_buf.template get_access<access::mode::read_write>(cgh);
       cgh.parallel_for<atomic_memory_order_seq_cst_kernel<T>>(
           range<1>(N), [=](item<1> it) {
-            auto aar =
-                atomic_ref<T, memory_order::seq_cst, memory_scope::device,
-                           access::address_space::global_space>(a_acc[0]);
-            auto bar =
-                atomic_ref<T, memory_order::seq_cst, memory_scope::device,
-                           access::address_space::global_space>(b_acc[0]);
+            auto aar = ::sycl::ext::oneapi::atomic_ref<
+                T, memory_order::seq_cst, memory_scope::device,
+                access::address_space::global_space>(a_acc[0]);
+            auto bar = ::sycl::ext::oneapi::atomic_ref<
+                T, memory_order::seq_cst, memory_scope::device,
+                access::address_space::global_space>(b_acc[0]);
             auto ald = aar.load();
             auto bld = bar.load();
             ald += 1;

SYCL/AtomicRef/compare_exchange.h

@@ -29,9 +29,9 @@ template <typename T> void compare_exchange_test(queue q, size_t N) {
       cgh.parallel_for<compare_exchange_kernel<T>>(
           range<1>(N), [=](item<1> it) {
             size_t gid = it.get_id(0);
-            auto atm =
-                atomic_ref<T, memory_order::relaxed, memory_scope::device,
-                           access::address_space::global_space>(exc[0]);
+            auto atm = ::sycl::ext::oneapi::atomic_ref<
+                T, memory_order::relaxed, memory_scope::device,
+                access::address_space::global_space>(exc[0]);
             T result = T(N); // Avoid copying pointer
             bool success = atm.compare_exchange_strong(result, (T)gid);
             if (success) {

SYCL/AtomicRef/exchange.h

@@ -27,8 +27,9 @@ template <typename T> void exchange_test(queue q, size_t N) {
           output_buf.template get_access<access::mode::discard_write>(cgh);
       cgh.parallel_for<exchange_kernel<T>>(range<1>(N), [=](item<1> it) {
         size_t gid = it.get_id(0);
-        auto atm = atomic_ref<T, memory_order::relaxed, memory_scope::device,
-                              access::address_space::global_space>(exc[0]);
+        auto atm = ::sycl::ext::oneapi::atomic_ref<
+            T, memory_order::relaxed, memory_scope::device,
+            access::address_space::global_space>(exc[0]);
         out[gid] = atm.exchange(T(gid));
       });
     });

SYCL/AtomicRef/load.h

@@ -26,8 +26,9 @@ template <typename T> void load_test(queue q, size_t N) {
           output_buf.template get_access<access::mode::discard_write>(cgh);
       cgh.parallel_for<load_kernel<T>>(range<1>(N), [=](item<1> it) {
         size_t gid = it.get_id(0);
-        auto atm = atomic_ref<T, memory_order::relaxed, memory_scope::device,
-                              access::address_space::global_space>(ld[0]);
+        auto atm = ::sycl::ext::oneapi::atomic_ref<
+            T, memory_order::relaxed, memory_scope::device,
+            access::address_space::global_space>(ld[0]);
         out[gid] = atm.load();
       });
     });

SYCL/AtomicRef/max.h

@@ -24,9 +24,9 @@ template <typename T> void max_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>(val[0]);
-
+        auto atm = ::sycl::ext::oneapi::atomic_ref<
+            T, memory_order::relaxed, memory_scope::device,
+            access::address_space::global_space>(val[0]);
         // +1 accounts for lowest() returning 0 for unsigned types
         out[gid] = atm.fetch_max(T(gid) + 1);
       });

SYCL/AtomicRef/min.h

@@ -24,8 +24,9 @@ template <typename T> void min_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>(val[0]);
+        auto atm = ::sycl::ext::oneapi::atomic_ref<
+            T, memory_order::relaxed, memory_scope::device,
+            access::address_space::global_space>(val[0]);
         out[gid] = atm.fetch_min(T(gid));
       });
     });

SYCL/AtomicRef/store.h

@@ -20,8 +20,9 @@ template <typename T> void store_test(queue q, size_t N) {
       auto st = store_buf.template get_access<access::mode::read_write>(cgh);
       cgh.parallel_for<store_kernel<T>>(range<1>(N), [=](item<1> it) {
         size_t gid = it.get_id(0);
-        auto atm = atomic_ref<T, memory_order::relaxed, memory_scope::device,
-                              access::address_space::global_space>(st[0]);
+        auto atm = ::sycl::ext::oneapi::atomic_ref<
+            T, memory_order::relaxed, memory_scope::device,
+            access::address_space::global_space>(st[0]);
         atm.store(T(gid));
       });
     });

SYCL/AtomicRef/sub.h

@@ -24,8 +24,9 @@ void sub_fetch_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>(val[0]);
+        auto atm = ::sycl::ext::oneapi::atomic_ref<
+            T, memory_order::relaxed, memory_scope::device,
+            access::address_space::global_space>(val[0]);
         out[gid] = atm.fetch_sub(Difference(1));
       });
     });
@@ -59,8 +60,9 @@ void sub_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>(val[0]);
+        auto atm = ::sycl::ext::oneapi::atomic_ref<
+            T, memory_order::relaxed, memory_scope::device,
+            access::address_space::global_space>(val[0]);
         out[gid] = atm -= Difference(1);
       });
     });
@@ -94,8 +96,9 @@ void sub_pre_dec_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>(val[0]);
+        auto atm = ::sycl::ext::oneapi::atomic_ref<
+            T, memory_order::relaxed, memory_scope::device,
+            access::address_space::global_space>(val[0]);
         out[gid] = --atm;
       });
     });
@@ -129,8 +132,9 @@ void sub_post_dec_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>(val[0]);
+        auto atm = ::sycl::ext::oneapi::atomic_ref<
+            T, memory_order::relaxed, memory_scope::device,
+            access::address_space::global_space>(val[0]);
         out[gid] = atm--;
       });
     });