diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp
index 2917774314e1f..66df8d9f97a05 100644
--- a/sycl/include/sycl/accessor.hpp
+++ b/sycl/include/sycl/accessor.hpp
@@ -8,8 +8,8 @@
 
 #pragma once
 
-#include <sycl/access/access.hpp>                     // for target, mode
-#include <sycl/atomic.hpp>                            // for atomic
+#include <sycl/access/access.hpp> // for target, mode
+#include <sycl/detail/defines.hpp>
 #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 +215,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 +1757,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/include/sycl/detail/assume_int.hpp b/sycl/include/sycl/detail/assume_int.hpp
new file mode 100644
index 0000000000000..f653c1cdcf104
--- /dev/null
+++ b/sycl/include/sycl/detail/assume_int.hpp
@@ -0,0 +1,21 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <sycl/detail/defines_elementary.hpp> // for __SYCL_ID_QUERIES_FIT_...
+
+#if __SYCL_ID_QUERIES_FIT_IN_INT__ && __has_builtin(__builtin_assume)
+#include <climits>
+#define __SYCL_ASSUME_INT(x) __builtin_assume((x) <= INT_MAX)
+#else
+#define __SYCL_ASSUME_INT(x)
+#if __SYCL_ID_QUERIES_FIT_IN_INT__ && !__has_builtin(__builtin_assume)
+#warning "No assumptions will be emitted due to no __builtin_assume available"
+#endif
+#endif
diff --git a/sycl/include/sycl/detail/defines.hpp b/sycl/include/sycl/detail/defines.hpp
index 3502465df101d..99063bfe2623a 100644
--- a/sycl/include/sycl/detail/defines.hpp
+++ b/sycl/include/sycl/detail/defines.hpp
@@ -8,17 +8,7 @@
 
 #pragma once
 
-#include <sycl/detail/defines_elementary.hpp> // for __SYCL_ID_QUERIES_FIT_...
-
-#if __SYCL_ID_QUERIES_FIT_IN_INT__ && __has_builtin(__builtin_assume)
-#include <climits>
-#define __SYCL_ASSUME_INT(x) __builtin_assume((x) <= INT_MAX)
-#else
-#define __SYCL_ASSUME_INT(x)
-#if __SYCL_ID_QUERIES_FIT_IN_INT__ && !__has_builtin(__builtin_assume)
-#warning "No assumptions will be emitted due to no __builtin_assume available"
-#endif
-#endif
+#include <sycl/detail/defines_elementary.hpp>
 
 // FIXME Check for  __SYCL_DEVICE_ONLY__ can be removed if implementation of
 // __has_attribute is fixed to consider LangOpts when generating attributes in
diff --git a/sycl/include/sycl/id.hpp b/sycl/include/sycl/id.hpp
index bc17ff460e0eb..2f502efec73c0 100644
--- a/sycl/include/sycl/id.hpp
+++ b/sycl/include/sycl/id.hpp
@@ -10,7 +10,7 @@
 
 #include <sycl/detail/array.hpp>              // for array
 #include <sycl/detail/common.hpp>             // for InitializedVal
-#include <sycl/detail/defines.hpp>            // for __SYCL_ASSUME_INT
+#include <sycl/detail/assume_int.hpp>            // for __SYCL_ASSUME_INT
 #include <sycl/detail/defines_elementary.hpp> // for __SYCL_DEPRECATED, __SYCL_A...
 #include <sycl/exception.hpp> // for make_error_code, errc, exce...
 #include <sycl/range.hpp>     // for range
diff --git a/sycl/include/sycl/item.hpp b/sycl/include/sycl/item.hpp
index 12aed3dc072f1..0d2005a020300 100644
--- a/sycl/include/sycl/item.hpp
+++ b/sycl/include/sycl/item.hpp
@@ -8,7 +8,7 @@
 
 #pragma once
 
-#include <sycl/detail/defines.hpp>            // for __SYCL_ASSUME_INT
+#include <sycl/detail/assume_int.hpp>            // for __SYCL_ASSUME_INT
 #include <sycl/detail/defines_elementary.hpp> // for __SYCL_ALWAYS_INLINE, __SYC...
 #include <sycl/detail/helpers.hpp>            // for Builder
 #include <sycl/detail/item_base.hpp>          // for id, range, ItemBase
diff --git a/sycl/include/sycl/nd_item.hpp b/sycl/include/sycl/nd_item.hpp
index 76ec89f2bba70..b67085caabb2e 100644
--- a/sycl/include/sycl/nd_item.hpp
+++ b/sycl/include/sycl/nd_item.hpp
@@ -8,6 +8,7 @@
 
 #pragma once
 
+#include <sycl/detail/defines.hpp>            // for __SYCL_ASSUME_INT
 #include <sycl/__spirv/spirv_types.hpp>       // for Scope
 #include <sycl/__spirv/spirv_vars.hpp>        // for initLocalInvocationId
 #include <sycl/access/access.hpp>             // for mode, fence_space
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..65e5f2c227877 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,10 @@
 // This test checks edge cases handling for std::exp(std::complex<T>) used
 // in SYCL kernels.
 
+#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);
+}