[SYCL] Cherry-pick UR patches related to bindless images (#20031)

This is a joined cherry-pick of:

---

[UR][L0] Remove Driver Exp Implementation of External Semaphore
(#19835)

- Intel L0 GPU Driver no longer supports the Driver Exp Implementation
of External Semaphore and the code has been removed from the codebase,
therefore the support needs to be removed from the adapter to allow
compiling with newer ze_intel_gpu.h headers.
- L0 Spec implementation is the only version required for customer
support.

Patch-by: Neil R. Spruit <[email protected]>

---

[UR][L0] urBindlessImagesGetImageMemoryHandleTypeSupportExp correction
(#19667)

VK_FORMAT_R8G8B8A8_UNORM is supported on L0 urt, correction to
verifyCommonImagePropertiesSupport.

Patch-by: Zhang, Winston <[email protected]>

---

[UR][L0][V2] Fixed supported logic for external semaphore
(#19863)

Patch-by: Neil R. Spruit <[email protected]>

---

[UR][L0] Fix Implict Event sync during external semaphore wait/signal
(#19859)

- When SYCl Calls the L0 adapter without a signal event, L0 creates an
internal event.
- Given in order command queue, the internal event for the wait/signal
is implicitly added to the wait list of the next command if
executeCommandList is called.
- This ensures that the internal event is properly synchronized with the
command queue during that next call given SYCL is failing to create a UR
event for tracking the wait/signal.

Patch-by: Neil R. Spruit <[email protected]>

Author: AlexeySachkov

sycl/test-e2e/bindless_images/dx11_interop/read_write_unsampled.cpp

@@ -213,7 +213,10 @@ int runTest(D3D11ProgramState &d3d11ProgramState, sycl::queue syclQueue,
   // E.g. LevelZero does not support `unorm` channel types.
   if (!bindless_helpers::memoryAllocationSupported(
           syclImageDesc, syclexp::image_memory_handle_type::opaque_handle,
-          syclQueue)) {
+          syclQueue) ||
+      (channelType == sycl::image_channel_type::unorm_int8 &&
+       syclQueue.get_device().get_backend() ==
+           sycl::backend::ext_oneapi_level_zero)) {
     // We cannot allocate the image memory, skip the test.
 #ifdef VERBOSE_PRINT
     std::cout << "Memory allocation unsupported. Skipping test.\n";

sycl/test-e2e/bindless_images/dx12_interop/read_write_unsampled.cpp

@@ -523,6 +523,14 @@ static bool
 runTest(DX12SYCLDevice &device, sycl::image_channel_type channelType,
         sycl::range<NDims> globalSize, sycl::range<NDims> localSize) {
 
+  // Skip unorm_int8 tests for Level Zero backend
+  if (channelType == sycl::image_channel_type::unorm_int8 &&
+      device.getSyclQueue().get_device().get_backend() ==
+          sycl::backend::ext_oneapi_level_zero) {
+    std::cout << "Skipping unorm_int8 test for Level Zero backend.\n";
+    return true;
+  }
+
   syclexp::image_descriptor syclImageDesc{globalSize, NChannels, channelType};
 
   // Verify ability to allocate the above image descriptor.

unified-runtime/source/adapters/level_zero/image.cpp

@@ -136,44 +136,17 @@ ur_result_t urBindlessImagesWaitExternalSemaphoreExp(
   const auto &ZeCommandList = CommandList->first;
   const auto &WaitList = (*Event)->WaitList;
 
-  if (UrPlatform->ZeExternalSemaphoreExt.LoaderExtension) {
-    ze_external_semaphore_wait_params_ext_t WaitParams = {
-        ZE_STRUCTURE_TYPE_EXTERNAL_SEMAPHORE_WAIT_PARAMS_EXT, nullptr, 0};
-    WaitParams.value = hasValue ? waitValue : 0;
-    ze_external_semaphore_ext_handle_t hExtSemaphore =
-        reinterpret_cast<ze_external_semaphore_ext_handle_t>(hSemaphore);
-    ZE2UR_CALL(UrPlatform->ZeExternalSemaphoreExt
-                   .zexCommandListAppendWaitExternalSemaphoresExp,
-               (ZeCommandList, 1, &hExtSemaphore, &WaitParams, ZeEvent,
-                WaitList.Length, WaitList.ZeEventList));
-  } else {
-    ze_command_list_handle_t translatedCommandList;
-    ZE2UR_CALL(zelLoaderTranslateHandle,
-               (ZEL_HANDLE_COMMAND_LIST, ZeCommandList,
-                (void **)&translatedCommandList));
-    ze_event_handle_t translatedEvent = ZeEvent;
-    if (ZeEvent) {
-      ZE2UR_CALL(zelLoaderTranslateHandle,
-                 (ZEL_HANDLE_EVENT, ZeEvent, (void **)&translatedEvent));
-    }
-    std::vector<ze_event_handle_t> EventHandles(WaitList.Length + 1, nullptr);
-    if (WaitList.Length > 0) {
-      for (size_t i = 0; i < WaitList.Length; i++) {
-        ze_event_handle_t ZeEvent = WaitList.ZeEventList[i];
-        ZE2UR_CALL(zelLoaderTranslateHandle,
-                   (ZEL_HANDLE_EVENT, ZeEvent, (void **)&EventHandles[i + 1]));
-      }
-    }
-    ze_intel_external_semaphore_wait_params_exp_t WaitParams = {
-        ZE_INTEL_STRUCTURE_TYPE_EXTERNAL_SEMAPHORE_WAIT_PARAMS_EXP, nullptr, 0};
-    WaitParams.value = hasValue ? waitValue : 0;
-    const ze_intel_external_semaphore_exp_handle_t hExtSemaphore =
-        reinterpret_cast<ze_intel_external_semaphore_exp_handle_t>(hSemaphore);
-    ZE2UR_CALL(UrPlatform->ZeExternalSemaphoreExt
-                   .zexExpCommandListAppendWaitExternalSemaphoresExp,
-               (translatedCommandList, 1, &hExtSemaphore, &WaitParams,
-                translatedEvent, WaitList.Length, EventHandles.data()));
-  }
+  ze_external_semaphore_wait_params_ext_t WaitParams = {
+      ZE_STRUCTURE_TYPE_EXTERNAL_SEMAPHORE_WAIT_PARAMS_EXT, nullptr, 0};
+  WaitParams.value = hasValue ? waitValue : 0;
+  ze_external_semaphore_ext_handle_t hExtSemaphore =
+      reinterpret_cast<ze_external_semaphore_ext_handle_t>(hSemaphore);
+  ZE2UR_CALL(UrPlatform->ZeExternalSemaphoreExt
+                 .zexCommandListAppendWaitExternalSemaphoresExp,
+             (ZeCommandList, 1, &hExtSemaphore, &WaitParams, ZeEvent,
+              WaitList.Length, WaitList.ZeEventList));
+
+  UR_CALL(hQueue->executeCommandList(CommandList, false, OkToBatch));
 
   return UR_RESULT_SUCCESS;
 }
@@ -221,47 +194,18 @@ ur_result_t urBindlessImagesSignalExternalSemaphoreExp(
   const auto &ZeCommandList = CommandList->first;
   const auto &WaitList = (*Event)->WaitList;
 
-  if (UrPlatform->ZeExternalSemaphoreExt.LoaderExtension) {
-    ze_external_semaphore_signal_params_ext_t SignalParams = {
-        ZE_STRUCTURE_TYPE_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS_EXT, nullptr, 0};
-    SignalParams.value = hasValue ? signalValue : 0;
-    ze_external_semaphore_ext_handle_t hExtSemaphore =
-        reinterpret_cast<ze_external_semaphore_ext_handle_t>(hSemaphore);
-
-    ZE2UR_CALL(UrPlatform->ZeExternalSemaphoreExt
-                   .zexCommandListAppendSignalExternalSemaphoresExp,
-               (ZeCommandList, 1, &hExtSemaphore, &SignalParams, ZeEvent,
-                WaitList.Length, WaitList.ZeEventList));
-  } else {
-    ze_intel_external_semaphore_signal_params_exp_t SignalParams = {
-        ZE_INTEL_STRUCTURE_TYPE_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS_EXP, nullptr,
-        0};
-    SignalParams.value = hasValue ? signalValue : 0;
-    const ze_intel_external_semaphore_exp_handle_t hExtSemaphore =
-        reinterpret_cast<ze_intel_external_semaphore_exp_handle_t>(hSemaphore);
-
-    ze_command_list_handle_t translatedCommandList;
-    ZE2UR_CALL(zelLoaderTranslateHandle,
-               (ZEL_HANDLE_COMMAND_LIST, ZeCommandList,
-                (void **)&translatedCommandList));
-    ze_event_handle_t translatedEvent = ZeEvent;
-    if (ZeEvent) {
-      ZE2UR_CALL(zelLoaderTranslateHandle,
-                 (ZEL_HANDLE_EVENT, ZeEvent, (void **)&translatedEvent));
-    }
-    std::vector<ze_event_handle_t> EventHandles(WaitList.Length + 1, nullptr);
-    if (WaitList.Length > 0) {
-      for (size_t i = 0; i < WaitList.Length; i++) {
-        ze_event_handle_t ZeEvent = WaitList.ZeEventList[i];
-        ZE2UR_CALL(zelLoaderTranslateHandle,
-                   (ZEL_HANDLE_EVENT, ZeEvent, (void **)&EventHandles[i + 1]));
-      }
-    }
-    ZE2UR_CALL(UrPlatform->ZeExternalSemaphoreExt
-                   .zexExpCommandListAppendSignalExternalSemaphoresExp,
-               (translatedCommandList, 1, &hExtSemaphore, &SignalParams,
-                translatedEvent, WaitList.Length, EventHandles.data()));
-  }
+  ze_external_semaphore_signal_params_ext_t SignalParams = {
+      ZE_STRUCTURE_TYPE_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS_EXT, nullptr, 0};
+  SignalParams.value = hasValue ? signalValue : 0;
+  ze_external_semaphore_ext_handle_t hExtSemaphore =
+      reinterpret_cast<ze_external_semaphore_ext_handle_t>(hSemaphore);
+
+  ZE2UR_CALL(UrPlatform->ZeExternalSemaphoreExt
+                 .zexCommandListAppendSignalExternalSemaphoresExp,
+             (ZeCommandList, 1, &hExtSemaphore, &SignalParams, ZeEvent,
+              WaitList.Length, WaitList.ZeEventList));
+
+  UR_CALL(hQueue->executeCommandList(CommandList, false, OkToBatch));
 
   return UR_RESULT_SUCCESS;
 }

unified-runtime/source/adapters/level_zero/image_common.cpp

@@ -1016,16 +1016,6 @@ bool verifyCommonImagePropertiesSupport(
     }
   }
 
-  // Verify unnormalized channel type support.
-  // LevelZero currently doesn't support unnormalized channel types.
-  switch (pImageFormat->channelType) {
-  default:
-    break;
-  case UR_IMAGE_CHANNEL_TYPE_UNORM_INT8:
-  case UR_IMAGE_CHANNEL_TYPE_UNORM_INT16:
-    return false;
-  }
-
   return supported;
 }
 
@@ -1358,126 +1348,60 @@ ur_result_t urBindlessImagesImportExternalSemaphoreExp(
                   " {} function not supported!", __FUNCTION__);
     return UR_RESULT_ERROR_UNSUPPORTED_FEATURE;
   }
-  if (UrPlatform->ZeExternalSemaphoreExt.LoaderExtension) {
-    ze_external_semaphore_ext_desc_t SemDesc = {
-        ZE_STRUCTURE_TYPE_EXTERNAL_SEMAPHORE_EXT_DESC, nullptr,
-        ZE_EXTERNAL_SEMAPHORE_EXT_FLAG_OPAQUE_FD};
-    ze_external_semaphore_ext_handle_t ExtSemaphoreHandle;
-    ze_external_semaphore_fd_ext_desc_t FDExpDesc = {
-        ZE_STRUCTURE_TYPE_EXTERNAL_SEMAPHORE_FD_EXT_DESC, nullptr, 0};
-    ze_external_semaphore_win32_ext_desc_t Win32ExpDesc = {
-        ZE_STRUCTURE_TYPE_EXTERNAL_SEMAPHORE_WIN32_EXT_DESC, nullptr, nullptr,
-        nullptr};
-    void *pNext = const_cast<void *>(pExternalSemaphoreDesc->pNext);
-    while (pNext != nullptr) {
-      const ur_base_desc_t *BaseDesc =
-          static_cast<const ur_base_desc_t *>(pNext);
-      if (BaseDesc->stype == UR_STRUCTURE_TYPE_EXP_FILE_DESCRIPTOR) {
-        auto FileDescriptor =
-            static_cast<const ur_exp_file_descriptor_t *>(pNext);
-        FDExpDesc.fd = FileDescriptor->fd;
-        SemDesc.pNext = &FDExpDesc;
-        switch (semHandleType) {
-        case UR_EXP_EXTERNAL_SEMAPHORE_TYPE_OPAQUE_FD:
-          SemDesc.flags = ZE_EXTERNAL_SEMAPHORE_EXT_FLAG_OPAQUE_FD;
-          break;
-        case UR_EXP_EXTERNAL_SEMAPHORE_TYPE_TIMELINE_FD:
-          SemDesc.flags =
-              ZE_EXTERNAL_SEMAPHORE_EXT_FLAG_VK_TIMELINE_SEMAPHORE_FD;
-          break;
-        default:
-          return UR_RESULT_ERROR_INVALID_VALUE;
-        }
-      } else if (BaseDesc->stype == UR_STRUCTURE_TYPE_EXP_WIN32_HANDLE) {
-        SemDesc.pNext = &Win32ExpDesc;
-        auto Win32Handle = static_cast<const ur_exp_win32_handle_t *>(pNext);
-        switch (semHandleType) {
-        case UR_EXP_EXTERNAL_SEMAPHORE_TYPE_WIN32_NT:
-          SemDesc.flags = ZE_EXTERNAL_SEMAPHORE_EXT_FLAG_OPAQUE_WIN32;
-          break;
-        case UR_EXP_EXTERNAL_SEMAPHORE_TYPE_WIN32_NT_DX12_FENCE:
-          SemDesc.flags = ZE_EXTERNAL_SEMAPHORE_EXT_FLAG_D3D12_FENCE;
-          break;
-        case UR_EXP_EXTERNAL_SEMAPHORE_TYPE_TIMELINE_WIN32_NT:
-          SemDesc.flags =
-              ZE_EXTERNAL_SEMAPHORE_EXT_FLAG_VK_TIMELINE_SEMAPHORE_WIN32;
-          break;
-        default:
-          return UR_RESULT_ERROR_INVALID_VALUE;
-        }
-        Win32ExpDesc.handle = Win32Handle->handle;
+  ze_external_semaphore_ext_desc_t SemDesc = {
+      ZE_STRUCTURE_TYPE_EXTERNAL_SEMAPHORE_EXT_DESC, nullptr,
+      ZE_EXTERNAL_SEMAPHORE_EXT_FLAG_OPAQUE_FD};
+  ze_external_semaphore_ext_handle_t ExtSemaphoreHandle;
+  ze_external_semaphore_fd_ext_desc_t FDExpDesc = {
+      ZE_STRUCTURE_TYPE_EXTERNAL_SEMAPHORE_FD_EXT_DESC, nullptr, 0};
+  ze_external_semaphore_win32_ext_desc_t Win32ExpDesc = {
+      ZE_STRUCTURE_TYPE_EXTERNAL_SEMAPHORE_WIN32_EXT_DESC, nullptr, nullptr,
+      nullptr};
+  void *pNext = const_cast<void *>(pExternalSemaphoreDesc->pNext);
+  while (pNext != nullptr) {
+    const ur_base_desc_t *BaseDesc = static_cast<const ur_base_desc_t *>(pNext);
+    if (BaseDesc->stype == UR_STRUCTURE_TYPE_EXP_FILE_DESCRIPTOR) {
+      auto FileDescriptor =
+          static_cast<const ur_exp_file_descriptor_t *>(pNext);
+      FDExpDesc.fd = FileDescriptor->fd;
+      SemDesc.pNext = &FDExpDesc;
+      switch (semHandleType) {
+      case UR_EXP_EXTERNAL_SEMAPHORE_TYPE_OPAQUE_FD:
+        SemDesc.flags = ZE_EXTERNAL_SEMAPHORE_EXT_FLAG_OPAQUE_FD;
+        break;
+      case UR_EXP_EXTERNAL_SEMAPHORE_TYPE_TIMELINE_FD:
+        SemDesc.flags = ZE_EXTERNAL_SEMAPHORE_EXT_FLAG_VK_TIMELINE_SEMAPHORE_FD;
+        break;
+      default:
+        return UR_RESULT_ERROR_INVALID_VALUE;
       }
-      pNext = const_cast<void *>(BaseDesc->pNext);
-    }
-    ZE2UR_CALL(UrPlatform->ZeExternalSemaphoreExt.zexImportExternalSemaphoreExp,
-               (hDevice->ZeDevice, &SemDesc, &ExtSemaphoreHandle));
-    *phExternalSemaphoreHandle =
-        (ur_exp_external_semaphore_handle_t)ExtSemaphoreHandle;
-
-  } else {
-    ze_intel_external_semaphore_exp_desc_t SemDesc = {
-        ZE_INTEL_STRUCTURE_TYPE_EXTERNAL_SEMAPHORE_EXP_DESC, nullptr,
-        ZE_EXTERNAL_SEMAPHORE_EXP_FLAGS_OPAQUE_FD};
-    ze_intel_external_semaphore_exp_handle_t ExtSemaphoreHandle;
-    ze_intel_external_semaphore_desc_fd_exp_desc_t FDExpDesc = {
-        ZE_INTEL_STRUCTURE_TYPE_EXTERNAL_SEMAPHORE_FD_EXP_DESC, nullptr, 0};
-    _ze_intel_external_semaphore_win32_exp_desc_t Win32ExpDesc = {
-        ZE_INTEL_STRUCTURE_TYPE_EXTERNAL_SEMAPHORE_WIN32_EXP_DESC, nullptr,
-        nullptr, nullptr};
-    void *pNext = const_cast<void *>(pExternalSemaphoreDesc->pNext);
-    while (pNext != nullptr) {
-      const ur_base_desc_t *BaseDesc =
-          static_cast<const ur_base_desc_t *>(pNext);
-      if (BaseDesc->stype == UR_STRUCTURE_TYPE_EXP_FILE_DESCRIPTOR) {
-        auto FileDescriptor =
-            static_cast<const ur_exp_file_descriptor_t *>(pNext);
-        FDExpDesc.fd = FileDescriptor->fd;
-        SemDesc.pNext = &FDExpDesc;
-        switch (semHandleType) {
-        case UR_EXP_EXTERNAL_SEMAPHORE_TYPE_OPAQUE_FD:
-          SemDesc.flags = ZE_EXTERNAL_SEMAPHORE_EXP_FLAGS_OPAQUE_FD;
-          break;
-        case UR_EXP_EXTERNAL_SEMAPHORE_TYPE_TIMELINE_FD:
-          SemDesc.flags = ZE_EXTERNAL_SEMAPHORE_EXP_FLAGS_TIMELINE_SEMAPHORE_FD;
-          break;
-        default:
-          return UR_RESULT_ERROR_INVALID_VALUE;
-        }
-      } else if (BaseDesc->stype == UR_STRUCTURE_TYPE_EXP_WIN32_HANDLE) {
-        SemDesc.pNext = &Win32ExpDesc;
-        auto Win32Handle = static_cast<const ur_exp_win32_handle_t *>(pNext);
-        switch (semHandleType) {
-        case UR_EXP_EXTERNAL_SEMAPHORE_TYPE_WIN32_NT:
-          SemDesc.flags = ZE_EXTERNAL_SEMAPHORE_EXP_FLAGS_OPAQUE_WIN32;
-          break;
-        case UR_EXP_EXTERNAL_SEMAPHORE_TYPE_WIN32_NT_DX12_FENCE:
-          SemDesc.flags = ZE_EXTERNAL_SEMAPHORE_EXP_FLAGS_D3D12_FENCE;
-          break;
-        case UR_EXP_EXTERNAL_SEMAPHORE_TYPE_TIMELINE_WIN32_NT:
-          SemDesc.flags =
-              ZE_EXTERNAL_SEMAPHORE_EXP_FLAGS_TIMELINE_SEMAPHORE_WIN32;
-          break;
-        default:
-          return UR_RESULT_ERROR_INVALID_VALUE;
-        }
-        Win32ExpDesc.handle = Win32Handle->handle;
+    } else if (BaseDesc->stype == UR_STRUCTURE_TYPE_EXP_WIN32_HANDLE) {
+      SemDesc.pNext = &Win32ExpDesc;
+      auto Win32Handle = static_cast<const ur_exp_win32_handle_t *>(pNext);
+      switch (semHandleType) {
+      case UR_EXP_EXTERNAL_SEMAPHORE_TYPE_WIN32_NT:
+        SemDesc.flags = ZE_EXTERNAL_SEMAPHORE_EXT_FLAG_OPAQUE_WIN32;
+        break;
+      case UR_EXP_EXTERNAL_SEMAPHORE_TYPE_WIN32_NT_DX12_FENCE:
+        SemDesc.flags = ZE_EXTERNAL_SEMAPHORE_EXT_FLAG_D3D12_FENCE;
+        break;
+      case UR_EXP_EXTERNAL_SEMAPHORE_TYPE_TIMELINE_WIN32_NT:
+        SemDesc.flags =
+            ZE_EXTERNAL_SEMAPHORE_EXT_FLAG_VK_TIMELINE_SEMAPHORE_WIN32;
+        break;
+      default:
+        return UR_RESULT_ERROR_INVALID_VALUE;
       }
-      pNext = const_cast<void *>(BaseDesc->pNext);
+      Win32ExpDesc.handle = Win32Handle->handle;
     }
-
-    ze_device_handle_t translatedDevice;
-    ZE2UR_CALL(zelLoaderTranslateHandle, (ZEL_HANDLE_DEVICE, hDevice->ZeDevice,
-                                          (void **)&translatedDevice));
-    // If the L0 loader is not aware of the extension, the handles need to be
-    // translated
-    ZE2UR_CALL(
-        UrPlatform->ZeExternalSemaphoreExt.zexExpImportExternalSemaphoreExp,
-        (translatedDevice, &SemDesc, &ExtSemaphoreHandle));
-
-    *phExternalSemaphoreHandle =
-        (ur_exp_external_semaphore_handle_t)ExtSemaphoreHandle;
+    pNext = const_cast<void *>(BaseDesc->pNext);
   }
 
+  ZE2UR_CALL(UrPlatform->ZeExternalSemaphoreExt.zexImportExternalSemaphoreExp,
+             (hDevice->ZeDevice, &SemDesc, &ExtSemaphoreHandle));
+  *phExternalSemaphoreHandle =
+      (ur_exp_external_semaphore_handle_t)ExtSemaphoreHandle;
+
   return UR_RESULT_SUCCESS;
 }
 
@@ -1490,15 +1414,9 @@ ur_result_t urBindlessImagesReleaseExternalSemaphoreExp(
                   " {} function not supported!", __FUNCTION__);
     return UR_RESULT_ERROR_UNSUPPORTED_FEATURE;
   }
-  if (UrPlatform->ZeExternalSemaphoreExt.LoaderExtension) {
-    ZE2UR_CALL(
-        UrPlatform->ZeExternalSemaphoreExt.zexDeviceReleaseExternalSemaphoreExp,
-        ((ze_external_semaphore_ext_handle_t)hExternalSemaphore));
-  } else {
-    ZE2UR_CALL(UrPlatform->ZeExternalSemaphoreExt
-                   .zexExpDeviceReleaseExternalSemaphoreExp,
-               ((ze_intel_external_semaphore_exp_handle_t)hExternalSemaphore));
-  }
+  ZE2UR_CALL(
+      UrPlatform->ZeExternalSemaphoreExt.zexDeviceReleaseExternalSemaphoreExp,
+      ((ze_external_semaphore_ext_handle_t)hExternalSemaphore));
 
   return UR_RESULT_SUCCESS;
 }