[SYCL] Implement SYCL_INTEL_mem_channel_property extension

On some targets manual assignment of buffers to memory regions can
improve memory bandwidth. This extension adds a buffer property to
indicate in which memory channel a particular buffer should be allocated.
This information is an optimization hint to the runtime and thus it is
legal to ignore.

Spec:
#2688

Signed-off-by: Dmitry Sidorov <[email protected]>

Author: MrSidims

sycl/include/CL/sycl/aspects.hpp

@@ -30,7 +30,8 @@ enum class aspect {
   usm_host_allocations,
   usm_shared_allocations,
   usm_restricted_shared_allocations,
-  usm_system_allocator
+  usm_system_allocator,
+  ext_intel_mem_channel
 };
 
 } // namespace sycl

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

@@ -503,9 +503,8 @@ constexpr pi_map_flags PI_MAP_WRITE_INVALIDATE_REGION =
 
 // NOTE: this is made 64-bit to match the size of cl_mem_properties_intel to
 // make the translation to OpenCL transparent.
-// TODO: populate
-//
 using pi_mem_properties = pi_bitfield;
+constexpr pi_mem_properties PI_MEM_CHANNEL_INTEL = CL_MEM_CHANNEL_INTEL;
 
 // NOTE: queue properties are implemented this way to better support bit
 // manipulations

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

@@ -40,6 +40,7 @@ enum PropWithDataKind {
   BufferContextBound,
   ImageUseMutex,
   ImageContextBound,
+  BufferMemChannel,
   PropWithDataKindSize
 };
 

sycl/include/CL/sycl/info/device_traits.def

@@ -85,3 +85,4 @@ __SYCL_PARAM_TRAITS_SPEC(device, usm_host_allocations, bool)
 __SYCL_PARAM_TRAITS_SPEC(device, usm_shared_allocations, bool)
 __SYCL_PARAM_TRAITS_SPEC(device, usm_restricted_shared_allocations, bool)
 __SYCL_PARAM_TRAITS_SPEC(device, usm_system_allocator, bool)
+__SYCL_PARAM_TRAITS_SPEC(device, ext_intel_mem_channel, bool)

sycl/include/CL/sycl/info/info_desc.hpp

@@ -130,7 +130,9 @@ enum class device : cl_device_info {
   usm_host_allocations = PI_USM_HOST_SUPPORT,
   usm_shared_allocations = PI_USM_SINGLE_SHARED_SUPPORT,
   usm_restricted_shared_allocations = PI_USM_CROSS_SHARED_SUPPORT,
-  usm_system_allocator = PI_USM_SYSTEM_SHARED_SUPPORT
+  usm_system_allocator = PI_USM_SYSTEM_SHARED_SUPPORT,
+
+  ext_intel_mem_channel = PI_MEM_CHANNEL_INTEL
 };
 
 enum class device_type : pi_uint64 {

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

@@ -39,6 +39,17 @@ class context_bound
 private:
   sycl::context MCtx;
 };
+
+class mem_channel : public detail::PropertyWithData<
+                        detail::PropWithDataKind::BufferMemChannel> {
+public:
+  mem_channel(cl_uint Channel) : m_Channel(Channel) {}
+  cl_uint get_channel() const { return m_Channel; }
+
+private:
+  cl_uint m_Channel;
+};
+
 } // namespace buffer
 } // namespace property
 

sycl/plugins/opencl/pi_opencl.cpp

@@ -523,22 +523,62 @@ pi_result piMemBufferCreate(pi_context context, pi_mem_flags flags, size_t size,
                             void *host_ptr, pi_mem *ret_mem,
                             const pi_mem_properties *properties) {
   pi_result ret_err = PI_INVALID_OPERATION;
-  clCreateBufferWithPropertiesINTEL_fn FuncPtr = nullptr;
-
-  if (properties)
+  if (properties) {
+    clCreateBufferWithPropertiesINTEL_fn FuncPtr = nullptr;
+    const size_t propSize = sizeof(properties) / sizeof(pi_mem_properties);
     // First we need to look up the function pointer
     ret_err = getExtFuncFromContext<clCreateBufferWithPropertiesName,
                                     clCreateBufferWithPropertiesINTEL_fn>(
         context, &FuncPtr);
+    if (FuncPtr) {
+      std::vector<pi_mem_properties> supported(properties,
+                                               properties + propSize);
+      // Go through buffer properties. If there is one, that shall be propagated
+      // to an OpenCL runtime - check if this property is being supported.
+      for (auto prop = supported.begin(); prop != supported.end(); ++prop) {
+        // Check if PI_MEM_CHANNEL_INTEL property is supported. If it's not -
+        // just ignore it, as it's an optimization hint.
+        if (*prop == PI_MEM_CHANNEL_INTEL) {
+          size_t deviceCount;
+          cl_int ret_err =
+              clGetContextInfo(cast<cl_context>(context), CL_CONTEXT_DEVICES, 0,
+                               nullptr, &deviceCount);
+          if (ret_err != CL_SUCCESS || deviceCount < 1)
+            return PI_INVALID_CONTEXT;
+          std::vector<cl_device_id> devicesInCtx(deviceCount);
+          ret_err = clGetContextInfo(
+              cast<cl_context>(context), CL_CONTEXT_DEVICES,
+              deviceCount * sizeof(cl_device_id), devicesInCtx.data(), nullptr);
+
+          size_t retSize;
+          ret_err = clGetDeviceInfo(devicesInCtx[0], CL_DEVICE_EXTENSIONS, 0,
+                                    nullptr, &retSize);
+          if (ret_err != CL_SUCCESS)
+            return PI_INVALID_DEVICE;
+          std::string extensions(retSize, '\0');
+          ret_err = clGetDeviceInfo(devicesInCtx[0], CL_DEVICE_EXTENSIONS,
+                                    retSize, &extensions[0], nullptr);
+          if (ret_err != CL_SUCCESS)
+            return PI_INVALID_DEVICE;
+
+          size_t pos = extensions.find("cl_intel_mem_channel_property");
+          if (pos == std::string::npos)
+            supported.erase(prop);
+        }
+      }
+      if (!supported.empty()) {
+        *ret_mem =
+            cast<pi_mem>(FuncPtr(cast<cl_context>(context), supported.data(),
+                                 cast<cl_mem_flags>(flags), size, host_ptr,
+                                 cast<cl_int *>(&ret_err)));
+        return ret_err;
+      }
+    }
+  }
 
-  if (FuncPtr)
-    *ret_mem = cast<pi_mem>(FuncPtr(cast<cl_context>(context), properties,
-                                    cast<cl_mem_flags>(flags), size, host_ptr,
-                                    cast<cl_int *>(&ret_err)));
-  else
-    *ret_mem = cast<pi_mem>(clCreateBuffer(cast<cl_context>(context),
-                                           cast<cl_mem_flags>(flags), size,
-                                           host_ptr, cast<cl_int *>(&ret_err)));
+  *ret_mem = cast<pi_mem>(clCreateBuffer(cast<cl_context>(context),
+                                         cast<cl_mem_flags>(flags), size,
+                                         host_ptr, cast<cl_int *>(&ret_err)));
   return ret_err;
 }
 

sycl/source/detail/device_info.hpp

@@ -936,6 +936,11 @@ inline bool get_device_info_host<info::device::usm_system_allocator>() {
   return true;
 }
 
+template <>
+inline bool get_device_info_host<info::device::ext_intel_mem_channel>() {
+  return false;
+}
+
 cl_uint get_native_vector_width(size_t idx);
 
 // USM
@@ -1003,6 +1008,17 @@ template <> struct get_device_info<bool, info::device::usm_system_allocator> {
   }
 };
 
+// Specialization for memory channel query
+template <> struct get_device_info<bool, info::device::ext_intel_mem_channel> {
+  static bool get(RT::PiDevice dev, const plugin &Plugin) {
+    pi_mem_properties caps;
+    pi_result Err = Plugin.call_nocheck<PiApiKind::piDeviceGetInfo>(
+        dev, pi::cast<RT::PiDeviceInfo>(info::device::ext_intel_mem_channel),
+        sizeof(pi_mem_properties), &caps, nullptr);
+    return (Err != PI_SUCCESS) ? false : (caps & PI_MEM_CHANNEL_INTEL);
+  }
+};
+
 } // namespace detail
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

sycl/test/abi/sycl_symbols_linux.dump

@@ -4036,6 +4036,7 @@ _ZNK2cl4sycl6device8get_infoILNS0_4info6deviceE4168EEENS3_12param_traitsIS4_XT_E
 _ZNK2cl4sycl6device8get_infoILNS0_4info6deviceE4169EEENS3_12param_traitsIS4_XT_EE11return_typeEv
 _ZNK2cl4sycl6device8get_infoILNS0_4info6deviceE4188EEENS3_12param_traitsIS4_XT_EE11return_typeEv
 _ZNK2cl4sycl6device8get_infoILNS0_4info6deviceE4189EEENS3_12param_traitsIS4_XT_EE11return_typeEv
+_ZNK2cl4sycl6device8get_infoILNS0_4info6deviceE16915EEENS3_12param_traitsIS4_XT_EE11return_typeEv
 _ZNK2cl4sycl6device9getNativeEv
 _ZNK2cl4sycl6kernel11get_contextEv
 _ZNK2cl4sycl6kernel11get_programEv

sycl/test/basic_tests/property_list.cpp

@@ -64,6 +64,23 @@ int main() {
     }
   }
 
+  {
+    cl::sycl::property_list MemChannelProp{
+        sycl_property::buffer::mem_channel(2)};
+    if (!MemChannelProp.has_property<sycl_property::buffer::mem_channel>()) {
+      std::cerr << "Error: property list has no property while should have."
+                << std::endl;
+      Failed = true;
+    }
+    auto Prop =
+        MemChannelProp.get_property<sycl_property::buffer::mem_channel>();
+    if (Prop.get_channel() != 2) {
+      std::cerr << "Error: mem_channel property is not equal to 2."
+                << std::endl;
+      Failed = true;
+    }
+  }
+
   std::cerr << "Test status : " << (Failed ? "FAILED" : "PASSED") << std::endl;
 
   return Failed;

sycl/test/on-device/basic_tests/buffer/buffer.cpp

@@ -40,6 +40,25 @@ int main() {
       assert(data1[i] == 0);
   }
 
+  {
+    int data1[10] = {-1, -1, -1, -1, -1, -1, -1, -1, -1, -1};
+    {
+      buffer<int, 1> b(data1, range<1>(10), {property::buffer::mem_channel{3}});
+      queue myQueue;
+      myQueue.submit([&](handler &cgh) {
+        auto B = b.get_access<access::mode::read_write>(cgh);
+        cgh.parallel_for<class init_a_2>(range<1>{10},
+                                         [=](id<1> index) { B[index] = 0; });
+      });
+      assert(b.has_property<property::buffer::mem_channel>());
+      auto prop = b.get_property<property::buffer::mem_channel>();
+      assert(prop.get_channel() == 3 && "oops it's not 3");
+
+    } // Data is copied back because there is a user side shared_ptr
+    for (int i = 0; i < 10; i++)
+      assert(data1[i] == 0);
+  }
+
   {
     std::vector<int> data1(10, -1);
     {