[SYCL] Pass buffer_location property to buffer

sycl/include/CL/sycl/accessor.hpp

@@ -1195,7 +1195,9 @@ class __SYCL_SPECIAL_CLASS accessor :
       buffer<T, Dims, AllocatorT> &BufferRef, TagT,
       const property_list &PropertyList = {},
       const detail::code_location CodeLoc = detail::code_location::current())
-      : accessor(BufferRef, PropertyList, CodeLoc) {}
+      : accessor(BufferRef, PropertyList, CodeLoc) {
+    adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
+  }
 
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
             typename TagT, typename... PropTypes,
@@ -1208,7 +1210,9 @@ class __SYCL_SPECIAL_CLASS accessor :
       const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
           {},
       const detail::code_location CodeLoc = detail::code_location::current())
-      : accessor(BufferRef, PropertyList, CodeLoc) {}
+      : accessor(BufferRef, PropertyList, CodeLoc) {
+    adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
+  }
 #endif
 
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
@@ -1285,7 +1289,9 @@ class __SYCL_SPECIAL_CLASS accessor :
       buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
       TagT, const property_list &PropertyList = {},
       const detail::code_location CodeLoc = detail::code_location::current())
-      : accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {}
+      : accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {
+    adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
+  }
 
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
             typename TagT, typename... PropTypes,
@@ -1299,7 +1305,9 @@ class __SYCL_SPECIAL_CLASS accessor :
       const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
           {},
       const detail::code_location CodeLoc = detail::code_location::current())
-      : accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {}
+      : accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {
+    adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
+  }
 
 #endif
 
@@ -1341,7 +1349,9 @@ class __SYCL_SPECIAL_CLASS accessor :
       buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
       TagT, const property_list &PropertyList = {},
       const detail::code_location CodeLoc = detail::code_location::current())
-      : accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
+      : accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {
+    adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
+  }
 
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
             typename TagT, typename... PropTypes,
@@ -1355,7 +1365,9 @@ class __SYCL_SPECIAL_CLASS accessor :
       const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
           {},
       const detail::code_location CodeLoc = detail::code_location::current())
-      : accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
+      : accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {
+    adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
+  }
 #endif
 
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
@@ -1399,7 +1411,9 @@ class __SYCL_SPECIAL_CLASS accessor :
       const property_list &PropertyList = {},
       const detail::code_location CodeLoc = detail::code_location::current())
       : accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
-                 CodeLoc) {}
+                 CodeLoc) {
+    adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
+  }
 
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
             typename TagT, typename... PropTypes,
@@ -1414,7 +1428,9 @@ class __SYCL_SPECIAL_CLASS accessor :
           {},
       const detail::code_location CodeLoc = detail::code_location::current())
       : accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
-                 CodeLoc) {}
+                 CodeLoc) {
+    adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
+  }
 #endif
 
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
@@ -1508,7 +1524,9 @@ class __SYCL_SPECIAL_CLASS accessor :
       buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
       id<Dimensions> AccessOffset, TagT, const property_list &PropertyList = {},
       const detail::code_location CodeLoc = detail::code_location::current())
-      : accessor(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {}
+      : accessor(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {
+    adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
+  }
 
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
             typename TagT, typename... PropTypes,
@@ -1522,7 +1540,9 @@ class __SYCL_SPECIAL_CLASS accessor :
       const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
           {},
       const detail::code_location CodeLoc = detail::code_location::current())
-      : accessor(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {}
+      : accessor(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {
+    adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
+  }
 #endif
 
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
@@ -1617,7 +1637,9 @@ class __SYCL_SPECIAL_CLASS accessor :
       const property_list &PropertyList = {},
       const detail::code_location CodeLoc = detail::code_location::current())
       : accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
-                 PropertyList, CodeLoc) {}
+                 PropertyList, CodeLoc) {
+    adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
+  }
 
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
             typename TagT, typename... PropTypes,
@@ -1632,7 +1654,9 @@ class __SYCL_SPECIAL_CLASS accessor :
           {},
       const detail::code_location CodeLoc = detail::code_location::current())
       : accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
-                 PropertyList, CodeLoc) {}
+                 PropertyList, CodeLoc) {
+    adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
+  }
 #endif
 
   template <typename... NewPropsT>
@@ -1802,6 +1826,32 @@ class __SYCL_SPECIAL_CLASS accessor :
           PI_INVALID_VALUE);
     }
   }
+
+#if __cplusplus >= 201703L
+  template <typename... PropTypes>
+  void adjustAccPropsInBuf(detail::SYCLMemObjI *SYCLMemObject) {
+    if constexpr (PropertyListT::template has_property<
+                      sycl::ext::intel::property::buffer_location>()) {
+      auto location = (PropertyListT::template get_property<
+                           sycl::ext::intel::property::buffer_location>())
+                          .get_location();
+      property_list PropList{
+          sycl::property::buffer::detail::buffer_location(location)};
+      detail::SYCLMemObjT *SYCLMemObjectT =
+          dynamic_cast<detail::SYCLMemObjT *>(SYCLMemObject);
+      SYCLMemObjectT->addOrReplaceAccessorProperties(PropList);
+    } else {
+      deleteAccPropsFromBuf(SYCLMemObject);
+    }
+  }
+
+  void deleteAccPropsFromBuf(detail::SYCLMemObjI *SYCLMemObject) {
+    detail::SYCLMemObjT *SYCLMemObjectT =
+        dynamic_cast<detail::SYCLMemObjT *>(SYCLMemObject);
+    SYCLMemObjectT->deleteAccessorProperty(
+        sycl::detail::PropWithDataKind::AccPropBufferLocation);
+  }
+#endif
 };
 
 #if __cplusplus >= 201703L

sycl/include/CL/sycl/detail/pi.h

@@ -591,6 +591,8 @@ constexpr pi_map_flags PI_MAP_WRITE_INVALIDATE_REGION =
 // make the translation to OpenCL transparent.
 using pi_mem_properties = pi_bitfield;
 constexpr pi_mem_properties PI_MEM_PROPERTIES_CHANNEL = CL_MEM_CHANNEL_INTEL;
+constexpr pi_mem_properties PI_MEM_PROPERTIES_ALLOC_BUFFER_LOCATION =
+    CL_MEM_ALLOC_BUFFER_LOCATION_INTEL;
 
 // NOTE: this is made 64-bit to match the size of cl_mem_properties_intel to
 // make the translation to OpenCL transparent.

sycl/include/CL/sycl/detail/property_helper.hpp

@@ -48,7 +48,7 @@ enum PropWithDataKind {
   ImageContextBound = 3,
   BufferMemChannel = 4,
   AccPropBufferLocation = 5,
-  PropWithDataKindSize = 6
+  PropWithDataKindSize = 6,
 };
 
 // Base class for dataless properties, needed to check that the type of an

sycl/include/CL/sycl/detail/property_list_base.hpp

@@ -103,6 +103,30 @@ class PropertyListBase {
                                      PI_INVALID_VALUE);
   }
 
+  void add_or_replace_accessor_properties_helper(
+      const std::vector<std::shared_ptr<PropertyWithDataBase>> &PropsWithData) {
+    for (auto &Prop : PropsWithData) {
+      if (Prop->isSame(sycl::detail::PropWithDataKind::AccPropBufferLocation)) {
+        delete_accessor_property_helper(
+            sycl::detail::PropWithDataKind::AccPropBufferLocation);
+        MPropsWithData.push_back(Prop);
+        break;
+      }
+    }
+  }
+
+  void delete_accessor_property_helper(const PropWithDataKind &Kind) {
+    auto It = MPropsWithData.begin();
+    for (; It != MPropsWithData.end(); ++It) {
+      if ((*It)->isSame(Kind))
+        break;
+    }
+    if (It != MPropsWithData.end()) {
+      std::iter_swap(It, MPropsWithData.end() - 1);
+      MPropsWithData.pop_back();
+    }
+  }
+
   // Stores enabled/disabled for simple properties
   std::bitset<DataLessPropKind::DataLessPropKindSize> MDataLessProps;
   // Stores shared_ptrs to complex properties

sycl/include/CL/sycl/detail/sycl_mem_obj_t.hpp

@@ -111,6 +111,15 @@ class __SYCL_EXPORT SYCLMemObjT : public SYCLMemObjI {
     return MProps.get_property<propertyT>();
   }
 
+  __SYCL_DLL_LOCAL void
+  addOrReplaceAccessorProperties(const property_list &PropertyList) {
+    MProps.add_or_replace_accessor_properties(PropertyList);
+  }
+
+  __SYCL_DLL_LOCAL void deleteAccessorProperty(const PropWithDataKind &Kind) {
+    MProps.delete_accessor_property(Kind);
+  }
+
   template <typename AllocatorT>
   __SYCL_DLL_LOCAL AllocatorT get_allocator() const {
     return MAllocator->getAllocator<AllocatorT>();

sycl/include/CL/sycl/properties/accessor_properties.hpp

@@ -58,6 +58,7 @@ struct buffer_location {
     constexpr bool operator!=(const buffer_location::instance<B> &) const {
       return A != B;
     }
+    int get_location() { return A; }
   };
 };
 } // namespace property

sycl/include/CL/sycl/properties/buffer_properties.hpp

@@ -51,6 +51,18 @@ class mem_channel : public detail::PropertyWithData<
   uint32_t MChannel;
 };
 
+namespace detail {
+class buffer_location
+    : public sycl::detail::PropertyWithData<
+          sycl::detail::PropWithDataKind::AccPropBufferLocation> {
+public:
+  buffer_location(uint64_t Location) : MLocation(Location) {}
+  uint64_t get_buffer_location() const { return MLocation; }
+
+private:
+  uint64_t MLocation;
+};
+} // namespace detail
 } // namespace buffer
 } // namespace property
 
@@ -75,6 +87,9 @@ template <>
 struct is_property<property::buffer::use_host_ptr> : std::true_type {};
 template <> struct is_property<property::buffer::use_mutex> : std::true_type {};
 template <>
+struct is_property<property::buffer::detail::buffer_location> : std::true_type {
+};
+template <>
 struct is_property<property::buffer::context_bound> : std::true_type {};
 template <>
 struct is_property<property::buffer::mem_channel> : std::true_type {};
@@ -91,6 +106,10 @@ struct is_property_of<property::buffer::use_mutex,
                       buffer<T, Dimensions, AllocatorT, void>>
     : std::true_type {};
 template <typename T, int Dimensions, typename AllocatorT>
+struct is_property_of<property::buffer::detail::buffer_location,
+                      buffer<T, Dimensions, AllocatorT, void>>
+    : std::true_type {};
+template <typename T, int Dimensions, typename AllocatorT>
 struct is_property_of<property::buffer::context_bound,
                       buffer<T, Dimensions, AllocatorT, void>>
     : std::true_type {};

sycl/include/CL/sycl/property_list.hpp

@@ -51,6 +51,13 @@ class property_list : protected detail::PropertyListBase {
     return has_property_helper<PropT>();
   }
 
+  void add_or_replace_accessor_properties(const property_list &PropertyList) {
+    add_or_replace_accessor_properties_helper(PropertyList.MPropsWithData);
+  }
+  void delete_accessor_property(const sycl::detail::PropWithDataKind &Kind) {
+    delete_accessor_property_helper(Kind);
+  }
+
   template <typename... T> operator ext::oneapi::accessor_property_list<T...>();
 
 private:

sycl/source/detail/context_impl.cpp

@@ -28,15 +28,17 @@ namespace detail {
 context_impl::context_impl(const device &Device, async_handler AsyncHandler,
                            const property_list &PropList)
     : MAsyncHandler(AsyncHandler), MDevices(1, Device), MContext(nullptr),
-      MPlatform(), MPropList(PropList), MHostContext(Device.is_host()) {
+      MPlatform(), MPropList(PropList), MHostContext(Device.is_host()),
+      MSupportBufferLocationByDevices(NotChecked) {
   MKernelProgramCache.setContextPtr(this);
 }
 
 context_impl::context_impl(const std::vector<cl::sycl::device> Devices,
                            async_handler AsyncHandler,
                            const property_list &PropList)
     : MAsyncHandler(AsyncHandler), MDevices(Devices), MContext(nullptr),
-      MPlatform(), MPropList(PropList), MHostContext(false) {
+      MPlatform(), MPropList(PropList), MHostContext(false),
+      MSupportBufferLocationByDevices(NotChecked) {
   MPlatform = detail::getSyclObjImpl(MDevices[0].get_platform());
   std::vector<RT::PiDevice> DeviceIds;
   for (const auto &D : MDevices) {
@@ -66,7 +68,7 @@ context_impl::context_impl(const std::vector<cl::sycl::device> Devices,
 context_impl::context_impl(RT::PiContext PiContext, async_handler AsyncHandler,
                            const plugin &Plugin)
     : MAsyncHandler(AsyncHandler), MDevices(), MContext(PiContext), MPlatform(),
-      MHostContext(false) {
+      MHostContext(false), MSupportBufferLocationByDevices(NotChecked) {
 
   std::vector<RT::PiDevice> DeviceIds;
   size_t DevicesNum = 0;
@@ -206,6 +208,20 @@ pi_native_handle context_impl::getNative() const {
   return Handle;
 }
 
+bool context_impl::isBufferLocationSupported() const {
+  if (MSupportBufferLocationByDevices != NotChecked)
+    return MSupportBufferLocationByDevices == Supported ? true : false;
+  // Check that devices within context have support of buffer location
+  MSupportBufferLocationByDevices = Supported;
+  for (auto &Device : MDevices) {
+    if (!Device.has_extension("cl_intel_mem_alloc_buffer_location")) {
+      MSupportBufferLocationByDevices = NotSupported;
+      break;
+    }
+  }
+  return MSupportBufferLocationByDevices == Supported ? true : false;
+}
+
 } // namespace detail
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

sycl/source/detail/context_impl.hpp

@@ -167,6 +167,11 @@ class context_impl {
   /// \return a native handle.
   pi_native_handle getNative() const;
 
+  // Returns true if buffer_location property is supported by devices
+  bool isBufferLocationSupported() const;
+
+  enum PropertySupport { NotSupported = 0, Supported = 1, NotChecked = 2 };
+
 private:
   async_handler MAsyncHandler;
   std::vector<device> MDevices;
@@ -177,6 +182,7 @@ class context_impl {
   std::map<std::pair<DeviceLibExt, RT::PiDevice>, RT::PiProgram>
       MCachedLibPrograms;
   mutable KernelProgramCache MKernelProgramCache;
+  mutable PropertySupport MSupportBufferLocationByDevices;
 };
 
 } // namespace detail

sycl/source/detail/memory_manager.cpp

@@ -362,6 +362,18 @@ MemoryManager::allocateBufferObject(ContextImplPtr TargetContext, void *UserPtr,
 
   RT::PiMem NewMem = nullptr;
   const detail::plugin &Plugin = TargetContext->getPlugin();
+
+  if (PropsList.has_property<property::buffer::detail::buffer_location>())
+    if (TargetContext->isBufferLocationSupported()) {
+      auto location =
+          PropsList.get_property<property::buffer::detail::buffer_location>()
+              .get_buffer_location();
+      pi_mem_properties props[3] = {PI_MEM_PROPERTIES_ALLOC_BUFFER_LOCATION,
+                                    location, 0};
+      memBufferCreateHelper(Plugin, TargetContext->getHandleRef(),
+                            CreationFlags, Size, UserPtr, &NewMem, props);
+      return NewMem;
+    }
   memBufferCreateHelper(Plugin, TargetContext->getHandleRef(), CreationFlags,
                         Size, UserPtr, &NewMem, nullptr);
   return NewMem;

sycl/test/abi/sycl_symbols_windows.dump

@@ -1088,6 +1088,7 @@
 ?acospi@__host_std@cl@@YAMM@Z
 ?acospi@__host_std@cl@@YANN@Z
 ?addHostAccessorAndWait@detail@sycl@cl@@YAXPEAVAccessorImplHost@123@@Z
+?addOrReplaceAccessorProperties@SYCLMemObjT@detail@sycl@cl@@QEAAXAEBVproperty_list@34@@Z
 ?addReduction@handler@sycl@cl@@AEAAXAEBV?$shared_ptr@$$CBX@std@@@Z
 ?addStream@handler@sycl@cl@@AEAAXAEBV?$shared_ptr@Vstream_impl@detail@sycl@cl@@@std@@@Z
 ?advise_usm@MemoryManager@detail@sycl@cl@@SAXPEBXV?$shared_ptr@Vqueue_impl@detail@sycl@cl@@@std@@_KW4_pi_mem_advice@@V?$vector@PEAU_pi_event@@V?$allocator@PEAU_pi_event@@@std@@@6@AEAPEAU_pi_event@@@Z
@@ -1686,6 +1687,7 @@
 ?degrees@__host_std@cl@@YA?AVhalf@half_impl@detail@sycl@2@V34562@@Z
 ?degrees@__host_std@cl@@YAMM@Z
 ?degrees@__host_std@cl@@YANN@Z
+?deleteAccessorProperty@SYCLMemObjT@detail@sycl@cl@@QEAAXAEBW4PropWithDataKind@234@@Z
 ?depends_on@handler@sycl@cl@@QEAAXAEBV?$vector@Vevent@sycl@cl@@V?$allocator@Vevent@sycl@cl@@@std@@@std@@@Z
 ?depends_on@handler@sycl@cl@@QEAAXVevent@23@@Z
 ?destructorNotification@buffer_impl@detail@sycl@cl@@QEAAXPEAX@Z