diff --git a/sycl/gdb/libsycl.so-gdb.py b/sycl/gdb/libsycl.so-gdb.py
index 64d1962be7b13..5b5d1bbe4d81b 100644
--- a/sycl/gdb/libsycl.so-gdb.py
+++ b/sycl/gdb/libsycl.so-gdb.py
@@ -51,37 +51,35 @@ def value(self, arg):
 class HostAccessor(Accessor):
     """For Host device memory layout"""
 
-    def payload(self):
-        return self.obj["impl"]["_M_ptr"].dereference()
-
     def memory_range(self, dim):
-        return self.payload()["MMemoryRange"]["common_array"][dim]
+        eval_string = "((" + str(self.obj.type) + ")" + str(self.obj) + ")->getMemoryRange()"
+        return gdb.parse_and_eval(eval_string)["common_array"][dim];
 
     def offset(self, dim):
-        return self.payload()["MOffset"]["common_array"][dim]
+        eval_string = "((" + str(self.obj.type) + ")" + str(self.obj) + ")->getOffset()"
+        return gdb.parse_and_eval(eval_string)["common_array"][dim];
 
     def data(self):
-        return self.payload()["MData"]
-
+        eval_string = "((" + str(self.obj.type) + ")" + str(self.obj) + ")->getPtr()"
+        return gdb.parse_and_eval(eval_string);
 
 class HostAccessorLocal(HostAccessor):
     """For Host device memory layout"""
 
+    def memory_range(self, dim):
+        eval_string = "((" + str(self.obj.type) + ")" + str(self.obj) + ")->getSize()"
+        return gdb.parse_and_eval(eval_string)["common_array"][dim];
+
     def index(self, arg):
         if arg.type.code == gdb.TYPE_CODE_INT:
             return int(arg)
         result = 0
         for dim in range(self.depth):
             result = (
-                result * self.payload()["MSize"]["common_array"][dim]
-                + arg["common_array"][dim]
+                result * self.memory_range() + arg["common_array"][dim]
             )
         return result
 
-    def data(self):
-        return self.payload()["MMem"]
-
-
 class DeviceAccessor(Accessor):
     """For CPU/GPU memory layout"""
 
@@ -104,7 +102,11 @@ def __init__(self, class_type, result_type, depth):
         self.depth = depth
 
     def get_arg_types(self):
-        return gdb.lookup_type("sycl::_V1::id<%s>" % self.depth)
+        try:
+            return gdb.lookup_type("sycl::_V1::id<%s>" % self.depth)
+        except:
+            pass
+        return None
 
     def get_result_type(self, *args):
         return self.result_type
diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp
index e6cf64c25a0d8..65b1afaebea4a 100644
--- a/sycl/include/sycl/accessor.hpp
+++ b/sycl/include/sycl/accessor.hpp
@@ -11,7 +11,6 @@
 #include <CL/__spirv/spirv_types.hpp>
 #include <sycl/atomic.hpp>
 #include <sycl/buffer.hpp>
-#include <sycl/detail/accessor_impl.hpp>
 #include <sycl/detail/cl.h>
 #include <sycl/detail/common.hpp>
 #include <sycl/detail/export.hpp>
@@ -246,6 +245,7 @@ void __SYCL_EXPORT constructorNotification(void *BufferObj, void *AccessorObj,
                                            access::target Target,
                                            access::mode Mode,
                                            const code_location &CodeLoc);
+
 template <typename T>
 using IsPropertyListT = typename std::is_base_of<PropertyListBase, T>;
 
@@ -357,6 +357,160 @@ class accessor_common {
   };
 };
 
+#if __cplusplus >= 201703L
+
+template <typename MayBeTag1, typename MayBeTag2>
+constexpr access::mode deduceAccessMode() {
+  // property_list = {} is not properly detected by deduction guide,
+  // when parameter is passed without curly braces: access(buffer, no_init)
+  // thus simplest approach is to check 2 last arguments for being a tag
+  if constexpr (std::is_same<MayBeTag1,
+                             mode_tag_t<access::mode::read>>::value ||
+                std::is_same<MayBeTag2,
+                             mode_tag_t<access::mode::read>>::value) {
+    return access::mode::read;
+  }
+
+  if constexpr (std::is_same<MayBeTag1,
+                             mode_tag_t<access::mode::write>>::value ||
+                std::is_same<MayBeTag2,
+                             mode_tag_t<access::mode::write>>::value) {
+    return access::mode::write;
+  }
+
+  if constexpr (
+      std::is_same<MayBeTag1,
+                   mode_target_tag_t<access::mode::read,
+                                     access::target::constant_buffer>>::value ||
+      std::is_same<MayBeTag2,
+                   mode_target_tag_t<access::mode::read,
+                                     access::target::constant_buffer>>::value) {
+    return access::mode::read;
+  }
+
+  return access::mode::read_write;
+}
+
+template <typename MayBeTag1, typename MayBeTag2>
+constexpr access::target deduceAccessTarget(access::target defaultTarget) {
+  if constexpr (
+      std::is_same<MayBeTag1,
+                   mode_target_tag_t<access::mode::read,
+                                     access::target::constant_buffer>>::value ||
+      std::is_same<MayBeTag2,
+                   mode_target_tag_t<access::mode::read,
+                                     access::target::constant_buffer>>::value) {
+    return access::target::constant_buffer;
+  }
+
+  return defaultTarget;
+}
+
+#endif
+
+template <int Dims> class LocalAccessorBaseDevice {
+public:
+  LocalAccessorBaseDevice(sycl::range<Dims> Size)
+      : AccessRange(Size),
+        MemRange(InitializedVal<Dims, range>::template get<0>()) {}
+  // TODO: Actually we need only one field here, but currently compiler requires
+  // all of them.
+  range<Dims> AccessRange;
+  range<Dims> MemRange;
+  id<Dims> Offset;
+
+  bool operator==(const LocalAccessorBaseDevice &Rhs) const {
+    return (AccessRange == Rhs.AccessRange);
+  }
+};
+
+// The class describes a requirement to access a SYCL memory object such as
+// sycl::buffer and sycl::image. For example, each accessor used in a kernel,
+// except one with access target "local", adds such requirement for the command
+// group.
+
+template <int Dims> class AccessorImplDevice {
+public:
+  AccessorImplDevice() = default;
+  AccessorImplDevice(id<Dims> Offset, range<Dims> AccessRange,
+                     range<Dims> MemoryRange)
+      : Offset(Offset), AccessRange(AccessRange), MemRange(MemoryRange) {}
+
+  id<Dims> Offset;
+  range<Dims> AccessRange;
+  range<Dims> MemRange;
+
+  bool operator==(const AccessorImplDevice &Rhs) const {
+    return (Offset == Rhs.Offset && AccessRange == Rhs.AccessRange &&
+            MemRange == Rhs.MemRange);
+  }
+};
+
+class AccessorImplHost;
+
+void __SYCL_EXPORT addHostAccessorAndWait(AccessorImplHost *Req);
+
+class SYCLMemObjI;
+
+using AccessorImplPtr = std::shared_ptr<AccessorImplHost>;
+
+class __SYCL_EXPORT AccessorBaseHost {
+public:
+  AccessorBaseHost(id<3> Offset, range<3> AccessRange, range<3> MemoryRange,
+                   access::mode AccessMode, void *SYCLMemObject, int Dims,
+                   int ElemSize, int OffsetInBytes = 0,
+                   bool IsSubBuffer = false,
+                   const property_list &PropertyList = {});
+
+public:
+  id<3> &getOffset();
+  range<3> &getAccessRange();
+  range<3> &getMemoryRange();
+  void *getPtr();
+  unsigned int getElemSize() const;
+
+  const id<3> &getOffset() const;
+  const range<3> &getAccessRange() const;
+  const range<3> &getMemoryRange() const;
+  void *getPtr() const;
+
+  const property_list &getPropList() const;
+
+  void *getMemoryObject() const;
+
+  template <class Obj>
+  friend decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject);
+
+  template <typename, int, access::mode, access::target, access::placeholder,
+            typename>
+  friend class accessor;
+
+  AccessorImplPtr impl;
+
+private:
+  friend class sycl::ext::intel::esimd::detail::AccessorPrivateProxy;
+};
+
+class LocalAccessorImplHost;
+using LocalAccessorImplPtr = std::shared_ptr<LocalAccessorImplHost>;
+
+class __SYCL_EXPORT LocalAccessorBaseHost {
+public:
+  LocalAccessorBaseHost(sycl::range<3> Size, int Dims, int ElemSize);
+  sycl::range<3> &getSize();
+  const sycl::range<3> &getSize() const;
+  void *getPtr();
+  void *getPtr() const;
+  int getNumOfDims();
+  int getElementSize();
+
+protected:
+  template <class Obj>
+  friend decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject);
+
+  LocalAccessorImplPtr impl;
+};
+
 template <int Dim, typename T> struct IsValidCoordDataT;
 template <typename T> struct IsValidCoordDataT<1, T> {
   constexpr static bool value =
@@ -950,9 +1104,30 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
              detail::InitializedVal<AdjustedDim, range>::template get<0>()) {}
 
 #else
-  using AccessorBaseHost::getAccessRange;
-  using AccessorBaseHost::getMemoryRange;
-  using AccessorBaseHost::getOffset;
+  id<3> &getOffset() { return AccessorBaseHost::getOffset(); }
+  range<3> &getAccessRange() { return AccessorBaseHost::getAccessRange(); }
+  range<3> &getMemoryRange() { return AccessorBaseHost::getMemoryRange(); }
+  void *getPtr() { return AccessorBaseHost::getPtr(); }
+
+  const id<3> &getOffset() const { return AccessorBaseHost::getOffset(); }
+  const range<3> &getAccessRange() const {
+    return AccessorBaseHost::getAccessRange();
+  }
+  const range<3> &getMemoryRange() const {
+    return AccessorBaseHost::getMemoryRange();
+  }
+
+  void *getPtr() const { return AccessorBaseHost::getPtr(); }
+
+  // The function references helper methods required by GDB pretty-printers
+  void GDBMethodsAnchor() {
+#ifndef NDEBUG
+    (void)getMemoryRange();
+    (void)getOffset();
+    (void)getPtr();
+    (void)getAccessRange();
+#endif
+  }
 
   char padding[sizeof(detail::AccessorImplDevice<AdjustedDim>) +
                sizeof(PtrType) - sizeof(detail::AccessorBaseHost)];
@@ -1138,6 +1313,7 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
             getAdjustedMode(PropertyList),
             detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
             BufferRef.OffsetInBytes, BufferRef.IsSubBuffer, PropertyList) {
+    GDBMethodsAnchor();
     preScreenAccessor(BufferRef.size(), PropertyList);
     if (!IsPlaceH)
       addHostAccessorAndWait(AccessorBaseHost::impl.get());
@@ -1171,6 +1347,7 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
             getAdjustedMode(PropertyList),
             detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
             BufferRef.OffsetInBytes, BufferRef.IsSubBuffer, PropertyList) {
+    GDBMethodsAnchor();
     preScreenAccessor(BufferRef.size(), PropertyList);
     if (!IsPlaceH)
       addHostAccessorAndWait(AccessorBaseHost::impl.get());
@@ -1234,6 +1411,7 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
             getAdjustedMode(PropertyList),
             detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
             BufferRef.OffsetInBytes, BufferRef.IsSubBuffer, PropertyList) {
+    GDBMethodsAnchor();
     preScreenAccessor(BufferRef.size(), PropertyList);
     detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
     detail::constructorNotification(detail::getSyclObjImpl(BufferRef).get(),
@@ -1266,6 +1444,7 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
             getAdjustedMode(PropertyList),
             detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
             BufferRef.OffsetInBytes, BufferRef.IsSubBuffer, PropertyList) {
+    GDBMethodsAnchor();
     preScreenAccessor(BufferRef.size(), PropertyList);
     detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
     detail::constructorNotification(detail::getSyclObjImpl(BufferRef).get(),
@@ -1452,6 +1631,7 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
                          detail::getSyclObjImpl(BufferRef).get(), Dimensions,
                          sizeof(DataT), BufferRef.OffsetInBytes,
                          BufferRef.IsSubBuffer, PropertyList) {
+    GDBMethodsAnchor();
     preScreenAccessor(BufferRef.size(), PropertyList);
     if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
                                 BufferRef.get_range()))
@@ -1493,6 +1673,7 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
                          detail::getSyclObjImpl(BufferRef).get(), Dimensions,
                          sizeof(DataT), BufferRef.OffsetInBytes,
                          BufferRef.IsSubBuffer, PropertyList) {
+    GDBMethodsAnchor();
     preScreenAccessor(BufferRef.size(), PropertyList);
     if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
                                 BufferRef.get_range()))
@@ -1565,6 +1746,7 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
                          detail::getSyclObjImpl(BufferRef).get(), Dimensions,
                          sizeof(DataT), BufferRef.OffsetInBytes,
                          BufferRef.IsSubBuffer, PropertyList) {
+    GDBMethodsAnchor();
     preScreenAccessor(BufferRef.size(), PropertyList);
     if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
                                 BufferRef.get_range()))
@@ -1605,6 +1787,7 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
                          detail::getSyclObjImpl(BufferRef).get(), Dimensions,
                          sizeof(DataT), BufferRef.OffsetInBytes,
                          BufferRef.IsSubBuffer, PropertyList) {
+    GDBMethodsAnchor();
     preScreenAccessor(BufferRef.size(), PropertyList);
     if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
                                 BufferRef.get_range()))
@@ -1673,8 +1856,8 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
         PropertyListT::template areSameCompileTimeProperties<NewPropsT...>(),
         "Compile-time-constant properties must be the same");
 #ifndef __SYCL_DEVICE_ONLY__
-    detail::constructorNotification(impl.get()->MSYCLMemObj, impl.get(),
-                                    AccessTarget, AccessMode, CodeLoc);
+    detail::constructorNotification(getMemoryObject(), impl.get(), AccessTarget,
+                                    AccessMode, CodeLoc);
 #endif
   }
 
@@ -1796,7 +1979,7 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
       !ext::oneapi::is_compile_time_property<Property>::value, bool>
   has_property() const noexcept {
 #ifndef __SYCL_DEVICE_ONLY__
-    return AccessorBaseHost::impl->MPropertyList.has_property<Property>();
+    return getPropList().template has_property<Property>();
 #else
     return false;
 #endif
@@ -1810,7 +1993,7 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
                 !ext::oneapi::is_compile_time_property<Property>::value>>
   Property get_property() const {
 #ifndef __SYCL_DEVICE_ONLY__
-    return AccessorBaseHost::impl->MPropertyList.get_property<Property>();
+    return getPropList().template get_property<Property>();
 #else
     return Property();
 #endif
@@ -2130,6 +2313,21 @@ class __SYCL_SPECIAL_CLASS local_accessor_base :
     return reinterpret_cast<PtrType>(LocalAccessorBaseHost::getPtr());
   }
 
+  void *getPtr() { return detail::LocalAccessorBaseHost::getPtr(); }
+  void *getPtr() const { return detail::LocalAccessorBaseHost::getPtr(); }
+  const range<3> &getSize() const {
+    return detail::LocalAccessorBaseHost::getSize();
+  }
+  range<3> &getSize() { return detail::LocalAccessorBaseHost::getSize(); }
+
+  // The function references helper methods required by GDB pretty-printers
+  void GDBMethodsAnchor() {
+#ifndef NDEBUG
+    (void)getSize();
+    (void)getPtr();
+#endif
+  }
+
 #endif // __SYCL_DEVICE_ONLY__
 
   // Method which calculates linear offset for the ID using Range and Offset.
@@ -2154,6 +2352,7 @@ class __SYCL_SPECIAL_CLASS local_accessor_base :
       : LocalAccessorBaseHost(range<3>{1, 1, 1}, AdjustedDim, sizeof(DataT)) {
     detail::constructorNotification(nullptr, LocalAccessorBaseHost::impl.get(),
                                     access::target::local, AccessMode, CodeLoc);
+    GDBMethodsAnchor();
   }
 #endif
 
@@ -2171,6 +2370,7 @@ class __SYCL_SPECIAL_CLASS local_accessor_base :
     (void)propList;
     detail::constructorNotification(nullptr, LocalAccessorBaseHost::impl.get(),
                                     access::target::local, AccessMode, CodeLoc);
+    GDBMethodsAnchor();
   }
 #endif
 
@@ -2185,6 +2385,7 @@ class __SYCL_SPECIAL_CLASS local_accessor_base :
                               AdjustedDim, sizeof(DataT)) {
     detail::constructorNotification(nullptr, LocalAccessorBaseHost::impl.get(),
                                     access::target::local, AccessMode, CodeLoc);
+    GDBMethodsAnchor();
   }
 #endif
 
@@ -2204,6 +2405,7 @@ class __SYCL_SPECIAL_CLASS local_accessor_base :
     (void)propList;
     detail::constructorNotification(nullptr, LocalAccessorBaseHost::impl.get(),
                                     access::target::local, AccessMode, CodeLoc);
+    GDBMethodsAnchor();
   }
 #endif
 
diff --git a/sycl/include/sycl/detail/accessor_impl.hpp b/sycl/include/sycl/detail/accessor_impl.hpp
deleted file mode 100644
index a820b9dc478de..0000000000000
--- a/sycl/include/sycl/detail/accessor_impl.hpp
+++ /dev/null
@@ -1,286 +0,0 @@
-//==------------ accessor_impl.hpp - SYCL standard header file -------------==//
-//
-// 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/access/access.hpp>
-#include <sycl/detail/export.hpp>
-#include <sycl/id.hpp>
-#include <sycl/property_list.hpp>
-#include <sycl/range.hpp>
-#include <sycl/stl.hpp>
-
-namespace sycl {
-__SYCL_INLINE_VER_NAMESPACE(_V1) {
-template <typename DataT, int Dimensions, access::mode AccessMode,
-          access::target AccessTarget, access::placeholder IsPlaceholder,
-          typename PropertyListT>
-class accessor;
-
-namespace ext {
-namespace intel {
-namespace esimd {
-namespace detail {
-// Forward declare a "back-door" access class to support ESIMD.
-class AccessorPrivateProxy;
-} // namespace detail
-} // namespace esimd
-} // namespace intel
-} // namespace ext
-
-namespace detail {
-
-class SYCLMemObjI;
-
-class Command;
-
-// The class describes a requirement to access a SYCL memory object such as
-// sycl::buffer and sycl::image. For example, each accessor used in a kernel,
-// except one with access target "local", adds such requirement for the command
-// group.
-
-template <int Dims> class AccessorImplDevice {
-public:
-  AccessorImplDevice() = default;
-  AccessorImplDevice(id<Dims> Offset, range<Dims> AccessRange,
-                     range<Dims> MemoryRange)
-      : Offset(Offset), AccessRange(AccessRange), MemRange(MemoryRange) {}
-
-  id<Dims> Offset;
-  range<Dims> AccessRange;
-  range<Dims> MemRange;
-
-  bool operator==(const AccessorImplDevice &Rhs) const {
-    return (Offset == Rhs.Offset && AccessRange == Rhs.AccessRange &&
-            MemRange == Rhs.MemRange);
-  }
-};
-
-template <int Dims> class LocalAccessorBaseDevice {
-public:
-  LocalAccessorBaseDevice(sycl::range<Dims> Size)
-      : AccessRange(Size),
-        MemRange(InitializedVal<Dims, range>::template get<0>()) {}
-  // TODO: Actually we need only one field here, but currently compiler requires
-  // all of them.
-  range<Dims> AccessRange;
-  range<Dims> MemRange;
-  id<Dims> Offset;
-
-  bool operator==(const LocalAccessorBaseDevice &Rhs) const {
-    return (AccessRange == Rhs.AccessRange);
-  }
-};
-
-class __SYCL_EXPORT AccessorImplHost {
-public:
-  AccessorImplHost(id<3> Offset, range<3> AccessRange, range<3> MemoryRange,
-                   access::mode AccessMode, void *SYCLMemObject, int Dims,
-                   int ElemSize, int OffsetInBytes = 0,
-                   bool IsSubBuffer = false,
-                   const property_list &PropertyList = {})
-      : MOffset(Offset), MAccessRange(AccessRange), MMemoryRange(MemoryRange),
-        MAccessMode(AccessMode),
-        MSYCLMemObj((detail::SYCLMemObjI *)SYCLMemObject), MDims(Dims),
-        MElemSize(ElemSize), MOffsetInBytes(OffsetInBytes),
-        MIsSubBuffer(IsSubBuffer), MPropertyList(PropertyList) {}
-
-  ~AccessorImplHost();
-
-  AccessorImplHost(const AccessorImplHost &Other)
-      : MOffset(Other.MOffset), MAccessRange(Other.MAccessRange),
-        MMemoryRange(Other.MMemoryRange), MAccessMode(Other.MAccessMode),
-        MSYCLMemObj(Other.MSYCLMemObj), MDims(Other.MDims),
-        MElemSize(Other.MElemSize), MOffsetInBytes(Other.MOffsetInBytes),
-        MIsSubBuffer(Other.MIsSubBuffer) {}
-
-  // The resize method provides a way to change the size of the
-  // allocated memory and corresponding properties for the accessor.
-  // These are normally fixed for the accessor, but this capability
-  // is needed to support the stream class.
-  // Stream implementation creates an accessor with initial size for
-  // work item. But the number of work items is not available during
-  // stream construction. The resize method allows to update the accessor
-  // as the information becomes available to the handler.
-
-  void resize(size_t GlobalSize);
-
-  id<3> MOffset;
-  // The size of accessing region.
-  range<3> MAccessRange;
-  // The size of memory object this requirement is created for.
-  range<3> MMemoryRange;
-  access::mode MAccessMode;
-
-  detail::SYCLMemObjI *MSYCLMemObj;
-
-  unsigned int MDims;
-  unsigned int MElemSize;
-  unsigned int MOffsetInBytes;
-  bool MIsSubBuffer;
-
-  void *MData = nullptr;
-
-  Command *MBlockedCmd = nullptr;
-
-  bool PerWI = false;
-
-  // To preserve runtime properties
-  property_list MPropertyList;
-};
-
-using AccessorImplPtr = std::shared_ptr<AccessorImplHost>;
-
-class AccessorBaseHost {
-public:
-  template <typename PropertyListT = property_list>
-  AccessorBaseHost(id<3> Offset, range<3> AccessRange, range<3> MemoryRange,
-                   access::mode AccessMode, void *SYCLMemObject, int Dims,
-                   int ElemSize, int OffsetInBytes = 0,
-                   bool IsSubBuffer = false,
-                   const PropertyListT &PropertyList = {}) {
-    impl = std::shared_ptr<AccessorImplHost>(new AccessorImplHost(
-        Offset, AccessRange, MemoryRange, AccessMode,
-        (detail::SYCLMemObjI *)SYCLMemObject, Dims, ElemSize, OffsetInBytes,
-        IsSubBuffer, PropertyList));
-  }
-
-protected:
-  id<3> &getOffset() { return impl->MOffset; }
-  range<3> &getAccessRange() { return impl->MAccessRange; }
-  range<3> &getMemoryRange() { return impl->MMemoryRange; }
-  void *getPtr() { return impl->MData; }
-  unsigned int getElemSize() const { return impl->MElemSize; }
-
-  const id<3> &getOffset() const { return impl->MOffset; }
-  const range<3> &getAccessRange() const { return impl->MAccessRange; }
-  const range<3> &getMemoryRange() const { return impl->MMemoryRange; }
-  void *getPtr() const { return const_cast<void *>(impl->MData); }
-
-  template <class Obj>
-  friend decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject);
-
-  template <typename DataT, int Dimensions, access::mode AccessMode,
-            access::target AccessTarget, access::placeholder IsPlaceholder,
-            typename PropertyListT>
-  friend class accessor;
-
-  AccessorImplPtr impl;
-
-private:
-  friend class sycl::ext::intel::esimd::detail::AccessorPrivateProxy;
-};
-
-class __SYCL_EXPORT LocalAccessorImplHost {
-public:
-  // Allocate ElemSize more data to have sufficient padding to enforce
-  // alignment.
-  LocalAccessorImplHost(sycl::range<3> Size, int Dims, int ElemSize)
-      : MSize(Size), MDims(Dims), MElemSize(ElemSize),
-        MMem(Size[0] * Size[1] * Size[2] * ElemSize + ElemSize) {}
-
-  sycl::range<3> MSize;
-  int MDims;
-  int MElemSize;
-  std::vector<char> MMem;
-};
-
-using LocalAccessorImplPtr = std::shared_ptr<LocalAccessorImplHost>;
-
-class LocalAccessorBaseHost {
-public:
-  LocalAccessorBaseHost(sycl::range<3> Size, int Dims, int ElemSize) {
-    impl = std::shared_ptr<LocalAccessorImplHost>(
-        new LocalAccessorImplHost(Size, Dims, ElemSize));
-  }
-  sycl::range<3> &getSize() { return impl->MSize; }
-  const sycl::range<3> &getSize() const { return impl->MSize; }
-  void *getPtr() {
-    // Const cast this in order to call the const getPtr.
-    return const_cast<const LocalAccessorBaseHost *>(this)->getPtr();
-  }
-  void *getPtr() const {
-    char *ptr = impl->MMem.data();
-
-    // Align the pointer to MElemSize.
-    size_t val = reinterpret_cast<size_t>(ptr);
-    if (val % impl->MElemSize != 0) {
-      ptr += impl->MElemSize - val % impl->MElemSize;
-    }
-
-    return ptr;
-  }
-
-  int getNumOfDims() { return impl->MDims; }
-  int getElementSize() { return impl->MElemSize; }
-
-protected:
-  template <class Obj>
-  friend decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject);
-
-  std::shared_ptr<LocalAccessorImplHost> impl;
-};
-
-using Requirement = AccessorImplHost;
-
-void __SYCL_EXPORT addHostAccessorAndWait(Requirement *Req);
-
-#if __cplusplus >= 201703L
-
-template <typename MayBeTag1, typename MayBeTag2>
-constexpr access::mode deduceAccessMode() {
-  // property_list = {} is not properly detected by deduction guide,
-  // when parameter is passed without curly braces: access(buffer, no_init)
-  // thus simplest approach is to check 2 last arguments for being a tag
-  if constexpr (std::is_same<MayBeTag1,
-                             mode_tag_t<access::mode::read>>::value ||
-                std::is_same<MayBeTag2,
-                             mode_tag_t<access::mode::read>>::value) {
-    return access::mode::read;
-  }
-
-  if constexpr (std::is_same<MayBeTag1,
-                             mode_tag_t<access::mode::write>>::value ||
-                std::is_same<MayBeTag2,
-                             mode_tag_t<access::mode::write>>::value) {
-    return access::mode::write;
-  }
-
-  if constexpr (
-      std::is_same<MayBeTag1,
-                   mode_target_tag_t<access::mode::read,
-                                     access::target::constant_buffer>>::value ||
-      std::is_same<MayBeTag2,
-                   mode_target_tag_t<access::mode::read,
-                                     access::target::constant_buffer>>::value) {
-    return access::mode::read;
-  }
-
-  return access::mode::read_write;
-}
-
-template <typename MayBeTag1, typename MayBeTag2>
-constexpr access::target deduceAccessTarget(access::target defaultTarget) {
-  if constexpr (
-      std::is_same<MayBeTag1,
-                   mode_target_tag_t<access::mode::read,
-                                     access::target::constant_buffer>>::value ||
-      std::is_same<MayBeTag2,
-                   mode_target_tag_t<access::mode::read,
-                                     access::target::constant_buffer>>::value) {
-    return access::target::constant_buffer;
-  }
-
-  return defaultTarget;
-}
-
-#endif
-
-} // namespace detail
-} // __SYCL_INLINE_VER_NAMESPACE(_V1)
-} // namespace sycl
diff --git a/sycl/include/sycl/detail/cg.hpp b/sycl/include/sycl/detail/cg.hpp
index 829e6411a5c8d..1dea9222d4ceb 100644
--- a/sycl/include/sycl/detail/cg.hpp
+++ b/sycl/include/sycl/detail/cg.hpp
@@ -10,7 +10,6 @@
 
 #include <sycl/accessor.hpp>
 #include <sycl/backend_types.hpp>
-#include <sycl/detail/accessor_impl.hpp>
 #include <sycl/detail/cg_types.hpp>
 #include <sycl/detail/common.hpp>
 #include <sycl/detail/export.hpp>
@@ -170,7 +169,7 @@ class CG {
   CG(CGTYPE Type, std::vector<std::vector<char>> ArgsStorage,
      std::vector<detail::AccessorImplPtr> AccStorage,
      std::vector<std::shared_ptr<const void>> SharedPtrStorage,
-     std::vector<Requirement *> Requirements,
+     std::vector<AccessorImplHost *> Requirements,
      std::vector<detail::EventImplPtr> Events, detail::code_location loc = {})
       : MType(Type), MArgsStorage(std::move(ArgsStorage)),
         MAccStorage(std::move(AccStorage)),
@@ -221,7 +220,7 @@ class CG {
 public:
   /// List of requirements that specify which memory is needed for the command
   /// group to be executed.
-  std::vector<Requirement *> MRequirements;
+  std::vector<AccessorImplHost *> MRequirements;
   /// List of events that order the execution of this CG
   std::vector<detail::EventImplPtr> MEvents;
   // Member variables to capture the user code-location
@@ -252,7 +251,7 @@ class CGExecKernel : public CG {
                std::vector<std::vector<char>> ArgsStorage,
                std::vector<detail::AccessorImplPtr> AccStorage,
                std::vector<std::shared_ptr<const void>> SharedPtrStorage,
-               std::vector<Requirement *> Requirements,
+               std::vector<AccessorImplHost *> Requirements,
                std::vector<detail::EventImplPtr> Events,
                std::vector<ArgDesc> Args, std::string KernelName,
                detail::OSModuleHandle OSModuleHandle,
@@ -303,7 +302,7 @@ class CGCopy : public CG {
          std::vector<std::vector<char>> ArgsStorage,
          std::vector<detail::AccessorImplPtr> AccStorage,
          std::vector<std::shared_ptr<const void>> SharedPtrStorage,
-         std::vector<Requirement *> Requirements,
+         std::vector<AccessorImplHost *> Requirements,
          std::vector<detail::EventImplPtr> Events,
          detail::code_location loc = {})
       : CG(CopyType, std::move(ArgsStorage), std::move(AccStorage),
@@ -318,39 +317,39 @@ class CGCopy : public CG {
 class CGFill : public CG {
 public:
   std::vector<char> MPattern;
-  Requirement *MPtr;
+  AccessorImplHost *MPtr;
 
   CGFill(std::vector<char> Pattern, void *Ptr,
          std::vector<std::vector<char>> ArgsStorage,
          std::vector<detail::AccessorImplPtr> AccStorage,
          std::vector<std::shared_ptr<const void>> SharedPtrStorage,
-         std::vector<Requirement *> Requirements,
+         std::vector<AccessorImplHost *> Requirements,
          std::vector<detail::EventImplPtr> Events,
          detail::code_location loc = {})
       : CG(Fill, std::move(ArgsStorage), std::move(AccStorage),
            std::move(SharedPtrStorage), std::move(Requirements),
            std::move(Events), std::move(loc)),
-        MPattern(std::move(Pattern)), MPtr((Requirement *)Ptr) {}
-  Requirement *getReqToFill() { return MPtr; }
+        MPattern(std::move(Pattern)), MPtr((AccessorImplHost *)Ptr) {}
+  AccessorImplHost *getReqToFill() { return MPtr; }
 };
 
 /// "Update host" command group class.
 class CGUpdateHost : public CG {
-  Requirement *MPtr;
+  AccessorImplHost *MPtr;
 
 public:
   CGUpdateHost(void *Ptr, std::vector<std::vector<char>> ArgsStorage,
                std::vector<detail::AccessorImplPtr> AccStorage,
                std::vector<std::shared_ptr<const void>> SharedPtrStorage,
-               std::vector<Requirement *> Requirements,
+               std::vector<AccessorImplHost *> Requirements,
                std::vector<detail::EventImplPtr> Events,
                detail::code_location loc = {})
       : CG(UpdateHost, std::move(ArgsStorage), std::move(AccStorage),
            std::move(SharedPtrStorage), std::move(Requirements),
            std::move(Events), std::move(loc)),
-        MPtr((Requirement *)Ptr) {}
+        MPtr((AccessorImplHost *)Ptr) {}
 
-  Requirement *getReqToUpdate() { return MPtr; }
+  AccessorImplHost *getReqToUpdate() { return MPtr; }
 };
 
 /// "Copy USM" command group class.
@@ -364,7 +363,7 @@ class CGCopyUSM : public CG {
             std::vector<std::vector<char>> ArgsStorage,
             std::vector<detail::AccessorImplPtr> AccStorage,
             std::vector<std::shared_ptr<const void>> SharedPtrStorage,
-            std::vector<Requirement *> Requirements,
+            std::vector<AccessorImplHost *> Requirements,
             std::vector<detail::EventImplPtr> Events,
             detail::code_location loc = {})
       : CG(CopyUSM, std::move(ArgsStorage), std::move(AccStorage),
@@ -388,7 +387,7 @@ class CGFillUSM : public CG {
             std::vector<std::vector<char>> ArgsStorage,
             std::vector<detail::AccessorImplPtr> AccStorage,
             std::vector<std::shared_ptr<const void>> SharedPtrStorage,
-            std::vector<Requirement *> Requirements,
+            std::vector<AccessorImplHost *> Requirements,
             std::vector<detail::EventImplPtr> Events,
             detail::code_location loc = {})
       : CG(FillUSM, std::move(ArgsStorage), std::move(AccStorage),
@@ -410,7 +409,7 @@ class CGPrefetchUSM : public CG {
                 std::vector<std::vector<char>> ArgsStorage,
                 std::vector<detail::AccessorImplPtr> AccStorage,
                 std::vector<std::shared_ptr<const void>> SharedPtrStorage,
-                std::vector<Requirement *> Requirements,
+                std::vector<AccessorImplHost *> Requirements,
                 std::vector<detail::EventImplPtr> Events,
                 detail::code_location loc = {})
       : CG(PrefetchUSM, std::move(ArgsStorage), std::move(AccStorage),
@@ -432,7 +431,7 @@ class CGAdviseUSM : public CG {
               std::vector<std::vector<char>> ArgsStorage,
               std::vector<detail::AccessorImplPtr> AccStorage,
               std::vector<std::shared_ptr<const void>> SharedPtrStorage,
-              std::vector<Requirement *> Requirements,
+              std::vector<AccessorImplHost *> Requirements,
               std::vector<detail::EventImplPtr> Events, CGTYPE Type,
               detail::code_location loc = {})
       : CG(Type, std::move(ArgsStorage), std::move(AccStorage),
@@ -452,7 +451,7 @@ class CGInteropTask : public CG {
                 std::vector<std::vector<char>> ArgsStorage,
                 std::vector<detail::AccessorImplPtr> AccStorage,
                 std::vector<std::shared_ptr<const void>> SharedPtrStorage,
-                std::vector<Requirement *> Requirements,
+                std::vector<AccessorImplHost *> Requirements,
                 std::vector<detail::EventImplPtr> Events, CGTYPE Type,
                 detail::code_location loc = {})
       : CG(Type, std::move(ArgsStorage), std::move(AccStorage),
@@ -477,7 +476,7 @@ class CGHostTask : public CG {
              std::vector<std::vector<char>> ArgsStorage,
              std::vector<detail::AccessorImplPtr> AccStorage,
              std::vector<std::shared_ptr<const void>> SharedPtrStorage,
-             std::vector<Requirement *> Requirements,
+             std::vector<AccessorImplHost *> Requirements,
              std::vector<detail::EventImplPtr> Events, CGTYPE Type,
              detail::code_location loc = {})
       : CG(Type, std::move(ArgsStorage), std::move(AccStorage),
@@ -495,7 +494,7 @@ class CGBarrier : public CG {
             std::vector<std::vector<char>> ArgsStorage,
             std::vector<detail::AccessorImplPtr> AccStorage,
             std::vector<std::shared_ptr<const void>> SharedPtrStorage,
-            std::vector<Requirement *> Requirements,
+            std::vector<AccessorImplHost *> Requirements,
             std::vector<detail::EventImplPtr> Events, CGTYPE Type,
             detail::code_location loc = {})
       : CG(Type, std::move(ArgsStorage), std::move(AccStorage),
diff --git a/sycl/include/sycl/ext/intel/esimd/detail/memory_intrin.hpp b/sycl/include/sycl/ext/intel/esimd/detail/memory_intrin.hpp
index b7a49d424997d..b21e848b289fe 100644
--- a/sycl/include/sycl/ext/intel/esimd/detail/memory_intrin.hpp
+++ b/sycl/include/sycl/ext/intel/esimd/detail/memory_intrin.hpp
@@ -13,7 +13,7 @@
 
 #pragma once
 
-#include <sycl/detail/accessor_impl.hpp>
+#include <sycl/accessor.hpp>
 #include <sycl/ext/intel/esimd/common.hpp>
 #include <sycl/ext/intel/esimd/detail/types.hpp>
 #include <sycl/ext/intel/esimd/detail/util.hpp>
diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp
index 52082f56e076f..85f8098e15efa 100644
--- a/sycl/include/sycl/handler.hpp
+++ b/sycl/include/sycl/handler.hpp
@@ -521,7 +521,7 @@ class __SYCL_EXPORT handler {
       accessor<DataT, Dims, AccessMode, AccessTarget, IsPlaceholder> &&Arg) {
     detail::AccessorBaseHost *AccBase = (detail::AccessorBaseHost *)&Arg;
     detail::AccessorImplPtr AccImpl = detail::getSyclObjImpl(*AccBase);
-    detail::Requirement *Req = AccImpl.get();
+    detail::AccessorImplHost *Req = AccImpl.get();
     // Add accessor to the list of requirements.
     MRequirements.push_back(Req);
     // Store copy of the accessor.
@@ -2546,7 +2546,7 @@ class __SYCL_EXPORT handler {
   /// have become required for this handler via require method.
   std::vector<detail::ArgDesc> MAssociatedAccesors;
   /// The list of requirements to the memory objects for the scheduling.
-  std::vector<detail::Requirement *> MRequirements;
+  std::vector<detail::AccessorImplHost *> MRequirements;
   /// Struct that encodes global size, local size, ...
   detail::NDRDescT MNDRDesc;
   std::string MKernelName;
diff --git a/sycl/include/sycl/interop_handle.hpp b/sycl/include/sycl/interop_handle.hpp
index fdc056c97454f..8804073f827a9 100644
--- a/sycl/include/sycl/interop_handle.hpp
+++ b/sycl/include/sycl/interop_handle.hpp
@@ -11,7 +11,6 @@
 #include <sycl/access/access.hpp>
 #include <sycl/accessor.hpp>
 #include <sycl/backend_types.hpp>
-#include <sycl/detail/accessor_impl.hpp>
 #include <sycl/detail/common.hpp>
 #include <sycl/detail/defines.hpp>
 #include <sycl/detail/pi.hpp>
@@ -146,7 +145,7 @@ class interop_handle {
 private:
   friend class detail::ExecCGCommand;
   friend class detail::DispatchHostTask;
-  using ReqToMem = std::pair<detail::Requirement *, pi_mem>;
+  using ReqToMem = std::pair<detail::AccessorImplHost *, pi_mem>;
 
   interop_handle(std::vector<ReqToMem> MemObjs,
                  const std::shared_ptr<detail::queue_impl> &Queue,
@@ -157,13 +156,14 @@ class interop_handle {
 
   template <backend Backend, typename DataT, int Dims>
   backend_return_t<Backend, buffer<DataT, Dims>>
-  getMemImpl(detail::Requirement *Req) const {
+  getMemImpl(detail::AccessorImplHost *Req) const {
     std::vector<pi_native_handle> NativeHandles{getNativeMem(Req)};
     return detail::BufferInterop<Backend, DataT, Dims>::GetNativeObjs(
         NativeHandles);
   }
 
-  __SYCL_EXPORT pi_native_handle getNativeMem(detail::Requirement *Req) const;
+  __SYCL_EXPORT pi_native_handle
+  getNativeMem(detail::AccessorImplHost *Req) const;
   __SYCL_EXPORT pi_native_handle getNativeQueue() const;
   __SYCL_EXPORT pi_native_handle getNativeDevice() const;
   __SYCL_EXPORT pi_native_handle getNativeContext() const;
diff --git a/sycl/include/sycl/interop_handler.hpp b/sycl/include/sycl/interop_handler.hpp
index dafa9b84ccbf9..ec185dd715f22 100644
--- a/sycl/include/sycl/interop_handler.hpp
+++ b/sycl/include/sycl/interop_handler.hpp
@@ -26,7 +26,7 @@ class __SYCL_DEPRECATED("interop_handler class is deprecated, use"
 
 public:
   using QueueImplPtr = std::shared_ptr<detail::queue_impl>;
-  using ReqToMem = std::pair<detail::Requirement *, pi_mem>;
+  using ReqToMem = std::pair<detail::AccessorImplHost *, pi_mem>;
 
   interop_handler(std::vector<ReqToMem> MemObjs, QueueImplPtr Queue)
       : MQueue(std::move(Queue)), MMemObjs(std::move(MemObjs)) {}
@@ -59,7 +59,8 @@ class __SYCL_DEPRECATED("interop_handler class is deprecated, use"
   template <backend BackendName, typename DataT, int Dims,
             access::mode AccessMode, access::target AccessTarget,
             access::placeholder IsPlaceholder>
-  auto getMemImpl(detail::Requirement *Req) const -> typename detail::interop<
+  auto
+  getMemImpl(detail::AccessorImplHost *Req) const -> typename detail::interop<
       BackendName,
       accessor<DataT, Dims, AccessMode, AccessTarget, IsPlaceholder>>::type {
     return (typename detail::interop<
@@ -67,7 +68,8 @@ class __SYCL_DEPRECATED("interop_handler class is deprecated, use"
                                   IsPlaceholder>>::type)GetNativeMem(Req);
   }
 
-  __SYCL_EXPORT pi_native_handle GetNativeMem(detail::Requirement *Req) const;
+  __SYCL_EXPORT pi_native_handle
+  GetNativeMem(detail::AccessorImplHost *Req) const;
   __SYCL_EXPORT pi_native_handle GetNativeQueue() const;
 };
 
diff --git a/sycl/plugins/esimd_emulator/CMakeLists.txt b/sycl/plugins/esimd_emulator/CMakeLists.txt
index b033a5e3e5199..0b32c39259a13 100755
--- a/sycl/plugins/esimd_emulator/CMakeLists.txt
+++ b/sycl/plugins/esimd_emulator/CMakeLists.txt
@@ -6,6 +6,7 @@
 include(ExternalProject)
 
 include_directories("${sycl_inc_dir}")
+include_directories("${SYCL_SOURCE_DIR}/source/")
 # FIXME/TODO: 'pi.h' is included in 'pi_esimd_emulator.cpp', and CL_*_INTEL
 # and CL_*_KHR definitions in 'pi.h' are from
 # ${OPENCL_INCLUDE}. Remove build dependency on OpenCL
diff --git a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp
index 84725166b0d1e..f27f2c6eecd08 100644
--- a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp
+++ b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp
@@ -14,8 +14,8 @@
 
 #include <stdint.h>
 
+#include <detail/accessor_impl.hpp>
 #include <sycl/backend_types.hpp>
-#include <sycl/detail/accessor_impl.hpp>
 #include <sycl/detail/common.hpp>
 #include <sycl/detail/export.hpp>
 #include <sycl/detail/helpers.hpp>
diff --git a/sycl/source/accessor.cpp b/sycl/source/accessor.cpp
index fd49b21e799b8..502eff15518c9 100644
--- a/sycl/source/accessor.cpp
+++ b/sycl/source/accessor.cpp
@@ -15,6 +15,71 @@ namespace detail {
 device getDeviceFromHandler(handler &CommandGroupHandlerRef) {
   return CommandGroupHandlerRef.MQueue->get_device();
 }
+
+AccessorBaseHost::AccessorBaseHost(id<3> Offset, range<3> AccessRange,
+                                   range<3> MemoryRange,
+                                   access::mode AccessMode, void *SYCLMemObject,
+                                   int Dims, int ElemSize, int OffsetInBytes,
+                                   bool IsSubBuffer,
+                                   const property_list &PropertyList) {
+  impl = std::shared_ptr<AccessorImplHost>(
+      new AccessorImplHost(Offset, AccessRange, MemoryRange, AccessMode,
+                           (detail::SYCLMemObjI *)SYCLMemObject, Dims, ElemSize,
+                           OffsetInBytes, IsSubBuffer, PropertyList));
+}
+
+id<3> &AccessorBaseHost::getOffset() { return impl->MOffset; }
+range<3> &AccessorBaseHost::getAccessRange() { return impl->MAccessRange; }
+range<3> &AccessorBaseHost::getMemoryRange() { return impl->MMemoryRange; }
+void *AccessorBaseHost::getPtr() { return impl->MData; }
+
+const property_list &AccessorBaseHost::getPropList() const {
+  return impl->MPropertyList;
+}
+
+unsigned int AccessorBaseHost::getElemSize() const { return impl->MElemSize; }
+
+const id<3> &AccessorBaseHost::getOffset() const { return impl->MOffset; }
+const range<3> &AccessorBaseHost::getAccessRange() const {
+  return impl->MAccessRange;
+}
+const range<3> &AccessorBaseHost::getMemoryRange() const {
+  return impl->MMemoryRange;
+}
+void *AccessorBaseHost::getPtr() const {
+  return const_cast<void *>(impl->MData);
+}
+
+void *AccessorBaseHost::getMemoryObject() const { return impl->MSYCLMemObj; }
+
+LocalAccessorBaseHost::LocalAccessorBaseHost(sycl::range<3> Size, int Dims,
+                                             int ElemSize) {
+  impl = std::shared_ptr<LocalAccessorImplHost>(
+      new LocalAccessorImplHost(Size, Dims, ElemSize));
+}
+sycl::range<3> &LocalAccessorBaseHost::getSize() { return impl->MSize; }
+const sycl::range<3> &LocalAccessorBaseHost::getSize() const {
+  return impl->MSize;
+}
+void *LocalAccessorBaseHost::getPtr() {
+  // Const cast this in order to call the const getPtr.
+  return const_cast<const LocalAccessorBaseHost *>(this)->getPtr();
+}
+void *LocalAccessorBaseHost::getPtr() const {
+  char *ptr = impl->MMem.data();
+
+  // Align the pointer to MElemSize.
+  size_t val = reinterpret_cast<size_t>(ptr);
+  if (val % impl->MElemSize != 0) {
+    ptr += impl->MElemSize - val % impl->MElemSize;
+  }
+
+  return ptr;
+}
+
+int LocalAccessorBaseHost::getNumOfDims() { return impl->MDims; }
+int LocalAccessorBaseHost::getElementSize() { return impl->MElemSize; }
+
 } // namespace detail
 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
 } // namespace sycl
diff --git a/sycl/source/detail/accessor_impl.cpp b/sycl/source/detail/accessor_impl.cpp
index 41f3858ec3176..651dba7360783 100644
--- a/sycl/source/detail/accessor_impl.cpp
+++ b/sycl/source/detail/accessor_impl.cpp
@@ -6,11 +6,11 @@
 //
 //===----------------------------------------------------------------------===//
 
+#include <detail/accessor_impl.hpp>
 #include <detail/buffer_impl.hpp>
 #include <detail/event_impl.hpp>
 #include <detail/scheduler/scheduler.hpp>
 #include <detail/xpti_registry.hpp>
-#include <sycl/detail/accessor_impl.hpp>
 
 namespace sycl {
 __SYCL_INLINE_VER_NAMESPACE(_V1) {
diff --git a/sycl/source/detail/accessor_impl.hpp b/sycl/source/detail/accessor_impl.hpp
new file mode 100644
index 0000000000000..ea6e2616d0b2e
--- /dev/null
+++ b/sycl/source/detail/accessor_impl.hpp
@@ -0,0 +1,120 @@
+//==------------ accessor_impl.hpp - SYCL standard header file -------------==//
+//
+// 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/access/access.hpp>
+#include <sycl/detail/export.hpp>
+#include <sycl/id.hpp>
+#include <sycl/property_list.hpp>
+#include <sycl/range.hpp>
+#include <sycl/stl.hpp>
+
+namespace sycl {
+__SYCL_INLINE_VER_NAMESPACE(_V1) {
+template <typename, int, access::mode, access::target, access::placeholder,
+          typename>
+class accessor;
+
+namespace ext {
+namespace intel {
+namespace esimd {
+namespace detail {
+// Forward declare a "back-door" access class to support ESIMD.
+class AccessorPrivateProxy;
+} // namespace detail
+} // namespace esimd
+} // namespace intel
+} // namespace ext
+
+namespace detail {
+
+class SYCLMemObjI;
+
+class Command;
+
+class __SYCL_EXPORT AccessorImplHost {
+public:
+  AccessorImplHost(id<3> Offset, range<3> AccessRange, range<3> MemoryRange,
+                   access::mode AccessMode, void *SYCLMemObject, int Dims,
+                   int ElemSize, int OffsetInBytes = 0,
+                   bool IsSubBuffer = false,
+                   const property_list &PropertyList = {})
+      : MOffset(Offset), MAccessRange(AccessRange), MMemoryRange(MemoryRange),
+        MAccessMode(AccessMode),
+        MSYCLMemObj((detail::SYCLMemObjI *)SYCLMemObject), MDims(Dims),
+        MElemSize(ElemSize), MOffsetInBytes(OffsetInBytes),
+        MIsSubBuffer(IsSubBuffer), MPropertyList(PropertyList) {}
+
+  ~AccessorImplHost();
+
+  AccessorImplHost(const AccessorImplHost &Other)
+      : MOffset(Other.MOffset), MAccessRange(Other.MAccessRange),
+        MMemoryRange(Other.MMemoryRange), MAccessMode(Other.MAccessMode),
+        MSYCLMemObj(Other.MSYCLMemObj), MDims(Other.MDims),
+        MElemSize(Other.MElemSize), MOffsetInBytes(Other.MOffsetInBytes),
+        MIsSubBuffer(Other.MIsSubBuffer), MPropertyList(Other.MPropertyList) {}
+
+  // The resize method provides a way to change the size of the
+  // allocated memory and corresponding properties for the accessor.
+  // These are normally fixed for the accessor, but this capability
+  // is needed to support the stream class.
+  // Stream implementation creates an accessor with initial size for
+  // work item. But the number of work items is not available during
+  // stream construction. The resize method allows to update the accessor
+  // as the information becomes available to the handler.
+
+  void resize(size_t GlobalSize);
+
+  id<3> MOffset;
+  // The size of accessing region.
+  range<3> MAccessRange;
+  // The size of memory object this requirement is created for.
+  range<3> MMemoryRange;
+  access::mode MAccessMode;
+
+  detail::SYCLMemObjI *MSYCLMemObj;
+
+  unsigned int MDims;
+  unsigned int MElemSize;
+  unsigned int MOffsetInBytes;
+  bool MIsSubBuffer;
+
+  void *MData = nullptr;
+
+  Command *MBlockedCmd = nullptr;
+
+  bool PerWI = false;
+
+  // To preserve runtime properties
+  property_list MPropertyList;
+};
+
+using AccessorImplPtr = std::shared_ptr<AccessorImplHost>;
+
+class __SYCL_EXPORT LocalAccessorImplHost {
+public:
+  // Allocate ElemSize more data to have sufficient padding to enforce
+  // alignment.
+  LocalAccessorImplHost(sycl::range<3> Size, int Dims, int ElemSize)
+      : MSize(Size), MDims(Dims), MElemSize(ElemSize),
+        MMem(Size[0] * Size[1] * Size[2] * ElemSize + ElemSize) {}
+
+  sycl::range<3> MSize;
+  int MDims;
+  int MElemSize;
+  std::vector<char> MMem;
+};
+
+using LocalAccessorImplPtr = std::shared_ptr<LocalAccessorImplHost>;
+
+using Requirement = AccessorImplHost;
+
+} // namespace detail
+} // __SYCL_INLINE_VER_NAMESPACE(_V1)
+} // namespace sycl
diff --git a/sycl/source/detail/scheduler/commands.hpp b/sycl/source/detail/scheduler/commands.hpp
index d95f0e307456a..f77d90d7a630a 100644
--- a/sycl/source/detail/scheduler/commands.hpp
+++ b/sycl/source/detail/scheduler/commands.hpp
@@ -17,10 +17,10 @@
 #include <unordered_set>
 #include <vector>
 
+#include <detail/accessor_impl.hpp>
 #include <detail/event_impl.hpp>
 #include <detail/program_manager/program_manager.hpp>
 #include <sycl/access/access.hpp>
-#include <sycl/detail/accessor_impl.hpp>
 #include <sycl/detail/cg.hpp>
 
 namespace sycl {
diff --git a/sycl/source/interop_handle.cpp b/sycl/source/interop_handle.cpp
index d0a560777b7b0..4900329a2eece 100644
--- a/sycl/source/interop_handle.cpp
+++ b/sycl/source/interop_handle.cpp
@@ -6,11 +6,11 @@
 //
 //===----------------------------------------------------------------------===//
 
+#include <detail/accessor_impl.hpp>
 #include <detail/backend_impl.hpp>
 #include <detail/context_impl.hpp>
 #include <detail/device_impl.hpp>
 #include <detail/queue_impl.hpp>
-#include <sycl/detail/accessor_impl.hpp>
 #include <sycl/exception.hpp>
 #include <sycl/interop_handle.hpp>
 
diff --git a/sycl/test/abi/layout_accessors_host.cpp b/sycl/test/abi/layout_accessors_host.cpp
index 09f2a088bbb98..ef03fac69c89d 100644
--- a/sycl/test/abi/layout_accessors_host.cpp
+++ b/sycl/test/abi/layout_accessors_host.cpp
@@ -8,59 +8,6 @@
 
 using namespace sycl;
 
-// CHECK:       0 | class sycl::detail::AccessorImplHost
-// CHECK-NEXT:  0 |   class sycl::id<3> MOffset
-// CHECK-NEXT:  0 |     class sycl::detail::array<3> (base)
-// CHECK-NEXT:  0 |       size_t[3] common_array
-// CHECK-NEXT: 24 |   class sycl::range<3> MAccessRange
-// CHECK-NEXT: 24 |     class sycl::detail::array<3> (base)
-// CHECK-NEXT: 24 |       size_t[3] common_array
-// CHECK-NEXT: 48 |   class sycl::range<3> MMemoryRange
-// CHECK-NEXT: 48 |     class sycl::detail::array<3> (base)
-// CHECK-NEXT: 48 |       size_t[3] common_array
-// CHECK-NEXT: 72 |   access::mode MAccessMode
-// CHECK-NEXT: 80 |   detail::SYCLMemObjI * MSYCLMemObj
-// CHECK-NEXT: 88 |   unsigned int MDims
-// CHECK-NEXT: 92 |   unsigned int MElemSize
-// CHECK-NEXT: 96 |   unsigned int MOffsetInBytes
-// CHECK-NEXT: 00 |   _Bool MIsSubBuffer
-// CHECK-NEXT: 04 |   void * MData
-// CHECK-NEXT: 12 |   Command * MBlockedCmd
-// CHECK-NEXT: 20 |   _Bool PerWI
-// CHECK-NEXT: 28 |   class sycl::property_list MPropertyList
-// CHECK-NEXT: 28 |     class sycl::detail::PropertyListBase (base)
-// CHECK-NEXT: 28 |       class std::bitset<32> MDataLessProps
-// CHECK-NEXT: 28 |         struct std::_Base_bitset<1> (base)
-// CHECK-NEXT: 28 |           _WordT _M_w
-// CHECK-NEXT: 36 |       class std::vector<class std::shared_ptr<class sycl::detail::PropertyWithDataBase> > MPropsWithData
-// CHECK-NEXT: 36 |         struct std::_Vector_base<class std::shared_ptr<class sycl::detail::PropertyWithDataBase>, class std::allocator<class std::shared_ptr<class sycl::detail::PropertyWithDataBase> > > (base)
-// CHECK-NEXT: 36 |           struct std::_Vector_base<class std::shared_ptr<class sycl::detail::PropertyWithDataBase>, class std::allocator<class std::shared_ptr<class sycl::detail::PropertyWithDataBase> > >::_Vector_impl _M_impl
-// CHECK-NEXT: 36 |             class std::allocator<class std::shared_ptr<class sycl::detail::PropertyWithDataBase> > (base) (empty)
-// CHECK-NEXT: 36 |               class __gnu_cxx::new_allocator<class std::shared_ptr<class sycl::detail::PropertyWithDataBase> > (base) (empty)
-// CHECK-NEXT: 36 |             struct std::_Vector_base<class std::shared_ptr<class sycl::detail::PropertyWithDataBase>, class std::allocator<class std::shared_ptr<class sycl::detail::PropertyWithDataBase> > >::_Vector_impl_data (base)
-// CHECK-NEXT: 36 |               pointer _M_start
-// CHECK-NEXT: 44 |               pointer _M_finish
-// CHECK-NEXT: 52 |               pointer _M_end_of_storage
-// CHECK-NEXT:    | [sizeof=160, dsize=160, align=8,
-// CHECK-NEXT:    |  nvsize=160, nvalign=8]
-
-// CHECK:       0 | class sycl::detail::LocalAccessorImplHost
-// CHECK-NEXT:  0 |   class sycl::range<3> MSize
-// CHECK-NEXT:  0 |     class sycl::detail::array<3> (base)
-// CHECK-NEXT:  0 |       size_t[3] common_array
-// CHECK-NEXT: 24 |   int MDims
-// CHECK-NEXT: 28 |   int MElemSize
-// CHECK-NEXT: 32 |   class std::vector<char> MMem
-// CHECK-NEXT: 32 |     struct std::_Vector_base<char, class std::allocator<char> > (base)
-// CHECK-NEXT: 32 |       struct std::_Vector_base<char, class std::allocator<char> >::_Vector_impl _M_impl
-// CHECK-NEXT: 32 |         class std::allocator<char> (base) (empty)
-// CHECK-NEXT: 32 |           class __gnu_cxx::new_allocator<char> (base) (empty)
-// CHECK:      32 |           pointer _M_start
-// CHECK-NEXT: 40 |           pointer _M_finish
-// CHECK-NEXT: 48 |           pointer _M_end_of_storage
-// CHECK-NEXT:    | [sizeof=56, dsize=56, align=8,
-// CHECK-NEXT:    |  nvsize=56, nvalign=8]
-
 //----------------------------------------------------------------------------//
 // Host buffer accessor.
 //----------------------------------------------------------------------------//
diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump
index cc6cc3f545a0c..ca6b049b964bf 100644
--- a/sycl/test/abi/sycl_symbols_linux.dump
+++ b/sycl/test/abi/sycl_symbols_linux.dump
@@ -3812,6 +3812,12 @@ _ZN4sycl3_V16detail13select_deviceERKSt8functionIFiRKNS0_6deviceEEE
 _ZN4sycl3_V16detail13select_deviceERKSt8functionIFiRKNS0_6deviceEEERKNS0_7contextE
 _ZN4sycl3_V16detail14getBorderColorENS0_19image_channel_orderE
 _ZN4sycl3_V16detail15getOrWaitEventsESt6vectorINS0_5eventESaIS3_EESt10shared_ptrINS1_12context_implEE
+_ZN4sycl3_V16detail16AccessorBaseHost14getAccessRangeEv
+_ZN4sycl3_V16detail16AccessorBaseHost14getMemoryRangeEv
+_ZN4sycl3_V16detail16AccessorBaseHost6getPtrEv
+_ZN4sycl3_V16detail16AccessorBaseHost9getOffsetEv
+_ZN4sycl3_V16detail16AccessorBaseHostC1ENS0_2idILi3EEENS0_5rangeILi3EEES6_NS0_6access4modeEPviiibRKNS0_13property_listE
+_ZN4sycl3_V16detail16AccessorBaseHostC2ENS0_2idILi3EEENS0_5rangeILi3EEES6_NS0_6access4modeEPviiibRKNS0_13property_listE
 _ZN4sycl3_V16detail16AccessorImplHost6resizeEm
 _ZN4sycl3_V16detail16AccessorImplHostD1Ev
 _ZN4sycl3_V16detail16AccessorImplHostD2Ev
@@ -3833,6 +3839,12 @@ _ZN4sycl3_V16detail19getPluginOpaqueDataILNS0_7backendE5EEEPvS4_
 _ZN4sycl3_V16detail19kernel_bundle_plain32set_specialization_constant_implEPKcPvm
 _ZN4sycl3_V16detail20associateWithHandlerERNS0_7handlerEPNS1_16AccessorBaseHostENS0_6access6targetE
 _ZN4sycl3_V16detail20getDeviceFromHandlerERNS0_7handlerE
+_ZN4sycl3_V16detail21LocalAccessorBaseHost12getNumOfDimsEv
+_ZN4sycl3_V16detail21LocalAccessorBaseHost14getElementSizeEv
+_ZN4sycl3_V16detail21LocalAccessorBaseHost6getPtrEv
+_ZN4sycl3_V16detail21LocalAccessorBaseHost7getSizeEv
+_ZN4sycl3_V16detail21LocalAccessorBaseHostC1ENS0_5rangeILi3EEEii
+_ZN4sycl3_V16detail21LocalAccessorBaseHostC2ENS0_5rangeILi3EEEii
 _ZN4sycl3_V16detail22addHostAccessorAndWaitEPNS1_16AccessorImplHostE
 _ZN4sycl3_V16detail22getImageNumberChannelsENS0_19image_channel_orderE
 _ZN4sycl3_V16detail22get_kernel_bundle_implERKNS0_7contextERKSt6vectorINS0_6deviceESaIS6_EENS0_12bundle_stateE
@@ -4099,6 +4111,13 @@ _ZNK4sycl3_V16detail12buffer_plain7getSizeEv
 _ZNK4sycl3_V16detail12sampler_impl18get_filtering_modeEv
 _ZNK4sycl3_V16detail12sampler_impl19get_addressing_modeEv
 _ZNK4sycl3_V16detail12sampler_impl33get_coordinate_normalization_modeEv
+_ZNK4sycl3_V16detail16AccessorBaseHost11getElemSizeEv
+_ZNK4sycl3_V16detail16AccessorBaseHost11getPropListEv
+_ZNK4sycl3_V16detail16AccessorBaseHost14getAccessRangeEv
+_ZNK4sycl3_V16detail16AccessorBaseHost14getMemoryRangeEv
+_ZNK4sycl3_V16detail16AccessorBaseHost15getMemoryObjectEv
+_ZNK4sycl3_V16detail16AccessorBaseHost6getPtrEv
+_ZNK4sycl3_V16detail16AccessorBaseHost9getOffsetEv
 _ZNK4sycl3_V16detail18device_image_plain10has_kernelERKNS0_9kernel_idE
 _ZNK4sycl3_V16detail18device_image_plain10has_kernelERKNS0_9kernel_idERKNS0_6deviceE
 _ZNK4sycl3_V16detail18device_image_plain9getNativeEv
@@ -4117,6 +4136,8 @@ _ZNK4sycl3_V16detail19kernel_bundle_plain33contains_specialization_constantsEv
 _ZNK4sycl3_V16detail19kernel_bundle_plain3endEv
 _ZNK4sycl3_V16detail19kernel_bundle_plain5beginEv
 _ZNK4sycl3_V16detail19kernel_bundle_plain5emptyEv
+_ZNK4sycl3_V16detail21LocalAccessorBaseHost6getPtrEv
+_ZNK4sycl3_V16detail21LocalAccessorBaseHost7getSizeEv
 _ZNK4sycl3_V16device11get_backendEv
 _ZNK4sycl3_V16device12get_platformEv
 _ZNK4sycl3_V16device13has_extensionERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE
diff --git a/sycl/test/abi/symbol_size_alignment.cpp b/sycl/test/abi/symbol_size_alignment.cpp
index c681df3434c90..2372c13d65e06 100644
--- a/sycl/test/abi/symbol_size_alignment.cpp
+++ b/sycl/test/abi/symbol_size_alignment.cpp
@@ -6,7 +6,6 @@
 
 #include <sycl/accessor.hpp>
 #include <sycl/buffer.hpp>
-#include <sycl/detail/accessor_impl.hpp>
 #include <sycl/device.hpp>
 #include <sycl/device_event.hpp>
 #include <sycl/device_selector.hpp>
@@ -43,9 +42,7 @@ int main() {
   check<accessor_t, 32, 8>();
   check<detail::AccessorImplDevice<1>, 24, 8>();
   check<detail::LocalAccessorBaseDevice<1>, 24, 8>();
-  check<detail::AccessorImplHost, 160, 8>();
   check<detail::AccessorBaseHost, 16, 8>();
-  check<detail::LocalAccessorImplHost, 56, 8>();
   check<buffer<int>, 40, 8>();
   check<context, 16, 8>();
   check<cpu_selector, 8, 8>();
diff --git a/sycl/test/gdb/accessors.cpp b/sycl/test/gdb/accessors.cpp
index 671a4100833ac..3649eb39d6c30 100644
--- a/sycl/test/gdb/accessors.cpp
+++ b/sycl/test/gdb/accessors.cpp
@@ -5,36 +5,23 @@
 
 void foo(sycl::buffer<int, 1> &BufA) {
   auto HostAcc = BufA.get_access<sycl::access_mode::read>();
-}
-
-// AccessorImplHost must have MMemoryRange, MOffset and MData fields
-
-// CHECK: CXXRecordDecl {{.*}} class AccessorImplHost definition
-// CHECK-NOT: CXXRecordDecl {{.*}} definition
-// CHECK: FieldDecl {{.*}} referenced MOffset
-// CHECK-NOT: CXXRecordDecl {{.*}} definition
-// CHECK: FieldDecl {{.*}} referenced MMemoryRange
-// CHECK-NOT: CXXRecordDecl {{.*}} definition
-// CHECK: FieldDecl {{.*}} referenced MData
-
-// accessor.impl must be present and of shared_ptr<AccessorImplHost> type
 
-// CHECK: CXXRecordDecl {{.*}} class AccessorBaseHost definition
-// CHECK-NOT: CXXRecordDecl {{.*}} definition
-// CHECK: FieldDecl {{.*}} referenced impl {{.*}}:'std::shared_ptr<sycl::detail::AccessorImplHost>'
-
-// LocalAccessorImplHost must have MSize and MMem fields
-
-// CHECK: CXXRecordDecl {{.*}} class LocalAccessorImplHost definition
-// CHECK-NOT: CXXRecordDecl {{.*}} definition
-// CHECK: FieldDecl {{.*}} referenced MSize
-// CHECK-NOT: CXXRecordDecl {{.*}} definition
-// CHECK: FieldDecl {{.*}} referenced MMem
+  sycl::accessor<char, 1, sycl::access::mode::read_write, sycl::target::local>
+      *LocalAcc;
+}
 
+// Host accessors should have the following methods which are used by gdb
+// pretty-printers
+//
 // CHECK: CXXRecordDecl {{.*}} class accessor definition
-// CHECK-NOT: CXXRecordDecl {{.*}} definition
-// CHECK: public {{.*}}:'sycl::detail::AccessorBaseHost'
-
+// CHECK: CXXMethodDecl {{.*}}getOffset
+// CHECK: CXXMethodDecl {{.*}}getAccessRange
+// CHECK: CXXMethodDecl {{.*}}getMemoryRange
+// CHECK: CXXMethodDecl {{.*}}getPtr
+
+// CHECK: CXXRecordDecl {{.*}} class local_accessor_base definition
+// CHECK: CXXMethodDecl {{.*}}getSize
+// CHECK: CXXMethodDecl {{.*}}getPtr
 
 // CHECK-DEBUG-INFO: !DICompositeType(tag: DW_TAG_class_type, name: "accessor<int, 1, (sycl::_V1::access::mode)1024, (sycl::_V1::access::target)2018, (sycl::_V1::access::placeholder)0, sycl::_V1::ext::oneapi::accessor_property_list<> >", {{.*}}, templateParams: ![[TEMPL_METADATA:[0-9]+]]
 // CHECK-DEBUG-INFO: ![[TEMPL_METADATA]] = !{![[DATA_T:[0-9]+]], ![[Dims:[0-9]+]], ![[AccMode:[0-9]+]], ![[AccTarget:[0-9]+]], ![[IsPlh:[0-9]+]], ![[PropListT:[0-9]+]]}