[SYCL] Fix handling of mem_channel buffer property

sycl/include/sycl/detail/pi.h

@@ -89,9 +89,10 @@
 // 12.28 Added piextMemImageCreateWithNativeHandle for creating images from
 // native handles.
 // 12.29 Support PI_EXT_PLATFORM_INFO_BACKEND query in piPlatformGetInfo
+// 12.30 Added PI_EXT_INTEL_DEVICE_INFO_MEM_CHANNEL_SUPPORT device info query.
 
 #define _PI_H_VERSION_MAJOR 12
-#define _PI_H_VERSION_MINOR 28
+#define _PI_H_VERSION_MINOR 30
 
 #define _PI_STRING_HELPER(a) #a
 #define _PI_CONCAT(a, b) _PI_STRING_HELPER(a.b)
@@ -348,6 +349,7 @@ typedef enum {
   PI_EXT_CODEPLAY_DEVICE_INFO_SUPPORTS_FUSION = 0x20005,
   PI_EXT_DEVICE_INFO_ATOMIC_FENCE_ORDER_CAPABILITIES = 0x20006,
   PI_EXT_DEVICE_INFO_ATOMIC_FENCE_SCOPE_CAPABILITIES = 0x20007,
+  PI_EXT_INTEL_DEVICE_INFO_MEM_CHANNEL_SUPPORT = 0x20008,
 } _pi_device_info;
 
 typedef enum {

sycl/include/sycl/info/device_traits.def

@@ -202,7 +202,7 @@ __SYCL_PARAM_TRAITS_SPEC(device, sub_group_independent_forward_progress, bool,
 __SYCL_PARAM_TRAITS_SPEC(device, ext_oneapi_srgb, bool,
                          PI_DEVICE_INFO_IMAGE_SRGB)
 __SYCL_PARAM_TRAITS_SPEC(device, ext_intel_mem_channel, bool,
-                         PI_MEM_PROPERTIES_CHANNEL)
+                         PI_EXT_INTEL_DEVICE_INFO_MEM_CHANNEL_SUPPORT)
 __SYCL_PARAM_TRAITS_SPEC(device, ext_oneapi_bfloat16_math_functions, bool,
                          PI_EXT_ONEAPI_DEVICE_INFO_BFLOAT16_MATH_FUNCTIONS)
 

sycl/plugins/cuda/pi_cuda.cpp

@@ -2056,6 +2056,11 @@ pi_result cuda_piDeviceGetInfo(pi_device device, pi_device_info param_name,
     return getInfo(param_value_size, param_value, param_value_size_ret,
                    memory_bandwidth);
   }
+  case PI_EXT_INTEL_DEVICE_INFO_MEM_CHANNEL_SUPPORT: {
+    // The mem-channel buffer property is not supported on CUDA devices.
+    return getInfo<pi_bool>(param_value_size, param_value, param_value_size_ret,
+                            false);
+  }
 
     // TODO: Investigate if this information is available on CUDA.
   case PI_DEVICE_INFO_PCI_ADDRESS:

sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp

@@ -813,6 +813,9 @@ pi_result piDeviceGetInfo(pi_device Device, pi_device_info ParamName,
     return ReturnValue(pi_int32{1});
   case PI_DEVICE_INFO_MAX_NUM_SUB_GROUPS:
     return ReturnValue(pi_uint32{1}); // Minimum required by SYCL 2020 spec
+  case PI_EXT_INTEL_DEVICE_INFO_MEM_CHANNEL_SUPPORT:
+    // The mem-channel buffer property is not supported on the ESIMD emulator.
+    return ReturnValue(pi_bool{false});
 
     CASE_PI_UNSUPPORTED(PI_DEVICE_INFO_SUB_GROUP_INDEPENDENT_FORWARD_PROGRESS)
     CASE_PI_UNSUPPORTED(PI_DEVICE_INFO_IL_VERSION)

sycl/plugins/hip/pi_hip.cpp

@@ -1955,6 +1955,11 @@ pi_result hip_piDeviceGetInfo(pi_device device, pi_device_info param_name,
 #endif
     return PI_ERROR_INVALID_VALUE;
   }
+  case PI_EXT_INTEL_DEVICE_INFO_MEM_CHANNEL_SUPPORT: {
+    // The mem-channel buffer property is not supported on HIP devices.
+    return getInfo<pi_bool>(param_value_size, param_value, param_value_size_ret,
+                            false);
+  }
 
   // TODO: Investigate if this information is available on HIP.
   case PI_DEVICE_INFO_PCI_ADDRESS:

sycl/plugins/opencl/pi_opencl.cpp

@@ -644,6 +644,21 @@ pi_result piDeviceGetInfo(pi_device device, pi_device_info paramName,
 
     return static_cast<pi_result>(CL_SUCCESS);
   }
+  case PI_EXT_INTEL_DEVICE_INFO_MEM_CHANNEL_SUPPORT: {
+    cl_int ret_err = CL_SUCCESS;
+    cl_bool result = CL_FALSE;
+    bool supported = false;
+
+    ret_err =
+        checkDeviceExtensions(cast<cl_device_id>(device),
+                              {"cl_intel_mem_channel_property"}, supported);
+    if (ret_err != CL_SUCCESS)
+      return static_cast<pi_result>(ret_err);
+
+    result = supported;
+    std::memcpy(paramValue, &result, sizeof(cl_bool));
+    return PI_SUCCESS;
+  }
   default:
     cl_int result = clGetDeviceInfo(
         cast<cl_device_id>(device), cast<cl_device_info>(paramName),

sycl/plugins/unified_runtime/pi2ur.hpp

@@ -785,6 +785,9 @@ inline pi_result piDeviceGetInfo(pi_device Device, pi_device_info ParamName,
   case PI_EXT_DEVICE_INFO_ATOMIC_FENCE_SCOPE_CAPABILITIES:
     InfoType = (ur_device_info_t)UR_DEVICE_INFO_ATOMIC_FENCE_SCOPE_CAPABILITIES;
     break;
+  case PI_EXT_INTEL_DEVICE_INFO_MEM_CHANNEL_SUPPORT:
+    InfoType = (ur_device_info_t)UR_EXT_DEVICE_INFO_MEM_CHANNEL_SUPPORT;
+    break;
   default:
     return PI_ERROR_UNKNOWN;
   };

sycl/plugins/unified_runtime/ur/adapters/level_zero/ur_level_zero.cpp

@@ -1013,6 +1013,8 @@ ur_result_t urDeviceGetInfo(
         UR_MEMORY_ORDER_CAPABILITY_FLAG_SEQ_CST;
     return ReturnValue(capabilities);
   }
+  case UR_EXT_DEVICE_INFO_MEM_CHANNEL_SUPPORT:
+    return ReturnValue(pi_bool{false});
 
   // TODO: Implement.
   default:

sycl/plugins/unified_runtime/ur/ur.hpp

@@ -43,6 +43,7 @@ const int UR_EXT_DEVICE_INFO_FREE_MEMORY = UR_EXT_DEVICE_INFO_END - 13;
 // const int ZER_EXT_DEVICE_INFO_DEVICE_ID = UR_EXT_DEVICE_INFO_END - 14;
 // const int ZER_EXT_DEVICE_INFO_IMAGE_MAX_ARRAY_SIZE =
 //     UR_DEVICE_INFO_IMAGE_MAX_ARRAY_SIZE;
+const int UR_EXT_DEVICE_INFO_MEM_CHANNEL_SUPPORT = UR_EXT_DEVICE_INFO_END - 15;
 
 const ur_device_info_t UR_EXT_DEVICE_INFO_OPENCL_C_VERSION =
     (ur_device_info_t)0x103D;

sycl/source/detail/device_impl.cpp

@@ -351,6 +351,8 @@ bool device_impl::has(aspect Aspect) const {
     return get_info<info::device::usm_device_allocations>();
   case aspect::usm_host_allocations:
     return get_info<info::device::usm_host_allocations>();
+  case aspect::ext_intel_mem_channel:
+    return get_info<info::device::ext_intel_mem_channel>();
   case aspect::usm_atomic_host_allocations:
     return is_host() ||
            (get_device_info_impl<pi_usm_capabilities,
@@ -442,10 +444,9 @@ bool device_impl::has(aspect Aspect) const {
             &async_barrier_supported, nullptr) == PI_SUCCESS;
     return call_successful && async_barrier_supported;
   }
-  default:
-    throw runtime_error("This device aspect has not been implemented yet.",
-                        PI_ERROR_INVALID_DEVICE);
   }
+  throw runtime_error("This device aspect has not been implemented yet.",
+                      PI_ERROR_INVALID_DEVICE);
 }
 
 std::shared_ptr<device_impl> device_impl::getHostDeviceImpl() {

sycl/source/detail/device_info.hpp

@@ -762,19 +762,6 @@ struct get_device_info_impl<bool, info::device::usm_system_allocations> {
   }
 };
 
-// Specialization for memory channel query
-template <>
-struct get_device_info_impl<bool, info::device::ext_intel_mem_channel> {
-  static bool get(const DeviceImplPtr &Dev) {
-    pi_mem_properties caps;
-    pi_result Err = Dev->getPlugin().call_nocheck<PiApiKind::piDeviceGetInfo>(
-        Dev->getHandleRef(),
-        PiInfoCode<info::device::ext_intel_mem_channel>::value,
-        sizeof(pi_mem_properties), &caps, nullptr);
-    return (Err != PI_SUCCESS) ? false : (caps & PI_MEM_PROPERTIES_CHANNEL);
-  }
-};
-
 // Specialization for kernel fusion support
 template <>
 struct get_device_info_impl<

sycl/source/detail/memory_manager.cpp

@@ -346,19 +346,36 @@ 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;
-    }
+  std::vector<pi_mem_properties> AllocProps;
+
+  if (PropsList.has_property<property::buffer::detail::buffer_location>() &&
+      TargetContext->isBufferLocationSupported()) {
+    auto Location =
+        PropsList.get_property<property::buffer::detail::buffer_location>()
+            .get_buffer_location();
+    AllocProps.reserve(AllocProps.size() + 2);
+    AllocProps.push_back(PI_MEM_PROPERTIES_ALLOC_BUFFER_LOCATION);
+    AllocProps.push_back(Location);
+  }
+
+  if (PropsList.has_property<property::buffer::mem_channel>()) {
+    auto Channel =
+        PropsList.get_property<property::buffer::mem_channel>().get_channel();
+    AllocProps.reserve(AllocProps.size() + 2);
+    AllocProps.push_back(PI_MEM_PROPERTIES_CHANNEL);
+    AllocProps.push_back(Channel);
+  }
+
+  pi_mem_properties *AllocPropsPtr = nullptr;
+  if (!AllocProps.empty()) {
+    // If there are allocation properties, push an end to the list and update
+    // the properties pointer.
+    AllocProps.push_back(0);
+    AllocPropsPtr = AllocProps.data();
+  }
+
   memBufferCreateHelper(Plugin, TargetContext->getHandleRef(), CreationFlags,
-                        Size, UserPtr, &NewMem, nullptr);
+                        Size, UserPtr, &NewMem, AllocPropsPtr);
   return NewMem;
 }
 

sycl/unittests/buffer/CMakeLists.txt

@@ -2,4 +2,5 @@ add_sycl_unittest(BufferTests OBJECT
   BufferLocation.cpp
   Image.cpp
   BufferDestructionCheck.cpp
+  MemChannel.cpp
 )

sycl/unittests/buffer/MemChannel.cpp

@@ -0,0 +1,106 @@
+//==------------ MemChannel.cpp --- check mem_channel property -------------==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#include <helpers/PiMock.hpp>
+#include <helpers/TestKernel.hpp>
+
+#include <sycl/accessor.hpp>
+#include <sycl/sycl.hpp>
+
+#include <gtest/gtest.h>
+
+#include <detail/buffer_impl.hpp>
+
+constexpr uint32_t DEFAULT_VALUE = 7777;
+static uint32_t PassedChannel = DEFAULT_VALUE;
+
+static pi_result
+redefinedMemBufferCreateBefore(pi_context, pi_mem_flags, size_t size, void *,
+                               pi_mem *, const pi_mem_properties *properties) {
+  PassedChannel = DEFAULT_VALUE;
+  if (!properties)
+    return PI_SUCCESS;
+
+  // properties must ended by 0
+  size_t I = 0;
+  while (properties[I] != 0) {
+    if (properties[I] == PI_MEM_PROPERTIES_CHANNEL) {
+      PassedChannel = properties[I + 1];
+      break;
+    }
+    I += 2;
+  }
+
+  return PI_SUCCESS;
+}
+
+template <bool RetVal>
+static pi_result
+redefinedDeviceGetInfoAfter(pi_device device, pi_device_info param_name,
+                            size_t param_value_size, void *param_value,
+                            size_t *param_value_size_ret) {
+  if (param_name == PI_EXT_INTEL_DEVICE_INFO_MEM_CHANNEL_SUPPORT) {
+    if (param_value)
+      *reinterpret_cast<pi_bool *>(param_value) = RetVal;
+    if (param_value_size_ret)
+      *param_value_size_ret = sizeof(pi_bool);
+  }
+  return PI_SUCCESS;
+}
+
+class BufferMemChannelTest : public ::testing::Test {
+public:
+  BufferMemChannelTest() : Mock{}, Plt{Mock.getPlatform()} {}
+
+protected:
+  void SetUp() override {}
+
+protected:
+  sycl::unittest::PiMock Mock;
+  sycl::platform Plt;
+};
+
+// Test that the mem channel aspect and info query correctly reports true when
+// device supports it.
+TEST_F(BufferMemChannelTest, MemChannelAspectTrue) {
+  Mock.redefineAfter<sycl::detail::PiApiKind::piDeviceGetInfo>(
+      redefinedDeviceGetInfoAfter<true>);
+
+  sycl::device Dev = Plt.get_devices()[0];
+  EXPECT_TRUE(Dev.get_info<sycl::info::device::ext_intel_mem_channel>());
+  EXPECT_TRUE(Dev.has(sycl::aspect::ext_intel_mem_channel));
+}
+
+// Test that the mem channel aspect and info query correctly reports false when
+// device supports it.
+TEST_F(BufferMemChannelTest, MemChannelAspectFalse) {
+  Mock.redefineAfter<sycl::detail::PiApiKind::piDeviceGetInfo>(
+      redefinedDeviceGetInfoAfter<false>);
+
+  sycl::device Dev = Plt.get_devices()[0];
+  EXPECT_FALSE(Dev.get_info<sycl::info::device::ext_intel_mem_channel>());
+  EXPECT_FALSE(Dev.has(sycl::aspect::ext_intel_mem_channel));
+}
+
+// Tests that the right buffer property identifier and values are passed to
+// buffer creation.
+TEST_F(BufferMemChannelTest, MemChannelProp) {
+  Mock.redefineAfter<sycl::detail::PiApiKind::piDeviceGetInfo>(
+      redefinedDeviceGetInfoAfter<true>);
+  Mock.redefineBefore<sycl::detail::PiApiKind::piMemBufferCreate>(
+      redefinedMemBufferCreateBefore);
+
+  sycl::queue Q{Plt.get_devices()[0]};
+  sycl::buffer<int, 1> Buf(3, sycl::property::buffer::mem_channel{42});
+  Q.submit([&](sycl::handler &CGH) {
+     sycl::accessor Acc{Buf, CGH, sycl::read_write};
+     constexpr size_t KS = sizeof(decltype(Acc));
+     CGH.single_task<TestKernel<KS>>([=]() { Acc[0] = 4; });
+   }).wait();
+  EXPECT_EQ(PassedChannel, (uint32_t)42);
+}