[SYCL][Bindless][UR] Add support for timeline semaphores (#17395)

Add support for timeline semaphores on linux and windows for CUDA
backend and preliminary support in L0 backend. Update
vulkan_interop/sampled_images.cpp to remove redundant wait when using
semaphores.

Author: DBDuncan

sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc

@@ -2083,6 +2083,8 @@ enum class external_semaphore_handle_type {
   opaque_fd = 0,
   win32_nt_handle = 1,
   win32_nt_dx12_fence = 2,
+  timeline_fd = 3,
+  timeline_win32_nt_handle = 4,
 };
 
 // Descriptor templated on specific resource type
@@ -2134,8 +2136,9 @@ The resulting `external_semaphore` can then be used in a SYCL command
 group, to either wait until the semaphore signalled, or signal the semaphore.
 
 If the type of semaphore imported supports setting the state of discrete 
-semaphore value (the semaphore type is `win32_nt_dx12_fence`), then the user 
-can specify which value the semaphore operation should wait on, or signal.
+semaphore value (the semaphore type is `win32_nt_dx12_fence`, `timeline_fd` or
+`timeline_win32_nt_handle`), then the user can specify which value the semaphore
+operation should wait on, or signal.
 
 We propose to extend the SYCL queue and handler classes with semaphore waiting 
 and signalling operations.
@@ -2226,17 +2229,19 @@ public:
 The behaviour of waiting on a semaphore will depend on the type of the 
 semaphore which was imported.
 
-If the semaphore does not support setting of a discrete state value (the 
-semaphore type is not `win32_nt_dx12_fence`), then any operations submitted to 
-the queue after a `ext_oneapi_wait_external_semaphore` call will not begin 
-until the imported semaphore is in a signalled state. After this, the semaphore 
-will be reset to a non-signalled state.
-
-If the semaphore does support setting of a discrete state value (the semaphore 
-type is `win32_nt_dx12_fence`), then any operations submitted to the queue 
-after a `ext_oneapi_wait_external_semaphore` call will not begin until the 
-imported semaphore is in a state greater than or equal to the `wait_value`. The 
-state of this type of semaphore will not be altered by the call to 
+If the semaphore does not support setting of a discrete state value (the
+semaphore type is not `win32_nt_dx12_fence`, `timeline_fd` or
+`timeline_win32_nt_handle`), then any operations submitted to the queue after a
+`ext_oneapi_wait_external_semaphore` call will not begin until the imported
+semaphore is in a signalled state. After this, the semaphore will be reset to a
+non-signalled state.
+
+If the semaphore does support setting of a discrete state value (the semaphore
+type is `win32_nt_dx12_fence`, `timeline_fd` or `timeline_win32_nt_handle`),
+then any operations submitted to the queue after a
+`ext_oneapi_wait_external_semaphore` call will not begin until the imported
+semaphore is in a state greater than or equal to the `wait_value`. The state of
+this type of semaphore will not be altered by the call to
 `ext_oneapi_wait_external_semaphore`.
 
 When `ext_oneapi_signal_external_semaphore` is called, the external semaphore 
@@ -2570,4 +2575,5 @@ These features still need to be handled:
                    sub-copies.
                  - Add support for USM to USM copies and sub-copies.
                  - Add support for host to host copies and sub-copies.
+|6.8|2025-03-13| - Add support for importing timeline semaphores.
 |======================

sycl/include/sycl/ext/oneapi/bindless_images_interop.hpp

@@ -28,6 +28,8 @@ enum class external_semaphore_handle_type {
   opaque_fd = 0,
   win32_nt_handle = 1,
   win32_nt_dx12_fence = 2,
+  timeline_fd = 3,
+  timeline_win32_nt_handle = 4,
 };
 
 /// Opaque external memory handle type

sycl/source/detail/bindless_images.cpp

@@ -639,16 +639,31 @@ __SYCL_EXPORT external_semaphore import_external_semaphore(
   urExternalSemDesc.stype = UR_STRUCTURE_TYPE_EXP_EXTERNAL_SEMAPHORE_DESC;
   urExternalSemDesc.pNext = &urFileDescriptor;
 
-  // For this specialization of `import_external_semaphore` the handleType is
-  // always `OPAQUE_FD`.
+  // This specialization handles timeline and binary semaphores which both
+  // have the same underlying external semaphore handle type of opaque file
+  // descriptor.
+
+  // Select appropriate semaphore handle type.
+  ur_exp_external_semaphore_type_t urHandleType;
+  switch (externalSemaphoreDesc.handle_type) {
+  case external_semaphore_handle_type::opaque_fd:
+    urHandleType = UR_EXP_EXTERNAL_SEMAPHORE_TYPE_OPAQUE_FD;
+    break;
+  case external_semaphore_handle_type::timeline_fd:
+    urHandleType = UR_EXP_EXTERNAL_SEMAPHORE_TYPE_TIMELINE_FD;
+    break;
+  default:
+    throw sycl::exception(sycl::make_error_code(sycl::errc::invalid),
+                          "Invalid semaphore handle type");
+  }
+
   Adapter->call<
       sycl::errc::invalid,
       sycl::detail::UrApiKind::urBindlessImagesImportExternalSemaphoreExp>(
-      C, Device, UR_EXP_EXTERNAL_SEMAPHORE_TYPE_OPAQUE_FD, &urExternalSemDesc,
-      &urExternalSemaphore);
+      C, Device, urHandleType, &urExternalSemDesc, &urExternalSemaphore);
 
   return external_semaphore{urExternalSemaphore,
-                                  external_semaphore_handle_type::opaque_fd};
+                            externalSemaphoreDesc.handle_type};
 }
 
 template <>
@@ -688,6 +703,9 @@ __SYCL_EXPORT external_semaphore import_external_semaphore(
   case external_semaphore_handle_type::win32_nt_dx12_fence:
     urHandleType = UR_EXP_EXTERNAL_SEMAPHORE_TYPE_WIN32_NT_DX12_FENCE;
     break;
+  case external_semaphore_handle_type::timeline_win32_nt_handle:
+    urHandleType = UR_EXP_EXTERNAL_SEMAPHORE_TYPE_TIMELINE_WIN32_NT;
+    break;
   default:
     throw sycl::exception(sycl::make_error_code(sycl::errc::invalid),
                           "Invalid semaphore handle type");

sycl/source/handler.cpp

@@ -1555,17 +1555,21 @@ void handler::ext_oneapi_wait_external_semaphore(
   throwIfGraphAssociated<
       ext::oneapi::experimental::detail::UnsupportedGraphFeatures::
           sycl_ext_oneapi_bindless_images>();
-  if (ExtSemaphore.handle_type !=
-          sycl::ext::oneapi::experimental::external_semaphore_handle_type::
-              opaque_fd &&
-      ExtSemaphore.handle_type !=
-          sycl::ext::oneapi::experimental::external_semaphore_handle_type::
-              win32_nt_handle) {
+
+  switch (ExtSemaphore.handle_type) {
+  case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
+      opaque_fd:
+  case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
+      win32_nt_handle:
+    break;
+  default:
     throw sycl::exception(
         make_error_code(errc::invalid),
         "Invalid type of semaphore for this operation. The "
         "type of semaphore used needs a user passed wait value.");
+    break;
   }
+
   impl->MExternalSemaphore =
       (ur_exp_external_semaphore_handle_t)ExtSemaphore.raw_handle;
   impl->MWaitValue = {};
@@ -1578,14 +1582,23 @@ void handler::ext_oneapi_wait_external_semaphore(
   throwIfGraphAssociated<
       ext::oneapi::experimental::detail::UnsupportedGraphFeatures::
           sycl_ext_oneapi_bindless_images>();
-  if (ExtSemaphore.handle_type !=
-      sycl::ext::oneapi::experimental::external_semaphore_handle_type::
-          win32_nt_dx12_fence) {
+
+  switch (ExtSemaphore.handle_type) {
+  case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
+      win32_nt_dx12_fence:
+  case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
+      timeline_fd:
+  case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
+      timeline_win32_nt_handle:
+    break;
+  default:
     throw sycl::exception(
         make_error_code(errc::invalid),
         "Invalid type of semaphore for this operation. The "
         "type of semaphore does not support user passed wait values.");
+    break;
   }
+
   impl->MExternalSemaphore =
       (ur_exp_external_semaphore_handle_t)ExtSemaphore.raw_handle;
   impl->MWaitValue = WaitValue;
@@ -1597,17 +1610,21 @@ void handler::ext_oneapi_signal_external_semaphore(
   throwIfGraphAssociated<
       ext::oneapi::experimental::detail::UnsupportedGraphFeatures::
           sycl_ext_oneapi_bindless_images>();
-  if (ExtSemaphore.handle_type !=
-          sycl::ext::oneapi::experimental::external_semaphore_handle_type::
-              opaque_fd &&
-      ExtSemaphore.handle_type !=
-          sycl::ext::oneapi::experimental::external_semaphore_handle_type::
-              win32_nt_handle) {
+
+  switch (ExtSemaphore.handle_type) {
+  case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
+      opaque_fd:
+  case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
+      win32_nt_handle:
+    break;
+  default:
     throw sycl::exception(
         make_error_code(errc::invalid),
         "Invalid type of semaphore for this operation. The "
         "type of semaphore used needs a user passed signal value.");
+    break;
   }
+
   impl->MExternalSemaphore =
       (ur_exp_external_semaphore_handle_t)ExtSemaphore.raw_handle;
   impl->MSignalValue = {};
@@ -1620,14 +1637,23 @@ void handler::ext_oneapi_signal_external_semaphore(
   throwIfGraphAssociated<
       ext::oneapi::experimental::detail::UnsupportedGraphFeatures::
           sycl_ext_oneapi_bindless_images>();
-  if (ExtSemaphore.handle_type !=
-      sycl::ext::oneapi::experimental::external_semaphore_handle_type::
-          win32_nt_dx12_fence) {
+
+  switch (ExtSemaphore.handle_type) {
+  case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
+      win32_nt_dx12_fence:
+  case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
+      timeline_fd:
+  case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
+      timeline_win32_nt_handle:
+    break;
+  default:
     throw sycl::exception(
         make_error_code(errc::invalid),
         "Invalid type of semaphore for this operation. The "
         "type of semaphore does not support user passed signal values.");
+    break;
   }
+
   impl->MExternalSemaphore =
       (ur_exp_external_semaphore_handle_t)ExtSemaphore.raw_handle;
   impl->MSignalValue = SignalValue;

sycl/test-e2e/bindless_images/vulkan_interop/sampled_images.cpp

@@ -413,7 +413,11 @@ bool run_test(sycl::range<NDims> dims, sycl::range<NDims> localSize,
 
     VK_CHECK_CALL(vkQueueSubmit(vk_transfer_queue, 1 /*submitCount*/,
                                 &submission, VK_NULL_HANDLE /*fence*/));
+// Do not wait when using semaphores as they can handle the kernel execution
+// order.
+#ifndef TEST_SEMAPHORE_IMPORT
     VK_CHECK_CALL(vkQueueWaitIdle(vk_transfer_queue));
+#endif
   }
 
   printString("Getting memory file descriptors\n");