[SYCL] Implement SYCL 2020 interoperability API part 3 (#3508)

Add support for sycl::kernel and sycl::kernel_bundle construction

Author: Alexander Batashev

sycl/include/CL/sycl/backend.hpp

@@ -13,10 +13,14 @@
 #include <CL/sycl/buffer.hpp>
 #include <CL/sycl/context.hpp>
 #include <CL/sycl/detail/backend_traits.hpp>
+#include <CL/sycl/detail/common.hpp>
+#include <CL/sycl/detail/export.hpp>
+#include <CL/sycl/detail/pi.h>
 #include <CL/sycl/detail/pi.hpp>
 #include <CL/sycl/device.hpp>
 #include <CL/sycl/event.hpp>
 #include <CL/sycl/exception.hpp>
+#include <CL/sycl/kernel_bundle.hpp>
 #include <CL/sycl/platform.hpp>
 #include <CL/sycl/queue.hpp>
 
@@ -25,11 +29,25 @@
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
 
+namespace detail {
+template <backend Backend, typename T> struct BackendInput {
+  // TODO replace usage of interop with specializations.
+  using type = typename interop<Backend, T>::type;
+};
+
+template <backend Backend, typename T> struct BackendReturn {
+  // TODO replace usage of interop with specializations.
+  using type = typename interop<Backend, T>::type;
+};
+} // namespace detail
+
 template <backend Backend> class backend_traits {
 public:
-  template <class T> using input_type = typename interop<Backend, T>::type;
+  template <class T>
+  using input_type = typename detail::BackendInput<Backend, T>::type;
 
-  template <class T> using return_type = typename interop<Backend, T>::type;
+  template <class T>
+  using return_type = typename detail::BackendReturn<Backend, T>::type;
 
   // TODO define errc once SYCL2020-style exceptions are supported.
 };
@@ -54,6 +72,9 @@ auto get_native(const accessor<DataT, Dimensions, AccessMode, AccessTarget,
     delete;
 
 namespace detail {
+// Forward declaration
+class kernel_bundle_impl;
+
 __SYCL_EXPORT platform make_platform(pi_native_handle NativeHandle,
                                      backend Backend);
 __SYCL_EXPORT device make_device(pi_native_handle NativeHandle,
@@ -66,6 +87,11 @@ __SYCL_EXPORT queue make_queue(pi_native_handle NativeHandle,
                                const async_handler &Handler, backend Backend);
 __SYCL_EXPORT event make_event(pi_native_handle NativeHandle,
                                const context &TargetContext, backend Backend);
+__SYCL_EXPORT kernel make_kernel(pi_native_handle NativeHandle,
+                                 const context &TargetContext, backend Backend);
+__SYCL_EXPORT std::shared_ptr<detail::kernel_bundle_impl>
+make_kernel_bundle(pi_native_handle NativeHandle, const context &TargetContext,
+                   bundle_state State, backend Backend);
 } // namespace detail
 
 template <backend Backend>
@@ -131,5 +157,28 @@ make_buffer(const typename backend_traits<Backend>::template input_type<
   return buffer<T, Dimensions, AllocatorT>(
       reinterpret_cast<cl_mem>(BackendObject), TargetContext, AvailableEvent);
 }
+
+template <backend Backend>
+kernel
+make_kernel(const typename backend_traits<Backend>::template input_type<kernel>
+                &BackendObject,
+            const context &TargetContext) {
+  return detail::make_kernel(detail::pi::cast<pi_native_handle>(BackendObject),
+                             TargetContext, Backend);
+}
+
+template <backend Backend, bundle_state State>
+typename std::enable_if<
+    detail::InteropFeatureSupportMap<Backend>::MakeKernelBundle == true,
+    kernel_bundle<State>>::type
+make_kernel_bundle(const typename backend_traits<Backend>::template input_type<
+                       kernel_bundle<State>> &BackendObject,
+                   const context &TargetContext) {
+  std::shared_ptr<detail::kernel_bundle_impl> KBImpl =
+      detail::make_kernel_bundle(
+          detail::pi::cast<pi_native_handle>(BackendObject), TargetContext,
+          State, Backend);
+  return detail::createSyclObjFromImpl<kernel_bundle<State>>(KBImpl);
+}
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

sycl/include/CL/sycl/backend/level_zero.hpp

@@ -50,6 +50,10 @@ struct interop<backend::level_zero, accessor<DataT, Dimensions, AccessMode,
 };
 
 namespace detail {
+template <> class BackendReturn<backend::level_zero, kernel> {
+  using type = ze_kernel_handle_t;
+};
+
 template <> struct InteropFeatureSupportMap<backend::level_zero> {
   static constexpr bool MakePlatform = true;
   static constexpr bool MakeDevice = false;
@@ -58,6 +62,7 @@ template <> struct InteropFeatureSupportMap<backend::level_zero> {
   static constexpr bool MakeEvent = false;
   static constexpr bool MakeBuffer = false;
   static constexpr bool MakeKernel = false;
+  static constexpr bool MakeKernelBundle = false;
 };
 } // namespace detail
 

sycl/include/CL/sycl/backend/opencl.hpp

@@ -10,9 +10,11 @@
 #pragma once
 
 #include <CL/sycl/accessor.hpp>
+#include <CL/sycl/backend.hpp>
 #include <CL/sycl/backend_types.hpp>
 #include <CL/sycl/detail/backend_traits.hpp>
 #include <CL/sycl/detail/cl.h>
+#include <CL/sycl/kernel_bundle.hpp>
 
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
@@ -59,6 +61,24 @@ struct interop<backend::opencl, buffer<DataT, Dimensions, AllocatorT>> {
 };
 
 namespace detail {
+template <bundle_state State>
+struct BackendInput<backend::opencl, kernel_bundle<State>> {
+  using type = cl_program;
+};
+
+template <bundle_state State>
+struct BackendReturn<backend::opencl, kernel_bundle<State>> {
+  using type = cl_program;
+};
+
+template <> struct BackendInput<backend::opencl, kernel> {
+  using type = cl_kernel;
+};
+
+template <> struct BackendReturn<backend::opencl, kernel> {
+  using type = cl_kernel;
+};
+
 template <> struct InteropFeatureSupportMap<backend::opencl> {
   static constexpr bool MakePlatform = true;
   static constexpr bool MakeDevice = true;
@@ -67,6 +87,7 @@ template <> struct InteropFeatureSupportMap<backend::opencl> {
   static constexpr bool MakeEvent = true;
   static constexpr bool MakeBuffer = true;
   static constexpr bool MakeKernel = true;
+  static constexpr bool MakeKernelBundle = true;
 };
 } // namespace detail
 

sycl/include/CL/sycl/detail/pi.def

@@ -79,6 +79,8 @@ _PI_API(piKernelRetain)
 _PI_API(piKernelRelease)
 _PI_API(piextKernelSetArgPointer)
 _PI_API(piKernelSetExecInfo)
+_PI_API(piextKernelCreateWithNativeHandle)
+_PI_API(piextKernelGetNativeHandle)
 // Event
 _PI_API(piEventCreate)
 _PI_API(piEventGetInfo)

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

@@ -35,9 +35,11 @@
 //   pi_device_binary_property_set PropertySetsEnd;
 // 2. A number of types needed to define pi_device_binary_property_set added.
 // 3. Added new ownership argument to piextContextCreateWithNativeHandle.
+// 4. Add interoperability interfaces for kernel.
 //
+#include "CL/cl.h"
 #define _PI_H_VERSION_MAJOR 3
-#define _PI_H_VERSION_MINOR 4
+#define _PI_H_VERSION_MINOR 5
 
 #define _PI_STRING_HELPER(a) #a
 #define _PI_CONCAT(a, b) _PI_STRING_HELPER(a.b)
@@ -125,7 +127,8 @@ typedef enum {
 typedef enum {
   PI_PROGRAM_BUILD_INFO_STATUS = CL_PROGRAM_BUILD_STATUS,
   PI_PROGRAM_BUILD_INFO_OPTIONS = CL_PROGRAM_BUILD_OPTIONS,
-  PI_PROGRAM_BUILD_INFO_LOG = CL_PROGRAM_BUILD_LOG
+  PI_PROGRAM_BUILD_INFO_LOG = CL_PROGRAM_BUILD_LOG,
+  PI_PROGRAM_BUILD_INFO_BINARY_TYPE = CL_PROGRAM_BINARY_TYPE
 } _pi_program_build_info;
 
 typedef enum {
@@ -135,6 +138,14 @@ typedef enum {
   PI_PROGRAM_BUILD_STATUS_IN_PROGRESS = CL_BUILD_IN_PROGRESS
 } _pi_program_build_status;
 
+typedef enum {
+  PI_PROGRAM_BINARY_TYPE_NONE = CL_PROGRAM_BINARY_TYPE_NONE,
+  PI_PROGRAM_BINARY_TYPE_COMPILED_OBJECT =
+      CL_PROGRAM_BINARY_TYPE_COMPILED_OBJECT,
+  PI_PROGRAM_BINARY_TYPE_LIBRARY = CL_PROGRAM_BINARY_TYPE_LIBRARY,
+  PI_PROGRAM_BINARY_TYPE_EXECUTABLE = CL_PROGRAM_BINARY_TYPE_EXECUTABLE
+} _pi_program_binary_type;
+
 // NOTE: this is made 64-bit to match the size of cl_device_type to
 // make the translation to OpenCL transparent.
 //
@@ -1236,6 +1247,25 @@ __SYCL_EXPORT pi_result piKernelSetExecInfo(pi_kernel kernel,
                                             size_t param_value_size,
                                             const void *param_value);
 
+/// Creates PI kernel object from a native handle.
+/// NOTE: The created PI object takes ownership of the native handle.
+///
+/// \param nativeHandle is the native handle to create PI kernel from.
+/// \param context is the PI context of the kernel.
+/// \param ownNativeHandle tells if SYCL RT should assume the ownership of
+///        the native handle, if it can.
+/// \param kernel is the PI kernel created from the native handle.
+__SYCL_EXPORT pi_result piextKernelCreateWithNativeHandle(
+    pi_native_handle nativeHandle, pi_context context, bool ownNativeHandle,
+    pi_kernel *kernel);
+
+/// Gets the native handle of a PI kernel object.
+///
+/// \param kernel is the PI kernel to get the native handle of.
+/// \param nativeHandle is the native handle of kernel.
+__SYCL_EXPORT pi_result
+piextKernelGetNativeHandle(pi_kernel kernel, pi_native_handle *nativeHandle);
+
 //
 // Events
 //

sycl/include/CL/sycl/kernel.hpp

@@ -10,6 +10,7 @@
 
 #include <CL/sycl/detail/common.hpp>
 #include <CL/sycl/detail/export.hpp>
+#include <CL/sycl/detail/pi.h>
 #include <CL/sycl/info/info_desc.hpp>
 #include <CL/sycl/stl.hpp>
 
@@ -20,6 +21,8 @@ namespace sycl {
 // Forward declaration
 class program;
 class context;
+template <backend Backend> class backend_traits;
+
 namespace detail {
 class kernel_impl;
 
@@ -57,6 +60,8 @@ class __SYCL_EXPORT kernel {
   ///
   /// \param ClKernel is a valid OpenCL cl_kernel instance
   /// \param SyclContext is a valid SYCL context
+  __SYCL2020_DEPRECATED(
+      "OpenCL interop constructors are deprecated, use make_kernel() instead")
   kernel(cl_kernel ClKernel, const context &SyclContext);
 
   kernel(const kernel &RHS) = default;
@@ -78,6 +83,8 @@ class __SYCL_EXPORT kernel {
   /// an invalid_object_error exception will be thrown.
   ///
   /// \return a valid cl_kernel instance
+  __SYCL2020_DEPRECATED(
+      "OpenCL interop get() functions are deprecated, use get_native() instead")
   cl_kernel get() const;
 
   /// Check if the associated SYCL context is a SYCL host context.
@@ -166,10 +173,19 @@ class __SYCL_EXPORT kernel {
                      param>::input_type Value) const;
   // clang-format on
 
+  template <backend Backend>
+  typename backend_traits<Backend>::template return_type<kernel>
+  get_native() const {
+    return detail::pi::cast<
+        backend_traits<Backend>::template return_type<kernel>>(getNativeImpl());
+  }
+
 private:
   /// Constructs a SYCL kernel object from a valid kernel_impl instance.
   kernel(std::shared_ptr<detail::kernel_impl> Impl);
 
+  pi_native_handle getNativeImpl() const;
+
   shared_ptr_class<detail::kernel_impl> impl;
 
   template <class Obj>

sycl/include/CL/sycl/kernel_bundle.hpp

@@ -11,6 +11,8 @@
 #include <CL/sycl/context.hpp>
 #include <CL/sycl/detail/common.hpp>
 #include <CL/sycl/detail/kernel_desc.hpp>
+#include <CL/sycl/detail/pi.h>
+#include <CL/sycl/detail/pi.hpp>
 #include <CL/sycl/device.hpp>
 #include <CL/sycl/kernel.hpp>
 
@@ -20,6 +22,8 @@
 
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
+// Forward declaration
+template <backend Backend> class backend_traits;
 
 enum class bundle_state : char { input = 0, object = 1, executable = 2 };
 
@@ -81,6 +85,8 @@ class __SYCL_EXPORT device_image_plain {
 
   bool has_kernel(const kernel_id &KernelID, const device &Dev) const noexcept;
 
+  pi_native_handle getNative() const;
+
 protected:
   detail::DeviceImageImplPtr impl;
 
@@ -303,6 +309,25 @@ class kernel_bundle : public detail::kernel_bundle_plain {
     return reinterpret_cast<device_image_iterator>(kernel_bundle_plain::end());
   }
 
+  template <backend Backend>
+  std::vector<typename backend_traits<Backend>::template return_type<
+      kernel_bundle<State>>>
+  get_native() {
+    std::vector<typename backend_traits<Backend>::template return_type<
+        kernel_bundle<State>>>
+        ReturnValue;
+    ReturnValue.reserve(std::distance(begin(), end()));
+
+    for (const device_image<State> &DevImg : *this) {
+      ReturnValue.push_back(
+          detail::pi::cast<typename backend_traits<
+              Backend>::template return_type<kernel_bundle<State>>>(
+              DevImg.getNative()));
+    }
+
+    return ReturnValue;
+  }
+
 private:
   kernel_bundle(detail::KernelBundleImplPtr Impl)
       : kernel_bundle_plain(std::move(Impl)) {}

sycl/include/CL/sycl/program.hpp

@@ -36,7 +36,8 @@ enum class program_state { none = 0, compiled = 1, linked = 2 };
 /// \sa queue
 ///
 /// \ingroup sycl_api
-class __SYCL_EXPORT program {
+class __SYCL_EXPORT __SYCL2020_DEPRECATED(
+    "program class is deprecated, use kernel_bundle instead") program {
 public:
   program() = delete;
 

sycl/plugins/cuda/pi_cuda.cpp

@@ -2492,6 +2492,12 @@ pi_result cuda_piEnqueueNativeKernel(
   return {};
 }
 
+pi_result cuda_piextKernelCreateWithNativeHandle(pi_native_handle, pi_context,
+                                                 pi_kernel *) {
+  sycl::detail::pi::die("Unsupported operation");
+  return PI_SUCCESS;
+}
+
 /// \TODO Not implemented
 pi_result cuda_piMemImageCreate(pi_context context, pi_mem_flags flags,
                                 const pi_image_format *image_format,
@@ -4711,6 +4717,8 @@ pi_result piPluginInit(pi_plugin *PluginInit) {
   _PI_CL(piKernelRelease, cuda_piKernelRelease)
   _PI_CL(piKernelSetExecInfo, cuda_piKernelSetExecInfo)
   _PI_CL(piextKernelSetArgPointer, cuda_piextKernelSetArgPointer)
+  _PI_CL(piextKernelCreateWithNativeHandle,
+         cuda_piextKernelCreateWithNativeHandle)
   // Event
   _PI_CL(piEventCreate, cuda_piEventCreate)
   _PI_CL(piEventGetInfo, cuda_piEventGetInfo)

sycl/plugins/level_zero/pi_level_zero.cpp

@@ -3878,6 +3878,22 @@ piEnqueueKernelLaunch(pi_queue Queue, pi_kernel Kernel, pi_uint32 WorkDim,
   return PI_SUCCESS;
 }
 
+pi_result piextKernelCreateWithNativeHandle(pi_native_handle, pi_context, bool,
+                                            pi_kernel *) {
+  die("Unsupported operation");
+  return PI_SUCCESS;
+}
+
+pi_result piextKernelGetNativeHandle(pi_kernel Kernel,
+                                     pi_native_handle *NativeHandle) {
+  PI_ASSERT(Kernel, PI_INVALID_KERNEL);
+  PI_ASSERT(NativeHandle, PI_INVALID_VALUE);
+
+  auto *ZeKernel = pi_cast<ze_kernel_handle_t *>(NativeHandle);
+  *ZeKernel = Kernel->ZeKernel;
+  return PI_SUCCESS;
+}
+
 //
 // Events
 //