[SYCL] Interoperability API to create a context from a Level Zero handler (#2772)

Implement the missing API to create a SYCL context from a Level Zero
handler. Also add a test to check the interop API for the Level Zero.
This commit also includes changes to properly copy the Level Zero
loader and headers to the build directory.

Author: againull

sycl/CMakeLists.txt

@@ -332,6 +332,7 @@ set( SYCL_TOOLCHAIN_DEPLOY_COMPONENTS
      clang-offload-bundler
      file-table-tform
      level-zero-loader
+     level-zero-headers
      llc
      llvm-ar
      llvm-foreach

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

@@ -52,10 +52,15 @@ struct interop<backend::level_zero, accessor<DataT, Dimensions, AccessMode,
 namespace level_zero {
 
 // 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);
+__SYCL_EXPORT platform make_platform(pi_native_handle NativeHandle);
+__SYCL_EXPORT device make_device(const platform &Platform,
+                                 pi_native_handle NativeHandle);
+__SYCL_EXPORT context make_context(const vector_class<device> &DeviceList,
+                                   pi_native_handle NativeHandle);
+__SYCL_EXPORT program make_program(const context &Context,
+                                   pi_native_handle NativeHandle);
+__SYCL_EXPORT queue make_queue(const context &Context,
+                               pi_native_handle InteropHandle);
 
 // Construction of SYCL platform.
 template <typename T, typename detail::enable_if_t<
@@ -72,6 +77,18 @@ T make(const platform &Platform,
   return make_device(Platform, reinterpret_cast<pi_native_handle>(Interop));
 }
 
+/// Construction of SYCL context.
+/// \param DeviceList is a vector of devices which must be encapsulated by
+///        created SYCL context. Provided devices and native context handle must
+///        be associated with the same platform.
+/// \param Interop is a Level Zero native context handle.
+template <typename T, typename std::enable_if<
+                          std::is_same<T, context>::value>::type * = nullptr>
+T make(const vector_class<device> &DeviceList,
+       typename interop<backend::level_zero, T>::type Interop) {
+  return make_context(DeviceList, detail::pi::cast<pi_native_handle>(Interop));
+}
+
 // Construction of SYCL program.
 template <typename T, typename detail::enable_if_t<
                           std::is_same<T, program>::value> * = nullptr>

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

@@ -35,8 +35,8 @@
 //   pi_device_binary_property_set PropertySetsEnd;
 // 2. A number of types needed to define pi_device_binary_property_set added.
 //
-#define _PI_H_VERSION_MAJOR 1
-#define _PI_H_VERSION_MINOR 2
+#define _PI_H_VERSION_MAJOR 2
+#define _PI_H_VERSION_MINOR 3
 
 #define _PI_STRING_HELPER(a) #a
 #define _PI_CONCAT(a, b) _PI_STRING_HELPER(a.b)
@@ -944,11 +944,27 @@ piextContextGetNativeHandle(pi_context context, pi_native_handle *nativeHandle);
 
 /// Creates PI context object from a native handle.
 /// NOTE: The created PI object takes ownership of the native handle.
+/// NOTE: The number of devices and the list of devices is needed for Level Zero
+/// backend because there is no possilibity to query this information from
+/// context handle for Level Zero. If backend has API to query a list of devices
+/// from the context native handle then these parameters are ignored.
 ///
 /// \param nativeHandle is the native handle to create PI context from.
+/// \param numDevices is the number of devices in the context. Parameter is
+///        ignored if number of devices can be queried from the context native
+///        handle for a backend.
+/// \param devices is the list of devices in the context. Parameter is ignored
+///        if devices can be queried from the context native handle for a
+///        backend.
 /// \param context is the PI context created from the native handle.
+/// \return PI_SUCCESS if successfully created pi_context from the handle.
+///         PI_OUT_OF_HOST_MEMORY if can't allocate memory for the pi_context
+///         object. PI_INVALID_VALUE if numDevices == 0 or devices is NULL but
+///         backend doesn't have API to query a list of devices from the context
+///         native handle. PI_UNKNOWN_ERROR in case of another error.
 __SYCL_EXPORT pi_result piextContextCreateWithNativeHandle(
-    pi_native_handle nativeHandle, pi_context *context);
+    pi_native_handle nativeHandle, pi_uint32 numDevices,
+    const pi_device *devices, pi_context *context);
 
 //
 // Queue

sycl/plugins/cuda/pi_cuda.cpp

@@ -1645,6 +1645,8 @@ pi_result cuda_piextContextGetNativeHandle(pi_context context,
 ///
 /// \return TBD
 pi_result cuda_piextContextCreateWithNativeHandle(pi_native_handle nativeHandle,
+                                                  pi_uint32 num_devices,
+                                                  const pi_device *devices,
                                                   pi_context *context) {
   cl::sycl::detail::pi::die(
       "Creation of PI context from native handle not implemented");

sycl/plugins/level_zero/CMakeLists.txt

@@ -44,7 +44,8 @@ if (NOT DEFINED LEVEL_ZERO_LIBRARY OR NOT DEFINED LEVEL_ZERO_INCLUDE_DIR)
     BUILD_BYPRODUCTS ${LEVEL_ZERO_LOADER}
   )
   ExternalProject_Add_Step(level-zero-loader llvminstall
-    COMMAND ${CMAKE_COMMAND} -E copy_directory <INSTALL_DIR>/ ${LLVM_BINARY_DIR}
+    COMMAND ${CMAKE_COMMAND} -E copy_directory <INSTALL_DIR>/lib/ ${LLVM_BINARY_DIR}/lib
+    COMMAND ${CMAKE_COMMAND} -E copy_directory <INSTALL_DIR>/include/ ${LLVM_BINARY_DIR}/include/sycl
     COMMENT "Installing level-zero-loader into the LLVM binary directory"
     DEPENDEES install
   )
@@ -54,14 +55,35 @@ if (NOT DEFINED LEVEL_ZERO_LIBRARY OR NOT DEFINED LEVEL_ZERO_INCLUDE_DIR)
     COMPONENT level-zero-loader
   )
 
-  list(APPEND SYCL_TOOLCHAIN_DEPLOY_COMPONENTS level-zero-loader)
+  set(LEVEL_ZERO_INCLUDE_DIR ${CMAKE_CURRENT_BINARY_DIR}/level_zero_loader_install/include/)
 else()
-  include_directories("${LEVEL_ZERO_INCLUDE_DIR}")
   file(GLOB LEVEL_ZERO_LIBRARY_SRC "${LEVEL_ZERO_LIBRARY}*")
-  file(COPY ${LEVEL_ZERO_LIBRARY_SRC} DESTINATION ${LLVM_LIBRARY_OUTPUT_INTDIR})
-  add_custom_target(level-zero-loader DEPENDS ${LEVEL_ZERO_LIBRARY} COMMENT "Copying Level Zero Loader ...")
+  get_filename_component(LEVEL_ZERO_LIB_NAME ${LEVEL_ZERO_LIBRARY} NAME)
+  add_custom_target(level-zero-loader
+      DEPENDS
+	  ${LLVM_LIBRARY_OUTPUT_INTDIR}/${LEVEL_ZERO_LIB_NAME}
+	  ${LLVM_BINARY_DIR}/include/sycl/level_zero
+  )
+  add_custom_command(
+      OUTPUT
+	  ${LLVM_LIBRARY_OUTPUT_INTDIR}/${LEVEL_ZERO_LIB_NAME}
+	  ${LLVM_BINARY_DIR}/include/sycl/level_zero
+      COMMENT
+	  "Copying Level Zero loader and headers"
+      COMMAND
+	  ${CMAKE_COMMAND} -E copy ${LEVEL_ZERO_LIBRARY_SRC} ${LLVM_LIBRARY_OUTPUT_INTDIR}
+      COMMAND
+	  ${CMAKE_COMMAND} -E copy_directory ${LEVEL_ZERO_INCLUDE_DIR} ${LLVM_BINARY_DIR}/include/sycl
+      DEPENDS
+	  ${LEVEL_ZERO_LIBRARY}
+          ${LEVEL_ZERO_INCLUDE_DIR}
+  )
 endif()
 
+list(APPEND SYCL_TOOLCHAIN_DEPLOY_COMPONENTS level-zero-loader level-zero-headers)
+
+include_directories("${LEVEL_ZERO_INCLUDE_DIR}")
+
 add_library (LevelZeroLoader-Headers INTERFACE)
 add_library (LevelZeroLoader::Headers ALIAS LevelZeroLoader-Headers)
 target_include_directories(LevelZeroLoader-Headers

sycl/plugins/level_zero/pi_level_zero.cpp

@@ -423,6 +423,20 @@ pi_result _pi_device::initialize() {
   return PI_SUCCESS;
 }
 
+pi_result _pi_context::initialize() {
+  // Create the immediate command list to be used for initializations
+  // Created as synchronous so level-zero performs implicit synchronization and
+  // there is no need to query for completion in the plugin
+  ze_command_queue_desc_t ZeCommandQueueDesc = {};
+  ZeCommandQueueDesc.ordinal = Devices[0]->ZeComputeQueueGroupIndex;
+  ZeCommandQueueDesc.index = 0;
+  ZeCommandQueueDesc.mode = ZE_COMMAND_QUEUE_MODE_SYNCHRONOUS;
+  ZE_CALL(zeCommandListCreateImmediate(ZeContext, Devices[0]->ZeDevice,
+                                       &ZeCommandQueueDesc,
+                                       &ZeCommandListInit));
+  return PI_SUCCESS;
+}
+
 pi_result
 _pi_queue::resetCommandListFenceEntry(ze_command_list_handle_t ZeCommandList,
                                       bool MakeAvailable) {
@@ -1703,29 +1717,19 @@ pi_result piContextCreate(const pi_context_properties *Properties,
 
   assert(RetContext);
 
+  ze_context_desc_t ContextDesc = {ZE_STRUCTURE_TYPE_CONTEXT_DESC, nullptr, 0};
+  ze_context_handle_t ZeContext;
+  ZE_CALL(zeContextCreate((*Devices)->Platform->ZeDriver, &ContextDesc,
+                          &ZeContext));
   try {
-    *RetContext = new _pi_context(NumDevices, Devices);
+    *RetContext = new _pi_context(ZeContext, NumDevices, Devices);
+    (*RetContext)->initialize();
   } catch (const std::bad_alloc &) {
     return PI_OUT_OF_HOST_MEMORY;
   } catch (...) {
     return PI_ERROR_UNKNOWN;
   }
 
-  ze_context_desc_t ContextDesc = {ZE_STRUCTURE_TYPE_CONTEXT_DESC, nullptr, 0};
-  ZE_CALL(zeContextCreate((*Devices)->Platform->ZeDriver, &ContextDesc,
-                          &((*RetContext)->ZeContext)));
-
-  // Create the immediate command list to be used for initializations
-  // Created as synchronous so level-zero performs implicit synchronization and
-  // there is no need to query for completion in the plugin
-  ze_command_queue_desc_t ZeCommandQueueDesc = {};
-  ZeCommandQueueDesc.ordinal = (*Devices)->ZeComputeQueueGroupIndex;
-  ZeCommandQueueDesc.index = 0;
-  ZeCommandQueueDesc.mode = ZE_COMMAND_QUEUE_MODE_SYNCHRONOUS;
-  ZE_CALL(zeCommandListCreateImmediate(
-      (*RetContext)->ZeContext, (*Devices)->ZeDevice, &ZeCommandQueueDesc,
-      (&(*RetContext)->ZeCommandListInit)));
-
   return PI_SUCCESS;
 }
 
@@ -1772,8 +1776,26 @@ pi_result piextContextGetNativeHandle(pi_context Context,
 }
 
 pi_result piextContextCreateWithNativeHandle(pi_native_handle NativeHandle,
-                                             pi_context *Context) {
-  die("piextContextCreateWithNativeHandle: not supported");
+                                             pi_uint32 NumDevices,
+                                             const pi_device *Devices,
+                                             pi_context *RetContext) {
+  assert(NativeHandle);
+  assert(RetContext);
+
+  if (!Devices || !NumDevices) {
+    return PI_INVALID_VALUE;
+  }
+
+  try {
+    *RetContext = new _pi_context(pi_cast<ze_context_handle_t>(NativeHandle),
+                                  NumDevices, Devices);
+    (*RetContext)->initialize();
+  } catch (const std::bad_alloc &) {
+    return PI_OUT_OF_HOST_MEMORY;
+  } catch (...) {
+    return PI_ERROR_UNKNOWN;
+  }
+
   return PI_SUCCESS;
 }
 

sycl/plugins/level_zero/pi_level_zero.hpp

@@ -192,10 +192,11 @@ struct _pi_device : _pi_object {
 };
 
 struct _pi_context : _pi_object {
-  _pi_context(pi_uint32 NumDevices, const pi_device *Devs)
-      : Devices{Devs, Devs + NumDevices}, ZeCommandListInit{nullptr},
-        ZeEventPool{nullptr}, NumEventsAvailableInEventPool{},
-        NumEventsLiveInEventPool{} {
+  _pi_context(ze_context_handle_t ZeContext, pi_uint32 NumDevices,
+              const pi_device *Devs)
+      : ZeContext{ZeContext}, Devices{Devs, Devs + NumDevices},
+        ZeCommandListInit{nullptr}, ZeEventPool{nullptr},
+        NumEventsAvailableInEventPool{}, NumEventsLiveInEventPool{} {
     // Create USM allocator context for each pair (device, context).
     for (uint32_t I = 0; I < NumDevices; I++) {
       pi_device Device = Devs[I];
@@ -207,9 +208,14 @@ struct _pi_context : _pi_object {
           std::piecewise_construct, std::make_tuple(Device),
           std::make_tuple(std::unique_ptr<SystemMemory>(
               new USMDeviceMemoryAlloc(this, Device))));
+      // NOTE: one must additionally call initialize() to complete
+      // PI context creation.
     }
   }
 
+  // Initialize the PI context.
+  pi_result initialize();
+
   // A L0 context handle is primarily used during creation and management of
   // resources that may be used by multiple devices.
   ze_context_handle_t ZeContext;

sycl/plugins/opencl/pi_opencl.cpp

@@ -511,6 +511,8 @@ pi_result piContextCreate(const pi_context_properties *properties,
 }
 
 pi_result piextContextCreateWithNativeHandle(pi_native_handle nativeHandle,
+                                             pi_uint32 num_devices,
+                                             const pi_device *devices,
                                              pi_context *piContext) {
   assert(piContext != nullptr);
   *piContext = reinterpret_cast<pi_context>(nativeHandle);

sycl/source/backend/level_zero.cpp

@@ -45,6 +45,24 @@ __SYCL_EXPORT device make_device(const platform &Platform,
       PlatformImpl->getOrMakeDeviceImpl(PiDevice, PlatformImpl));
 }
 
+//----------------------------------------------------------------------------
+// Implementation of level_zero::make<context>
+__SYCL_EXPORT context make_context(const vector_class<device> &DeviceList,
+                                   pi_native_handle NativeHandle) {
+  const auto &Plugin = pi::getPlugin<backend::level_zero>();
+  // Create PI context first.
+  pi_context PiContext;
+  vector_class<pi_device> DeviceHandles;
+  for (auto Dev : DeviceList) {
+    DeviceHandles.push_back(detail::getSyclObjImpl(Dev)->getHandleRef());
+  }
+  Plugin.call<PiApiKind::piextContextCreateWithNativeHandle>(
+      NativeHandle, DeviceHandles.size(), DeviceHandles.data(), &PiContext);
+  // Construct the SYCL context from PI context.
+  return detail::createSyclObjFromImpl<context>(
+      std::make_shared<context_impl>(PiContext, async_handler{}, Plugin));
+}
+
 //----------------------------------------------------------------------------
 // Implementation of level_zero::make<program>
 __SYCL_EXPORT program make_program(const context &Context,

sycl/source/backend/opencl.cpp

@@ -50,8 +50,8 @@ __SYCL_EXPORT context make_context(pi_native_handle NativeHandle) {
   const auto &Plugin = pi::getPlugin<backend::opencl>();
   // Create PI context first.
   pi::PiContext PiContext;
-  Plugin.call<PiApiKind::piextContextCreateWithNativeHandle>(NativeHandle,
-                                                             &PiContext);
+  Plugin.call<PiApiKind::piextContextCreateWithNativeHandle>(
+      NativeHandle, 0, nullptr, &PiContext);
   // Construct the SYCL context from PI context.
   return detail::createSyclObjFromImpl<context>(
       std::make_shared<context_impl>(PiContext, async_handler{}, Plugin));

sycl/test/abi/sycl_symbols_linux.dump

@@ -3591,6 +3591,7 @@ _ZN2cl10__host_std9u_sub_satEmm
 _ZN2cl10__host_std9u_sub_satEtt
 _ZN2cl4sycl10level_zero10make_queueERKNS0_7contextEm
 _ZN2cl4sycl10level_zero11make_deviceERKNS0_8platformEm
+_ZN2cl4sycl10level_zero12make_contextERKSt6vectorINS0_6deviceESaIS3_EEm
 _ZN2cl4sycl10level_zero12make_programERKNS0_7contextEm
 _ZN2cl4sycl10level_zero13make_platformEm
 _ZN2cl4sycl11malloc_hostEmRKNS0_5queueE

sycl/test/on-device/plugins/interop-level-zero.cpp

@@ -0,0 +1,52 @@
+// REQUIRES: level_zero
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -I %sycl_source_dir %s -o %t.out
+// RUN: env SYCL_BE=PI_LEVEL_ZERO %GPU_RUN_PLACEHOLDER %t.out
+
+// Test for Level Zero interop API
+
+#include <CL/sycl.hpp>
+// clang-format off
+#include <level_zero/ze_api.h>
+#include <CL/sycl/backend/level_zero.hpp>
+// clang-format on
+
+using namespace cl::sycl;
+
+int main() {
+  queue Queue{};
+  auto Context = Queue.get_info<info::queue::context>();
+  auto Device = Queue.get_info<info::queue::device>();
+  auto Platform = Device.get_info<info::device::platform>();
+
+  // Get native Level Zero handles
+  auto ZePlatform = Platform.get_native<backend::level_zero>();
+  auto ZeDevice = Device.get_native<backend::level_zero>();
+  auto ZeContext = Context.get_native<backend::level_zero>();
+  auto ZeQueue = Queue.get_native<backend::level_zero>();
+
+  // Re-create SYCL objects from native Level Zero handles
+  auto PlatformInterop = level_zero::make<platform>(ZePlatform);
+  auto DeviceInterop = level_zero::make<device>(PlatformInterop, ZeDevice);
+  auto ContextInterop =
+      level_zero::make<context>(PlatformInterop.get_devices(), ZeContext);
+  auto QueueInterop = level_zero::make<queue>(ContextInterop, ZeQueue);
+
+  // Check native handles
+  assert(ZePlatform == PlatformInterop.get_native<backend::level_zero>());
+  assert(ZeDevice == DeviceInterop.get_native<backend::level_zero>());
+  assert(ZeContext == ContextInterop.get_native<backend::level_zero>());
+  assert(ZeQueue == QueueInterop.get_native<backend::level_zero>());
+
+  // Verify re-created objects
+  int Arr[] = {2};
+  {
+    cl::sycl::buffer<int, 1> Buf(Arr, 1);
+    QueueInterop.submit([&](cl::sycl::handler &CGH) {
+      auto Acc = Buf.get_access<cl::sycl::access::mode::read_write>(CGH);
+      CGH.single_task<class SimpleKernel>([=]() { Acc[0] *= 3; });
+    });
+  }
+  assert(Arr[0] == 6);
+
+  return 0;
+}