diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp
index 834fc0cbcc246..3f3b0da4bcc26 100644
--- a/sycl/include/sycl/accessor.hpp
+++ b/sycl/include/sycl/accessor.hpp
@@ -19,12 +19,14 @@
 #include <sycl/detail/handler_proxy.hpp>
 #include <sycl/detail/image_accessor_util.hpp>
 #include <sycl/detail/image_ocl_types.hpp>
+#include <sycl/device.hpp>
 #include <sycl/exception.hpp>
 #include <sycl/ext/oneapi/accessor_property_list.hpp>
 #include <sycl/id.hpp>
 #include <sycl/image.hpp>
 #include <sycl/pointers.hpp>
 #include <sycl/properties/accessor_properties.hpp>
+#include <sycl/properties/buffer_properties.hpp>
 #include <sycl/property_list.hpp>
 #include <sycl/property_list_conversion.hpp>
 #include <sycl/sampler.hpp>
@@ -264,20 +266,6 @@ template <> struct IsCxPropertyList<ext::oneapi::accessor_property_list<>> {
   constexpr static bool value = false;
 };
 
-// The function extends or truncates number of dimensions of objects of id
-// or ranges classes. When extending the new values are filled with
-// DefaultValue, truncation just removes extra values.
-template <int NewDim, int DefaultValue, template <int> class T, int OldDim>
-static T<NewDim> convertToArrayOfN(T<OldDim> OldObj) {
-  T<NewDim> NewObj = InitializedVal<NewDim, T>::template get<0>();
-  const int CopyDims = NewDim > OldDim ? OldDim : NewDim;
-  for (int I = 0; I < CopyDims; ++I)
-    NewObj[I] = OldObj[I];
-  for (int I = CopyDims; I < NewDim; ++I)
-    NewObj[I] = DefaultValue;
-  return NewObj;
-}
-
 __SYCL_EXPORT device getDeviceFromHandler(handler &CommandGroupHandlerRef);
 
 template <typename DataT, int Dimensions, access::mode AccessMode,
@@ -526,15 +514,14 @@ class image_accessor
     // host.
   }
 #else
-      : AccessorBaseHost({detail::getSyclObjImpl(ImageRef)->getRowPitch(),
-                          detail::getSyclObjImpl(ImageRef)->getSlicePitch(), 0},
+      : AccessorBaseHost({ImageRef.getRowPitch(), ImageRef.getSlicePitch(), 0},
                          detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
                          detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
                          AccessMode, detail::getSyclObjImpl(ImageRef).get(),
                          Dimensions, ImageElementSize),
         MImageCount(ImageRef.size()),
-        MImgChannelOrder(detail::getSyclObjImpl(ImageRef)->getChannelOrder()),
-        MImgChannelType(detail::getSyclObjImpl(ImageRef)->getChannelType()) {
+        MImgChannelOrder(ImageRef.getChannelOrder()),
+        MImgChannelType(ImageRef.getChannelType()) {
     addHostAccessorAndWait(AccessorBaseHost::impl.get());
   }
 #endif
@@ -557,15 +544,14 @@ class image_accessor
     // host.
   }
 #else
-      : AccessorBaseHost({detail::getSyclObjImpl(ImageRef)->getRowPitch(),
-                          detail::getSyclObjImpl(ImageRef)->getSlicePitch(), 0},
+      : AccessorBaseHost({ImageRef.getRowPitch(), ImageRef.getSlicePitch(), 0},
                          detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
                          detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
                          AccessMode, detail::getSyclObjImpl(ImageRef).get(),
                          Dimensions, ImageElementSize),
         MImageCount(ImageRef.size()),
-        MImgChannelOrder(detail::getSyclObjImpl(ImageRef)->getChannelOrder()),
-        MImgChannelType(detail::getSyclObjImpl(ImageRef)->getChannelType()) {
+        MImgChannelOrder(ImageRef.getChannelOrder()),
+        MImgChannelType(ImageRef.getChannelType()) {
     checkDeviceFeatureSupported<info::device::image_support>(
         getDeviceFromHandler(CommandGroupHandlerRef));
   }
@@ -1203,7 +1189,7 @@ class __SYCL_SPECIAL_CLASS accessor :
       const property_list &PropertyList = {},
       const detail::code_location CodeLoc = detail::code_location::current())
       : accessor(BufferRef, PropertyList, CodeLoc) {
-    adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
+    adjustAccPropsInBuf(BufferRef);
   }
 
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
@@ -1218,7 +1204,7 @@ class __SYCL_SPECIAL_CLASS accessor :
           {},
       const detail::code_location CodeLoc = detail::code_location::current())
       : accessor(BufferRef, PropertyList, CodeLoc) {
-    adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
+    adjustAccPropsInBuf(BufferRef);
   }
 #endif
 
@@ -1297,7 +1283,7 @@ class __SYCL_SPECIAL_CLASS accessor :
       TagT, const property_list &PropertyList = {},
       const detail::code_location CodeLoc = detail::code_location::current())
       : accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {
-    adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
+    adjustAccPropsInBuf(BufferRef);
   }
 
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
@@ -1313,7 +1299,7 @@ class __SYCL_SPECIAL_CLASS accessor :
           {},
       const detail::code_location CodeLoc = detail::code_location::current())
       : accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {
-    adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
+    adjustAccPropsInBuf(BufferRef);
   }
 
 #endif
@@ -1357,7 +1343,7 @@ class __SYCL_SPECIAL_CLASS accessor :
       TagT, const property_list &PropertyList = {},
       const detail::code_location CodeLoc = detail::code_location::current())
       : accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {
-    adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
+    adjustAccPropsInBuf(BufferRef);
   }
 
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
@@ -1373,7 +1359,7 @@ class __SYCL_SPECIAL_CLASS accessor :
           {},
       const detail::code_location CodeLoc = detail::code_location::current())
       : accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {
-    adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
+    adjustAccPropsInBuf(BufferRef);
   }
 #endif
 
@@ -1419,7 +1405,7 @@ class __SYCL_SPECIAL_CLASS accessor :
       const detail::code_location CodeLoc = detail::code_location::current())
       : accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
                  CodeLoc) {
-    adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
+    adjustAccPropsInBuf(BufferRef);
   }
 
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
@@ -1436,7 +1422,7 @@ class __SYCL_SPECIAL_CLASS accessor :
       const detail::code_location CodeLoc = detail::code_location::current())
       : accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
                  CodeLoc) {
-    adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
+    adjustAccPropsInBuf(BufferRef);
   }
 #endif
 
@@ -1532,7 +1518,7 @@ class __SYCL_SPECIAL_CLASS accessor :
       id<Dimensions> AccessOffset, TagT, const property_list &PropertyList = {},
       const detail::code_location CodeLoc = detail::code_location::current())
       : accessor(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {
-    adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
+    adjustAccPropsInBuf(BufferRef);
   }
 
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
@@ -1548,7 +1534,7 @@ class __SYCL_SPECIAL_CLASS accessor :
           {},
       const detail::code_location CodeLoc = detail::code_location::current())
       : accessor(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {
-    adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
+    adjustAccPropsInBuf(BufferRef);
   }
 #endif
 
@@ -1645,7 +1631,7 @@ class __SYCL_SPECIAL_CLASS accessor :
       const detail::code_location CodeLoc = detail::code_location::current())
       : accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
                  PropertyList, CodeLoc) {
-    adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
+    adjustAccPropsInBuf(BufferRef);
   }
 
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
@@ -1662,7 +1648,7 @@ class __SYCL_SPECIAL_CLASS accessor :
       const detail::code_location CodeLoc = detail::code_location::current())
       : accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
                  PropertyList, CodeLoc) {
-    adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
+    adjustAccPropsInBuf(BufferRef);
   }
 #endif
 
@@ -1854,8 +1840,8 @@ class __SYCL_SPECIAL_CLASS accessor :
   }
 
 #if __cplusplus >= 201703L
-  template <typename... PropTypes>
-  void adjustAccPropsInBuf(detail::SYCLMemObjI *SYCLMemObject) {
+  template <typename BufT, typename... PropTypes>
+  void adjustAccPropsInBuf(BufT &Buffer) {
     if constexpr (PropertyListT::template has_property<
                       sycl::ext::intel::property::buffer_location>()) {
       auto location = (PropertyListT::template get_property<
@@ -1863,18 +1849,14 @@ class __SYCL_SPECIAL_CLASS accessor :
                           .get_location();
       property_list PropList{
           sycl::property::buffer::detail::buffer_location(location)};
-      detail::SYCLMemObjT *SYCLMemObjectT =
-          dynamic_cast<detail::SYCLMemObjT *>(SYCLMemObject);
-      SYCLMemObjectT->addOrReplaceAccessorProperties(PropList);
+      Buffer.addOrReplaceAccessorProperties(PropList);
     } else {
-      deleteAccPropsFromBuf(SYCLMemObject);
+      deleteAccPropsFromBuf(Buffer);
     }
   }
 
-  void deleteAccPropsFromBuf(detail::SYCLMemObjI *SYCLMemObject) {
-    detail::SYCLMemObjT *SYCLMemObjectT =
-        dynamic_cast<detail::SYCLMemObjT *>(SYCLMemObject);
-    SYCLMemObjectT->deleteAccessorProperty(
+  template <typename BufT> void deleteAccPropsFromBuf(BufT &Buffer) {
+    Buffer.deleteAccProps(
         sycl::detail::PropWithDataKind::AccPropBufferLocation);
   }
 #endif
@@ -2261,8 +2243,7 @@ class __SYCL_SPECIAL_CLASS accessor<DataT, Dimensions, AccessMode,
            handler &CommandGroupHandler)
       : detail::image_accessor<DataT, Dimensions, AccessMode,
                                access::target::image, IsPlaceholder>(
-            Image, CommandGroupHandler,
-            (detail::getSyclObjImpl(Image))->getElementSize()) {
+            Image, CommandGroupHandler, Image.getElementSize()) {
 #ifndef __SYCL_DEVICE_ONLY__
     detail::associateWithHandler(CommandGroupHandler, this,
                                  access::target::image);
@@ -2274,8 +2255,7 @@ class __SYCL_SPECIAL_CLASS accessor<DataT, Dimensions, AccessMode,
            handler &CommandGroupHandler, const property_list &propList)
       : detail::image_accessor<DataT, Dimensions, AccessMode,
                                access::target::image, IsPlaceholder>(
-            Image, CommandGroupHandler,
-            (detail::getSyclObjImpl(Image))->getElementSize()) {
+            Image, CommandGroupHandler, Image.getElementSize()) {
     (void)propList;
 #ifndef __SYCL_DEVICE_ONLY__
     detail::associateWithHandler(CommandGroupHandler, this,
@@ -2319,14 +2299,14 @@ class accessor<DataT, Dimensions, AccessMode, access::target::host_image,
   accessor(sycl::image<Dimensions, AllocatorT> &Image)
       : detail::image_accessor<DataT, Dimensions, AccessMode,
                                access::target::host_image, IsPlaceholder>(
-            Image, (detail::getSyclObjImpl(Image))->getElementSize()) {}
+            Image, Image.getElementSize()) {}
 
   template <typename AllocatorT>
   accessor(sycl::image<Dimensions, AllocatorT> &Image,
            const property_list &propList)
       : detail::image_accessor<DataT, Dimensions, AccessMode,
                                access::target::host_image, IsPlaceholder>(
-            Image, (detail::getSyclObjImpl(Image))->getElementSize()) {
+            Image, Image.getElementSize()) {
     (void)propList;
   }
 };
@@ -2368,8 +2348,7 @@ class __SYCL_SPECIAL_CLASS accessor<DataT, Dimensions, AccessMode,
            handler &CommandGroupHandler)
       : detail::image_accessor<DataT, Dimensions + 1, AccessMode,
                                access::target::image, IsPlaceholder>(
-            Image, CommandGroupHandler,
-            (detail::getSyclObjImpl(Image))->getElementSize()) {
+            Image, CommandGroupHandler, Image.getElementSize()) {
 #ifndef __SYCL_DEVICE_ONLY__
     detail::associateWithHandler(CommandGroupHandler, this,
                                  access::target::image_array);
@@ -2381,8 +2360,7 @@ class __SYCL_SPECIAL_CLASS accessor<DataT, Dimensions, AccessMode,
            handler &CommandGroupHandler, const property_list &propList)
       : detail::image_accessor<DataT, Dimensions + 1, AccessMode,
                                access::target::image, IsPlaceholder>(
-            Image, CommandGroupHandler,
-            (detail::getSyclObjImpl(Image))->getElementSize()) {
+            Image, CommandGroupHandler, Image.getElementSize()) {
     (void)propList;
 #ifndef __SYCL_DEVICE_ONLY__
     detail::associateWithHandler(CommandGroupHandler, this,
diff --git a/sycl/include/sycl/buffer.hpp b/sycl/include/sycl/buffer.hpp
index 39308e515251e..3225f7827737c 100644
--- a/sycl/include/sycl/buffer.hpp
+++ b/sycl/include/sycl/buffer.hpp
@@ -8,12 +8,14 @@
 
 #pragma once
 
-#include <sycl/detail/buffer_impl.hpp>
 #include <sycl/detail/common.hpp>
 #include <sycl/detail/stl_type_traits.hpp>
+#include <sycl/detail/sycl_mem_obj_allocator.hpp>
+#include <sycl/event.hpp>
 #include <sycl/exception.hpp>
 #include <sycl/ext/oneapi/accessor_property_list.hpp>
 #include <sycl/property_list.hpp>
+#include <sycl/range.hpp>
 #include <sycl/stl.hpp>
 
 namespace sycl {
@@ -26,8 +28,16 @@ template <int dimensions> class range;
 template <typename DataT>
 using buffer_allocator = detail::sycl_memory_object_allocator<DataT>;
 
+template <typename DataT, int Dimensions, access::mode AccessMode>
+class host_accessor;
+
+template <typename T, int Dimensions, typename AllocatorT, typename Enable>
+class buffer;
+
 namespace detail {
 
+class buffer_impl;
+
 template <typename T, int Dimensions, typename AllocatorT>
 buffer<T, Dimensions, AllocatorT, void>
 make_buffer_helper(pi_native_handle Handle, const context &Ctx, event Evt = {},
@@ -45,6 +55,71 @@ auto get_native_buffer(const buffer<DataT, Dimensions, Allocator, void> &Obj)
 template <backend Backend, typename DataT, int Dimensions,
           typename AllocatorT = buffer_allocator<std::remove_const_t<DataT>>>
 struct BufferInterop;
+
+// The non-template base for the sycl::buffer class
+class __SYCL_EXPORT buffer_plain {
+protected:
+  buffer_plain(size_t SizeInBytes, size_t, const property_list &Props,
+               std::unique_ptr<detail::SYCLMemObjAllocator> Allocator);
+
+  buffer_plain(void *HostData, size_t SizeInBytes, size_t RequiredAlign,
+               const property_list &Props,
+               std::unique_ptr<detail::SYCLMemObjAllocator> Allocator);
+
+  buffer_plain(const void *HostData, size_t SizeInBytes, size_t RequiredAlign,
+               const property_list &Props,
+               std::unique_ptr<detail::SYCLMemObjAllocator> Allocator);
+
+  buffer_plain(const std::shared_ptr<const void> &HostData,
+               const size_t SizeInBytes, size_t RequiredAlign,
+               const property_list &Props,
+               std::unique_ptr<detail::SYCLMemObjAllocator> Allocator,
+               bool IsConstPtr);
+
+  buffer_plain(const std::function<void(void *)>
+                   &CopyFromInput, // EnableIfNotConstIterator<InputIterator>
+                                   // First, InputIterator Last,
+               const size_t SizeInBytes, size_t RequiredAlign,
+               const property_list &Props,
+               std::unique_ptr<detail::SYCLMemObjAllocator> Allocator,
+               bool IsConstPtr);
+
+  buffer_plain(pi_native_handle MemObject, context SyclContext,
+               std::unique_ptr<detail::SYCLMemObjAllocator> Allocator,
+               bool OwnNativeHandle, event AvailableEvent);
+
+  buffer_plain(const std::shared_ptr<detail::buffer_impl> &impl) : impl(impl) {}
+
+  void set_final_data_internal();
+
+  void set_final_data_internal(
+      const std::function<void(const std::function<void(void *const Ptr)> &)>
+          &FinalDataFunc);
+
+  void set_write_back(bool NeedWriteBack);
+
+  void constructorNotification(const detail::code_location &CodeLoc,
+                               void *UserObj, const void *HostObj,
+                               const void *Type, uint32_t Dim,
+                               uint32_t ElemType, size_t Range[3]);
+
+  template <typename propertyT> bool has_property() const noexcept;
+
+  template <typename propertyT> propertyT get_property() const;
+
+  std::vector<pi_native_handle> getNativeVector(backend BackendName) const;
+
+  const std::unique_ptr<SYCLMemObjAllocator> &get_allocator_internal() const;
+
+  void deleteAccProps(const sycl::detail::PropWithDataKind &Kind);
+
+  void addOrReplaceAccessorProperties(const property_list &PropertyList);
+
+  size_t getSize() const;
+
+  std::shared_ptr<detail::buffer_impl> impl;
+};
+
 } // namespace detail
 
 /// Defines a shared array that can be used by kernels in queues.
@@ -59,7 +134,7 @@ template <typename T, int dimensions = 1,
           typename AllocatorT = buffer_allocator<std::remove_const_t<T>>,
           typename __Enabled = typename detail::enable_if_t<(dimensions > 0) &&
                                                             (dimensions <= 3)>>
-class buffer {
+class buffer : public detail::buffer_plain {
   // TODO check is_device_copyable<T>::value after converting sycl::vec into a
   // trivially copyable class.
   static_assert(!std::is_same<T, std::string>::value,
@@ -97,53 +172,55 @@ class buffer {
   buffer(const range<dimensions> &bufferRange,
          const property_list &propList = {},
          const detail::code_location CodeLoc = detail::code_location::current())
-      : Range(bufferRange) {
-    impl = std::make_shared<detail::buffer_impl>(
-        size() * sizeof(T), detail::getNextPowerOfTwo(sizeof(T)), propList,
-        make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, T>>());
-    impl->constructorNotification(CodeLoc, (void *)impl.get(), nullptr,
-                                  (const void *)typeid(T).name(), dimensions,
-                                  sizeof(T), rangeToArray(Range).data());
+      : buffer_plain(bufferRange.size() * sizeof(T),
+                     detail::getNextPowerOfTwo(sizeof(T)), propList,
+                     make_unique_ptr<
+                         detail::SYCLMemObjAllocatorHolder<AllocatorT, T>>()),
+        Range(bufferRange) {
+    buffer_plain::constructorNotification(
+        CodeLoc, (void *)impl.get(), nullptr, (const void *)typeid(T).name(),
+        dimensions, sizeof(T), rangeToArray(Range).data());
   }
 
   buffer(const range<dimensions> &bufferRange, AllocatorT allocator,
          const property_list &propList = {},
          const detail::code_location CodeLoc = detail::code_location::current())
-      : Range(bufferRange) {
-    impl = std::make_shared<detail::buffer_impl>(
-        size() * sizeof(T), detail::getNextPowerOfTwo(sizeof(T)), propList,
-        make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, T>>(
-            allocator));
-    impl->constructorNotification(CodeLoc, (void *)impl.get(), nullptr,
-                                  (const void *)typeid(T).name(), dimensions,
-                                  sizeof(T), rangeToArray(Range).data());
+      : buffer_plain(
+            bufferRange.size() * sizeof(T),
+            detail::getNextPowerOfTwo(sizeof(T)), propList,
+            make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, T>>(
+                allocator)),
+        Range(bufferRange) {
+    buffer_plain::constructorNotification(
+        CodeLoc, (void *)impl.get(), nullptr, (const void *)typeid(T).name(),
+        dimensions, sizeof(T), rangeToArray(Range).data());
   }
 
   buffer(T *hostData, const range<dimensions> &bufferRange,
          const property_list &propList = {},
          const detail::code_location CodeLoc = detail::code_location::current())
-      : Range(bufferRange) {
-    impl = std::make_shared<detail::buffer_impl>(
-        hostData, size() * sizeof(T), detail::getNextPowerOfTwo(sizeof(T)),
-        propList,
-        make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, T>>());
-    impl->constructorNotification(CodeLoc, (void *)impl.get(), hostData,
-                                  (const void *)typeid(T).name(), dimensions,
-                                  sizeof(T), rangeToArray(Range).data());
+      : buffer_plain(hostData, bufferRange.size() * sizeof(T),
+                     detail::getNextPowerOfTwo(sizeof(T)), propList,
+                     make_unique_ptr<
+                         detail::SYCLMemObjAllocatorHolder<AllocatorT, T>>()),
+        Range(bufferRange) {
+    buffer_plain::constructorNotification(
+        CodeLoc, (void *)impl.get(), hostData, (const void *)typeid(T).name(),
+        dimensions, sizeof(T), rangeToArray(Range).data());
   }
 
   buffer(T *hostData, const range<dimensions> &bufferRange,
          AllocatorT allocator, const property_list &propList = {},
          const detail::code_location CodeLoc = detail::code_location::current())
-      : Range(bufferRange) {
-    impl = std::make_shared<detail::buffer_impl>(
-        hostData, size() * sizeof(T), detail::getNextPowerOfTwo(sizeof(T)),
-        propList,
-        make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, T>>(
-            allocator));
-    impl->constructorNotification(CodeLoc, (void *)impl.get(), hostData,
-                                  (const void *)typeid(T).name(), dimensions,
-                                  sizeof(T), rangeToArray(Range).data());
+      : buffer_plain(
+            hostData, bufferRange.size() * sizeof(T),
+            detail::getNextPowerOfTwo(sizeof(T)), propList,
+            make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, T>>(
+                allocator)),
+        Range(bufferRange) {
+    buffer_plain::constructorNotification(
+        CodeLoc, (void *)impl.get(), hostData, (const void *)typeid(T).name(),
+        dimensions, sizeof(T), rangeToArray(Range).data());
   }
 
   template <typename _T = T>
@@ -151,14 +228,14 @@ class buffer {
          const range<dimensions> &bufferRange,
          const property_list &propList = {},
          const detail::code_location CodeLoc = detail::code_location::current())
-      : Range(bufferRange) {
-    impl = std::make_shared<detail::buffer_impl>(
-        hostData, size() * sizeof(T), detail::getNextPowerOfTwo(sizeof(T)),
-        propList,
-        make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, T>>());
-    impl->constructorNotification(CodeLoc, (void *)impl.get(), hostData,
-                                  (const void *)typeid(T).name(), dimensions,
-                                  sizeof(T), rangeToArray(Range).data());
+      : buffer_plain(hostData, bufferRange.size() * sizeof(T),
+                     detail::getNextPowerOfTwo(sizeof(T)), propList,
+                     make_unique_ptr<
+                         detail::SYCLMemObjAllocatorHolder<AllocatorT, T>>()),
+        Range(bufferRange) {
+    buffer_plain::constructorNotification(
+        CodeLoc, (void *)impl.get(), hostData, (const void *)typeid(T).name(),
+        dimensions, sizeof(T), rangeToArray(Range).data());
   }
 
   template <typename _T = T>
@@ -166,77 +243,81 @@ class buffer {
          const range<dimensions> &bufferRange, AllocatorT allocator,
          const property_list &propList = {},
          const detail::code_location CodeLoc = detail::code_location::current())
-      : Range(bufferRange) {
-    impl = std::make_shared<detail::buffer_impl>(
-        hostData, size() * sizeof(T), detail::getNextPowerOfTwo(sizeof(T)),
-        propList,
-        make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, T>>(
-            allocator));
-    impl->constructorNotification(CodeLoc, (void *)impl.get(), hostData,
-                                  (const void *)typeid(T).name(), dimensions,
-                                  sizeof(T), rangeToArray(Range).data());
+      : buffer_plain(
+            bufferRange.size() * sizeof(T),
+            detail::getNextPowerOfTwo(sizeof(T)), propList,
+            make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, T>>(
+                allocator)),
+        Range(bufferRange) {
+    buffer_plain::constructorNotification(
+        CodeLoc, (void *)impl.get(), hostData, (const void *)typeid(T).name(),
+        dimensions, sizeof(T), rangeToArray(Range).data());
   }
 
   buffer(const std::shared_ptr<T> &hostData,
          const range<dimensions> &bufferRange, AllocatorT allocator,
          const property_list &propList = {},
          const detail::code_location CodeLoc = detail::code_location::current())
-      : Range(bufferRange) {
-    impl = std::make_shared<detail::buffer_impl>(
-        hostData, size() * sizeof(T), detail::getNextPowerOfTwo(sizeof(T)),
-        propList,
-        make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, T>>(
-            allocator));
-    impl->constructorNotification(CodeLoc, (void *)impl.get(),
-                                  (void *)hostData.get(),
-                                  (const void *)typeid(T).name(), dimensions,
-                                  sizeof(T), rangeToArray(Range).data());
+      : buffer_plain(
+            hostData, bufferRange.size() * sizeof(T),
+            detail::getNextPowerOfTwo(sizeof(T)), propList,
+            make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, T>>(
+                allocator),
+            std::is_const<T>::value),
+        Range(bufferRange) {
+    buffer_plain::constructorNotification(
+        CodeLoc, (void *)impl.get(), (void *)hostData.get(),
+        (const void *)typeid(T).name(), dimensions, sizeof(T),
+        rangeToArray(Range).data());
   }
 
   buffer(const std::shared_ptr<T[]> &hostData,
          const range<dimensions> &bufferRange, AllocatorT allocator,
          const property_list &propList = {},
          const detail::code_location CodeLoc = detail::code_location::current())
-      : Range(bufferRange) {
-    impl = std::make_shared<detail::buffer_impl>(
-        hostData, size() * sizeof(T), detail::getNextPowerOfTwo(sizeof(T)),
-        propList,
-        make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, T>>(
-            allocator));
-    impl->constructorNotification(CodeLoc, (void *)impl.get(),
-                                  (void *)hostData.get(),
-                                  (const void *)typeid(T).name(), dimensions,
-                                  sizeof(T), rangeToArray(Range).data());
+      : buffer_plain(
+            hostData, bufferRange.size() * sizeof(T),
+            detail::getNextPowerOfTwo(sizeof(T)), propList,
+            make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, T>>(
+                allocator),
+            std::is_const<T>::value),
+        Range(bufferRange) {
+    buffer_plain::constructorNotification(
+        CodeLoc, (void *)impl.get(), (void *)hostData.get(),
+        (const void *)typeid(T).name(), dimensions, sizeof(T),
+        rangeToArray(Range).data());
   }
 
   buffer(const std::shared_ptr<T> &hostData,
          const range<dimensions> &bufferRange,
          const property_list &propList = {},
          const detail::code_location CodeLoc = detail::code_location::current())
-      : Range(bufferRange) {
-    impl = std::make_shared<detail::buffer_impl>(
-        hostData, size() * sizeof(T), detail::getNextPowerOfTwo(sizeof(T)),
-        propList,
-        make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, T>>());
-    impl->constructorNotification(CodeLoc, (void *)impl.get(),
-                                  (void *)hostData.get(),
-                                  (const void *)typeid(T).name(), dimensions,
-                                  sizeof(T), rangeToArray(Range).data());
+      : buffer_plain(
+            hostData, bufferRange.size() * sizeof(T),
+            detail::getNextPowerOfTwo(sizeof(T)), propList,
+            make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, T>>(),
+            std::is_const<T>::value),
+        Range(bufferRange) {
+    buffer_plain::constructorNotification(
+        CodeLoc, (void *)impl.get(), (void *)hostData.get(),
+        (const void *)typeid(T).name(), dimensions, sizeof(T),
+        rangeToArray(Range).data());
   }
 
   buffer(const std::shared_ptr<T[]> &hostData,
          const range<dimensions> &bufferRange,
          const property_list &propList = {},
          const detail::code_location CodeLoc = detail::code_location::current())
-      : Range(bufferRange) {
-    impl = std::make_shared<detail::buffer_impl>(
-        hostData, size() * sizeof(T), detail::getNextPowerOfTwo(sizeof(T)),
-        propList,
-        make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, T>>());
-    impl->constructorNotification(CodeLoc, (void *)impl.get(),
-                                  (void *)hostData.get(),
-                                  (const void *)typeid(T).name(), dimensions,
-                                  sizeof(T), rangeToArray(Range).data());
+      : buffer_plain(
+            hostData, bufferRange.size() * sizeof(T),
+            detail::getNextPowerOfTwo(sizeof(T)), propList,
+            make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, T>>(),
+            std::is_const<T>::value),
+        Range(bufferRange) {
+    buffer_plain::constructorNotification(
+        CodeLoc, (void *)impl.get(), (void *)hostData.get(),
+        (const void *)typeid(T).name(), dimensions, sizeof(T),
+        rangeToArray(Range).data());
   }
 
   template <class InputIterator, int N = dimensions,
@@ -245,16 +326,31 @@ class buffer {
   buffer(InputIterator first, InputIterator last, AllocatorT allocator,
          const property_list &propList = {},
          const detail::code_location CodeLoc = detail::code_location::current())
-      : Range(range<1>(std::distance(first, last))) {
-    impl = std::make_shared<detail::buffer_impl>(
-        first, last, size() * sizeof(T), detail::getNextPowerOfTwo(sizeof(T)),
-        propList,
-        make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, T>>(
-            allocator));
+      : buffer_plain(
+            // The functor which will be used to initialize the data
+            [first, last](void *ToPtr) {
+              // We need to cast MUserPtr to pointer to the iteration type to
+              // get correct offset in std::copy when it will increment
+              // destination pointer.
+              using IteratorValueType =
+                  detail::iterator_value_type_t<InputIterator>;
+              using IteratorNonConstValueType =
+                  detail::remove_const_t<IteratorValueType>;
+              using IteratorPointerToNonConstValueType =
+                  detail::add_pointer_t<IteratorNonConstValueType>;
+              std::copy(first, last,
+                        static_cast<IteratorPointerToNonConstValueType>(ToPtr));
+            },
+            std::distance(first, last) * sizeof(T),
+            detail::getNextPowerOfTwo(sizeof(T)), propList,
+            make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, T>>(
+                allocator),
+            detail::iterator_to_const_type_t<InputIterator>::value),
+        Range(range<1>(std::distance(first, last))) {
     size_t r[3] = {Range[0], 0, 0};
-    impl->constructorNotification(CodeLoc, (void *)impl.get(), &first,
-                                  (const void *)typeid(T).name(), dimensions,
-                                  sizeof(T), r);
+    buffer_plain::constructorNotification(CodeLoc, (void *)impl.get(), &first,
+                                          (const void *)typeid(T).name(),
+                                          dimensions, sizeof(T), r);
   }
 
   template <class InputIterator, int N = dimensions,
@@ -263,15 +359,30 @@ class buffer {
   buffer(InputIterator first, InputIterator last,
          const property_list &propList = {},
          const detail::code_location CodeLoc = detail::code_location::current())
-      : Range(range<1>(std::distance(first, last))) {
-    impl = std::make_shared<detail::buffer_impl>(
-        first, last, size() * sizeof(T), detail::getNextPowerOfTwo(sizeof(T)),
-        propList,
-        make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, T>>());
+      : buffer_plain(
+            // The functor which will be used to initialize the data
+            [first, last](void *ToPtr) {
+              // We need to cast MUserPtr to pointer to the iteration type to
+              // get correct offset in std::copy when it will increment
+              // destination pointer.
+              using IteratorValueType =
+                  detail::iterator_value_type_t<InputIterator>;
+              using IteratorNonConstValueType =
+                  detail::remove_const_t<IteratorValueType>;
+              using IteratorPointerToNonConstValueType =
+                  detail::add_pointer_t<IteratorNonConstValueType>;
+              std::copy(first, last,
+                        static_cast<IteratorPointerToNonConstValueType>(ToPtr));
+            },
+            std::distance(first, last) * sizeof(T),
+            detail::getNextPowerOfTwo(sizeof(T)), propList,
+            make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, T>>(),
+            detail::iterator_to_const_type_t<InputIterator>::value),
+        Range(range<1>(std::distance(first, last))) {
     size_t r[3] = {Range[0], 0, 0};
-    impl->constructorNotification(CodeLoc, (void *)impl.get(), &first,
-                                  (const void *)typeid(T).name(), dimensions,
-                                  sizeof(T), r);
+    buffer_plain::constructorNotification(CodeLoc, (void *)impl.get(), &first,
+                                          (const void *)typeid(T).name(),
+                                          dimensions, sizeof(T), r);
   }
 
   // This constructor is a prototype for a future SYCL specification
@@ -281,16 +392,16 @@ class buffer {
   buffer(Container &container, AllocatorT allocator,
          const property_list &propList = {},
          const detail::code_location CodeLoc = detail::code_location::current())
-      : Range(range<1>(container.size())) {
-    impl = std::make_shared<detail::buffer_impl>(
-        container.data(), size() * sizeof(T),
-        detail::getNextPowerOfTwo(sizeof(T)), propList,
-        make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, T>>(
-            allocator));
+      : buffer_plain(
+            container.data(), container.size() * sizeof(T),
+            detail::getNextPowerOfTwo(sizeof(T)), propList,
+            make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, T>>(
+                allocator)),
+        Range(range<1>(container.size())) {
     size_t r[3] = {Range[0], 0, 0};
-    impl->constructorNotification(CodeLoc, (void *)impl.get(), container.data(),
-                                  (const void *)typeid(T).name(), dimensions,
-                                  sizeof(T), r);
+    buffer_plain::constructorNotification(
+        CodeLoc, (void *)impl.get(), container.data(),
+        (const void *)typeid(T).name(), dimensions, sizeof(T), r);
   }
 
   // This constructor is a prototype for a future SYCL specification
@@ -304,12 +415,12 @@ class buffer {
   buffer(buffer<T, dimensions, AllocatorT> &b, const id<dimensions> &baseIndex,
          const range<dimensions> &subRange,
          const detail::code_location CodeLoc = detail::code_location::current())
-      : impl(b.impl), Range(subRange),
+      : buffer_plain(b.impl), Range(subRange),
         OffsetInBytes(getOffsetInBytes<T>(baseIndex, b.Range)),
         IsSubBuffer(true) {
-    impl->constructorNotification(CodeLoc, (void *)impl.get(), impl.get(),
-                                  (const void *)typeid(T).name(), dimensions,
-                                  sizeof(T), rangeToArray(Range).data());
+    buffer_plain::constructorNotification(
+        CodeLoc, (void *)impl.get(), impl.get(), (const void *)typeid(T).name(),
+        dimensions, sizeof(T), rangeToArray(Range).data());
 
     if (b.is_sub_buffer())
       throw sycl::invalid_object_error(
@@ -324,40 +435,22 @@ class buffer {
           PI_ERROR_INVALID_VALUE);
   }
 
-#ifdef __SYCL_INTERNAL_API
-  template <int N = dimensions, typename = EnableIfOneDimension<N>>
-  buffer(cl_mem MemObject, const context &SyclContext,
-         event AvailableEvent = {},
-         const detail::code_location CodeLoc = detail::code_location::current())
-      : Range{0} {
-
-    impl = std::make_shared<detail::buffer_impl>(
-        detail::pi::cast<pi_native_handle>(MemObject), SyclContext,
-        make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, T>>(),
-        /* OwnNativeHandle */ true, AvailableEvent);
-    Range[0] = impl->getSize() / sizeof(T);
-    impl->constructorNotification(CodeLoc, (void *)impl.get(), &MemObject,
-                                  (const void *)typeid(T).name(), dimensions,
-                                  sizeof(T), rangeToArray(Range).data());
-  }
-#endif
-
   buffer(const buffer &rhs,
          const detail::code_location CodeLoc = detail::code_location::current())
-      : impl(rhs.impl), Range(rhs.Range), OffsetInBytes(rhs.OffsetInBytes),
-        IsSubBuffer(rhs.IsSubBuffer) {
-    impl->constructorNotification(CodeLoc, (void *)impl.get(), impl.get(),
-                                  (const void *)typeid(T).name(), dimensions,
-                                  sizeof(T), rangeToArray(Range).data());
+      : buffer_plain(rhs.impl), Range(rhs.Range),
+        OffsetInBytes(rhs.OffsetInBytes), IsSubBuffer(rhs.IsSubBuffer) {
+    buffer_plain::constructorNotification(
+        CodeLoc, (void *)impl.get(), impl.get(), (const void *)typeid(T).name(),
+        dimensions, sizeof(T), rangeToArray(Range).data());
   }
 
   buffer(buffer &&rhs,
          const detail::code_location CodeLoc = detail::code_location::current())
-      : impl(std::move(rhs.impl)), Range(rhs.Range),
+      : buffer_plain(std::move(rhs.impl)), Range(rhs.Range),
         OffsetInBytes(rhs.OffsetInBytes), IsSubBuffer(rhs.IsSubBuffer) {
-    impl->constructorNotification(CodeLoc, (void *)impl.get(), impl.get(),
-                                  (const void *)typeid(T).name(), dimensions,
-                                  sizeof(T), rangeToArray(Range).data());
+    buffer_plain::constructorNotification(
+        CodeLoc, (void *)impl.get(), impl.get(), (const void *)typeid(T).name(),
+        dimensions, sizeof(T), rangeToArray(Range).data());
   }
 
   buffer &operator=(const buffer &rhs) = default;
@@ -386,7 +479,8 @@ class buffer {
   size_t byte_size() const noexcept { return size() * sizeof(T); }
 
   AllocatorT get_allocator() const {
-    return impl->template get_allocator<AllocatorT>();
+    return buffer_plain::get_allocator_internal()
+        ->template getAllocator<AllocatorT>();
   }
 
   template <access::mode Mode, access::target Target = access::target::device>
@@ -468,10 +562,65 @@ class buffer {
 
   template <typename Destination = std::nullptr_t>
   void set_final_data(Destination finalData = nullptr) {
-    impl->set_final_data(finalData);
+    this->set_final_data_internal(finalData);
+  }
+
+  void set_final_data_internal(std::nullptr_t) {
+    buffer_plain::set_final_data_internal();
   }
 
-  void set_write_back(bool flag = true) { impl->set_write_back(flag); }
+  template <template <typename WeakT> class WeakPtrT, typename WeakT>
+  detail::enable_if_t<
+      std::is_convertible<WeakPtrT<WeakT>, std::weak_ptr<WeakT>>::value>
+  set_final_data_internal(WeakPtrT<WeakT> FinalData) {
+    std::weak_ptr<WeakT> TempFinalData(FinalData);
+    this->set_final_data_internal(TempFinalData);
+  }
+
+  template <typename WeakT>
+  void set_final_data_internal(std::weak_ptr<WeakT> FinalData) {
+    buffer_plain::set_final_data_internal(
+        [FinalData](const std::function<void(void *const Ptr)> &F) {
+          if (std::shared_ptr<WeakT> LockedFinalData = FinalData.lock())
+            F(LockedFinalData.get());
+        });
+  }
+
+  template <typename Destination>
+  detail::EnableIfOutputPointerT<Destination>
+  set_final_data_internal(Destination FinalData) {
+    if (!FinalData)
+      buffer_plain::set_final_data_internal();
+    else
+      buffer_plain::set_final_data_internal(
+          [FinalData](const std::function<void(void *const Ptr)> &F) {
+            F(FinalData);
+          });
+  }
+
+  template <typename Destination>
+  detail::EnableIfOutputIteratorT<Destination>
+  set_final_data_internal(Destination FinalData) {
+    const size_t Size = size();
+    buffer_plain::set_final_data_internal(
+        [FinalData, Size](const std::function<void(void *const Ptr)> &F) {
+          using DestinationValueT = detail::iterator_value_type_t<Destination>;
+          // TODO if Destination is ContiguousIterator then don't create
+          // ContiguousStorage. updateHostMemory works only with pointer to
+          // continuous data.
+          std::unique_ptr<DestinationValueT[]> ContiguousStorage(
+              new DestinationValueT[Size]);
+          F(ContiguousStorage.get());
+          std::copy(ContiguousStorage.get(), ContiguousStorage.get() + Size,
+                    FinalData);
+        });
+  }
+
+  void set_final_data(std::nullptr_t) {
+    buffer_plain::set_final_data_internal();
+  }
+
+  void set_write_back(bool flag = true) { buffer_plain::set_write_back(flag); }
 
   bool is_sub_buffer() const { return IsSubBuffer; }
 
@@ -524,11 +673,11 @@ class buffer {
   }
 
   template <typename propertyT> bool has_property() const noexcept {
-    return impl->template has_property<propertyT>();
+    return buffer_plain::template has_property<propertyT>();
   }
 
   template <typename propertyT> propertyT get_property() const {
-    return impl->template get_property<propertyT>();
+    return buffer_plain::template get_property<propertyT>();
   }
 
 protected:
@@ -543,7 +692,6 @@ class buffer {
   }
 
 private:
-  std::shared_ptr<detail::buffer_impl> impl;
   template <class Obj>
   friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
   template <typename A, int dims, typename C, typename Enable>
@@ -565,28 +713,36 @@ class buffer {
   buffer(pi_native_handle MemObject, const context &SyclContext,
          bool OwnNativeHandle, event AvailableEvent = {},
          const detail::code_location CodeLoc = detail::code_location::current())
-      : Range{0} {
+      : buffer_plain(
+            MemObject, SyclContext,
+            make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, T>>(),
+            OwnNativeHandle, std::move(AvailableEvent)),
+        Range{0} {
+
+    Range[0] = buffer_plain::getSize() / sizeof(T);
+    buffer_plain::constructorNotification(
+        CodeLoc, (void *)impl.get(), &MemObject, (const void *)typeid(T).name(),
+        dimensions, sizeof(T), rangeToArray(Range).data());
+  }
+
+  void addOrReplaceAccessorProperties(const property_list &PropertyList) {
+    buffer_plain::addOrReplaceAccessorProperties(PropertyList);
+  }
 
-    impl = std::make_shared<detail::buffer_impl>(
-        MemObject, SyclContext,
-        make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, T>>(),
-        OwnNativeHandle, AvailableEvent);
-    Range[0] = impl->getSize() / sizeof(T);
-    impl->constructorNotification(CodeLoc, (void *)impl.get(), &MemObject,
-                                  (const void *)typeid(T).name(), dimensions,
-                                  sizeof(T), rangeToArray(Range).data());
+  void deleteAccProps(const sycl::detail::PropWithDataKind &Kind) {
+    buffer_plain::deleteAccProps(Kind);
   }
 
   // Reinterpret contructor
-  buffer(std::shared_ptr<detail::buffer_impl> Impl,
+  buffer(const std::shared_ptr<detail::buffer_impl> &Impl,
          range<dimensions> reinterpretRange, size_t reinterpretOffset,
          bool isSubBuffer,
          const detail::code_location CodeLoc = detail::code_location::current())
-      : impl(Impl), Range(reinterpretRange), OffsetInBytes(reinterpretOffset),
-        IsSubBuffer(isSubBuffer) {
-    impl->constructorNotification(CodeLoc, (void *)impl.get(), Impl.get(),
-                                  (const void *)typeid(T).name(), dimensions,
-                                  sizeof(T), rangeToArray(Range).data());
+      : buffer_plain(Impl), Range(reinterpretRange),
+        OffsetInBytes(reinterpretOffset), IsSubBuffer(isSubBuffer) {
+    buffer_plain::constructorNotification(
+        CodeLoc, (void *)impl.get(), Impl.get(), (const void *)typeid(T).name(),
+        dimensions, sizeof(T), rangeToArray(Range).data());
   }
 
   template <typename Type, int N>
@@ -639,7 +795,7 @@ class buffer {
   template <backend BackendName>
   backend_return_t<BackendName, buffer<T, dimensions, AllocatorT>>
   getNative() const {
-    auto NativeHandles = impl->getNativeVector(BackendName);
+    auto NativeHandles = buffer_plain::getNativeVector(BackendName);
     return detail::BufferInterop<BackendName, T, dimensions,
                                  AllocatorT>::GetNativeObjs(NativeHandles);
   }
diff --git a/sycl/include/sycl/detail/accessor_impl.hpp b/sycl/include/sycl/detail/accessor_impl.hpp
index 8dec9653035e1..47af8f3b10a81 100644
--- a/sycl/include/sycl/detail/accessor_impl.hpp
+++ b/sycl/include/sycl/detail/accessor_impl.hpp
@@ -10,7 +10,6 @@
 
 #include <sycl/access/access.hpp>
 #include <sycl/detail/export.hpp>
-#include <sycl/detail/sycl_mem_obj_i.hpp>
 #include <sycl/id.hpp>
 #include <sycl/range.hpp>
 #include <sycl/stl.hpp>
@@ -34,6 +33,8 @@ class AccessorPrivateProxy;
 
 namespace detail {
 
+class SYCLMemObjI;
+
 class Command;
 
 // The class describes a requirement to access a SYCL memory object such as
@@ -77,11 +78,12 @@ template <int Dims> class LocalAccessorBaseDevice {
 class __SYCL_EXPORT AccessorImplHost {
 public:
   AccessorImplHost(id<3> Offset, range<3> AccessRange, range<3> MemoryRange,
-                   access::mode AccessMode, detail::SYCLMemObjI *SYCLMemObject,
-                   int Dims, int ElemSize, int OffsetInBytes = 0,
+                   access::mode AccessMode, void *SYCLMemObject, int Dims,
+                   int ElemSize, int OffsetInBytes = 0,
                    bool IsSubBuffer = false, bool IsESIMDAcc = false)
       : MOffset(Offset), MAccessRange(AccessRange), MMemoryRange(MemoryRange),
-        MAccessMode(AccessMode), MSYCLMemObj(SYCLMemObject), MDims(Dims),
+        MAccessMode(AccessMode),
+        MSYCLMemObj((detail::SYCLMemObjI *)SYCLMemObject), MDims(Dims),
         MElemSize(ElemSize), MOffsetInBytes(OffsetInBytes),
         MIsSubBuffer(IsSubBuffer), MIsESIMDAcc(IsESIMDAcc) {}
 
@@ -135,12 +137,13 @@ using AccessorImplPtr = std::shared_ptr<AccessorImplHost>;
 class AccessorBaseHost {
 public:
   AccessorBaseHost(id<3> Offset, range<3> AccessRange, range<3> MemoryRange,
-                   access::mode AccessMode, detail::SYCLMemObjI *SYCLMemObject,
-                   int Dims, int ElemSize, int OffsetInBytes = 0,
+                   access::mode AccessMode, void *SYCLMemObject, int Dims,
+                   int ElemSize, int OffsetInBytes = 0,
                    bool IsSubBuffer = false) {
-    impl = std::shared_ptr<AccessorImplHost>(new AccessorImplHost(
-        Offset, AccessRange, MemoryRange, AccessMode, SYCLMemObject, Dims,
-        ElemSize, OffsetInBytes, IsSubBuffer));
+    impl = std::shared_ptr<AccessorImplHost>(
+        new AccessorImplHost(Offset, AccessRange, MemoryRange, AccessMode,
+                             (detail::SYCLMemObjI *)SYCLMemObject, Dims,
+                             ElemSize, OffsetInBytes, IsSubBuffer));
   }
 
 protected:
diff --git a/sycl/include/sycl/detail/cg.hpp b/sycl/include/sycl/detail/cg.hpp
index 93ab48e6995bc..829e6411a5c8d 100644
--- a/sycl/include/sycl/detail/cg.hpp
+++ b/sycl/include/sycl/detail/cg.hpp
@@ -39,6 +39,9 @@ class queue;
 
 namespace detail {
 
+class event_impl;
+using EventImplPtr = std::shared_ptr<event_impl>;
+
 // Periodically there is a need to extend handler and CG classes to hold more
 // data(members) than it has now. But any modification of the layout of those
 // classes is an ABI break. To have an ability to have more data the following
diff --git a/sycl/include/sycl/detail/common.hpp b/sycl/include/sycl/detail/common.hpp
index 72249f860067c..d66dc2cf273ef 100644
--- a/sycl/include/sycl/detail/common.hpp
+++ b/sycl/include/sycl/detail/common.hpp
@@ -25,6 +25,18 @@ namespace sycl {
 __SYCL_INLINE_VER_NAMESPACE(_V1) {
 namespace detail {
 
+// The check for output iterator is commented out as it blocks set_final_data
+// with void * argument to be used.
+// TODO: Align these checks with the SYCL specification when the behaviour
+// with void * is clarified.
+template <typename DataT>
+using EnableIfOutputPointerT = detail::enable_if_t<
+    /*is_output_iterator<DataT>::value &&*/ std::is_pointer<DataT>::value>;
+
+template <typename DataT>
+using EnableIfOutputIteratorT = detail::enable_if_t<
+    /*is_output_iterator<DataT>::value &&*/ !std::is_pointer<DataT>::value>;
+
 #if !defined(NDEBUG) && (_MSC_VER > 1929 || __has_builtin(__builtin_FILE))
 #define __CODELOC_FILE_NAME __builtin_FILE()
 #else
@@ -378,6 +390,21 @@ template <typename T> struct InlineVariableHelper {
 };
 
 template <typename T> constexpr T InlineVariableHelper<T>::value;
+
+// The function extends or truncates number of dimensions of objects of id
+// or ranges classes. When extending the new values are filled with
+// DefaultValue, truncation just removes extra values.
+template <int NewDim, int DefaultValue, template <int> class T, int OldDim>
+static T<NewDim> convertToArrayOfN(T<OldDim> OldObj) {
+  T<NewDim> NewObj = detail::InitializedVal<NewDim, T>::template get<0>();
+  const int CopyDims = NewDim > OldDim ? OldDim : NewDim;
+  for (int I = 0; I < CopyDims; ++I)
+    NewObj[I] = OldObj[I];
+  for (int I = CopyDims; I < NewDim; ++I)
+    NewObj[I] = DefaultValue;
+  return NewObj;
+}
+
 } // namespace detail
 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
 } // namespace sycl
diff --git a/sycl/include/sycl/image.hpp b/sycl/include/sycl/image.hpp
index 755e084869b63..682291eb06f54 100644
--- a/sycl/include/sycl/image.hpp
+++ b/sycl/include/sycl/image.hpp
@@ -8,18 +8,22 @@
 
 #pragma once
 
-#include <cstddef>
+#include <sycl/detail/aligned_allocator.hpp>
 #include <sycl/detail/common.hpp>
 #include <sycl/detail/generic_type_traits.hpp>
-#include <sycl/detail/image_impl.hpp>
+#include <sycl/detail/sycl_mem_obj_allocator.hpp>
 #include <sycl/event.hpp>
 #include <sycl/ext/oneapi/accessor_property_list.hpp>
 #include <sycl/stl.hpp>
 #include <sycl/types.hpp>
 
+#include <cstddef>
+
 namespace sycl {
 __SYCL_INLINE_VER_NAMESPACE(_V1) {
 
+class handler;
+
 enum class image_channel_order : unsigned int {
   a = 0,
   r = 1,
@@ -60,6 +64,110 @@ using byte = unsigned char;
 
 using image_allocator = detail::aligned_allocator<byte>;
 
+namespace detail {
+
+class image_impl;
+
+// validImageDataT: cl_int4, cl_uint4, cl_float4, cl_half4
+template <typename T>
+using is_validImageDataT = typename detail::is_contained<
+    T, type_list<cl_int4, cl_uint4, cl_float4, cl_half4>>::type;
+
+template <typename DataT>
+using EnableIfImgAccDataT =
+    typename detail::enable_if_t<is_validImageDataT<DataT>::value, DataT>;
+
+// The non-template base for the sycl::image class
+class __SYCL_EXPORT image_plain {
+protected:
+  image_plain(image_channel_order Order, image_channel_type Type,
+              const range<3> &Range,
+              std::unique_ptr<SYCLMemObjAllocator> Allocator,
+              uint8_t Dimensions, const property_list &PropList);
+
+  image_plain(image_channel_order Order, image_channel_type Type,
+              const range<3> &Range, const range<2> &Pitch,
+              std::unique_ptr<SYCLMemObjAllocator> Allocator,
+              uint8_t Dimensions, const property_list &PropList);
+
+  image_plain(void *HostPointer, image_channel_order Order,
+              image_channel_type Type, const range<3> &Range,
+              std::unique_ptr<SYCLMemObjAllocator> Allocator,
+              uint8_t Dimensions, const property_list &PropList);
+
+  image_plain(const void *HostPointer, image_channel_order Order,
+              image_channel_type Type, const range<3> &Range,
+              std::unique_ptr<SYCLMemObjAllocator> Allocator,
+              uint8_t Dimensions, const property_list &PropList);
+
+  image_plain(void *HostPointer, image_channel_order Order,
+              image_channel_type Type, const range<3> &Range,
+              const range<2> &Pitch,
+              std::unique_ptr<SYCLMemObjAllocator> Allocator,
+              uint8_t Dimensions, const property_list &PropList);
+
+  image_plain(const std::shared_ptr<const void> &HostPointer,
+              image_channel_order Order, image_channel_type Type,
+              const range<3> &Range,
+              std::unique_ptr<SYCLMemObjAllocator> Allocator,
+              uint8_t Dimensions, const property_list &PropList,
+              bool IsConstPtr);
+
+  image_plain(const std::shared_ptr<const void> &HostPointer,
+              image_channel_order Order, image_channel_type Type,
+              const range<3> &Range, const range<2> &Pitch,
+              std::unique_ptr<SYCLMemObjAllocator> Allocator,
+              uint8_t Dimensions, const property_list &PropList,
+              bool IsConstPtr);
+
+#ifdef __SYCL_INTERNAL_API
+  image_plain(cl_mem ClMemObject, const context &SyclContext,
+              event AvailableEvent,
+              std::unique_ptr<SYCLMemObjAllocator> Allocator,
+              uint8_t Dimensions);
+#endif
+
+  template <typename propertyT> bool has_property() const noexcept;
+
+  template <typename propertyT> propertyT get_property() const;
+
+  range<3> get_range() const;
+
+  range<2> get_pitch() const;
+
+  size_t get_size() const;
+
+  size_t get_count() const;
+
+  void set_final_data_internal();
+
+  void set_final_data_internal(
+      const std::function<void(const std::function<void(void *const Ptr)> &)>
+          &FinalDataFunc);
+
+  void set_write_back(bool flag);
+
+  const std::unique_ptr<SYCLMemObjAllocator> &get_allocator_internal() const;
+
+  size_t getElementSize() const;
+
+  size_t getRowPitch() const;
+
+  size_t getSlicePitch() const;
+
+  image_channel_order getChannelOrder() const;
+
+  image_channel_type getChannelType() const;
+
+  std::shared_ptr<detail::image_impl> impl;
+};
+
+template <typename DataT, int Dims, access::mode AccMode,
+          access::target AccTarget, access::placeholder IsPlaceholder>
+class image_accessor;
+
+} // namespace detail
+
 /// Defines a shared image data.
 ///
 /// Images can be 1-, 2-, and 3-dimensional. They have to be accessed using the
@@ -70,164 +178,162 @@ using image_allocator = detail::aligned_allocator<byte>;
 ///
 /// \ingroup sycl_api
 template <int Dimensions = 1, typename AllocatorT = sycl::image_allocator>
-class image {
+class image : public detail::image_plain {
+
 public:
   image(image_channel_order Order, image_channel_type Type,
-        const range<Dimensions> &Range, const property_list &PropList = {}) {
-    impl = std::make_shared<detail::image_impl<Dimensions>>(
-        Order, Type, Range,
-        make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, byte>>(),
-        PropList);
-  }
+        const range<Dimensions> &Range, const property_list &PropList = {})
+      : image_plain(Order, Type, detail::convertToArrayOfN<3, 1>(Range),
+                    make_unique_ptr<
+                        detail::SYCLMemObjAllocatorHolder<AllocatorT, byte>>(),
+                    Dimensions, PropList) {}
 
   image(image_channel_order Order, image_channel_type Type,
         const range<Dimensions> &Range, AllocatorT Allocator,
-        const property_list &PropList = {}) {
-    impl = std::make_shared<detail::image_impl<Dimensions>>(
-        Order, Type, Range,
-        make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, byte>>(
-            Allocator),
-        PropList);
-  }
+        const property_list &PropList = {})
+      : image_plain(
+            Order, Type, detail::convertToArrayOfN<3, 1>(Range),
+            make_unique_ptr<
+                detail::SYCLMemObjAllocatorHolder<AllocatorT, byte>>(Allocator),
+            Dimensions, PropList) {}
 
   /* Available only when: dimensions >1 */
   template <bool B = (Dimensions > 1)>
   image(image_channel_order Order, image_channel_type Type,
         const range<Dimensions> &Range,
         const typename detail::enable_if_t<B, range<Dimensions - 1>> &Pitch,
-        const property_list &PropList = {}) {
-    impl = std::make_shared<detail::image_impl<Dimensions>>(
-        Order, Type, Range, Pitch,
-        make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, byte>>(),
-        PropList);
-  }
+        const property_list &PropList = {})
+      : image_plain(Order, Type, detail::convertToArrayOfN<3, 1>(Range),
+                    detail::convertToArrayOfN<2, 0>(Pitch),
+                    make_unique_ptr<
+                        detail::SYCLMemObjAllocatorHolder<AllocatorT, byte>>(),
+                    Dimensions, PropList) {}
 
   /* Available only when: dimensions >1 */
   template <bool B = (Dimensions > 1)>
   image(image_channel_order Order, image_channel_type Type,
         const range<Dimensions> &Range,
         const typename detail::enable_if_t<B, range<Dimensions - 1>> &Pitch,
-        AllocatorT Allocator, const property_list &PropList = {}) {
-    impl = std::make_shared<detail::image_impl<Dimensions>>(
-        Order, Type, Range, Pitch,
-        make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, byte>>(
-            Allocator),
-        PropList);
-  }
+        AllocatorT Allocator, const property_list &PropList = {})
+      : image_plain(
+            Order, Type, detail::convertToArrayOfN<3, 1>(Range),
+            detail::convertToArrayOfN<2, 0>(Pitch),
+            make_unique_ptr<
+                detail::SYCLMemObjAllocatorHolder<AllocatorT, byte>>(Allocator),
+            Dimensions, PropList) {}
 
   image(void *HostPointer, image_channel_order Order, image_channel_type Type,
-        const range<Dimensions> &Range, const property_list &PropList = {}) {
-    impl = std::make_shared<detail::image_impl<Dimensions>>(
-        HostPointer, Order, Type, Range,
-        make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, byte>>(),
-        PropList);
-  }
+        const range<Dimensions> &Range, const property_list &PropList = {})
+      : image_plain(HostPointer, Order, Type,
+                    detail::convertToArrayOfN<3, 1>(Range),
+                    make_unique_ptr<
+                        detail::SYCLMemObjAllocatorHolder<AllocatorT, byte>>(),
+                    Dimensions, PropList) {}
 
   image(void *HostPointer, image_channel_order Order, image_channel_type Type,
         const range<Dimensions> &Range, AllocatorT Allocator,
-        const property_list &PropList = {}) {
-    impl = std::make_shared<detail::image_impl<Dimensions>>(
-        HostPointer, Order, Type, Range,
-        make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, byte>>(
-            Allocator),
-        PropList);
-  }
+        const property_list &PropList = {})
+      : image_plain(
+            HostPointer, Order, Type, detail::convertToArrayOfN<3, 1>(Range),
+            make_unique_ptr<
+                detail::SYCLMemObjAllocatorHolder<AllocatorT, byte>>(Allocator),
+            Dimensions, PropList) {}
 
   image(const void *HostPointer, image_channel_order Order,
         image_channel_type Type, const range<Dimensions> &Range,
-        const property_list &PropList = {}) {
-    impl = std::make_shared<detail::image_impl<Dimensions>>(
-        HostPointer, Order, Type, Range,
-        make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, byte>>(),
-        PropList);
-  }
+        const property_list &PropList = {})
+      : image_plain(HostPointer, Order, Type,
+                    detail::convertToArrayOfN<3, 1>(Range),
+                    make_unique_ptr<
+                        detail::SYCLMemObjAllocatorHolder<AllocatorT, byte>>(),
+                    Dimensions, PropList) {}
 
   image(const void *HostPointer, image_channel_order Order,
         image_channel_type Type, const range<Dimensions> &Range,
-        AllocatorT Allocator, const property_list &PropList = {}) {
-    impl = std::make_shared<detail::image_impl<Dimensions>>(
-        HostPointer, Order, Type, Range,
-        make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, byte>>(
-            Allocator),
-        PropList);
-  }
+        AllocatorT Allocator, const property_list &PropList = {})
+      : image_plain(
+            HostPointer, Order, Type, detail::convertToArrayOfN<3, 1>(Range),
+            make_unique_ptr<
+                detail::SYCLMemObjAllocatorHolder<AllocatorT, byte>>(Allocator),
+            Dimensions, PropList) {}
 
   /* Available only when: dimensions >1 */
   template <bool B = (Dimensions > 1)>
   image(void *HostPointer, image_channel_order Order, image_channel_type Type,
         const range<Dimensions> &Range,
         const typename detail::enable_if_t<B, range<Dimensions - 1>> &Pitch,
-        const property_list &PropList = {}) {
-    impl = std::make_shared<detail::image_impl<Dimensions>>(
-        HostPointer, Order, Type, Range, Pitch,
-        make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, byte>>(),
-        PropList);
-  }
+        const property_list &PropList = {})
+      : image_plain(HostPointer, Order, Type,
+                    detail::convertToArrayOfN<3, 1>(Range),
+                    detail::convertToArrayOfN<2, 0>(Pitch),
+                    make_unique_ptr<
+                        detail::SYCLMemObjAllocatorHolder<AllocatorT, byte>>(),
+                    Dimensions, PropList) {}
 
   /* Available only when: dimensions >1 */
   template <bool B = (Dimensions > 1)>
   image(void *HostPointer, image_channel_order Order, image_channel_type Type,
         const range<Dimensions> &Range,
         const typename detail::enable_if_t<B, range<Dimensions - 1>> &Pitch,
-        AllocatorT Allocator, const property_list &PropList = {}) {
-    impl = std::make_shared<detail::image_impl<Dimensions>>(
-        HostPointer, Order, Type, Range, Pitch,
-        make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, byte>>(
-            Allocator),
-        PropList);
-  }
+        AllocatorT Allocator, const property_list &PropList = {})
+      : image_plain(
+            HostPointer, Order, Type, detail::convertToArrayOfN<3, 1>(Range),
+            detail::convertToArrayOfN<2, 0>(Pitch),
+            make_unique_ptr<
+                detail::SYCLMemObjAllocatorHolder<AllocatorT, byte>>(Allocator),
+            Dimensions, PropList) {}
 
   image(std::shared_ptr<void> &HostPointer, image_channel_order Order,
         image_channel_type Type, const range<Dimensions> &Range,
-        const property_list &PropList = {}) {
-    impl = std::make_shared<detail::image_impl<Dimensions>>(
-        HostPointer, Order, Type, Range,
-        make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, byte>>(),
-        PropList);
-  }
+        const property_list &PropList = {})
+      : image_plain(HostPointer, Order, Type,
+                    detail::convertToArrayOfN<3, 1>(Range),
+                    make_unique_ptr<
+                        detail::SYCLMemObjAllocatorHolder<AllocatorT, byte>>(),
+                    Dimensions, PropList, /*IsConstPtr*/ false) {}
 
   image(std::shared_ptr<void> &HostPointer, image_channel_order Order,
         image_channel_type Type, const range<Dimensions> &Range,
-        AllocatorT Allocator, const property_list &PropList = {}) {
-    impl = std::make_shared<detail::image_impl<Dimensions>>(
-        HostPointer, Order, Type, Range,
-        make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, byte>>(
-            Allocator),
-        PropList);
-  }
+        AllocatorT Allocator, const property_list &PropList = {})
+      : image_plain(
+            HostPointer, Order, Type, detail::convertToArrayOfN<3, 1>(Range),
+            make_unique_ptr<
+                detail::SYCLMemObjAllocatorHolder<AllocatorT, byte>>(Allocator),
+            Dimensions, PropList, /*IsConstPtr*/ false) {}
 
   /* Available only when: dimensions >1 */
   template <bool B = (Dimensions > 1)>
   image(std::shared_ptr<void> &HostPointer, image_channel_order Order,
         image_channel_type Type, const range<Dimensions> &Range,
         const typename detail::enable_if_t<B, range<Dimensions - 1>> &Pitch,
-        const property_list &PropList = {}) {
-    impl = std::make_shared<detail::image_impl<Dimensions>>(
-        HostPointer, Order, Type, Range, Pitch,
-        make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, byte>>(),
-        PropList);
-  }
+        const property_list &PropList = {})
+      : image_plain(HostPointer, Order, Type,
+                    detail::convertToArrayOfN<3, 1>(Range),
+                    detail::convertToArrayOfN<2, 0>(Pitch),
+                    make_unique_ptr<
+                        detail::SYCLMemObjAllocatorHolder<AllocatorT, byte>>(),
+                    Dimensions, PropList, /*IsConstPtr*/ false) {}
 
   /* Available only when: dimensions >1 */
   template <bool B = (Dimensions > 1)>
   image(std::shared_ptr<void> &HostPointer, image_channel_order Order,
         image_channel_type Type, const range<Dimensions> &Range,
         const typename detail::enable_if_t<B, range<Dimensions - 1>> &Pitch,
-        AllocatorT Allocator, const property_list &PropList = {}) {
-    impl = std::make_shared<detail::image_impl<Dimensions>>(
-        HostPointer, Order, Type, Range, Pitch,
-        make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, byte>>(
-            Allocator),
-        PropList);
-  }
+        AllocatorT Allocator, const property_list &PropList = {})
+      : image_plain(
+            HostPointer, Order, Type, detail::convertToArrayOfN<3, 1>(Range),
+            detail::convertToArrayOfN<2, 0>(Pitch),
+            make_unique_ptr<
+                detail::SYCLMemObjAllocatorHolder<AllocatorT, byte>>(Allocator),
+            Dimensions, PropList, /*IsConstPtr*/ false) {}
 
 #ifdef __SYCL_INTERNAL_API
   image(cl_mem ClMemObject, const context &SyclContext,
-        event AvailableEvent = {}) {
-    impl = std::make_shared<detail::image_impl<Dimensions>>(
-        ClMemObject, SyclContext, AvailableEvent,
-        make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT, byte>>());
-  }
+        event AvailableEvent = {})
+      : image_plain(ClMemObject, SyclContext, AvailableEvent,
+                    make_unique_ptr<
+                        detail::SYCLMemObjAllocatorHolder<AllocatorT, byte>>(),
+                    Dimensions) {}
 #endif
 
   /* -- common interface members -- */
@@ -248,32 +354,36 @@ class image {
 
   /* -- property interface members -- */
   template <typename propertyT> bool has_property() const noexcept {
-    return impl->template has_property<propertyT>();
+    return image_plain::template has_property<propertyT>();
   }
 
   template <typename propertyT> propertyT get_property() const {
-    return impl->template get_property<propertyT>();
+    return image_plain::get_property<propertyT>();
   }
 
-  range<Dimensions> get_range() const { return impl->get_range(); }
+  range<Dimensions> get_range() const {
+    return detail::convertToArrayOfN<Dimensions, 0>(image_plain::get_range());
+  }
 
   /* Available only when: dimensions >1 */
   template <bool B = (Dimensions > 1)>
   typename detail::enable_if_t<B, range<Dimensions - 1>> get_pitch() const {
-    return impl->get_pitch();
+    return detail::convertToArrayOfN<Dimensions - 1, 0>(
+        image_plain::get_pitch());
   }
 
   // Returns the size of the image storage in bytes
-  size_t get_size() const { return impl->getSize(); }
+  size_t get_size() const { return image_plain::get_size(); }
 
   // Returns the total number of elements in the image
   __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
   size_t get_count() const { return size(); }
-  size_t size() const noexcept { return impl->get_count(); }
+  size_t size() const noexcept { return image_plain::get_count(); }
 
   // Returns the allocator provided to the image
   AllocatorT get_allocator() const {
-    return impl->template get_allocator<AllocatorT>();
+    return image_plain::get_allocator_internal()
+        ->template getAllocator<AllocatorT>();
   }
 
   template <typename DataT, access::mode AccessMode>
@@ -299,16 +409,91 @@ class image {
 
   template <typename Destination = std::nullptr_t>
   void set_final_data(Destination finalData = nullptr) {
-    impl->set_final_data(finalData);
+    this->set_final_data_internal(finalData);
+  }
+
+  void set_final_data_internal(std::nullptr_t) {
+    image_plain::set_final_data_internal();
   }
 
-  void set_write_back(bool flag = true) { impl->set_write_back(flag); }
+  template <template <typename WeakT> class WeakPtrT, typename WeakT>
+  detail::enable_if_t<
+      std::is_convertible<WeakPtrT<WeakT>, std::weak_ptr<WeakT>>::value>
+  set_final_data_internal(WeakPtrT<WeakT> FinalData) {
+    std::weak_ptr<WeakT> TempFinalData(FinalData);
+    this->set_final_data_internal(TempFinalData);
+  }
+
+  template <typename WeakT>
+  void set_final_data_internal(std::weak_ptr<WeakT> FinalData) {
+    image_plain::set_final_data_internal(
+        [FinalData](const std::function<void(void *const Ptr)> &F) {
+          if (std::shared_ptr<WeakT> LockedFinalData = FinalData.lock())
+            F(LockedFinalData.get());
+        });
+  }
+
+  template <typename Destination>
+  detail::EnableIfOutputPointerT<Destination>
+  set_final_data_internal(Destination FinalData) {
+    if (!FinalData)
+      image_plain::set_final_data_internal();
+    else
+      image_plain::set_final_data_internal(
+          [FinalData](const std::function<void(void *const Ptr)> &F) {
+            F(FinalData);
+          });
+  }
+
+  template <typename Destination>
+  detail::EnableIfOutputIteratorT<Destination>
+  set_final_data_internal(Destination FinalData) {
+    const size_t Size = size();
+    image_plain::set_final_data_internal(
+        [FinalData, Size](const std::function<void(void *const Ptr)> &F) {
+          using DestinationValueT = detail::iterator_value_type_t<Destination>;
+          // TODO if Destination is ContiguousIterator then don't create
+          // ContiguousStorage. updateHostMemory works only with pointer to
+          // continuous data.
+          std::unique_ptr<DestinationValueT[]> ContiguousStorage(
+              new DestinationValueT[Size]);
+          F(ContiguousStorage.get());
+          std::copy(ContiguousStorage.get(), ContiguousStorage.get() + Size,
+                    FinalData);
+        });
+  }
+
+  void set_write_back(bool flag = true) { image_plain::set_write_back(flag); }
 
 private:
-  std::shared_ptr<detail::image_impl<Dimensions>> impl;
+  // This utility api is currently used by accessor to get the element size of
+  // the image. Element size is dependent on num of channels and channel type.
+  // This information is not accessible from the image using any public API.
+  size_t getElementSize() const { return image_plain::getElementSize(); }
+
+  size_t getRowPitch() const { return image_plain::getRowPitch(); }
+
+  size_t getSlicePitch() const { return image_plain::getSlicePitch(); }
+
+  image_channel_order getChannelOrder() const {
+    return image_plain::getChannelOrder();
+  }
+
+  image_channel_type getChannelType() const {
+    return image_plain::getChannelType();
+  }
 
   template <class Obj>
   friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
+
+  template <typename DataT, int Dims, access::mode AccMode,
+            access::target AccTarget, access::placeholder IsPlaceholder,
+            typename PropertyListT>
+  friend class accessor;
+
+  template <typename DataT, int Dims, access::mode AccMode,
+            access::target AccTarget, access::placeholder IsPlaceholder>
+  friend class detail::image_accessor;
 };
 
 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
@@ -318,7 +503,7 @@ namespace std {
 template <int Dimensions, typename AllocatorT>
 struct hash<sycl::image<Dimensions, AllocatorT>> {
   size_t operator()(const sycl::image<Dimensions, AllocatorT> &I) const {
-    return hash<std::shared_ptr<sycl::detail::image_impl<Dimensions>>>()(
+    return hash<std::shared_ptr<sycl::detail::image_impl>>()(
         sycl::detail::getSyclObjImpl(I));
   }
 };
diff --git a/sycl/source/CMakeLists.txt b/sycl/source/CMakeLists.txt
index 3150b00d76f41..11f997a543fc8 100644
--- a/sycl/source/CMakeLists.txt
+++ b/sycl/source/CMakeLists.txt
@@ -164,6 +164,7 @@ set(SYCL_SOURCES
     "detail/util.cpp"
     "detail/xpti_registry.cpp"
     "accessor.cpp"
+    "buffer.cpp"
     "context.cpp"
     "device.cpp"
     "device_selector.cpp"
@@ -172,6 +173,7 @@ set(SYCL_SOURCES
     "exception_list.cpp"
     "function_pointer.cpp"
     "handler.cpp"
+    "image.cpp"
     "interop_handle.cpp"
     "interop_handler.cpp"
     "kernel.cpp"
diff --git a/sycl/source/buffer.cpp b/sycl/source/buffer.cpp
new file mode 100644
index 0000000000000..3152346ee8249
--- /dev/null
+++ b/sycl/source/buffer.cpp
@@ -0,0 +1,126 @@
+#include <detail/buffer_impl.hpp>
+#include <sycl/buffer.hpp>
+#include <sycl/ext/oneapi/accessor_property_list.hpp>
+#include <sycl/properties/all_properties.hpp>
+#include <sycl/property_list.hpp>
+#include <sycl/range.hpp>
+
+#include <memory>
+
+namespace sycl {
+__SYCL_INLINE_VER_NAMESPACE(_V1) {
+namespace detail {
+
+buffer_plain::buffer_plain(
+    size_t SizeInBytes, size_t RequiredAlign, const property_list &Props,
+    std::unique_ptr<detail::SYCLMemObjAllocator> Allocator) {
+
+  impl = std::make_shared<detail::buffer_impl>(SizeInBytes, RequiredAlign,
+                                               Props, std::move(Allocator));
+}
+
+buffer_plain::buffer_plain(
+    void *HostData, size_t SizeInBytes, size_t RequiredAlign,
+    const sycl::property_list &Props,
+    std::unique_ptr<sycl::detail::SYCLMemObjAllocator> Allocator) {
+  impl = std::make_shared<detail::buffer_impl>(
+      HostData, SizeInBytes, RequiredAlign, Props, std::move(Allocator));
+}
+
+buffer_plain::buffer_plain(
+    const void *HostData, size_t SizeInBytes, size_t RequiredAlign,
+    const property_list &Props,
+    std::unique_ptr<detail::SYCLMemObjAllocator> Allocator) {
+  impl = std::make_shared<detail::buffer_impl>(
+      HostData, SizeInBytes, RequiredAlign, Props, std::move(Allocator));
+}
+
+buffer_plain::buffer_plain(
+    const std::shared_ptr<const void> &HostData, const size_t SizeInBytes,
+    size_t RequiredAlign, const property_list &Props,
+    std::unique_ptr<detail::SYCLMemObjAllocator> Allocator, bool IsConstPtr) {
+  impl = std::make_shared<detail::buffer_impl>(
+      HostData, SizeInBytes, RequiredAlign, Props, std::move(Allocator),
+      IsConstPtr);
+}
+
+buffer_plain::buffer_plain(
+    const std::function<void(void *)> &CopyFromInput, const size_t SizeInBytes,
+    size_t RequiredAlign, const property_list &Props,
+    std::unique_ptr<detail::SYCLMemObjAllocator> Allocator, bool IsConstPtr) {
+  impl = std::make_shared<detail::buffer_impl>(
+      CopyFromInput, SizeInBytes, RequiredAlign, Props, std::move(Allocator),
+      IsConstPtr);
+}
+
+buffer_plain::buffer_plain(
+    pi_native_handle MemObject, context SyclContext,
+    std::unique_ptr<detail::SYCLMemObjAllocator> Allocator,
+    bool OwnNativeHandle, event AvailableEvent) {
+  impl = std::make_shared<detail::buffer_impl>(
+      MemObject, std::move(SyclContext), std::move(Allocator), OwnNativeHandle,
+      std::move(AvailableEvent));
+}
+
+void buffer_plain::set_final_data_internal() { impl->set_final_data(nullptr); }
+
+void buffer_plain::set_final_data_internal(
+    const std::function<void(const std::function<void(void *const Ptr)> &)>
+        &FinalDataFunc) {
+  impl->set_final_data(FinalDataFunc);
+}
+
+void buffer_plain::constructorNotification(const detail::code_location &CodeLoc,
+                                           void *UserObj, const void *HostObj,
+                                           const void *Type, uint32_t Dim,
+                                           uint32_t ElemType, size_t Range[3]) {
+  impl->constructorNotification(CodeLoc, UserObj, HostObj, Type, Dim, ElemType,
+                                Range);
+}
+
+void buffer_plain::set_write_back(bool NeedWriteBack) {
+  impl->set_write_back(NeedWriteBack);
+}
+
+#define __SYCL_PARAM_TRAITS_SPEC(param_type)                                   \
+  template <>                                                                  \
+  __SYCL_EXPORT bool buffer_plain::has_property<param_type>() const noexcept { \
+    return impl->has_property<param_type>();                                   \
+  }
+#include <sycl/detail/properties_traits.def>
+
+#undef __SYCL_PARAM_TRAITS_SPEC
+
+#define __SYCL_PARAM_TRAITS_SPEC(param_type)                                   \
+  template <>                                                                  \
+  __SYCL_EXPORT param_type buffer_plain::get_property<param_type>() const {    \
+    return impl->get_property<param_type>();                                   \
+  }
+#include <sycl/detail/properties_traits.def>
+
+#undef __SYCL_PARAM_TRAITS_SPEC
+
+std::vector<pi_native_handle>
+buffer_plain::getNativeVector(backend BackendName) const {
+  return impl->getNativeVector(BackendName);
+}
+
+const std::unique_ptr<SYCLMemObjAllocator> &
+buffer_plain::get_allocator_internal() const {
+  return impl->get_allocator_internal();
+}
+
+void buffer_plain::deleteAccProps(const sycl::detail::PropWithDataKind &Kind) {
+  impl->deleteAccessorProperty(Kind);
+}
+
+void buffer_plain::addOrReplaceAccessorProperties(
+    const property_list &PropertyList) {
+  impl->addOrReplaceAccessorProperties(PropertyList);
+}
+
+size_t buffer_plain::getSize() const { return impl->getSize(); }
+
+} // 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 add26f5e73d32..41f3858ec3176 100644
--- a/sycl/source/detail/accessor_impl.cpp
+++ b/sycl/source/detail/accessor_impl.cpp
@@ -6,11 +6,11 @@
 //
 //===----------------------------------------------------------------------===//
 
+#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>
-#include <sycl/detail/buffer_impl.hpp>
 
 namespace sycl {
 __SYCL_INLINE_VER_NAMESPACE(_V1) {
diff --git a/sycl/source/detail/buffer_impl.cpp b/sycl/source/detail/buffer_impl.cpp
index 5c4eba9493477..0a29b153ef293 100644
--- a/sycl/source/detail/buffer_impl.cpp
+++ b/sycl/source/detail/buffer_impl.cpp
@@ -6,12 +6,12 @@
 //
 //===----------------------------------------------------------------------===//
 
+#include <detail/buffer_impl.hpp>
 #include <detail/context_impl.hpp>
 #include <detail/global_handler.hpp>
 #include <detail/memory_manager.hpp>
 #include <detail/scheduler/scheduler.hpp>
 #include <detail/xpti_registry.hpp>
-#include <sycl/detail/buffer_impl.hpp>
 
 namespace sycl {
 __SYCL_INLINE_VER_NAMESPACE(_V1) {
@@ -38,12 +38,6 @@ void buffer_impl::constructorNotification(const detail::code_location &CodeLoc,
   XPTIRegistry::bufferConstructorNotification(UserObj, CodeLoc, HostObj, Type,
                                               Dim, ElemSize, Range);
 }
-// TODO: remove once ABI break is allowed
-void buffer_impl::constructorNotification(const detail::code_location &CodeLoc,
-                                          void *UserObj) {
-  size_t r[3] = {0, 0, 0};
-  constructorNotification(CodeLoc, UserObj, nullptr, "", 0, 0, r);
-}
 
 void buffer_impl::destructorNotification(void *UserObj) {
   XPTIRegistry::bufferDestructorNotification(UserObj);
diff --git a/sycl/include/sycl/detail/buffer_impl.hpp b/sycl/source/detail/buffer_impl.hpp
similarity index 81%
rename from sycl/include/sycl/detail/buffer_impl.hpp
rename to sycl/source/detail/buffer_impl.hpp
index fa740e463fa62..353befc1b8145 100644
--- a/sycl/include/sycl/detail/buffer_impl.hpp
+++ b/sycl/source/detail/buffer_impl.hpp
@@ -9,12 +9,12 @@
 #pragma once
 
 #include "sycl/detail/pi.h"
+#include <detail/sycl_mem_obj_t.hpp>
 #include <sycl/access/access.hpp>
 #include <sycl/context.hpp>
 #include <sycl/detail/common.hpp>
 #include <sycl/detail/export.hpp>
 #include <sycl/detail/helpers.hpp>
-#include <sycl/detail/sycl_mem_obj_t.hpp>
 #include <sycl/property_list.hpp>
 #include <sycl/stl.hpp>
 #include <sycl/types.hpp>
@@ -81,10 +81,10 @@ class __SYCL_EXPORT buffer_impl final : public SYCLMemObjT {
     BaseT::handleHostData(HostData, RequiredAlign);
   }
 
-  template <typename T>
-  buffer_impl(const std::shared_ptr<T> &HostData, const size_t SizeInBytes,
-              size_t RequiredAlign, const property_list &Props,
-              std::unique_ptr<SYCLMemObjAllocator> Allocator)
+  buffer_impl(const std::shared_ptr<const void> &HostData,
+              const size_t SizeInBytes, size_t RequiredAlign,
+              const property_list &Props,
+              std::unique_ptr<SYCLMemObjAllocator> Allocator, bool IsConstPtr)
       : BaseT(SizeInBytes, Props, std::move(Allocator)) {
 
     if (Props.has_property<
@@ -93,48 +93,28 @@ class __SYCL_EXPORT buffer_impl final : public SYCLMemObjT {
           "The use_pinned_host_memory cannot be used with host pointer",
           PI_ERROR_INVALID_OPERATION);
 
-    BaseT::handleHostData(HostData, RequiredAlign);
+    BaseT::handleHostData(std::const_pointer_cast<void>(HostData),
+                          RequiredAlign, IsConstPtr);
   }
 
-  template <typename T>
-  using EnableIfNotConstIterator =
-      enable_if_t<!iterator_to_const_type_t<T>::value, T>;
-
-  template <class InputIterator>
-  buffer_impl(EnableIfNotConstIterator<InputIterator> First, InputIterator Last,
+  buffer_impl(const std::function<void(void *)> &CopyFromInput,
               const size_t SizeInBytes, size_t RequiredAlign,
               const property_list &Props,
-              std::unique_ptr<SYCLMemObjAllocator> Allocator)
+              std::unique_ptr<detail::SYCLMemObjAllocator> Allocator,
+              bool IsConstPtr)
       : BaseT(SizeInBytes, Props, std::move(Allocator)) {
-
-    if (Props.has_property<sycl::property::buffer::use_host_ptr>())
+    if (Props.has_property<
+            sycl::ext::oneapi::property::buffer::use_pinned_host_memory>())
       throw sycl::invalid_object_error(
-          "Buffer constructor from a pair of iterator values cannot have the "
-          "use_host_ptr property.",
+          "The use_pinned_host_memory cannot be used with host pointer",
           PI_ERROR_INVALID_OPERATION);
 
-    BaseT::handleHostData(First, Last, RequiredAlign);
+    BaseT::handleHostData(CopyFromInput, RequiredAlign, IsConstPtr);
   }
 
   template <typename T>
-  using EnableIfConstIterator =
-      enable_if_t<iterator_to_const_type_t<T>::value, T>;
-
-  template <class InputIterator>
-  buffer_impl(EnableIfConstIterator<InputIterator> First, InputIterator Last,
-              const size_t SizeInBytes, size_t RequiredAlign,
-              const property_list &Props,
-              std::unique_ptr<SYCLMemObjAllocator> Allocator)
-      : BaseT(SizeInBytes, Props, std::move(Allocator)) {
-
-    if (Props.has_property<sycl::property::buffer::use_host_ptr>())
-      throw sycl::invalid_object_error(
-          "Buffer constructor from a pair of iterator values cannot have the "
-          "use_host_ptr property.",
-          PI_ERROR_INVALID_OPERATION);
-
-    BaseT::handleHostData(First, Last, RequiredAlign);
-  }
+  using EnableIfNotConstIterator =
+      enable_if_t<!iterator_to_const_type_t<T>::value, T>;
 
   buffer_impl(cl_mem MemObject, const context &SyclContext,
               std::unique_ptr<SYCLMemObjAllocator> Allocator,
diff --git a/sycl/source/detail/image_impl.cpp b/sycl/source/detail/image_impl.cpp
index 5a583eb893697..324dbf574985b 100644
--- a/sycl/source/detail/image_impl.cpp
+++ b/sycl/source/detail/image_impl.cpp
@@ -7,9 +7,8 @@
 //===----------------------------------------------------------------------===//
 
 #include <detail/context_impl.hpp>
+#include <detail/image_impl.hpp>
 #include <detail/memory_manager.hpp>
-#include <sycl/detail/image_impl.hpp>
-#include <sycl/image.hpp>
 
 #include <algorithm>
 #include <vector>
@@ -261,13 +260,13 @@ static void getImageInfo(const ContextImplPtr Context, RT::PiMemImageInfo Info,
                                             nullptr);
 }
 
-template <int Dimensions>
-image_impl<Dimensions>::image_impl(
-    cl_mem MemObject, const context &SyclContext, event AvailableEvent,
-    std::unique_ptr<SYCLMemObjAllocator> Allocator)
+image_impl::image_impl(cl_mem MemObject, const context &SyclContext,
+                       event AvailableEvent,
+                       std::unique_ptr<SYCLMemObjAllocator> Allocator,
+                       uint8_t Dimensions)
     : BaseT(MemObject, SyclContext, std::move(AvailableEvent),
             std::move(Allocator)),
-      MRange(InitializedVal<Dimensions, range>::template get<0>()) {
+      MDimensions(Dimensions), MRange({0, 0, 0}) {
   RT::PiMem Mem = pi::cast<RT::PiMem>(BaseT::MInteropMemObject);
   const ContextImplPtr Context = getSyclObjImpl(SyclContext);
   const detail::plugin &Plugin = Context->getPlugin();
@@ -286,7 +285,7 @@ image_impl<Dimensions>::image_impl(
   getImageInfo(Context, PI_IMAGE_INFO_ROW_PITCH, MRowPitch, Mem);
   getImageInfo(Context, PI_IMAGE_INFO_SLICE_PITCH, MSlicePitch, Mem);
 
-  switch (Dimensions) {
+  switch (MDimensions) {
   case 3:
     getImageInfo(Context, PI_IMAGE_INFO_DEPTH, MRange[2], Mem);
     __SYCL_FALLTHROUGH;
@@ -298,10 +297,8 @@ image_impl<Dimensions>::image_impl(
   }
 }
 
-template <int Dimensions>
-void *image_impl<Dimensions>::allocateMem(ContextImplPtr Context,
-                                          bool InitFromUserData, void *HostPtr,
-                                          RT::PiEvent &OutEventToWait) {
+void *image_impl::allocateMem(ContextImplPtr Context, bool InitFromUserData,
+                              void *HostPtr, RT::PiEvent &OutEventToWait) {
   bool HostPtrReadOnly = false;
   BaseT::determineHostPtr(Context, InitFromUserData, HostPtr, HostPtrReadOnly);
 
@@ -319,10 +316,8 @@ void *image_impl<Dimensions>::allocateMem(ContextImplPtr Context,
       OutEventToWait);
 }
 
-template <int Dimensions>
-bool image_impl<Dimensions>::checkImageDesc(const RT::PiMemImageDesc &Desc,
-                                            ContextImplPtr Context,
-                                            void *UserPtr) {
+bool image_impl::checkImageDesc(const RT::PiMemImageDesc &Desc,
+                                ContextImplPtr Context, void *UserPtr) {
   if (checkAny(Desc.image_type, PI_MEM_TYPE_IMAGE1D, PI_MEM_TYPE_IMAGE1D_ARRAY,
                PI_MEM_TYPE_IMAGE2D_ARRAY, PI_MEM_TYPE_IMAGE2D) &&
       !checkImageValueRange<info::device::image2d_max_width>(
@@ -401,9 +396,8 @@ bool image_impl<Dimensions>::checkImageDesc(const RT::PiMemImageDesc &Desc,
   return true;
 }
 
-template <int Dimensions>
-bool image_impl<Dimensions>::checkImageFormat(
-    const RT::PiMemImageFormat &Format, ContextImplPtr Context) {
+bool image_impl::checkImageFormat(const RT::PiMemImageFormat &Format,
+                                  ContextImplPtr Context) {
   (void)Context;
   if (checkAny(Format.image_channel_order, PI_IMAGE_CHANNEL_ORDER_INTENSITY,
                PI_IMAGE_CHANNEL_ORDER_LUMINANCE) &&
@@ -446,16 +440,10 @@ bool image_impl<Dimensions>::checkImageFormat(
   return true;
 }
 
-template <int Dimensions>
-std::vector<device>
-image_impl<Dimensions>::getDevices(const ContextImplPtr Context) {
+std::vector<device> image_impl::getDevices(const ContextImplPtr Context) {
   return Context->get_info<info::context::devices>();
 }
 
-template class image_impl<1>;
-template class image_impl<2>;
-template class image_impl<3>;
-
 } // namespace detail
 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
 } // namespace sycl
diff --git a/sycl/include/sycl/detail/image_impl.hpp b/sycl/source/detail/image_impl.hpp
similarity index 65%
rename from sycl/include/sycl/detail/image_impl.hpp
rename to sycl/source/detail/image_impl.hpp
index 303ebf3444ec2..f8a345302082f 100644
--- a/sycl/include/sycl/detail/image_impl.hpp
+++ b/sycl/source/detail/image_impl.hpp
@@ -8,13 +8,14 @@
 
 #pragma once
 
+#include <detail/sycl_mem_obj_t.hpp>
 #include <sycl/detail/aligned_allocator.hpp>
 #include <sycl/detail/common.hpp>
 #include <sycl/detail/export.hpp>
 #include <sycl/detail/generic_type_traits.hpp>
-#include <sycl/detail/sycl_mem_obj_t.hpp>
 #include <sycl/device.hpp>
 #include <sycl/event.hpp>
+#include <sycl/image.hpp>
 #include <sycl/property_list.hpp>
 #include <sycl/range.hpp>
 #include <sycl/stl.hpp>
@@ -57,29 +58,14 @@ convertChannelType(image_channel_type Type);
 __SYCL_EXPORT image_channel_type
 convertChannelType(RT::PiMemImageChannelType Type);
 
-// validImageDataT: cl_int4, cl_uint4, cl_float4, cl_half4
-template <typename T>
-using is_validImageDataT = typename detail::is_contained<
-    T, type_list<cl_int4, cl_uint4, cl_float4, cl_half4>>::type;
-
-template <typename DataT>
-using EnableIfImgAccDataT =
-    typename detail::enable_if_t<is_validImageDataT<DataT>::value, DataT>;
-
-template <int Dimensions>
 class __SYCL_EXPORT image_impl final : public SYCLMemObjT {
   using BaseT = SYCLMemObjT;
   using typename BaseT::MemObjType;
 
 private:
-  template <bool B>
-  using EnableIfPitchT = typename detail::enable_if_t<B, range<Dimensions - 1>>;
-  static_assert(Dimensions >= 1 || Dimensions <= 3,
-                "Dimensions of sycl::image can be 1, 2 or 3");
-
   void setPitches() {
     size_t WHD[3] = {1, 1, 1}; // Width, Height, Depth.
-    for (int I = 0; I < Dimensions; I++)
+    for (int I = 0; I < MDimensions; I++)
       WHD[I] = MRange[I];
 
     MRowPitch = MElementSize * WHD[0];
@@ -87,41 +73,38 @@ class __SYCL_EXPORT image_impl final : public SYCLMemObjT {
     BaseT::MSizeInBytes = MSlicePitch * WHD[2];
   }
 
-  template <bool B = (Dimensions > 1)>
-  void setPitches(const EnableIfPitchT<B> Pitch) {
+  void setPitches(const range<2> &Pitch) {
     MRowPitch = Pitch[0];
     MSlicePitch =
-        (Dimensions == 3) ? Pitch[1] : MRowPitch; // Dimensions will be 2/3.
+        (MDimensions == 3) ? Pitch[1] : MRowPitch; // Dimensions will be 2/3.
     // NumSlices is depth when dim==3, and height when dim==2.
     size_t NumSlices =
-        (Dimensions == 3) ? MRange[2] : MRange[1]; // Dimensions will be 2/3.
+        (MDimensions == 3) ? MRange[2] : MRange[1]; // Dimensions will be 2/3.
 
     BaseT::MSizeInBytes = MSlicePitch * NumSlices;
   }
 
 public:
   image_impl(image_channel_order Order, image_channel_type Type,
-             const range<Dimensions> &ImageRange,
-             std::unique_ptr<SYCLMemObjAllocator> Allocator,
+             const range<3> &ImageRange,
+             std::unique_ptr<SYCLMemObjAllocator> Allocator, uint8_t Dimensions,
              const property_list &PropList = {})
       : image_impl((void *)nullptr, Order, Type, ImageRange,
-                   std::move(Allocator), PropList) {}
+                   std::move(Allocator), Dimensions, PropList) {}
 
-  template <bool B = (Dimensions > 1)>
   image_impl(image_channel_order Order, image_channel_type Type,
-             const range<Dimensions> &ImageRange,
-             const EnableIfPitchT<B> &Pitch,
-             std::unique_ptr<SYCLMemObjAllocator> Allocator,
+             const range<3> &ImageRange, const range<2> &Pitch,
+             std::unique_ptr<SYCLMemObjAllocator> Allocator, uint8_t Dimensions,
              const property_list &PropList = {})
       : image_impl((void *)nullptr, Order, Type, ImageRange, Pitch,
-                   std::move(Allocator), PropList) {}
+                   std::move(Allocator), Dimensions, PropList) {}
 
   image_impl(void *HData, image_channel_order Order, image_channel_type Type,
-             const range<Dimensions> &ImageRange,
-             std::unique_ptr<SYCLMemObjAllocator> Allocator,
+             const range<3> &ImageRange,
+             std::unique_ptr<SYCLMemObjAllocator> Allocator, uint8_t Dimensions,
              const property_list &PropList = {})
-      : BaseT(PropList, std::move(Allocator)), MRange(ImageRange),
-        MOrder(Order), MType(Type),
+      : BaseT(PropList, std::move(Allocator)), MDimensions(Dimensions),
+        MRange(ImageRange), MOrder(Order), MType(Type),
         MNumChannels(getImageNumberChannels(MOrder)),
         MElementSize(getImageElementSize(MNumChannels, MType)) {
     setPitches();
@@ -129,80 +112,67 @@ class __SYCL_EXPORT image_impl final : public SYCLMemObjT {
   }
 
   image_impl(const void *HData, image_channel_order Order,
-             image_channel_type Type, const range<Dimensions> &ImageRange,
-             std::unique_ptr<SYCLMemObjAllocator> Allocator,
+             image_channel_type Type, const range<3> &ImageRange,
+             std::unique_ptr<SYCLMemObjAllocator> Allocator, uint8_t Dimensions,
              const property_list &PropList = {})
-      : BaseT(PropList, std::move(Allocator)), MRange(ImageRange),
-        MOrder(Order), MType(Type),
+      : BaseT(PropList, std::move(Allocator)), MDimensions(Dimensions),
+        MRange(ImageRange), MOrder(Order), MType(Type),
         MNumChannels(getImageNumberChannels(MOrder)),
         MElementSize(getImageElementSize(MNumChannels, MType)) {
     setPitches();
     BaseT::handleHostData(HData, detail::getNextPowerOfTwo(MElementSize));
   }
 
-  template <bool B = (Dimensions > 1)>
   image_impl(void *HData, image_channel_order Order, image_channel_type Type,
-             const range<Dimensions> &ImageRange,
-             const EnableIfPitchT<B> &Pitch,
-             std::unique_ptr<SYCLMemObjAllocator> Allocator,
+             const range<3> &ImageRange, const range<2> &Pitch,
+             std::unique_ptr<SYCLMemObjAllocator> Allocator, uint8_t Dimensions,
              const property_list &PropList = {})
-      : BaseT(PropList, std::move(Allocator)), MRange(ImageRange),
-        MOrder(Order), MType(Type),
+      : BaseT(PropList, std::move(Allocator)), MDimensions(Dimensions),
+        MRange(ImageRange), MOrder(Order), MType(Type),
         MNumChannels(getImageNumberChannels(MOrder)),
         MElementSize(getImageElementSize(MNumChannels, MType)) {
     setPitches(Pitch);
     BaseT::handleHostData(HData, detail::getNextPowerOfTwo(MElementSize));
   }
 
-  image_impl(std::shared_ptr<void> &HData, image_channel_order Order,
-             image_channel_type Type, const range<Dimensions> &ImageRange,
-             std::unique_ptr<SYCLMemObjAllocator> Allocator,
-             const property_list &PropList = {})
-      : BaseT(PropList, std::move(Allocator)), MRange(ImageRange),
-        MOrder(Order), MType(Type),
+  image_impl(const std::shared_ptr<const void> &HData,
+             image_channel_order Order, image_channel_type Type,
+             const range<3> &ImageRange,
+             std::unique_ptr<SYCLMemObjAllocator> Allocator, uint8_t Dimensions,
+             const property_list &PropList, bool IsConstPtr)
+      : BaseT(PropList, std::move(Allocator)), MDimensions(Dimensions),
+        MRange(ImageRange), MOrder(Order), MType(Type),
         MNumChannels(getImageNumberChannels(MOrder)),
         MElementSize(getImageElementSize(MNumChannels, MType)) {
     setPitches();
-    BaseT::handleHostData(HData, detail::getNextPowerOfTwo(MElementSize));
+    BaseT::handleHostData(std::const_pointer_cast<void>(HData),
+                          detail::getNextPowerOfTwo(MElementSize), IsConstPtr);
   }
 
-  /* Available only when: Dimensions > 1 */
-  template <bool B = (Dimensions > 1)>
-  image_impl(std::shared_ptr<void> &HData, image_channel_order Order,
-             image_channel_type Type, const range<Dimensions> &ImageRange,
-             const EnableIfPitchT<B> &Pitch,
-             std::unique_ptr<SYCLMemObjAllocator> Allocator,
-             const property_list &PropList = {})
-      : BaseT(PropList, std::move(Allocator)), MRange(ImageRange),
-        MOrder(Order), MType(Type),
+  image_impl(const std::shared_ptr<const void> &HData,
+             image_channel_order Order, image_channel_type Type,
+             const range<3> &ImageRange, const range<2> &Pitch,
+             std::unique_ptr<SYCLMemObjAllocator> Allocator, uint8_t Dimensions,
+             const property_list &PropList, bool IsConstPtr)
+      : BaseT(PropList, std::move(Allocator)), MDimensions(Dimensions),
+        MRange(ImageRange), MOrder(Order), MType(Type),
         MNumChannels(getImageNumberChannels(MOrder)),
         MElementSize(getImageElementSize(MNumChannels, MType)) {
     setPitches(Pitch);
-    BaseT::handleHostData(HData, detail::getNextPowerOfTwo(MElementSize));
+    BaseT::handleHostData(std::const_pointer_cast<void>(HData),
+                          detail::getNextPowerOfTwo(MElementSize), IsConstPtr);
   }
 
   image_impl(cl_mem MemObject, const context &SyclContext, event AvailableEvent,
-             std::unique_ptr<SYCLMemObjAllocator> Allocator);
+             std::unique_ptr<SYCLMemObjAllocator> Allocator,
+             uint8_t Dimensions);
 
   // Return a range object representing the size of the image in terms of the
   // number of elements in each dimension as passed to the constructor
-  range<Dimensions> get_range() const { return MRange; }
+  range<3> get_range() const { return MRange; }
 
   // Return a range object representing the pitch of the image in bytes.
-  // Available only when: Dimensions == 2.
-  template <bool B = (Dimensions == 2)>
-  typename detail::enable_if_t<B, range<1>> get_pitch() const {
-    range<1> Temp = range<1>(MRowPitch);
-    return Temp;
-  }
-
-  // Return a range object representing the pitch of the image in bytes.
-  // Available only when: Dimensions == 3.
-  template <bool B = (Dimensions == 3)>
-  typename detail::enable_if_t<B, range<2>> get_pitch() const {
-    range<2> Temp = range<2>(MRowPitch, MSlicePitch);
-    return Temp;
-  }
+  range<2> get_pitch() const { return {MRowPitch, MSlicePitch}; }
 
   // Returns the total number of elements in the image
   size_t get_count() const { return size(); }
@@ -237,9 +207,9 @@ class __SYCL_EXPORT image_impl final : public SYCLMemObjT {
   std::vector<device> getDevices(const ContextImplPtr Context);
 
   RT::PiMemObjectType getImageType() {
-    if (Dimensions == 1)
+    if (MDimensions == 1)
       return (MIsArrayImage ? PI_MEM_TYPE_IMAGE1D_ARRAY : PI_MEM_TYPE_IMAGE1D);
-    if (Dimensions == 2)
+    if (MDimensions == 2)
       return (MIsArrayImage ? PI_MEM_TYPE_IMAGE2D_ARRAY : PI_MEM_TYPE_IMAGE2D);
     return PI_MEM_TYPE_IMAGE3D;
   }
@@ -252,8 +222,8 @@ class __SYCL_EXPORT image_impl final : public SYCLMemObjT {
     // is different than MAccessRange, etc in bufffers)
     static constexpr int XTermPos = 0, YTermPos = 1, ZTermPos = 2;
     Desc.image_width = MRange[XTermPos];
-    Desc.image_height = Dimensions > 1 ? MRange[YTermPos] : 1;
-    Desc.image_depth = Dimensions > 2 ? MRange[ZTermPos] : 1;
+    Desc.image_height = MDimensions > 1 ? MRange[YTermPos] : 1;
+    Desc.image_depth = MDimensions > 2 ? MRange[ZTermPos] : 1;
 
     // TODO handle cases with IMAGE1D_ARRAY and IMAGE2D_ARRAY
     Desc.image_array_size = 0;
@@ -279,8 +249,9 @@ class __SYCL_EXPORT image_impl final : public SYCLMemObjT {
   bool checkImageFormat(const RT::PiMemImageFormat &Format,
                         ContextImplPtr Context);
 
+  uint8_t MDimensions = 0;
   bool MIsArrayImage = false;
-  range<Dimensions> MRange;
+  range<3> MRange;
   image_channel_order MOrder;
   image_channel_type MType;
   uint8_t MNumChannels = 0; // Maximum Value - 4
diff --git a/sycl/source/detail/memory_manager.hpp b/sycl/source/detail/memory_manager.hpp
index 6ad9f89ac3e2d..8025cfa94ef44 100644
--- a/sycl/source/detail/memory_manager.hpp
+++ b/sycl/source/detail/memory_manager.hpp
@@ -8,9 +8,9 @@
 
 #pragma once
 
+#include <detail/sycl_mem_obj_i.hpp>
 #include <sycl/access/access.hpp>
 #include <sycl/detail/export.hpp>
-#include <sycl/detail/sycl_mem_obj_i.hpp>
 #include <sycl/id.hpp>
 #include <sycl/property_list.hpp>
 #include <sycl/range.hpp>
diff --git a/sycl/source/detail/scheduler/graph_builder.cpp b/sycl/source/detail/scheduler/graph_builder.cpp
index 58899cbe2c668..10ae8d5aa0d57 100644
--- a/sycl/source/detail/scheduler/graph_builder.cpp
+++ b/sycl/source/detail/scheduler/graph_builder.cpp
@@ -12,6 +12,7 @@
 #include <detail/memory_manager.hpp>
 #include <detail/queue_impl.hpp>
 #include <detail/scheduler/scheduler.hpp>
+#include <detail/sycl_mem_obj_t.hpp>
 #include <sycl/access/access.hpp>
 #include <sycl/exception.hpp>
 
diff --git a/sycl/source/detail/scheduler/scheduler.cpp b/sycl/source/detail/scheduler/scheduler.cpp
index 74564413cb377..7eed957f1afa7 100644
--- a/sycl/source/detail/scheduler/scheduler.cpp
+++ b/sycl/source/detail/scheduler/scheduler.cpp
@@ -6,7 +6,7 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include "sycl/detail/sycl_mem_obj_i.hpp"
+#include "detail/sycl_mem_obj_i.hpp"
 #include <detail/global_handler.hpp>
 #include <detail/queue_impl.hpp>
 #include <detail/scheduler/scheduler.hpp>
diff --git a/sycl/source/detail/scheduler/scheduler.hpp b/sycl/source/detail/scheduler/scheduler.hpp
index e2903693c17f5..b6676fa4b0aa9 100644
--- a/sycl/source/detail/scheduler/scheduler.hpp
+++ b/sycl/source/detail/scheduler/scheduler.hpp
@@ -10,8 +10,8 @@
 
 #include <detail/scheduler/commands.hpp>
 #include <detail/scheduler/leaves_collection.hpp>
+#include <detail/sycl_mem_obj_i.hpp>
 #include <sycl/detail/cg.hpp>
-#include <sycl/detail/sycl_mem_obj_i.hpp>
 
 #include <cstddef>
 #include <memory>
diff --git a/sycl/source/detail/scheduler/scheduler_helpers.cpp b/sycl/source/detail/scheduler/scheduler_helpers.cpp
index 75a79eb8909fd..175684342c60b 100644
--- a/sycl/source/detail/scheduler/scheduler_helpers.cpp
+++ b/sycl/source/detail/scheduler/scheduler_helpers.cpp
@@ -6,6 +6,7 @@
 //
 //===----------------------------------------------------------------------===//
 
+#include <detail/buffer_impl.hpp>
 #include <detail/queue_impl.hpp>
 #include <detail/scheduler/scheduler.hpp>
 #include <detail/scheduler/scheduler_helpers.hpp>
diff --git a/sycl/include/sycl/detail/sycl_mem_obj_i.hpp b/sycl/source/detail/sycl_mem_obj_i.hpp
similarity index 100%
rename from sycl/include/sycl/detail/sycl_mem_obj_i.hpp
rename to sycl/source/detail/sycl_mem_obj_i.hpp
diff --git a/sycl/source/detail/sycl_mem_obj_t.cpp b/sycl/source/detail/sycl_mem_obj_t.cpp
index 052733d056990..3636e5d5e0ac6 100644
--- a/sycl/source/detail/sycl_mem_obj_t.cpp
+++ b/sycl/source/detail/sycl_mem_obj_t.cpp
@@ -11,7 +11,7 @@
 #include <detail/memory_manager.hpp>
 #include <detail/plugin.hpp>
 #include <detail/scheduler/scheduler.hpp>
-#include <sycl/detail/sycl_mem_obj_t.hpp>
+#include <detail/sycl_mem_obj_t.hpp>
 
 namespace sycl {
 __SYCL_INLINE_VER_NAMESPACE(_V1) {
diff --git a/sycl/include/sycl/detail/sycl_mem_obj_t.hpp b/sycl/source/detail/sycl_mem_obj_t.hpp
similarity index 63%
rename from sycl/include/sycl/detail/sycl_mem_obj_t.hpp
rename to sycl/source/detail/sycl_mem_obj_t.hpp
index 8afdad9271630..cc49509c98975 100644
--- a/sycl/include/sycl/detail/sycl_mem_obj_t.hpp
+++ b/sycl/source/detail/sycl_mem_obj_t.hpp
@@ -8,10 +8,10 @@
 
 #pragma once
 
+#include <detail/sycl_mem_obj_i.hpp>
 #include <sycl/detail/common.hpp>
 #include <sycl/detail/export.hpp>
 #include <sycl/detail/sycl_mem_obj_allocator.hpp>
-#include <sycl/detail/sycl_mem_obj_i.hpp>
 #include <sycl/detail/type_traits.hpp>
 #include <sycl/event.hpp>
 #include <sycl/properties/buffer_properties.hpp>
@@ -82,79 +82,52 @@ class __SYCL_EXPORT SYCLMemObjT : public SYCLMemObjI {
 
   const plugin &getPlugin() const;
 
-  __SYCL_DLL_LOCAL size_t getSize() const override { return MSizeInBytes; }
+  size_t getSize() const override { return MSizeInBytes; }
   __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
-  __SYCL_DLL_LOCAL size_t get_count() const { return size(); }
-  __SYCL_DLL_LOCAL size_t size() const noexcept {
+  size_t get_count() const { return size(); }
+  size_t size() const noexcept {
     size_t AllocatorValueSize = MAllocator->getValueSize();
     return (getSize() + AllocatorValueSize - 1) / AllocatorValueSize;
   }
 
-  template <typename propertyT>
-  __SYCL_DLL_LOCAL bool has_property() const noexcept {
+  template <typename propertyT> bool has_property() const noexcept {
     return MProps.has_property<propertyT>();
   }
 
-  template <typename propertyT>
-  __SYCL_DLL_LOCAL propertyT get_property() const {
+  template <typename propertyT> propertyT get_property() const {
     return MProps.get_property<propertyT>();
   }
 
-  __SYCL_DLL_LOCAL void
-  addOrReplaceAccessorProperties(const property_list &PropertyList) {
+  void addOrReplaceAccessorProperties(const property_list &PropertyList) {
     MProps.add_or_replace_accessor_properties(PropertyList);
   }
 
-  __SYCL_DLL_LOCAL void deleteAccessorProperty(const PropWithDataKind &Kind) {
+  void deleteAccessorProperty(const PropWithDataKind &Kind) {
     MProps.delete_accessor_property(Kind);
   }
 
-  template <typename AllocatorT>
-  __SYCL_DLL_LOCAL AllocatorT get_allocator() const {
-    return MAllocator->getAllocator<AllocatorT>();
+  const std::unique_ptr<SYCLMemObjAllocator> &get_allocator_internal() const {
+    return MAllocator;
   }
 
-  __SYCL_DLL_LOCAL void *allocateHostMem() override {
-    return MAllocator->allocate(size());
-  }
+  void *allocateHostMem() override { return MAllocator->allocate(size()); }
 
-  __SYCL_DLL_LOCAL void releaseHostMem(void *Ptr) override {
+  void releaseHostMem(void *Ptr) override {
     if (Ptr)
       MAllocator->deallocate(Ptr, size());
   }
 
   void releaseMem(ContextImplPtr Context, void *MemAllocation) override;
 
-  __SYCL_DLL_LOCAL void *getUserPtr() const {
+  void *getUserPtr() const {
     return MOpenCLInterop ? static_cast<void *>(MInteropMemObject) : MUserPtr;
   }
 
-  __SYCL_DLL_LOCAL void set_write_back(bool NeedWriteBack) {
-    MNeedWriteBack = NeedWriteBack;
-  }
-
-  __SYCL_DLL_LOCAL void set_final_data(std::nullptr_t) {
-    MUploadDataFunctor = nullptr;
-  }
+  void set_write_back(bool NeedWriteBack) { MNeedWriteBack = NeedWriteBack; }
 
-  template <template <typename T> class PtrT, typename T>
-  __SYCL_DLL_LOCAL
-      enable_if_t<std::is_convertible<PtrT<T>, std::weak_ptr<T>>::value>
-      set_final_data(PtrT<T> FinalData) {
-    std::weak_ptr<T> TempFinalData(FinalData);
-    set_final_data(TempFinalData);
-  }
-
-  template <typename T>
-  __SYCL_DLL_LOCAL void set_final_data(std::weak_ptr<T> FinalData) {
-    MUploadDataFunctor = [this, FinalData]() {
-      if (std::shared_ptr<T> LockedFinalData = FinalData.lock()) {
-        updateHostMemory(LockedFinalData.get());
-      }
-    };
-  }
+  void set_final_data(std::nullptr_t) { MUploadDataFunctor = nullptr; }
 
-  __SYCL_DLL_LOCAL void set_final_data_from_storage() {
+  void set_final_data_from_storage() {
     MUploadDataFunctor = [this]() {
       if (MSharedPtrStorage.use_count() > 1) {
         void *FinalData = const_cast<void *>(MSharedPtrStorage.get());
@@ -163,29 +136,13 @@ class __SYCL_EXPORT SYCLMemObjT : public SYCLMemObjI {
     };
   }
 
-  template <typename Destination>
-  __SYCL_DLL_LOCAL EnableIfOutputPointerT<Destination>
-  set_final_data(Destination FinalData) {
-    if (!FinalData)
-      MUploadDataFunctor = nullptr;
-    else
-      MUploadDataFunctor = [this, FinalData]() { updateHostMemory(FinalData); };
-  }
+  void set_final_data(
+      const std::function<void(const std::function<void(void *const Ptr)> &)>
+          &FinalDataFunc) {
 
-  template <typename Destination>
-  __SYCL_DLL_LOCAL EnableIfOutputIteratorT<Destination>
-  set_final_data(Destination FinalData) {
-    MUploadDataFunctor = [this, FinalData]() {
-      using DestinationValueT = iterator_value_type_t<Destination>;
-      // TODO if Destination is ContiguousIterator then don't create
-      // ContiguousStorage. updateHostMemory works only with pointer to
-      // continuous data.
-      const size_t Size = MSizeInBytes / sizeof(DestinationValueT);
-      std::unique_ptr<DestinationValueT[]> ContiguousStorage(
-          new DestinationValueT[Size]);
-      updateHostMemory(ContiguousStorage.get());
-      std::copy(ContiguousStorage.get(), ContiguousStorage.get() + Size,
-                FinalData);
+    auto UpdateFunc = [this](void *const Ptr) { updateHostMemory(Ptr); };
+    MUploadDataFunctor = [FinalDataFunc, UpdateFunc]() {
+      FinalDataFunc(UpdateFunc);
     };
   }
 
@@ -200,22 +157,23 @@ class __SYCL_EXPORT SYCLMemObjT : public SYCLMemObjI {
   void updateHostMemory();
 
 public:
-  __SYCL_DLL_LOCAL bool useHostPtr() {
+  bool useHostPtr() {
     return has_property<property::buffer::use_host_ptr>() ||
            has_property<property::image::use_host_ptr>();
   }
 
-  __SYCL_DLL_LOCAL bool canReuseHostPtr(void *HostPtr,
-                                        const size_t RequiredAlign) {
+  bool canReuseHostPtr(void *HostPtr, const size_t RequiredAlign) {
     bool Aligned =
         (reinterpret_cast<std::uintptr_t>(HostPtr) % RequiredAlign) == 0;
     return Aligned || useHostPtr();
   }
 
-  __SYCL_DLL_LOCAL void handleHostData(void *HostPtr,
-                                       const size_t RequiredAlign) {
-    if (!MHostPtrReadOnly)
-      set_final_data(reinterpret_cast<char *>(HostPtr));
+  void handleHostData(void *HostPtr, const size_t RequiredAlign) {
+    if (!MHostPtrReadOnly && HostPtr) {
+      set_final_data([HostPtr](const std::function<void(void *const Ptr)> &F) {
+        F(HostPtr);
+      });
+    }
 
     if (canReuseHostPtr(HostPtr, RequiredAlign)) {
       MUserPtr = HostPtr;
@@ -227,17 +185,15 @@ class __SYCL_EXPORT SYCLMemObjT : public SYCLMemObjI {
     }
   }
 
-  __SYCL_DLL_LOCAL void handleHostData(const void *HostPtr,
-                                       const size_t RequiredAlign) {
+  void handleHostData(const void *HostPtr, const size_t RequiredAlign) {
     MHostPtrReadOnly = true;
     handleHostData(const_cast<void *>(HostPtr), RequiredAlign);
   }
 
-  template <typename T>
-  __SYCL_DLL_LOCAL void handleHostData(const std::shared_ptr<T> &HostPtr,
-                                       const size_t RequiredAlign) {
+  void handleHostData(const std::shared_ptr<void> &HostPtr,
+                      const size_t RequiredAlign, bool IsConstPtr) {
     MSharedPtrStorage = HostPtr;
-    MHostPtrReadOnly = std::is_const<T>::value;
+    MHostPtrReadOnly = IsConstPtr;
     if (HostPtr) {
       if (!MHostPtrReadOnly)
         set_final_data_from_storage();
@@ -253,10 +209,9 @@ class __SYCL_EXPORT SYCLMemObjT : public SYCLMemObjI {
     }
   }
 
-  template <class InputIterator>
-  __SYCL_DLL_LOCAL void handleHostData(InputIterator First, InputIterator Last,
-                                       const size_t RequiredAlign) {
-    MHostPtrReadOnly = iterator_to_const_type_t<InputIterator>::value;
+  void handleHostData(const std::function<void(void *)> &CopyFromInput,
+                      const size_t RequiredAlign, bool IsConstPtr) {
+    MHostPtrReadOnly = IsConstPtr;
     setAlign(RequiredAlign);
     if (useHostPtr())
       throw runtime_error(
@@ -268,26 +223,18 @@ class __SYCL_EXPORT SYCLMemObjT : public SYCLMemObjI {
     MShadowCopy = allocateHostMem();
     MUserPtr = MShadowCopy;
 
-    // We need to cast MUserPtr to pointer to the iteration type to get correct
-    // offset in std::copy when it will increment destination pointer.
-    using IteratorValueType = iterator_value_type_t<InputIterator>;
-    using IteratorNonConstValueType = remove_const_t<IteratorValueType>;
-    using IteratorPointerToNonConstValueType =
-        add_pointer_t<IteratorNonConstValueType>;
-    std::copy(First, Last,
-              static_cast<IteratorPointerToNonConstValueType>(MUserPtr));
+    CopyFromInput(MUserPtr);
   }
 
-  __SYCL_DLL_LOCAL void setAlign(size_t RequiredAlign) {
+  void setAlign(size_t RequiredAlign) {
     MAllocator->setAlignment(RequiredAlign);
   }
 
   static size_t getBufSizeForContext(const ContextImplPtr &Context,
                                      pi_native_handle MemObject);
 
-  __SYCL_DLL_LOCAL void *allocateMem(ContextImplPtr Context,
-                                     bool InitFromUserData, void *HostPtr,
-                                     RT::PiEvent &InteropEvent) override {
+  void *allocateMem(ContextImplPtr Context, bool InitFromUserData,
+                    void *HostPtr, RT::PiEvent &InteropEvent) override {
     (void)Context;
     (void)InitFromUserData;
     (void)HostPtr;
@@ -295,9 +242,7 @@ class __SYCL_EXPORT SYCLMemObjT : public SYCLMemObjI {
     throw runtime_error("Not implemented", PI_ERROR_INVALID_OPERATION);
   }
 
-  __SYCL_DLL_LOCAL MemObjType getType() const override {
-    return MemObjType::Undefined;
-  }
+  MemObjType getType() const override { return MemObjType::Undefined; }
 
   ContextImplPtr getInteropContext() const override { return MInteropContext; }
 
diff --git a/sycl/source/image.cpp b/sycl/source/image.cpp
new file mode 100644
index 0000000000000..f8f6cb53e2531
--- /dev/null
+++ b/sycl/source/image.cpp
@@ -0,0 +1,151 @@
+#include <detail/image_impl.hpp>
+#include <sycl/event.hpp>
+#include <sycl/ext/oneapi/accessor_property_list.hpp>
+#include <sycl/image.hpp>
+#include <sycl/properties/all_properties.hpp>
+#include <sycl/property_list.hpp>
+#include <sycl/range.hpp>
+
+#include <memory>
+
+namespace sycl {
+__SYCL_INLINE_VER_NAMESPACE(_V1) {
+namespace detail {
+
+image_plain::image_plain(image_channel_order Order, image_channel_type Type,
+                         const range<3> &Range,
+                         std::unique_ptr<SYCLMemObjAllocator> Allocator,
+                         uint8_t Dimensions, const property_list &PropList) {
+  impl = std::make_shared<detail::image_impl>(
+      Order, Type, Range, std::move(Allocator), Dimensions, PropList);
+}
+
+image_plain::image_plain(image_channel_order Order, image_channel_type Type,
+                         const range<3> &Range, const range<2> &Pitch,
+                         std::unique_ptr<SYCLMemObjAllocator> Allocator,
+                         uint8_t Dimensions, const property_list &PropList) {
+  impl = std::make_shared<detail::image_impl>(
+      Order, Type, Range, Pitch, std::move(Allocator), Dimensions, PropList);
+}
+
+image_plain::image_plain(void *HostPointer, image_channel_order Order,
+                         image_channel_type Type, const range<3> &Range,
+                         std::unique_ptr<SYCLMemObjAllocator> Allocator,
+                         uint8_t Dimensions, const property_list &PropList) {
+  impl = std::make_shared<detail::image_impl>(HostPointer, Order, Type, Range,
+                                              std::move(Allocator), Dimensions,
+                                              PropList);
+}
+
+image_plain::image_plain(const void *HostPointer, image_channel_order Order,
+                         image_channel_type Type, const range<3> &Range,
+                         std::unique_ptr<SYCLMemObjAllocator> Allocator,
+                         uint8_t Dimensions, const property_list &PropList) {
+  impl = std::make_shared<detail::image_impl>(HostPointer, Order, Type, Range,
+                                              std::move(Allocator), Dimensions,
+                                              PropList);
+}
+
+image_plain::image_plain(void *HostPointer, image_channel_order Order,
+                         image_channel_type Type, const range<3> &Range,
+                         const range<2> &Pitch,
+                         std::unique_ptr<SYCLMemObjAllocator> Allocator,
+                         uint8_t Dimensions, const property_list &PropList) {
+  impl = std::make_shared<detail::image_impl>(HostPointer, Order, Type, Range,
+                                              Pitch, std::move(Allocator),
+                                              Dimensions, PropList);
+}
+
+image_plain::image_plain(const std::shared_ptr<const void> &HostPointer,
+                         image_channel_order Order, image_channel_type Type,
+                         const range<3> &Range,
+                         std::unique_ptr<SYCLMemObjAllocator> Allocator,
+                         uint8_t Dimensions, const property_list &PropList,
+                         bool IsConstPtr) {
+  impl = std::make_shared<detail::image_impl>(HostPointer, Order, Type, Range,
+                                              std::move(Allocator), Dimensions,
+                                              PropList, IsConstPtr);
+}
+
+image_plain::image_plain(const std::shared_ptr<const void> &HostPointer,
+                         image_channel_order Order, image_channel_type Type,
+                         const range<3> &Range, const range<2> &Pitch,
+                         std::unique_ptr<SYCLMemObjAllocator> Allocator,
+                         uint8_t Dimensions, const property_list &PropList,
+                         bool IsConstPtr) {
+  impl = std::make_shared<detail::image_impl>(HostPointer, Order, Type, Range,
+                                              Pitch, std::move(Allocator),
+                                              Dimensions, PropList, IsConstPtr);
+}
+
+#ifdef __SYCL_INTERNAL_API
+image_plain::image_plain(cl_mem ClMemObject, const context &SyclContext,
+                         event AvailableEvent,
+                         std::unique_ptr<SYCLMemObjAllocator> Allocator,
+                         uint8_t Dimensions) {
+  impl = std::make_shared<detail::image_impl>(ClMemObject, SyclContext,
+                                              AvailableEvent,
+                                              std::move(Allocator), Dimensions);
+}
+#endif
+
+#define __SYCL_PARAM_TRAITS_SPEC(param_type)                                   \
+  template <>                                                                  \
+  __SYCL_EXPORT bool image_plain::has_property<param_type>() const noexcept {  \
+    return impl->has_property<param_type>();                                   \
+  }
+#include <sycl/detail/properties_traits.def>
+
+#undef __SYCL_PARAM_TRAITS_SPEC
+
+#define __SYCL_PARAM_TRAITS_SPEC(param_type)                                   \
+  template <>                                                                  \
+  __SYCL_EXPORT param_type image_plain::get_property<param_type>() const {     \
+    return impl->get_property<param_type>();                                   \
+  }
+#include <sycl/detail/properties_traits.def>
+
+#undef __SYCL_PARAM_TRAITS_SPEC
+
+range<3> image_plain::get_range() const { return impl->get_range(); }
+
+range<2> image_plain::get_pitch() const { return impl->get_pitch(); }
+
+size_t image_plain::get_size() const { return impl->size(); }
+
+size_t image_plain::get_count() const { return impl->get_count(); }
+
+void image_plain::set_final_data_internal() { impl->set_final_data(nullptr); }
+
+void image_plain::set_final_data_internal(
+    const std::function<void(const std::function<void(void *const Ptr)> &)>
+        &FinalDataFunc) {
+  impl->set_final_data(FinalDataFunc);
+}
+
+void image_plain::set_write_back(bool NeedWriteBack) {
+  impl->set_write_back(NeedWriteBack);
+}
+
+const std::unique_ptr<SYCLMemObjAllocator> &
+image_plain::get_allocator_internal() const {
+  return impl->get_allocator_internal();
+}
+
+size_t image_plain::getElementSize() const { return impl->getElementSize(); }
+
+size_t image_plain::getRowPitch() const { return impl->getRowPitch(); }
+
+size_t image_plain::getSlicePitch() const { return impl->getSlicePitch(); }
+
+image_channel_order image_plain::getChannelOrder() const {
+  return impl->getChannelOrder();
+}
+
+image_channel_type image_plain::getChannelType() const {
+  return impl->getChannelType();
+}
+
+} // namespace detail
+} // __SYCL_INLINE_VER_NAMESPACE(_V1)
+} // namespace sycl
diff --git a/sycl/test/abi/layout_buffer.cpp b/sycl/test/abi/layout_buffer.cpp
index c7c076a82d025..43af08dc0d3fb 100644
--- a/sycl/test/abi/layout_buffer.cpp
+++ b/sycl/test/abi/layout_buffer.cpp
@@ -7,79 +7,9 @@
 #include <sycl/buffer.hpp>
 
 void foo(sycl::buffer<int, 2>) {}
-// CHECK: 0 | class sycl::detail::buffer_impl
-// CHECK-NEXT: 0 |   class sycl::detail::SYCLMemObjT (primary base)
-// CHECK-NEXT:  0 |     class sycl::detail::SYCLMemObjI (primary base)
-// CHECK-NEXT:  0 |       (SYCLMemObjI vtable pointer)
-// CHECK-NEXT:  8 |       class std::shared_ptr<struct sycl::detail::MemObjRecord> MRecord
-// CHECK-NEXT:  8 |         class std::__shared_ptr<struct sycl::detail::MemObjRecord, __gnu_cxx::_S_atomic> (base)
-// CHECK-NEXT:  8 |           class std::__shared_ptr_access<struct sycl::detail::MemObjRecord, __gnu_cxx::_S_atomic, false, false> (base) (empty)
-// CHECK-NEXT:  8 |           element_type * _M_ptr
-// CHECK-NEXT:  16 |           class std::__shared_count<__gnu_cxx::_S_atomic> _M_refcount
-// CHECK-NEXT:  16 |             _Sp_counted_base<(_Lock_policy)2U> * _M_pi
-// CHECK-NEXT:  24 |     class std::unique_ptr<class sycl::detail::SYCLMemObjAllocator> MAllocator
-// CHECK:  24 |         class std::__uniq_ptr_impl<class sycl::detail::SYCLMemObjAllocator, struct std::default_delete<class sycl::detail::SYCLMemObjAllocator> >
-// CHECK-NEXT:  24 |           class std::tuple<class sycl::detail::SYCLMemObjAllocator *, struct std::default_delete<class sycl::detail::SYCLMemObjAllocator> > _M_t
-// CHECK-NEXT:  24 |             struct std::_Tuple_impl<0, class sycl::detail::SYCLMemObjAllocator *, struct std::default_delete<class sycl::detail::SYCLMemObjAllocator> > (base)
-// CHECK-NEXT:  24 |               struct std::_Tuple_impl<1, struct std::default_delete<class sycl::detail::SYCLMemObjAllocator> > (base) (empty)
-// CHECK-NEXT:  24 |                 struct std::_Head_base<1, struct std::default_delete<class sycl::detail::SYCLMemObjAllocator>, true> (base) (empty)
-// CHECK-NEXT:  24 |                   struct std::default_delete<class sycl::detail::SYCLMemObjAllocator>
-// CHECK-NEXT:  24 |               struct std::_Head_base<0, class sycl::detail::SYCLMemObjAllocator *, false> (base)
-// CHECK-NEXT:  24 |                 class sycl::detail::SYCLMemObjAllocator * _M_head_impl
-// CHECK-NEXT:  32 |     class sycl::property_list MProps
-// CHECK-NEXT:  32 |       class sycl::detail::PropertyListBase (base)
-// CHECK-NEXT:  32 |         class std::bitset<32> MDataLessProps
-// CHECK-NEXT:  32 |           struct std::_Base_bitset<1> (base)
-// CHECK-NEXT:  32 |             _WordT _M_w
-// CHECK-NEXT:  40 |         class std::vector<class std::shared_ptr<class sycl::detail::PropertyWithDataBase> > MPropsWithData
-// CHECK-NEXT:  40 |           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:  40 |             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:  40 |               class std::allocator<class std::shared_ptr<class sycl::detail::PropertyWithDataBase> > (base) (empty)
-// CHECK-NEXT:  40 |                 class __gnu_cxx::new_allocator<class std::shared_ptr<class sycl::detail::PropertyWithDataBase> > (base) (empty)
-// CHECK:  40 |                 pointer _M_start
-// CHECK-NEXT:  48 |                 pointer _M_finish
-// CHECK-NEXT:  56 |                 pointer _M_end_of_storage
-// CHECK-NEXT:  64 |     class std::shared_ptr<class sycl::detail::event_impl> MInteropEvent
-// CHECK-NEXT:  64 |       class std::__shared_ptr<class sycl::detail::event_impl, __gnu_cxx::_S_atomic> (base)
-// CHECK-NEXT:  64 |         class std::__shared_ptr_access<class sycl::detail::event_impl, __gnu_cxx::_S_atomic, false, false> (base) (empty)
-// CHECK-NEXT:  64 |         element_type * _M_ptr
-// CHECK-NEXT:  72 |         class std::__shared_count<__gnu_cxx::_S_atomic> _M_refcount
-// CHECK-NEXT:  72 |           _Sp_counted_base<(_Lock_policy)2U> * _M_pi
-// CHECK-NEXT:  80 |     class std::shared_ptr<class sycl::detail::context_impl> MInteropContext
-// CHECK-NEXT:  80 |       class std::__shared_ptr<class sycl::detail::context_impl, __gnu_cxx::_S_atomic> (base)
-// CHECK-NEXT:  80 |         class std::__shared_ptr_access<class sycl::detail::context_impl, __gnu_cxx::_S_atomic, false, false> (base) (empty)
-// CHECK-NEXT:  80 |         element_type * _M_ptr
-// CHECK-NEXT:  88 |         class std::__shared_count<__gnu_cxx::_S_atomic> _M_refcount
-// CHECK-NEXT:  88 |           _Sp_counted_base<(_Lock_policy)2U> * _M_pi
-// CHECK-NEXT:  96 |     RT::PiMem MInteropMemObject
-// CHECK-NEXT:  104 |     _Bool MOpenCLInterop
-// CHECK-NEXT:  105 |     _Bool MHostPtrReadOnly
-// CHECK-NEXT:  106 |     _Bool MNeedWriteBack
-// CHECK-NEXT:  112 |     size_t MSizeInBytes
-// CHECK-NEXT:  120 |     void * MUserPtr
-// CHECK-NEXT:  128 |     void * MShadowCopy
-// CHECK-NEXT:  136 |     class std::function<void (void)> MUploadDataFunctor
-// CHECK-NEXT:  136 |       struct std::_Maybe_unary_or_binary_function<void> (base) (empty)
-// CHECK-NEXT:  136 |       class std::_Function_base (base)
-// CHECK-NEXT:  136 |         union std::_Any_data _M_functor
-// CHECK-NEXT:  136 |           union std::_Nocopy_types _M_unused
-// CHECK-NEXT:  136 |             void * _M_object
-// CHECK-NEXT:  136 |             const void * _M_const_object
-// CHECK-NEXT:  136 |             void (*)(void) _M_function_pointer
-// CHECK-NEXT:  136 |             void (class std::_Undefined_class::*)(void) _M_member_pointer
-// CHECK-NEXT:  136 |           char[16] _M_pod_data
-// CHECK-NEXT:  152 |         _Manager_type _M_manager
-// CHECK-NEXT:  160 |       _Invoker_type _M_invoker
-// CHECK-NEXT:  168 |     class std::shared_ptr<const void> MSharedPtrStorage
-// CHECK-NEXT:  168 |       class std::__shared_ptr<const void, __gnu_cxx::_S_atomic> (base)
-// CHECK-NEXT:  168 |         class std::__shared_ptr_access<const void, __gnu_cxx::_S_atomic, false, true> (base) (empty)
-// CHECK-NEXT:  168 |         element_type * _M_ptr
-// CHECK-NEXT:  176 |         class std::__shared_count<__gnu_cxx::_S_atomic> _M_refcount
-// CHECK-NEXT:  176 |           _Sp_counted_base<(_Lock_policy)2U> * _M_pi
-// CHECK-NEXT:      | [sizeof=184, dsize=184, align=8,
-// CHECK-NEXT:      |  nvsize=184, nvalign=8]
 
 // CHECK: 0 | class sycl::buffer<int, 2, class sycl::detail::aligned_allocator<int>, void>
+// CHECK: 0 | class sycl::detail::buffer_plain (base)
 // CHECK-NEXT: 0 |   class std::shared_ptr<class sycl::detail::buffer_impl> impl
 // CHECK-NEXT: 0 |     class std::__shared_ptr<class sycl::detail::buffer_impl, __gnu_cxx::_S_atomic> (base)
 // CHECK-NEXT: 0 |       class std::__shared_ptr_access<class sycl::detail::buffer_impl, __gnu_cxx::_S_atomic, false, false> (base) (empty)
diff --git a/sycl/test/abi/layout_image.cpp b/sycl/test/abi/layout_image.cpp
index e514ca0c622ff..854cd361d45d8 100644
--- a/sycl/test/abi/layout_image.cpp
+++ b/sycl/test/abi/layout_image.cpp
@@ -8,92 +8,11 @@
 
 sycl::image<2> Img{sycl::image_channel_order::rgba, sycl::image_channel_type::fp16, sycl::range<2>{10, 10}};
 
-// CHECK: 0 | class sycl::detail::image_impl<2>
-// CHECK-NEXT: 0 |   class sycl::detail::SYCLMemObjT (primary base)
-// CHECK-NEXT: 0 |     class sycl::detail::SYCLMemObjI (primary base)
-// CHECK-NEXT: 0 |       (SYCLMemObjI vtable pointer)
-// CHECK-NEXT: 8 |       class std::shared_ptr<struct sycl::detail::MemObjRecord> MRecord
-// CHECK-NEXT: 8 |         class std::__shared_ptr<struct sycl::detail::MemObjRecord, __gnu_cxx::_S_atomic> (base)
-// CHECK-NEXT: 8 |           class std::__shared_ptr_access<struct sycl::detail::MemObjRecord, __gnu_cxx::_S_atomic, false, false> (base) (empty)
-// CHECK-NEXT: 8 |           element_type * _M_ptr
-// CHECK-NEXT: 16 |           class std::__shared_count<__gnu_cxx::_S_atomic> _M_refcount
-// CHECK-NEXT: 16 |             _Sp_counted_base<(_Lock_policy)2U> * _M_pi
-// CHECK-NEXT: 24 |     class std::unique_ptr<class sycl::detail::SYCLMemObjAllocator> MAllocator
-// CHECK: 24 |         class std::__uniq_ptr_impl<class sycl::detail::SYCLMemObjAllocator, struct std::default_delete<class sycl::detail::SYCLMemObjAllocator> >
-// CHECK-NEXT: 24 |           class std::tuple<class sycl::detail::SYCLMemObjAllocator *, struct std::default_delete<class sycl::detail::SYCLMemObjAllocator> > _M_t
-// CHECK-NEXT: 24 |             struct std::_Tuple_impl<0, class sycl::detail::SYCLMemObjAllocator *, struct std::default_delete<class sycl::detail::SYCLMemObjAllocator> > (base)
-// CHECK-NEXT: 24 |               struct std::_Tuple_impl<1, struct std::default_delete<class sycl::detail::SYCLMemObjAllocator> > (base) (empty)
-// CHECK-NEXT: 24 |                 struct std::_Head_base<1, struct std::default_delete<class sycl::detail::SYCLMemObjAllocator>, true> (base) (empty)
-// CHECK-NEXT: 24 |                   struct std::default_delete<class sycl::detail::SYCLMemObjAllocator>
-// CHECK-NEXT: 24 |               struct std::_Head_base<0, class sycl::detail::SYCLMemObjAllocator *, false> (base)
-// CHECK-NEXT: 24 |                 class sycl::detail::SYCLMemObjAllocator * _M_head_impl
-// CHECK-NEXT: 32 |     class sycl::property_list MProps
-// CHECK-NEXT: 32 |       class sycl::detail::PropertyListBase (base)
-// CHECK-NEXT: 32 |         class std::bitset<32> MDataLessProps
-// CHECK-NEXT: 32 |           struct std::_Base_bitset<1> (base)
-// CHECK-NEXT: 32 |             _WordT _M_w
-// CHECK-NEXT: 40 |         class std::vector<class std::shared_ptr<class sycl::detail::PropertyWithDataBase> > MPropsWithData
-// CHECK-NEXT: 40 |           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: 40 |             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: 40 |               class std::allocator<class std::shared_ptr<class sycl::detail::PropertyWithDataBase> > (base) (empty)
-// CHECK-NEXT: 40 |                 class __gnu_cxx::new_allocator<class std::shared_ptr<class sycl::detail::PropertyWithDataBase> > (base) (empty)
-// CHECK: 40 |                 pointer _M_start
-// CHECK-NEXT: 48 |                 pointer _M_finish
-// CHECK-NEXT: 56 |                 pointer _M_end_of_storage
-// CHECK-NEXT: 64 |     class std::shared_ptr<class sycl::detail::event_impl> MInteropEvent
-// CHECK-NEXT: 64 |       class std::__shared_ptr<class sycl::detail::event_impl, __gnu_cxx::_S_atomic> (base)
-// CHECK-NEXT: 64 |         class std::__shared_ptr_access<class sycl::detail::event_impl, __gnu_cxx::_S_atomic, false, false> (base) (empty)
-// CHECK-NEXT: 64 |         element_type * _M_ptr
-// CHECK-NEXT: 72 |         class std::__shared_count<__gnu_cxx::_S_atomic> _M_refcount
-// CHECK-NEXT: 72 |           _Sp_counted_base<(_Lock_policy)2U> * _M_pi
-// CHECK-NEXT: 80 |     class std::shared_ptr<class sycl::detail::context_impl> MInteropContext
-// CHECK-NEXT: 80 |       class std::__shared_ptr<class sycl::detail::context_impl, __gnu_cxx::_S_atomic> (base)
-// CHECK-NEXT: 80 |         class std::__shared_ptr_access<class sycl::detail::context_impl, __gnu_cxx::_S_atomic, false, false> (base) (empty)
-// CHECK-NEXT: 80 |         element_type * _M_ptr
-// CHECK-NEXT: 88 |         class std::__shared_count<__gnu_cxx::_S_atomic> _M_refcount
-// CHECK-NEXT: 88 |           _Sp_counted_base<(_Lock_policy)2U> * _M_pi
-// CHECK-NEXT: 96 |     RT::PiMem MInteropMemObject
-// CHECK-NEXT: 104 |     _Bool MOpenCLInterop
-// CHECK-NEXT: 105 |     _Bool MHostPtrReadOnly
-// CHECK-NEXT: 106 |     _Bool MNeedWriteBack
-// CHECK-NEXT: 112 |     size_t MSizeInBytes
-// CHECK-NEXT: 120 |     void * MUserPtr
-// CHECK-NEXT: 128 |     void * MShadowCopy
-// CHECK-NEXT: 136 |     class std::function<void (void)> MUploadDataFunctor
-// CHECK-NEXT: 136 |       struct std::_Maybe_unary_or_binary_function<void> (base) (empty)
-// CHECK-NEXT: 136 |       class std::_Function_base (base)
-// CHECK-NEXT: 136 |         union std::_Any_data _M_functor
-// CHECK-NEXT: 136 |           union std::_Nocopy_types _M_unused
-// CHECK-NEXT: 136 |             void * _M_object
-// CHECK-NEXT: 136 |             const void * _M_const_object
-// CHECK-NEXT: 136 |             void (*)(void) _M_function_pointer
-// CHECK-NEXT: 136 |             void (class std::_Undefined_class::*)(void) _M_member_pointer
-// CHECK-NEXT: 136 |           char[16] _M_pod_data
-// CHECK-NEXT: 152 |         _Manager_type _M_manager
-// CHECK-NEXT: 160 |       _Invoker_type _M_invoker
-// CHECK-NEXT: 168 |     class std::shared_ptr<const void> MSharedPtrStorage
-// CHECK-NEXT: 168 |       class std::__shared_ptr<const void, __gnu_cxx::_S_atomic> (base)
-// CHECK-NEXT: 168 |         class std::__shared_ptr_access<const void, __gnu_cxx::_S_atomic, false, true> (base) (empty)
-// CHECK-NEXT: 168 |         element_type * _M_ptr
-// CHECK-NEXT: 176 |         class std::__shared_count<__gnu_cxx::_S_atomic> _M_refcount
-// CHECK-NEXT: 176 |           _Sp_counted_base<(_Lock_policy)2U> * _M_pi
-// CHECK-NEXT: 184 |   _Bool MIsArrayImage
-// CHECK-NEXT: 192 |   class sycl::range<2> MRange
-// CHECK-NEXT: 192 |     class sycl::detail::array<2> (base)
-// CHECK-NEXT: 192 |       size_t[2] common_array
-// CHECK-NEXT: 208 |   image_channel_order MOrder
-// CHECK-NEXT: 212 |   image_channel_type MType
-// CHECK-NEXT: 216 |   uint8_t MNumChannels
-// CHECK-NEXT: 224 |   size_t MElementSize
-// CHECK-NEXT: 232 |   size_t MRowPitch
-// CHECK-NEXT: 240 |   size_t MSlicePitch
-// CHECK-NEXT:     | [sizeof=248, dsize=248, align=8,
-// CHECK-NEXT:     |  nvsize=248, nvalign=8]
-
 // CHECK: 0 | class sycl::image<2>
-// CHECK-NEXT: 0 |   class std::shared_ptr<class sycl::detail::image_impl<2> > impl
-// CHECK-NEXT: 0 |     class std::__shared_ptr<class sycl::detail::image_impl<2>, __gnu_cxx::_S_atomic> (base)
-// CHECK-NEXT: 0 |       class std::__shared_ptr_access<class sycl::detail::image_impl<2>, __gnu_cxx::_S_atomic, false, false> (base) (empty)
+// CHECK-NEXT: 0 | class sycl::detail::image_plain (base)
+// CHECK-NEXT: 0 |   class std::shared_ptr<class sycl::detail::image_impl> impl
+// CHECK-NEXT: 0 |     class std::__shared_ptr<class sycl::detail::image_impl, __gnu_cxx::_S_atomic> (base)
+// CHECK-NEXT: 0 |       class std::__shared_ptr_access<class sycl::detail::image_impl, __gnu_cxx::_S_atomic, false, false> (base) (empty)
 // CHECK-NEXT: 0 |       element_type * _M_ptr
 // CHECK-NEXT: 8 |       class std::__shared_count<__gnu_cxx::_S_atomic> _M_refcount
 // CHECK-NEXT: 8 |         _Sp_counted_base<(_Lock_policy)2U> * _M_pi
diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump
index 6a62770fca038..67ad8af4dec16 100644
--- a/sycl/test/abi/sycl_symbols_linux.dump
+++ b/sycl/test/abi/sycl_symbols_linux.dump
@@ -3715,69 +3715,12 @@ _ZN4sycl3_V16ONEAPI15filter_selectorC2ERKNSt7__cxx1112basic_stringIcSt11char_tra
 _ZN4sycl3_V16ONEAPI6detail16reduGetMaxWGSizeESt10shared_ptrINS0_6detail10queue_implEEm
 _ZN4sycl3_V16ONEAPI6detail17reduComputeWGSizeEmmRm
 _ZN4sycl3_V16detail10build_implERKNS0_13kernel_bundleILNS0_12bundle_stateE0EEERKSt6vectorINS0_6deviceESaIS8_EERKNS0_13property_listE
-_ZN4sycl3_V16detail10image_implILi1EE10getDevicesESt10shared_ptrINS1_12context_implEE
-_ZN4sycl3_V16detail10image_implILi1EE10setPitchesEv
-_ZN4sycl3_V16detail10image_implILi1EE11allocateMemESt10shared_ptrINS1_12context_implEEbPvRP9_pi_event
-_ZN4sycl3_V16detail10image_implILi1EE12getImageDescEb
-_ZN4sycl3_V16detail10image_implILi1EE12getImageTypeEv
-_ZN4sycl3_V16detail10image_implILi1EE14checkImageDescERK14_pi_image_descSt10shared_ptrINS1_12context_implEEPv
-_ZN4sycl3_V16detail10image_implILi1EE14getImageFormatEv
-_ZN4sycl3_V16detail10image_implILi1EE16checkImageFormatERK16_pi_image_formatSt10shared_ptrINS1_12context_implEE
-_ZN4sycl3_V16detail10image_implILi1EEC1ENS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi1EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISB_EERKNS0_13property_listE
-_ZN4sycl3_V16detail10image_implILi1EEC1EP7_cl_memRKNS0_7contextENS0_5eventESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISB_EE
-_ZN4sycl3_V16detail10image_implILi1EEC1EPKvNS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi1EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISD_EERKNS0_13property_listE
-_ZN4sycl3_V16detail10image_implILi1EEC1EPvNS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi1EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISC_EERKNS0_13property_listE
-_ZN4sycl3_V16detail10image_implILi1EEC1ERSt10shared_ptrIvENS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi1EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISE_EERKNS0_13property_listE
-_ZN4sycl3_V16detail10image_implILi1EEC2ENS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi1EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISB_EERKNS0_13property_listE
-_ZN4sycl3_V16detail10image_implILi1EEC2EP7_cl_memRKNS0_7contextENS0_5eventESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISB_EE
-_ZN4sycl3_V16detail10image_implILi1EEC2EPKvNS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi1EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISD_EERKNS0_13property_listE
-_ZN4sycl3_V16detail10image_implILi1EEC2EPvNS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi1EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISC_EERKNS0_13property_listE
-_ZN4sycl3_V16detail10image_implILi1EEC2ERSt10shared_ptrIvENS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi1EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISE_EERKNS0_13property_listE
-_ZN4sycl3_V16detail10image_implILi1EED0Ev
-_ZN4sycl3_V16detail10image_implILi1EED1Ev
-_ZN4sycl3_V16detail10image_implILi1EED2Ev
-_ZN4sycl3_V16detail10image_implILi2EE10getDevicesESt10shared_ptrINS1_12context_implEE
-_ZN4sycl3_V16detail10image_implILi2EE10setPitchesEv
-_ZN4sycl3_V16detail10image_implILi2EE11allocateMemESt10shared_ptrINS1_12context_implEEbPvRP9_pi_event
-_ZN4sycl3_V16detail10image_implILi2EE12getImageDescEb
-_ZN4sycl3_V16detail10image_implILi2EE12getImageTypeEv
-_ZN4sycl3_V16detail10image_implILi2EE14checkImageDescERK14_pi_image_descSt10shared_ptrINS1_12context_implEEPv
-_ZN4sycl3_V16detail10image_implILi2EE14getImageFormatEv
-_ZN4sycl3_V16detail10image_implILi2EE16checkImageFormatERK16_pi_image_formatSt10shared_ptrINS1_12context_implEE
-_ZN4sycl3_V16detail10image_implILi2EEC1ENS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi2EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISB_EERKNS0_13property_listE
-_ZN4sycl3_V16detail10image_implILi2EEC1EP7_cl_memRKNS0_7contextENS0_5eventESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISB_EE
-_ZN4sycl3_V16detail10image_implILi2EEC1EPKvNS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi2EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISD_EERKNS0_13property_listE
-_ZN4sycl3_V16detail10image_implILi2EEC1EPvNS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi2EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISC_EERKNS0_13property_listE
-_ZN4sycl3_V16detail10image_implILi2EEC1ERSt10shared_ptrIvENS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi2EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISE_EERKNS0_13property_listE
-_ZN4sycl3_V16detail10image_implILi2EEC2ENS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi2EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISB_EERKNS0_13property_listE
-_ZN4sycl3_V16detail10image_implILi2EEC2EP7_cl_memRKNS0_7contextENS0_5eventESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISB_EE
-_ZN4sycl3_V16detail10image_implILi2EEC2EPKvNS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi2EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISD_EERKNS0_13property_listE
-_ZN4sycl3_V16detail10image_implILi2EEC2EPvNS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi2EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISC_EERKNS0_13property_listE
-_ZN4sycl3_V16detail10image_implILi2EEC2ERSt10shared_ptrIvENS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi2EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISE_EERKNS0_13property_listE
-_ZN4sycl3_V16detail10image_implILi2EED0Ev
-_ZN4sycl3_V16detail10image_implILi2EED1Ev
-_ZN4sycl3_V16detail10image_implILi2EED2Ev
-_ZN4sycl3_V16detail10image_implILi3EE10getDevicesESt10shared_ptrINS1_12context_implEE
-_ZN4sycl3_V16detail10image_implILi3EE10setPitchesEv
-_ZN4sycl3_V16detail10image_implILi3EE11allocateMemESt10shared_ptrINS1_12context_implEEbPvRP9_pi_event
-_ZN4sycl3_V16detail10image_implILi3EE12getImageDescEb
-_ZN4sycl3_V16detail10image_implILi3EE12getImageTypeEv
-_ZN4sycl3_V16detail10image_implILi3EE14checkImageDescERK14_pi_image_descSt10shared_ptrINS1_12context_implEEPv
-_ZN4sycl3_V16detail10image_implILi3EE14getImageFormatEv
-_ZN4sycl3_V16detail10image_implILi3EE16checkImageFormatERK16_pi_image_formatSt10shared_ptrINS1_12context_implEE
-_ZN4sycl3_V16detail10image_implILi3EEC1ENS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi3EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISB_EERKNS0_13property_listE
-_ZN4sycl3_V16detail10image_implILi3EEC1EP7_cl_memRKNS0_7contextENS0_5eventESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISB_EE
-_ZN4sycl3_V16detail10image_implILi3EEC1EPKvNS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi3EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISD_EERKNS0_13property_listE
-_ZN4sycl3_V16detail10image_implILi3EEC1EPvNS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi3EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISC_EERKNS0_13property_listE
-_ZN4sycl3_V16detail10image_implILi3EEC1ERSt10shared_ptrIvENS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi3EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISE_EERKNS0_13property_listE
-_ZN4sycl3_V16detail10image_implILi3EEC2ENS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi3EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISB_EERKNS0_13property_listE
-_ZN4sycl3_V16detail10image_implILi3EEC2EP7_cl_memRKNS0_7contextENS0_5eventESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISB_EE
-_ZN4sycl3_V16detail10image_implILi3EEC2EPKvNS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi3EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISD_EERKNS0_13property_listE
-_ZN4sycl3_V16detail10image_implILi3EEC2EPvNS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi3EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISC_EERKNS0_13property_listE
-_ZN4sycl3_V16detail10image_implILi3EEC2ERSt10shared_ptrIvENS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi3EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISE_EERKNS0_13property_listE
-_ZN4sycl3_V16detail10image_implILi3EED0Ev
-_ZN4sycl3_V16detail10image_implILi3EED1Ev
-_ZN4sycl3_V16detail10image_implILi3EED2Ev
+_ZN4sycl3_V16detail10image_impl10getDevicesESt10shared_ptrINS1_12context_implEE
+_ZN4sycl3_V16detail10image_impl11allocateMemESt10shared_ptrINS1_12context_implEEbPvRP9_pi_event
+_ZN4sycl3_V16detail10image_impl14checkImageDescERK14_pi_image_descSt10shared_ptrINS1_12context_implEEPv
+_ZN4sycl3_V16detail10image_impl16checkImageFormatERK16_pi_image_formatSt10shared_ptrINS1_12context_implEE
+_ZN4sycl3_V16detail10image_implC1EP7_cl_memRKNS0_7contextENS0_5eventESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISA_EEh
+_ZN4sycl3_V16detail10image_implC2EP7_cl_memRKNS0_7contextENS0_5eventESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISA_EEh
 _ZN4sycl3_V16detail10make_eventEmRKNS0_7contextENS0_7backendE
 _ZN4sycl3_V16detail10make_eventEmRKNS0_7contextEbNS0_7backendE
 _ZN4sycl3_V16detail10make_queueEmRKNS0_7contextERKNS0_6deviceEbRKSt8functionIFvNS0_14exception_listEEENS0_7backendE
@@ -3795,8 +3738,26 @@ _ZN4sycl3_V16detail11SYCLMemObjTC2EmRKNS0_7contextEbNS0_5eventESt10unique_ptrINS
 _ZN4sycl3_V16detail11SYCLMemObjTC2EmRKNS0_7contextEmNS0_5eventESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteIS8_EE
 _ZN4sycl3_V16detail11buffer_impl11allocateMemESt10shared_ptrINS1_12context_implEEbPvRP9_pi_event
 _ZN4sycl3_V16detail11buffer_impl22destructorNotificationEPv
-_ZN4sycl3_V16detail11buffer_impl23constructorNotificationERKNS1_13code_locationEPv
 _ZN4sycl3_V16detail11buffer_impl23constructorNotificationERKNS1_13code_locationEPvPKvS8_jjPm
+_ZN4sycl3_V16detail11image_plain14set_write_backEb
+_ZN4sycl3_V16detail11image_plain23set_final_data_internalERKSt8functionIFvRKS3_IFvPvEEEE
+_ZN4sycl3_V16detail11image_plain23set_final_data_internalEv
+_ZN4sycl3_V16detail11image_plainC1ENS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi3EEERKNS5_ILi2EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISD_EEhRKNS0_13property_listE
+_ZN4sycl3_V16detail11image_plainC1ENS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi3EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISA_EEhRKNS0_13property_listE
+_ZN4sycl3_V16detail11image_plainC1EP7_cl_memRKNS0_7contextENS0_5eventESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISA_EEh
+_ZN4sycl3_V16detail11image_plainC1EPKvNS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi3EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISC_EEhRKNS0_13property_listE
+_ZN4sycl3_V16detail11image_plainC1EPvNS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi3EEERKNS6_ILi2EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISE_EEhRKNS0_13property_listE
+_ZN4sycl3_V16detail11image_plainC1EPvNS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi3EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISB_EEhRKNS0_13property_listE
+_ZN4sycl3_V16detail11image_plainC1ERKSt10shared_ptrIKvENS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi3EEERKNSA_ILi2EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISI_EEhRKNS0_13property_listEb
+_ZN4sycl3_V16detail11image_plainC1ERKSt10shared_ptrIKvENS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi3EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISF_EEhRKNS0_13property_listEb
+_ZN4sycl3_V16detail11image_plainC2ENS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi3EEERKNS5_ILi2EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISD_EEhRKNS0_13property_listE
+_ZN4sycl3_V16detail11image_plainC2ENS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi3EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISA_EEhRKNS0_13property_listE
+_ZN4sycl3_V16detail11image_plainC2EP7_cl_memRKNS0_7contextENS0_5eventESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISA_EEh
+_ZN4sycl3_V16detail11image_plainC2EPKvNS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi3EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISC_EEhRKNS0_13property_listE
+_ZN4sycl3_V16detail11image_plainC2EPvNS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi3EEERKNS6_ILi2EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISE_EEhRKNS0_13property_listE
+_ZN4sycl3_V16detail11image_plainC2EPvNS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi3EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISB_EEhRKNS0_13property_listE
+_ZN4sycl3_V16detail11image_plainC2ERKSt10shared_ptrIKvENS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi3EEERKNSA_ILi2EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISI_EEhRKNS0_13property_listEb
+_ZN4sycl3_V16detail11image_plainC2ERKSt10shared_ptrIKvENS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi3EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISF_EEhRKNS0_13property_listEb
 _ZN4sycl3_V16detail11make_deviceEmNS0_7backendE
 _ZN4sycl3_V16detail11make_kernelERKNS0_7contextERKNS0_13kernel_bundleILNS0_12bundle_stateE2EEEmbNS0_7backendE
 _ZN4sycl3_V16detail11make_kernelEmRKNS0_7contextENS0_7backendE
@@ -3808,6 +3769,24 @@ _ZN4sycl3_V16detail11stream_implC1EmmRKNS0_13property_listE
 _ZN4sycl3_V16detail11stream_implC1EmmRNS0_7handlerE
 _ZN4sycl3_V16detail11stream_implC2EmmRKNS0_13property_listE
 _ZN4sycl3_V16detail11stream_implC2EmmRNS0_7handlerE
+_ZN4sycl3_V16detail12buffer_plain14deleteAccPropsERKNS1_16PropWithDataKindE
+_ZN4sycl3_V16detail12buffer_plain14set_write_backEb
+_ZN4sycl3_V16detail12buffer_plain23constructorNotificationERKNS1_13code_locationEPvPKvS8_jjPm
+_ZN4sycl3_V16detail12buffer_plain23set_final_data_internalERKSt8functionIFvRKS3_IFvPvEEEE
+_ZN4sycl3_V16detail12buffer_plain23set_final_data_internalEv
+_ZN4sycl3_V16detail12buffer_plain30addOrReplaceAccessorPropertiesERKNS0_13property_listE
+_ZN4sycl3_V16detail12buffer_plainC1EPKvmmRKNS0_13property_listESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteIS9_EE
+_ZN4sycl3_V16detail12buffer_plainC1EPvmmRKNS0_13property_listESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteIS8_EE
+_ZN4sycl3_V16detail12buffer_plainC1ERKSt10shared_ptrIKvEmmRKNS0_13property_listESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISC_EEb
+_ZN4sycl3_V16detail12buffer_plainC1ERKSt8functionIFvPvEEmmRKNS0_13property_listESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISD_EEb
+_ZN4sycl3_V16detail12buffer_plainC1EmNS0_7contextESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteIS5_EEbNS0_5eventE
+_ZN4sycl3_V16detail12buffer_plainC1EmmRKNS0_13property_listESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteIS7_EE
+_ZN4sycl3_V16detail12buffer_plainC2EPKvmmRKNS0_13property_listESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteIS9_EE
+_ZN4sycl3_V16detail12buffer_plainC2EPvmmRKNS0_13property_listESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteIS8_EE
+_ZN4sycl3_V16detail12buffer_plainC2ERKSt10shared_ptrIKvEmmRKNS0_13property_listESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISC_EEb
+_ZN4sycl3_V16detail12buffer_plainC2ERKSt8functionIFvPvEEmmRKNS0_13property_listESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISD_EEb
+_ZN4sycl3_V16detail12buffer_plainC2EmNS0_7contextESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteIS5_EEbNS0_5eventE
+_ZN4sycl3_V16detail12buffer_plainC2EmmRKNS0_13property_listESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteIS7_EE
 _ZN4sycl3_V16detail12compile_implERKNS0_13kernel_bundleILNS0_12bundle_stateE0EEERKSt6vectorINS0_6deviceESaIS8_EERKNS0_13property_listE
 _ZN4sycl3_V16detail12isOutOfRangeENS0_3vecIiLi4EEENS0_15addressing_modeENS0_5rangeILi3EEE
 _ZN4sycl3_V16detail12make_contextEmRKSt8functionIFvNS0_14exception_listEEENS0_7backendE
@@ -4070,39 +4049,77 @@ _ZNK4sycl3_V15queue9getNativeEv
 _ZNK4sycl3_V16ONEAPI15filter_selector13select_deviceEv
 _ZNK4sycl3_V16ONEAPI15filter_selector5resetEv
 _ZNK4sycl3_V16ONEAPI15filter_selectorclERKNS0_6deviceE
-_ZNK4sycl3_V16detail10image_implILi1EE11getRowPitchEv
-_ZNK4sycl3_V16detail10image_implILi1EE13getSlicePitchEv
-_ZNK4sycl3_V16detail10image_implILi1EE14getChannelTypeEv
-_ZNK4sycl3_V16detail10image_implILi1EE14getElementSizeEv
-_ZNK4sycl3_V16detail10image_implILi1EE15getChannelOrderEv
-_ZNK4sycl3_V16detail10image_implILi1EE4sizeEv
-_ZNK4sycl3_V16detail10image_implILi1EE7getTypeEv
-_ZNK4sycl3_V16detail10image_implILi1EE9get_countEv
-_ZNK4sycl3_V16detail10image_implILi1EE9get_rangeEv
-_ZNK4sycl3_V16detail10image_implILi2EE11getRowPitchEv
-_ZNK4sycl3_V16detail10image_implILi2EE13getSlicePitchEv
-_ZNK4sycl3_V16detail10image_implILi2EE14getChannelTypeEv
-_ZNK4sycl3_V16detail10image_implILi2EE14getElementSizeEv
-_ZNK4sycl3_V16detail10image_implILi2EE15getChannelOrderEv
-_ZNK4sycl3_V16detail10image_implILi2EE4sizeEv
-_ZNK4sycl3_V16detail10image_implILi2EE7getTypeEv
-_ZNK4sycl3_V16detail10image_implILi2EE9get_countEv
-_ZNK4sycl3_V16detail10image_implILi2EE9get_rangeEv
-_ZNK4sycl3_V16detail10image_implILi3EE11getRowPitchEv
-_ZNK4sycl3_V16detail10image_implILi3EE13getSlicePitchEv
-_ZNK4sycl3_V16detail10image_implILi3EE14getChannelTypeEv
-_ZNK4sycl3_V16detail10image_implILi3EE14getElementSizeEv
-_ZNK4sycl3_V16detail10image_implILi3EE15getChannelOrderEv
-_ZNK4sycl3_V16detail10image_implILi3EE4sizeEv
-_ZNK4sycl3_V16detail10image_implILi3EE7getTypeEv
-_ZNK4sycl3_V16detail10image_implILi3EE9get_countEv
-_ZNK4sycl3_V16detail10image_implILi3EE9get_rangeEv
 _ZNK4sycl3_V16detail11SYCLMemObjT9getPluginEv
 _ZNK4sycl3_V16detail11SYCLMemObjT9isInteropEv
 _ZNK4sycl3_V16detail11buffer_impl15getNativeVectorENS0_7backendE
 _ZNK4sycl3_V16detail11buffer_impl16addInteropObjectERSt6vectorImSaImEE
+_ZNK4sycl3_V16detail11image_plain11getRowPitchEv
+_ZNK4sycl3_V16detail11image_plain12get_propertyINS0_3ext6oneapi4cuda8property7context19use_primary_contextEEET_v
+_ZNK4sycl3_V16detail11image_plain12get_propertyINS0_3ext6oneapi8property6buffer22use_pinned_host_memoryEEET_v
+_ZNK4sycl3_V16detail11image_plain12get_propertyINS0_8property5image12use_host_ptrEEET_v
+_ZNK4sycl3_V16detail11image_plain12get_propertyINS0_8property5image13context_boundEEET_v
+_ZNK4sycl3_V16detail11image_plain12get_propertyINS0_8property5image9use_mutexEEET_v
+_ZNK4sycl3_V16detail11image_plain12get_propertyINS0_8property5queue8in_orderEEET_v
+_ZNK4sycl3_V16detail11image_plain12get_propertyINS0_8property6buffer12use_host_ptrEEET_v
+_ZNK4sycl3_V16detail11image_plain12get_propertyINS0_8property6buffer13context_boundEEET_v
+_ZNK4sycl3_V16detail11image_plain12get_propertyINS0_8property6buffer9use_mutexEEET_v
+_ZNK4sycl3_V16detail11image_plain12get_propertyINS0_8property6noinitEEET_v
+_ZNK4sycl3_V16detail11image_plain12get_propertyINS0_8property7context4cuda19use_primary_contextEEET_v
+_ZNK4sycl3_V16detail11image_plain12get_propertyINS0_8property7no_initEEET_v
+_ZNK4sycl3_V16detail11image_plain12get_propertyINS0_8property9reduction22initialize_to_identityEEET_v
+_ZNK4sycl3_V16detail11image_plain12has_propertyINS0_3ext6oneapi4cuda8property7context19use_primary_contextEEEbv
+_ZNK4sycl3_V16detail11image_plain12has_propertyINS0_3ext6oneapi8property6buffer22use_pinned_host_memoryEEEbv
+_ZNK4sycl3_V16detail11image_plain12has_propertyINS0_8property5image12use_host_ptrEEEbv
+_ZNK4sycl3_V16detail11image_plain12has_propertyINS0_8property5image13context_boundEEEbv
+_ZNK4sycl3_V16detail11image_plain12has_propertyINS0_8property5image9use_mutexEEEbv
+_ZNK4sycl3_V16detail11image_plain12has_propertyINS0_8property5queue8in_orderEEEbv
+_ZNK4sycl3_V16detail11image_plain12has_propertyINS0_8property6buffer12use_host_ptrEEEbv
+_ZNK4sycl3_V16detail11image_plain12has_propertyINS0_8property6buffer13context_boundEEEbv
+_ZNK4sycl3_V16detail11image_plain12has_propertyINS0_8property6buffer9use_mutexEEEbv
+_ZNK4sycl3_V16detail11image_plain12has_propertyINS0_8property6noinitEEEbv
+_ZNK4sycl3_V16detail11image_plain12has_propertyINS0_8property7context4cuda19use_primary_contextEEEbv
+_ZNK4sycl3_V16detail11image_plain12has_propertyINS0_8property7no_initEEEbv
+_ZNK4sycl3_V16detail11image_plain12has_propertyINS0_8property9reduction22initialize_to_identityEEEbv
+_ZNK4sycl3_V16detail11image_plain13getSlicePitchEv
+_ZNK4sycl3_V16detail11image_plain14getChannelTypeEv
+_ZNK4sycl3_V16detail11image_plain14getElementSizeEv
+_ZNK4sycl3_V16detail11image_plain15getChannelOrderEv
+_ZNK4sycl3_V16detail11image_plain22get_allocator_internalEv
+_ZNK4sycl3_V16detail11image_plain8get_sizeEv
+_ZNK4sycl3_V16detail11image_plain9get_countEv
+_ZNK4sycl3_V16detail11image_plain9get_pitchEv
+_ZNK4sycl3_V16detail11image_plain9get_rangeEv
 _ZNK4sycl3_V16detail11stream_impl22get_max_statement_sizeEv
 _ZNK4sycl3_V16detail11stream_impl8get_sizeEv
+_ZNK4sycl3_V16detail12buffer_plain12get_propertyINS0_3ext6oneapi4cuda8property7context19use_primary_contextEEET_v
+_ZNK4sycl3_V16detail12buffer_plain12get_propertyINS0_3ext6oneapi8property6buffer22use_pinned_host_memoryEEET_v
+_ZNK4sycl3_V16detail12buffer_plain12get_propertyINS0_8property5image12use_host_ptrEEET_v
+_ZNK4sycl3_V16detail12buffer_plain12get_propertyINS0_8property5image13context_boundEEET_v
+_ZNK4sycl3_V16detail12buffer_plain12get_propertyINS0_8property5image9use_mutexEEET_v
+_ZNK4sycl3_V16detail12buffer_plain12get_propertyINS0_8property5queue8in_orderEEET_v
+_ZNK4sycl3_V16detail12buffer_plain12get_propertyINS0_8property6buffer12use_host_ptrEEET_v
+_ZNK4sycl3_V16detail12buffer_plain12get_propertyINS0_8property6buffer13context_boundEEET_v
+_ZNK4sycl3_V16detail12buffer_plain12get_propertyINS0_8property6buffer9use_mutexEEET_v
+_ZNK4sycl3_V16detail12buffer_plain12get_propertyINS0_8property6noinitEEET_v
+_ZNK4sycl3_V16detail12buffer_plain12get_propertyINS0_8property7context4cuda19use_primary_contextEEET_v
+_ZNK4sycl3_V16detail12buffer_plain12get_propertyINS0_8property7no_initEEET_v
+_ZNK4sycl3_V16detail12buffer_plain12get_propertyINS0_8property9reduction22initialize_to_identityEEET_v
+_ZNK4sycl3_V16detail12buffer_plain12has_propertyINS0_3ext6oneapi4cuda8property7context19use_primary_contextEEEbv
+_ZNK4sycl3_V16detail12buffer_plain12has_propertyINS0_3ext6oneapi8property6buffer22use_pinned_host_memoryEEEbv
+_ZNK4sycl3_V16detail12buffer_plain12has_propertyINS0_8property5image12use_host_ptrEEEbv
+_ZNK4sycl3_V16detail12buffer_plain12has_propertyINS0_8property5image13context_boundEEEbv
+_ZNK4sycl3_V16detail12buffer_plain12has_propertyINS0_8property5image9use_mutexEEEbv
+_ZNK4sycl3_V16detail12buffer_plain12has_propertyINS0_8property5queue8in_orderEEEbv
+_ZNK4sycl3_V16detail12buffer_plain12has_propertyINS0_8property6buffer12use_host_ptrEEEbv
+_ZNK4sycl3_V16detail12buffer_plain12has_propertyINS0_8property6buffer13context_boundEEEbv
+_ZNK4sycl3_V16detail12buffer_plain12has_propertyINS0_8property6buffer9use_mutexEEEbv
+_ZNK4sycl3_V16detail12buffer_plain12has_propertyINS0_8property6noinitEEEbv
+_ZNK4sycl3_V16detail12buffer_plain12has_propertyINS0_8property7context4cuda19use_primary_contextEEEbv
+_ZNK4sycl3_V16detail12buffer_plain12has_propertyINS0_8property7no_initEEEbv
+_ZNK4sycl3_V16detail12buffer_plain12has_propertyINS0_8property9reduction22initialize_to_identityEEEbv
+_ZNK4sycl3_V16detail12buffer_plain15getNativeVectorENS0_7backendE
+_ZNK4sycl3_V16detail12buffer_plain22get_allocator_internalEv
+_ZNK4sycl3_V16detail12buffer_plain7getSizeEv
 _ZNK4sycl3_V16detail12sampler_impl18get_filtering_modeEv
 _ZNK4sycl3_V16detail12sampler_impl19get_addressing_modeEv
 _ZNK4sycl3_V16detail12sampler_impl33get_coordinate_normalization_modeEv
diff --git a/sycl/test/abi/symbol_size_alignment.cpp b/sycl/test/abi/symbol_size_alignment.cpp
index 3ef86e5acdc56..b1eae36d543e9 100644
--- a/sycl/test/abi/symbol_size_alignment.cpp
+++ b/sycl/test/abi/symbol_size_alignment.cpp
@@ -7,8 +7,6 @@
 #include <sycl/accessor.hpp>
 #include <sycl/buffer.hpp>
 #include <sycl/detail/accessor_impl.hpp>
-#include <sycl/detail/buffer_impl.hpp>
-#include <sycl/detail/image_impl.hpp>
 #include <sycl/device.hpp>
 #include <sycl/device_event.hpp>
 #include <sycl/device_selector.hpp>
@@ -59,12 +57,8 @@ int main() {
   check<gpu_selector, 8, 8>();
 #ifdef _MSC_VER
   check<handler, 568, 8>();
-  check<detail::buffer_impl, 216, 8>();
-  check<detail::image_impl<1>, 272, 8>();
 #else
   check<handler, 576, 8>();
-  check<detail::buffer_impl, 184, 8>();
-  check<detail::image_impl<1>, 240, 8>();
 #endif
   check<image<1>, 16, 8>();
   check<kernel, 16, 8>();
diff --git a/sycl/test/abi/vtable.cpp b/sycl/test/abi/vtable.cpp
index dabb4c14e1673..008a510496db7 100644
--- a/sycl/test/abi/vtable.cpp
+++ b/sycl/test/abi/vtable.cpp
@@ -23,22 +23,6 @@ void foo(sycl::detail::HostKernelBase &HKB) {
 // CHECK-NEXT:   4 | sycl::detail::HostKernelBase::~HostKernelBase() [complete]
 // CHECK-NEXT:   5 | sycl::detail::HostKernelBase::~HostKernelBase() [deleting]
 
-void foo(sycl::detail::SYCLMemObjI &MemObj) { (void)MemObj.getType(); }
-
-// CHECK:    Vtable for 'sycl::detail::SYCLMemObjI' (11 entries).
-// CHECK-NEXT:   0 | offset_to_top (0)
-// CHECK-NEXT:   1 | sycl::detail::SYCLMemObjI RTTI
-// CHECK-NEXT:       -- (sycl::detail::SYCLMemObjI, 0) vtable address --
-// CHECK-NEXT:   2 | sycl::detail::SYCLMemObjI::~SYCLMemObjI() [complete]
-// CHECK-NEXT:   3 | sycl::detail::SYCLMemObjI::~SYCLMemObjI() [deleting]
-// CHECK-NEXT:   4 | MemObjType sycl::detail::SYCLMemObjI::getType() const [pure]
-// CHECK-NEXT:   5 | void *sycl::detail::SYCLMemObjI::allocateMem(ContextImplPtr, bool, void *, RT::PiEvent &) [pure]
-// CHECK-NEXT:   6 | void *sycl::detail::SYCLMemObjI::allocateHostMem() [pure]
-// CHECK-NEXT:   7 | void sycl::detail::SYCLMemObjI::releaseMem(ContextImplPtr, void *) [pure]
-// CHECK-NEXT:   8 | void sycl::detail::SYCLMemObjI::releaseHostMem(void *) [pure]
-// CHECK-NEXT:   9 | size_t sycl::detail::SYCLMemObjI::getSize() const [pure]
-// CHECK-NEXT:  10 | ContextImplPtr sycl::detail::SYCLMemObjI::getInteropContext() const [pure]
-
 void foo(sycl::detail::pi::DeviceBinaryImage &Img) { Img.print(); }
 // CHECK:    Vtable for 'sycl::detail::pi::DeviceBinaryImage' (6 entries).
 // CHECK-NEXT:   0 | offset_to_top (0)
diff --git a/sycl/test/gdb/printers.cpp b/sycl/test/gdb/printers.cpp
index 7784f05e648e0..7b229cbcc0bb3 100644
--- a/sycl/test/gdb/printers.cpp
+++ b/sycl/test/gdb/printers.cpp
@@ -12,8 +12,14 @@ typedef sycl::buffer<int> dummy_buffer;
 // CHECK-NOT: CXXRecordDecl {{.*}} definition
 // CHECK: FieldDecl {{.*}} referenced common_array
 
-// buffer must have impl field
+// buffer_plain must have impl field
 
-// CHECK: CXXRecordDecl {{.*}} class buffer definition
+// CHECK: CXXRecordDecl {{.*}} class buffer_plain definition
 // CHECK-NOT: CXXRecordDecl {{.*}} definition
 // CHECK: FieldDecl {{.*}} referenced impl
+
+// Check buffer inherits from buffer_plain
+
+// CHECK: CXXRecordDecl {{.*}} class buffer definition
+// CHECK-NOT: CXXRecordDecl {{.*}} definition
+// CHECK: detail::buffer_plain':'sycl::detail::buffer_plain
diff --git a/sycl/unittests/buffer/BufferLocation.cpp b/sycl/unittests/buffer/BufferLocation.cpp
index 965e2404c7432..9994e8a93bc09 100644
--- a/sycl/unittests/buffer/BufferLocation.cpp
+++ b/sycl/unittests/buffer/BufferLocation.cpp
@@ -16,6 +16,8 @@
 
 #include <gtest/gtest.h>
 
+#include <detail/buffer_impl.hpp>
+
 const uint64_t DEFAULT_VALUE = 7777;
 static uint64_t PassedLocation = DEFAULT_VALUE;
 
diff --git a/sycl/unittests/scheduler/FinishedCmdCleanup.cpp b/sycl/unittests/scheduler/FinishedCmdCleanup.cpp
index 978e842eae416..908d0ed97fa9f 100644
--- a/sycl/unittests/scheduler/FinishedCmdCleanup.cpp
+++ b/sycl/unittests/scheduler/FinishedCmdCleanup.cpp
@@ -9,6 +9,7 @@
 #include "SchedulerTest.hpp"
 #include "SchedulerTestUtils.hpp"
 
+#include <detail/buffer_impl.hpp>
 #include <detail/event_impl.hpp>
 
 using namespace sycl;
diff --git a/sycl/unittests/scheduler/MemObjCommandCleanup.cpp b/sycl/unittests/scheduler/MemObjCommandCleanup.cpp
index 39f2c845f4e86..9b5d196cde339 100644
--- a/sycl/unittests/scheduler/MemObjCommandCleanup.cpp
+++ b/sycl/unittests/scheduler/MemObjCommandCleanup.cpp
@@ -9,6 +9,8 @@
 #include "SchedulerTest.hpp"
 #include "SchedulerTestUtils.hpp"
 
+#include <detail/buffer_impl.hpp>
+
 using namespace sycl;
 
 TEST_F(SchedulerTest, MemObjCommandCleanupAllocaUsers) {
diff --git a/sycl/unittests/scheduler/NoHostUnifiedMemory.cpp b/sycl/unittests/scheduler/NoHostUnifiedMemory.cpp
index 13c97a39a74f0..765b57640823a 100644
--- a/sycl/unittests/scheduler/NoHostUnifiedMemory.cpp
+++ b/sycl/unittests/scheduler/NoHostUnifiedMemory.cpp
@@ -11,6 +11,8 @@
 
 #include <helpers/PiMock.hpp>
 
+#include <detail/buffer_impl.hpp>
+
 #include <iostream>
 #include <memory>
 
diff --git a/sycl/unittests/scheduler/PostEnqueueCleanup.cpp b/sycl/unittests/scheduler/PostEnqueueCleanup.cpp
index 8386af6394628..101f4506f47a7 100644
--- a/sycl/unittests/scheduler/PostEnqueueCleanup.cpp
+++ b/sycl/unittests/scheduler/PostEnqueueCleanup.cpp
@@ -13,6 +13,8 @@
 #include <helpers/PiMock.hpp>
 #include <helpers/ScopedEnvVar.hpp>
 
+#include <detail/buffer_impl.hpp>
+
 #include <algorithm>
 #include <cstddef>
 #include <memory>