diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp
index 2917774314e1f..3e9bfd250b11d 100644
--- a/sycl/include/sycl/accessor.hpp
+++ b/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));
diff --git a/sycl/include/sycl/atomic.hpp b/sycl/include/sycl/atomic.hpp
index 8cae0c047c0a1..c2bedd9989fcd 100644
--- a/sycl/include/sycl/atomic.hpp
+++ b/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
diff --git a/sycl/test-e2e/Basic/accessor/accessor.cpp b/sycl/test-e2e/Basic/accessor/accessor.cpp
index 3ab25cd68405f..617f9a39585d6 100644
--- a/sycl/test-e2e/Basic/accessor/accessor.cpp
+++ b/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 {
diff --git a/sycl/test-e2e/Basic/compare_exchange_strong.cpp b/sycl/test-e2e/Basic/compare_exchange_strong.cpp
index c95b7797f0749..92e0a30f04d98 100644
--- a/sycl/test-e2e/Basic/compare_exchange_strong.cpp
+++ b/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;
 
diff --git a/sycl/test-e2e/Basic/parallel_for_disable_range_roundup.cpp b/sycl/test-e2e/Basic/parallel_for_disable_range_roundup.cpp
index 509172de63f03..676188b39e9e0 100644
--- a/sycl/test-e2e/Basic/parallel_for_disable_range_roundup.cpp
+++ b/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};
diff --git a/sycl/test-e2e/Basic/parallel_for_range_roundup.cpp b/sycl/test-e2e/Basic/parallel_for_range_roundup.cpp
index 0866e2dfe32f8..67123d393269d 100644
--- a/sycl/test-e2e/Basic/parallel_for_range_roundup.cpp
+++ b/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;
diff --git a/sycl/test-e2e/DeviceLib/exp/exp-std-complex-edge-cases.hpp b/sycl/test-e2e/DeviceLib/exp/exp-std-complex-edge-cases.hpp
index c0c4d4cc85e91..66a429f23dc4b 100644
--- a/sycl/test-e2e/DeviceLib/exp/exp-std-complex-edge-cases.hpp
+++ b/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>
 
diff --git a/sycl/test-e2e/Regression/atomic_load.cpp b/sycl/test-e2e/Regression/atomic_load.cpp
index 8771db61db935..3f5c31b9ad0f6 100644
--- a/sycl/test-e2e/Regression/atomic_load.cpp
+++ b/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;
diff --git a/sycl/test-e2e/Regression/implicit_atomic_conversion.cpp b/sycl/test-e2e/Regression/implicit_atomic_conversion.cpp
index 8adfc00d37048..00ea968a5b53f 100644
--- a/sycl/test-e2e/Regression/implicit_atomic_conversion.cpp
+++ b/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;
diff --git a/sycl/test-e2e/XPTI/buffer/accessors.cpp b/sycl/test-e2e/XPTI/buffer/accessors.cpp
index ed04243675a73..2bac9a54b2403 100644
--- a/sycl/test-e2e/XPTI/buffer/accessors.cpp
+++ b/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;
diff --git a/sycl/test/include_deps/sycl_accessor.hpp.cpp b/sycl/test/include_deps/sycl_accessor.hpp.cpp
index 7749d8dd4a79f..fcd8fe3e5e904 100644
--- a/sycl/test/include_deps/sycl_accessor.hpp.cpp
+++ b/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
diff --git a/sycl/test/include_deps/sycl_detail_core.hpp.cpp b/sycl/test/include_deps/sycl_detail_core.hpp.cpp
index d33c354e88af9..113ded48b7526 100644
--- a/sycl/test/include_deps/sycl_detail_core.hpp.cpp
+++ b/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
diff --git a/sycl/test/regression/atomic.cpp b/sycl/test/regression/atomic.cpp
new file mode 100644
index 0000000000000..987fd7a0fbb2a
--- /dev/null
+++ b/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);
+}