[SYCL] Implements getNative() interoperability for Level Zero

sycl/CMakeLists.txt

@@ -9,7 +9,7 @@ option(SYCL_ENABLE_WERROR "Treat all warnings as errors in SYCL project" OFF)
 option(SYCL_ADD_DEV_VERSION_POSTFIX "Adds -V postfix to version string" ON)
 
 set(SYCL_MAJOR_VERSION 2)
-set(SYCL_MINOR_VERSION 0)
+set(SYCL_MINOR_VERSION 1)
 set(SYCL_PATCH_VERSION 0)
 set(SYCL_DEV_ABI_VERSION 0)
 if (SYCL_ADD_DEV_VERSION_POSTFIX)

sycl/include/CL/sycl/backend/cuda.hpp

@@ -21,6 +21,7 @@ typedef int CUdevice;
 typedef struct CUctx_st *CUcontext;
 typedef struct CUstream_st *CUstream;
 typedef struct CUevent_st *CUevent;
+typedef struct CUmod_st *CUmodule;
 
 // As defined in the CUDA 10.1 header file. This requires CUDA version > 3.2
 #if defined(_WIN64) || defined(__LP64__)
@@ -40,6 +41,8 @@ template <> struct interop<backend::cuda, queue> { using type = CUstream; };
 
 template <> struct interop<backend::cuda, event> { using type = CUevent; };
 
+template <> struct interop<backend::cuda, program> { using type = CUmodule; };
+
 template <typename DataT, int Dimensions, access::mode AccessMode>
 struct interop<backend::cuda, accessor<DataT, Dimensions, AccessMode,
                                        access::target::global_buffer,

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

@@ -0,0 +1,81 @@
+//==--------- level_zero.hpp - SYCL Level-Zero backend ---------------------==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <CL/sycl.hpp>
+#include <level_zero/ze_api.h>
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+
+template <> struct interop<backend::level0, platform> {
+  using type = ze_driver_handle_t;
+};
+
+template <> struct interop<backend::level0, device> {
+  using type = ze_device_handle_t;
+};
+
+template <> struct interop<backend::level0, queue> {
+  using type = ze_command_queue_handle_t;
+};
+
+template <> struct interop<backend::level0, program> {
+  using type = ze_module_handle_t;
+};
+
+template <typename DataT, int Dimensions, access::mode AccessMode>
+struct interop<backend::level0, accessor<DataT, Dimensions, AccessMode,
+                                         access::target::global_buffer,
+                                         access::placeholder::false_t>> {
+  using type = char *;
+};
+
+namespace level0 {
+
+// Implementation of various "make" functions resides in libsycl.so
+platform make_platform(pi_native_handle NativeHandle);
+device make_device(const platform &Platform, pi_native_handle NativeHandle);
+program make_program(const context &Context, pi_native_handle NativeHandle);
+queue make_queue(const context &Context, pi_native_handle InteropHandle);
+
+// Construction of SYCL platform.
+template <typename T, typename std::enable_if<
+                          std::is_same<T, platform>::value>::type * = nullptr>
+T make(typename interop<backend::level0, T>::type Interop) {
+  return make_platform(reinterpret_cast<pi_native_handle>(Interop));
+}
+
+// Construction of SYCL device.
+template <typename T, typename std::enable_if<
+                          std::is_same<T, device>::value>::type * = nullptr>
+T make(const platform &Platform,
+       typename interop<backend::level0, T>::type Interop) {
+  return make_device(Platform, reinterpret_cast<pi_native_handle>(Interop));
+}
+
+// Construction of SYCL program.
+template <typename T, typename std::enable_if<
+                          std::is_same<T, program>::value>::type * = nullptr>
+T make(const context &Context,
+       typename interop<backend::level0, T>::type Interop) {
+  return make_program(Context, reinterpret_cast<pi_native_handle>(Interop));
+}
+
+// Construction of SYCL queue.
+template <typename T, typename std::enable_if<
+                          std::is_same<T, queue>::value>::type * = nullptr>
+T make(const context &Context,
+       typename interop<backend::level0, T>::type Interop) {
+  return make_queue(Context, reinterpret_cast<pi_native_handle>(Interop));
+}
+
+} // namespace level0
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)

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

@@ -850,9 +850,10 @@ piextDeviceGetNativeHandle(pi_device device, pi_native_handle *nativeHandle);
 /// NOTE: The created PI object takes ownership of the native handle.
 ///
 /// \param nativeHandle is the native handle to create PI device from.
+/// \param platform is the platform of the device.
 /// \param device is the PI device created from the native handle.
 __SYCL_EXPORT pi_result piextDeviceCreateWithNativeHandle(
-    pi_native_handle nativeHandle, pi_device *device);
+    pi_native_handle nativeHandle, pi_platform platform, pi_device *device);
 
 /// Selects the most appropriate device binary based on runtime information
 /// and the IR characteristics.
@@ -944,9 +945,10 @@ piextQueueGetNativeHandle(pi_queue queue, pi_native_handle *nativeHandle);
 /// NOTE: The created PI object takes ownership of the native handle.
 ///
 /// \param nativeHandle is the native handle to create PI queue from.
+/// \param context is the PI context of the queue.
 /// \param queue is the PI queue created from the native handle.
 __SYCL_EXPORT pi_result piextQueueCreateWithNativeHandle(
-    pi_native_handle nativeHandle, pi_queue *queue);
+    pi_native_handle nativeHandle, pi_context context, pi_queue *queue);
 
 //
 // Memory
@@ -1066,9 +1068,10 @@ piextProgramGetNativeHandle(pi_program program, pi_native_handle *nativeHandle);
 /// NOTE: The created PI object takes ownership of the native handle.
 ///
 /// \param nativeHandle is the native handle to create PI program from.
+/// \param context is the PI context of the program.
 /// \param program is the PI program created from the native handle.
 __SYCL_EXPORT pi_result piextProgramCreateWithNativeHandle(
-    pi_native_handle nativeHandle, pi_program *program);
+    pi_native_handle nativeHandle, pi_context context, pi_program *program);
 
 //
 // Kernel

sycl/include/CL/sycl/detail/pi.hpp

@@ -334,14 +334,15 @@ template <class To, class From> inline To cast(From value) {
 
 // These conversions should use PI interop API.
 template <> inline pi::PiProgram cast(cl_program) {
-  RT::assertion(false, "pi::cast -> use piextProgramFromNative");
+  RT::assertion(false, "pi::cast -> use piextCreateProgramWithNativeHandle");
   return {};
 }
 
 template <> inline pi::PiDevice cast(cl_device_id) {
-  RT::assertion(false, "pi::cast -> use piextDeviceFromNative");
+  RT::assertion(false, "pi::cast -> use piextCreateDeviceWithNativeHandle");
   return {};
 }
+
 } // namespace pi
 } // namespace detail
 

sycl/include/CL/sycl/platform.hpp

@@ -112,7 +112,7 @@ class __SYCL_EXPORT platform {
   /// \return a native handle, the type of which defined by the backend.
   template <backend BackendName>
   auto get_native() const -> typename interop<BackendName, platform>::type {
-    return detail::pi::cast<typename interop<BackendName, platform>::type>(
+    return reinterpret_cast<typename interop<BackendName, platform>::type>(
         getNative());
   }
 

sycl/include/CL/sycl/program.hpp

@@ -322,7 +322,17 @@ class __SYCL_EXPORT program {
 #endif // __SYCL_DEVICE_ONLY__
   }
 
+  /// Gets the native handle of the SYCL platform.
+  ///
+  /// \return a native handle, the type of which defined by the backend.
+  template <backend BackendName>
+  auto get_native() const -> typename interop<BackendName, program>::type {
+    return reinterpret_cast<typename interop<BackendName, program>::type>(
+        getNative());
+  }
+
 private:
+  pi_native_handle getNative() const;
   program(shared_ptr_class<detail::program_impl> impl);
 
   /// Template-free version of get_kernel.

sycl/plugins/cuda/pi_cuda.cpp

@@ -1338,10 +1338,12 @@ pi_result cuda_piextDeviceGetNativeHandle(pi_device device,
 /// NOTE: The created PI object takes ownership of the native handle.
 ///
 /// \param[in] nativeHandle The native handle to create PI device object from.
+/// \param[in] platform is the PI platform of the device.
 /// \param[out] device Set to the PI device object created from native handle.
 ///
 /// \return TBD
 pi_result cuda_piextDeviceCreateWithNativeHandle(pi_native_handle nativeHandle,
+                                                 pi_platform platform,
                                                  pi_device *device) {
   cl::sycl::detail::pi::die(
       "Creation of PI device from native handle not implemented");
@@ -1879,10 +1881,12 @@ pi_result cuda_piextQueueGetNativeHandle(pi_queue queue,
 /// NOTE: The created PI object takes ownership of the native handle.
 ///
 /// \param[in] nativeHandle The native handle to create PI queue object from.
+/// \param[in] context is the PI context of the queue.
 /// \param[out] queue Set to the PI queue object created from native handle.
 ///
 /// \return TBD
 pi_result cuda_piextQueueCreateWithNativeHandle(pi_native_handle nativeHandle,
+                                                pi_context context,
                                                 pi_queue *queue) {
   cl::sycl::detail::pi::die(
       "Creation of PI queue from native handle not implemented");
@@ -2489,10 +2493,12 @@ pi_result cuda_piextProgramGetNativeHandle(pi_program program,
 /// NOTE: The created PI object takes ownership of the native handle.
 ///
 /// \param[in] nativeHandle The native handle to create PI program object from.
+/// \param[in] context The PI context of the program.
 /// \param[out] program Set to the PI program object created from native handle.
 ///
 /// \return TBD
 pi_result cuda_piextProgramCreateWithNativeHandle(pi_native_handle nativeHandle,
+                                                  pi_context context,
                                                   pi_program *program) {
   cl::sycl::detail::pi::die(
       "Creation of PI program from native handle not implemented");

sycl/plugins/level_zero/pi_level0.cpp

@@ -1182,10 +1182,16 @@ pi_result piextDeviceGetNativeHandle(pi_device Device,
 }
 
 pi_result piextDeviceCreateWithNativeHandle(pi_native_handle NativeHandle,
+                                            pi_platform Platform,
                                             pi_device *Device) {
+  assert(NativeHandle);
+  assert(Device);
+  assert(Platform);
+
   // Create PI device from the given L0 device handle.
-  die("piextDeviceCreateWithNativeHandle: not supported");
-  return PI_SUCCESS;
+  auto ZeDevice = pi_cast<ze_device_handle_t>(NativeHandle);
+  *Device = new _pi_device(ZeDevice, Platform);
+  return (*Device)->initialize();
 }
 
 pi_result piContextCreate(const pi_context_properties *Properties,
@@ -1370,13 +1376,24 @@ pi_result piQueueFinish(pi_queue Queue) {
 
 pi_result piextQueueGetNativeHandle(pi_queue Queue,
                                     pi_native_handle *NativeHandle) {
-  die("piextQueueGetNativeHandle: not supported");
+  assert(Queue);
+  assert(NativeHandle);
+
+  auto ZeQueue = pi_cast<ze_command_queue_handle_t *>(NativeHandle);
+  // Extract the L0 queue handle from the given PI queue
+  *ZeQueue = Queue->ZeCommandQueue;
   return PI_SUCCESS;
 }
 
 pi_result piextQueueCreateWithNativeHandle(pi_native_handle NativeHandle,
+                                           pi_context Context,
                                            pi_queue *Queue) {
-  die("piextQueueCreateWithNativeHandle: not supported");
+  assert(NativeHandle);
+  assert(Context);
+  assert(Queue);
+
+  auto ZeQueue = pi_cast<ze_command_queue_handle_t>(NativeHandle);
+  *Queue = new _pi_queue(ZeQueue, Context);
   return PI_SUCCESS;
 }
 
@@ -1873,13 +1890,43 @@ pi_result piProgramRelease(pi_program Program) {
 
 pi_result piextProgramGetNativeHandle(pi_program Program,
                                       pi_native_handle *NativeHandle) {
-  die("piextProgramGetNativeHandle: not supported");
+  assert(Program);
+  assert(NativeHandle);
+
+  auto ZeModule = pi_cast<ze_module_handle_t *>(NativeHandle);
+  // Extract the L0 module handle from the given PI program
+  *ZeModule = Program->ZeModule;
   return PI_SUCCESS;
 }
 
 pi_result piextProgramCreateWithNativeHandle(pi_native_handle NativeHandle,
+                                             pi_context Context,
                                              pi_program *Program) {
-  die("piextProgramCreateWithNativeHandle: not supported");
+  assert(NativeHandle);
+  assert(Context);
+  assert(Program);
+
+  auto ZeModule = pi_cast<ze_module_handle_t>(NativeHandle);
+
+  // Create PI program from the given L0 module handle.
+  //
+  // TODO: We don't have the real L0 module descriptor with
+  // which it was created, but that's only needed for zeModuleCreate,
+  // which we don't expect to be called on the interop program.
+  //
+  ze_module_desc_t ZeModuleDesc = {};
+  ZeModuleDesc.version = ZE_MODULE_DESC_VERSION_CURRENT;
+  ZeModuleDesc.format = ZE_MODULE_FORMAT_NATIVE;
+  ZeModuleDesc.inputSize = 0;
+  ZeModuleDesc.pInputModule = nullptr;
+
+  try {
+    *Program = new _pi_program(ZeModule, ZeModuleDesc, Context);
+  } catch (const std::bad_alloc &) {
+    return PI_OUT_OF_HOST_MEMORY;
+  } catch (...) {
+    return PI_ERROR_UNKNOWN;
+  }
   return PI_SUCCESS;
 }
 

sycl/plugins/opencl/pi_opencl.cpp

@@ -160,7 +160,6 @@ static pi_result USMSetIndirectAccess(pi_kernel kernel) {
 
 extern "C" {
 
-// Example of a PI interface that does not map exactly to an OpenCL one.
 pi_result piPlatformsGet(pi_uint32 num_entries, pi_platform *platforms,
                          pi_uint32 *num_platforms) {
   cl_int result = clGetPlatformIDs(cast<cl_uint>(num_entries),
@@ -184,7 +183,6 @@ pi_result piextPlatformCreateWithNativeHandle(pi_native_handle nativeHandle,
   return PI_SUCCESS;
 }
 
-// Example of a PI interface that does not map exactly to an OpenCL one.
 pi_result piDevicesGet(pi_platform platform, pi_device_type device_type,
                        pi_uint32 num_entries, pi_device *devices,
                        pi_uint32 *num_devices) {
@@ -274,7 +272,7 @@ pi_result piextDeviceSelectBinary(pi_device device, pi_device_binary *images,
 }
 
 pi_result piextDeviceCreateWithNativeHandle(pi_native_handle nativeHandle,
-                                            pi_device *piDevice) {
+                                            pi_platform, pi_device *piDevice) {
   assert(piDevice != nullptr);
   *piDevice = reinterpret_cast<pi_device>(nativeHandle);
   return PI_SUCCESS;
@@ -321,7 +319,7 @@ pi_result piQueueCreate(pi_context context, pi_device device,
 }
 
 pi_result piextQueueCreateWithNativeHandle(pi_native_handle nativeHandle,
-                                           pi_queue *piQueue) {
+                                           pi_context, pi_queue *piQueue) {
   assert(piQueue != nullptr);
   *piQueue = reinterpret_cast<pi_queue>(nativeHandle);
   return PI_SUCCESS;
@@ -406,6 +404,7 @@ pi_result piProgramCreate(pi_context context, const void *il, size_t length,
 }
 
 pi_result piextProgramCreateWithNativeHandle(pi_native_handle nativeHandle,
+                                             pi_context,
                                              pi_program *piProgram) {
   assert(piProgram != nullptr);
   *piProgram = reinterpret_cast<pi_program>(nativeHandle);

sycl/source/CMakeLists.txt

@@ -95,6 +95,7 @@ endfunction(add_sycl_rt_library)
 set(SYCL_SOURCES
     "${sycl_inc_dir}/CL/sycl.hpp"
     "backend/opencl.cpp"
+    "backend/level_zero.cpp"
     "detail/accessor_impl.cpp"
     "detail/buffer_impl.cpp"
     "detail/builtins_common.cpp"