[SYCL] Fix sycl::atomic regression (#19015)

This is a follow-up for #18839 which broke host compilation of
`sycl::atomic`.

The actual fix is just an addition of `#include
<sycl/__spirv/spirv_ops.hpp>` to `sycl/atomic.hpp`.

However, I didn't want `spirv_ops.hpp` to appear again in
`sycl/detail/core.hpp` and therefore I went further to make sure that
`sycl/atomic.hpp` isn't used by `sycl/accessor.hpp` which prompted other
changes in this PR.

Author: AlexeySachkov

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/DeviceLib/exp/exp-std-complex-edge-cases.hpp

@@ -1,9 +1,12 @@
 // This test checks edge cases handling for std::exp(std::complex<T>) used
 // in SYCL kernels.
 
+// This include should happen before <sycl/detail/core.hpp> or otherwise NAN
+// may not be constexpr on some Windows configurations. See intel/llvm#19114
+#include <cmath>
+
 #include <sycl/detail/core.hpp>
 
-#include <cmath>
 #include <complex>
 #include <type_traits>
 

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;