From d0d2a8651663eebce9397ed390d73dcaf9d6c01e Mon Sep 17 00:00:00 2001
From: Alexander Batashev <alexander.batashev@intel.com>
Date: Fri, 7 Aug 2020 10:50:39 +0300
Subject: [PATCH 01/19] [SYCL] Implement SYCL_ONEAPI_accessor_properties
 extension

Signed-off-by: Alexander Batashev <alexander.batashev@intel.com>
---
 sycl/include/CL/sycl/accessor.hpp             | 448 +++++++++++++++++-
 .../CL/sycl/accessor_property_list.hpp        | 176 +++++++
 sycl/include/CL/sycl/buffer.hpp               |  29 +-
 sycl/include/CL/sycl/detail/accessor_impl.hpp |   4 +
 sycl/include/CL/sycl/detail/buffer_impl.hpp   |   3 +-
 sycl/include/CL/sycl/detail/image_impl.hpp    |   3 +-
 .../CL/sycl/detail/property_list_base.hpp     | 114 +++++
 sycl/include/CL/sycl/handler.hpp              |   8 +-
 sycl/include/CL/sycl/image.hpp                |  14 +-
 sycl/include/CL/sycl/interop_handler.hpp      |   3 +-
 sycl/include/CL/sycl/multi_ptr.hpp            |  88 ++--
 .../sycl/properties/accessor_properties.hpp   |  84 ++++
 sycl/include/CL/sycl/property_list.hpp        |  85 +---
 sycl/include/CL/sycl/types.hpp                |  18 +-
 sycl/test/abi/layout_accessors.cpp            |  10 +-
 .../accessor/accessor_property_list.cpp       | 107 +++++
 16 files changed, 1018 insertions(+), 176 deletions(-)
 create mode 100644 sycl/include/CL/sycl/accessor_property_list.hpp
 create mode 100644 sycl/include/CL/sycl/detail/property_list_base.hpp
 create mode 100644 sycl/test/basic_tests/accessor/accessor_property_list.cpp

diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp
index dd4f7fc09be01..8c101a48aea24 100755
--- a/sycl/include/CL/sycl/accessor.hpp
+++ b/sycl/include/CL/sycl/accessor.hpp
@@ -9,6 +9,7 @@
 #pragma once
 
 #include <CL/__spirv/spirv_types.hpp>
+#include <CL/sycl/accessor_property_list.hpp>
 #include <CL/sycl/atomic.hpp>
 #include <CL/sycl/buffer.hpp>
 #include <CL/sycl/detail/accessor_impl.hpp>
@@ -26,6 +27,8 @@
 #include <CL/sycl/property_list.hpp>
 #include <CL/sycl/sampler.hpp>
 
+#include <type_traits>
+
 /// \file accessor.hpp
 /// The file contains implementations of accessor class.
 ///
@@ -214,10 +217,25 @@ namespace sycl {
 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>
+          access::placeholder IsPlaceholder = access::placeholder::false_t,
+          typename PropertyListT = property_list>
 class accessor;
 
 namespace detail {
+template <typename T>
+using IsPropertyListT = typename std::is_base_of<PropertyListBase, T>;
+
+template <typename T>
+using IsRunTimePropertyListT = typename std::is_same<property_list, T>;
+
+template <typename T> struct IsOneapiPropertyListT {
+  constexpr static bool value = false;
+};
+
+template <typename... Props>
+struct IsOneapiPropertyListT<ext::ONEAPI::accessor_property_list<Props...>> {
+  constexpr static bool value = true;
+};
 
 // The function extends or truncates number of dimensions of objects of id
 // or ranges classes. When extending the new values are filled with
@@ -732,7 +750,8 @@ class __image_array_slice__ {
 ///
 /// \ingroup sycl_api_acc
 template <typename DataT, int Dimensions, access::mode AccessMode,
-          access::target AccessTarget, access::placeholder IsPlaceholder>
+          access::target AccessTarget, access::placeholder IsPlaceholder,
+          typename PropertyListT>
 class accessor :
 #ifndef __SYCL_DEVICE_ONLY__
     public detail::AccessorBaseHost,
@@ -751,6 +770,8 @@ class accessor :
                  AccessMode == access::mode::read),
                 "Access mode can be only read for constant buffers");
 
+  static_assert(detail::IsPropertyListT<PropertyListT>::value, "PropertyListT must be property_list or accessor_property_list");
+
   using AccessorCommonT = detail::accessor_common<DataT, Dimensions, AccessMode,
                                                   AccessTarget, IsPlaceholder>;
 
@@ -791,10 +812,10 @@ class accessor :
     return std::is_same<T, DataT>::value && (Dims > 0) && (Dims == Dimensions);
   }
 
-  static access::mode getAdjustedMode(const property_list &PropertyList) {
+  static access::mode getAdjustedMode(const PropertyListT &PropertyList) {
     access::mode AdjustedMode = AccessMode;
 
-    if (PropertyList.has_property<property::noinit>()) {
+    if (PropertyList.template has_property<property::noinit>()) {
       if (AdjustedMode == access::mode::write) {
         AdjustedMode = access::mode::discard_write;
       } else if (AdjustedMode == access::mode::read_write) {
@@ -864,6 +885,14 @@ class accessor :
 #endif // __SYCL_EXPLICIT_SIMD__
   ConcreteASPtrType getQualifiedPtr() const { return MData; }
 
+  template <typename, int, access::mode, access::target, access::placeholder,
+            typename>
+  friend class accessor;
+
+#ifndef __SYCL_DEVICE_ONLY__
+  using AccessorBaseHost::impl;
+#endif
+
 public:
   // Default constructor for objects later initialized with __init member.
   accessor()
@@ -918,6 +947,7 @@ class accessor :
 public:
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
             typename detail::enable_if_t<
+                detail::IsRunTimePropertyListT<PropertyListT>::value &&
                 std::is_same<T, DataT>::value && Dims == 0 &&
                 ((!IsPlaceH && IsHostBuf) ||
                  (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))> * = nullptr>
@@ -939,8 +969,35 @@ class accessor :
 #endif
   }
 
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            typename... PropTypes,
+            typename detail::enable_if_t<
+                detail::IsOneapiPropertyListT<PropertyListT>::value &&
+                std::is_same<T, DataT>::value && Dims == 0 &&
+                ((!IsPlaceH && IsHostBuf) ||
+                 (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))> * = nullptr>
+  accessor(buffer<T, 1, AllocatorT> &BufferRef,
+           const ext::ONEAPI::accessor_property_list<PropTypes...>
+               &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()),
+            getAdjustedMode(PropertyList),
+            detail::getSyclObjImpl(BufferRef).get(), AdjustedDim, sizeof(DataT),
+            BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) {
+    checkDeviceAccessorBufferSize(BufferRef.get_count());
+    if (!IsPlaceH)
+      addHostAccessorAndWait(AccessorBaseHost::impl.get());
+#endif
+  }
+
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
             typename = typename detail::enable_if_t<
+                detail::IsRunTimePropertyListT<PropertyListT>::value &&
                 std::is_same<T, DataT>::value && (Dims == 0) &&
                 (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
   accessor(buffer<T, 1, AllocatorT> &BufferRef, handler &CommandGroupHandler,
@@ -963,10 +1020,37 @@ class accessor :
 #endif
 
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
-            typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>() &&
-                                           ((!IsPlaceH && IsHostBuf) ||
-                                            (IsPlaceH &&
-                                             (IsGlobalBuf || IsConstantBuf)))>>
+            typename... PropTypes,
+            typename = typename detail::enable_if_t<
+                detail::IsOneapiPropertyListT<PropertyListT>::value &&
+                std::is_same<T, DataT>::value && (Dims == 0) &&
+                (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
+  accessor(buffer<T, 1, AllocatorT> &BufferRef, handler &CommandGroupHandler,
+           const ext::ONEAPI::accessor_property_list<PropTypes...>
+               &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()),
+            getAdjustedMode(PropertyList),
+            detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
+            BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) {
+    checkDeviceAccessorBufferSize(BufferRef.get_count());
+    detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
+  }
+#endif
+
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            typename = detail::enable_if_t<
+                detail::IsRunTimePropertyListT<PropertyListT>::value &&
+                IsSameAsBuffer<T, Dims>() &&
+                ((!IsPlaceH && IsHostBuf) ||
+                 (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>>
   accessor(buffer<T, Dims, AllocatorT> &BufferRef,
            const property_list &PropertyList = {})
 #ifdef __SYCL_DEVICE_ONLY__
@@ -987,21 +1071,61 @@ class accessor :
   }
 #endif
 
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            typename... PropTypes,
+            typename = detail::enable_if_t<
+                detail::IsOneapiPropertyListT<PropertyListT>::value &&
+                IsSameAsBuffer<T, Dims>() &&
+                ((!IsPlaceH && IsHostBuf) ||
+                 (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>>
+  accessor(buffer<T, Dims, AllocatorT> &BufferRef,
+           const ext::ONEAPI::accessor_property_list<PropTypes...>
+               &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()),
+            getAdjustedMode(PropertyList),
+            detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
+            BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) {
+    checkDeviceAccessorBufferSize(BufferRef.get_count());
+    if (!IsPlaceH)
+      addHostAccessorAndWait(AccessorBaseHost::impl.get());
+  }
+#endif
+
 #if __cplusplus > 201402L
 
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
             typename TagT,
             typename = detail::enable_if_t<
+                detail::IsRunTimePropertyListT<PropertyListT>::value &&
                 IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && IsPlaceH &&
                 (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
   accessor(buffer<T, Dims, AllocatorT> &BufferRef, TagT,
            const property_list &PropertyList = {})
       : accessor(BufferRef, PropertyList) {}
 
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            typename TagT, typename... PropTypes,
+            typename = detail::enable_if_t<
+                detail::IsOneapiPropertyListT<PropertyListT>::value &&
+                IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && IsPlaceH &&
+                (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
+  accessor(buffer<T, Dims, AllocatorT> &BufferRef, TagT,
+           const ext::ONEAPI::accessor_property_list<PropTypes...>
+               &PropertyList = {})
+      : accessor(BufferRef, PropertyList) {}
 #endif
 
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
             typename = detail::enable_if_t<
+                detail::IsRunTimePropertyListT<PropertyListT>::value &&
                 IsSameAsBuffer<T, Dims>() &&
                 (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
   accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
@@ -1024,45 +1148,112 @@ class accessor :
   }
 #endif
 
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            typename... PropTypes,
+            typename = detail::enable_if_t<
+                detail::IsOneapiPropertyListT<PropertyListT>::value &&
+                IsSameAsBuffer<T, Dims>() &&
+                (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
+  accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
+           const ext::ONEAPI::accessor_property_list<PropTypes...>
+               &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()),
+            getAdjustedMode(PropertyList),
+            detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
+            BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) {
+    checkDeviceAccessorBufferSize(BufferRef.get_count());
+    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<
+                detail::IsRunTimePropertyListT<PropertyListT>::value &&
                 IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && !IsPlaceH &&
                 (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
   accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
            TagT, const property_list &PropertyList = {})
       : accessor(BufferRef, CommandGroupHandler, PropertyList) {}
 
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            typename TagT, typename... PropTypes,
+            typename = detail::enable_if_t<
+                detail::IsOneapiPropertyListT<PropertyListT>::value &&
+                IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && !IsPlaceH &&
+                (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
+  accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
+           TagT,
+           const ext::ONEAPI::accessor_property_list<PropTypes...>
+               &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)))>>
+            typename = detail::enable_if_t<
+                detail::IsRunTimePropertyListT<PropertyListT>::value &&
+                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) {}
 
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            typename... PropTypes,
+            typename = detail::enable_if_t<
+                detail::IsOneapiPropertyListT<PropertyListT>::value &&
+                IsSameAsBuffer<T, Dims>() &&
+                ((!IsPlaceH && IsHostBuf) ||
+                 (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>>
+  accessor(buffer<T, Dims, AllocatorT> &BufferRef,
+           range<Dimensions> AccessRange,
+           const ext::ONEAPI::accessor_property_list<PropTypes...>
+               &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)>>
+            typename = detail::enable_if_t<
+                detail::IsRunTimePropertyListT<PropertyListT>::value &&
+                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) {}
 
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            typename TagT, typename... PropTypes,
+            typename = detail::enable_if_t<
+                detail::IsOneapiPropertyListT<PropertyListT>::value &&
+                IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && IsPlaceH &&
+                (IsGlobalBuf || IsConstantBuf)>>
+  accessor(buffer<T, Dims, AllocatorT> &BufferRef,
+           range<Dimensions> AccessRange, TagT,
+           const ext::ONEAPI::accessor_property_list<PropTypes...>
+               &PropertyList = {})
+      : accessor(BufferRef, AccessRange, {}, PropertyList) {}
 #endif
 
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
             typename = detail::enable_if_t<
+                detail::IsRunTimePropertyListT<PropertyListT>::value &&
                 IsSameAsBuffer<T, Dims>() &&
                 (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
   accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
@@ -1071,11 +1262,25 @@ class accessor :
       : accessor(BufferRef, CommandGroupHandler, AccessRange, {},
                  PropertyList) {}
 
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            typename... PropTypes,
+            typename = detail::enable_if_t<
+                detail::IsOneapiPropertyListT<PropertyListT>::value &&
+                IsSameAsBuffer<T, Dims>() &&
+                (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
+  accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
+           range<Dimensions> AccessRange,
+           const ext::ONEAPI::accessor_property_list<PropTypes...>
+               &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<
+                detail::IsRunTimePropertyListT<PropertyListT>::value &&
                 IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && !IsPlaceH &&
                 (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
   accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
@@ -1084,13 +1289,26 @@ class accessor :
       : accessor(BufferRef, CommandGroupHandler, AccessRange, {},
                  PropertyList) {}
 
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            typename TagT, typename... PropTypes,
+            typename = detail::enable_if_t<
+                detail::IsOneapiPropertyListT<PropertyListT>::value &&
+                IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && !IsPlaceH &&
+                (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
+  accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
+           range<Dimensions> AccessRange, TagT,
+           const ext::ONEAPI::accessor_property_list<PropTypes...>
+               &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)))>>
+            typename = detail::enable_if_t<
+                detail::IsRunTimePropertyListT<PropertyListT>::value &&
+                IsSameAsBuffer<T, Dims>() &&
+                ((!IsPlaceH && IsHostBuf) ||
+                 (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>>
   accessor(buffer<T, Dims, AllocatorT> &BufferRef,
            range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
            const property_list &PropertyList = {})
@@ -1112,22 +1330,64 @@ class accessor :
   }
 #endif
 
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            typename... PropTypes,
+            typename = detail::enable_if_t<
+                detail::IsOneapiPropertyListT<PropertyListT>::value &&
+                IsSameAsBuffer<T, Dims>() &&
+                ((!IsPlaceH && IsHostBuf) ||
+                 (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>>
+  accessor(buffer<T, Dims, AllocatorT> &BufferRef,
+           range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
+           const ext::ONEAPI::accessor_property_list<PropTypes...>
+               &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()),
+                         getAdjustedMode(PropertyList),
+                         detail::getSyclObjImpl(BufferRef).get(), Dimensions,
+                         sizeof(DataT), BufferRef.OffsetInBytes,
+                         BufferRef.IsSubBuffer) {
+    checkDeviceAccessorBufferSize(BufferRef.get_count());
+    if (!IsPlaceH)
+      addHostAccessorAndWait(AccessorBaseHost::impl.get());
+  }
+#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)>>
+            typename = detail::enable_if_t<
+                detail::IsRunTimePropertyListT<PropertyListT>::value &&
+                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) {}
 
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            typename TagT, typename... PropTypes,
+            typename = detail::enable_if_t<
+                detail::IsOneapiPropertyListT<PropertyListT>::value &&
+                IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && IsPlaceH &&
+                (IsGlobalBuf || IsConstantBuf)>>
+  accessor(buffer<T, Dims, AllocatorT> &BufferRef,
+           range<Dimensions> AccessRange, id<Dimensions> AccessOffset, TagT,
+           const ext::ONEAPI::accessor_property_list<PropTypes...>
+               &PropertyList = {})
+      : accessor(BufferRef, AccessRange, AccessOffset, PropertyList) {}
 #endif
 
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
             typename = detail::enable_if_t<
+                detail::IsRunTimePropertyListT<PropertyListT>::value &&
                 IsSameAsBuffer<T, Dims>() &&
                 (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
   accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
@@ -1151,11 +1411,40 @@ class accessor :
   }
 #endif
 
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            typename... PropTypes,
+            typename = detail::enable_if_t<
+                detail::IsOneapiPropertyListT<PropertyListT>::value &&
+                IsSameAsBuffer<T, Dims>() &&
+                (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
+  accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
+           range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
+           const ext::ONEAPI::accessor_property_list<PropTypes...>
+               &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()),
+                         getAdjustedMode(PropertyList),
+                         detail::getSyclObjImpl(BufferRef).get(), Dimensions,
+                         sizeof(DataT), BufferRef.OffsetInBytes,
+                         BufferRef.IsSubBuffer) {
+    checkDeviceAccessorBufferSize(BufferRef.get_count());
+    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<
+                detail::IsRunTimePropertyListT<PropertyListT>::value &&
                 IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && !IsPlaceH &&
                 (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
   accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
@@ -1164,7 +1453,35 @@ class accessor :
       : accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
                  PropertyList) {}
 
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            typename TagT, typename... PropTypes,
+            typename = detail::enable_if_t<
+                detail::IsOneapiPropertyListT<PropertyListT>::value &&
+                IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && !IsPlaceH &&
+                (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
+  accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
+           range<Dimensions> AccessRange, id<Dimensions> AccessOffset, TagT,
+           const ext::ONEAPI::accessor_property_list<PropTypes...>
+               &PropertyList = {})
+      : accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
+                 PropertyList) {}
+#endif
+
+  template <typename... NewPropsT>
+  accessor(
+      const accessor<DataT, Dimensions, AccessMode, AccessTarget, IsPlaceholder,
+                     ext::ONEAPI::accessor_property_list<NewPropsT...>> &Other)
+#ifdef __SYCL_DEVICE_ONLY__
+      : impl(Other.impl)
+#else
+      : detail::AccessorBaseHost(Other)
 #endif
+  {
+    static_assert(detail::IsOneapiPropertyListT<PropertyListT>::value,
+                  "Conversion is only available for accessor_property_list");
+    static_assert(PropertyListT::template areSameCxProperties<NewPropsT...>(),
+                  "Compile-time-constant properties must be the same");
+  }
 
   constexpr bool is_placeholder() const { return IsPlaceH; }
 
@@ -1287,12 +1604,29 @@ accessor(buffer<DataT, Dimensions, AllocatorT>)
     ->accessor<DataT, Dimensions, access::mode::read_write,
                target::global_buffer, access::placeholder::true_t>;
 
+template <typename DataT, int Dimensions, typename AllocatorT,
+          typename... PropsT>
+accessor(buffer<DataT, Dimensions, AllocatorT>,
+         const ext::ONEAPI::accessor_property_list<PropsT...> &)
+    ->accessor<DataT, Dimensions, access::mode::read_write,
+               target::global_buffer, access::placeholder::true_t,
+               ext::ONEAPI::accessor_property_list<PropsT...>>;
+
 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1>
 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1)
     ->accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type1>(),
                detail::deduceAccessTarget<Type1, Type1>(target::global_buffer),
                access::placeholder::true_t>;
 
+template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
+          typename... PropsT>
+accessor(buffer<DataT, Dimensions, AllocatorT>, Type1,
+         const ext::ONEAPI::accessor_property_list<PropsT...> &)
+    ->accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type1>(),
+               detail::deduceAccessTarget<Type1, Type1>(target::global_buffer),
+               access::placeholder::true_t,
+               ext::ONEAPI::accessor_property_list<PropsT...>>;
+
 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
           typename Type2>
 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2)
@@ -1300,6 +1634,15 @@ accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2)
                detail::deduceAccessTarget<Type1, Type2>(target::global_buffer),
                access::placeholder::true_t>;
 
+template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
+          typename Type2, typename... PropsT>
+accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2,
+         const ext::ONEAPI::accessor_property_list<PropsT...> &)
+    ->accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type2>(),
+               detail::deduceAccessTarget<Type1, Type2>(target::global_buffer),
+               access::placeholder::true_t,
+               ext::ONEAPI::accessor_property_list<PropsT...>>;
+
 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
           typename Type2, typename Type3>
 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3)
@@ -1307,6 +1650,15 @@ accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3)
                detail::deduceAccessTarget<Type2, Type3>(target::global_buffer),
                access::placeholder::true_t>;
 
+template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
+          typename Type2, typename Type3, typename... PropsT>
+accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3,
+         const ext::ONEAPI::accessor_property_list<PropsT...> &)
+    ->accessor<DataT, Dimensions, detail::deduceAccessMode<Type2, Type3>(),
+               detail::deduceAccessTarget<Type2, Type3>(target::global_buffer),
+               access::placeholder::true_t,
+               ext::ONEAPI::accessor_property_list<PropsT...>>;
+
 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
           typename Type2, typename Type3, typename Type4>
 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3, Type4)
@@ -1314,17 +1666,43 @@ accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3, Type4)
                detail::deduceAccessTarget<Type3, Type4>(target::global_buffer),
                access::placeholder::true_t>;
 
+template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
+          typename Type2, typename Type3, typename Type4, typename... PropsT>
+accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3, Type4,
+         const ext::ONEAPI::accessor_property_list<PropsT...> &)
+    ->accessor<DataT, Dimensions, detail::deduceAccessMode<Type3, Type4>(),
+               detail::deduceAccessTarget<Type3, Type4>(target::global_buffer),
+               access::placeholder::true_t,
+               ext::ONEAPI::accessor_property_list<PropsT...>>;
+
 template <typename DataT, int Dimensions, typename AllocatorT>
 accessor(buffer<DataT, Dimensions, AllocatorT>, handler)
     ->accessor<DataT, Dimensions, access::mode::read_write,
                target::global_buffer, access::placeholder::false_t>;
 
+template <typename DataT, int Dimensions, typename AllocatorT,
+          typename... PropsT>
+accessor(buffer<DataT, Dimensions, AllocatorT>, handler,
+         const ext::ONEAPI::accessor_property_list<PropsT...> &)
+    ->accessor<DataT, Dimensions, access::mode::read_write,
+               target::global_buffer, access::placeholder::false_t,
+               ext::ONEAPI::accessor_property_list<PropsT...>>;
+
 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1>
 accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1)
     ->accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type1>(),
                detail::deduceAccessTarget<Type1, Type1>(target::global_buffer),
                access::placeholder::false_t>;
 
+template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
+          typename... PropsT>
+accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1,
+         const ext::ONEAPI::accessor_property_list<PropsT...> &)
+    ->accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type1>(),
+               detail::deduceAccessTarget<Type1, Type1>(target::global_buffer),
+               access::placeholder::false_t,
+               ext::ONEAPI::accessor_property_list<PropsT...>>;
+
 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
           typename Type2>
 accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2)
@@ -1332,6 +1710,15 @@ accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2)
                detail::deduceAccessTarget<Type1, Type2>(target::global_buffer),
                access::placeholder::false_t>;
 
+template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
+          typename Type2, typename... PropsT>
+accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2,
+         const ext::ONEAPI::accessor_property_list<PropsT...> &)
+    ->accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type2>(),
+               detail::deduceAccessTarget<Type1, Type2>(target::global_buffer),
+               access::placeholder::false_t,
+               ext::ONEAPI::accessor_property_list<PropsT...>>;
+
 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
           typename Type2, typename Type3>
 accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2, Type3)
@@ -1339,6 +1726,15 @@ accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2, Type3)
                detail::deduceAccessTarget<Type2, Type3>(target::global_buffer),
                access::placeholder::false_t>;
 
+template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
+          typename Type2, typename Type3, typename... PropsT>
+accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2, Type3,
+         const ext::ONEAPI::accessor_property_list<PropsT...> &)
+    ->accessor<DataT, Dimensions, detail::deduceAccessMode<Type2, Type3>(),
+               detail::deduceAccessTarget<Type2, Type3>(target::global_buffer),
+               access::placeholder::false_t,
+               ext::ONEAPI::accessor_property_list<PropsT...>>;
+
 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
           typename Type2, typename Type3, typename Type4>
 accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2, Type3,
@@ -1347,6 +1743,14 @@ accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2, Type3,
                detail::deduceAccessTarget<Type3, Type4>(target::global_buffer),
                access::placeholder::false_t>;
 
+template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
+          typename Type2, typename Type3, typename Type4, typename... PropsT>
+accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2, Type3,
+         Type4, const ext::ONEAPI::accessor_property_list<PropsT...> &)
+    ->accessor<DataT, Dimensions, detail::deduceAccessMode<Type3, Type4>(),
+               detail::deduceAccessTarget<Type3, Type4>(target::global_buffer),
+               access::placeholder::false_t,
+               ext::ONEAPI::accessor_property_list<PropsT...>>;
 #endif
 
 /// Local accessor
diff --git a/sycl/include/CL/sycl/accessor_property_list.hpp b/sycl/include/CL/sycl/accessor_property_list.hpp
new file mode 100644
index 0000000000000..84b366ed7a2c0
--- /dev/null
+++ b/sycl/include/CL/sycl/accessor_property_list.hpp
@@ -0,0 +1,176 @@
+//==----- accessor_property_list.hpp --- SYCL accessor property list -------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <CL/sycl/access/access.hpp>
+#include <CL/sycl/detail/common.hpp>
+#include <CL/sycl/detail/property_list_base.hpp>
+#include <CL/sycl/property_list.hpp>
+
+#include <type_traits>
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace detail {} // namespace detail
+namespace ext {
+namespace ONEAPI {
+
+template <typename T> struct is_property : std::false_type {};
+
+template <typename T> struct is_compile_time_property : std::false_type {};
+
+/// Objects of the accessor_property_list class are containers for the SYCL
+/// properties.
+///
+/// Unlike \c property_list, accessor_property_list can take
+/// compile-time-constant properties.
+///
+/// \sa accessor
+/// \sa property_list
+///
+/// \ingroup sycl_api
+template <typename... PropsT>
+class accessor_property_list : protected detail::PropertyListBase {
+  // These structures check if compile-time-constant property is present in
+  // list. For runtime properties this check is always true.
+  template <class T, class U> struct AreSameTemplate : std::is_same<T, U> {};
+  template <template <class...> class T, class T1, class T2>
+  struct AreSameTemplate<T<T1>, T<T2>> : std::true_type {};
+#if __cplusplus >= 201703L
+  template <template <auto...> class T, auto...T1, auto...T2>
+  struct AreSameTemplate<T<T1...>, T<T2...>> : std::true_type {};
+#endif
+  template <typename PropT, typename... PropListT>
+  struct ContainsProperty;
+  template <typename PropT> struct ContainsProperty<PropT> : std::false_type {};
+  template <typename PropT, typename Head, typename... Tail>
+  struct ContainsProperty<PropT, Head, Tail...>
+      : std::conditional<AreSameTemplate<PropT, Head>::value ||
+                             !is_compile_time_property<PropT>::value,
+                         std::true_type,
+                         ContainsProperty<PropT, Tail...>>::type {};
+
+  // The following structures help to check if two property lists contain the
+  // same compile-time-constant properties.
+  template <typename...> struct PropertyContainer {
+    using Head = void;
+    using Rest = void;
+  };
+  template <typename T, typename... Other>
+  struct PropertyContainer<T, Other...> {
+    using Head = T;
+    using Rest = PropertyContainer<Other...>;
+  };
+  template <typename T> struct PropertyContainer<T> {
+    using Head = T;
+    using Rest = void;
+  };
+
+  template <typename ContainerT, typename... OtherProps>
+  struct ContainsSameProperties
+      : std::conditional<
+            ContainsProperty<typename ContainerT::Head, OtherProps...>::value,
+            ContainsSameProperties<typename ContainerT::Rest, OtherProps...>,
+            std::false_type>::type {};
+  template <typename... OtherProps>
+  struct ContainsSameProperties<void, OtherProps...> : std::true_type {};
+
+#if __cplusplus >= 201703L
+  template <typename ContainerT, template <auto...> class PropT, auto... Args>
+  struct GetCxPropertyHelper {
+    using type = typename std::conditional_t<
+        AreSameTemplate<typename ContainerT::Head, PropT<Args...>>::value,
+        typename ContainerT::Head,
+        typename GetCxPropertyHelper<typename ContainerT::Rest, PropT,
+                                     Args...>::type>;
+  };
+  template <typename Head, template <auto...> class PropT, auto... Args>
+  struct GetCxPropertyHelper<PropertyContainer<Head>, PropT, Args...> {
+    using type = typename std::conditional_t<
+        AreSameTemplate<Head, PropT<Args...>>::value, Head, void>;
+  };
+#endif
+
+  // The structs validate that all objects passed are SYCL properties
+  template <typename... Tail> struct AllProperties : std::true_type {};
+  template <typename T, typename... Tail>
+  struct AllProperties<T, Tail...>
+      : std::conditional<
+            std::is_base_of<detail::DataLessPropertyBase, T>::value ||
+                std::is_base_of<detail::PropertyWithDataBase, T>::value ||
+                is_compile_time_property<T>::value,
+            AllProperties<Tail...>, std::false_type>::type {};
+
+public:
+  template <
+      typename = typename std::enable_if<AllProperties<PropsT...>::value>::type>
+  accessor_property_list(PropsT... Props) : detail::PropertyListBase(false) {
+    ctorHelper(Props...);
+  }
+
+  template <typename... OtherProps,
+            typename = typename std::enable_if<
+                ContainsSameProperties<PropertyContainer<PropsT...>,
+                                       OtherProps...>::value &&
+                ContainsSameProperties<PropertyContainer<OtherProps...>,
+                                       PropsT...>::value>::type>
+  accessor_property_list(const accessor_property_list<OtherProps...> &OtherList)
+      : detail::PropertyListBase(OtherList.MDataLessProps,
+                                 OtherList.MPropsWithData) {}
+
+  template <typename PropT,
+            typename = std::enable_if<!is_compile_time_property<PropT>::value>>
+  PropT get_property() const {
+    if (!has_property<PropT>())
+      throw sycl::invalid_object_error("The property is not found",
+                                       PI_INVALID_VALUE);
+
+    return get_property_helper<PropT>();
+  }
+
+  template <class PropT>
+  typename std::enable_if<!is_compile_time_property<PropT>::value, bool>::type
+  has_property() const {
+    return has_property_helper<PropT>();
+  }
+
+#if __cplusplus >= 201703L
+  template <template <auto...> class PropT, auto... Args>
+  static constexpr
+      typename std::enable_if<is_compile_time_property<PropT<Args...>>::value,
+                              bool>::type
+      has_property() {
+    return ContainsProperty<PropT<Args...>, PropsT...>::value;
+  }
+  template <template <auto...> class PropT, auto... Args,
+            typename = typename std::enable_if<
+                is_compile_time_property<PropT<Args...>>::value>::type>
+  static constexpr auto get_property() {
+    return typename GetCxPropertyHelper<PropertyContainer<PropsT...>, PropT,
+                                        Args...>::type{};
+  }
+#endif
+
+private:
+  template <typename... OtherProps> friend class accessor_property_list;
+  template <typename, int, sycl::access::mode, sycl::access::target,
+            sycl::access::placeholder, typename>
+  friend class accessor;
+
+public:
+  template <typename... OtherPropsT>
+  static constexpr bool areSameCxProperties() {
+    return ContainsSameProperties<PropertyContainer<OtherPropsT...>,
+                                  PropsT...>::value;
+  }
+};
+} // namespace ONEAPI
+} // namespace ext
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)
diff --git a/sycl/include/CL/sycl/buffer.hpp b/sycl/include/CL/sycl/buffer.hpp
index 7d5c8a8b59e1b..251a962d4340f 100755
--- a/sycl/include/CL/sycl/buffer.hpp
+++ b/sycl/include/CL/sycl/buffer.hpp
@@ -8,6 +8,7 @@
 
 #pragma once
 
+#include <CL/sycl/accessor_property_list.hpp>
 #include <CL/sycl/detail/buffer_impl.hpp>
 #include <CL/sycl/detail/common.hpp>
 #include <CL/sycl/detail/stl_type_traits.hpp>
@@ -247,36 +248,42 @@ class buffer {
 
   template <access::mode Mode,
             access::target Target = access::target::global_buffer>
-  accessor<T, dimensions, Mode, Target, access::placeholder::false_t>
+  accessor<T, dimensions, Mode, Target, access::placeholder::false_t,
+           property_list>
   get_access(handler &CommandGroupHandler) {
-    return accessor<T, dimensions, Mode, Target, access::placeholder::false_t>(
-        *this, CommandGroupHandler);
+    return accessor<T, dimensions, Mode, Target, access::placeholder::false_t,
+                    property_list>(*this,
+                                                      CommandGroupHandler);
   }
 
   template <access::mode mode>
   accessor<T, dimensions, mode, access::target::host_buffer,
-           access::placeholder::false_t>
+           access::placeholder::false_t, property_list>
   get_access() {
     return accessor<T, dimensions, mode, access::target::host_buffer,
-                    access::placeholder::false_t>(*this);
+                    access::placeholder::false_t,
+                    property_list>(*this);
   }
 
   template <access::mode mode,
             access::target target = access::target::global_buffer>
-  accessor<T, dimensions, mode, target, access::placeholder::false_t>
+  accessor<T, dimensions, mode, target, access::placeholder::false_t,
+           property_list>
   get_access(handler &commandGroupHandler, range<dimensions> accessRange,
              id<dimensions> accessOffset = {}) {
-    return accessor<T, dimensions, mode, target, access::placeholder::false_t>(
+    return accessor<T, dimensions, mode, target, access::placeholder::false_t,
+                    property_list>(
         *this, commandGroupHandler, accessRange, accessOffset);
   }
 
   template <access::mode mode>
   accessor<T, dimensions, mode, access::target::host_buffer,
-           access::placeholder::false_t>
+           access::placeholder::false_t, property_list>
   get_access(range<dimensions> accessRange, id<dimensions> accessOffset = {}) {
     return accessor<T, dimensions, mode, access::target::host_buffer,
-                    access::placeholder::false_t>(*this, accessRange,
-                                                  accessOffset);
+                    access::placeholder::false_t,
+                    property_list>(*this, accessRange,
+                                                      accessOffset);
   }
 
 #if __cplusplus > 201402L
@@ -339,7 +346,7 @@ class buffer {
   template <typename A, int dims, typename C, typename Enable>
   friend class buffer;
   template <typename DataT, int dims, access::mode mode, access::target target,
-            access::placeholder isPlaceholder>
+            access::placeholder isPlaceholder, typename PropertyListT>
   friend class accessor;
   range<dimensions> Range;
   // Offset field specifies the origin of the sub buffer inside the parent
diff --git a/sycl/include/CL/sycl/detail/accessor_impl.hpp b/sycl/include/CL/sycl/detail/accessor_impl.hpp
index 277b740e6a229..792e848e636b5 100644
--- a/sycl/include/CL/sycl/detail/accessor_impl.hpp
+++ b/sycl/include/CL/sycl/detail/accessor_impl.hpp
@@ -167,6 +167,10 @@ class AccessorBaseHost {
   template <class Obj>
   friend decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject);
 
+  template <typename, int, access::mode, access::target, access::placeholder,
+            typename>
+  friend class accessor;
+
   AccessorImplPtr impl;
 
 private:
diff --git a/sycl/include/CL/sycl/detail/buffer_impl.hpp b/sycl/include/CL/sycl/detail/buffer_impl.hpp
index 0ff488cf97363..df6600af3ec86 100644
--- a/sycl/include/CL/sycl/detail/buffer_impl.hpp
+++ b/sycl/include/CL/sycl/detail/buffer_impl.hpp
@@ -27,7 +27,8 @@ __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
 // Forward declarations
 template <typename DataT, int Dimensions, access::mode AccessMode,
-          access::target AccessTarget, access::placeholder IsPlaceholder>
+          access::target AccessTarget, access::placeholder IsPlaceholder,
+          typename PropertyListT>
 class accessor;
 template <typename T, int Dimensions, typename AllocatorT, typename Enable>
 class buffer;
diff --git a/sycl/include/CL/sycl/detail/image_impl.hpp b/sycl/include/CL/sycl/detail/image_impl.hpp
index b5def72028ae8..d3ef0b7e2af59 100644
--- a/sycl/include/CL/sycl/detail/image_impl.hpp
+++ b/sycl/include/CL/sycl/detail/image_impl.hpp
@@ -28,7 +28,8 @@ enum class image_channel_type : unsigned int;
 
 template <int Dimensions, typename AllocatorT> class image;
 template <typename DataT, int Dimensions, access::mode AccessMode,
-          access::target AccessTarget, access::placeholder IsPlaceholder>
+          access::target AccessTarget, access::placeholder IsPlaceholder,
+          typename property_listT>
 class accessor;
 class handler;
 
diff --git a/sycl/include/CL/sycl/detail/property_list_base.hpp b/sycl/include/CL/sycl/detail/property_list_base.hpp
new file mode 100644
index 0000000000000..b02d1cd80c0b1
--- /dev/null
+++ b/sycl/include/CL/sycl/detail/property_list_base.hpp
@@ -0,0 +1,114 @@
+//==------- property_list_base.hpp --- Base for SYCL property lists --------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <CL/sycl/detail/common.hpp>
+#include <CL/sycl/detail/property_helper.hpp>
+
+#include <bitset>
+#include <memory>
+#include <type_traits>
+#include <vector>
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace detail {
+class PropertyListBase {
+protected:
+  explicit PropertyListBase(
+      std::bitset<DataLessPropKind::DataLessPropKindSize> DataLessProps)
+      : MDataLessProps(DataLessProps) {}
+  PropertyListBase(
+      std::bitset<DataLessPropKind::DataLessPropKindSize> DataLessProps,
+      std::vector<std::shared_ptr<PropertyWithDataBase>> PropsWithData)
+      : MDataLessProps(DataLessProps),
+        MPropsWithData(std::move(PropsWithData)) {}
+  void ctorHelper() {}
+
+  template <typename... PropsT, class PropT>
+  typename std::enable_if<
+      std::is_base_of<DataLessPropertyBase, PropT>::value>::type
+  ctorHelper(PropT &, PropsT... Props) {
+    const int PropKind = static_cast<int>(PropT::getKind());
+    MDataLessProps[PropKind] = true;
+    ctorHelper(Props...);
+  }
+
+  template <typename... PropsT, class PropT>
+  typename std::enable_if<
+      std::is_base_of<PropertyWithDataBase, PropT>::value>::type
+  ctorHelper(PropT &Prop, PropsT... Props) {
+    MPropsWithData.emplace_back(new PropT(Prop));
+    ctorHelper(Props...);
+  }
+
+  // Compile-time-constant properties are simply skipped
+  template <typename... PropsT, class PropT>
+  typename std::enable_if<
+      !std::is_base_of<PropertyWithDataBase, PropT>::value &&
+      !std::is_base_of<DataLessPropertyBase, PropT>::value>::type
+  ctorHelper(PropT &, PropsT... Props) {
+    ctorHelper(Props...);
+  }
+
+  template <typename PropT>
+  typename std::enable_if<std::is_base_of<DataLessPropertyBase, PropT>::value,
+                          bool>::type
+  has_property_helper() const {
+    const int PropKind = static_cast<int>(PropT::getKind());
+    if (PropKind >= detail::DataLessPropKind::DataLessPropKindSize)
+      return false;
+    return MDataLessProps[PropKind];
+  }
+
+  template <typename PropT>
+  typename std::enable_if<std::is_base_of<PropertyWithDataBase, PropT>::value,
+                          bool>::type
+  has_property_helper() const {
+    const int PropKind = static_cast<int>(PropT::getKind());
+    for (const std::shared_ptr<PropertyWithDataBase> &Prop : MPropsWithData)
+      if (Prop->isSame(PropKind))
+        return true;
+    return false;
+  }
+
+  template <typename PropT>
+  typename std::enable_if<std::is_base_of<DataLessPropertyBase, PropT>::value,
+                          PropT>::type
+  get_property_helper() const {
+    // In case of simple property we can just construct it
+    return PropT{};
+  }
+
+  template <typename PropT>
+  typename std::enable_if<std::is_base_of<PropertyWithDataBase, PropT>::value,
+                          PropT>::type
+  get_property_helper() const {
+    const int PropKind = static_cast<int>(PropT::getKind());
+    if (PropKind >= PropWithDataKind::PropWithDataKindSize)
+      throw sycl::invalid_object_error("The property is not found",
+                                       PI_INVALID_VALUE);
+
+    for (const std::shared_ptr<PropertyWithDataBase> &Prop : MPropsWithData)
+      if (Prop->isSame(PropKind))
+        return *static_cast<PropT *>(Prop.get());
+
+    throw sycl::invalid_object_error("The property is not found",
+                                     PI_INVALID_VALUE);
+  }
+
+private:
+  // Stores enable/not enabled for simple properties
+  std::bitset<DataLessPropKind::DataLessPropKindSize> MDataLessProps;
+  // Stores shared_ptrs to complex properties
+  std::vector<std::shared_ptr<PropertyWithDataBase>> MPropsWithData;
+};
+} // namespace detail
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)
\ No newline at end of file
diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp
index 7506ccacd19ef..e304aab921237 100644
--- a/sycl/include/CL/sycl/handler.hpp
+++ b/sycl/include/CL/sycl/handler.hpp
@@ -1767,8 +1767,9 @@ class __SYCL_EXPORT handler {
   /// \param Pattern is a value to be used to fill the memory.
   template <typename T, int Dims, access::mode AccessMode,
             access::target AccessTarget,
-            access::placeholder IsPlaceholder = access::placeholder::false_t>
-  void fill(accessor<T, Dims, AccessMode, AccessTarget, IsPlaceholder> Dst,
+            access::placeholder IsPlaceholder = access::placeholder::false_t,
+            typename PropertyListT = property_list>
+  void fill(accessor<T, Dims, AccessMode, AccessTarget, IsPlaceholder, PropertyListT> Dst,
             const T &Pattern) {
     throwIfActionIsCreated();
     // TODO add check:T must be an integral scalar value or a SYCL vector type
@@ -1918,7 +1919,8 @@ class __SYCL_EXPORT handler {
   friend class detail::queue_impl;
   // Make accessor class friend to keep the list of associated accessors.
   template <typename DataT, int Dims, access::mode AccMode,
-            access::target AccTarget, access::placeholder isPlaceholder>
+            access::target AccTarget, access::placeholder isPlaceholder,
+            typename PropertyListT>
   friend class accessor;
   friend device detail::getDeviceFromHandler(handler &);
 
diff --git a/sycl/include/CL/sycl/image.hpp b/sycl/include/CL/sycl/image.hpp
index 128ad8b861479..221448250f02c 100644
--- a/sycl/include/CL/sycl/image.hpp
+++ b/sycl/include/CL/sycl/image.hpp
@@ -8,6 +8,7 @@
 
 #pragma once
 
+#include "CL/sycl/accessor_property_list.hpp"
 #include <CL/sycl/detail/common.hpp>
 #include <CL/sycl/detail/generic_type_traits.hpp>
 #include <CL/sycl/detail/image_impl.hpp>
@@ -272,18 +273,23 @@ class image {
 
   template <typename DataT, access::mode AccessMode>
   accessor<detail::EnableIfImgAccDataT<DataT>, Dimensions, AccessMode,
-           access::target::image, access::placeholder::false_t>
+           access::target::image, access::placeholder::false_t,
+           property_list>
   get_access(handler &commandGroupHandler) {
     return accessor<DataT, Dimensions, AccessMode, access::target::image,
-                    access::placeholder::false_t>(*this, commandGroupHandler);
+                    access::placeholder::false_t,
+                    property_list>(*this,
+                                                      commandGroupHandler);
   }
 
   template <typename DataT, access::mode AccessMode>
   accessor<detail::EnableIfImgAccDataT<DataT>, Dimensions, AccessMode,
-           access::target::host_image, access::placeholder::false_t>
+           access::target::host_image, access::placeholder::false_t,
+           property_list>
   get_access() {
     return accessor<DataT, Dimensions, AccessMode, access::target::host_image,
-                    access::placeholder::false_t>(*this);
+                    access::placeholder::false_t,
+                    property_list>(*this);
   }
 
   template <typename Destination = std::nullptr_t>
diff --git a/sycl/include/CL/sycl/interop_handler.hpp b/sycl/include/CL/sycl/interop_handler.hpp
index 84f9f7c8932d9..35c6cf88c6760 100644
--- a/sycl/include/CL/sycl/interop_handler.hpp
+++ b/sycl/include/CL/sycl/interop_handler.hpp
@@ -18,7 +18,8 @@ namespace sycl {
 class interop_handler {
   // Make accessor class friend to access the detail mem objects
   template <typename DataT, int Dims, access::mode AccMode,
-            access::target AccTarget, access::placeholder isPlaceholder>
+            access::target AccTarget, access::placeholder isPlaceholder,
+            typename PropertyListT>
   friend class accessor;
 
 public:
diff --git a/sycl/include/CL/sycl/multi_ptr.hpp b/sycl/include/CL/sycl/multi_ptr.hpp
index a3c6bc4768c38..f15d03d5bc5a5 100644
--- a/sycl/include/CL/sycl/multi_ptr.hpp
+++ b/sycl/include/CL/sycl/multi_ptr.hpp
@@ -19,7 +19,8 @@ __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
 // Forward declaration
 template <typename dataT, int dimensions, access::mode accessMode,
-          access::target accessTarget, access::placeholder isPlaceholder>
+          access::target accessTarget, access::placeholder isPlaceholder,
+          typename PropertyListT>
 class accessor;
 
 /// Provides constructors for address space qualified and non address space
@@ -110,39 +111,41 @@ template <typename ElementType, access::address_space Space> class multi_ptr {
 
   // Only if Space == global_space || global_device_space
   template <int dimensions, access::mode Mode,
-            access::placeholder isPlaceholder,
+            access::placeholder isPlaceholder, typename PropertyListT,
             access::address_space _Space = Space,
             typename = typename std::enable_if<
                 _Space == Space &&
                 (Space == access::address_space::global_space ||
                  Space == access::address_space::global_device_space)>::type>
-  multi_ptr(accessor<ElementType, dimensions, Mode,
-                     access::target::global_buffer, isPlaceholder>
-                Accessor) {
+  multi_ptr(
+      accessor<ElementType, dimensions, Mode, access::target::global_buffer,
+               isPlaceholder, PropertyListT>
+          Accessor) {
     m_Pointer = (pointer_t)(Accessor.get_pointer().get());
   }
 
   // Only if Space == local_space
   template <
       int dimensions, access::mode Mode, access::placeholder isPlaceholder,
-      access::address_space _Space = Space,
+      typename PropertyListT, access::address_space _Space = Space,
       typename = typename std::enable_if<
           _Space == Space && Space == access::address_space::local_space>::type>
   multi_ptr(accessor<ElementType, dimensions, Mode, access::target::local,
-                     isPlaceholder>
+                     isPlaceholder, PropertyListT>
                 Accessor)
       : multi_ptr(Accessor.get_pointer()) {}
 
   // Only if Space == constant_space
   template <int dimensions, access::mode Mode,
-            access::placeholder isPlaceholder,
+            access::placeholder isPlaceholder, typename PropertyListT,
             access::address_space _Space = Space,
             typename = typename std::enable_if<
                 _Space == Space &&
                 Space == access::address_space::constant_space>::type>
-  multi_ptr(accessor<ElementType, dimensions, Mode,
-                     access::target::constant_buffer, isPlaceholder>
-                Accessor)
+  multi_ptr(
+      accessor<ElementType, dimensions, Mode, access::target::constant_buffer,
+               isPlaceholder, PropertyListT>
+          Accessor)
       : multi_ptr(Accessor.get_pointer()) {}
 
   // The following constructors are necessary to create multi_ptr<const
@@ -156,7 +159,7 @@ template <typename ElementType, access::address_space Space> class multi_ptr {
   // Only if Space == global_space || global_device_space and element type is
   // const
   template <int dimensions, access::mode Mode,
-            access::placeholder isPlaceholder,
+            access::placeholder isPlaceholder, typename PropertyListT,
             access::address_space _Space = Space, typename ET = ElementType,
             typename = typename std::enable_if<
                 _Space == Space &&
@@ -164,35 +167,39 @@ template <typename ElementType, access::address_space Space> class multi_ptr {
                  Space == access::address_space::global_device_space) &&
                 std::is_const<ET>::value &&
                 std::is_same<ET, ElementType>::value>::type>
-  multi_ptr(accessor<typename std::remove_const<ET>::type, dimensions, Mode,
-                     access::target::global_buffer, isPlaceholder>
-                Accessor)
+  multi_ptr(
+      accessor<typename std::remove_const<ET>::type, dimensions, Mode,
+               access::target::global_buffer, isPlaceholder, PropertyListT>
+          Accessor)
       : multi_ptr(Accessor.get_pointer()) {}
 
   // Only if Space == local_space and element type is const
   template <
       int dimensions, access::mode Mode, access::placeholder isPlaceholder,
-      access::address_space _Space = Space, typename ET = ElementType,
+      typename PropertyListT, access::address_space _Space = Space,
+      typename ET = ElementType,
       typename = typename std::enable_if<
           _Space == Space && Space == access::address_space::local_space &&
           std::is_const<ET>::value &&
           std::is_same<ET, ElementType>::value>::type>
   multi_ptr(accessor<typename std::remove_const<ET>::type, dimensions, Mode,
-                     access::target::local, isPlaceholder>
+                     access::target::local, isPlaceholder, PropertyListT>
                 Accessor)
       : multi_ptr(Accessor.get_pointer()) {}
 
   // Only if Space == constant_space and element type is const
   template <
       int dimensions, access::mode Mode, access::placeholder isPlaceholder,
-      access::address_space _Space = Space, typename ET = ElementType,
+      typename PropertyListT, access::address_space _Space = Space,
+      typename ET = ElementType,
       typename = typename std::enable_if<
           _Space == Space && Space == access::address_space::constant_space &&
           std::is_const<ET>::value &&
           std::is_same<ET, ElementType>::value>::type>
-  multi_ptr(accessor<typename std::remove_const<ET>::type, dimensions, Mode,
-                     access::target::constant_buffer, isPlaceholder>
-                Accessor)
+  multi_ptr(
+      accessor<typename std::remove_const<ET>::type, dimensions, Mode,
+               access::target::constant_buffer, isPlaceholder, PropertyListT>
+          Accessor)
       : multi_ptr(Accessor.get_pointer()) {}
 
   // TODO: This constructor is the temporary solution for the existing problem
@@ -369,37 +376,37 @@ template <access::address_space Space> class multi_ptr<void, Space> {
 
   // Only if Space == global_space || global_device_space
   template <typename ElementType, int dimensions, access::mode Mode,
-            access::address_space _Space = Space,
+            typename PropertyListT, access::address_space _Space = Space,
             typename = typename std::enable_if<
                 _Space == Space &&
                 (Space == access::address_space::global_space ||
                  Space == access::address_space::global_device_space)>::type>
   multi_ptr(
       accessor<ElementType, dimensions, Mode, access::target::global_buffer,
-               access::placeholder::false_t>
+               access::placeholder::false_t, PropertyListT>
           Accessor)
       : multi_ptr(Accessor.get_pointer()) {}
 
   // Only if Space == local_space
   template <
       typename ElementType, int dimensions, access::mode Mode,
-      access::address_space _Space = Space,
+      typename PropertyListT, access::address_space _Space = Space,
       typename = typename std::enable_if<
           _Space == Space && Space == access::address_space::local_space>::type>
   multi_ptr(accessor<ElementType, dimensions, Mode, access::target::local,
-                     access::placeholder::false_t>
+                     access::placeholder::false_t, PropertyListT>
                 Accessor)
       : multi_ptr(Accessor.get_pointer()) {}
 
   // Only if Space == constant_space
   template <typename ElementType, int dimensions, access::mode Mode,
-            access::address_space _Space = Space,
+            typename PropertyListT, access::address_space _Space = Space,
             typename = typename std::enable_if<
                 _Space == Space &&
                 Space == access::address_space::constant_space>::type>
   multi_ptr(
       accessor<ElementType, dimensions, Mode, access::target::constant_buffer,
-               access::placeholder::false_t>
+               access::placeholder::false_t, PropertyListT>
           Accessor)
       : multi_ptr(Accessor.get_pointer()) {}
 
@@ -491,37 +498,37 @@ class multi_ptr<const void, Space> {
 
   // Only if Space == global_space || global_device_space
   template <typename ElementType, int dimensions, access::mode Mode,
-            access::address_space _Space = Space,
+            typename PropertyListT, access::address_space _Space = Space,
             typename = typename std::enable_if<
                 _Space == Space &&
                 (Space == access::address_space::global_space ||
                  Space == access::address_space::global_device_space)>::type>
   multi_ptr(
       accessor<ElementType, dimensions, Mode, access::target::global_buffer,
-               access::placeholder::false_t>
+               access::placeholder::false_t, PropertyListT>
           Accessor)
       : multi_ptr(Accessor.get_pointer()) {}
 
   // Only if Space == local_space
   template <
       typename ElementType, int dimensions, access::mode Mode,
-      access::address_space _Space = Space,
+      typename PropertyListT, access::address_space _Space = Space,
       typename = typename std::enable_if<
           _Space == Space && Space == access::address_space::local_space>::type>
   multi_ptr(accessor<ElementType, dimensions, Mode, access::target::local,
-                     access::placeholder::false_t>
+                     access::placeholder::false_t, PropertyListT>
                 Accessor)
       : multi_ptr(Accessor.get_pointer()) {}
 
   // Only if Space == constant_space
   template <typename ElementType, int dimensions, access::mode Mode,
-            access::address_space _Space = Space,
+            typename PropertyListT, access::address_space _Space = Space,
             typename = typename std::enable_if<
                 _Space == Space &&
                 Space == access::address_space::constant_space>::type>
   multi_ptr(
       accessor<ElementType, dimensions, Mode, access::target::constant_buffer,
-               access::placeholder::false_t>
+               access::placeholder::false_t, PropertyListT>
           Accessor)
       : multi_ptr(Accessor.get_pointer()) {}
 
@@ -550,18 +557,19 @@ class multi_ptr<const void, Space> {
 
 #ifdef __cpp_deduction_guides
 template <int dimensions, access::mode Mode, access::placeholder isPlaceholder,
-          class T>
-multi_ptr(
-    accessor<T, dimensions, Mode, access::target::global_buffer, isPlaceholder>)
+          typename PropertyListT, class T>
+multi_ptr(accessor<T, dimensions, Mode, access::target::global_buffer,
+                   isPlaceholder, PropertyListT>)
     ->multi_ptr<T, access::address_space::global_space>;
 template <int dimensions, access::mode Mode, access::placeholder isPlaceholder,
-          class T>
+          typename PropertyListT, class T>
 multi_ptr(accessor<T, dimensions, Mode, access::target::constant_buffer,
-                   isPlaceholder>)
+                   isPlaceholder, PropertyListT>)
     ->multi_ptr<T, access::address_space::constant_space>;
 template <int dimensions, access::mode Mode, access::placeholder isPlaceholder,
-          class T>
-multi_ptr(accessor<T, dimensions, Mode, access::target::local, isPlaceholder>)
+          typename PropertyListT, class T>
+multi_ptr(accessor<T, dimensions, Mode, access::target::local, isPlaceholder,
+                   PropertyListT>)
     ->multi_ptr<T, access::address_space::local_space>;
 #endif
 
diff --git a/sycl/include/CL/sycl/properties/accessor_properties.hpp b/sycl/include/CL/sycl/properties/accessor_properties.hpp
index 2aea424846ee6..b1272b48dc8ec 100644
--- a/sycl/include/CL/sycl/properties/accessor_properties.hpp
+++ b/sycl/include/CL/sycl/properties/accessor_properties.hpp
@@ -8,8 +8,10 @@
 
 #pragma once
 
+#include "CL/sycl/accessor_property_list.hpp"
 #include <CL/sycl/detail/common.hpp>
 #include <CL/sycl/detail/property_helper.hpp>
+#include <type_traits>
 
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
@@ -33,5 +35,87 @@ constexpr const auto &noinit =
 
 #endif
 
+namespace ext {
+namespace INTEL {
+namespace property {
+template <int A> struct buffer_location {
+  constexpr bool operator==(const buffer_location<A> &) const { return true; }
+  constexpr bool operator!=(const buffer_location<A> &) const { return false; }
+  template <bool B>
+  constexpr bool operator==(const buffer_location<B> &) const {
+    return false;
+  }
+  template <bool B>
+  constexpr bool operator!=(const buffer_location<B> &) const {
+    return true;
+  }
+};
+
+} // namespace property
+} // namespace INTEL
+namespace ONEAPI {
+namespace property {
+template <bool A = true> struct no_offset {
+  constexpr bool operator==(const no_offset<A> &) const {
+    return true;
+  }
+  constexpr bool operator!=(const no_offset<A> &) const {
+    return false;
+  }
+  template <bool B>
+  constexpr bool operator==(const no_offset<B> &) const {
+    return false;
+  }
+  template <bool B>
+  constexpr bool operator!=(const no_offset<B> &) const {
+    return true;
+  }
+};
+template <bool A = true> struct no_alias {
+  constexpr bool operator==(const no_alias<A> &) const {
+    return true;
+  }
+  constexpr bool operator!=(const no_alias<A> &) const {
+    return false;
+  }
+  template <bool B>
+  constexpr bool operator==(const no_alias<B> &) const {
+    return false;
+  }
+  template <bool B>
+  constexpr bool operator!=(const no_alias<B> &) const {
+    return true;
+  }
+
+};
+} // namespace property
+
+#if __cplusplus > 201402L
+
+inline constexpr property::no_offset no_offset;
+inline constexpr property::no_alias no_alias;
+
+#else
+
+namespace {
+constexpr const auto &no_offset =
+    sycl::detail::InlineVariableHelper<property::no_offset<>>::value;
+constexpr const auto &no_alias =
+    sycl::detail::InlineVariableHelper<property::no_alias<>>::value;
+} // namespace
+
+#endif
+
+template <bool B>
+struct is_compile_time_property<ext::ONEAPI::property::no_offset<B>>
+    : std::true_type {};
+template <bool B>
+struct is_compile_time_property<ext::ONEAPI::property::no_alias<B>>
+    : std::true_type {};
+template <int I>
+struct is_compile_time_property<ext::INTEL::property::buffer_location<I>>
+    : std::true_type {};
+} // namespace ONEAPI
+} // namespace ext
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)
diff --git a/sycl/include/CL/sycl/property_list.hpp b/sycl/include/CL/sycl/property_list.hpp
index 5edaa636692f6..61b58bfa07580 100644
--- a/sycl/include/CL/sycl/property_list.hpp
+++ b/sycl/include/CL/sycl/property_list.hpp
@@ -9,12 +9,7 @@
 #pragma once
 
 #include <CL/sycl/detail/common.hpp>
-#include <CL/sycl/detail/property_helper.hpp>
-
-#include <bitset>
-#include <memory>
-#include <type_traits>
-#include <vector>
+#include <CL/sycl/detail/property_list_base.hpp>
 
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
@@ -22,7 +17,7 @@ namespace sycl {
 /// Objects of the property_list class are containers for the SYCL properties
 ///
 /// \ingroup sycl_api
-class property_list {
+class property_list : protected detail::PropertyListBase {
 
   // The structs validate that all objects passed are SYCL properties
   template <typename... Tail> struct AllProperties : std::true_type {};
@@ -36,7 +31,7 @@ class property_list {
 public:
   template <typename... PropsT, typename = typename std::enable_if<
                                     AllProperties<PropsT...>::value>::type>
-  property_list(PropsT... Props) : MDataLessProps(false) {
+  property_list(PropsT... Props) : detail::PropertyListBase(false) {
     ctorHelper(Props...);
   }
 
@@ -51,80 +46,6 @@ class property_list {
   template <typename PropT> bool has_property() const {
     return has_property_helper<PropT>();
   }
-
-private:
-  void ctorHelper() {}
-
-  template <typename... PropsT, class PropT>
-  typename std::enable_if<
-      std::is_base_of<detail::DataLessPropertyBase, PropT>::value>::type
-  ctorHelper(PropT &, PropsT... Props) {
-    const int PropKind = static_cast<int>(PropT::getKind());
-    MDataLessProps[PropKind] = true;
-    ctorHelper(Props...);
-  }
-
-  template <typename... PropsT, class PropT>
-  typename std::enable_if<
-      std::is_base_of<detail::PropertyWithDataBase, PropT>::value>::type
-  ctorHelper(PropT &Prop, PropsT... Props) {
-    MPropsWithData.emplace_back(new PropT(Prop));
-    ctorHelper(Props...);
-  }
-
-  template <typename PropT>
-  typename std::enable_if<
-      std::is_base_of<detail::DataLessPropertyBase, PropT>::value, bool>::type
-  has_property_helper() const {
-    const int PropKind = static_cast<int>(PropT::getKind());
-    if (PropKind >= detail::DataLessPropKind::DataLessPropKindSize)
-      return false;
-    return MDataLessProps[PropKind];
-  }
-
-  template <typename PropT>
-  typename std::enable_if<
-      std::is_base_of<detail::PropertyWithDataBase, PropT>::value, bool>::type
-  has_property_helper() const {
-    const int PropKind = static_cast<int>(PropT::getKind());
-    for (const std::shared_ptr<detail::PropertyWithDataBase> &Prop :
-         MPropsWithData)
-      if (Prop->isSame(PropKind))
-        return true;
-    return false;
-  }
-
-  template <typename PropT>
-  typename std::enable_if<
-      std::is_base_of<detail::DataLessPropertyBase, PropT>::value, PropT>::type
-  get_property_helper() const {
-    // In case of simple property we can just construct it
-    return PropT{};
-  }
-
-  template <typename PropT>
-  typename std::enable_if<
-      std::is_base_of<detail::PropertyWithDataBase, PropT>::value, PropT>::type
-  get_property_helper() const {
-    const int PropKind = static_cast<int>(PropT::getKind());
-    if (PropKind >= detail::PropWithDataKind::PropWithDataKindSize)
-      throw sycl::invalid_object_error("The property is not found",
-                                       PI_INVALID_VALUE);
-
-    for (const std::shared_ptr<detail::PropertyWithDataBase> &Prop :
-         MPropsWithData)
-      if (Prop->isSame(PropKind))
-        return *static_cast<PropT *>(Prop.get());
-
-    throw sycl::invalid_object_error("The property is not found",
-                                     PI_INVALID_VALUE);
-  }
-
-private:
-  // Stores enable/not enabled for simple properties
-  std::bitset<detail::DataLessPropKind::DataLessPropKindSize> MDataLessProps;
-  // Stores shared_ptrs to complex properties
-  std::vector<std::shared_ptr<detail::PropertyWithDataBase>> MPropsWithData;
 };
 
 } // namespace sycl
diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp
index 21c8cec78f072..63d783299a06b 100644
--- a/sycl/include/CL/sycl/types.hpp
+++ b/sycl/include/CL/sycl/types.hpp
@@ -899,9 +899,12 @@ template <typename Type, int NumElements> class vec {
     load(Offset, ConstPtr);
   }
   template <int Dimensions, access::mode Mode,
-            access::placeholder IsPlaceholder, access::target Target>
-  void load(size_t Offset,
-            accessor<DataT, Dimensions, Mode, Target, IsPlaceholder> Acc) {
+            access::placeholder IsPlaceholder, access::target Target,
+            typename PropertyListT>
+  void
+  load(size_t Offset,
+       accessor<DataT, Dimensions, Mode, Target, IsPlaceholder, PropertyListT>
+           Acc) {
     multi_ptr<const DataT, detail::TargetToAS<Target>::AS> MultiPtr(Acc);
     load(Offset, MultiPtr);
   }
@@ -912,9 +915,12 @@ template <typename Type, int NumElements> class vec {
     }
   }
   template <int Dimensions, access::mode Mode,
-            access::placeholder IsPlaceholder, access::target Target>
-  void store(size_t Offset,
-             accessor<DataT, Dimensions, Mode, Target, IsPlaceholder> Acc) {
+            access::placeholder IsPlaceholder, access::target Target,
+            typename PropertyListT>
+  void
+  store(size_t Offset,
+        accessor<DataT, Dimensions, Mode, Target, IsPlaceholder, PropertyListT>
+            Acc) {
     multi_ptr<DataT, detail::TargetToAS<Target>::AS> MultiPtr(Acc);
     store(Offset, MultiPtr);
   }
diff --git a/sycl/test/abi/layout_accessors.cpp b/sycl/test/abi/layout_accessors.cpp
index bde079ac8f614..8fcb3ac47f8b8 100644
--- a/sycl/test/abi/layout_accessors.cpp
+++ b/sycl/test/abi/layout_accessors.cpp
@@ -13,7 +13,7 @@ void hostAcc(accessor<int, 1, access::mode::read, access::target::host_buffer> A
   (void)Acc.get_size();
 }
 
-// CHECK: 0 | class cl::sycl::accessor<int, 1, cl::sycl::access::mode::read, cl::sycl::access::target::host_buffer, cl::sycl::access::placeholder::false_t>
+// CHECK: 0 | class cl::sycl::accessor<int, 1, cl::sycl::access::mode::read, cl::sycl::access::target::host_buffer, cl::sycl::access::placeholder::false_t, class cl::sycl::property_list>
 // CHECK-NEXT:  0 |   class cl::sycl::detail::AccessorBaseHost (base)
 // CHECK-NEXT:  0 |     class std::shared_ptr<class cl::sycl::detail::AccessorImplHost> impl
 // CHECK-NEXT:  0 |       class std::__shared_ptr<class cl::sycl::detail::AccessorImplHost, __gnu_cxx::_S_atomic> (base)
@@ -34,7 +34,7 @@ void hostAcc(accessor<int, 1, access::mode::read, access::target::global_buffer>
   (void)Acc.get_size();
 }
 
-// CHECK: 0 | class cl::sycl::accessor<int, 1, cl::sycl::access::mode::read, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>
+// CHECK: 0 | class cl::sycl::accessor<int, 1, cl::sycl::access::mode::read, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t, class cl::sycl::property_list>
 // CHECK-NEXT: 0 |   class cl::sycl::detail::AccessorBaseHost (base)
 // CHECK-NEXT: 0 |     class std::shared_ptr<class cl::sycl::detail::AccessorImplHost> impl
 // CHECK-NEXT: 0 |       class std::__shared_ptr<class cl::sycl::detail::AccessorImplHost, __gnu_cxx::_S_atomic> (base)
@@ -55,7 +55,7 @@ void hostAcc(accessor<int, 1, access::mode::read_write, access::target::local> A
   (void)Acc.get_size();
 }
 
-// CHECK: 0 | class cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::local, cl::sycl::access::placeholder::false_t>
+// CHECK: 0 | class cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::local, cl::sycl::access::placeholder::false_t, class cl::sycl::property_list>
 // CHECK-NEXT: 0 |   class cl::sycl::detail::LocalAccessorBaseHost (base)
 // CHECK-NEXT: 0 |     class std::shared_ptr<class cl::sycl::detail::LocalAccessorImplHost> impl
 // CHECK-NEXT: 0 |       class std::__shared_ptr<class cl::sycl::detail::LocalAccessorImplHost, __gnu_cxx::_S_atomic> (base)
@@ -75,7 +75,7 @@ void hostAcc(accessor<int, 1, access::mode::read_write, access::target::local> A
 void hostAcc(accessor<int4, 1, access::mode::read_write, access::target::host_image> Acc) {
   (void)Acc.get_count();
 }
-// CHECK: 0 | class cl::sycl::accessor<class cl::sycl::vec<int, 4>, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::host_image, cl::sycl::access::placeholder::false_t>
+// CHECK: 0 | class cl::sycl::accessor<class cl::sycl::vec<int, 4>, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::host_image, cl::sycl::access::placeholder::false_t, class cl::sycl::property_list>
 // CHECK-NEXT: 0 |   class cl::sycl::detail::image_accessor<class cl::sycl::vec<int, 4>, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::host_image, cl::sycl::access::placeholder::false_t> (base)
 // CHECK-NEXT: 0 |     class cl::sycl::detail::AccessorBaseHost (base)
 // CHECK-NEXT: 0 |       class std::shared_ptr<class cl::sycl::detail::AccessorImplHost> impl
@@ -97,7 +97,7 @@ void hostAcc(accessor<int4, 1, access::mode::read_write, access::target::host_im
 void hostAcc(accessor<int4, 1, access::mode::read, access::target::image> Acc) {
   (void)Acc.get_count();
 }
-// CHECK: 0 | class cl::sycl::accessor<class cl::sycl::vec<int, 4>, 1, cl::sycl::access::mode::read, cl::sycl::access::target::image, cl::sycl::access::placeholder::false_t>
+// CHECK: 0 | class cl::sycl::accessor<class cl::sycl::vec<int, 4>, 1, cl::sycl::access::mode::read, cl::sycl::access::target::image, cl::sycl::access::placeholder::false_t, class cl::sycl::property_list>
 // CHECK-NEXT: 0 |   class cl::sycl::detail::image_accessor<class cl::sycl::vec<int, 4>, 1, cl::sycl::access::mode::read, cl::sycl::access::target::image, cl::sycl::access::placeholder::false_t> (base)
 // CHECK-NEXT: 0 |     class cl::sycl::detail::AccessorBaseHost (base)
 // CHECK-NEXT: 0 |       class std::shared_ptr<class cl::sycl::detail::AccessorImplHost> impl
diff --git a/sycl/test/basic_tests/accessor/accessor_property_list.cpp b/sycl/test/basic_tests/accessor/accessor_property_list.cpp
new file mode 100644
index 0000000000000..51a980342af4a
--- /dev/null
+++ b/sycl/test/basic_tests/accessor/accessor_property_list.cpp
@@ -0,0 +1,107 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: env SYCL_DEVICE_TYPE=HOST %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+#include <CL/sycl.hpp>
+
+#include <type_traits>
+
+using namespace sycl::ext::ONEAPI;
+
+int main() {
+  {
+    // Create empty property list
+    accessor_property_list PL;
+  }
+
+  {
+    // Single CT property
+    accessor_property_list PL{no_alias};
+    static_assert(PL.has_property<property::no_alias>(), "Property not found");
+    static_assert(!PL.has_property<property::no_offset>(), "Property is found");
+  }
+
+  {
+    // Single RT property
+    accessor_property_list PL{sycl::noinit};
+    assert(PL.has_property<sycl::property::noinit>() && "Property not found");
+  }
+
+  {
+    // Compile time and runtime properties
+    accessor_property_list PL{sycl::noinit, no_alias};
+    assert(PL.has_property<property::no_alias>() && "Property not found");
+    assert(PL.has_property<sycl::property::noinit>() && "Property not found");
+  }
+
+  {
+    // Multiple CT properties
+    accessor_property_list PL{no_alias, no_offset};
+    static_assert(PL.has_property<property::no_alias>(), "Property not found");
+    static_assert(PL.has_property<property::no_offset>(), "Property not found");
+
+    static_assert(PL.get_property<property::no_alias>() == no_alias,
+                  "Properties are not equal");
+    property::no_alias<false> NAfalse;
+    static_assert(PL.get_property<property::no_alias>() != NAfalse,
+                  "Properties are equal");
+  }
+
+  {
+    int data[1] = {0};
+
+    sycl::buffer<int, 1> buf_data(data, sycl::range<1>(1),
+                                  {sycl::property::buffer::use_host_ptr()});
+    sycl::queue Queue;
+
+    Queue.submit([&](sycl::handler &cgh) {
+      accessor_property_list PL{no_alias};
+      sycl::accessor acc_1(buf_data, cgh, PL);
+      sycl::accessor acc_2(buf_data, cgh, sycl::range<1>(0), PL);
+      sycl::accessor acc_3(buf_data, cgh, sycl::range<1>(0), sycl::id<1>(1),
+                           PL);
+      sycl::accessor acc_4(buf_data, cgh, sycl::read_only, PL);
+      sycl::accessor acc_5(buf_data, cgh, sycl::range<1>(0), sycl::read_only,
+                           PL);
+      sycl::accessor acc_6(buf_data, cgh, sycl::range<1>(0), sycl::id<1>(0),
+                           sycl::read_only, PL);
+      sycl::accessor acc_7(buf_data, cgh, sycl::write_only, PL);
+      sycl::accessor acc_8(buf_data, cgh, sycl::range<1>(0), sycl::write_only,
+                           PL);
+      sycl::accessor acc_9(buf_data, cgh, sycl::range<1>(0), sycl::id<1>(0),
+                           sycl::write_only, PL);
+      cgh.single_task<class NullKernel>([]() {});
+    });
+    accessor_property_list PL{no_alias};
+    sycl::accessor acc_1(buf_data, PL);
+    sycl::accessor acc_2(buf_data, sycl::range<1>(0), PL);
+    sycl::accessor acc_3(buf_data, sycl::range<1>(0), sycl::id<1>(0), PL);
+    sycl::accessor acc_4(buf_data, sycl::read_only, PL);
+    sycl::accessor acc_5(buf_data, sycl::range<1>(0), sycl::read_only, PL);
+    sycl::accessor acc_6(buf_data, sycl::range<1>(0), sycl::id<1>(0),
+                         sycl::read_only, PL);
+    sycl::accessor acc_7(buf_data, sycl::write_only, PL);
+    sycl::accessor acc_8(buf_data, sycl::range<1>(0), sycl::write_only, PL);
+    sycl::accessor acc_9(buf_data, sycl::range<1>(0), sycl::id<1>(0),
+                         sycl::write_only, PL);
+  }
+
+  {
+    int data[1] = {0};
+
+    sycl::buffer<int, 1> buf_data(data, sycl::range<1>(1),
+                                  {sycl::property::buffer::use_host_ptr()});
+
+    accessor_property_list PL{sycl::noinit, no_alias};
+    sycl::accessor acc_1(buf_data, PL);
+    sycl::accessor<int, 1, sycl::access::mode::read_write,
+                   sycl::access::target::global_buffer,
+                   sycl::access::placeholder::true_t,
+                   accessor_property_list<property::no_alias<>>>
+        acc_2(acc_1);
+  }
+
+  return 0;
+}
\ No newline at end of file

From 415c939c09f23e10bd969ed98a5b66806063f300 Mon Sep 17 00:00:00 2001
From: Alexander Batashev <alexander.batashev@intel.com>
Date: Wed, 9 Sep 2020 10:35:46 +0300
Subject: [PATCH 02/19] Minor fix

---
 .../CL/sycl/{ => ONEAPI}/accessor_property_list.hpp        | 0
 sycl/include/CL/sycl/property_list.hpp                     | 7 +++++++
 sycl/include/CL/sycl/property_list_conversion.hpp          | 0
 3 files changed, 7 insertions(+)
 rename sycl/include/CL/sycl/{ => ONEAPI}/accessor_property_list.hpp (100%)
 create mode 100644 sycl/include/CL/sycl/property_list_conversion.hpp

diff --git a/sycl/include/CL/sycl/accessor_property_list.hpp b/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
similarity index 100%
rename from sycl/include/CL/sycl/accessor_property_list.hpp
rename to sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
diff --git a/sycl/include/CL/sycl/property_list.hpp b/sycl/include/CL/sycl/property_list.hpp
index 61b58bfa07580..abd066b127d0e 100644
--- a/sycl/include/CL/sycl/property_list.hpp
+++ b/sycl/include/CL/sycl/property_list.hpp
@@ -13,6 +13,10 @@
 
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
+namespace ONEAPI {
+template <typename... PropsT>
+class accessor_property_list;
+}
 
 /// Objects of the property_list class are containers for the SYCL properties
 ///
@@ -46,6 +50,9 @@ class property_list : protected detail::PropertyListBase {
   template <typename PropT> bool has_property() const {
     return has_property_helper<PropT>();
   }
+
+  template <typename... PropsT>
+  operator ONEAPI::accessor_property_list<PropsT...>();
 };
 
 } // namespace sycl
diff --git a/sycl/include/CL/sycl/property_list_conversion.hpp b/sycl/include/CL/sycl/property_list_conversion.hpp
new file mode 100644
index 0000000000000..e69de29bb2d1d

From 41365d75ece09c75fa69b5e3d37cc9ad6338e694 Mon Sep 17 00:00:00 2001
From: Alexander Batashev <alexander.batashev@intel.com>
Date: Thu, 10 Sep 2020 14:07:41 +0300
Subject: [PATCH 03/19] Update for latest revision

---
 .../CL/sycl/ONEAPI/accessor_property_list.hpp |  94 ++++++---
 sycl/include/CL/sycl/accessor.hpp             | 192 +++++++++---------
 sycl/include/CL/sycl/buffer.hpp               |  18 +-
 .../CL/sycl/detail/property_list_base.hpp     |   3 +-
 sycl/include/CL/sycl/image.hpp                |  10 +-
 .../sycl/properties/accessor_properties.hpp   | 123 ++++++-----
 sycl/include/CL/sycl/property_list.hpp        |   6 +-
 .../CL/sycl/property_list_conversion.hpp      |  23 +++
 .../accessor/accessor_property_list.cpp       |   9 +-
 9 files changed, 274 insertions(+), 204 deletions(-)

diff --git a/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp b/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
index 84b366ed7a2c0..0adfd51f2051e 100644
--- a/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
+++ b/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
@@ -17,12 +17,13 @@
 
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
-namespace detail {} // namespace detail
-namespace ext {
+namespace detail {
+// This helper template must be specialized for nestend instance template
+// of each compile-time-constant property.
+template <typename T> struct IsCxPropertyInstance : std::false_type {};
+} // namespace detail
 namespace ONEAPI {
 
-template <typename T> struct is_property : std::false_type {};
-
 template <typename T> struct is_compile_time_property : std::false_type {};
 
 /// Objects of the accessor_property_list class are containers for the SYCL
@@ -46,13 +47,16 @@ class accessor_property_list : protected detail::PropertyListBase {
   template <template <auto...> class T, auto...T1, auto...T2>
   struct AreSameTemplate<T<T1...>, T<T2...>> : std::true_type {};
 #endif
+  // This template helps to identify if PropListT parameter pack contains
+  // property of PropT type, where PropT is a nested instance template of
+  // compile-time-constant property.
   template <typename PropT, typename... PropListT>
   struct ContainsProperty;
   template <typename PropT> struct ContainsProperty<PropT> : std::false_type {};
   template <typename PropT, typename Head, typename... Tail>
   struct ContainsProperty<PropT, Head, Tail...>
       : std::conditional<AreSameTemplate<PropT, Head>::value ||
-                             !is_compile_time_property<PropT>::value,
+                             !detail::IsCxPropertyInstance<PropT>::value,
                          std::true_type,
                          ContainsProperty<PropT, Tail...>>::type {};
 
@@ -72,6 +76,25 @@ class accessor_property_list : protected detail::PropertyListBase {
     using Rest = void;
   };
 
+#if __cplusplus >= 201703L
+  // This template serves the same purpose as ContainsProperty, but operates on
+  // template template arguments.
+  template <typename ContainerT, template <auto...> typename PropT,
+            auto... Args>
+  struct ContainsPropertyInstance
+      : std::conditional_t<
+            AreSameTemplate<PropT<Args...>, typename ContainerT::Head>::value ||
+                !detail::IsCxPropertyInstance<typename ContainerT::Head>::value,
+            std::true_type,
+            ContainsPropertyInstance<typename ContainerT::Rest, PropT,
+                                     Args...>> {};
+
+  template <template <auto...> typename PropT, auto... Args>
+  struct ContainsPropertyInstance<void, PropT, Args...> : std::false_type {};
+#endif
+
+  // This template checks if to lists of properties contain the same set of
+  // compile-time-constant properties in any order.
   template <typename ContainerT, typename... OtherProps>
   struct ContainsSameProperties
       : std::conditional<
@@ -82,6 +105,8 @@ class accessor_property_list : protected detail::PropertyListBase {
   struct ContainsSameProperties<void, OtherProps...> : std::true_type {};
 
 #if __cplusplus >= 201703L
+  // This template helps to extract exact property instance type based on
+  // template template argument.
   template <typename ContainerT, template <auto...> class PropT, auto... Args>
   struct GetCxPropertyHelper {
     using type = typename std::conditional_t<
@@ -104,9 +129,14 @@ class accessor_property_list : protected detail::PropertyListBase {
       : std::conditional<
             std::is_base_of<detail::DataLessPropertyBase, T>::value ||
                 std::is_base_of<detail::PropertyWithDataBase, T>::value ||
-                is_compile_time_property<T>::value,
+                detail::IsCxPropertyInstance<T>::value,
             AllProperties<Tail...>, std::false_type>::type {};
 
+  accessor_property_list(
+      std::bitset<detail::DataLessPropKind::DataLessPropKindSize> DataLessProps,
+      std::vector<std::shared_ptr<detail::PropertyWithDataBase>> PropsWithData)
+      : detail::PropertyListBase(DataLessProps, PropsWithData) {}
+
 public:
   template <
       typename = typename std::enable_if<AllProperties<PropsT...>::value>::type>
@@ -114,6 +144,9 @@ class accessor_property_list : protected detail::PropertyListBase {
     ctorHelper(Props...);
   }
 
+  accessor_property_list(const sycl::property_list &Props)
+      : detail::PropertyListBase(Props.MDataLessProps, Props.MPropsWithData) {}
+
   template <typename... OtherProps,
             typename = typename std::enable_if<
                 ContainsSameProperties<PropertyContainer<PropsT...>,
@@ -124,8 +157,8 @@ class accessor_property_list : protected detail::PropertyListBase {
       : detail::PropertyListBase(OtherList.MDataLessProps,
                                  OtherList.MPropsWithData) {}
 
-  template <typename PropT,
-            typename = std::enable_if<!is_compile_time_property<PropT>::value>>
+  template <typename PropT, typename = typename std::enable_if<
+                                !is_compile_time_property<PropT>::value>::type>
   PropT get_property() const {
     if (!has_property<PropT>())
       throw sycl::invalid_object_error("The property is not found",
@@ -141,36 +174,51 @@ class accessor_property_list : protected detail::PropertyListBase {
   }
 
 #if __cplusplus >= 201703L
-  template <template <auto...> class PropT, auto... Args>
+  template <typename T>
   static constexpr
-      typename std::enable_if<is_compile_time_property<PropT<Args...>>::value,
-                              bool>::type
+      typename std::enable_if_t<is_compile_time_property<T>::value, bool>
       has_property() {
-    return ContainsProperty<PropT<Args...>, PropsT...>::value;
+    return ContainsPropertyInstance<PropertyContainer<PropsT...>,
+                                    T::template instance>::value;
+  }
+
+  template <typename T, typename = typename std::enable_if_t<
+                            is_compile_time_property<T>::value>>
+  static constexpr auto get_property() {
+    return typename GetCxPropertyHelper<PropertyContainer<PropsT...>,
+                                        T::template instance>::type{};
+  }
+  /*
+  template <template <auto...> class PropT, auto... Args>
+  static constexpr
+      typename std::enable_if<detail::IsCxPropertyInstance<typename
+  PropT::instance<Args...>>::value, bool>::type has_property() { return
+  ContainsProperty<typename PropT::instance<Args...>, PropsT...>::value;
   }
   template <template <auto...> class PropT, auto... Args,
             typename = typename std::enable_if<
-                is_compile_time_property<PropT<Args...>>::value>::type>
-  static constexpr auto get_property() {
-    return typename GetCxPropertyHelper<PropertyContainer<PropsT...>, PropT,
-                                        Args...>::type{};
+                detail::IsCxPropertyInstance<typename
+  PropT::instance<Args...>>::value>::type> static constexpr auto get_property()
+  { return typename GetCxPropertyHelper<PropertyContainer<PropsT...>, typename
+  PropT::instance, Args...>::type{};
   }
+  */
 #endif
 
+  template <typename... OtherPropsT>
+  static constexpr bool areSameCxProperties() {
+    return ContainsSameProperties<PropertyContainer<OtherPropsT...>,
+                                  PropsT...>::value;
+  }
+
 private:
   template <typename... OtherProps> friend class accessor_property_list;
   template <typename, int, sycl::access::mode, sycl::access::target,
             sycl::access::placeholder, typename>
   friend class accessor;
 
-public:
-  template <typename... OtherPropsT>
-  static constexpr bool areSameCxProperties() {
-    return ContainsSameProperties<PropertyContainer<OtherPropsT...>,
-                                  PropsT...>::value;
-  }
+  friend class sycl::property_list;
 };
 } // namespace ONEAPI
-} // namespace ext
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)
diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp
index 8c101a48aea24..feb0e1966549e 100755
--- a/sycl/include/CL/sycl/accessor.hpp
+++ b/sycl/include/CL/sycl/accessor.hpp
@@ -9,7 +9,7 @@
 #pragma once
 
 #include <CL/__spirv/spirv_types.hpp>
-#include <CL/sycl/accessor_property_list.hpp>
+#include <CL/sycl/ONEAPI/accessor_property_list.hpp>
 #include <CL/sycl/atomic.hpp>
 #include <CL/sycl/buffer.hpp>
 #include <CL/sycl/detail/accessor_impl.hpp>
@@ -25,6 +25,7 @@
 #include <CL/sycl/pointers.hpp>
 #include <CL/sycl/properties/accessor_properties.hpp>
 #include <CL/sycl/property_list.hpp>
+#include <CL/sycl/property_list_conversion.hpp>
 #include <CL/sycl/sampler.hpp>
 
 #include <type_traits>
@@ -218,7 +219,7 @@ 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,
-          typename PropertyListT = property_list>
+          typename PropertyListT = ONEAPI::accessor_property_list<>>
 class accessor;
 
 namespace detail {
@@ -226,17 +227,22 @@ template <typename T>
 using IsPropertyListT = typename std::is_base_of<PropertyListBase, T>;
 
 template <typename T>
-using IsRunTimePropertyListT = typename std::is_same<property_list, T>;
+using IsRunTimePropertyListT =
+    typename std::is_same<ONEAPI::accessor_property_list<>, T>;
 
-template <typename T> struct IsOneapiPropertyListT {
+template <typename T> struct IsCxPropertyList {
   constexpr static bool value = false;
 };
 
 template <typename... Props>
-struct IsOneapiPropertyListT<ext::ONEAPI::accessor_property_list<Props...>> {
+struct IsCxPropertyList<ONEAPI::accessor_property_list<Props...>> {
   constexpr static bool value = true;
 };
 
+template <> struct IsCxPropertyList<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.
@@ -770,7 +776,8 @@ class accessor :
                  AccessMode == access::mode::read),
                 "Access mode can be only read for constant buffers");
 
-  static_assert(detail::IsPropertyListT<PropertyListT>::value, "PropertyListT must be property_list or accessor_property_list");
+  static_assert(detail::IsPropertyListT<PropertyListT>::value,
+                "PropertyListT must be accessor_property_list");
 
   using AccessorCommonT = detail::accessor_common<DataT, Dimensions, AccessMode,
                                                   AccessTarget, IsPlaceholder>;
@@ -972,13 +979,13 @@ class accessor :
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
             typename... PropTypes,
             typename detail::enable_if_t<
-                detail::IsOneapiPropertyListT<PropertyListT>::value &&
+                detail::IsCxPropertyList<PropertyListT>::value &&
                 std::is_same<T, DataT>::value && Dims == 0 &&
                 ((!IsPlaceH && IsHostBuf) ||
                  (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))> * = nullptr>
-  accessor(buffer<T, 1, AllocatorT> &BufferRef,
-           const ext::ONEAPI::accessor_property_list<PropTypes...>
-               &PropertyList = {})
+  accessor(
+      buffer<T, 1, AllocatorT> &BufferRef,
+      const ONEAPI::accessor_property_list<PropTypes...> &PropertyList = {})
 #ifdef __SYCL_DEVICE_ONLY__
       : impl(id<AdjustedDim>(), range<1>{1}, BufferRef.get_range()) {
     (void)PropertyList;
@@ -1022,12 +1029,12 @@ class accessor :
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
             typename... PropTypes,
             typename = typename detail::enable_if_t<
-                detail::IsOneapiPropertyListT<PropertyListT>::value &&
+                detail::IsCxPropertyList<PropertyListT>::value &&
                 std::is_same<T, DataT>::value && (Dims == 0) &&
                 (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
-  accessor(buffer<T, 1, AllocatorT> &BufferRef, handler &CommandGroupHandler,
-           const ext::ONEAPI::accessor_property_list<PropTypes...>
-               &PropertyList = {})
+  accessor(
+      buffer<T, 1, AllocatorT> &BufferRef, handler &CommandGroupHandler,
+      const ONEAPI::accessor_property_list<PropTypes...> &PropertyList = {})
 #ifdef __SYCL_DEVICE_ONLY__
       : impl(id<AdjustedDim>(), range<1>{1}, BufferRef.get_range()) {
     (void)CommandGroupHandler;
@@ -1074,13 +1081,13 @@ class accessor :
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
             typename... PropTypes,
             typename = detail::enable_if_t<
-                detail::IsOneapiPropertyListT<PropertyListT>::value &&
+                detail::IsCxPropertyList<PropertyListT>::value &&
                 IsSameAsBuffer<T, Dims>() &&
                 ((!IsPlaceH && IsHostBuf) ||
                  (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>>
-  accessor(buffer<T, Dims, AllocatorT> &BufferRef,
-           const ext::ONEAPI::accessor_property_list<PropTypes...>
-               &PropertyList = {})
+  accessor(
+      buffer<T, Dims, AllocatorT> &BufferRef,
+      const ONEAPI::accessor_property_list<PropTypes...> &PropertyList = {})
 #ifdef __SYCL_DEVICE_ONLY__
       : impl(id<Dimensions>(), BufferRef.get_range(), BufferRef.get_range()) {
     (void)PropertyList;
@@ -1114,12 +1121,12 @@ class accessor :
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
             typename TagT, typename... PropTypes,
             typename = detail::enable_if_t<
-                detail::IsOneapiPropertyListT<PropertyListT>::value &&
+                detail::IsCxPropertyList<PropertyListT>::value &&
                 IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && IsPlaceH &&
                 (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
-  accessor(buffer<T, Dims, AllocatorT> &BufferRef, TagT,
-           const ext::ONEAPI::accessor_property_list<PropTypes...>
-               &PropertyList = {})
+  accessor(
+      buffer<T, Dims, AllocatorT> &BufferRef, TagT,
+      const ONEAPI::accessor_property_list<PropTypes...> &PropertyList = {})
       : accessor(BufferRef, PropertyList) {}
 #endif
 
@@ -1151,12 +1158,12 @@ class accessor :
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
             typename... PropTypes,
             typename = detail::enable_if_t<
-                detail::IsOneapiPropertyListT<PropertyListT>::value &&
+                detail::IsCxPropertyList<PropertyListT>::value &&
                 IsSameAsBuffer<T, Dims>() &&
                 (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
-  accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
-           const ext::ONEAPI::accessor_property_list<PropTypes...>
-               &PropertyList = {})
+  accessor(
+      buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
+      const ONEAPI::accessor_property_list<PropTypes...> &PropertyList = {})
 #ifdef __SYCL_DEVICE_ONLY__
       : impl(id<AdjustedDim>(), BufferRef.get_range(), BufferRef.get_range()) {
     (void)CommandGroupHandler;
@@ -1190,13 +1197,13 @@ class accessor :
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
             typename TagT, typename... PropTypes,
             typename = detail::enable_if_t<
-                detail::IsOneapiPropertyListT<PropertyListT>::value &&
+                detail::IsCxPropertyList<PropertyListT>::value &&
                 IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && !IsPlaceH &&
                 (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
-  accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
-           TagT,
-           const ext::ONEAPI::accessor_property_list<PropTypes...>
-               &PropertyList = {})
+  accessor(
+      buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
+      TagT,
+      const ONEAPI::accessor_property_list<PropTypes...> &PropertyList = {})
       : accessor(BufferRef, CommandGroupHandler, PropertyList) {}
 
 #endif
@@ -1215,14 +1222,13 @@ class accessor :
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
             typename... PropTypes,
             typename = detail::enable_if_t<
-                detail::IsOneapiPropertyListT<PropertyListT>::value &&
+                detail::IsCxPropertyList<PropertyListT>::value &&
                 IsSameAsBuffer<T, Dims>() &&
                 ((!IsPlaceH && IsHostBuf) ||
                  (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>>
-  accessor(buffer<T, Dims, AllocatorT> &BufferRef,
-           range<Dimensions> AccessRange,
-           const ext::ONEAPI::accessor_property_list<PropTypes...>
-               &PropertyList = {})
+  accessor(
+      buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
+      const ONEAPI::accessor_property_list<PropTypes...> &PropertyList = {})
       : accessor(BufferRef, AccessRange, {}, PropertyList) {}
 
 #if __cplusplus > 201402L
@@ -1241,13 +1247,13 @@ class accessor :
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
             typename TagT, typename... PropTypes,
             typename = detail::enable_if_t<
-                detail::IsOneapiPropertyListT<PropertyListT>::value &&
+                detail::IsCxPropertyList<PropertyListT>::value &&
                 IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && IsPlaceH &&
                 (IsGlobalBuf || IsConstantBuf)>>
-  accessor(buffer<T, Dims, AllocatorT> &BufferRef,
-           range<Dimensions> AccessRange, TagT,
-           const ext::ONEAPI::accessor_property_list<PropTypes...>
-               &PropertyList = {})
+  accessor(
+      buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
+      TagT,
+      const ONEAPI::accessor_property_list<PropTypes...> &PropertyList = {})
       : accessor(BufferRef, AccessRange, {}, PropertyList) {}
 #endif
 
@@ -1265,13 +1271,13 @@ class accessor :
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
             typename... PropTypes,
             typename = detail::enable_if_t<
-                detail::IsOneapiPropertyListT<PropertyListT>::value &&
+                detail::IsCxPropertyList<PropertyListT>::value &&
                 IsSameAsBuffer<T, Dims>() &&
                 (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
-  accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
-           range<Dimensions> AccessRange,
-           const ext::ONEAPI::accessor_property_list<PropTypes...>
-               &PropertyList = {})
+  accessor(
+      buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
+      range<Dimensions> AccessRange,
+      const ONEAPI::accessor_property_list<PropTypes...> &PropertyList = {})
       : accessor(BufferRef, CommandGroupHandler, AccessRange, {},
                  PropertyList) {}
 
@@ -1292,13 +1298,13 @@ class accessor :
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
             typename TagT, typename... PropTypes,
             typename = detail::enable_if_t<
-                detail::IsOneapiPropertyListT<PropertyListT>::value &&
+                detail::IsCxPropertyList<PropertyListT>::value &&
                 IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && !IsPlaceH &&
                 (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
-  accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
-           range<Dimensions> AccessRange, TagT,
-           const ext::ONEAPI::accessor_property_list<PropTypes...>
-               &PropertyList = {})
+  accessor(
+      buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
+      range<Dimensions> AccessRange, TagT,
+      const ONEAPI::accessor_property_list<PropTypes...> &PropertyList = {})
       : accessor(BufferRef, CommandGroupHandler, AccessRange, {},
                  PropertyList) {}
 #endif
@@ -1333,14 +1339,14 @@ class accessor :
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
             typename... PropTypes,
             typename = detail::enable_if_t<
-                detail::IsOneapiPropertyListT<PropertyListT>::value &&
+                detail::IsCxPropertyList<PropertyListT>::value &&
                 IsSameAsBuffer<T, Dims>() &&
                 ((!IsPlaceH && IsHostBuf) ||
                  (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>>
-  accessor(buffer<T, Dims, AllocatorT> &BufferRef,
-           range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
-           const ext::ONEAPI::accessor_property_list<PropTypes...>
-               &PropertyList = {})
+  accessor(
+      buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
+      id<Dimensions> AccessOffset,
+      const ONEAPI::accessor_property_list<PropTypes...> &PropertyList = {})
 #ifdef __SYCL_DEVICE_ONLY__
       : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
     (void)PropertyList;
@@ -1375,13 +1381,13 @@ class accessor :
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
             typename TagT, typename... PropTypes,
             typename = detail::enable_if_t<
-                detail::IsOneapiPropertyListT<PropertyListT>::value &&
+                detail::IsCxPropertyList<PropertyListT>::value &&
                 IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && IsPlaceH &&
                 (IsGlobalBuf || IsConstantBuf)>>
-  accessor(buffer<T, Dims, AllocatorT> &BufferRef,
-           range<Dimensions> AccessRange, id<Dimensions> AccessOffset, TagT,
-           const ext::ONEAPI::accessor_property_list<PropTypes...>
-               &PropertyList = {})
+  accessor(
+      buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
+      id<Dimensions> AccessOffset, TagT,
+      const ONEAPI::accessor_property_list<PropTypes...> &PropertyList = {})
       : accessor(BufferRef, AccessRange, AccessOffset, PropertyList) {}
 #endif
 
@@ -1414,13 +1420,13 @@ class accessor :
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
             typename... PropTypes,
             typename = detail::enable_if_t<
-                detail::IsOneapiPropertyListT<PropertyListT>::value &&
+                detail::IsCxPropertyList<PropertyListT>::value &&
                 IsSameAsBuffer<T, Dims>() &&
                 (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
-  accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
-           range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
-           const ext::ONEAPI::accessor_property_list<PropTypes...>
-               &PropertyList = {})
+  accessor(
+      buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
+      range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
+      const ONEAPI::accessor_property_list<PropTypes...> &PropertyList = {})
 #ifdef __SYCL_DEVICE_ONLY__
       : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
     (void)CommandGroupHandler;
@@ -1456,13 +1462,13 @@ class accessor :
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
             typename TagT, typename... PropTypes,
             typename = detail::enable_if_t<
-                detail::IsOneapiPropertyListT<PropertyListT>::value &&
+                detail::IsCxPropertyList<PropertyListT>::value &&
                 IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && !IsPlaceH &&
                 (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
-  accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
-           range<Dimensions> AccessRange, id<Dimensions> AccessOffset, TagT,
-           const ext::ONEAPI::accessor_property_list<PropTypes...>
-               &PropertyList = {})
+  accessor(
+      buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
+      range<Dimensions> AccessRange, id<Dimensions> AccessOffset, TagT,
+      const ONEAPI::accessor_property_list<PropTypes...> &PropertyList = {})
       : accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
                  PropertyList) {}
 #endif
@@ -1470,14 +1476,14 @@ class accessor :
   template <typename... NewPropsT>
   accessor(
       const accessor<DataT, Dimensions, AccessMode, AccessTarget, IsPlaceholder,
-                     ext::ONEAPI::accessor_property_list<NewPropsT...>> &Other)
+                     ONEAPI::accessor_property_list<NewPropsT...>> &Other)
 #ifdef __SYCL_DEVICE_ONLY__
       : impl(Other.impl)
 #else
       : detail::AccessorBaseHost(Other)
 #endif
   {
-    static_assert(detail::IsOneapiPropertyListT<PropertyListT>::value,
+    static_assert(detail::IsCxPropertyList<PropertyListT>::value,
                   "Conversion is only available for accessor_property_list");
     static_assert(PropertyListT::template areSameCxProperties<NewPropsT...>(),
                   "Compile-time-constant properties must be the same");
@@ -1607,10 +1613,10 @@ accessor(buffer<DataT, Dimensions, AllocatorT>)
 template <typename DataT, int Dimensions, typename AllocatorT,
           typename... PropsT>
 accessor(buffer<DataT, Dimensions, AllocatorT>,
-         const ext::ONEAPI::accessor_property_list<PropsT...> &)
+         const ONEAPI::accessor_property_list<PropsT...> &)
     ->accessor<DataT, Dimensions, access::mode::read_write,
                target::global_buffer, access::placeholder::true_t,
-               ext::ONEAPI::accessor_property_list<PropsT...>>;
+               ONEAPI::accessor_property_list<PropsT...>>;
 
 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1>
 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1)
@@ -1621,11 +1627,11 @@ accessor(buffer<DataT, Dimensions, AllocatorT>, Type1)
 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
           typename... PropsT>
 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1,
-         const ext::ONEAPI::accessor_property_list<PropsT...> &)
+         const ONEAPI::accessor_property_list<PropsT...> &)
     ->accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type1>(),
                detail::deduceAccessTarget<Type1, Type1>(target::global_buffer),
                access::placeholder::true_t,
-               ext::ONEAPI::accessor_property_list<PropsT...>>;
+               ONEAPI::accessor_property_list<PropsT...>>;
 
 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
           typename Type2>
@@ -1637,11 +1643,11 @@ accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2)
 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
           typename Type2, typename... PropsT>
 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2,
-         const ext::ONEAPI::accessor_property_list<PropsT...> &)
+         const ONEAPI::accessor_property_list<PropsT...> &)
     ->accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type2>(),
                detail::deduceAccessTarget<Type1, Type2>(target::global_buffer),
                access::placeholder::true_t,
-               ext::ONEAPI::accessor_property_list<PropsT...>>;
+               ONEAPI::accessor_property_list<PropsT...>>;
 
 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
           typename Type2, typename Type3>
@@ -1653,11 +1659,11 @@ accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3)
 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
           typename Type2, typename Type3, typename... PropsT>
 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3,
-         const ext::ONEAPI::accessor_property_list<PropsT...> &)
+         const ONEAPI::accessor_property_list<PropsT...> &)
     ->accessor<DataT, Dimensions, detail::deduceAccessMode<Type2, Type3>(),
                detail::deduceAccessTarget<Type2, Type3>(target::global_buffer),
                access::placeholder::true_t,
-               ext::ONEAPI::accessor_property_list<PropsT...>>;
+               ONEAPI::accessor_property_list<PropsT...>>;
 
 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
           typename Type2, typename Type3, typename Type4>
@@ -1669,11 +1675,11 @@ accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3, Type4)
 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
           typename Type2, typename Type3, typename Type4, typename... PropsT>
 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3, Type4,
-         const ext::ONEAPI::accessor_property_list<PropsT...> &)
+         const ONEAPI::accessor_property_list<PropsT...> &)
     ->accessor<DataT, Dimensions, detail::deduceAccessMode<Type3, Type4>(),
                detail::deduceAccessTarget<Type3, Type4>(target::global_buffer),
                access::placeholder::true_t,
-               ext::ONEAPI::accessor_property_list<PropsT...>>;
+               ONEAPI::accessor_property_list<PropsT...>>;
 
 template <typename DataT, int Dimensions, typename AllocatorT>
 accessor(buffer<DataT, Dimensions, AllocatorT>, handler)
@@ -1683,10 +1689,10 @@ accessor(buffer<DataT, Dimensions, AllocatorT>, handler)
 template <typename DataT, int Dimensions, typename AllocatorT,
           typename... PropsT>
 accessor(buffer<DataT, Dimensions, AllocatorT>, handler,
-         const ext::ONEAPI::accessor_property_list<PropsT...> &)
+         const ONEAPI::accessor_property_list<PropsT...> &)
     ->accessor<DataT, Dimensions, access::mode::read_write,
                target::global_buffer, access::placeholder::false_t,
-               ext::ONEAPI::accessor_property_list<PropsT...>>;
+               ONEAPI::accessor_property_list<PropsT...>>;
 
 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1>
 accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1)
@@ -1697,11 +1703,11 @@ accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1)
 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
           typename... PropsT>
 accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1,
-         const ext::ONEAPI::accessor_property_list<PropsT...> &)
+         const ONEAPI::accessor_property_list<PropsT...> &)
     ->accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type1>(),
                detail::deduceAccessTarget<Type1, Type1>(target::global_buffer),
                access::placeholder::false_t,
-               ext::ONEAPI::accessor_property_list<PropsT...>>;
+               ONEAPI::accessor_property_list<PropsT...>>;
 
 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
           typename Type2>
@@ -1713,11 +1719,11 @@ accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2)
 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
           typename Type2, typename... PropsT>
 accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2,
-         const ext::ONEAPI::accessor_property_list<PropsT...> &)
+         const ONEAPI::accessor_property_list<PropsT...> &)
     ->accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type2>(),
                detail::deduceAccessTarget<Type1, Type2>(target::global_buffer),
                access::placeholder::false_t,
-               ext::ONEAPI::accessor_property_list<PropsT...>>;
+               ONEAPI::accessor_property_list<PropsT...>>;
 
 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
           typename Type2, typename Type3>
@@ -1729,11 +1735,11 @@ accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2, Type3)
 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
           typename Type2, typename Type3, typename... PropsT>
 accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2, Type3,
-         const ext::ONEAPI::accessor_property_list<PropsT...> &)
+         const ONEAPI::accessor_property_list<PropsT...> &)
     ->accessor<DataT, Dimensions, detail::deduceAccessMode<Type2, Type3>(),
                detail::deduceAccessTarget<Type2, Type3>(target::global_buffer),
                access::placeholder::false_t,
-               ext::ONEAPI::accessor_property_list<PropsT...>>;
+               ONEAPI::accessor_property_list<PropsT...>>;
 
 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
           typename Type2, typename Type3, typename Type4>
@@ -1746,11 +1752,11 @@ accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2, Type3,
 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
           typename Type2, typename Type3, typename Type4, typename... PropsT>
 accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2, Type3,
-         Type4, const ext::ONEAPI::accessor_property_list<PropsT...> &)
+         Type4, const ONEAPI::accessor_property_list<PropsT...> &)
     ->accessor<DataT, Dimensions, detail::deduceAccessMode<Type3, Type4>(),
                detail::deduceAccessTarget<Type3, Type4>(target::global_buffer),
                access::placeholder::false_t,
-               ext::ONEAPI::accessor_property_list<PropsT...>>;
+               ONEAPI::accessor_property_list<PropsT...>>;
 #endif
 
 /// Local accessor
diff --git a/sycl/include/CL/sycl/buffer.hpp b/sycl/include/CL/sycl/buffer.hpp
index 251a962d4340f..14e844f3fb75b 100755
--- a/sycl/include/CL/sycl/buffer.hpp
+++ b/sycl/include/CL/sycl/buffer.hpp
@@ -8,7 +8,7 @@
 
 #pragma once
 
-#include <CL/sycl/accessor_property_list.hpp>
+#include <CL/sycl/ONEAPI/accessor_property_list.hpp>
 #include <CL/sycl/detail/buffer_impl.hpp>
 #include <CL/sycl/detail/common.hpp>
 #include <CL/sycl/detail/stl_type_traits.hpp>
@@ -249,40 +249,40 @@ class buffer {
   template <access::mode Mode,
             access::target Target = access::target::global_buffer>
   accessor<T, dimensions, Mode, Target, access::placeholder::false_t,
-           property_list>
+           ONEAPI::accessor_property_list<>>
   get_access(handler &CommandGroupHandler) {
     return accessor<T, dimensions, Mode, Target, access::placeholder::false_t,
-                    property_list>(*this,
+                    ONEAPI::accessor_property_list<>>(*this,
                                                       CommandGroupHandler);
   }
 
   template <access::mode mode>
   accessor<T, dimensions, mode, access::target::host_buffer,
-           access::placeholder::false_t, property_list>
+           access::placeholder::false_t, ONEAPI::accessor_property_list<>>
   get_access() {
     return accessor<T, dimensions, mode, access::target::host_buffer,
                     access::placeholder::false_t,
-                    property_list>(*this);
+                    ONEAPI::accessor_property_list<>>(*this);
   }
 
   template <access::mode mode,
             access::target target = access::target::global_buffer>
   accessor<T, dimensions, mode, target, access::placeholder::false_t,
-           property_list>
+           ONEAPI::accessor_property_list<>>
   get_access(handler &commandGroupHandler, range<dimensions> accessRange,
              id<dimensions> accessOffset = {}) {
     return accessor<T, dimensions, mode, target, access::placeholder::false_t,
-                    property_list>(
+                    ONEAPI::accessor_property_list<>>(
         *this, commandGroupHandler, accessRange, accessOffset);
   }
 
   template <access::mode mode>
   accessor<T, dimensions, mode, access::target::host_buffer,
-           access::placeholder::false_t, property_list>
+           access::placeholder::false_t, ONEAPI::accessor_property_list<>>
   get_access(range<dimensions> accessRange, id<dimensions> accessOffset = {}) {
     return accessor<T, dimensions, mode, access::target::host_buffer,
                     access::placeholder::false_t,
-                    property_list>(*this, accessRange,
+                    ONEAPI::accessor_property_list<>>(*this, accessRange,
                                                       accessOffset);
   }
 
diff --git a/sycl/include/CL/sycl/detail/property_list_base.hpp b/sycl/include/CL/sycl/detail/property_list_base.hpp
index b02d1cd80c0b1..8fc13e50dd0d0 100644
--- a/sycl/include/CL/sycl/detail/property_list_base.hpp
+++ b/sycl/include/CL/sycl/detail/property_list_base.hpp
@@ -103,7 +103,6 @@ class PropertyListBase {
                                      PI_INVALID_VALUE);
   }
 
-private:
   // Stores enable/not enabled for simple properties
   std::bitset<DataLessPropKind::DataLessPropKindSize> MDataLessProps;
   // Stores shared_ptrs to complex properties
@@ -111,4 +110,4 @@ class PropertyListBase {
 };
 } // namespace detail
 } // namespace sycl
-} // __SYCL_INLINE_NAMESPACE(cl)
\ No newline at end of file
+} // __SYCL_INLINE_NAMESPACE(cl)
diff --git a/sycl/include/CL/sycl/image.hpp b/sycl/include/CL/sycl/image.hpp
index 221448250f02c..bd8216a2f5dd1 100644
--- a/sycl/include/CL/sycl/image.hpp
+++ b/sycl/include/CL/sycl/image.hpp
@@ -8,7 +8,7 @@
 
 #pragma once
 
-#include "CL/sycl/accessor_property_list.hpp"
+#include "CL/sycl/ONEAPI/accessor_property_list.hpp"
 #include <CL/sycl/detail/common.hpp>
 #include <CL/sycl/detail/generic_type_traits.hpp>
 #include <CL/sycl/detail/image_impl.hpp>
@@ -274,22 +274,22 @@ class image {
   template <typename DataT, access::mode AccessMode>
   accessor<detail::EnableIfImgAccDataT<DataT>, Dimensions, AccessMode,
            access::target::image, access::placeholder::false_t,
-           property_list>
+           ONEAPI::accessor_property_list<>>
   get_access(handler &commandGroupHandler) {
     return accessor<DataT, Dimensions, AccessMode, access::target::image,
                     access::placeholder::false_t,
-                    property_list>(*this,
+                    ONEAPI::accessor_property_list<>>(*this,
                                                       commandGroupHandler);
   }
 
   template <typename DataT, access::mode AccessMode>
   accessor<detail::EnableIfImgAccDataT<DataT>, Dimensions, AccessMode,
            access::target::host_image, access::placeholder::false_t,
-           property_list>
+           ONEAPI::accessor_property_list<>>
   get_access() {
     return accessor<DataT, Dimensions, AccessMode, access::target::host_image,
                     access::placeholder::false_t,
-                    property_list>(*this);
+                    ONEAPI::accessor_property_list<>>(*this);
   }
 
   template <typename Destination = std::nullptr_t>
diff --git a/sycl/include/CL/sycl/properties/accessor_properties.hpp b/sycl/include/CL/sycl/properties/accessor_properties.hpp
index b1272b48dc8ec..2c9cd04dcce3d 100644
--- a/sycl/include/CL/sycl/properties/accessor_properties.hpp
+++ b/sycl/include/CL/sycl/properties/accessor_properties.hpp
@@ -8,7 +8,7 @@
 
 #pragma once
 
-#include "CL/sycl/accessor_property_list.hpp"
+#include <CL/sycl/ONEAPI/accessor_property_list.hpp>
 #include <CL/sycl/detail/common.hpp>
 #include <CL/sycl/detail/property_helper.hpp>
 #include <type_traits>
@@ -35,87 +35,82 @@ constexpr const auto &noinit =
 
 #endif
 
-namespace ext {
 namespace INTEL {
 namespace property {
-template <int A> struct buffer_location {
-  constexpr bool operator==(const buffer_location<A> &) const { return true; }
-  constexpr bool operator!=(const buffer_location<A> &) const { return false; }
-  template <bool B>
-  constexpr bool operator==(const buffer_location<B> &) const {
-    return false;
-  }
-  template <bool B>
-  constexpr bool operator!=(const buffer_location<B> &) const {
-    return true;
-  }
+struct buffer_location {
+  template <int A> struct instance {
+    constexpr bool operator==(const buffer_location::instance<A> &) const {
+      return true;
+    }
+    constexpr bool operator!=(const buffer_location::instance<A> &) const {
+      return false;
+    }
+    template <int B>
+    constexpr bool operator==(const buffer_location::instance<B> &) const {
+      return false;
+    }
+    template <int B>
+    constexpr bool operator!=(const buffer_location::instance<B> &) const {
+      return true;
+    }
+  };
 };
-
 } // namespace property
+#if __cplusplus > 201402L
+template <int A>
+inline constexpr property::buffer_location::instance<A> buffer_location;
+#endif
 } // namespace INTEL
 namespace ONEAPI {
 namespace property {
-template <bool A = true> struct no_offset {
-  constexpr bool operator==(const no_offset<A> &) const {
-    return true;
-  }
-  constexpr bool operator!=(const no_offset<A> &) const {
-    return false;
-  }
-  template <bool B>
-  constexpr bool operator==(const no_offset<B> &) const {
-    return false;
-  }
-  template <bool B>
-  constexpr bool operator!=(const no_offset<B> &) const {
-    return true;
-  }
+struct no_offset {
+  template <bool B = true> struct instance {
+    constexpr bool operator==(const no_offset::instance<B> &) const {
+      return true;
+    }
+    constexpr bool operator!=(const no_offset::instance<B> &) const {
+      return false;
+    }
+  };
 };
-template <bool A = true> struct no_alias {
-  constexpr bool operator==(const no_alias<A> &) const {
-    return true;
-  }
-  constexpr bool operator!=(const no_alias<A> &) const {
-    return false;
-  }
-  template <bool B>
-  constexpr bool operator==(const no_alias<B> &) const {
-    return false;
-  }
-  template <bool B>
-  constexpr bool operator!=(const no_alias<B> &) const {
-    return true;
-  }
-
+struct no_alias {
+  template <bool B = true> struct instance {
+    constexpr bool operator==(const no_alias::instance<B> &) const {
+      return true;
+    }
+    constexpr bool operator!=(const no_alias::instance<B> &) const {
+      return false;
+    }
+  };
 };
 } // namespace property
 
 #if __cplusplus > 201402L
 
-inline constexpr property::no_offset no_offset;
-inline constexpr property::no_alias no_alias;
-
-#else
-
-namespace {
-constexpr const auto &no_offset =
-    sycl::detail::InlineVariableHelper<property::no_offset<>>::value;
-constexpr const auto &no_alias =
-    sycl::detail::InlineVariableHelper<property::no_alias<>>::value;
-} // namespace
+inline constexpr property::no_offset::instance no_offset;
+inline constexpr property::no_alias::instance no_alias;
 
 #endif
 
-template <bool B>
-struct is_compile_time_property<ext::ONEAPI::property::no_offset<B>>
-    : std::true_type {};
-template <bool B>
-struct is_compile_time_property<ext::ONEAPI::property::no_alias<B>>
+template <>
+struct is_compile_time_property<ONEAPI::property::no_offset> : std::true_type {
+};
+template <>
+struct is_compile_time_property<ONEAPI::property::no_alias> : std::true_type {};
+template <>
+struct is_compile_time_property<INTEL::property::buffer_location>
     : std::true_type {};
+} // namespace ONEAPI
+namespace detail {
 template <int I>
-struct is_compile_time_property<ext::INTEL::property::buffer_location<I>>
+struct IsCxPropertyInstance<INTEL::property::buffer_location::instance<I>>
     : std::true_type {};
-} // namespace ONEAPI
-} // namespace ext
+template <>
+struct IsCxPropertyInstance<ONEAPI::property::no_alias::instance<>>
+    : std::true_type {};
+template <>
+struct IsCxPropertyInstance<ONEAPI::property::no_offset::instance<>>
+    : std::true_type {};
+} // namespace detail
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)
diff --git a/sycl/include/CL/sycl/property_list.hpp b/sycl/include/CL/sycl/property_list.hpp
index abd066b127d0e..d6ce4279f9e7b 100644
--- a/sycl/include/CL/sycl/property_list.hpp
+++ b/sycl/include/CL/sycl/property_list.hpp
@@ -51,8 +51,10 @@ class property_list : protected detail::PropertyListBase {
     return has_property_helper<PropT>();
   }
 
-  template <typename... PropsT>
-  operator ONEAPI::accessor_property_list<PropsT...>();
+  template <typename... T> operator ONEAPI::accessor_property_list<T...>();
+
+private:
+  template <typename... PropsT> friend class ONEAPI::accessor_property_list;
 };
 
 } // namespace sycl
diff --git a/sycl/include/CL/sycl/property_list_conversion.hpp b/sycl/include/CL/sycl/property_list_conversion.hpp
index e69de29bb2d1d..f8195a9f315f1 100644
--- a/sycl/include/CL/sycl/property_list_conversion.hpp
+++ b/sycl/include/CL/sycl/property_list_conversion.hpp
@@ -0,0 +1,23 @@
+//==---- property_list_conversion.hpp --- SYCL property list conversion ----==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+// This file contains conversion routines from property_list to
+// accessor_property_list. A separate file helps to avoid cyclic dependencies
+// between header files.
+
+#include <CL/sycl/ONEAPI/accessor_property_list.hpp>
+#include <CL/sycl/property_list.hpp>
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+template <typename... T>
+inline property_list::operator ONEAPI::accessor_property_list<T...>() {
+  return ONEAPI::accessor_property_list<T...>(MDataLessProps, MPropsWithData);
+}
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)
diff --git a/sycl/test/basic_tests/accessor/accessor_property_list.cpp b/sycl/test/basic_tests/accessor/accessor_property_list.cpp
index 51a980342af4a..4bb43681f4d4a 100644
--- a/sycl/test/basic_tests/accessor/accessor_property_list.cpp
+++ b/sycl/test/basic_tests/accessor/accessor_property_list.cpp
@@ -8,7 +8,7 @@
 
 #include <type_traits>
 
-using namespace sycl::ext::ONEAPI;
+using namespace sycl::ONEAPI;
 
 int main() {
   {
@@ -44,9 +44,6 @@ int main() {
 
     static_assert(PL.get_property<property::no_alias>() == no_alias,
                   "Properties are not equal");
-    property::no_alias<false> NAfalse;
-    static_assert(PL.get_property<property::no_alias>() != NAfalse,
-                  "Properties are equal");
   }
 
   {
@@ -99,9 +96,9 @@ int main() {
     sycl::accessor<int, 1, sycl::access::mode::read_write,
                    sycl::access::target::global_buffer,
                    sycl::access::placeholder::true_t,
-                   accessor_property_list<property::no_alias<>>>
+                   accessor_property_list<property::no_alias::instance<>>>
         acc_2(acc_1);
   }
 
   return 0;
-}
\ No newline at end of file
+}

From ee55a6a2a08f645b959b56c1b1abc051b889420e Mon Sep 17 00:00:00 2001
From: Alexander Batashev <alexander.batashev@intel.com>
Date: Thu, 10 Sep 2020 14:12:05 +0300
Subject: [PATCH 04/19] Fix tests

---
 sycl/test/abi/layout_accessors.cpp | 12 +++++++-----
 1 file changed, 7 insertions(+), 5 deletions(-)

diff --git a/sycl/test/abi/layout_accessors.cpp b/sycl/test/abi/layout_accessors.cpp
index 8fcb3ac47f8b8..65812402b87a7 100644
--- a/sycl/test/abi/layout_accessors.cpp
+++ b/sycl/test/abi/layout_accessors.cpp
@@ -1,6 +1,8 @@
 // RUN: %clangxx -fsycl -c -fno-color-diagnostics -Xclang -fdump-record-layouts %s | FileCheck %s
 // REQUIRES: linux
 
+// clang-format off
+
 #include <CL/sycl/accessor.hpp>
 
 using namespace cl::sycl;
@@ -13,7 +15,7 @@ void hostAcc(accessor<int, 1, access::mode::read, access::target::host_buffer> A
   (void)Acc.get_size();
 }
 
-// CHECK: 0 | class cl::sycl::accessor<int, 1, cl::sycl::access::mode::read, cl::sycl::access::target::host_buffer, cl::sycl::access::placeholder::false_t, class cl::sycl::property_list>
+// CHECK: 0 | class cl::sycl::accessor<int, 1, cl::sycl::access::mode::read, cl::sycl::access::target::host_buffer, cl::sycl::access::placeholder::false_t, class cl::sycl::ONEAPI::accessor_property_list<> >
 // CHECK-NEXT:  0 |   class cl::sycl::detail::AccessorBaseHost (base)
 // CHECK-NEXT:  0 |     class std::shared_ptr<class cl::sycl::detail::AccessorImplHost> impl
 // CHECK-NEXT:  0 |       class std::__shared_ptr<class cl::sycl::detail::AccessorImplHost, __gnu_cxx::_S_atomic> (base)
@@ -34,7 +36,7 @@ void hostAcc(accessor<int, 1, access::mode::read, access::target::global_buffer>
   (void)Acc.get_size();
 }
 
-// CHECK: 0 | class cl::sycl::accessor<int, 1, cl::sycl::access::mode::read, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t, class cl::sycl::property_list>
+// CHECK: 0 | class cl::sycl::accessor<int, 1, cl::sycl::access::mode::read, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t, class cl::sycl::ONEAPI::accessor_property_list<> >
 // CHECK-NEXT: 0 |   class cl::sycl::detail::AccessorBaseHost (base)
 // CHECK-NEXT: 0 |     class std::shared_ptr<class cl::sycl::detail::AccessorImplHost> impl
 // CHECK-NEXT: 0 |       class std::__shared_ptr<class cl::sycl::detail::AccessorImplHost, __gnu_cxx::_S_atomic> (base)
@@ -55,7 +57,7 @@ void hostAcc(accessor<int, 1, access::mode::read_write, access::target::local> A
   (void)Acc.get_size();
 }
 
-// CHECK: 0 | class cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::local, cl::sycl::access::placeholder::false_t, class cl::sycl::property_list>
+// CHECK: 0 | class cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::local, cl::sycl::access::placeholder::false_t, class cl::sycl::ONEAPI::accessor_property_list<> >
 // CHECK-NEXT: 0 |   class cl::sycl::detail::LocalAccessorBaseHost (base)
 // CHECK-NEXT: 0 |     class std::shared_ptr<class cl::sycl::detail::LocalAccessorImplHost> impl
 // CHECK-NEXT: 0 |       class std::__shared_ptr<class cl::sycl::detail::LocalAccessorImplHost, __gnu_cxx::_S_atomic> (base)
@@ -75,7 +77,7 @@ void hostAcc(accessor<int, 1, access::mode::read_write, access::target::local> A
 void hostAcc(accessor<int4, 1, access::mode::read_write, access::target::host_image> Acc) {
   (void)Acc.get_count();
 }
-// CHECK: 0 | class cl::sycl::accessor<class cl::sycl::vec<int, 4>, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::host_image, cl::sycl::access::placeholder::false_t, class cl::sycl::property_list>
+// CHECK: 0 | class cl::sycl::accessor<class cl::sycl::vec<int, 4>, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::host_image, cl::sycl::access::placeholder::false_t, class cl::sycl::ONEAPI::accessor_property_list<> >
 // CHECK-NEXT: 0 |   class cl::sycl::detail::image_accessor<class cl::sycl::vec<int, 4>, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::host_image, cl::sycl::access::placeholder::false_t> (base)
 // CHECK-NEXT: 0 |     class cl::sycl::detail::AccessorBaseHost (base)
 // CHECK-NEXT: 0 |       class std::shared_ptr<class cl::sycl::detail::AccessorImplHost> impl
@@ -97,7 +99,7 @@ void hostAcc(accessor<int4, 1, access::mode::read_write, access::target::host_im
 void hostAcc(accessor<int4, 1, access::mode::read, access::target::image> Acc) {
   (void)Acc.get_count();
 }
-// CHECK: 0 | class cl::sycl::accessor<class cl::sycl::vec<int, 4>, 1, cl::sycl::access::mode::read, cl::sycl::access::target::image, cl::sycl::access::placeholder::false_t, class cl::sycl::property_list>
+// CHECK: 0 | class cl::sycl::accessor<class cl::sycl::vec<int, 4>, 1, cl::sycl::access::mode::read, cl::sycl::access::target::image, cl::sycl::access::placeholder::false_t, class cl::sycl::ONEAPI::accessor_property_list<> >
 // CHECK-NEXT: 0 |   class cl::sycl::detail::image_accessor<class cl::sycl::vec<int, 4>, 1, cl::sycl::access::mode::read, cl::sycl::access::target::image, cl::sycl::access::placeholder::false_t> (base)
 // CHECK-NEXT: 0 |     class cl::sycl::detail::AccessorBaseHost (base)
 // CHECK-NEXT: 0 |       class std::shared_ptr<class cl::sycl::detail::AccessorImplHost> impl

From 976dfec76bab39c774c189c188fc60f088221518 Mon Sep 17 00:00:00 2001
From: Alexander Batashev <alexander.batashev@intel.com>
Date: Thu, 10 Sep 2020 14:13:34 +0300
Subject: [PATCH 05/19] More clang-format

---
 sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp | 5 ++---
 sycl/include/CL/sycl/handler.hpp                       | 6 ++++--
 sycl/include/CL/sycl/property_list.hpp                 | 3 +--
 3 files changed, 7 insertions(+), 7 deletions(-)

diff --git a/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp b/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
index 0adfd51f2051e..e23a5e3370eb2 100644
--- a/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
+++ b/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
@@ -44,14 +44,13 @@ class accessor_property_list : protected detail::PropertyListBase {
   template <template <class...> class T, class T1, class T2>
   struct AreSameTemplate<T<T1>, T<T2>> : std::true_type {};
 #if __cplusplus >= 201703L
-  template <template <auto...> class T, auto...T1, auto...T2>
+  template <template <auto...> class T, auto... T1, auto... T2>
   struct AreSameTemplate<T<T1...>, T<T2...>> : std::true_type {};
 #endif
   // This template helps to identify if PropListT parameter pack contains
   // property of PropT type, where PropT is a nested instance template of
   // compile-time-constant property.
-  template <typename PropT, typename... PropListT>
-  struct ContainsProperty;
+  template <typename PropT, typename... PropListT> struct ContainsProperty;
   template <typename PropT> struct ContainsProperty<PropT> : std::false_type {};
   template <typename PropT, typename Head, typename... Tail>
   struct ContainsProperty<PropT, Head, Tail...>
diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp
index e304aab921237..77db78115de05 100644
--- a/sycl/include/CL/sycl/handler.hpp
+++ b/sycl/include/CL/sycl/handler.hpp
@@ -1769,8 +1769,10 @@ class __SYCL_EXPORT handler {
             access::target AccessTarget,
             access::placeholder IsPlaceholder = access::placeholder::false_t,
             typename PropertyListT = property_list>
-  void fill(accessor<T, Dims, AccessMode, AccessTarget, IsPlaceholder, PropertyListT> Dst,
-            const T &Pattern) {
+  void
+  fill(accessor<T, Dims, AccessMode, AccessTarget, IsPlaceholder, PropertyListT>
+           Dst,
+       const T &Pattern) {
     throwIfActionIsCreated();
     // TODO add check:T must be an integral scalar value or a SYCL vector type
     static_assert(isValidTargetForExplicitOp(AccessTarget),
diff --git a/sycl/include/CL/sycl/property_list.hpp b/sycl/include/CL/sycl/property_list.hpp
index d6ce4279f9e7b..b3bcb413a324d 100644
--- a/sycl/include/CL/sycl/property_list.hpp
+++ b/sycl/include/CL/sycl/property_list.hpp
@@ -14,8 +14,7 @@
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
 namespace ONEAPI {
-template <typename... PropsT>
-class accessor_property_list;
+template <typename... PropsT> class accessor_property_list;
 }
 
 /// Objects of the property_list class are containers for the SYCL properties

From 33da077f948d0719b2169afd483119232584a9de Mon Sep 17 00:00:00 2001
From: Alexander Batashev <alexander.batashev@intel.com>
Date: Thu, 10 Sep 2020 16:02:37 +0300
Subject: [PATCH 06/19] Fix template arguments

---
 sycl/include/CL/sycl/ONEAPI/atomic_accessor.hpp | 7 ++++---
 1 file changed, 4 insertions(+), 3 deletions(-)

diff --git a/sycl/include/CL/sycl/ONEAPI/atomic_accessor.hpp b/sycl/include/CL/sycl/ONEAPI/atomic_accessor.hpp
index ed415a90891c7..123cd222046ab 100644
--- a/sycl/include/CL/sycl/ONEAPI/atomic_accessor.hpp
+++ b/sycl/include/CL/sycl/ONEAPI/atomic_accessor.hpp
@@ -45,10 +45,11 @@ template <typename DataT, int Dimensions, memory_order DefaultOrder,
           access::placeholder IsPlaceholder = access::placeholder::false_t>
 class atomic_accessor
     : public accessor<DataT, Dimensions, access::mode::read_write, AccessTarget,
-                      IsPlaceholder> {
+                      IsPlaceholder, ONEAPI::accessor_property_list<>> {
 
-  using AccessorT = accessor<DataT, Dimensions, access::mode::read_write,
-                             AccessTarget, IsPlaceholder>;
+  using AccessorT =
+      accessor<DataT, Dimensions, access::mode::read_write, AccessTarget,
+               IsPlaceholder, ONEAPI::accessor_property_list<>>;
 
 private:
   using AccessorT::getLinearIndex;

From 5dba5833bc2565d6ec50d3fddf55eff7366c63f1 Mon Sep 17 00:00:00 2001
From: Alexander Batashev <alexander.batashev@intel.com>
Date: Fri, 11 Sep 2020 09:46:49 +0300
Subject: [PATCH 07/19] Fix more missing templates

---
 sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp | 3 ---
 sycl/include/CL/sycl/ONEAPI/reduction.hpp              | 8 +++++---
 2 files changed, 5 insertions(+), 6 deletions(-)

diff --git a/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp b/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
index e23a5e3370eb2..95cb1f2802e82 100644
--- a/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
+++ b/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
@@ -212,9 +212,6 @@ class accessor_property_list : protected detail::PropertyListBase {
 
 private:
   template <typename... OtherProps> friend class accessor_property_list;
-  template <typename, int, sycl::access::mode, sycl::access::target,
-            sycl::access::placeholder, typename>
-  friend class accessor;
 
   friend class sycl::property_list;
 };
diff --git a/sycl/include/CL/sycl/ONEAPI/reduction.hpp b/sycl/include/CL/sycl/ONEAPI/reduction.hpp
index e33db4783156b..1296d2a90e3d5 100644
--- a/sycl/include/CL/sycl/ONEAPI/reduction.hpp
+++ b/sycl/include/CL/sycl/ONEAPI/reduction.hpp
@@ -8,6 +8,7 @@
 
 #pragma once
 
+#include "CL/sycl/ONEAPI/accessor_property_list.hpp"
 #include <CL/sycl/ONEAPI/group_algorithm.hpp>
 #include <CL/sycl/accessor.hpp>
 #include <CL/sycl/handler.hpp>
@@ -348,11 +349,12 @@ class reduction_impl {
   using reducer_type = reducer<T, BinaryOperation>;
   using result_type = T;
   using binary_operation = BinaryOperation;
-  using accessor_type =
-      accessor<T, Dims, AccMode, access::target::global_buffer, IsPlaceholder>;
+  using accessor_type = accessor < T, Dims, AccMode,
+        access::target::global_buffer, IsPlaceholder,
+        ONEAPI::accessor_property_list<>>;
   using rw_accessor_type =
       accessor<T, Dims, access::mode::read_write, access::target::global_buffer,
-               IsPlaceholder>;
+               IsPlaceholder, ONEAPI::accessor_property_list<>>;
   static constexpr access::mode accessor_mode = AccMode;
   static constexpr int accessor_dim = Dims;
   static constexpr int buffer_dim = (Dims == 0) ? 1 : Dims;

From 809061d5071a39b917b15b0c3056d4dc2dbca89b Mon Sep 17 00:00:00 2001
From: Alexander Batashev <alexander.batashev@intel.com>
Date: Fri, 11 Sep 2020 09:49:39 +0300
Subject: [PATCH 08/19] clang-format

---
 sycl/include/CL/sycl/ONEAPI/reduction.hpp | 6 +++---
 1 file changed, 3 insertions(+), 3 deletions(-)

diff --git a/sycl/include/CL/sycl/ONEAPI/reduction.hpp b/sycl/include/CL/sycl/ONEAPI/reduction.hpp
index 1296d2a90e3d5..a11d3d89fb513 100644
--- a/sycl/include/CL/sycl/ONEAPI/reduction.hpp
+++ b/sycl/include/CL/sycl/ONEAPI/reduction.hpp
@@ -349,9 +349,9 @@ class reduction_impl {
   using reducer_type = reducer<T, BinaryOperation>;
   using result_type = T;
   using binary_operation = BinaryOperation;
-  using accessor_type = accessor < T, Dims, AccMode,
-        access::target::global_buffer, IsPlaceholder,
-        ONEAPI::accessor_property_list<>>;
+  using accessor_type =
+      accessor<T, Dims, AccMode, access::target::global_buffer, IsPlaceholder,
+               ONEAPI::accessor_property_list<>>;
   using rw_accessor_type =
       accessor<T, Dims, access::mode::read_write, access::target::global_buffer,
                IsPlaceholder, ONEAPI::accessor_property_list<>>;

From e9b9246e033711d3a9e1cc0d98f6856cdbb0e67b Mon Sep 17 00:00:00 2001
From: Alexander Batashev <alexander.batashev@intel.com>
Date: Fri, 11 Sep 2020 10:32:01 +0300
Subject: [PATCH 09/19] Propagate template arguments to accessor_common

---
 sycl/include/CL/sycl/accessor.hpp  | 15 +++++++++------
 sycl/test/abi/layout_accessors.cpp |  6 +++---
 2 files changed, 12 insertions(+), 9 deletions(-)

diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp
index feb0e1966549e..debba0f9cc02d 100755
--- a/sycl/include/CL/sycl/accessor.hpp
+++ b/sycl/include/CL/sycl/accessor.hpp
@@ -260,7 +260,8 @@ static T<NewDim> convertToArrayOfN(T<OldDim> OldObj) {
 __SYCL_EXPORT device getDeviceFromHandler(handler &CommandGroupHandlerRef);
 
 template <typename DataT, int Dimensions, access::mode AccessMode,
-          access::target AccessTarget, access::placeholder IsPlaceholder>
+          access::target AccessTarget, access::placeholder IsPlaceholder,
+          typename PropertyListT = ONEAPI::accessor_property_list<>>
 class accessor_common {
 protected:
   constexpr static bool IsPlaceH = IsPlaceholder == access::placeholder::true_t;
@@ -289,8 +290,9 @@ class accessor_common {
   using ConstRefType = const DataT &;
   using PtrType = detail::const_if_const_AS<AS, DataT> *;
 
-  using AccType =
-      accessor<DataT, Dimensions, AccessMode, AccessTarget, IsPlaceholder>;
+  using AccType = accessor < DataT, Dimensions, AccessMode, AccessTarget,
+        IsPlaceholder,
+        PropertyListT>;
 
   // The class which allows to access value of N dimensional accessor using N
   // subscript operators, e.g. accessor[2][2][3]
@@ -763,7 +765,7 @@ class accessor :
     public detail::AccessorBaseHost,
 #endif
     public detail::accessor_common<DataT, Dimensions, AccessMode, AccessTarget,
-                                   IsPlaceholder> {
+                                   IsPlaceholder, PropertyListT> {
 protected:
   static_assert((AccessTarget == access::target::global_buffer ||
                  AccessTarget == access::target::constant_buffer ||
@@ -779,8 +781,9 @@ class accessor :
   static_assert(detail::IsPropertyListT<PropertyListT>::value,
                 "PropertyListT must be accessor_property_list");
 
-  using AccessorCommonT = detail::accessor_common<DataT, Dimensions, AccessMode,
-                                                  AccessTarget, IsPlaceholder>;
+  using AccessorCommonT = detail::accessor_common < DataT, Dimensions,
+        AccessMode, AccessTarget, IsPlaceholder,
+        PropertyListT>;
 
   constexpr static int AdjustedDim = Dimensions == 0 ? 1 : Dimensions;
 
diff --git a/sycl/test/abi/layout_accessors.cpp b/sycl/test/abi/layout_accessors.cpp
index 65812402b87a7..38dc8c6156957 100644
--- a/sycl/test/abi/layout_accessors.cpp
+++ b/sycl/test/abi/layout_accessors.cpp
@@ -23,7 +23,7 @@ void hostAcc(accessor<int, 1, access::mode::read, access::target::host_buffer> A
 // CHECK-NEXT:  0 |         std::__shared_ptr<class cl::sycl::detail::AccessorImplHost, __gnu_cxx::_S_atomic>::element_type * _M_ptr
 // CHECK-NEXT:  8 |         class std::__shared_count<__gnu_cxx::_S_atomic> _M_refcount
 // CHECK-NEXT:  8 |           _Sp_counted_base<(enum __gnu_cxx::_Lock_policy)2U> * _M_pi
-// CHECK-NEXT:  0 |   class cl::sycl::detail::accessor_common<int, 1, cl::sycl::access::mode::read, cl::sycl::access::target::host_buffer, cl::sycl::access::placeholder::false_t> (base) (empty)
+// CHECK-NEXT:  0 |   class cl::sycl::detail::accessor_common<int, 1, cl::sycl::access::mode::read, cl::sycl::access::target::host_buffer, cl::sycl::access::placeholder::false_t, class cl::sycl::ONEAPI::accessor_property_list<> > (base) (empty)
 // CHECK-NEXT: 16 |   char [16] padding
 // CHECK-NEXT: [sizeof=32, dsize=32, align=8,
 // CHECK-NEXT: nvsize=32, nvalign=8]
@@ -44,7 +44,7 @@ void hostAcc(accessor<int, 1, access::mode::read, access::target::global_buffer>
 // CHECK-NEXT: 0 |         std::__shared_ptr<class cl::sycl::detail::AccessorImplHost, __gnu_cxx::_S_atomic>::element_type * _M_ptr
 // CHECK-NEXT: 8 |         class std::__shared_count<__gnu_cxx::_S_atomic> _M_refcount
 // CHECK-NEXT: 8 |           _Sp_counted_base<(enum __gnu_cxx::_Lock_policy)2U> * _M_pi
-// CHECK-NEXT: 0 |   class cl::sycl::detail::accessor_common<int, 1, cl::sycl::access::mode::read, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t> (base) (empty)
+// CHECK-NEXT: 0 |   class cl::sycl::detail::accessor_common<int, 1, cl::sycl::access::mode::read, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t, class cl::sycl::ONEAPI::accessor_property_list<> > (base) (empty)
 // CHECK-NEXT: 16 |   char [16] padding
 // CHECK-NEXT: [sizeof=32, dsize=32, align=8,
 // CHECK-NEXT: nvsize=32, nvalign=8]
@@ -65,7 +65,7 @@ void hostAcc(accessor<int, 1, access::mode::read_write, access::target::local> A
 // CHECK-NEXT: 0 |         std::__shared_ptr<class cl::sycl::detail::LocalAccessorImplHost, __gnu_cxx::_S_atomic>::element_type * _M_ptr
 // CHECK-NEXT: 8 |         class std::__shared_count<__gnu_cxx::_S_atomic> _M_refcount
 // CHECK-NEXT: 8 |           _Sp_counted_base<(enum __gnu_cxx::_Lock_policy)2U> * _M_pi
-// CHECK-NEXT: 0 |   class cl::sycl::detail::accessor_common<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::local, cl::sycl::access::placeholder::false_t> (base) (empty)
+// CHECK-NEXT: 0 |   class cl::sycl::detail::accessor_common<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::local, cl::sycl::access::placeholder::false_t, class cl::sycl::ONEAPI::accessor_property_list<> > (base) (empty)
 // CHECK-NEXT: 16 |   char [16] padding
 // CHECK-NEXT: [sizeof=32, dsize=32, align=8,
 // CHECK-NEXT: nvsize=32, nvalign=8]

From ccb172fca85bd0d58312f9d69e7ada5be6831cd3 Mon Sep 17 00:00:00 2001
From: Alexander Batashev <alexander.batashev@intel.com>
Date: Fri, 11 Sep 2020 14:15:01 +0300
Subject: [PATCH 10/19] More clang-format

---
 sycl/include/CL/sycl/accessor.hpp | 22 +++++++++++-----------
 1 file changed, 11 insertions(+), 11 deletions(-)

diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp
index debba0f9cc02d..d59326391e0ad 100755
--- a/sycl/include/CL/sycl/accessor.hpp
+++ b/sycl/include/CL/sycl/accessor.hpp
@@ -290,9 +290,8 @@ class accessor_common {
   using ConstRefType = const DataT &;
   using PtrType = detail::const_if_const_AS<AS, DataT> *;
 
-  using AccType = accessor < DataT, Dimensions, AccessMode, AccessTarget,
-        IsPlaceholder,
-        PropertyListT>;
+  using AccType = accessor<DataT, Dimensions, AccessMode, AccessTarget,
+                           IsPlaceholder, PropertyListT>;
 
   // The class which allows to access value of N dimensional accessor using N
   // subscript operators, e.g. accessor[2][2][3]
@@ -688,10 +687,11 @@ class __image_array_slice__ {
   }
 
 public:
-  __image_array_slice__(accessor<DataT, Dimensions, AccessMode,
-                                 access::target::image_array, IsPlaceholder>
-                            BaseAcc,
-                        size_t Idx)
+  __image_array_slice__(
+      accessor<DataT, Dimensions, AccessMode, access::target::image_array,
+               IsPlaceholder, ONEAPI::accessor_property_list<>>
+          BaseAcc,
+      size_t Idx)
       : MBaseAcc(BaseAcc), MIdx(Idx) {}
 
   template <typename CoordT, int Dims = Dimensions,
@@ -746,7 +746,7 @@ class __image_array_slice__ {
 private:
   size_t MIdx;
   accessor<DataT, Dimensions, AccessMode, access::target::image_array,
-           IsPlaceholder>
+           IsPlaceholder, ONEAPI::accessor_property_list<>>
       MBaseAcc;
 };
 
@@ -781,9 +781,9 @@ class accessor :
   static_assert(detail::IsPropertyListT<PropertyListT>::value,
                 "PropertyListT must be accessor_property_list");
 
-  using AccessorCommonT = detail::accessor_common < DataT, Dimensions,
-        AccessMode, AccessTarget, IsPlaceholder,
-        PropertyListT>;
+  using AccessorCommonT =
+      detail::accessor_common<DataT, Dimensions, AccessMode, AccessTarget,
+                              IsPlaceholder, PropertyListT>;
 
   constexpr static int AdjustedDim = Dimensions == 0 ? 1 : Dimensions;
 

From ff28dc28239883ca2e7ad044ae53f34d4bb1ccd4 Mon Sep 17 00:00:00 2001
From: Alexander Batashev <alexander.batashev@intel.com>
Date: Fri, 11 Sep 2020 14:59:40 +0300
Subject: [PATCH 11/19] Remove unused code

---
 .../CL/sycl/ONEAPI/accessor_property_list.hpp   | 17 +----------------
 1 file changed, 1 insertion(+), 16 deletions(-)

diff --git a/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp b/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
index 95cb1f2802e82..8b7fe91c252b8 100644
--- a/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
+++ b/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
@@ -37,7 +37,7 @@ template <typename T> struct is_compile_time_property : std::false_type {};
 ///
 /// \ingroup sycl_api
 template <typename... PropsT>
-class accessor_property_list : protected detail::PropertyListBase {
+class accessor_property_list : protected sycl::detail::PropertyListBase {
   // These structures check if compile-time-constant property is present in
   // list. For runtime properties this check is always true.
   template <class T, class U> struct AreSameTemplate : std::is_same<T, U> {};
@@ -187,21 +187,6 @@ class accessor_property_list : protected detail::PropertyListBase {
     return typename GetCxPropertyHelper<PropertyContainer<PropsT...>,
                                         T::template instance>::type{};
   }
-  /*
-  template <template <auto...> class PropT, auto... Args>
-  static constexpr
-      typename std::enable_if<detail::IsCxPropertyInstance<typename
-  PropT::instance<Args...>>::value, bool>::type has_property() { return
-  ContainsProperty<typename PropT::instance<Args...>, PropsT...>::value;
-  }
-  template <template <auto...> class PropT, auto... Args,
-            typename = typename std::enable_if<
-                detail::IsCxPropertyInstance<typename
-  PropT::instance<Args...>>::value>::type> static constexpr auto get_property()
-  { return typename GetCxPropertyHelper<PropertyContainer<PropsT...>, typename
-  PropT::instance, Args...>::type{};
-  }
-  */
 #endif
 
   template <typename... OtherPropsT>

From 49028d5742213031d28daa255fbac37c2d8aeb97 Mon Sep 17 00:00:00 2001
From: Alexander Batashev <alexander.batashev@intel.com>
Date: Fri, 11 Sep 2020 16:04:09 +0300
Subject: [PATCH 12/19] Fix empty prop list lookup

---
 .../CL/sycl/ONEAPI/accessor_property_list.hpp | 38 +++++++++++--------
 .../accessor/accessor_property_list.cpp       |  7 ++++
 2 files changed, 30 insertions(+), 15 deletions(-)

diff --git a/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp b/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
index 8b7fe91c252b8..6257e28bb6a43 100644
--- a/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
+++ b/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
@@ -55,7 +55,7 @@ class accessor_property_list : protected sycl::detail::PropertyListBase {
   template <typename PropT, typename Head, typename... Tail>
   struct ContainsProperty<PropT, Head, Tail...>
       : std::conditional<AreSameTemplate<PropT, Head>::value ||
-                             !detail::IsCxPropertyInstance<PropT>::value,
+                             !sycl::detail::IsCxPropertyInstance<PropT>::value,
                          std::true_type,
                          ContainsProperty<PropT, Tail...>>::type {};
 
@@ -82,8 +82,11 @@ class accessor_property_list : protected sycl::detail::PropertyListBase {
             auto... Args>
   struct ContainsPropertyInstance
       : std::conditional_t<
-            AreSameTemplate<PropT<Args...>, typename ContainerT::Head>::value ||
-                !detail::IsCxPropertyInstance<typename ContainerT::Head>::value,
+            !std::is_same_v<typename ContainerT::Head, void> &&
+                (AreSameTemplate<PropT<Args...>,
+                                 typename ContainerT::Head>::value ||
+                 !sycl::detail::IsCxPropertyInstance<
+                     typename ContainerT::Head>::value),
             std::true_type,
             ContainsPropertyInstance<typename ContainerT::Rest, PropT,
                                      Args...>> {};
@@ -126,25 +129,29 @@ class accessor_property_list : protected sycl::detail::PropertyListBase {
   template <typename T, typename... Tail>
   struct AllProperties<T, Tail...>
       : std::conditional<
-            std::is_base_of<detail::DataLessPropertyBase, T>::value ||
-                std::is_base_of<detail::PropertyWithDataBase, T>::value ||
-                detail::IsCxPropertyInstance<T>::value,
+            std::is_base_of<sycl::detail::DataLessPropertyBase, T>::value ||
+                std::is_base_of<sycl::detail::PropertyWithDataBase, T>::value ||
+                sycl::detail::IsCxPropertyInstance<T>::value,
             AllProperties<Tail...>, std::false_type>::type {};
 
   accessor_property_list(
-      std::bitset<detail::DataLessPropKind::DataLessPropKindSize> DataLessProps,
-      std::vector<std::shared_ptr<detail::PropertyWithDataBase>> PropsWithData)
-      : detail::PropertyListBase(DataLessProps, PropsWithData) {}
+      std::bitset<sycl::detail::DataLessPropKind::DataLessPropKindSize>
+          DataLessProps,
+      std::vector<std::shared_ptr<sycl::detail::PropertyWithDataBase>>
+          PropsWithData)
+      : sycl::detail::PropertyListBase(DataLessProps, PropsWithData) {}
 
 public:
   template <
       typename = typename std::enable_if<AllProperties<PropsT...>::value>::type>
-  accessor_property_list(PropsT... Props) : detail::PropertyListBase(false) {
+  accessor_property_list(PropsT... Props)
+      : sycl::detail::PropertyListBase(false) {
     ctorHelper(Props...);
   }
 
   accessor_property_list(const sycl::property_list &Props)
-      : detail::PropertyListBase(Props.MDataLessProps, Props.MPropsWithData) {}
+      : sycl::detail::PropertyListBase(Props.MDataLessProps,
+                                       Props.MPropsWithData) {}
 
   template <typename... OtherProps,
             typename = typename std::enable_if<
@@ -153,8 +160,8 @@ class accessor_property_list : protected sycl::detail::PropertyListBase {
                 ContainsSameProperties<PropertyContainer<OtherProps...>,
                                        PropsT...>::value>::type>
   accessor_property_list(const accessor_property_list<OtherProps...> &OtherList)
-      : detail::PropertyListBase(OtherList.MDataLessProps,
-                                 OtherList.MPropsWithData) {}
+      : sycl::detail::PropertyListBase(OtherList.MDataLessProps,
+                                       OtherList.MPropsWithData) {}
 
   template <typename PropT, typename = typename std::enable_if<
                                 !is_compile_time_property<PropT>::value>::type>
@@ -181,8 +188,9 @@ class accessor_property_list : protected sycl::detail::PropertyListBase {
                                     T::template instance>::value;
   }
 
-  template <typename T, typename = typename std::enable_if_t<
-                            is_compile_time_property<T>::value>>
+  template <typename T,
+            typename = typename std::enable_if_t<
+                is_compile_time_property<T>::value && has_property<T>()>>
   static constexpr auto get_property() {
     return typename GetCxPropertyHelper<PropertyContainer<PropsT...>,
                                         T::template instance>::type{};
diff --git a/sycl/test/basic_tests/accessor/accessor_property_list.cpp b/sycl/test/basic_tests/accessor/accessor_property_list.cpp
index 4bb43681f4d4a..d9f714ce89246 100644
--- a/sycl/test/basic_tests/accessor/accessor_property_list.cpp
+++ b/sycl/test/basic_tests/accessor/accessor_property_list.cpp
@@ -16,6 +16,13 @@ int main() {
     accessor_property_list PL;
   }
 
+  {
+    // No properties
+    accessor_property_list PL{};
+    static_assert(!PL.has_property<property::no_alias>(), "Property is found");
+    static_assert(!PL.has_property<property::no_offset>(), "Property is found");
+  }
+
   {
     // Single CT property
     accessor_property_list PL{no_alias};

From c22997b2b51c583101279d7534eaeb7cf9edd30a Mon Sep 17 00:00:00 2001
From: Alexander Batashev <alexander.batashev@intel.com>
Date: Mon, 14 Sep 2020 11:47:41 +0300
Subject: [PATCH 13/19] Fix include

---
 sycl/include/CL/sycl/image.hpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/sycl/include/CL/sycl/image.hpp b/sycl/include/CL/sycl/image.hpp
index bd8216a2f5dd1..cf47c5d093778 100644
--- a/sycl/include/CL/sycl/image.hpp
+++ b/sycl/include/CL/sycl/image.hpp
@@ -8,7 +8,7 @@
 
 #pragma once
 
-#include "CL/sycl/ONEAPI/accessor_property_list.hpp"
+#include <CL/sycl/ONEAPI/accessor_property_list.hpp>
 #include <CL/sycl/detail/common.hpp>
 #include <CL/sycl/detail/generic_type_traits.hpp>
 #include <CL/sycl/detail/image_impl.hpp>

From c4cebbabe483018fcd7aa790eb3718af7340741e Mon Sep 17 00:00:00 2001
From: Alexander Batashev <alexbatashev@outlook.com>
Date: Mon, 14 Sep 2020 21:30:08 +0300
Subject: [PATCH 14/19] Fix small typo

Co-authored-by: Mikhail Lychkov <51128024+mlychkov@users.noreply.github.com>
---
 sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp b/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
index 6257e28bb6a43..b24944a48fcf8 100644
--- a/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
+++ b/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
@@ -18,7 +18,7 @@
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
 namespace detail {
-// This helper template must be specialized for nestend instance template
+// This helper template must be specialized for nested instance template
 // of each compile-time-constant property.
 template <typename T> struct IsCxPropertyInstance : std::false_type {};
 } // namespace detail

From 0091796669e99e68a25fc1f9bfc5a77ba28b5840 Mon Sep 17 00:00:00 2001
From: Alexander Batashev <alexander.batashev@intel.com>
Date: Tue, 15 Sep 2020 10:43:46 +0300
Subject: [PATCH 15/19] Split CT and RT tests

---
 ...list.cpp => accessor_property_list_ct.cpp} | 17 -------------
 .../accessor/accessor_property_list_rt.cpp    | 24 +++++++++++++++++++
 2 files changed, 24 insertions(+), 17 deletions(-)
 rename sycl/test/basic_tests/accessor/{accessor_property_list.cpp => accessor_property_list_ct.cpp} (85%)
 create mode 100644 sycl/test/basic_tests/accessor/accessor_property_list_rt.cpp

diff --git a/sycl/test/basic_tests/accessor/accessor_property_list.cpp b/sycl/test/basic_tests/accessor/accessor_property_list_ct.cpp
similarity index 85%
rename from sycl/test/basic_tests/accessor/accessor_property_list.cpp
rename to sycl/test/basic_tests/accessor/accessor_property_list_ct.cpp
index d9f714ce89246..5a40a4da476e9 100644
--- a/sycl/test/basic_tests/accessor/accessor_property_list.cpp
+++ b/sycl/test/basic_tests/accessor/accessor_property_list_ct.cpp
@@ -1,8 +1,4 @@
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
-// RUN: env SYCL_DEVICE_TYPE=HOST %t.out
-// RUN: %CPU_RUN_PLACEHOLDER %t.out
-// RUN: %GPU_RUN_PLACEHOLDER %t.out
-// RUN: %ACC_RUN_PLACEHOLDER %t.out
 
 #include <CL/sycl.hpp>
 
@@ -30,19 +26,6 @@ int main() {
     static_assert(!PL.has_property<property::no_offset>(), "Property is found");
   }
 
-  {
-    // Single RT property
-    accessor_property_list PL{sycl::noinit};
-    assert(PL.has_property<sycl::property::noinit>() && "Property not found");
-  }
-
-  {
-    // Compile time and runtime properties
-    accessor_property_list PL{sycl::noinit, no_alias};
-    assert(PL.has_property<property::no_alias>() && "Property not found");
-    assert(PL.has_property<sycl::property::noinit>() && "Property not found");
-  }
-
   {
     // Multiple CT properties
     accessor_property_list PL{no_alias, no_offset};
diff --git a/sycl/test/basic_tests/accessor/accessor_property_list_rt.cpp b/sycl/test/basic_tests/accessor/accessor_property_list_rt.cpp
new file mode 100644
index 0000000000000..9d53d082cc43a
--- /dev/null
+++ b/sycl/test/basic_tests/accessor/accessor_property_list_rt.cpp
@@ -0,0 +1,24 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: env SYCL_DEVICE_TYPE=HOST %t.out
+
+#include <CL/sycl.hpp>
+
+#include <type_traits>
+
+using namespace sycl::ONEAPI;
+
+int main() {
+  {
+    // Single RT property
+    accessor_property_list PL{sycl::noinit};
+    assert(PL.has_property<sycl::property::noinit>() && "Property not found");
+  }
+
+  {
+    // Compile time and runtime properties
+    accessor_property_list PL{sycl::noinit, no_alias};
+    assert(PL.has_property<property::no_alias>() && "Property not found");
+    assert(PL.has_property<sycl::property::noinit>() && "Property not found");
+  }
+  return 0;
+}

From d70980cbdf4d1db88b0c2e23f10e517abbb2d1eb Mon Sep 17 00:00:00 2001
From: Alexander Batashev <alexander.batashev@intel.com>
Date: Tue, 15 Sep 2020 11:48:06 +0300
Subject: [PATCH 16/19] Attempt to improve code readability

---
 .../CL/sycl/ONEAPI/accessor_property_list.hpp | 65 ++++++++++++-------
 sycl/include/CL/sycl/accessor.hpp             |  5 +-
 .../sycl/properties/accessor_properties.hpp   |  8 +--
 3 files changed, 49 insertions(+), 29 deletions(-)

diff --git a/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp b/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
index b24944a48fcf8..bc7a078b7ae7c 100644
--- a/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
+++ b/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
@@ -17,10 +17,14 @@
 
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
+// Forward declaration
+template <typename, int, access::mode, access::target, access::placeholder,
+          typename PropertyListT>
+class accessor;
 namespace detail {
 // This helper template must be specialized for nested instance template
 // of each compile-time-constant property.
-template <typename T> struct IsCxPropertyInstance : std::false_type {};
+template <typename T> struct IsCompileTimePropertyInstance : std::false_type {};
 } // namespace detail
 namespace ONEAPI {
 
@@ -44,6 +48,9 @@ class accessor_property_list : protected sycl::detail::PropertyListBase {
   template <template <class...> class T, class T1, class T2>
   struct AreSameTemplate<T<T1>, T<T2>> : std::true_type {};
 #if __cplusplus >= 201703L
+  // Declaring non-type template parameters with auto is a C++17 feature. Since
+  // the extension is written against SYCL 2020, which implies use of C++17,
+  // there's no need to provide alternative implementations for older standards.
   template <template <auto...> class T, auto... T1, auto... T2>
   struct AreSameTemplate<T<T1...>, T<T2...>> : std::true_type {};
 #endif
@@ -54,13 +61,13 @@ class accessor_property_list : protected sycl::detail::PropertyListBase {
   template <typename PropT> struct ContainsProperty<PropT> : std::false_type {};
   template <typename PropT, typename Head, typename... Tail>
   struct ContainsProperty<PropT, Head, Tail...>
-      : std::conditional<AreSameTemplate<PropT, Head>::value ||
-                             !sycl::detail::IsCxPropertyInstance<PropT>::value,
-                         std::true_type,
-                         ContainsProperty<PropT, Tail...>>::type {};
+      : std::conditional<
+            AreSameTemplate<PropT, Head>::value ||
+                !sycl::detail::IsCompileTimePropertyInstance<PropT>::value,
+            std::true_type, ContainsProperty<PropT, Tail...>>::type {};
 
-  // The following structures help to check if two property lists contain the
-  // same compile-time-constant properties.
+  // PropertyContainer is a helper structure, that holds list of properties.
+  // It is used to avoid multiple parameter packs in templates.
   template <typename...> struct PropertyContainer {
     using Head = void;
     using Rest = void;
@@ -85,7 +92,7 @@ class accessor_property_list : protected sycl::detail::PropertyListBase {
             !std::is_same_v<typename ContainerT::Head, void> &&
                 (AreSameTemplate<PropT<Args...>,
                                  typename ContainerT::Head>::value ||
-                 !sycl::detail::IsCxPropertyInstance<
+                 !sycl::detail::IsCompileTimePropertyInstance<
                      typename ContainerT::Head>::value),
             std::true_type,
             ContainsPropertyInstance<typename ContainerT::Rest, PropT,
@@ -108,30 +115,36 @@ class accessor_property_list : protected sycl::detail::PropertyListBase {
 
 #if __cplusplus >= 201703L
   // This template helps to extract exact property instance type based on
-  // template template argument.
+  // template template argument. If there's an instance of target property in
+  // ContainerT, find instance template and use it as type. Otherwise, just
+  // use void as return type.
   template <typename ContainerT, template <auto...> class PropT, auto... Args>
-  struct GetCxPropertyHelper {
+  struct GetCompileTimePropertyHelper {
     using type = typename std::conditional_t<
         AreSameTemplate<typename ContainerT::Head, PropT<Args...>>::value,
         typename ContainerT::Head,
-        typename GetCxPropertyHelper<typename ContainerT::Rest, PropT,
-                                     Args...>::type>;
+        typename GetCompileTimePropertyHelper<typename ContainerT::Rest, PropT,
+                                              Args...>::type>;
   };
   template <typename Head, template <auto...> class PropT, auto... Args>
-  struct GetCxPropertyHelper<PropertyContainer<Head>, PropT, Args...> {
+  struct GetCompileTimePropertyHelper<PropertyContainer<Head>, PropT, Args...> {
     using type = typename std::conditional_t<
         AreSameTemplate<Head, PropT<Args...>>::value, Head, void>;
   };
 #endif
 
-  // The structs validate that all objects passed are SYCL properties
+  // The structs validate that all objects passed are SYCL properties.
+  // Properties are either run time SYCL 1.2.1 properties, and thus derive from
+  // either DataLessPropertyBase or from PropertyWithDataBase, or
+  // compile-time-constant properties, and thus specialize
+  // IsCompileTimePropertyInstance template.
   template <typename... Tail> struct AllProperties : std::true_type {};
   template <typename T, typename... Tail>
   struct AllProperties<T, Tail...>
       : std::conditional<
             std::is_base_of<sycl::detail::DataLessPropertyBase, T>::value ||
                 std::is_base_of<sycl::detail::PropertyWithDataBase, T>::value ||
-                sycl::detail::IsCxPropertyInstance<T>::value,
+                sycl::detail::IsCompileTimePropertyInstance<T>::value,
             AllProperties<Tail...>, std::false_type>::type {};
 
   accessor_property_list(
@@ -192,21 +205,27 @@ class accessor_property_list : protected sycl::detail::PropertyListBase {
             typename = typename std::enable_if_t<
                 is_compile_time_property<T>::value && has_property<T>()>>
   static constexpr auto get_property() {
-    return typename GetCxPropertyHelper<PropertyContainer<PropsT...>,
-                                        T::template instance>::type{};
+    return typename GetCompileTimePropertyHelper<PropertyContainer<PropsT...>,
+                                                 T::template instance>::type{};
   }
 #endif
 
-  template <typename... OtherPropsT>
-  static constexpr bool areSameCxProperties() {
-    return ContainsSameProperties<PropertyContainer<OtherPropsT...>,
-                                  PropsT...>::value;
-  }
-
 private:
+  template <typename, int, access::mode, access::target, access::placeholder,
+            typename PropertyListT>
+  friend class sycl::accessor;
+
   template <typename... OtherProps> friend class accessor_property_list;
 
   friend class sycl::property_list;
+
+  // Helper method, used by accessor to restrict conversions to compatible
+  // property lists.
+  template <typename... OtherPropsT>
+  static constexpr bool areSameCompileTimeProperties() {
+    return ContainsSameProperties<PropertyContainer<OtherPropsT...>,
+                                  PropsT...>::value;
+  }
 };
 } // namespace ONEAPI
 } // namespace sycl
diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp
index d59326391e0ad..d5af074ed772a 100755
--- a/sycl/include/CL/sycl/accessor.hpp
+++ b/sycl/include/CL/sycl/accessor.hpp
@@ -1488,8 +1488,9 @@ class accessor :
   {
     static_assert(detail::IsCxPropertyList<PropertyListT>::value,
                   "Conversion is only available for accessor_property_list");
-    static_assert(PropertyListT::template areSameCxProperties<NewPropsT...>(),
-                  "Compile-time-constant properties must be the same");
+    static_assert(
+        PropertyListT::template areSameCompileTimeProperties<NewPropsT...>(),
+        "Compile-time-constant properties must be the same");
   }
 
   constexpr bool is_placeholder() const { return IsPlaceH; }
diff --git a/sycl/include/CL/sycl/properties/accessor_properties.hpp b/sycl/include/CL/sycl/properties/accessor_properties.hpp
index 2c9cd04dcce3d..8471f23a7c6be 100644
--- a/sycl/include/CL/sycl/properties/accessor_properties.hpp
+++ b/sycl/include/CL/sycl/properties/accessor_properties.hpp
@@ -103,13 +103,13 @@ struct is_compile_time_property<INTEL::property::buffer_location>
 } // namespace ONEAPI
 namespace detail {
 template <int I>
-struct IsCxPropertyInstance<INTEL::property::buffer_location::instance<I>>
-    : std::true_type {};
+struct IsCompileTimePropertyInstance<
+    INTEL::property::buffer_location::instance<I>> : std::true_type {};
 template <>
-struct IsCxPropertyInstance<ONEAPI::property::no_alias::instance<>>
+struct IsCompileTimePropertyInstance<ONEAPI::property::no_alias::instance<>>
     : std::true_type {};
 template <>
-struct IsCxPropertyInstance<ONEAPI::property::no_offset::instance<>>
+struct IsCompileTimePropertyInstance<ONEAPI::property::no_offset::instance<>>
     : std::true_type {};
 } // namespace detail
 } // namespace sycl

From 2c048818e273e210a0040dcee73892980c31290a Mon Sep 17 00:00:00 2001
From: Alexander Batashev <alexbatashev@outlook.com>
Date: Tue, 15 Sep 2020 14:25:49 +0300
Subject: [PATCH 17/19] Apply suggestions from code review

Co-authored-by: Mikhail Lychkov <51128024+mlychkov@users.noreply.github.com>
---
 sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp | 2 +-
 sycl/include/CL/sycl/detail/property_list_base.hpp     | 2 +-
 2 files changed, 2 insertions(+), 2 deletions(-)

diff --git a/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp b/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
index bc7a078b7ae7c..a05443bd9fbce 100644
--- a/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
+++ b/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
@@ -102,7 +102,7 @@ class accessor_property_list : protected sycl::detail::PropertyListBase {
   struct ContainsPropertyInstance<void, PropT, Args...> : std::false_type {};
 #endif
 
-  // This template checks if to lists of properties contain the same set of
+  // This template checks if two lists of properties contain the same set of
   // compile-time-constant properties in any order.
   template <typename ContainerT, typename... OtherProps>
   struct ContainsSameProperties
diff --git a/sycl/include/CL/sycl/detail/property_list_base.hpp b/sycl/include/CL/sycl/detail/property_list_base.hpp
index 8fc13e50dd0d0..4b048e9e8a017 100644
--- a/sycl/include/CL/sycl/detail/property_list_base.hpp
+++ b/sycl/include/CL/sycl/detail/property_list_base.hpp
@@ -103,7 +103,7 @@ class PropertyListBase {
                                      PI_INVALID_VALUE);
   }
 
-  // Stores enable/not enabled for simple properties
+  // Stores enabled/disabled for simple properties
   std::bitset<DataLessPropKind::DataLessPropKindSize> MDataLessProps;
   // Stores shared_ptrs to complex properties
   std::vector<std::shared_ptr<PropertyWithDataBase>> MPropsWithData;

From ed8efd7d55df4ebf0b7c214410be6babb3a8ff9a Mon Sep 17 00:00:00 2001
From: Alexander Batashev <alexander.batashev@intel.com>
Date: Wed, 16 Sep 2020 10:12:42 +0300
Subject: [PATCH 18/19] Address more comments

---
 .../CL/sycl/ONEAPI/accessor_property_list.hpp | 21 ++++-----
 .../sycl/properties/accessor_properties.hpp   | 10 +---
 .../accessor/accessor_property_list_rt.cpp    |  1 +
 sycl/test/fpga_tests/buffer_location.cpp      | 22 +++++++++
 .../fpga_tests/buffer_location_codegen.cpp    | 47 +++++++++++++++++++
 5 files changed, 82 insertions(+), 19 deletions(-)
 create mode 100644 sycl/test/fpga_tests/buffer_location.cpp
 create mode 100644 sycl/test/fpga_tests/buffer_location_codegen.cpp

diff --git a/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp b/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
index a05443bd9fbce..73f2af501badc 100644
--- a/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
+++ b/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
@@ -61,10 +61,8 @@ class accessor_property_list : protected sycl::detail::PropertyListBase {
   template <typename PropT> struct ContainsProperty<PropT> : std::false_type {};
   template <typename PropT, typename Head, typename... Tail>
   struct ContainsProperty<PropT, Head, Tail...>
-      : std::conditional<
-            AreSameTemplate<PropT, Head>::value ||
-                !sycl::detail::IsCompileTimePropertyInstance<PropT>::value,
-            std::true_type, ContainsProperty<PropT, Tail...>>::type {};
+      : std::conditional<AreSameTemplate<PropT, Head>::value, std::true_type,
+                         ContainsProperty<PropT, Tail...>>::type {};
 
   // PropertyContainer is a helper structure, that holds list of properties.
   // It is used to avoid multiple parameter packs in templates.
@@ -90,10 +88,8 @@ class accessor_property_list : protected sycl::detail::PropertyListBase {
   struct ContainsPropertyInstance
       : std::conditional_t<
             !std::is_same_v<typename ContainerT::Head, void> &&
-                (AreSameTemplate<PropT<Args...>,
-                                 typename ContainerT::Head>::value ||
-                 !sycl::detail::IsCompileTimePropertyInstance<
-                     typename ContainerT::Head>::value),
+                AreSameTemplate<PropT<Args...>,
+                                typename ContainerT::Head>::value,
             std::true_type,
             ContainsPropertyInstance<typename ContainerT::Rest, PropT,
                                      Args...>> {};
@@ -103,11 +99,14 @@ class accessor_property_list : protected sycl::detail::PropertyListBase {
 #endif
 
   // This template checks if two lists of properties contain the same set of
-  // compile-time-constant properties in any order.
-  template <typename ContainerT, typename... OtherProps>
+  // compile-time-constant properties in any order. Run time properties are
+  // skipped.
   struct ContainsSameProperties
       : std::conditional<
-            ContainsProperty<typename ContainerT::Head, OtherProps...>::value,
+            !detail::IsCompileTimePropertyInstance<
+                typename ContainerT::Head>::value ||
+                ContainsProperty<typename ContainerT::Head,
+                                 OtherProps...>::value,
             ContainsSameProperties<typename ContainerT::Rest, OtherProps...>,
             std::false_type>::type {};
   template <typename... OtherProps>
diff --git a/sycl/include/CL/sycl/properties/accessor_properties.hpp b/sycl/include/CL/sycl/properties/accessor_properties.hpp
index 8471f23a7c6be..45e5075b45ae4 100644
--- a/sycl/include/CL/sycl/properties/accessor_properties.hpp
+++ b/sycl/include/CL/sycl/properties/accessor_properties.hpp
@@ -39,19 +39,13 @@ namespace INTEL {
 namespace property {
 struct buffer_location {
   template <int A> struct instance {
-    constexpr bool operator==(const buffer_location::instance<A> &) const {
-      return true;
-    }
-    constexpr bool operator!=(const buffer_location::instance<A> &) const {
-      return false;
-    }
     template <int B>
     constexpr bool operator==(const buffer_location::instance<B> &) const {
-      return false;
+      return A == B;
     }
     template <int B>
     constexpr bool operator!=(const buffer_location::instance<B> &) const {
-      return true;
+      return A != B;
     }
   };
 };
diff --git a/sycl/test/basic_tests/accessor/accessor_property_list_rt.cpp b/sycl/test/basic_tests/accessor/accessor_property_list_rt.cpp
index 9d53d082cc43a..b794f6efdf158 100644
--- a/sycl/test/basic_tests/accessor/accessor_property_list_rt.cpp
+++ b/sycl/test/basic_tests/accessor/accessor_property_list_rt.cpp
@@ -11,6 +11,7 @@ int main() {
   {
     // Single RT property
     accessor_property_list PL{sycl::noinit};
+    static_assert(!PL.has_property<property::no_offset>(), "Property is found");
     assert(PL.has_property<sycl::property::noinit>() && "Property not found");
   }
 
diff --git a/sycl/test/fpga_tests/buffer_location.cpp b/sycl/test/fpga_tests/buffer_location.cpp
new file mode 100644
index 0000000000000..611b51d52d8bb
--- /dev/null
+++ b/sycl/test/fpga_tests/buffer_location.cpp
@@ -0,0 +1,22 @@
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUNx: %ACC_RUN_PLACEHOLDER %t.out
+
+#include <CL/sycl.hpp>
+
+int main() {
+  sycl::queue Queue;
+  sycl::buffer<int, 1> Buf{sycl::range{1}};
+
+  Queue.submit([&](sycl::handler &CGH) {
+    sycl::ONEAPI::accessor_property_list PL{sycl::INTEL::buffer_location<1>};
+    sycl::accessor Acc(Buf, CGH, sycl::write_only, PL);
+    CGH.single_task<class Test>([=]() { Acc[0] = 42; });
+  });
+
+  Queue.wait();
+
+  auto Acc = Buf.template get_access<sycl::access::mode::read_write>();
+  assert(Acc[0] == 42 && "Value mismatch");
+
+  return 0;
+}
diff --git a/sycl/test/fpga_tests/buffer_location_codegen.cpp b/sycl/test/fpga_tests/buffer_location_codegen.cpp
new file mode 100644
index 0000000000000..2edff0b8fcefd
--- /dev/null
+++ b/sycl/test/fpga_tests/buffer_location_codegen.cpp
@@ -0,0 +1,47 @@
+// RUN: %clangxx -fsycl -c -fsycl-device-only -S -emit-llvm %s -o - | FileCheck %s
+
+// CHECK: define {{.*}}spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE15kernel_function"{{.*}} !kernel_arg_buffer_location ![[MDBL:[0-9]+]]
+// CHECK: ![[MDBL]] = !{i32 3, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 2, i32 -1, i32 -1, i32 -1, i32 2, i32 -1, i32 -1, i32 -1, i32 -1}
+
+#include <CL/sycl.hpp>
+
+struct Base {
+  int A, B;
+  cl::sycl::accessor<
+      char, 1, cl::sycl::access::mode::read,
+      cl::sycl::access::target::global_buffer,
+      cl::sycl::access::placeholder::false_t,
+      cl::sycl::ONEAPI::accessor_property_list<
+          cl::sycl::INTEL::property::buffer_location::instance<2>>>
+      AccField;
+};
+
+struct Captured
+    : Base,
+      cl::sycl::accessor<
+          char, 1, cl::sycl::access::mode::read,
+          cl::sycl::access::target::global_buffer,
+          cl::sycl::access::placeholder::false_t,
+          cl::sycl::ONEAPI::accessor_property_list<
+              cl::sycl::INTEL::property::buffer_location::instance<2>>> {
+  int C;
+};
+
+int main() {
+  Captured Obj;
+  cl::sycl::accessor<
+      int, 1, cl::sycl::access::mode::read_write,
+      cl::sycl::access::target::global_buffer,
+      cl::sycl::access::placeholder::false_t,
+      cl::sycl::ONEAPI::accessor_property_list<
+          cl::sycl::INTEL::property::buffer_location::instance<3>>>
+      accessorA;
+  sycl::queue Queue;
+  Queue.submit([&](sycl::handler &CGH) {
+    CGH.single_task<class kernel_function>([=]() {
+      (int)accessorA[0];
+      (int)Obj[0];
+    });
+  });
+  return 0;
+}

From e982a29d53b1e2a891321c2582a50188ac1c09ad Mon Sep 17 00:00:00 2001
From: Alexander Batashev <alexander.batashev@intel.com>
Date: Wed, 16 Sep 2020 11:22:35 +0300
Subject: [PATCH 19/19] Fix merge artifacts

---
 sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp | 1 +
 1 file changed, 1 insertion(+)

diff --git a/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp b/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
index 73f2af501badc..0e96138d6a08a 100644
--- a/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
+++ b/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
@@ -101,6 +101,7 @@ class accessor_property_list : protected sycl::detail::PropertyListBase {
   // This template checks if two lists of properties contain the same set of
   // compile-time-constant properties in any order. Run time properties are
   // skipped.
+  template <typename ContainerT, typename... OtherProps>
   struct ContainsSameProperties
       : std::conditional<
             !detail::IsCompileTimePropertyInstance<