diff --git a/sycl/include/CL/sycl/access/access.hpp b/sycl/include/CL/sycl/access/access.hpp
index 240032c47ead8..10101d02435f5 100644
--- a/sycl/include/CL/sycl/access/access.hpp
+++ b/sycl/include/CL/sycl/access/access.hpp
@@ -7,6 +7,7 @@
 //===----------------------------------------------------------------------===//
 #pragma once
 
+#include <CL/sycl/detail/common.hpp>
 #include <CL/sycl/detail/defines.hpp>
 
 __SYCL_INLINE_NAMESPACE(cl) {
@@ -47,7 +48,43 @@ enum class address_space : int {
   local_space
 };
 
-}  // namespace access
+} // namespace access
+
+using access::target;
+using access_mode = access::mode;
+
+template <access_mode mode> struct mode_tag_t {
+  explicit mode_tag_t() = default;
+};
+
+template <access_mode mode, target trgt> struct mode_target_tag_t {
+  explicit mode_target_tag_t() = default;
+};
+
+#if __cplusplus > 201402L
+
+inline constexpr mode_tag_t<access_mode::read> read_only{};
+inline constexpr mode_tag_t<access_mode::read_write> read_write{};
+inline constexpr mode_tag_t<access_mode::write> write_only{};
+inline constexpr mode_target_tag_t<access_mode::read, target::constant_buffer>
+    read_constant{};
+
+#else
+
+namespace {
+
+constexpr const auto &read_only =
+    sycl::detail::InlineVariableHelper<mode_tag_t<access_mode::read>>::value;
+constexpr const auto &read_write = sycl::detail::InlineVariableHelper<
+    mode_tag_t<access_mode::read_write>>::value;
+constexpr const auto &write_only =
+    sycl::detail::InlineVariableHelper<mode_tag_t<access_mode::write>>::value;
+constexpr const auto &read_constant = sycl::detail::InlineVariableHelper<
+    mode_target_tag_t<access_mode::read, target::constant_buffer>>::value;
+
+} // namespace
+
+#endif
 
 namespace detail {
 
diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp
index 52fc535098c37..ad8e1cf4f2571 100644
--- a/sycl/include/CL/sycl/accessor.hpp
+++ b/sycl/include/CL/sycl/accessor.hpp
@@ -134,6 +134,7 @@
 ///    "image_accessor" -> a2;
 ///    "image_accessor" -> a4;
 ///    "image_accessor" -> a5;
+///    a1 -> "host_accessor";
 /// }
 /// \enddot
 ///
@@ -156,6 +157,13 @@
 //      |   |   |       | global_buffer   |   +-------------+
 //      |   |   |       | constant_buffer |
 //      |   |   |       +-----------------+
+//      |   |   |                 |
+//      |   |   |                 v
+//      |   |   |       +-----------------+
+//      |   |   |       |                 |
+//      |   |   |       |  host_accessor  |
+//      |   |   |       |                 |
+//      |   |   |       +-----------------+
 //      |   |   |
 //      |   |   +------------------------------------+
 //      |   |                                        |
@@ -190,7 +198,8 @@
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
 
-template <typename DataT, int Dimensions, access::mode AccessMode,
+template <typename DataT, int Dimensions = 1,
+          access::mode AccessMode = access::mode::read_write,
           access::target AccessTarget = access::target::global_buffer,
           access::placeholder IsPlaceholder = access::placeholder::false_t>
 class accessor;
@@ -709,7 +718,7 @@ class accessor :
 #endif
     public detail::accessor_common<DataT, Dimensions, AccessMode, AccessTarget,
                                    IsPlaceholder> {
-
+protected:
   static_assert((AccessTarget == access::target::global_buffer ||
                  AccessTarget == access::target::constant_buffer ||
                  AccessTarget == access::target::host_buffer),
@@ -756,6 +765,34 @@ class accessor :
     return Result;
   }
 
+  template <typename T, int Dims> static constexpr bool IsSameAsBuffer() {
+    return std::is_same<T, DataT>::value && (Dims > 0) && (Dims == Dimensions);
+  }
+
+  static access::mode getAdjustedMode(const property_list &PropertyList) {
+    access::mode AdjustedMode = AccessMode;
+
+    if (PropertyList.has_property<property::noinit>()) {
+      if (AdjustedMode == access::mode::write) {
+        AdjustedMode = access::mode::discard_write;
+      } else if (AdjustedMode == access::mode::read_write) {
+        AdjustedMode = access::mode::discard_read_write;
+      }
+    }
+
+    return AdjustedMode;
+  }
+
+#if __cplusplus > 201402L
+
+  template <typename TagT> static constexpr bool IsValidTag() {
+    return std::is_same<TagT, mode_tag_t<AccessMode>>::value ||
+           std::is_same<TagT,
+                        mode_target_tag_t<AccessMode, AccessTarget>>::value;
+  }
+
+#endif
+
 #ifdef __SYCL_DEVICE_ONLY__
 
   id<AdjustedDim> &getOffset() { return impl.Offset; }
@@ -811,17 +848,45 @@ class accessor :
   using reference = DataT &;
   using const_reference = const DataT &;
 
-  template <int Dims = Dimensions, typename AllocatorT,
+  // The list of accessor constructors with their arguments
+  // -------+---------+-------+----+-----+--------------
+  // Dimensions = 0
+  // -------+---------+-------+----+-----+--------------
+  // buffer |         |       |    |     | property_list
+  // buffer | handler |       |    |     | property_list
+  // -------+---------+-------+----+-----+--------------
+  // Dimensions >= 1
+  // -------+---------+-------+----+-----+--------------
+  // buffer |         |       |    |     | property_list
+  // buffer |         |       |    | tag | property_list
+  // buffer | handler |       |    |     | property_list
+  // buffer | handler |       |    | tag | property_list
+  // buffer |         | range |    |     | property_list
+  // buffer |         | range |    | tag | property_list
+  // buffer | handler | range |    |     | property_list
+  // buffer | handler | range |    | tag | property_list
+  // buffer |         | range | id |     | property_list
+  // buffer |         | range | id | tag | property_list
+  // buffer | handler | range | id |     | property_list
+  // buffer | handler | range | id | tag | property_list
+  // -------+---------+-------+----+-----+--------------
+
+public:
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
             typename detail::enable_if_t<
-                Dims == 0 && ((!IsPlaceH && IsHostBuf) ||
-                (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>* = nullptr>
-  accessor(buffer<DataT, 1, AllocatorT> &BufferRef)
+                std::is_same<T, DataT>::value && Dims == 0 &&
+                ((!IsPlaceH && IsHostBuf) ||
+                 (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))> * = nullptr>
+  accessor(buffer<T, 1, AllocatorT> &BufferRef,
+           const property_list &PropertyList = {})
 #ifdef __SYCL_DEVICE_ONLY__
       : impl(id<AdjustedDim>(), range<1>{1}, BufferRef.get_range()) {
+    (void)PropertyList;
 #else
       : AccessorBaseHost(
             /*Offset=*/{0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}),
-            detail::convertToArrayOfN<3, 1>(BufferRef.get_range()), AccessMode,
+            detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
+            getAdjustedMode(PropertyList),
             detail::getSyclObjImpl(BufferRef).get(), AdjustedDim, sizeof(DataT),
             BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) {
     if (!IsPlaceH)
@@ -829,41 +894,45 @@ class accessor :
 #endif
   }
 
-  template <int Dims = Dimensions, typename AllocatorT,
-	   typename = typename detail::enable_if_t<
-		   (Dims == 0) && 
-                    (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>
-                                      >
-  accessor(buffer<DataT,1,AllocatorT> &BufferRef,
-		  handler &CommandGroupHandler)
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            typename = typename detail::enable_if_t<
+                std::is_same<T, DataT>::value && (Dims == 0) &&
+                (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
+  accessor(buffer<T, 1, AllocatorT> &BufferRef, handler &CommandGroupHandler,
+           const property_list &PropertyList = {})
 #ifdef __SYCL_DEVICE_ONLY__
       : impl(id<AdjustedDim>(), range<1>{1}, BufferRef.get_range()) {
     (void)CommandGroupHandler;
+    (void)PropertyList;
   }
 #else
       : AccessorBaseHost(
             /*Offset=*/{0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}),
-            detail::convertToArrayOfN<3, 1>(BufferRef.get_range()), AccessMode,
+            detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
+            getAdjustedMode(PropertyList),
             detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
             BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) {
     detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
   }
 #endif
 
-  template <int Dims = Dimensions, typename AllocatorT,
-            typename = detail::enable_if_t<(Dims > 0) && (Dims == Dimensions) &&
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>() &&
                                            ((!IsPlaceH && IsHostBuf) ||
                                             (IsPlaceH &&
                                              (IsGlobalBuf || IsConstantBuf)))>>
-  accessor(buffer<DataT, Dims, AllocatorT> &BufferRef)
+  accessor(buffer<T, Dims, AllocatorT> &BufferRef,
+           const property_list &PropertyList = {})
 #ifdef __SYCL_DEVICE_ONLY__
       : impl(id<Dimensions>(), BufferRef.get_range(), BufferRef.get_range()) {
+    (void)PropertyList;
   }
 #else
       : AccessorBaseHost(
             /*Offset=*/{0, 0, 0},
             detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
-            detail::convertToArrayOfN<3, 1>(BufferRef.get_range()), AccessMode,
+            detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
+            getAdjustedMode(PropertyList),
             detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
             BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) {
     if (!IsPlaceH)
@@ -871,71 +940,182 @@ class accessor :
   }
 #endif
 
-  template <int Dims = Dimensions, typename AllocatorT,
+#if __cplusplus > 201402L
+
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            typename TagT,
+            typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>() &&
+                                           IsValidTag<TagT>() && IsPlaceH &&
+                                           (IsGlobalBuf || IsConstantBuf)>>
+  accessor(buffer<T, Dims, AllocatorT> &BufferRef, TagT,
+           const property_list &PropertyList = {})
+      : accessor(BufferRef, PropertyList) {}
+
+#endif
+
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
             typename = detail::enable_if_t<
-                (Dims > 0) && (Dims == Dimensions) &&
+                IsSameAsBuffer<T, Dims>() &&
                 (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
-  accessor(buffer<DataT, Dims, AllocatorT> &BufferRef,
-           handler &CommandGroupHandler)
+  accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
+           const property_list &PropertyList = {})
 #ifdef __SYCL_DEVICE_ONLY__
       : impl(id<AdjustedDim>(), BufferRef.get_range(), BufferRef.get_range()) {
     (void)CommandGroupHandler;
+    (void)PropertyList;
   }
 #else
       : AccessorBaseHost(
             /*Offset=*/{0, 0, 0},
             detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
-            detail::convertToArrayOfN<3, 1>(BufferRef.get_range()), AccessMode,
+            detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
+            getAdjustedMode(PropertyList),
             detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
             BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) {
     detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
   }
 #endif
 
-  template <int Dims = Dimensions, typename AllocatorT,
-            typename = detail::enable_if_t<(Dims > 0) && (Dims == Dimensions) &&
+#if __cplusplus > 201402L
+
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            typename TagT,
+            typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>() &&
+                                           IsValidTag<TagT>() && !IsPlaceH &&
+                                           (IsGlobalBuf || IsConstantBuf)>>
+  accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
+           TagT, const property_list &PropertyList = {})
+      : accessor(BufferRef, CommandGroupHandler, PropertyList) {}
+
+#endif
+
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>() &&
+                                           ((!IsPlaceH && IsHostBuf) ||
+                                            (IsPlaceH &&
+                                             (IsGlobalBuf || IsConstantBuf)))>>
+  accessor(buffer<T, Dims, AllocatorT> &BufferRef,
+           range<Dimensions> AccessRange,
+           const property_list &PropertyList = {})
+      : accessor(BufferRef, AccessRange, {}, PropertyList) {}
+
+#if __cplusplus > 201402L
+
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            typename TagT,
+            typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>() &&
+                                           IsValidTag<TagT>() && IsPlaceH &&
+                                           (IsGlobalBuf || IsConstantBuf)>>
+  accessor(buffer<T, Dims, AllocatorT> &BufferRef,
+           range<Dimensions> AccessRange, TagT,
+           const property_list &PropertyList = {})
+      : accessor(BufferRef, AccessRange, {}, PropertyList) {}
+
+#endif
+
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>() &&
+                                           (!IsPlaceH &&
+                                            (IsGlobalBuf || IsConstantBuf))>>
+  accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
+           range<Dimensions> AccessRange,
+           const property_list &PropertyList = {})
+      : accessor(BufferRef, CommandGroupHandler, AccessRange, {},
+                 PropertyList) {}
+
+#if __cplusplus > 201402L
+
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            typename TagT,
+            typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>() &&
+                                           IsValidTag<TagT>() && !IsPlaceH &&
+                                           (IsGlobalBuf || IsConstantBuf)>>
+  accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
+           range<Dimensions> AccessRange, TagT,
+           const property_list &PropertyList = {})
+      : accessor(BufferRef, CommandGroupHandler, AccessRange, {},
+                 PropertyList) {}
+
+#endif
+
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>() &&
                                            ((!IsPlaceH && IsHostBuf) ||
                                             (IsPlaceH &&
                                              (IsGlobalBuf || IsConstantBuf)))>>
-  accessor(buffer<DataT, Dims, AllocatorT> &BufferRef,
-           range<Dimensions> AccessRange, id<Dimensions> AccessOffset = {})
+  accessor(buffer<T, Dims, AllocatorT> &BufferRef,
+           range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
+           const property_list &PropertyList = {})
 #ifdef __SYCL_DEVICE_ONLY__
       : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
+    (void)PropertyList;
   }
 #else
       : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
                          detail::convertToArrayOfN<3, 1>(AccessRange),
                          detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
-                         AccessMode, detail::getSyclObjImpl(BufferRef).get(),
-                         Dimensions, sizeof(DataT), BufferRef.OffsetInBytes,
+                         getAdjustedMode(PropertyList),
+                         detail::getSyclObjImpl(BufferRef).get(), Dimensions,
+                         sizeof(DataT), BufferRef.OffsetInBytes,
                          BufferRef.IsSubBuffer) {
     if (!IsPlaceH)
       addHostAccessorAndWait(AccessorBaseHost::impl.get());
   }
 #endif
 
-  template <int Dims = Dimensions, typename AllocatorT,
-            typename = detail::enable_if_t<(Dims > 0) && (Dims == Dimensions) &&
+#if __cplusplus > 201402L
+
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            typename TagT,
+            typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>() &&
+                                           IsValidTag<TagT>() && IsPlaceH &&
+                                           (IsGlobalBuf || IsConstantBuf)>>
+  accessor(buffer<T, Dims, AllocatorT> &BufferRef,
+           range<Dimensions> AccessRange, id<Dimensions> AccessOffset, TagT,
+           const property_list &PropertyList = {})
+      : accessor(BufferRef, AccessRange, AccessOffset, PropertyList) {}
+
+#endif
+
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>() &&
                                            (!IsPlaceH &&
                                             (IsGlobalBuf || IsConstantBuf))>>
-  accessor(buffer<DataT, Dims, AllocatorT> &BufferRef,
-           handler &CommandGroupHandler, range<Dimensions> AccessRange,
-           id<Dimensions> AccessOffset = {})
+  accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
+           range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
+           const property_list &PropertyList = {})
 #ifdef __SYCL_DEVICE_ONLY__
       : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
     (void)CommandGroupHandler;
+    (void)PropertyList;
   }
 #else
       : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
                          detail::convertToArrayOfN<3, 1>(AccessRange),
                          detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
-                         AccessMode, detail::getSyclObjImpl(BufferRef).get(),
-                         Dimensions, sizeof(DataT), BufferRef.OffsetInBytes,
+                         getAdjustedMode(PropertyList),
+                         detail::getSyclObjImpl(BufferRef).get(), Dimensions,
+                         sizeof(DataT), BufferRef.OffsetInBytes,
                          BufferRef.IsSubBuffer) {
     detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
   }
 #endif
 
+#if __cplusplus > 201402L
+
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            typename TagT,
+            typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>() &&
+                                           IsValidTag<TagT>() && !IsPlaceH &&
+                                           (IsGlobalBuf || IsConstantBuf)>>
+  accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
+           range<Dimensions> AccessRange, id<Dimensions> AccessOffset, TagT,
+           const property_list &PropertyList = {})
+      : accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
+                 PropertyList) {}
+
+#endif
+
   constexpr bool is_placeholder() const { return IsPlaceH; }
 
   size_t get_size() const { return getAccessRange().size() * sizeof(DataT); }
@@ -1041,6 +1221,48 @@ class accessor :
   bool operator!=(const accessor &Rhs) const { return !(*this == Rhs); }
 };
 
+#if __cplusplus > 201402L
+
+template <typename DataT, int Dimensions, typename AllocatorT, typename... Ts>
+accessor(buffer<DataT, Dimensions, AllocatorT>, Ts...)
+    ->accessor<DataT, Dimensions, access::mode::read_write,
+               target::global_buffer, access::placeholder::true_t>;
+
+template <typename DataT, int Dimensions, typename AllocatorT, typename... Ts>
+accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Ts...)
+    ->accessor<DataT, Dimensions, access::mode::read_write,
+               target::global_buffer, access::placeholder::false_t>;
+
+template <typename DataT, int Dimensions, typename AllocatorT,
+          access_mode AccessMode, typename... Ts>
+accessor(buffer<DataT, Dimensions, AllocatorT>, Ts..., mode_tag_t<AccessMode>,
+         property_list = {})
+    ->accessor<DataT, Dimensions, AccessMode, target::global_buffer,
+               access::placeholder::true_t>;
+
+template <typename DataT, int Dimensions, typename AllocatorT,
+          access_mode AccessMode, typename... Ts>
+accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Ts...,
+         mode_tag_t<AccessMode>, property_list = {})
+    ->accessor<DataT, Dimensions, AccessMode, target::global_buffer,
+               access::placeholder::false_t>;
+
+template <typename DataT, int Dimensions, typename AllocatorT,
+          access_mode AccessMode, target AccessTarget, typename... Ts>
+accessor(buffer<DataT, Dimensions, AllocatorT>, Ts...,
+         mode_target_tag_t<AccessMode, AccessTarget>, property_list = {})
+    ->accessor<DataT, Dimensions, AccessMode, AccessTarget,
+               access::placeholder::true_t>;
+
+template <typename DataT, int Dimensions, typename AllocatorT,
+          access_mode AccessMode, target AccessTarget, typename... Ts>
+accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Ts...,
+         mode_target_tag_t<AccessMode, AccessTarget>, property_list = {})
+    ->accessor<DataT, Dimensions, AccessMode, AccessTarget,
+               access::placeholder::false_t>;
+
+#endif
+
 /// Local accessor
 ///
 /// \ingroup sycl_api_acc
@@ -1315,6 +1537,139 @@ class accessor<DataT, Dimensions, AccessMode, access::target::image_array,
   }
 };
 
+template <typename DataT, int Dimensions = 1,
+          access_mode AccessMode = access_mode::read_write>
+class host_accessor
+    : public accessor<DataT, Dimensions, AccessMode, target::host_buffer,
+                      access::placeholder::false_t> {
+protected:
+  using AccessorT = accessor<DataT, Dimensions, AccessMode, target::host_buffer,
+                             access::placeholder::false_t>;
+
+  constexpr static int AdjustedDim = Dimensions == 0 ? 1 : Dimensions;
+
+  template <typename T, int Dims> static constexpr bool IsSameAsBuffer() {
+    return std::is_same<T, DataT>::value && (Dims > 0) && (Dims == Dimensions);
+  }
+
+#if __cplusplus > 201402L
+
+  template <typename TagT> static constexpr bool IsValidTag() {
+    return std::is_same<TagT, mode_tag_t<AccessMode>>::value;
+  }
+
+#endif
+
+  void
+  __init(typename accessor<DataT, Dimensions, AccessMode, target::host_buffer,
+                           access::placeholder::false_t>::ConcreteASPtrType Ptr,
+         range<AdjustedDim> AccessRange, range<AdjustedDim> MemRange,
+         id<AdjustedDim> Offset) {
+    AccessorT::__init(Ptr, AccessRange, MemRange, Offset);
+  }
+
+public:
+  host_accessor() : AccessorT() {}
+
+  // The list of host_accessor constructors with their arguments
+  // -------+---------+-------+----+----------+--------------
+  // Dimensions = 0
+  // -------+---------+-------+----+----------+--------------
+  // buffer |         |       |    |          | property_list
+  // buffer | handler |       |    |          | property_list
+  // -------+---------+-------+----+----------+--------------
+  // Dimensions >= 1
+  // -------+---------+-------+----+----------+--------------
+  // buffer |         |       |    |          | property_list
+  // buffer |         |       |    | mode_tag | property_list
+  // buffer | handler |       |    |          | property_list
+  // buffer | handler |       |    | mode_tag | property_list
+  // buffer |         | range |    |          | property_list
+  // buffer |         | range |    | mode_tag | property_list
+  // buffer | handler | range |    |          | property_list
+  // buffer | handler | range |    | mode_tag | property_list
+  // buffer |         | range | id |          | property_list
+  // buffer |         | range | id | mode_tag | property_list
+  // buffer | handler | range | id |          | property_list
+  // buffer | handler | range | id | mode_tag | property_list
+  // -------+---------+-------+----+----------+--------------
+  // host_accessor with handler argument will be added later
+  // to facilitate non-blocking accessor use case
+
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            typename = typename detail::enable_if_t<
+                std::is_same<T, DataT>::value && Dims == 0>>
+  host_accessor(buffer<T, 1, AllocatorT> &BufferRef,
+                const property_list &PropertyList = {})
+      : AccessorT(BufferRef, PropertyList) {}
+
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
+  host_accessor(buffer<T, Dims, AllocatorT> &BufferRef,
+                const property_list &PropertyList = {})
+      : AccessorT(BufferRef, PropertyList) {}
+
+#if __cplusplus > 201402L
+
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
+  host_accessor(buffer<DataT, Dimensions, AllocatorT> &BufferRef,
+                mode_tag_t<AccessMode>, const property_list &PropertyList = {})
+      : host_accessor(BufferRef, PropertyList) {}
+
+#endif
+
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
+  host_accessor(buffer<DataT, Dimensions, AllocatorT> &BufferRef,
+                range<Dimensions> AccessRange,
+                const property_list &PropertyList = {})
+      : AccessorT(BufferRef, AccessRange, {}, PropertyList) {}
+
+#if __cplusplus > 201402L
+
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
+  host_accessor(buffer<DataT, Dimensions, AllocatorT> &BufferRef,
+                range<Dimensions> AccessRange, mode_tag_t<AccessMode>,
+                const property_list &PropertyList = {})
+      : host_accessor(BufferRef, AccessRange, {}, PropertyList) {}
+
+#endif
+
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
+  host_accessor(buffer<DataT, Dimensions, AllocatorT> &BufferRef,
+                range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
+                const property_list &PropertyList = {})
+      : AccessorT(BufferRef, AccessRange, AccessOffset, PropertyList) {}
+
+#if __cplusplus > 201402L
+
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
+  host_accessor(buffer<DataT, Dimensions, AllocatorT> &BufferRef,
+                range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
+                mode_tag_t<AccessMode>, const property_list &PropertyList = {})
+      : host_accessor(BufferRef, AccessRange, AccessOffset, PropertyList) {}
+
+#endif
+};
+
+#if __cplusplus > 201402L
+
+template <typename DataT, int Dimensions, typename AllocatorT, typename... Ts>
+host_accessor(buffer<DataT, Dimensions, AllocatorT>, Ts...)
+    ->host_accessor<DataT, Dimensions, access::mode::read_write>;
+
+template <typename DataT, int Dimensions, typename AllocatorT,
+          access_mode AccessMode, typename... Ts>
+host_accessor(buffer<DataT, Dimensions, AllocatorT>, Ts...,
+              mode_tag_t<AccessMode>, property_list = {})
+    ->host_accessor<DataT, Dimensions, AccessMode>;
+
+#endif
+
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)
 
diff --git a/sycl/include/CL/sycl/detail/common.hpp b/sycl/include/CL/sycl/detail/common.hpp
index e5d70adb7e829..11bb8f395a7c8 100644
--- a/sycl/include/CL/sycl/detail/common.hpp
+++ b/sycl/include/CL/sycl/detail/common.hpp
@@ -307,6 +307,11 @@ using KernelSetId = size_t;
 constexpr KernelSetId SpvFileKSId = 0;
 constexpr KernelSetId LastKSId = SpvFileKSId;
 
+template <typename T> struct InlineVariableHelper {
+  static constexpr T value{};
+};
+
+template <typename T> constexpr T InlineVariableHelper<T>::value;
 } // namespace detail
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)
diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp
index 90c60b89b2616..8a6b1f97cdab3 100644
--- a/sycl/include/CL/sycl/handler.hpp
+++ b/sycl/include/CL/sycl/handler.hpp
@@ -771,9 +771,8 @@ class __SYCL_EXPORT handler {
   /// \param Acc is a SYCL accessor describing required memory region.
   template <typename DataT, int Dims, access::mode AccMode,
             access::target AccTarget>
-  void
-  require(accessor<DataT, Dims, AccMode, AccTarget, access::placeholder::true_t>
-              Acc) {
+  void require(accessor<DataT, Dims, AccMode, AccTarget,
+                        access::placeholder::true_t> &Acc) {
 #ifndef __SYCL_DEVICE_ONLY__
     associateWithHandler(&Acc, AccTarget);
 #else
diff --git a/sycl/include/CL/sycl/property_list.hpp b/sycl/include/CL/sycl/property_list.hpp
index fbc3ef7bf5402..90aea2d60164a 100644
--- a/sycl/include/CL/sycl/property_list.hpp
+++ b/sycl/include/CL/sycl/property_list.hpp
@@ -9,6 +9,7 @@
 #pragma once
 
 #include <CL/sycl/context.hpp>
+#include <CL/sycl/detail/common.hpp>
 #include <tuple>
 #include <type_traits>
 #include <utility>
@@ -42,6 +43,8 @@ class enable_profiling;
 class in_order;
 } // namespace queue
 
+class noinit;
+
 namespace detail {
 
 // List of all properties' IDs.
@@ -60,6 +63,9 @@ enum PropKind {
   QueueEnableProfiling,
   InOrder,
 
+  // Accessor
+  NoInit,
+
   PropKindSize
 };
 
@@ -147,6 +153,9 @@ RegisterProp(PropKind::BufferContextBound, buffer::context_bound);
 RegisterProp(PropKind::QueueEnableProfiling, queue::enable_profiling);
 RegisterProp(PropKind::InOrder, queue::in_order);
 
+// Accessor
+RegisterProp(PropKind::NoInit, noinit);
+
 // Sentinel, needed for automatic build of tuple in property_list.
 RegisterProp(PropKind::PropKindSize, PropBase);
 
@@ -212,8 +221,25 @@ class enable_profiling
 class in_order : public detail::Prop<detail::PropKind::InOrder> {};
 } // namespace queue
 
+class noinit : public detail::Prop<detail::PropKind::NoInit> {};
+
 } // namespace property
 
+#if __cplusplus > 201402L
+
+inline constexpr property::noinit noinit;
+
+#else
+
+namespace {
+
+constexpr const auto &noinit =
+    sycl::detail::InlineVariableHelper<property::noinit>::value;
+
+}
+
+#endif
+
 class property_list {
 
   // The structs validate that all objects passed are base of PropBase class.
diff --git a/sycl/test/basic_tests/accessor/accessor.cpp b/sycl/test/basic_tests/accessor/accessor.cpp
index 1cbb4d983828a..b5ac254e91966 100644
--- a/sycl/test/basic_tests/accessor/accessor.cpp
+++ b/sycl/test/basic_tests/accessor/accessor.cpp
@@ -1,8 +1,13 @@
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple  %s -o %t.out
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -Dsimplification_test -std=c++17 %s -o %t.s.out
 // RUN: env SYCL_DEVICE_TYPE=HOST %t.out
+// RUN: env SYCL_DEVICE_TYPE=HOST %t.s.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.s.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.s.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.s.out
 
 //==----------------accessor.cpp - SYCL accessor basic test ----------------==//
 //
@@ -67,17 +72,24 @@ int main() {
                                  {cl::sycl::property::buffer::use_host_ptr()});
 
     sycl::id<1> id1(1);
+#ifndef simplification_test
     auto acc_src = buf_src.get_access<sycl::access::mode::read>();
     auto acc_dst = buf_dst.get_access<sycl::access::mode::read_write>();
-
+#else
+    sycl::host_accessor acc_src(buf_src, sycl::read_only);
+    sycl::host_accessor acc_dst(buf_dst);
+#endif
     assert(!acc_src.is_placeholder());
     assert(acc_src.get_size() == sizeof(src));
     assert(acc_src.get_count() == 2);
     assert(acc_src.get_range() == sycl::range<1>(2));
 
-    // Make sure that operator[] is defined for both size_t and id<1>.
+    // operator[] overload for size_t was intentionally removed
+    // to remove ambiguity, when passing item to operator[].
     // Implicit conversion from IdxSzT to size_t guarantees that no
-    // implicit conversion from size_t to id<1> will happen.
+    // implicit conversion from size_t to id<1> will happen,
+    // thus `acc_src[IdxSzT(0)]` will no longer compile.
+    // Replaced with acc_src[0].
     assert(acc_src[0] + acc_src[IdxID1(1)] == 10);
 
     acc_dst[0] = acc_src[0] + acc_src[IdxID1(0)];
@@ -92,7 +104,11 @@ int main() {
       data[i] = i;
     {
       sycl::buffer<int, 3> buf(data, sycl::range<3>(2, 3, 4));
+#ifndef simplification_test
       auto acc = buf.get_access<sycl::access::mode::read_write>();
+#else
+      sycl::host_accessor acc(buf);
+#endif
 
       assert(!acc.is_placeholder());
       assert(acc.get_size() == sizeof(data));
@@ -117,7 +133,11 @@ int main() {
                              {cl::sycl::property::buffer::use_host_ptr()});
 
     Queue.submit([&](sycl::handler &cgh) {
+#ifndef simplification_test
       auto acc = buf.get_access<sycl::access::mode::read_write>(cgh);
+#else
+      sycl::accessor acc(buf, cgh);
+#endif
       assert(!acc.is_placeholder());
       assert(acc.get_size() == sizeof(int));
       assert(acc.get_count() == 1);
@@ -140,7 +160,11 @@ int main() {
                                  {cl::sycl::property::buffer::use_host_ptr()});
 
         Queue.submit([&](sycl::handler &cgh) {
+#ifndef simplification_test
           auto acc = buf.get_access<sycl::access::mode::read_write>(cgh);
+#else
+          sycl::accessor acc(buf, cgh);
+#endif
           cgh.parallel_for<class dim2_subscr>(Range, [=](sycl::item<2> itemID) {
             acc[itemID.get_id(0)][itemID.get_id(1)] += itemID.get_linear_id();
           });
@@ -168,7 +192,11 @@ int main() {
                                  {cl::sycl::property::buffer::use_host_ptr()});
 
         Queue.submit([&](sycl::handler &cgh) {
+#ifndef simplification_test
           auto acc = buf.get_access<sycl::access::mode::read_write>(cgh);
+#else
+          sycl::accessor acc(buf, cgh);
+#endif
           cgh.parallel_for<class dim3_subscr>(Range, [=](sycl::item<3> itemID) {
             acc[itemID.get_id(0)][itemID.get_id(1)][itemID.get_id(2)] +=
                 itemID.get_linear_id();
@@ -195,14 +223,22 @@ int main() {
       sycl::buffer<int, 1> buf(sycl::range<1>(3));
 
       Queue.submit([&](sycl::handler& cgh) {
+#ifndef simplification_test
         auto dev_acc = buf.get_access<sycl::access::mode::discard_write>(cgh);
+#else
+        sycl::accessor dev_acc(buf, cgh, sycl::noinit);
+#endif
 
         cgh.parallel_for<class test_discard_write>(
             sycl::range<1>{3},
             [=](sycl::id<1> index) { dev_acc[index] = 42; });
       });
 
+#ifndef simplification_test
       auto host_acc = buf.get_access<sycl::access::mode::read>();
+#else
+      sycl::host_accessor host_acc(buf, sycl::read_only);
+#endif
       for (int i = 0; i != 3; ++i)
         assert(host_acc[i] == 42);
 
@@ -219,15 +255,23 @@ int main() {
       sycl::buffer<int, 1> buf(sycl::range<1>(3));
 
       Queue.submit([&](sycl::handler& cgh) {
+#ifndef simplification_test
         auto dev_acc = buf.get_access<sycl::access::mode::write>(cgh);
+#else
+        sycl::accessor dev_acc(buf, cgh, sycl::write_only);
+#endif
 
         cgh.parallel_for<class test_discard_read_write>(
             sycl::range<1>{3},
             [=](sycl::id<1> index) { dev_acc[index] = 42; });
       });
 
+#ifndef simplification_test
       auto host_acc =
         buf.get_access<sycl::access::mode::discard_read_write>();
+#else
+      sycl::host_accessor host_acc(buf, sycl::noinit);
+#endif
     } catch (cl::sycl::exception e) {
       std::cout << "SYCL exception caught: " << e.what();
       return 1;
@@ -243,7 +287,11 @@ int main() {
         sycl::buffer<int, 1> buf((int *)array, sycl::range<1>(10),
                                  {cl::sycl::property::buffer::use_host_ptr()});
         queue.submit([&](sycl::handler &cgh) {
+#ifndef simplification_test
           auto acc = buf.get_access<sycl::access::mode::read_write>(cgh);
+#else
+          sycl::accessor acc(buf, cgh);
+#endif
           auto acc_wrapped = AccWrapper<decltype(acc)>{acc};
           cgh.parallel_for<class wrapped_access1>(
               sycl::range<1>(buf.get_count()), [=](sycl::item<1> it) {
@@ -273,8 +321,13 @@ int main() {
         sycl::buffer<int, 1> buf2((int *)array2, sycl::range<1>(10),
                                   {cl::sycl::property::buffer::use_host_ptr()});
         queue.submit([&](sycl::handler &cgh) {
+#ifndef simplification_test
           auto acc1 = buf1.get_access<sycl::access::mode::read_write>(cgh);
           auto acc2 = buf2.get_access<sycl::access::mode::read_write>(cgh);
+#else
+          sycl::accessor acc1(buf1, cgh);
+          sycl::accessor acc2(buf2, cgh);
+#endif
           auto acc_wrapped =
               AccsWrapper<decltype(acc1), decltype(acc2)>{10, acc1, 5, acc2};
           cgh.parallel_for<class wrapped_access2>(
@@ -304,7 +357,11 @@ int main() {
         sycl::buffer<int, 1> buf((int *)array, sycl::range<1>(10),
                                  {cl::sycl::property::buffer::use_host_ptr()});
         queue.submit([&](sycl::handler &cgh) {
+#ifndef simplification_test
           auto acc = buf.get_access<sycl::access::mode::read_write>(cgh);
+#else
+          sycl::accessor acc(buf, cgh);
+#endif
           auto acc_wrapped = AccWrapper<decltype(acc)>{acc};
           Wrapper1 wr1;
           auto wr2 = Wrapper2<decltype(acc)>{wr1, acc_wrapped};
@@ -332,15 +389,24 @@ int main() {
       sycl::buffer<int, 1> buf(array, sycl::range<1>(3));
 
       queue.submit([&](sycl::handler& cgh) {
+#ifndef simplification_test
         auto acc1 = buf.get_access<sycl::access::mode::read>(cgh);
         auto acc2 = buf.get_access<sycl::access::mode::read_write>(cgh);
+#else
+        sycl::accessor acc1(buf, cgh, sycl::read_only);
+        sycl::accessor acc2(buf, cgh);
+#endif
 
         cgh.parallel_for<class two_accessors_to_buf>(
             sycl::range<1>{3},
             [=](sycl::id<1> index) { acc2[index] = 41 + acc1[index]; });
       });
 
+#ifndef simplification_test
       auto host_acc = buf.get_access<sycl::access::mode::read>();
+#else
+      sycl::host_accessor host_acc(buf, sycl::read_only);
+#endif
       for (int i = 0; i != 3; ++i)
         assert(host_acc[i] == 42);
 
@@ -393,13 +459,17 @@ int main() {
       sycl::accessor<int, 0, sycl::access::mode::read_write,
                      sycl::access::target::global_buffer>
           acc1(buf1, cgh);
+#ifndef simplification_test
       sycl::accessor<int, 1, sycl::access::mode::read_write,
                      sycl::access::target::global_buffer>
           acc2(buf2, cgh);
       sycl::accessor<int, 1, sycl::access::mode::read_write,
                      sycl::access::target::global_buffer>
           acc3(buf3, cgh, sycl::range<1>(1));
-
+#else
+      sycl::accessor acc2(buf2, cgh);
+      sycl::accessor acc3(buf3, cgh, sycl::range<1>(1));
+#endif
       cgh.single_task<class acc_alloc_buf>([=]() {
         acc1 *= 2;
         acc2[0] *= 2;
@@ -410,15 +480,111 @@ int main() {
     sycl::accessor<int, 0, sycl::access::mode::read,
                    sycl::access::target::host_buffer>
         acc4(buf1);
+#ifndef simplification_test
     sycl::accessor<int, 1, sycl::access::mode::read,
                    sycl::access::target::host_buffer>
         acc5(buf2);
     sycl::accessor<int, 1, sycl::access::mode::read,
                    sycl::access::target::host_buffer>
         acc6(buf3, sycl::range<1>(1));
+#else
+    sycl::host_accessor acc5(buf2, sycl::read_only);
+    sycl::host_accessor acc6(buf3, sycl::range<1>(1), sycl::read_only);
+#endif
 
     assert(acc4 == 2);
     assert(acc5[0] == 4);
     assert(acc6[0] == 6);
   }
+
+  // Constant buffer accessor
+  {
+    try {
+      int data = -1;
+      int cnst = 399;
+
+      {
+        sycl::buffer<int, 1> d(&data, sycl::range<1>(1));
+        sycl::buffer<int, 1> c(&cnst, sycl::range<1>(1));
+
+        sycl::queue queue;
+        queue.submit([&](sycl::handler &cgh) {
+#ifndef simplification_test
+          sycl::accessor<int, 1, sycl::access::mode::write,
+                         sycl::access::target::global_buffer>
+              D(d, cgh);
+          sycl::accessor<int, 1, sycl::access::mode::read,
+                         sycl::access::target::constant_buffer>
+              C(c, cgh);
+#else
+          sycl::accessor D(d, cgh, sycl::write_only);
+          sycl::accessor C(c, cgh, sycl::read_constant);
+#endif
+
+          cgh.single_task<class acc_with_const>([=]() {
+            D[0] = C[0];
+          });
+        });
+
+#ifndef simplification_test
+        auto host_acc = d.get_access<sycl::access::mode::read>();
+#else
+        sycl::host_accessor host_acc(d, sycl::read_only);
+#endif
+        assert(host_acc[0] == 399);
+      }
+
+    } catch (sycl::exception e) {
+      std::cout << "SYCL exception caught: " << e.what();
+      return 1;
+    }
+  }
+
+  // Placeholder accessor
+  {
+    try {
+      int data = -1;
+      int cnst = 399;
+
+      {
+        sycl::buffer<int, 1> d(&data, sycl::range<1>(1));
+        sycl::buffer<int, 1> c(&cnst, sycl::range<1>(1));
+
+#ifndef simplification_test
+        sycl::accessor<int, 1, sycl::access::mode::write,
+                       sycl::access::target::global_buffer,
+                       sycl::access::placeholder::true_t>
+            D(d);
+        sycl::accessor<int, 1, sycl::access::mode::read,
+                       sycl::access::target::constant_buffer,
+                       sycl::access::placeholder::true_t>
+            C(c);
+#else
+        sycl::accessor D(d, sycl::write_only);
+        sycl::accessor C(c, sycl::read_constant);
+#endif
+
+        sycl::queue queue;
+        queue.submit([&](sycl::handler &cgh) {
+          cgh.require(D);
+          cgh.require(C);
+
+          cgh.single_task<class placeholder_acc>([=]() {
+            D[0] = C[0];
+          });
+        });
+
+#ifndef simplification_test
+        auto host_acc = d.get_access<sycl::access::mode::read>();
+#else
+        sycl::host_accessor host_acc(d, sycl::read_only);
+#endif
+        assert(host_acc[0] == 399);
+      }
+
+    } catch (sycl::exception e) {
+      std::cout << "SYCL exception caught: " << e.what();
+      return 1;
+    }
+  }
 }