[SYCL] Fix sycl::atomic regression

sycl/include/sycl/accessor.hpp

@@ -9,7 +9,6 @@
 #pragma once
 
 #include <sycl/access/access.hpp>                     // for target, mode
-#include <sycl/atomic.hpp>                            // for atomic
 #include <sycl/buffer.hpp>                            // for range
 #include <sycl/detail/accessor_iterator.hpp>          // for accessor_iterator
 #include <sycl/detail/common.hpp>                     // for code_location
@@ -215,6 +214,8 @@
 namespace sycl {
 inline namespace _V1 {
 class stream;
+template <typename T, access::address_space addressSpace> class atomic;
+
 namespace ext::intel::esimd::detail {
 // Forward declare a "back-door" access class to support ESIMD.
 class AccessorPrivateProxy;
@@ -1755,14 +1756,14 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
   }
 
   template <int Dims = Dimensions>
-  operator typename std::enable_if_t<Dims == 0 &&
-                                         AccessMode == access::mode::atomic,
+  operator typename std::enable_if_t<
+      Dims == 0 && AccessMode == access::mode::atomic,
 #ifdef __ENABLE_USM_ADDR_SPACE__
-                                     atomic<DataT>
+      atomic<DataT, access::address_space::global_space>
 #else
-                                     atomic<DataT, AS>
+      atomic<DataT, AS>
 #endif
-                                     >() const {
+      >() const {
     const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
     return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
         getQualifiedPtr() + LinearIndex));

sycl/include/sycl/atomic.hpp

@@ -8,6 +8,7 @@
 
 #pragma once
 
+#include <sycl/__spirv/spirv_ops.hpp>
 #include <sycl/__spirv/spirv_types.hpp>       // for Scope, MemorySemanticsMask
 #include <sycl/access/access.hpp>             // for address_space, decorated
 #include <sycl/detail/defines_elementary.hpp> // for __SYCL2020_DEPRECATED

sycl/test-e2e/Basic/accessor/accessor.cpp

@@ -15,6 +15,8 @@
 //===----------------------------------------------------------------------===//
 #include <cassert>
 #include <iostream>
+
+#include <sycl/atomic.hpp>
 #include <sycl/detail/core.hpp>
 
 struct IdxID1 {

sycl/test-e2e/Basic/compare_exchange_strong.cpp

@@ -1,6 +1,7 @@
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 
+#include <sycl/atomic.hpp>
 #include <sycl/detail/core.hpp>
 using namespace sycl;
 

sycl/test-e2e/Basic/parallel_for_disable_range_roundup.cpp

@@ -6,8 +6,10 @@
 // RUN: %{build} -sycl-std=2020 -o %t2.out
 // RUN: env SYCL_PARALLEL_FOR_RANGE_ROUNDING_TRACE=1 %{run} %t2.out | FileCheck %s --check-prefix CHECK-ENABLED
 
-#include <iostream>
+#include <sycl/atomic.hpp>
 #include <sycl/detail/core.hpp>
+
+#include <iostream>
 using namespace sycl;
 
 range<1> Range1 = {0};

sycl/test-e2e/Basic/parallel_for_range_roundup.cpp

@@ -25,10 +25,12 @@
 // UNSUPPORTED: hip
 // UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/17077
 //
-#include <iostream>
+#include <sycl/atomic.hpp>
 #include <sycl/detail/core.hpp>
 #include <sycl/vector.hpp>
 
+#include <iostream>
+
 using namespace sycl;
 
 constexpr size_t MagicY = 33, MagicZ = 64;

sycl/test-e2e/Regression/atomic_load.cpp

@@ -1,6 +1,8 @@
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
+#include <sycl/atomic.hpp>
 #include <sycl/detail/core.hpp>
+
 using namespace sycl;
 
 template <typename T> class foo;

sycl/test-e2e/Regression/implicit_atomic_conversion.cpp

@@ -1,6 +1,7 @@
 // RUN: %{build} -Wno-error=deprecated-declarations -o %t.out
 // RUN: %{run} %t.out
 
+#include <sycl/atomic.hpp>
 #include <sycl/detail/core.hpp>
 
 using namespace sycl;

sycl/test-e2e/XPTI/buffer/accessors.cpp

@@ -9,6 +9,7 @@
 
 #else
 
+#include <sycl/atomic.hpp>
 #include <sycl/detail/core.hpp>
 
 using namespace sycl::access;

sycl/test/include_deps/sycl_accessor.hpp.cpp

@@ -7,32 +7,21 @@
 // CHECK-NEXT: accessor.hpp
 // CHECK-NEXT: access/access.hpp
 // CHECK-NEXT: detail/defines_elementary.hpp
-// CHECK-NEXT: atomic.hpp
-// CHECK-NEXT: __spirv/spirv_types.hpp
-// CHECK-NEXT: detail/defines.hpp
-// CHECK-NEXT: detail/helpers.hpp
-// CHECK-NEXT: detail/export.hpp
-// CHECK-NEXT: memory_enums.hpp
-// CHECK-NEXT: __spirv/spirv_vars.hpp
-// CHECK-NEXT: multi_ptr.hpp
-// CHECK-NEXT: aliases.hpp
-// CHECK-NEXT: detail/address_space_cast.hpp
-// CHECK-NEXT: detail/type_traits.hpp
-// CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
-// CHECK-NEXT: half_type.hpp
-// CHECK-NEXT: bit_cast.hpp
-// CHECK-NEXT: detail/iostream_proxy.hpp
-// CHECK-NEXT: aspects.hpp
-// CHECK-NEXT: info/aspects.def
-// CHECK-NEXT: info/aspects_deprecated.def
 // CHECK-NEXT: buffer.hpp
 // CHECK-NEXT: backend_types.hpp
 // CHECK-NEXT: detail/array.hpp
 // CHECK-NEXT: exception.hpp
+// CHECK-NEXT: detail/export.hpp
 // CHECK-NEXT: detail/string.hpp
 // CHECK-NEXT: detail/common.hpp
 // CHECK-NEXT: stl_wrappers/cassert
 // CHECK-NEXT: stl_wrappers/assert.h
+// CHECK-NEXT: __spirv/spirv_vars.hpp
+// CHECK-NEXT: __spirv/spirv_types.hpp
+// CHECK-NEXT: detail/defines.hpp
+// CHECK-NEXT: detail/helpers.hpp
+// CHECK-NEXT: memory_enums.hpp
+// CHECK-NEXT: detail/iostream_proxy.hpp
 // CHECK-NEXT: detail/is_device_copyable.hpp
 // CHECK-NEXT: detail/owner_less_base.hpp
 // CHECK-NEXT: detail/impl_utils.hpp
@@ -51,6 +40,16 @@
 // CHECK-NEXT: ur_api.h
 // CHECK-NEXT: detail/accessor_iterator.hpp
 // CHECK-NEXT: detail/generic_type_traits.hpp
+// CHECK-NEXT: aliases.hpp
+// CHECK-NEXT: detail/type_traits.hpp
+// CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
+// CHECK-NEXT: half_type.hpp
+// CHECK-NEXT: bit_cast.hpp
+// CHECK-NEXT: aspects.hpp
+// CHECK-NEXT: info/aspects.def
+// CHECK-NEXT: info/aspects_deprecated.def
+// CHECK-NEXT: multi_ptr.hpp
+// CHECK-NEXT: detail/address_space_cast.hpp
 // CHECK-NEXT: ext/oneapi/bfloat16.hpp
 // CHECK-NEXT: detail/handler_proxy.hpp
 // CHECK-NEXT: pointers.hpp

sycl/test/include_deps/sycl_detail_core.hpp.cpp

@@ -8,32 +8,21 @@
 // CHECK-NEXT: accessor.hpp
 // CHECK-NEXT: access/access.hpp
 // CHECK-NEXT: detail/defines_elementary.hpp
-// CHECK-NEXT: atomic.hpp
-// CHECK-NEXT: __spirv/spirv_types.hpp
-// CHECK-NEXT: detail/defines.hpp
-// CHECK-NEXT: detail/helpers.hpp
-// CHECK-NEXT: detail/export.hpp
-// CHECK-NEXT: memory_enums.hpp
-// CHECK-NEXT: __spirv/spirv_vars.hpp
-// CHECK-NEXT: multi_ptr.hpp
-// CHECK-NEXT: aliases.hpp
-// CHECK-NEXT: detail/address_space_cast.hpp
-// CHECK-NEXT: detail/type_traits.hpp
-// CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
-// CHECK-NEXT: half_type.hpp
-// CHECK-NEXT: bit_cast.hpp
-// CHECK-NEXT: detail/iostream_proxy.hpp
-// CHECK-NEXT: aspects.hpp
-// CHECK-NEXT: info/aspects.def
-// CHECK-NEXT: info/aspects_deprecated.def
 // CHECK-NEXT: buffer.hpp
 // CHECK-NEXT: backend_types.hpp
 // CHECK-NEXT: detail/array.hpp
 // CHECK-NEXT: exception.hpp
+// CHECK-NEXT: detail/export.hpp
 // CHECK-NEXT: detail/string.hpp
 // CHECK-NEXT: detail/common.hpp
 // CHECK-NEXT: stl_wrappers/cassert
 // CHECK-NEXT: stl_wrappers/assert.h
+// CHECK-NEXT: __spirv/spirv_vars.hpp
+// CHECK-NEXT: __spirv/spirv_types.hpp
+// CHECK-NEXT: detail/defines.hpp
+// CHECK-NEXT: detail/helpers.hpp
+// CHECK-NEXT: memory_enums.hpp
+// CHECK-NEXT: detail/iostream_proxy.hpp
 // CHECK-NEXT: detail/is_device_copyable.hpp
 // CHECK-NEXT: detail/owner_less_base.hpp
 // CHECK-NEXT: detail/impl_utils.hpp
@@ -52,6 +41,16 @@
 // CHECK-NEXT: ur_api.h
 // CHECK-NEXT: detail/accessor_iterator.hpp
 // CHECK-NEXT: detail/generic_type_traits.hpp
+// CHECK-NEXT: aliases.hpp
+// CHECK-NEXT: detail/type_traits.hpp
+// CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
+// CHECK-NEXT: half_type.hpp
+// CHECK-NEXT: bit_cast.hpp
+// CHECK-NEXT: aspects.hpp
+// CHECK-NEXT: info/aspects.def
+// CHECK-NEXT: info/aspects_deprecated.def
+// CHECK-NEXT: multi_ptr.hpp
+// CHECK-NEXT: detail/address_space_cast.hpp
 // CHECK-NEXT: ext/oneapi/bfloat16.hpp
 // CHECK-NEXT: detail/handler_proxy.hpp
 // CHECK-NEXT: pointers.hpp

sycl/test/regression/atomic.cpp

@@ -0,0 +1,74 @@
+// RUN: %clangxx -fsycl -fsyntax-only -Wno-deprecated-declarations -Xclang -verify %s
+// expected-no-diagnostics
+
+#include <sycl/atomic.hpp>
+#include <sycl/multi_ptr.hpp>
+
+SYCL_EXTERNAL void
+store(sycl::multi_ptr<int, sycl::access::address_space::global_space> mptr,
+      int value) {
+  sycl::atomic<int> a(mptr);
+  a.store(value);
+}
+
+SYCL_EXTERNAL int
+load(sycl::multi_ptr<int, sycl::access::address_space::global_space> mptr) {
+  sycl::atomic<int> a(mptr);
+  return a.load();
+}
+
+SYCL_EXTERNAL int
+exchange(sycl::multi_ptr<int, sycl::access::address_space::global_space> mptr,
+         int value) {
+  sycl::atomic<int> a(mptr);
+  return a.exchange(value);
+}
+
+SYCL_EXTERNAL int
+fetch_add(sycl::multi_ptr<int, sycl::access::address_space::global_space> mptr,
+          int value) {
+  sycl::atomic<int> a(mptr);
+  return a.fetch_add(value);
+}
+
+SYCL_EXTERNAL int
+fetch_sub(sycl::multi_ptr<int, sycl::access::address_space::global_space> mptr,
+          int value) {
+  sycl::atomic<int> a(mptr);
+  return a.fetch_sub(value);
+}
+
+SYCL_EXTERNAL int
+fetch_and(sycl::multi_ptr<int, sycl::access::address_space::global_space> mptr,
+          int value) {
+  sycl::atomic<int> a(mptr);
+  return a.fetch_and(value);
+}
+
+SYCL_EXTERNAL int
+fetch_or(sycl::multi_ptr<int, sycl::access::address_space::global_space> mptr,
+         int value) {
+  sycl::atomic<int> a(mptr);
+  return a.fetch_or(value);
+}
+
+SYCL_EXTERNAL int
+fetch_xor(sycl::multi_ptr<int, sycl::access::address_space::global_space> mptr,
+          int value) {
+  sycl::atomic<int> a(mptr);
+  return a.fetch_xor(value);
+}
+
+SYCL_EXTERNAL int
+fetch_min(sycl::multi_ptr<int, sycl::access::address_space::global_space> mptr,
+          int value) {
+  sycl::atomic<int> a(mptr);
+  return a.fetch_min(value);
+}
+
+SYCL_EXTERNAL int
+fetch_max(sycl::multi_ptr<int, sycl::access::address_space::global_space> mptr,
+          int value) {
+  sycl::atomic<int> a(mptr);
+  return a.fetch_max(value);
+}