[SYCL] Add support for get_native for buffer and fix backend_return_t (#5881)

Author: maximdimakov

sycl/doc/PreprocessorMacros.md

@@ -61,6 +61,10 @@ This file describes macros that have effect on SYCL compiler and run-time.
   will change the behavior of `sycl::get_native()` function and using types for 
   next structs: `interop<backend::opencl, event>`, `BackendInput<backend::opencl, event>`, 
   `BackendReturn<backend::opencl, event>` to be in line with the spec.
+  2) According to spec, `backend_return_t` for opencl buffer 
+  should be `std::vector<cl_mem>` instead of `cl_mem`. Defining this macro 
+  will change the behavior of `interop_handle::get_native_mem()` and `sycl::get_native()` functions
+  and using type for `BackendReturn<backend::opencl, buffer>` to be in line with the spec.
 
 ## Version macros
 

sycl/include/CL/sycl/backend.hpp

@@ -64,6 +64,54 @@ template <backend Backend, typename SyclType>
 using backend_return_t =
     typename backend_traits<Backend>::template return_type<SyclType>;
 
+namespace detail {
+template <backend Backend, typename DataT, int Dimensions, typename AllocatorT>
+struct BufferInterop {
+  using ReturnType =
+      backend_return_t<Backend, buffer<DataT, Dimensions, AllocatorT>>;
+
+  static ReturnType GetNativeObjs(const std::vector<pi_native_handle> &Handle) {
+    ReturnType ReturnValue = 0;
+    if (Handle.size()) {
+      ReturnValue = detail::pi::cast<ReturnType>(Handle[0]);
+    }
+    return ReturnValue;
+  }
+};
+
+#ifdef SYCL2020_CONFORMANT_APIS
+template <typename DataT, int Dimensions, typename AllocatorT>
+struct BufferInterop<backend::opencl, DataT, Dimensions, AllocatorT> {
+  using ReturnType =
+      backend_return_t<backend::opencl, buffer<DataT, Dimensions, AllocatorT>>;
+
+  static ReturnType GetNativeObjs(const std::vector<pi_native_handle> &Handle) {
+    ReturnType ReturnValue{};
+    for (auto &Obj : Handle) {
+      ReturnValue.push_back(
+          detail::pi::cast<typename decltype(ReturnValue)::value_type>(Obj));
+    }
+    return ReturnValue;
+  }
+};
+#endif
+
+template <backend BackendName, typename DataT, int Dimensions,
+          typename AllocatorT>
+auto get_native_buffer(const buffer<DataT, Dimensions, AllocatorT, void> &Obj)
+    -> backend_return_t<BackendName,
+                        buffer<DataT, Dimensions, AllocatorT, void>> {
+  // No check for backend mismatch because buffer can be allocated on different
+  // backends
+  if (BackendName == backend::ext_oneapi_level_zero)
+    throw sycl::runtime_error(
+        errc::feature_not_supported,
+        "Buffer interop is not supported by level zero yet",
+        PI_INVALID_OPERATION);
+  return Obj.template getNative<BackendName>();
+}
+} // namespace detail
+
 template <backend BackendName, class SyclObjectT>
 auto get_native(const SyclObjectT &Obj)
     -> backend_return_t<BackendName, SyclObjectT> {
@@ -75,6 +123,28 @@ auto get_native(const SyclObjectT &Obj)
   return Obj.template get_native<BackendName>();
 }
 
+template <backend BackendName, typename DataT, int Dimensions,
+          typename AllocatorT,
+          std::enable_if_t<BackendName == backend::opencl> * = nullptr>
+#ifndef SYCL2020_CONFORMANT_APIS
+__SYCL_DEPRECATED(
+    "get_native<backend::opencl, buffer>, which return type "
+    "cl_mem is deprecated. According to SYCL 2020 spec, please define "
+    "SYCL2020_CONFORMANT_APIS and use vector<cl_mem> instead.")
+#endif
+auto get_native(const buffer<DataT, Dimensions, AllocatorT> &Obj)
+    -> backend_return_t<BackendName, buffer<DataT, Dimensions, AllocatorT>> {
+  return detail::get_native_buffer<BackendName>(Obj);
+}
+
+template <backend BackendName, typename DataT, int Dimensions,
+          typename AllocatorT,
+          std::enable_if_t<BackendName != backend::opencl> * = nullptr>
+auto get_native(const buffer<DataT, Dimensions, AllocatorT> &Obj)
+    -> backend_return_t<BackendName, buffer<DataT, Dimensions, AllocatorT>> {
+  return detail::get_native_buffer<BackendName>(Obj);
+}
+
 // define SYCL2020_CONFORMANT_APIS to correspond SYCL 2020 spec and return
 // vector<cl_event> from get_native instead of just cl_event
 #ifdef SYCL2020_CONFORMANT_APIS

sycl/include/CL/sycl/buffer.hpp

@@ -31,6 +31,16 @@ make_buffer_helper(pi_native_handle Handle, const context &Ctx, event Evt = {},
   return buffer<T, Dimensions, AllocatorT, void>(Handle, Ctx, OwnNativeHandle,
                                                  Evt);
 }
+
+template <backend BackendName, typename DataT, int Dimensions,
+          typename Allocator>
+auto get_native_buffer(const buffer<DataT, Dimensions, Allocator, void> &Obj)
+    -> backend_return_t<BackendName,
+                        buffer<DataT, Dimensions, Allocator, void>>;
+
+template <backend Backend, typename DataT, int Dimensions,
+          typename AllocatorT = cl::sycl::buffer_allocator>
+struct BufferInterop;
 } // namespace detail
 
 /// Defines a shared array that can be used by kernels in queues.
@@ -605,6 +615,21 @@ class buffer {
       return newRange[0] == 1 && newRange[2] == parentRange[2];
     return newRange[1] == parentRange[1] && newRange[2] == parentRange[2];
   }
+
+  template <backend BackendName, typename DataT, int Dimensions,
+            typename Allocator>
+  friend auto detail::get_native_buffer(
+      const buffer<DataT, Dimensions, Allocator, void> &Obj)
+      -> backend_return_t<BackendName,
+                          buffer<DataT, Dimensions, Allocator, void>>;
+
+  template <backend BackendName>
+  backend_return_t<BackendName, buffer<T, dimensions, AllocatorT>>
+  getNative() const {
+    auto NativeHandles = impl->getNativeVector(BackendName);
+    return detail::BufferInterop<BackendName, T, dimensions,
+                                 AllocatorT>::GetNativeObjs(NativeHandles);
+  }
 };
 
 #ifdef __cpp_deduction_guides

sycl/include/CL/sycl/detail/backend_traits_opencl.hpp

@@ -83,10 +83,17 @@ struct BackendInput<backend::opencl, buffer<DataT, Dimensions, AllocatorT>> {
   using type = cl_mem;
 };
 
+#ifdef SYCL2020_CONFORMANT_APIS
+template <typename DataT, int Dimensions, typename AllocatorT>
+struct BackendReturn<backend::opencl, buffer<DataT, Dimensions, AllocatorT>> {
+  using type = std::vector<cl_mem>;
+};
+#else
 template <typename DataT, int Dimensions, typename AllocatorT>
 struct BackendReturn<backend::opencl, buffer<DataT, Dimensions, AllocatorT>> {
   using type = cl_mem;
 };
+#endif
 
 template <> struct BackendInput<backend::opencl, context> {
   using type = cl_context;

sycl/include/CL/sycl/detail/buffer_impl.hpp

@@ -189,6 +189,10 @@ class __SYCL_EXPORT buffer_impl final : public SYCLMemObjT {
   }
 
   void resize(size_t size) { BaseT::MSizeInBytes = size; }
+
+  void addInteropObject(std::vector<pi_native_handle> &Handles) const;
+
+  std::vector<pi_native_handle> getNativeVector(backend BackendName) const;
 };
 
 } // namespace detail

sycl/include/CL/sycl/interop_handle.hpp

@@ -67,7 +67,7 @@ class interop_handle {
 #else
     (void)Acc;
     // we believe this won't be ever called on device side
-    return 0;
+    return backend_return_t<Backend, buffer<DataT, Dims>>{0};
 #endif
   }
 
@@ -158,25 +158,9 @@ class interop_handle {
   template <backend Backend, typename DataT, int Dims>
   backend_return_t<Backend, buffer<DataT, Dims>>
   getMemImpl(detail::Requirement *Req) const {
-    /*
-      Do not update this cast: a C-style cast is required here.
-
-      This function tries to cast pi_native_handle to the native handle type.
-      pi_native_handle is a typedef of uintptr_t. It is used to store opaque
-      pointers, such as cl_device, and integer handles, such as CUdevice. To
-      convert a uintptr_t to a pointer type, such as cl_device, reinterpret_cast
-      must be used. However, reinterpret_cast cannot be used to convert
-      uintptr_t to a different integer type, such as CUdevice. For this,
-      static_cast must be used. This function must employ a cast that is capable
-      of reinterpret_cast and static_cast depending on the arguments passed to
-      it. A C-style cast will achieve this. The compiler will attempt to
-      interpret it as a static_cast, and will fall back to reinterpret_cast
-      where appropriate.
-
-      https://en.cppreference.com/w/cpp/language/reinterpret_cast
-      https://en.cppreference.com/w/cpp/language/explicit_cast
-      */
-    return (backend_return_t<Backend, buffer<DataT, Dims>>)(getNativeMem(Req));
+    std::vector<pi_native_handle> NativeHandles{getNativeMem(Req)};
+    return detail::BufferInterop<Backend, DataT, Dims>::GetNativeObjs(
+        NativeHandles);
   }
 
   __SYCL_EXPORT pi_native_handle getNativeMem(detail::Requirement *Req) const;

sycl/source/detail/buffer_impl.cpp

@@ -48,6 +48,53 @@ void buffer_impl::constructorNotification(const detail::code_location &CodeLoc,
 void buffer_impl::destructorNotification(void *UserObj) {
   XPTIRegistry::bufferDestructorNotification(UserObj);
 }
+
+void buffer_impl::addInteropObject(
+    std::vector<pi_native_handle> &Handles) const {
+  if (MOpenCLInterop) {
+    if (std::find(Handles.begin(), Handles.end(),
+                  pi::cast<pi_native_handle>(MInteropMemObject)) ==
+        Handles.end()) {
+      const plugin &Plugin = getPlugin();
+      Plugin.call<PiApiKind::piMemRetain>(
+          pi::cast<RT::PiMem>(MInteropMemObject));
+      Handles.push_back(pi::cast<pi_native_handle>(MInteropMemObject));
+    }
+  }
+}
+
+std::vector<pi_native_handle>
+buffer_impl::getNativeVector(backend BackendName) const {
+  std::vector<pi_native_handle> Handles{};
+  if (!MRecord) {
+    addInteropObject(Handles);
+    return Handles;
+  }
+
+  for (auto &Cmd : MRecord->MAllocaCommands) {
+    RT::PiMem NativeMem = pi::cast<RT::PiMem>(Cmd->getMemAllocation());
+    auto Ctx = Cmd->getWorkerContext();
+    auto Platform = Ctx->getPlatformImpl();
+    // If Host Shared Memory is not supported then there is alloca for host that
+    // doesn't have platform
+    if (!Platform)
+      continue;
+    auto Plugin = Platform->getPlugin();
+
+    if (Plugin.getBackend() != BackendName)
+      continue;
+    if (Plugin.getBackend() == backend::opencl) {
+      Plugin.call<PiApiKind::piMemRetain>(NativeMem);
+    }
+
+    pi_native_handle Handle;
+    Plugin.call<PiApiKind::piextMemGetNativeHandle>(NativeMem, &Handle);
+    Handles.push_back(Handle);
+  }
+
+  addInteropObject(Handles);
+  return Handles;
+}
 } // namespace detail
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

sycl/test/abi/sycl_symbols_linux.dump

@@ -4156,6 +4156,8 @@ _ZNK2cl4sycl6detail10image_implILi3EE9get_countEv
 _ZNK2cl4sycl6detail10image_implILi3EE9get_rangeEv
 _ZNK2cl4sycl6detail11SYCLMemObjT9getPluginEv
 _ZNK2cl4sycl6detail11SYCLMemObjT9isInteropEv
+_ZNK2cl4sycl6detail11buffer_impl15getNativeVectorENS0_7backendE
+_ZNK2cl4sycl6detail11buffer_impl16addInteropObjectERSt6vectorImSaImEE
 _ZNK2cl4sycl6detail11stream_impl22get_max_statement_sizeEv
 _ZNK2cl4sycl6detail11stream_impl8get_sizeEv
 _ZNK2cl4sycl6detail12sampler_impl18get_filtering_modeEv

sycl/test/abi/sycl_symbols_windows.dump

@@ -1095,6 +1095,7 @@
 ?acospi@__host_std@cl@@YANN@Z
 ?add@device_global_map@detail@sycl@cl@@YAXPEBXPEBD@Z
 ?addHostAccessorAndWait@detail@sycl@cl@@YAXPEAVAccessorImplHost@123@@Z
+?addInteropObject@buffer_impl@detail@sycl@cl@@QEBAXAEAV?$vector@_KV?$allocator@_K@std@@@std@@@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
@@ -2193,6 +2194,7 @@
 ?getNativeImpl@kernel@sycl@cl@@AEBA_KXZ
 ?getNativeMem@interop_handle@sycl@cl@@AEBA_KPEAVAccessorImplHost@detail@23@@Z
 ?getNativeQueue@interop_handle@sycl@cl@@AEBA_KXZ
+?getNativeVector@buffer_impl@detail@sycl@cl@@QEBA?AV?$vector@_KV?$allocator@_K@std@@@std@@W4backend@34@@Z
 ?getNativeVector@event@sycl@cl@@AEBA?AV?$vector@_KV?$allocator@_K@std@@@std@@XZ
 ?getOSMemSize@OSUtil@detail@sycl@cl@@SA_KXZ
 ?getOSModuleHandle@OSUtil@detail@sycl@cl@@SA_JPEBX@Z

sycl/unittests/SYCL2020/GetNativeOpenCL.cpp

@@ -15,6 +15,7 @@
 
 #include <helpers/CommonRedefinitions.hpp>
 #include <helpers/PiMock.hpp>
+#include <helpers/TestKernel.hpp>
 
 #include <gtest/gtest.h>
 
@@ -52,6 +53,17 @@ static pi_result redefinedEventRetain(pi_event c) {
   return PI_SUCCESS;
 }
 
+static pi_result redefinedMemRetain(pi_mem c) {
+  ++TestCounter;
+  return PI_SUCCESS;
+}
+
+pi_result redefinedMemBufferCreate(pi_context, pi_mem_flags, size_t size,
+                                   void *, pi_mem *,
+                                   const pi_mem_properties *) {
+  return PI_SUCCESS;
+}
+
 pi_result redefinedEventGetInfo(pi_event event, pi_event_info param_name,
                                 size_t param_value_size, void *param_value,
                                 size_t *param_value_size_ret) {
@@ -93,6 +105,9 @@ TEST(GetNative, GetNativeHandle) {
   Mock.redefine<detail::PiApiKind::piDeviceRetain>(redefinedDeviceRetain);
   Mock.redefine<detail::PiApiKind::piProgramRetain>(redefinedProgramRetain);
   Mock.redefine<detail::PiApiKind::piEventRetain>(redefinedEventRetain);
+  Mock.redefine<detail::PiApiKind::piMemRetain>(redefinedMemRetain);
+  Mock.redefine<sycl::detail::PiApiKind::piMemBufferCreate>(
+      redefinedMemBufferCreate);
   Mock.redefine<detail::PiApiKind::piextUSMEnqueueMemset>(
       redefinedUSMEnqueueMemset);
 
@@ -108,14 +123,22 @@ TEST(GetNative, GetNativeHandle) {
   unsigned char *HostAlloc = (unsigned char *)malloc_host(1, Context);
   auto Event = Queue.memset(HostAlloc, 42, 1);
 
+  int Data[1] = {0};
+  sycl::buffer<int, 1> Buffer(&Data[0], sycl::range<1>(1));
+  Queue.submit([&](sycl::handler &cgh) {
+    auto Acc = Buffer.get_access<sycl::access::mode::read_write>(cgh);
+    cgh.single_task<TestKernel>([=]() { (void)Acc; });
+  });
+
   get_native<backend::opencl>(Context);
   get_native<backend::opencl>(Queue);
   get_native<backend::opencl>(Program);
   get_native<backend::opencl>(Device);
   get_native<backend::opencl>(Event);
+  get_native<backend::opencl>(Buffer);
 
   // Depending on global caches state, piDeviceRetain is called either once or
-  // twice, so there'll be 5 or 6 calls.
-  ASSERT_EQ(TestCounter, 5 + DeviceRetainCounter - 1)
+  // twice, so there'll be 6 or 7 calls.
+  ASSERT_EQ(TestCounter, 6 + DeviceRetainCounter - 1)
       << "Not all the retain methods were called";
 }