[SYCL] USM Buffer Location Properties (#5634)

Lower level runtime's usm allocation API now supports pass in of buffer location property [1] defined in OpenCL spec [2]

For this feature to be accessible to users, sycl runtime now passes in usm::buffer_location property defined in sycl spec [3] into the opencl runtime calls, the property is only passed if the extension name occurs in the platform extension string. If the lower level runtime does not support such property, then the property will yield no effect.

This is a temporary solution to allow user to specify which memory location the device usm allocation should be in. The full solution will require retuning an annotated_ptr that carries compile time properties for further optimization. The full solution spec is in [4]

[1] https://github.com/intel/fpga-runtime-for-opencl/pull/46/files
[2] https://github.com/KhronosGroup/OpenCL-Docs/blob/master/extensions/cl_intel_mem_alloc_buffer_location.asciidoc
[3] #5665
[4] #5656

Author: sherry-yuan

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

@@ -592,6 +592,13 @@ constexpr pi_map_flags PI_MAP_WRITE_INVALIDATE_REGION =
 using pi_mem_properties = pi_bitfield;
 constexpr pi_mem_properties PI_MEM_PROPERTIES_CHANNEL = CL_MEM_CHANNEL_INTEL;
 
+// NOTE: this is made 64-bit to match the size of cl_mem_properties_intel to
+// make the translation to OpenCL transparent.
+using pi_usm_mem_properties = pi_bitfield;
+constexpr pi_usm_mem_properties PI_MEM_ALLOC_FLAGS = CL_MEM_ALLOC_FLAGS_INTEL;
+constexpr pi_usm_mem_properties PI_MEM_USM_ALLOC_BUFFER_LOCATION =
+    CL_MEM_ALLOC_BUFFER_LOCATION_INTEL;
+
 // NOTE: queue properties are implemented this way to better support bit
 // manipulations
 using pi_queue_properties = pi_bitfield;
@@ -1609,10 +1616,6 @@ typedef enum {
   PI_MEM_TYPE_SHARED = CL_MEM_TYPE_SHARED_INTEL
 } _pi_usm_type;
 
-typedef enum : pi_bitfield {
-  PI_MEM_ALLOC_FLAGS = CL_MEM_ALLOC_FLAGS_INTEL
-} _pi_usm_mem_properties;
-
 // Flag is used for piProgramUSMEnqueuePrefetch. PI_USM_MIGRATION_TBD0 is a
 // placeholder for future developments and should not change the behaviour of
 // piProgramUSMEnqueuePrefetch
@@ -1624,7 +1627,6 @@ using pi_usm_capability_query = _pi_usm_capability_query;
 using pi_usm_capabilities = _pi_usm_capabilities;
 using pi_mem_info = _pi_mem_info;
 using pi_usm_type = _pi_usm_type;
-using pi_usm_mem_properties = _pi_usm_mem_properties;
 using pi_usm_migration_flags = _pi_usm_migration_flags;
 
 /// Allocates host memory accessible by the device.

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

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

sycl/include/CL/sycl/feature_test.hpp.in

@@ -67,6 +67,7 @@ namespace sycl {
 #define SYCL_EXT_INTEL_KERNEL_ARGS_RESTRICT 1
 #define SYCL_EXT_INTEL_MEM_CHANNEL_PROPERTY 1
 #define SYCL_EXT_INTEL_USM_ADDRESS_SPACES 1
+#define SYCL_EXT_INTEL_RUNTIME_BUFFER_LOCATION 1
 #define SYCL_EXT_ONEAPI_BACKEND_LEVEL_ZERO 1
 #cmakedefine01 SYCL_BUILD_PI_CUDA
 #if SYCL_BUILD_PI_CUDA

sycl/include/CL/sycl/usm.hpp

@@ -10,6 +10,7 @@
 #include <CL/sycl/detail/export.hpp>
 #include <CL/sycl/usm/usm_allocator.hpp>
 #include <CL/sycl/usm/usm_enums.hpp>
+#include <sycl/ext/intel/experimental/usm_properties.hpp>
 
 #include <cstddef>
 

sycl/include/sycl/ext/intel/experimental/usm_properties.hpp

@@ -0,0 +1,37 @@
+#pragma once
+
+#include <CL/sycl/context.hpp>
+#include <CL/sycl/detail/property_helper.hpp>
+#include <CL/sycl/properties/property_traits.hpp>
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace ext {
+namespace intel {
+namespace experimental {
+namespace property {
+namespace usm {
+
+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 usm
+} // namespace property
+} // namespace experimental
+} // namespace intel
+} // namespace ext
+
+template <>
+struct is_property<ext::intel::experimental::property::usm::buffer_location>
+    : std::true_type {};
+
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)

sycl/source/detail/usm/usm_impl.cpp

@@ -93,11 +93,13 @@ void *alignedAllocHost(size_t Alignment, size_t Size, const context &Ctxt,
 
 void *alignedAlloc(size_t Alignment, size_t Size, const context &Ctxt,
                    const device &Dev, alloc Kind,
-                   const detail::code_location &CL) {
+                   const detail::code_location &CL,
+                   const property_list &PropList = {}) {
   XPTI_CREATE_TRACEPOINT(CL);
   void *RetVal = nullptr;
   if (Size == 0)
     return nullptr;
+
   if (Ctxt.is_host()) {
     if (Kind == alloc::unknown) {
       RetVal = nullptr;
@@ -125,8 +127,25 @@ void *alignedAlloc(size_t Alignment, size_t Size, const context &Ctxt,
     switch (Kind) {
     case alloc::device: {
       Id = detail::getSyclObjImpl(Dev)->getHandleRef();
-      Error = Plugin.call_nocheck<PiApiKind::piextUSMDeviceAlloc>(
-          &RetVal, C, Id, nullptr, Size, Alignment);
+      // Parse out buffer location property
+      // Buffer location is only supported on FPGA devices
+      bool IsBufferLocSupported =
+          Dev.has_extension("cl_intel_mem_alloc_buffer_location");
+      if (IsBufferLocSupported &&
+          PropList.has_property<cl::sycl::ext::intel::experimental::property::
+                                    usm::buffer_location>()) {
+        auto location = PropList
+                            .get_property<cl::sycl::ext::intel::experimental::
+                                              property::usm::buffer_location>()
+                            .get_buffer_location();
+        pi_usm_mem_properties props[3] = {PI_MEM_USM_ALLOC_BUFFER_LOCATION,
+                                          location, 0};
+        Error = Plugin.call_nocheck<PiApiKind::piextUSMDeviceAlloc>(
+            &RetVal, C, Id, props, Size, Alignment);
+      } else {
+        Error = Plugin.call_nocheck<PiApiKind::piextUSMDeviceAlloc>(
+            &RetVal, C, Id, nullptr, Size, Alignment);
+      }
       break;
     }
     case alloc::shared: {
@@ -193,8 +212,10 @@ void *malloc_device(size_t Size, const device &Dev, const context &Ctxt,
 }
 
 void *malloc_device(size_t Size, const device &Dev, const context &Ctxt,
-                    const property_list &, const detail::code_location CL) {
-  return malloc_device(Size, Dev, Ctxt, CL);
+                    const property_list &PropList,
+                    const detail::code_location CL) {
+  return detail::usm::alignedAlloc(0, Size, Ctxt, Dev, alloc::device, CL,
+                                   PropList);
 }
 
 void *malloc_device(size_t Size, const queue &Q,

sycl/test/extensions/usm/usm_alloc_utility.cpp

@@ -83,7 +83,10 @@ int main() {
     array = (int *)malloc_device(N * sizeof(int), q);
     check_and_free(array, dev, ctxt);
 
-    array = (int *)malloc_device(N * sizeof(int), q, property_list{});
+    array = malloc_device<int>(
+        N, q,
+        property_list{
+            ext::intel::experimental::property::usm::buffer_location(2)});
     check_and_free(array, dev, ctxt);
 
     array = (int *)aligned_alloc_device(alignof(long long), N * sizeof(int),