From 11ccf692902b6aa38e8e408c46ddb61f19dd5afb Mon Sep 17 00:00:00 2001 From: Georgij Tsarin Date: Wed, 7 Oct 2026 16:43:51 +0300 Subject: [PATCH 1/3] [SYCL][E2E][Doc] Test external semaphores on regular command lists --- .../sycl_ext_oneapi_bindless_images.asciidoc | 15 +- .../dx11_interop/read_write_unsampled.cpp | 4 +- .../D3D12_sycl_buffer_timeline_semaphore.cpp | 5 +- .../D3D12_sycl_interop_1D_read.cpp | 5 +- .../D3D12_sycl_interop_1D_write_unsampled.cpp | 5 +- .../D3D12_sycl_interop_2D_arithmetic.cpp | 5 +- .../D3D12_sycl_interop_2D_read.cpp | 5 +- .../D3D12_sycl_interop_2D_write_unsampled.cpp | 5 +- .../D3D12_sycl_interop_3D_read.cpp | 5 +- .../D3D12_sycl_interop_3D_write_unsampled.cpp | 5 +- .../D3D12_win32_named_semaphore.cpp | 3 +- .../external_semaphore_regular_cl.cpp | 145 ++++++++++++++++ .../external_semaphore_regular_cl_fails.cpp | 82 --------- ...example_6_import_memory_and_semaphores.cpp | 8 +- .../external_semaphore_regular_cl.cpp | 158 ++++++++++++++++++ .../external_semaphore_regular_cl_fails.cpp | 86 ---------- .../vulkan_sycl_2d_arithmetic.cpp | 5 +- .../vulkan_interop/vulkan_sycl_buffer.cpp | 5 +- .../vulkan_sycl_buffer_binary_semaphore.cpp | 5 +- .../vulkan_sycl_buffer_timeline_semaphore.cpp | 5 +- .../vulkan_sycl_image_interop_read_1d.cpp | 5 +- .../vulkan_sycl_image_interop_read_2d.cpp | 5 +- .../vulkan_sycl_image_interop_read_3d.cpp | 5 +- ..._sycl_image_interop_write_1d_unsampled.cpp | 5 +- ..._sycl_image_interop_write_2d_unsampled.cpp | 5 +- ..._sycl_image_interop_write_3d_unsampled.cpp | 5 +- ...ycl_image_unsampled_timeline_semaphore.cpp | 4 +- 27 files changed, 353 insertions(+), 242 deletions(-) create mode 100644 sycl/test-e2e/bindless_images/dx12_interop/external_semaphore_regular_cl.cpp delete mode 100644 sycl/test-e2e/bindless_images/dx12_interop/external_semaphore_regular_cl_fails.cpp create mode 100644 sycl/test-e2e/bindless_images/vulkan_interop/external_semaphore_regular_cl.cpp delete mode 100644 sycl/test-e2e/bindless_images/vulkan_interop/external_semaphore_regular_cl_fails.cpp diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc index 75631847b3311..5f194a52581f6 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc @@ -2333,17 +2333,12 @@ depending on the API and operating system, so do external semaphore resource handles. It is important to note that the use of imported external semaphore objects -within SYCL requires the SYCL queue to have been constructed with both of the -following properties: - -* `sycl::property::queue::in_order` -- the semaphore synchronization mechanism - is not supported on the default out-of-order queues. -* `sycl::ext::intel::property::queue::immediate_command_list` -- external - semaphore operations are only supported on queues backed by immediate - command lists. This restriction might be lifted in the future. +within SYCL requires the SYCL queue to have been constructed with the +`sycl::property::queue::in_order` property. The semaphore synchronization +mechanism is not supported on the default out-of-order queues. Use of the semaphore synchronization mechanism with a SYCL queue that was not -constructed with both of these properties will result in undefined behaviour. +constructed with this property will result in undefined behaviour. External semaphore import is facilitated through the following proposed descriptor struct. @@ -2850,4 +2845,6 @@ This query should be added in a later revision of the proposal. - Add `supports_importing_handle_type` |6.12|2025-07-03| - Add `win32_nt_dx11_resource` enum entry to `external_mem_handle_type` for Direct3D 11 memory interop. +|6.13|2026-10-07| - Allow external semaphore operations on in-order queues + backed by regular command lists. |====================== diff --git a/sycl/test-e2e/bindless_images/dx11_interop/read_write_unsampled.cpp b/sycl/test-e2e/bindless_images/dx11_interop/read_write_unsampled.cpp index bfab497821055..15748ef302ee3 100644 --- a/sycl/test-e2e/bindless_images/dx11_interop/read_write_unsampled.cpp +++ b/sycl/test-e2e/bindless_images/dx11_interop/read_write_unsampled.cpp @@ -475,8 +475,8 @@ int runTest(D3D11ProgramState &d3d11ProgramState, sycl::queue syclQueue, int main() { // Create SYCL queue, relying on SYCL device selection. #ifdef TEST_SEMAPHORE_IMPORT - // External semaphore ops require an in-order queue backed by immediate - // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). + // External semaphore operations require an in-order queue. This test also + // uses immediate command lists for its existing configuration. sycl::queue syclQueue{ {sycl::property::queue::in_order{}, sycl::ext::intel::property::queue::immediate_command_list{}}}; diff --git a/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_timeline_semaphore.cpp b/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_timeline_semaphore.cpp index 94d793d7f4255..eb7a755edc4c8 100644 --- a/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_timeline_semaphore.cpp +++ b/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_timeline_semaphore.cpp @@ -111,9 +111,8 @@ int main(int argc, char **argv) { // SYCL INTEROP try { - // Bindless image interop requires an in-order queue (per spec). External - // semaphore ops additionally require immediate command lists; see - // sycl_ext_oneapi_bindless_images.asciidoc. + // External semaphore operations require an in-order queue. This test also + // uses immediate command lists for its existing configuration. sycl::property_list qProps = useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, sycl::ext::intel::property::queue:: diff --git a/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_1D_read.cpp b/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_1D_read.cpp index 900b4e81fcb75..03b8d46dc392c 100644 --- a/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_1D_read.cpp +++ b/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_1D_read.cpp @@ -192,9 +192,8 @@ int runTest( // SYCL Import and Verification try { - // Bindless image interop requires an in-order queue (per spec). External - // semaphore ops additionally require immediate command lists; see - // sycl_ext_oneapi_bindless_images.asciidoc. + // External semaphore operations require an in-order queue. This test also + // uses immediate command lists for its existing configuration. sycl::property_list qProps = useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, sycl::ext::intel::property::queue:: diff --git a/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_1D_write_unsampled.cpp b/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_1D_write_unsampled.cpp index 62242d37917aa..6fcdd5d6305fc 100644 --- a/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_1D_write_unsampled.cpp +++ b/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_1D_write_unsampled.cpp @@ -149,9 +149,8 @@ int runTest( } try { - // Bindless image interop requires an in-order queue (per spec). External - // semaphore ops additionally require immediate command lists; see - // sycl_ext_oneapi_bindless_images.asciidoc. + // External semaphore operations require an in-order queue. This test also + // uses immediate command lists for its existing configuration. sycl::property_list qProps = useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, sycl::ext::intel::property::queue:: diff --git a/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_2D_arithmetic.cpp b/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_2D_arithmetic.cpp index 8479ad43ea030..8f11dcb0e6062 100644 --- a/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_2D_arithmetic.cpp +++ b/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_2D_arithmetic.cpp @@ -408,9 +408,8 @@ int runTest( signalExportableFence(ctx, extFenceB); try { - // Bindless image interop requires an in-order queue (per spec). External - // semaphore ops additionally require immediate command lists; see - // sycl_ext_oneapi_bindless_images.asciidoc. + // External semaphore operations require an in-order queue. This test also + // uses immediate command lists for its existing configuration. sycl::property_list qProps = useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, sycl::ext::intel::property::queue:: diff --git a/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_2D_read.cpp b/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_2D_read.cpp index cc986f87a57d1..c752e3db11ec0 100644 --- a/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_2D_read.cpp +++ b/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_2D_read.cpp @@ -206,9 +206,8 @@ int runTest( // SYCL Import and Verification try { - // Bindless image interop requires an in-order queue (per spec). External - // semaphore ops additionally require immediate command lists; see - // sycl_ext_oneapi_bindless_images.asciidoc. + // External semaphore operations require an in-order queue. This test also + // uses immediate command lists for its existing configuration. sycl::property_list qProps = useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, sycl::ext::intel::property::queue:: diff --git a/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_2D_write_unsampled.cpp b/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_2D_write_unsampled.cpp index 9f263e371059d..c8d57422696de 100644 --- a/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_2D_write_unsampled.cpp +++ b/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_2D_write_unsampled.cpp @@ -190,9 +190,8 @@ int runTest( } try { - // Bindless image interop requires an in-order queue (per spec). External - // semaphore ops additionally require immediate command lists; see - // sycl_ext_oneapi_bindless_images.asciidoc. + // External semaphore operations require an in-order queue. This test also + // uses immediate command lists for its existing configuration. sycl::property_list qProps = useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, sycl::ext::intel::property::queue:: diff --git a/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_3D_read.cpp b/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_3D_read.cpp index 1ccf550ab7787..e0a56632601f0 100644 --- a/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_3D_read.cpp +++ b/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_3D_read.cpp @@ -199,9 +199,8 @@ int runTest( // SYCL Import and Verification try { - // Bindless image interop requires an in-order queue (per spec). External - // semaphore ops additionally require immediate command lists; see - // sycl_ext_oneapi_bindless_images.asciidoc. + // External semaphore operations require an in-order queue. This test also + // uses immediate command lists for its existing configuration. sycl::property_list qProps = useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, sycl::ext::intel::property::queue:: diff --git a/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_3D_write_unsampled.cpp b/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_3D_write_unsampled.cpp index b9bb94de20b3a..91f45ed507f60 100644 --- a/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_3D_write_unsampled.cpp +++ b/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_3D_write_unsampled.cpp @@ -152,9 +152,8 @@ int runTest( } try { - // Bindless image interop requires an in-order queue (per spec). External - // semaphore ops additionally require immediate command lists; see - // sycl_ext_oneapi_bindless_images.asciidoc. + // External semaphore operations require an in-order queue. This test also + // uses immediate command lists for its existing configuration. sycl::property_list qProps = useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, sycl::ext::intel::property::queue:: diff --git a/sycl/test-e2e/bindless_images/dx12_interop/D3D12_win32_named_semaphore.cpp b/sycl/test-e2e/bindless_images/dx12_interop/D3D12_win32_named_semaphore.cpp index ffec32957a4c6..dab88e43d4b8e 100644 --- a/sycl/test-e2e/bindless_images/dx12_interop/D3D12_win32_named_semaphore.cpp +++ b/sycl/test-e2e/bindless_images/dx12_interop/D3D12_win32_named_semaphore.cpp @@ -114,7 +114,8 @@ int main(int argc, char **argv) { executeAndWait(d3dCtx); try { - // Semaphore ops require an in-order queue with immediate CLs. + // External semaphore operations require an in-order queue. This test also + // uses immediate command lists for its existing configuration. sycl::property_list qProps{ sycl::property::queue::in_order{}, sycl::ext::intel::property::queue::immediate_command_list{}}; diff --git a/sycl/test-e2e/bindless_images/dx12_interop/external_semaphore_regular_cl.cpp b/sycl/test-e2e/bindless_images/dx12_interop/external_semaphore_regular_cl.cpp new file mode 100644 index 0000000000000..70e32797027b3 --- /dev/null +++ b/sycl/test-e2e/bindless_images/dx12_interop/external_semaphore_regular_cl.cpp @@ -0,0 +1,145 @@ +// REQUIRES: aspect-ext_oneapi_external_semaphore_import, aspect-usm_shared_allocations +// REQUIRES: windows, level_zero + +// Regular command-list external semaphore support requires this driver +// version. +// REQUIRES-INTEL-DRIVER: win: 101.9030 + +// XFAIL: windows && run-mode && gpu-intel-gen12 +// XFAIL-TRACKER: https://github.com/intel/llvm/issues/23249 + +// RUN: %{build} %link-directx -o %t.exe %if target-spir %{ -Wno-ignored-attributes %} +// RUN: env UR_L0_BATCH_SIZE=8 %{run} %t.exe + +// Verify that external semaphore wait and signal operations work on an +// in-order queue backed by regular (non-immediate) command lists. + +#include "d3d12_setup.hpp" +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#define WIN32_LEAN_AND_MEAN +#include + +namespace syclexp = sycl::ext::oneapi::experimental; + +[[noreturn]] void failWithoutCleanup(const std::string &message) { + std::cerr << message << std::endl; + std::_Exit(EXIT_FAILURE); +} + +using MarkerAtomicRef = + sycl::atomic_ref; + +bool waitForMarker(int *marker) { + const auto deadline = + std::chrono::steady_clock::now() + std::chrono::seconds(5); + while (MarkerAtomicRef(*marker).load() == 0) { + if (std::chrono::steady_clock::now() >= deadline) + return false; + std::this_thread::sleep_for(std::chrono::milliseconds(1)); + } + return true; +} + +bool remainsIncomplete(const sycl::event &event) { + const auto deadline = + std::chrono::steady_clock::now() + std::chrono::milliseconds(100); + do { + if (event.get_info() == + sycl::info::event_command_status::complete) + return false; + std::this_thread::sleep_for(std::chrono::milliseconds(1)); + } while (std::chrono::steady_clock::now() < deadline); + return true; +} + +bool waitForCompletion(const sycl::event &event) { + const auto deadline = + std::chrono::steady_clock::now() + std::chrono::seconds(5); + while (event.get_info() != + sycl::info::event_command_status::complete) { + if (std::chrono::steady_clock::now() >= deadline) + return false; + std::this_thread::sleep_for(std::chrono::milliseconds(1)); + } + return true; +} + +int main() { + sycl::queue q{ + {sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue::no_immediate_command_list{}}}; + auto device = q.get_device(); + auto context = q.get_context(); + + D3D12Context d3dCtx = createD3D12Context(); + D3D12ExportableFence extFence = createExportableFence(d3dCtx); + + auto semDesc = + syclexp::external_semaphore_descriptor{ + extFence.sharedHandle, + syclexp::external_semaphore_handle_type::win32_nt_dx12_fence}; + syclexp::external_semaphore syclSem = + syclexp::import_external_semaphore(semDesc, device, context); + int *marker = sycl::malloc_shared(1, q); + MarkerAtomicRef(*marker).store(0); + + try { + constexpr uint64_t D3DSignalValue = 1; + q.single_task([=]() { MarkerAtomicRef(*marker).store(1); }); + sycl::event waitEvent = + q.ext_oneapi_wait_external_semaphore(syclSem, D3DSignalValue); + + if (!waitForMarker(marker)) + failWithoutCleanup( + "The external semaphore wait did not submit the preceding batch."); + if (!remainsIncomplete(waitEvent)) + failWithoutCleanup( + "The external semaphore wait completed before D3D12 signaled it."); + + signalExportableFence(d3dCtx, extFence); + if (!waitForCompletion(waitEvent)) + failWithoutCleanup( + "The external semaphore wait did not complete after D3D12 signaled " + "it."); + + constexpr uint64_t SyclSignalValue = 2; + q.ext_oneapi_signal_external_semaphore(syclSem, SyclSignalValue); + + ThrowIfFailed(extFence.fence->SetEventOnCompletion(SyclSignalValue, + d3dCtx.fenceEvent), + "Failed to wait for shared fence"); + constexpr DWORD TimeoutMs = 5000; + DWORD waitResult = WaitForSingleObject(d3dCtx.fenceEvent, TimeoutMs); + if (waitResult == WAIT_TIMEOUT) + failWithoutCleanup( + "Timed out waiting for the regular command-list semaphore signal."); + if (waitResult != WAIT_OBJECT_0) + failWithoutCleanup("Waiting for the external semaphore failed with " + "error " + + std::to_string(GetLastError())); + + q.wait_and_throw(); + } catch (const std::exception &e) { + failWithoutCleanup(e.what()); + } + + sycl::free(marker, q); + syclexp::release_external_semaphore(syclSem, device, context); + cleanupExportableFence(extFence); + if (d3dCtx.fenceEvent) + CloseHandle(d3dCtx.fenceEvent); + return 0; +} diff --git a/sycl/test-e2e/bindless_images/dx12_interop/external_semaphore_regular_cl_fails.cpp b/sycl/test-e2e/bindless_images/dx12_interop/external_semaphore_regular_cl_fails.cpp deleted file mode 100644 index 291605229d99b..0000000000000 --- a/sycl/test-e2e/bindless_images/dx12_interop/external_semaphore_regular_cl_fails.cpp +++ /dev/null @@ -1,82 +0,0 @@ - -// -// REQUIRES: aspect-ext_oneapi_external_semaphore_import, windows, level_zero -// -// RUN: %{build} %link-directx -o %t.exe %if target-spir %{ -Wno-ignored-attributes %} -// RUN: %{run} %t.exe - -// XFAIL: windows && run-mode && !gpu-intel-gen12 -// XFAIL-TRACKER: https://github.com/intel/llvm/issues/23249 - -// Waiting on a DX12-fence external semaphore from a queue backed by a -// regular (non-immediate) command list must throw sycl::exception. -// -// Mirrors vulkan_interop/external_semaphore_regular_cl_fails.cpp but uses -// the win32_nt_dx12_fence handle-type path. The Level Zero adapter -// rejects external_semaphore wait/signal at submission time when the -// queue is not using immediate command lists; this test verifies the -// rejection still fires for the DX12 fence handle type. -// -// Flow: -// 1. Create a real, exportable D3D12 timeline fence and signal it. -// 2. Import it into SYCL via a (lawful) immediate-CL queue. -// 3. Call ext_oneapi_wait_external_semaphore on a separate queue that -// explicitly opts into no_immediate_command_list, and expect a -// sycl::exception. - -#include "d3d12_setup.hpp" -#include -#include -#include -#include - -#define WIN32_LEAN_AND_MEAN -#include - -namespace syclexp = sycl::ext::oneapi::experimental; - -int main() { - D3D12Context d3dCtx = createD3D12Context(); - D3D12ExportableFence extFence = createExportableFence(d3dCtx); - - // Make the fence reach a known value so wait(value=1) is satisfiable - // if it ever gets that far. - signalExportableFence(d3dCtx, extFence); - - // Lawful queue: import the semaphore here. - sycl::queue immQ{ - {sycl::property::queue::in_order{}, - sycl::ext::intel::property::queue::immediate_command_list{}}}; - auto device = immQ.get_device(); - auto context = immQ.get_context(); - - auto semDesc = - syclexp::external_semaphore_descriptor{ - extFence.sharedHandle, - syclexp::external_semaphore_handle_type::win32_nt_dx12_fence}; - syclexp::external_semaphore syclSem = - syclexp::import_external_semaphore(semDesc, device, context); - - // The non-immediate-CL queue is what should trigger rejection on use. - sycl::queue regQ{ - context, device, - sycl::property_list{ - sycl::property::queue::in_order{}, - sycl::ext::intel::property::queue::no_immediate_command_list{}}}; - - int ret = 1; - try { - regQ.ext_oneapi_wait_external_semaphore(syclSem, extFence.fenceValue); - regQ.wait_and_throw(); - std::cerr << "FAIL: ext_oneapi_wait_external_semaphore (dx12_fence) on a " - "non-immediate-CL queue did not throw." - << std::endl; - } catch (const sycl::exception &e) { - std::cout << "Got expected sycl::exception: " << e.what() << std::endl; - ret = 0; - } - - syclexp::release_external_semaphore(syclSem, device, context); - cleanupExportableFence(extFence); - return ret; -} diff --git a/sycl/test-e2e/bindless_images/examples/example_6_import_memory_and_semaphores.cpp b/sycl/test-e2e/bindless_images/examples/example_6_import_memory_and_semaphores.cpp index 0b91824cf9bc8..658f6f903a78c 100644 --- a/sycl/test-e2e/bindless_images/examples/example_6_import_memory_and_semaphores.cpp +++ b/sycl/test-e2e/bindless_images/examples/example_6_import_memory_and_semaphores.cpp @@ -14,13 +14,9 @@ #include int main() { - // Set up queue. - // External semaphore ops require an in-order queue backed by immediate - // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). + // External semaphore operations require an in-order queue. sycl::device dev; - sycl::queue q(dev, - {sycl::property::queue::in_order{}, - sycl::ext::intel::property::queue::immediate_command_list{}}); + sycl::queue q(dev, sycl::property_list{sycl::property::queue::in_order{}}); size_t width = 123 /* passed from external API */; size_t height = 123 /* passed from external API */; diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/external_semaphore_regular_cl.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/external_semaphore_regular_cl.cpp new file mode 100644 index 0000000000000..81e83137c418e --- /dev/null +++ b/sycl/test-e2e/bindless_images/vulkan_interop/external_semaphore_regular_cl.cpp @@ -0,0 +1,158 @@ +// REQUIRES: aspect-ext_oneapi_external_semaphore_import +// REQUIRES: aspect-usm_shared_allocations +// REQUIRES: vulkan && level_zero + +// Regular command-list external semaphore support requires these driver +// versions. +// REQUIRES-INTEL-DRIVER: lin: 39758, win: 101.9030 + +// XFAIL: windows && run-mode && gpu-intel-gen12 +// XFAIL-TRACKER: https://github.com/intel/llvm/issues/23249 + +// RUN: %{build} %link-vulkan -o %t.out %if target-spir %{ -Wno-ignored-attributes %} +// RUN: env UR_L0_BATCH_SIZE=8 %{run} %t.out + +// Verify that external semaphore wait and signal operations work on an +// in-order queue backed by regular (non-immediate) command lists. A timeline +// semaphore provides equivalent coverage on both platforms while avoiding +// Linux binary semaphore sharing issue CMPLRLLVM-78008. + +#include "sycl_vulkan_setup.hpp" +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace syclexp = sycl::ext::oneapi::experimental; + +[[noreturn]] void failWithoutCleanup(const std::string &message) { + std::cerr << message << std::endl; + std::_Exit(EXIT_FAILURE); +} + +using MarkerAtomicRef = + sycl::atomic_ref; + +bool waitForMarker(int *marker) { + const auto deadline = + std::chrono::steady_clock::now() + std::chrono::seconds(5); + while (MarkerAtomicRef(*marker).load() == 0) { + if (std::chrono::steady_clock::now() >= deadline) + return false; + std::this_thread::sleep_for(std::chrono::milliseconds(1)); + } + return true; +} + +bool remainsIncomplete(const sycl::event &event) { + const auto deadline = + std::chrono::steady_clock::now() + std::chrono::milliseconds(100); + do { + if (event.get_info() == + sycl::info::event_command_status::complete) + return false; + std::this_thread::sleep_for(std::chrono::milliseconds(1)); + } while (std::chrono::steady_clock::now() < deadline); + return true; +} + +bool waitForCompletion(const sycl::event &event) { + const auto deadline = + std::chrono::steady_clock::now() + std::chrono::seconds(5); + while (event.get_info() != + sycl::info::event_command_status::complete) { + if (std::chrono::steady_clock::now() >= deadline) + return false; + std::this_thread::sleep_for(std::chrono::milliseconds(1)); + } + return true; +} + +int main() { + constexpr uint64_t VulkanSignalValue = 1; + constexpr uint64_t SyclSignalValue = 2; + + sycl::queue q{ + {sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue::no_immediate_command_list{}}}; + auto device = q.get_device(); + auto context = q.get_context(); + + VulkanContext vkCtx = createSyclVulkanContext(device); + VkSemaphore vkSem = createExportableTimelineSemaphore(vkCtx); +#ifdef _WIN32 + HANDLE semHandle = getSemaphoreHandle(vkCtx, vkSem); + syclexp::external_semaphore_descriptor desc{ + semHandle, + syclexp::external_semaphore_handle_type::timeline_win32_nt_handle}; +#else + int semFd = getSemaphoreFd(vkCtx, vkSem); + syclexp::external_semaphore_descriptor desc{ + semFd, syclexp::external_semaphore_handle_type::timeline_fd}; +#endif + syclexp::external_semaphore syclSem = + syclexp::import_external_semaphore(desc, device, context); + int *marker = sycl::malloc_shared(1, q); + MarkerAtomicRef(*marker).store(0); + + try { + q.single_task([=]() { MarkerAtomicRef(*marker).store(1); }); + sycl::event waitEvent = + q.ext_oneapi_wait_external_semaphore(syclSem, VulkanSignalValue); + + if (!waitForMarker(marker)) + failWithoutCleanup( + "The external semaphore wait did not submit the preceding batch."); + if (!remainsIncomplete(waitEvent)) + failWithoutCleanup( + "The external semaphore wait completed before Vulkan signaled it."); + + VkSemaphoreSignalInfo signalInfo = { + VK_STRUCTURE_TYPE_SEMAPHORE_SIGNAL_INFO}; + signalInfo.semaphore = vkSem; + signalInfo.value = VulkanSignalValue; + VkResult signalResult = vkSignalSemaphore(vkCtx.device, &signalInfo); + if (signalResult != VK_SUCCESS) + failWithoutCleanup("Failed to signal the Vulkan timeline semaphore: " + + std::to_string(signalResult)); + if (!waitForCompletion(waitEvent)) + failWithoutCleanup( + "The external semaphore wait did not complete after Vulkan signaled " + "it."); + + q.ext_oneapi_signal_external_semaphore(syclSem, SyclSignalValue); + + VkSemaphoreWaitInfo waitInfo = {VK_STRUCTURE_TYPE_SEMAPHORE_WAIT_INFO}; + waitInfo.semaphoreCount = 1; + waitInfo.pSemaphores = &vkSem; + waitInfo.pValues = &SyclSignalValue; + + constexpr uint64_t TimeoutNs = 5'000'000'000; + VkResult waitResult = vkWaitSemaphores(vkCtx.device, &waitInfo, TimeoutNs); + if (waitResult == VK_TIMEOUT) + failWithoutCleanup( + "Timed out waiting for the regular command-list semaphore signal."); + if (waitResult != VK_SUCCESS) + failWithoutCleanup("Waiting for the Vulkan timeline semaphore failed: " + + std::to_string(waitResult)); + + q.wait_and_throw(); + } catch (const std::exception &e) { + failWithoutCleanup(e.what()); + } + + sycl::free(marker, q); + syclexp::release_external_semaphore(syclSem, device, context); + vkDestroySemaphore(vkCtx.device, vkSem, nullptr); + cleanupVulkanContext(vkCtx); + return 0; +} diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/external_semaphore_regular_cl_fails.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/external_semaphore_regular_cl_fails.cpp deleted file mode 100644 index f57d06ab0ff36..0000000000000 --- a/sycl/test-e2e/bindless_images/vulkan_interop/external_semaphore_regular_cl_fails.cpp +++ /dev/null @@ -1,86 +0,0 @@ - -// -// REQUIRES: aspect-ext_oneapi_external_semaphore_import -// REQUIRES: vulkan && level_zero -// -// UNSUPPORTED: linux && run-mode -// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/23242 -// -// XFAIL: windows && run-mode && !gpu-intel-gen12 -// XFAIL-TRACKER: https://github.com/intel/llvm/issues/23249 -// -// RUN: %{build} %link-vulkan -o %t.out %if target-spir %{ -Wno-ignored-attributes %} -// RUN: %{run} %t.out - -// Waiting on an external semaphore from a queue backed by a regular -// (non-immediate) command list must throw sycl::exception. -// -// sycl_ext_oneapi_bindless_images requires queues used with external -// semaphore wait/signal to be constructed with BOTH -// - sycl::property::queue::in_order -// - sycl::ext::intel::property::queue::immediate_command_list -// The Level Zero adapter rejects external_semaphore wait/signal at the -// point of submission when the queue is not using immediate command -// lists. This test verifies that contract by: -// 1. Creating a real, exportable Vulkan binary semaphore. -// 2. Importing it into SYCL via a (lawful) immediate-CL queue. -// 3. Calling ext_oneapi_wait_external_semaphore on a separate queue -// that explicitly opts into no_immediate_command_list, and -// expecting a sycl::exception. - -#include "sycl_vulkan_setup.hpp" -#include -#include -#include -#include - -namespace syclexp = sycl::ext::oneapi::experimental; - -int main() { - VulkanContext vkCtx = createSyclVulkanContext(); - VkSemaphore vkSem = createExportableSemaphore(vkCtx); - - // Lawful queue: import the semaphore here. - sycl::queue immQ{ - {sycl::property::queue::in_order{}, - sycl::ext::intel::property::queue::immediate_command_list{}}}; - auto device = immQ.get_device(); - auto context = immQ.get_context(); - -#ifdef _WIN32 - HANDLE semHandle = getSemaphoreHandle(vkCtx, vkSem); - syclexp::external_semaphore_descriptor desc{ - semHandle, syclexp::external_semaphore_handle_type::win32_nt_handle}; -#else - int semFd = getSemaphoreFd(vkCtx, vkSem); - syclexp::external_semaphore_descriptor desc{ - semFd, syclexp::external_semaphore_handle_type::opaque_fd}; -#endif - - syclexp::external_semaphore syclSem = - syclexp::import_external_semaphore(desc, device, context); - - // The non-immediate-CL queue is what should trigger rejection on use. - sycl::queue regQ{ - context, device, - sycl::property_list{ - sycl::property::queue::in_order{}, - sycl::ext::intel::property::queue::no_immediate_command_list{}}}; - - int ret = 1; - try { - regQ.ext_oneapi_wait_external_semaphore(syclSem); - regQ.wait_and_throw(); - std::cerr << "FAIL: ext_oneapi_wait_external_semaphore on a " - "non-immediate-CL queue did not throw." - << std::endl; - } catch (const sycl::exception &e) { - std::cout << "Got expected sycl::exception: " << e.what() << std::endl; - ret = 0; - } - - syclexp::release_external_semaphore(syclSem, device, context); - vkDestroySemaphore(vkCtx.device, vkSem, nullptr); - cleanupVulkanContext(vkCtx); - return ret; -} diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_2d_arithmetic.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_2d_arithmetic.cpp index 84f224588dab0..5186e49358bdf 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_2d_arithmetic.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_2d_arithmetic.cpp @@ -240,9 +240,8 @@ int runTest( }); try { - // Bindless image interop requires an in-order queue (per spec). External - // semaphore ops additionally require immediate command lists; see - // sycl_ext_oneapi_bindless_images.asciidoc. + // External semaphore operations require an in-order queue. This test also + // uses immediate command lists for its existing configuration. sycl::property_list qProps = useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, sycl::ext::intel::property::queue:: diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_buffer.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_buffer.cpp index 6dcd2f21863a0..163cc67cc6e42 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_buffer.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_buffer.cpp @@ -271,9 +271,8 @@ int main(int argc, char **argv) { // SYCL INTEROP try { - // Bindless image interop requires an in-order queue (per spec). External - // semaphore ops additionally require immediate command lists; see - // sycl_ext_oneapi_bindless_images.asciidoc. + // External semaphore operations require an in-order queue. This test also + // uses immediate command lists for its existing configuration. sycl::property_list qProps = useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, sycl::ext::intel::property::queue:: diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_buffer_binary_semaphore.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_buffer_binary_semaphore.cpp index fd5bf9cf6becb..1bcaaced7f33d 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_buffer_binary_semaphore.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_buffer_binary_semaphore.cpp @@ -145,9 +145,8 @@ int main(int argc, char **argv) { // SYCL INTEROP try { - // Bindless image interop requires an in-order queue (per spec). External - // semaphore ops additionally require immediate command lists; see - // sycl_ext_oneapi_bindless_images.asciidoc. + // External semaphore operations require an in-order queue. This test also + // uses immediate command lists for its existing configuration. sycl::property_list qProps = useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, sycl::ext::intel::property::queue:: diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_buffer_timeline_semaphore.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_buffer_timeline_semaphore.cpp index 018e4f0a29458..4fb0ef3e14b84 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_buffer_timeline_semaphore.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_buffer_timeline_semaphore.cpp @@ -130,9 +130,8 @@ int main(int argc, char **argv) { // SYCL INTEROP try { - // Bindless image interop requires an in-order queue (per spec). External - // semaphore ops additionally require immediate command lists; see - // sycl_ext_oneapi_bindless_images.asciidoc. + // External semaphore operations require an in-order queue. This test also + // uses immediate command lists for its existing configuration. sycl::property_list qProps = useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, sycl::ext::intel::property::queue:: diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_1d.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_1d.cpp index 2c42b58421014..93fb6bff195ca 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_1d.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_1d.cpp @@ -269,9 +269,8 @@ int runTest( // SYCL Import and Verification namespace syclexp = sycl::ext::oneapi::experimental; try { - // Bindless image interop requires an in-order queue (per spec). External - // semaphore ops additionally require immediate command lists; see - // sycl_ext_oneapi_bindless_images.asciidoc. + // External semaphore operations require an in-order queue. This test also + // uses immediate command lists for its existing configuration. sycl::property_list qProps = useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, sycl::ext::intel::property::queue:: diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_2d.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_2d.cpp index b25a92378d7ec..16fc2ad92b451 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_2d.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_2d.cpp @@ -241,9 +241,8 @@ int runTest( // SYCL Import and Verification namespace syclexp = sycl::ext::oneapi::experimental; try { - // Bindless image interop requires an in-order queue (per spec). External - // semaphore ops additionally require immediate command lists; see - // sycl_ext_oneapi_bindless_images.asciidoc. + // External semaphore operations require an in-order queue. This test also + // uses immediate command lists for its existing configuration. sycl::property_list qProps = useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, sycl::ext::intel::property::queue:: diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_3d.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_3d.cpp index f570d5ac9a90f..d174f0d89a7a6 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_3d.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_3d.cpp @@ -140,9 +140,8 @@ int runTest( // SYCL Import and Verification namespace syclexp = sycl::ext::oneapi::experimental; try { - // Bindless image interop requires an in-order queue (per spec). External - // semaphore ops additionally require immediate command lists; see - // sycl_ext_oneapi_bindless_images.asciidoc. + // External semaphore operations require an in-order queue. This test also + // uses immediate command lists for its existing configuration. sycl::property_list qProps = useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, sycl::ext::intel::property::queue:: diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_1d_unsampled.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_1d_unsampled.cpp index b042ccd8e2ac9..9ac49f3d720ce 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_1d_unsampled.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_1d_unsampled.cpp @@ -217,9 +217,8 @@ int runTest( } try { - // Bindless image interop requires an in-order queue (per spec). External - // semaphore ops additionally require immediate command lists; see - // sycl_ext_oneapi_bindless_images.asciidoc. + // External semaphore operations require an in-order queue. This test also + // uses immediate command lists for its existing configuration. sycl::property_list qProps = useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, sycl::ext::intel::property::queue:: diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_2d_unsampled.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_2d_unsampled.cpp index 5d8200aad1f72..73e5376c48e7d 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_2d_unsampled.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_2d_unsampled.cpp @@ -203,9 +203,8 @@ int runTest( } try { - // Bindless image interop requires an in-order queue (per spec). External - // semaphore ops additionally require immediate command lists; see - // sycl_ext_oneapi_bindless_images.asciidoc. + // External semaphore operations require an in-order queue. This test also + // uses immediate command lists for its existing configuration. sycl::property_list qProps = useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, sycl::ext::intel::property::queue:: diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_3d_unsampled.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_3d_unsampled.cpp index 02d6944c2b591..cc36dfff0eca6 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_3d_unsampled.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_3d_unsampled.cpp @@ -154,9 +154,8 @@ int runTest( } try { - // Bindless image interop requires an in-order queue (per spec). External - // semaphore ops additionally require immediate command lists; see - // sycl_ext_oneapi_bindless_images.asciidoc. + // External semaphore operations require an in-order queue. This test also + // uses immediate command lists for its existing configuration. sycl::property_list qProps = useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, sycl::ext::intel::property::queue:: diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_unsampled_timeline_semaphore.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_unsampled_timeline_semaphore.cpp index 98a8cc4d50463..9cddaf9b5ea19 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_unsampled_timeline_semaphore.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_unsampled_timeline_semaphore.cpp @@ -167,8 +167,8 @@ int runTest( uint64_t syclSignalVal = 2; try { - // External semaphore ops require an in-order queue backed by immediate - // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). + // External semaphore operations require an in-order queue. This test also + // uses immediate command lists for its existing configuration. sycl::queue q{ {sycl::property::queue::in_order{}, sycl::ext::intel::property::queue::immediate_command_list{}}}; From db9eccf2bf7e2d97ce762b505f2332ac48ab419f Mon Sep 17 00:00:00 2001 From: Georgij Tsarin Date: Fri, 9 Oct 2026 14:42:30 +0300 Subject: [PATCH 2/3] [SYCL][E2E][Doc] Harden external semaphore testing on regular command lists --- .../sycl_ext_oneapi_bindless_images.asciidoc | 12 ++++++++- .../dx12_interop/d3d12_setup.hpp | 4 +-- .../external_semaphore_regular_cl.cpp | 25 ++++++++++++++++--- 3 files changed, 35 insertions(+), 6 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc index 5f194a52581f6..1afb201b8da30 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc @@ -2340,6 +2340,15 @@ mechanism is not supported on the default out-of-order queues. Use of the semaphore synchronization mechanism with a SYCL queue that was not constructed with this property will result in undefined behaviour. +[NOTE] +==== +Support for external semaphore operations on queues backed by regular command +lists is backend- and driver-dependent. The {dpcpp} Level Zero backend requires +Intel graphics driver version 39758 or newer on Linux and 101.9030 or newer on +Windows. On older drivers, an external semaphore wait submitted to a regular +command list might not complete. +==== + External semaphore import is facilitated through the following proposed descriptor struct. @@ -2846,5 +2855,6 @@ This query should be added in a later revision of the proposal. |6.12|2025-07-03| - Add `win32_nt_dx11_resource` enum entry to `external_mem_handle_type` for Direct3D 11 memory interop. |6.13|2026-10-07| - Allow external semaphore operations on in-order queues - backed by regular command lists. + backed by regular command lists and document the associated + backend and driver requirements. |====================== diff --git a/sycl/test-e2e/bindless_images/dx12_interop/d3d12_setup.hpp b/sycl/test-e2e/bindless_images/dx12_interop/d3d12_setup.hpp index 45f8a678baf31..5752458eef85e 100644 --- a/sycl/test-e2e/bindless_images/dx12_interop/d3d12_setup.hpp +++ b/sycl/test-e2e/bindless_images/dx12_interop/d3d12_setup.hpp @@ -148,10 +148,10 @@ struct D3D12ImageResources { // --------------------------------------------------------- // D3D12 LIFECYCLE MANAGEMENT // --------------------------------------------------------- -inline D3D12Context createD3D12Context() { +inline D3D12Context createD3D12Context(IDXGIAdapter *adapter = nullptr) { D3D12Context ctx; - ThrowIfFailed(D3D12CreateDevice(nullptr, D3D_FEATURE_LEVEL_11_0, + ThrowIfFailed(D3D12CreateDevice(adapter, D3D_FEATURE_LEVEL_11_0, IID_PPV_ARGS(&ctx.device)), "Failed to create D3D12 Device"); diff --git a/sycl/test-e2e/bindless_images/dx12_interop/external_semaphore_regular_cl.cpp b/sycl/test-e2e/bindless_images/dx12_interop/external_semaphore_regular_cl.cpp index 70e32797027b3..8aff5b5e34f17 100644 --- a/sycl/test-e2e/bindless_images/dx12_interop/external_semaphore_regular_cl.cpp +++ b/sycl/test-e2e/bindless_images/dx12_interop/external_semaphore_regular_cl.cpp @@ -1,4 +1,5 @@ // REQUIRES: aspect-ext_oneapi_external_semaphore_import, aspect-usm_shared_allocations +// REQUIRES: aspect-ext_intel_device_info_luid // REQUIRES: windows, level_zero // Regular command-list external semaphore support requires this driver @@ -18,10 +19,12 @@ #include #include #include +#include #include #include #include #include +#include #include #include #include @@ -32,6 +35,22 @@ namespace syclexp = sycl::ext::oneapi::experimental; +D3D12Context createSyclD3D12Context(const sycl::device &device) { + auto syclLuid = device.get_info(); + static_assert(sizeof(LUID) == sizeof(syclLuid)); + LUID adapterLuid{}; + std::memcpy(&adapterLuid, syclLuid.data(), sizeof(adapterLuid)); + + ComPtr factory; + ThrowIfFailed(CreateDXGIFactory1(IID_PPV_ARGS(&factory)), + "Failed to create DXGI factory"); + + ComPtr adapter; + ThrowIfFailed(factory->EnumAdapterByLuid(adapterLuid, IID_PPV_ARGS(&adapter)), + "Failed to find the DXGI adapter for the SYCL device"); + return createD3D12Context(adapter.Get()); +} + [[noreturn]] void failWithoutCleanup(const std::string &message) { std::cerr << message << std::endl; std::_Exit(EXIT_FAILURE); @@ -84,7 +103,7 @@ int main() { auto device = q.get_device(); auto context = q.get_context(); - D3D12Context d3dCtx = createD3D12Context(); + D3D12Context d3dCtx = createSyclD3D12Context(device); D3D12ExportableFence extFence = createExportableFence(d3dCtx); auto semDesc = @@ -97,7 +116,7 @@ int main() { MarkerAtomicRef(*marker).store(0); try { - constexpr uint64_t D3DSignalValue = 1; + const uint64_t D3DSignalValue = extFence.fenceValue + 1; q.single_task([=]() { MarkerAtomicRef(*marker).store(1); }); sycl::event waitEvent = q.ext_oneapi_wait_external_semaphore(syclSem, D3DSignalValue); @@ -115,7 +134,7 @@ int main() { "The external semaphore wait did not complete after D3D12 signaled " "it."); - constexpr uint64_t SyclSignalValue = 2; + const uint64_t SyclSignalValue = D3DSignalValue + 1; q.ext_oneapi_signal_external_semaphore(syclSem, SyclSignalValue); ThrowIfFailed(extFence.fence->SetEventOnCompletion(SyclSignalValue, From 8782f82440aa82064b44ec407d451d5775adee3e Mon Sep 17 00:00:00 2001 From: Georgij Tsarin Date: Mon, 12 Oct 2026 16:43:31 +0300 Subject: [PATCH 3/3] [UR][L0 v2] Isolate external semaphore operations in regular command lists --- .../adapters/level_zero/v2/queue_batched.cpp | 38 ++++++++++++++----- 1 file changed, 29 insertions(+), 9 deletions(-) diff --git a/unified-runtime/source/adapters/level_zero/v2/queue_batched.cpp b/unified-runtime/source/adapters/level_zero/v2/queue_batched.cpp index 0126e6e5d1659..61da3ff0bde4e 100644 --- a/unified-runtime/source/adapters/level_zero/v2/queue_batched.cpp +++ b/unified-runtime/source/adapters/level_zero/v2/queue_batched.cpp @@ -841,17 +841,27 @@ ur_result_t ur_queue_batched_t::bindlessImagesWaitExternalSemaphoreExp( wait_list_view(phEventWaitList, numEventsInWaitList, this); auto lockedBatch = currentCmdLists.lock(); + + // Keep the external wait in a separate regular command list. Some drivers + // block the whole list at an external wait, including commands recorded + // before it, so submit any preceding work first. + if (!lockedBatch->isGraphCaptureActive() && + !lockedBatch->isActiveBatchEmpty()) { + UR_CALL(renewBatchUnlocked(lockedBatch)); + } + markIssuedCommandInBatch(lockedBatch); UR_CALL(lockedBatch->getListManager().bindlessImagesWaitExternalSemaphoreExp( hSemaphore, hasWaitValue, waitValue, waitListView, getEvent(lockedBatch, phEvent))); - // External semaphore operations must not be deferred in a batch. Batching - // leaves the append sitting in the queue's open regular command list until - // some unrelated event flushes it, which means the external producer's - // signal cannot release a command list that was never submitted and can - // deadlock. + if (lockedBatch->isGraphCaptureActive()) { + return UR_RESULT_SUCCESS; + } + + // Submit the external wait immediately so the external producer can release + // it without relying on a later operation to flush the batch. return renewBatchUnlocked(lockedBatch); } @@ -863,6 +873,14 @@ ur_result_t ur_queue_batched_t::bindlessImagesSignalExternalSemaphoreExp( wait_list_view(phEventWaitList, numEventsInWaitList, this); auto lockedBatch = currentCmdLists.lock(); + + // Keep the external signal in a separate regular command list, ordered after + // all preceding work. + if (!lockedBatch->isGraphCaptureActive() && + !lockedBatch->isActiveBatchEmpty()) { + UR_CALL(renewBatchUnlocked(lockedBatch)); + } + markIssuedCommandInBatch(lockedBatch); UR_CALL( @@ -870,10 +888,12 @@ ur_result_t ur_queue_batched_t::bindlessImagesSignalExternalSemaphoreExp( hSemaphore, hasSignalValue, signalValue, waitListView, getEvent(lockedBatch, phEvent))); - // External semaphore operations must not be deferred in a batch. Batching - // leaves the append sitting in the queue's open regular command list until - // some unrelated event flushes it, which means the external consumer may - // never observe the signal. + if (lockedBatch->isGraphCaptureActive()) { + return UR_RESULT_SUCCESS; + } + + // Submit the external signal immediately so the external consumer can + // observe it without relying on a later operation to flush the batch. return renewBatchUnlocked(lockedBatch); }