From f9743a9522c5b3d230576e8ea4470129dd90d8ae Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 13 May 2026 14:32:36 -0700 Subject: [PATCH 01/17] bindless image semaphore usage requires queues that use immediate command lists --- .../sycl_ext_oneapi_bindless_images.asciidoc | 19 ++++--- .../dx11_interop/read_write_unsampled.cpp | 11 +++- .../D3D12_sycl_buffer_timeline_semaphore.cpp | 7 ++- .../D3D12_sycl_buffer_win32_name_native.cpp | 7 ++- .../D3D12_sycl_interop_1D_read.cpp | 7 ++- .../D3D12_sycl_interop_1D_write_unsampled.cpp | 7 ++- .../D3D12_sycl_interop_2D_arithmetic.cpp | 7 ++- .../D3D12_sycl_interop_2D_read.cpp | 7 ++- .../D3D12_sycl_interop_2D_write_unsampled.cpp | 7 ++- .../D3D12_sycl_interop_3D_read.cpp | 7 ++- .../D3D12_sycl_interop_3D_write_unsampled.cpp | 7 ++- .../external_semaphore_regular_cl_fails.cpp | 49 ++++++++++++++++ ...example_6_import_memory_and_semaphores.cpp | 9 ++- .../external_semaphore_regular_cl_fails.cpp | 57 +++++++++++++++++++ .../vulkan_sycl_2d_arithmetic.cpp | 7 ++- .../vulkan_interop/vulkan_sycl_buffer.cpp | 7 ++- .../vulkan_sycl_buffer_binary_semaphore.cpp | 7 ++- .../vulkan_sycl_buffer_timeline_semaphore.cpp | 7 ++- .../vulkan_sycl_image_interop_read_1d.cpp | 7 ++- .../vulkan_sycl_image_interop_read_2d.cpp | 7 ++- .../vulkan_sycl_image_interop_read_3d.cpp | 7 ++- ..._sycl_image_interop_write_1d_unsampled.cpp | 7 ++- ..._sycl_image_interop_write_2d_unsampled.cpp | 7 ++- ..._sycl_image_interop_write_3d_unsampled.cpp | 7 ++- ...ycl_image_unsampled_timeline_semaphore.cpp | 7 ++- 25 files changed, 255 insertions(+), 30 deletions(-) create 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_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 c55987f4bae35..bf9526bad3530 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc @@ -2259,13 +2259,18 @@ memory resources handles can take different forms of structure and type 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 has the restriction in that imported external semaphores can only -be used in conjuction with SYCL queues that have been constructed with the -`property::queue::in_order` property. The semaphore synchronization mechanism -is not supported for the default SYCL out-of-order queues. Use of the semaphore -synchronization mechanism with SYCL queues which were not constructed with the -`queue::in_order` property will result in undefined behaviour. +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. + +Use of the semaphore synchronization mechanism with a SYCL queue that was not +constructed with both of these properties will result in undefined behaviour. External semaphore import is facilitated through the following proposed descriptor struct. 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 a08cde973cd9e..3178013770ccf 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 @@ -17,6 +17,7 @@ #include "dx11_interop.h" #include +#include #ifdef TEST_SEMAPHORE_IMPORT #include // Used for ID3D11Device5 / ID3D11DeviceContext4 / ID3D11Fence @@ -438,8 +439,16 @@ int runTest(D3D11ProgramState &d3d11ProgramState, sycl::queue syclQueue, } int main() { - // Create SYCL queue, relying on SYCL device selection + // 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). + sycl::queue syclQueue{ + {sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue::immediate_command_list{}}}; +#else sycl::queue syclQueue; +#endif sycl::device syclDevice = syclQueue.get_device(); // Initialize D3D11 and create DX11 programs state from the SYCL device 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 e3115a72deb92..db61a4bd5daaf 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 @@ -50,6 +50,7 @@ #include #include #include +#include #include #define WIN32_LEAN_AND_MEAN @@ -120,7 +121,11 @@ int main(int argc, char **argv) { // SYCL INTEROP try { - sycl::queue q; + // External semaphore ops require an in-order queue backed by immediate + // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). + sycl::queue q{ + {sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue::immediate_command_list{}}}; auto device = q.get_device(); auto context = q.get_context(); diff --git a/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_win32_name_native.cpp b/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_win32_name_native.cpp index f96ffe8dd441e..96fad149bf1f3 100644 --- a/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_win32_name_native.cpp +++ b/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_win32_name_native.cpp @@ -32,6 +32,7 @@ #include #include #include +#include #include #define WIN32_LEAN_AND_MEAN @@ -180,7 +181,11 @@ int main(int argc, char **argv) { // SYCL INTEROP - using resource_win32_name NATIVELY try { - sycl::queue q; + // External semaphore ops require an in-order queue backed by immediate + // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). + sycl::queue q{ + {sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue::immediate_command_list{}}}; auto device = q.get_device(); auto context = q.get_context(); 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 8e8096461735f..aa7cb1bad8794 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 @@ -98,6 +98,7 @@ #include #include #include +#include namespace syclexp = sycl::ext::oneapi::experimental; @@ -187,7 +188,11 @@ int runTest( // SYCL Import and Verification try { - sycl::queue q; + // External semaphore ops require an in-order queue backed by immediate + // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). + sycl::queue q{ + {sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue::immediate_command_list{}}}; syclexp::external_mem_descriptor extMemDesc{ imgRes.sharedHandle, syclexp::external_mem_handle_type::win32_nt_handle, 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 6a65ebd906869..57abd7558a5ab 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 @@ -71,6 +71,7 @@ #include #include #include +#include namespace syclexp = sycl::ext::oneapi::experimental; @@ -148,7 +149,11 @@ int runTest( } try { - sycl::queue q; + // External semaphore ops require an in-order queue backed by immediate + // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). + sycl::queue q{ + {sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue::immediate_command_list{}}}; syclexp::external_mem_descriptor extMemDesc{ imgRes.sharedHandle, syclexp::external_mem_handle_type::win32_nt_handle, 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 4c4cdc594be42..740335b55f2a6 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 @@ -118,6 +118,7 @@ #include #include #include +#include #include namespace syclexp = sycl::ext::oneapi::experimental; @@ -415,7 +416,11 @@ int runTest( signalExportableFence(ctx, extFenceB); try { - sycl::queue q; + // External semaphore ops require an in-order queue backed by immediate + // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). + sycl::queue q{ + {sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue::immediate_command_list{}}}; auto extMemA = syclexp::import_external_memory( syclexp::external_mem_descriptor{ 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 63b9e979f2410..dd41ec397a87f 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 @@ -122,6 +122,7 @@ #include #include #include +#include namespace syclexp = sycl::ext::oneapi::experimental; @@ -208,7 +209,11 @@ int runTest( // SYCL Import and Verification try { - sycl::queue q; + // External semaphore ops require an in-order queue backed by immediate + // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). + sycl::queue q{ + {sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue::immediate_command_list{}}}; syclexp::external_mem_descriptor extMemDesc{ imgRes.sharedHandle, syclexp::external_mem_handle_type::win32_nt_handle, 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 8a9b8fd3bcfc3..d11b7ae8b0634 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 @@ -71,6 +71,7 @@ #include #include #include +#include namespace syclexp = sycl::ext::oneapi::experimental; @@ -189,7 +190,11 @@ int runTest( } try { - sycl::queue q; + // External semaphore ops require an in-order queue backed by immediate + // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). + sycl::queue q{ + {sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue::immediate_command_list{}}}; syclexp::external_mem_descriptor extMemDesc{ imgRes.sharedHandle, syclexp::external_mem_handle_type::win32_nt_handle, 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 3e3ec7ee0172c..6e2911833b7c5 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 @@ -116,6 +116,7 @@ #include #include #include +#include namespace syclexp = sycl::ext::oneapi::experimental; @@ -201,7 +202,11 @@ int runTest( // SYCL Import and Verification try { - sycl::queue q; + // External semaphore ops require an in-order queue backed by immediate + // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). + sycl::queue q{ + {sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue::immediate_command_list{}}}; syclexp::external_mem_descriptor extMemDesc{ imgRes.sharedHandle, syclexp::external_mem_handle_type::win32_nt_handle, 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 498f0cde6ecf9..184b02011835c 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 @@ -71,6 +71,7 @@ #include #include #include +#include namespace syclexp = sycl::ext::oneapi::experimental; @@ -148,7 +149,11 @@ int runTest( } try { - sycl::queue q; + // External semaphore ops require an in-order queue backed by immediate + // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). + sycl::queue q{ + {sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue::immediate_command_list{}}}; syclexp::external_mem_descriptor extMemDesc{ imgRes.sharedHandle, syclexp::external_mem_handle_type::win32_nt_handle, 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 new file mode 100644 index 0000000000000..75bfa4de693ef --- /dev/null +++ b/sycl/test-e2e/bindless_images/dx12_interop/external_semaphore_regular_cl_fails.cpp @@ -0,0 +1,49 @@ + +// +// REQUIRES: aspect-ext_oneapi_external_semaphore_import, windows +// +// RUN: %{build} -o %t.exe +// RUN: %{run} %t.exe + +// Importing a DX12 fence external semaphore on 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. Today the runtime rejects at +// import independent of handle type; this is belt-and-suspenders coverage +// in case a future adapter change treats handle types differently. +// +// This is a contract test, not an interop test: no real D3D12 device is +// created. The runtime rejects at import before inspecting the handle, so +// a null handle is enough. That's also why this test does not link against +// DirectX -- it lives here because the DX12 fence handle type is what +// readers will look for in this directory. + +#include +#include +#include +#include + +namespace syclexp = sycl::ext::oneapi::experimental; + +int main() { + sycl::queue q{ + {sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue::no_immediate_command_list{}}}; + + syclexp::external_semaphore_descriptor desc{ + /*handle=*/nullptr, + syclexp::external_semaphore_handle_type::win32_nt_dx12_fence}; + + try { + (void)syclexp::import_external_semaphore(desc, q); + } catch (const sycl::exception &e) { + std::cout << "Got expected sycl::exception: " << e.what() << std::endl; + return 0; + } + + std::cerr << "FAIL: import_external_semaphore (dx12_fence) on a " + "non-immediate-CL queue did not throw." + << std::endl; + return 1; +} 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 4451e0f59116f..f35419a5fb6c5 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 @@ -11,11 +11,16 @@ #include #include +#include int main() { - // Set up queue + // Set up queue. + // External semaphore ops require an in-order queue backed by immediate + // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). sycl::device dev; - sycl::queue q(dev); + sycl::queue q(dev, + {sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue::immediate_command_list{}}); 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_fails.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/external_semaphore_regular_cl_fails.cpp new file mode 100644 index 0000000000000..2bbd919523550 --- /dev/null +++ b/sycl/test-e2e/bindless_images/vulkan_interop/external_semaphore_regular_cl_fails.cpp @@ -0,0 +1,57 @@ + +// +// REQUIRES: aspect-ext_oneapi_external_semaphore_import +// +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +// Importing an external semaphore on a queue backed by a regular +// (non-immediate) command list must throw sycl::exception. +// +// sycl_ext_oneapi_bindless_images requires queues used with external +// semaphores to be constructed with BOTH +// - sycl::property::queue::in_order +// - sycl::ext::intel::property::queue::immediate_command_list +// This test violates the second requirement (explicitly requests +// no_immediate_command_list) and verifies the runtime rejects the import. +// +// This is a contract test, not an interop test: no real Vulkan context is +// created. The runtime rejects at import before inspecting the handle, so +// a bogus fd / handle is enough. That's also why this test does not gate +// on `vulkan` or link against the Vulkan loader -- it lives here because +// the Vulkan-flavored handle types are what most readers will look for. + +#include +#include +#include +#include + +namespace syclexp = sycl::ext::oneapi::experimental; + +int main() { + sycl::queue q{ + {sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue::no_immediate_command_list{}}}; + +#ifdef _WIN32 + syclexp::external_semaphore_descriptor desc{ + /*handle=*/nullptr, + syclexp::external_semaphore_handle_type::win32_nt_handle}; +#else + syclexp::external_semaphore_descriptor desc{ + /*file_descriptor=*/-1, + syclexp::external_semaphore_handle_type::opaque_fd}; +#endif + + try { + (void)syclexp::import_external_semaphore(desc, q); + } catch (const sycl::exception &e) { + std::cout << "Got expected sycl::exception: " << e.what() << std::endl; + return 0; + } + + std::cerr << "FAIL: import_external_semaphore on a non-immediate-CL queue " + "did not throw." + << std::endl; + return 1; +} 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 463dbe4044295..0b8d040c2e58b 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 @@ -130,6 +130,7 @@ #include #include #include +#include #include namespace syclexp = sycl::ext::oneapi::experimental; @@ -243,7 +244,11 @@ int runTest( }); try { - sycl::queue q; + // External semaphore ops require an in-order queue backed by immediate + // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). + sycl::queue q{ + {sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue::immediate_command_list{}}}; #ifdef _WIN32 auto extMemA = syclexp::import_external_memory( syclexp::external_mem_descriptor{ 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 3d9642bf8b559..35c2d353f739e 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 @@ -59,6 +59,7 @@ #include #include #include +#include #include namespace syclexp = sycl::ext::oneapi::experimental; @@ -277,7 +278,11 @@ int main(int argc, char **argv) { // SYCL INTEROP try { - sycl::queue q; + // External semaphore ops require an in-order queue backed by immediate + // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). + sycl::queue q{ + {sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue::immediate_command_list{}}}; auto device = q.get_device(); auto context = q.get_context(); 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 f192a25d3e66d..de03ea7fdc96f 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 @@ -55,6 +55,7 @@ #include #include #include +#include #include namespace syclexp = sycl::ext::oneapi::experimental; @@ -149,7 +150,11 @@ int main(int argc, char **argv) { // SYCL INTEROP try { - sycl::queue q; + // External semaphore ops require an in-order queue backed by immediate + // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). + sycl::queue q{ + {sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue::immediate_command_list{}}}; auto device = q.get_device(); auto context = q.get_context(); 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 a364380230efe..d5c32e475f190 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 @@ -59,6 +59,7 @@ #include #include #include +#include #include namespace syclexp = sycl::ext::oneapi::experimental; @@ -140,7 +141,11 @@ int main(int argc, char **argv) { // SYCL INTEROP try { - sycl::queue q; + // External semaphore ops require an in-order queue backed by immediate + // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). + sycl::queue q{ + {sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue::immediate_command_list{}}}; auto device = q.get_device(); auto context = q.get_context(); 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 43c0dfc76b4cb..37742fea9ceb0 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 @@ -164,6 +164,7 @@ VK_FORMAT_R8G8B8A8_UNORM #include #include #include +#include // --------------------------------------------------------- // SYCL TYPE MAPPING HELPERS @@ -257,7 +258,11 @@ int runTest( // SYCL Import and Verification namespace syclexp = sycl::ext::oneapi::experimental; try { - sycl::queue q; + // External semaphore ops require an in-order queue backed by immediate + // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). + sycl::queue q{ + {sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue::immediate_command_list{}}}; // Import Memory (Platform Specific) #ifdef _WIN32 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 ac59681b4071b..e5c2596ef91f3 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 @@ -135,6 +135,7 @@ #include #include #include +#include // --------------------------------------------------------- // SYCL TYPE MAPPING HELPERS @@ -245,7 +246,11 @@ int runTest( // SYCL Import and Verification namespace syclexp = sycl::ext::oneapi::experimental; try { - sycl::queue q; + // External semaphore ops require an in-order queue backed by immediate + // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). + sycl::queue q{ + {sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue::immediate_command_list{}}}; // Import Memory (Platform Specific) #ifdef _WIN32 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 9b19aba506ebc..3cd7f557ad5e8 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 @@ -44,6 +44,7 @@ #include #include #include +#include // --------------------------------------------------------- // SYCL TYPE MAPPING HELPERS @@ -138,7 +139,11 @@ int runTest( // SYCL Import and Verification namespace syclexp = sycl::ext::oneapi::experimental; try { - sycl::queue q; + // External semaphore ops require an in-order queue backed by immediate + // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). + sycl::queue q{ + {sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue::immediate_command_list{}}}; // Import Memory (Platform Specific) #ifdef _WIN32 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 aa8532eaf61a5..7e7be225781ee 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 @@ -99,6 +99,7 @@ #include #include #include +#include namespace syclexp = sycl::ext::oneapi::experimental; @@ -212,7 +213,11 @@ int runTest( } try { - sycl::queue q; + // External semaphore ops require an in-order queue backed by immediate + // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). + sycl::queue q{ + {sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue::immediate_command_list{}}}; // IMPORT MEMORY #ifdef _WIN32 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 3644fd78b7fd0..d4f1357e9373d 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 @@ -96,6 +96,7 @@ #include #include #include +#include namespace syclexp = sycl::ext::oneapi::experimental; @@ -208,7 +209,11 @@ int runTest( } try { - sycl::queue q; + // External semaphore ops require an in-order queue backed by immediate + // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). + sycl::queue q{ + {sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue::immediate_command_list{}}}; // IMPORT MEMORY #ifdef _WIN32 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 f3e1b5458e07b..bc1c6b1d878e5 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 @@ -40,6 +40,7 @@ #include #include #include +#include namespace syclexp = sycl::ext::oneapi::experimental; @@ -152,7 +153,11 @@ int runTest( } try { - sycl::queue q; + // External semaphore ops require an in-order queue backed by immediate + // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). + sycl::queue q{ + {sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue::immediate_command_list{}}}; // IMPORT MEMORY #ifdef _WIN32 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 e88d5d6ede809..4777448970135 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 @@ -60,6 +60,7 @@ #include #include #include +#include namespace syclexp = sycl::ext::oneapi::experimental; @@ -174,7 +175,11 @@ int runTest( uint64_t syclSignalVal = 2; try { - sycl::queue q; + // External semaphore ops require an in-order queue backed by immediate + // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). + sycl::queue q{ + {sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue::immediate_command_list{}}}; #ifdef _WIN32 HANDLE memHandle = getMemHandle(vkCtx, imgResc.memory); From ed6c1d3b665e4a4e28f936b43668936a16926b50 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Fri, 15 May 2026 11:35:58 -0700 Subject: [PATCH 02/17] narrower include --- .../bindless_images/dx11_interop/read_write_unsampled.cpp | 2 +- .../dx12_interop/D3D12_sycl_buffer_timeline_semaphore.cpp | 2 +- .../dx12_interop/D3D12_sycl_buffer_win32_name_native.cpp | 2 +- .../bindless_images/dx12_interop/D3D12_sycl_interop_1D_read.cpp | 2 +- .../dx12_interop/D3D12_sycl_interop_1D_write_unsampled.cpp | 2 +- .../dx12_interop/D3D12_sycl_interop_2D_arithmetic.cpp | 2 +- .../bindless_images/dx12_interop/D3D12_sycl_interop_2D_read.cpp | 2 +- .../dx12_interop/D3D12_sycl_interop_2D_write_unsampled.cpp | 2 +- .../bindless_images/dx12_interop/D3D12_sycl_interop_3D_read.cpp | 2 +- .../dx12_interop/D3D12_sycl_interop_3D_write_unsampled.cpp | 2 +- .../dx12_interop/external_semaphore_regular_cl_fails.cpp | 2 +- .../examples/example_6_import_memory_and_semaphores.cpp | 2 +- .../vulkan_interop/external_semaphore_regular_cl_fails.cpp | 2 +- .../vulkan_interop/vulkan_sycl_2d_arithmetic.cpp | 2 +- .../bindless_images/vulkan_interop/vulkan_sycl_buffer.cpp | 2 +- .../vulkan_interop/vulkan_sycl_buffer_binary_semaphore.cpp | 2 +- .../vulkan_interop/vulkan_sycl_buffer_timeline_semaphore.cpp | 2 +- .../vulkan_interop/vulkan_sycl_image_interop_read_1d.cpp | 2 +- .../vulkan_interop/vulkan_sycl_image_interop_read_2d.cpp | 2 +- .../vulkan_interop/vulkan_sycl_image_interop_read_3d.cpp | 2 +- .../vulkan_sycl_image_interop_write_1d_unsampled.cpp | 2 +- .../vulkan_sycl_image_interop_write_2d_unsampled.cpp | 2 +- .../vulkan_sycl_image_interop_write_3d_unsampled.cpp | 2 +- .../vulkan_sycl_image_unsampled_timeline_semaphore.cpp | 2 +- 24 files changed, 24 insertions(+), 24 deletions(-) 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 3178013770ccf..990d2bed635be 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 @@ -17,7 +17,7 @@ #include "dx11_interop.h" #include -#include +#include #ifdef TEST_SEMAPHORE_IMPORT #include // Used for ID3D11Device5 / ID3D11DeviceContext4 / ID3D11Fence 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 db61a4bd5daaf..28f65653b5161 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 @@ -50,7 +50,7 @@ #include #include #include -#include +#include #include #define WIN32_LEAN_AND_MEAN diff --git a/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_win32_name_native.cpp b/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_win32_name_native.cpp index 96fad149bf1f3..581e48d3188ec 100644 --- a/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_win32_name_native.cpp +++ b/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_win32_name_native.cpp @@ -32,7 +32,7 @@ #include #include #include -#include +#include #include #define WIN32_LEAN_AND_MEAN 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 aa7cb1bad8794..e1e2ef6301a52 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 @@ -98,7 +98,7 @@ #include #include #include -#include +#include namespace syclexp = sycl::ext::oneapi::experimental; 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 57abd7558a5ab..c3ba154eeaa0c 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 @@ -71,7 +71,7 @@ #include #include #include -#include +#include namespace syclexp = sycl::ext::oneapi::experimental; 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 740335b55f2a6..49e185d282f35 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 @@ -118,7 +118,7 @@ #include #include #include -#include +#include #include namespace syclexp = sycl::ext::oneapi::experimental; 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 dd41ec397a87f..e51b8172df4d8 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 @@ -122,7 +122,7 @@ #include #include #include -#include +#include namespace syclexp = sycl::ext::oneapi::experimental; 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 d11b7ae8b0634..fd59cdfcd0286 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 @@ -71,7 +71,7 @@ #include #include #include -#include +#include namespace syclexp = sycl::ext::oneapi::experimental; 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 6e2911833b7c5..5e2d44af8978f 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 @@ -116,7 +116,7 @@ #include #include #include -#include +#include namespace syclexp = sycl::ext::oneapi::experimental; 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 184b02011835c..220b201588a4e 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 @@ -71,7 +71,7 @@ #include #include #include -#include +#include namespace syclexp = sycl::ext::oneapi::experimental; 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 index 75bfa4de693ef..44d545e61a3c1 100644 --- 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 @@ -22,7 +22,7 @@ #include #include #include -#include +#include namespace syclexp = sycl::ext::oneapi::experimental; 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 f35419a5fb6c5..0b91824cf9bc8 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 @@ -11,7 +11,7 @@ #include #include -#include +#include int main() { // Set up queue. 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 index 2bbd919523550..645d1683d5cc1 100644 --- 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 @@ -24,7 +24,7 @@ #include #include #include -#include +#include namespace syclexp = sycl::ext::oneapi::experimental; 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 0b8d040c2e58b..d9cd21b0bfb2f 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 @@ -130,7 +130,7 @@ #include #include #include -#include +#include #include namespace syclexp = sycl::ext::oneapi::experimental; 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 35c2d353f739e..65572b382ec54 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 @@ -59,7 +59,7 @@ #include #include #include -#include +#include #include namespace syclexp = sycl::ext::oneapi::experimental; 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 de03ea7fdc96f..e4d4cb3e663c1 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 @@ -55,7 +55,7 @@ #include #include #include -#include +#include #include namespace syclexp = sycl::ext::oneapi::experimental; 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 d5c32e475f190..fcba2223ab136 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 @@ -59,7 +59,7 @@ #include #include #include -#include +#include #include namespace syclexp = sycl::ext::oneapi::experimental; 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 37742fea9ceb0..d0f429456f575 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 @@ -164,7 +164,7 @@ VK_FORMAT_R8G8B8A8_UNORM #include #include #include -#include +#include // --------------------------------------------------------- // SYCL TYPE MAPPING HELPERS 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 e5c2596ef91f3..ff9626d70a4a9 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 @@ -135,7 +135,7 @@ #include #include #include -#include +#include // --------------------------------------------------------- // SYCL TYPE MAPPING HELPERS 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 3cd7f557ad5e8..017e3a26a8b79 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 @@ -44,7 +44,7 @@ #include #include #include -#include +#include // --------------------------------------------------------- // SYCL TYPE MAPPING HELPERS 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 7e7be225781ee..441f1670c310f 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 @@ -99,7 +99,7 @@ #include #include #include -#include +#include namespace syclexp = sycl::ext::oneapi::experimental; 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 d4f1357e9373d..09a429424d541 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 @@ -96,7 +96,7 @@ #include #include #include -#include +#include namespace syclexp = sycl::ext::oneapi::experimental; 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 bc1c6b1d878e5..145eddb998043 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 @@ -40,7 +40,7 @@ #include #include #include -#include +#include namespace syclexp = sycl::ext::oneapi::experimental; 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 4777448970135..228b825ec45c1 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 @@ -60,7 +60,7 @@ #include #include #include -#include +#include namespace syclexp = sycl::ext::oneapi::experimental; From a58cfbc86920589967648d50fbb39b411ab8c232 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Fri, 15 May 2026 15:29:16 -0700 Subject: [PATCH 03/17] revise coupla tests --- .../external_semaphore_regular_cl_fails.cpp | 74 +++++++++++++------ .../external_semaphore_regular_cl_fails.cpp | 69 +++++++++++------ .../vulkan_interop/vulkan_sycl_buffer.cpp | 3 - 3 files changed, 98 insertions(+), 48 deletions(-) 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 index 44d545e61a3c1..31a3983d52b14 100644 --- 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 @@ -2,48 +2,78 @@ // // REQUIRES: aspect-ext_oneapi_external_semaphore_import, windows // -// RUN: %{build} -o %t.exe +// RUN: %{build} %link-directx -o %t.exe %if target-spir %{ -Wno-ignored-attributes %} // RUN: %{run} %t.exe -// Importing a DX12 fence external semaphore on a queue backed by a regular -// (non-immediate) command list must throw sycl::exception. +// 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. Today the runtime rejects at -// import independent of handle type; this is belt-and-suspenders coverage -// in case a future adapter change treats handle types differently. +// 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. // -// This is a contract test, not an interop test: no real D3D12 device is -// created. The runtime rejects at import before inspecting the handle, so -// a null handle is enough. That's also why this test does not link against -// DirectX -- it lives here because the DX12 fence handle type is what -// readers will look for in this directory. +// 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() { - sycl::queue q{ + 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::no_immediate_command_list{}}}; + 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); - syclexp::external_semaphore_descriptor desc{ - /*handle=*/nullptr, - syclexp::external_semaphore_handle_type::win32_nt_dx12_fence}; + // 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 { - (void)syclexp::import_external_semaphore(desc, q); + 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; - return 0; + ret = 0; } - std::cerr << "FAIL: import_external_semaphore (dx12_fence) on a " - "non-immediate-CL queue did not throw." - << std::endl; - return 1; + syclexp::release_external_semaphore(syclSem, device, context); + cleanupExportableFence(extFence); + return ret; } 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 index 645d1683d5cc1..d325d49defe5e 100644 --- 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 @@ -1,26 +1,28 @@ // // REQUIRES: aspect-ext_oneapi_external_semaphore_import +// REQUIRES: vulkan // -// RUN: %{build} -o %t.out +// RUN: %{build} %link-vulkan -o %t.out %if target-spir %{ -Wno-ignored-attributes %} // RUN: %{run} %t.out -// Importing an external semaphore on a queue backed by a regular +// 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 -// semaphores to be constructed with BOTH +// semaphore wait/signal to be constructed with BOTH // - sycl::property::queue::in_order // - sycl::ext::intel::property::queue::immediate_command_list -// This test violates the second requirement (explicitly requests -// no_immediate_command_list) and verifies the runtime rejects the import. -// -// This is a contract test, not an interop test: no real Vulkan context is -// created. The runtime rejects at import before inspecting the handle, so -// a bogus fd / handle is enough. That's also why this test does not gate -// on `vulkan` or link against the Vulkan loader -- it lives here because -// the Vulkan-flavored handle types are what most readers will look for. +// 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 "vulkan_setup.hpp" #include #include #include @@ -29,29 +31,50 @@ namespace syclexp = sycl::ext::oneapi::experimental; int main() { - sycl::queue q{ + VulkanContext vkCtx = createVulkanContext(); + VkSemaphore vkSem = createExportableSemaphore(vkCtx); + + // Lawful queue: import the semaphore here. + sycl::queue immQ{ {sycl::property::queue::in_order{}, - sycl::ext::intel::property::queue::no_immediate_command_list{}}}; + 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{ - /*handle=*/nullptr, - syclexp::external_semaphore_handle_type::win32_nt_handle}; + semHandle, syclexp::external_semaphore_handle_type::win32_nt_handle}; #else + int semFd = getSemaphoreFd(vkCtx, vkSem); syclexp::external_semaphore_descriptor desc{ - /*file_descriptor=*/-1, - syclexp::external_semaphore_handle_type::opaque_fd}; + 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 { - (void)syclexp::import_external_semaphore(desc, q); + 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; - return 0; + ret = 0; } - std::cerr << "FAIL: import_external_semaphore on a non-immediate-CL queue " - "did not throw." - << std::endl; - return 1; + 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_buffer.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_buffer.cpp index 65572b382ec54..903525509b9e8 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 @@ -1,9 +1,6 @@ // REQUIRES: aspect-ext_oneapi_external_memory_import // REQUIRES: vulkan -// UNSUPPORTED: windows && gpu-intel-gen12 -// UNSUPPORTED-TRACKER: URLZA-723 - // clang-format off // On Linux L0, there are problem with semaphores and latest drivers. From d49c119dea7952f1395e49baf514e048b534ee6e Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Fri, 15 May 2026 15:49:24 -0700 Subject: [PATCH 04/17] validation optional --- .../vulkan_interop/vulkan_setup.hpp | 23 +++++++++++++++---- 1 file changed, 18 insertions(+), 5 deletions(-) diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_setup.hpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_setup.hpp index 3bd402194cc94..6827e30727856 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_setup.hpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_setup.hpp @@ -334,6 +334,11 @@ inline VulkanContext createVulkanContext() { appInfo.apiVersion = VK_API_VERSION_1_2; std::vector instanceExtensions; + VkInstanceCreateInfo createInfo{}; + createInfo.sType = VK_STRUCTURE_TYPE_INSTANCE_CREATE_INFO; + createInfo.pApplicationInfo = &appInfo; + +#ifndef ONT_VALIDATE_FOR_THE_LOVE_OF_ALL_THATS_HOLY_WHAT_WERE_YOU_THINKING uint32_t instanceExtensionCount = 0; VK_CHECK(vkEnumerateInstanceExtensionProperties( nullptr, &instanceExtensionCount, nullptr)); @@ -373,13 +378,8 @@ inline VulkanContext createVulkanContext() { throw std::runtime_error("failed to find Vulkan validation layer!"); } - VkInstanceCreateInfo createInfo{}; - createInfo.sType = VK_STRUCTURE_TYPE_INSTANCE_CREATE_INFO; - createInfo.pApplicationInfo = &appInfo; createInfo.ppEnabledLayerNames = &validationLayerName; createInfo.enabledLayerCount = 1; - createInfo.enabledExtensionCount = (uint32_t)instanceExtensions.size(); - createInfo.ppEnabledExtensionNames = instanceExtensions.data(); VkDebugUtilsMessengerCreateInfoEXT debugUtilsMessengerInfo{}; debugUtilsMessengerInfo.sType = @@ -397,8 +397,16 @@ inline VulkanContext createVulkanContext() { // that occur while creating the instance itself. This debug util messangegr // will be alive only during instance creation. createInfo.pNext = &debugUtilsMessengerInfo; +#else + ctx.debugMessenger = VK_NULL_HANDLE; +#endif + + createInfo.enabledExtensionCount = (uint32_t)instanceExtensions.size(); + createInfo.ppEnabledExtensionNames = instanceExtensions.data(); + VK_CHECK(vkCreateInstance(&createInfo, nullptr, &ctx.instance)); +#ifndef ONT_VALIDATE_FOR_THE_LOVE_OF_ALL_THATS_HOLY_WHAT_WERE_YOU_THINKING // Create a persistent debug messenger that stays alive for the application's // lifetime to capture all subsequent events. auto vkCreateDebugUtilsMessengerFuncPtr = @@ -411,6 +419,7 @@ inline VulkanContext createVulkanContext() { throw std::runtime_error( "Failed to fetch vkCreateDebugUtilsMessengerEXT function pointer!"); } +#endif uint32_t deviceCount = 0; vkEnumeratePhysicalDevices(ctx.instance, &deviceCount, nullptr); @@ -466,6 +475,7 @@ inline VulkanContext createVulkanContext() { inline void cleanupVulkanContext(VulkanContext &ctx) { vkDestroyDevice(ctx.device, nullptr); +#ifndef ONT_VALIDATE_FOR_THE_LOVE_OF_ALL_THATS_HOLY_WHAT_WERE_YOU_THINKING auto vkDestroyDebugUtilsMessengerFuncPtr = (PFN_vkDestroyDebugUtilsMessengerEXT)vkGetInstanceProcAddr( ctx.instance, "vkDestroyDebugUtilsMessengerEXT"); @@ -473,6 +483,7 @@ inline void cleanupVulkanContext(VulkanContext &ctx) { vkDestroyDebugUtilsMessengerFuncPtr(ctx.instance, ctx.debugMessenger, nullptr); } +#endif vkDestroyInstance(ctx.instance, nullptr); } @@ -1230,6 +1241,7 @@ inline void cleanupVulkan(VulkanContext &ctx, ImageResources &res) { vkDestroyImage(ctx.device, res.image, nullptr); vkFreeMemory(ctx.device, res.memory, nullptr); vkDestroyDevice(ctx.device, nullptr); +#ifndef ONT_VALIDATE_FOR_THE_LOVE_OF_ALL_THATS_HOLY_WHAT_WERE_YOU_THINKING auto vkDestroyDebugUtilsMessengerFuncPtr = (PFN_vkDestroyDebugUtilsMessengerEXT)vkGetInstanceProcAddr( ctx.instance, "vkDestroyDebugUtilsMessengerEXT"); @@ -1237,5 +1249,6 @@ inline void cleanupVulkan(VulkanContext &ctx, ImageResources &res) { vkDestroyDebugUtilsMessengerFuncPtr(ctx.instance, ctx.debugMessenger, nullptr); } +#endif vkDestroyInstance(ctx.instance, nullptr); } From eb3f2f10c4e10f6367519272e832299f219141df Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Fri, 15 May 2026 15:57:19 -0700 Subject: [PATCH 05/17] spelling correction --- .../bindless_images/vulkan_interop/vulkan_setup.hpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_setup.hpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_setup.hpp index 6827e30727856..7fa71442bb168 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_setup.hpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_setup.hpp @@ -338,7 +338,7 @@ inline VulkanContext createVulkanContext() { createInfo.sType = VK_STRUCTURE_TYPE_INSTANCE_CREATE_INFO; createInfo.pApplicationInfo = &appInfo; -#ifndef ONT_VALIDATE_FOR_THE_LOVE_OF_ALL_THATS_HOLY_WHAT_WERE_YOU_THINKING +#ifndef ONT_VALIDATE uint32_t instanceExtensionCount = 0; VK_CHECK(vkEnumerateInstanceExtensionProperties( nullptr, &instanceExtensionCount, nullptr)); @@ -406,7 +406,7 @@ inline VulkanContext createVulkanContext() { VK_CHECK(vkCreateInstance(&createInfo, nullptr, &ctx.instance)); -#ifndef ONT_VALIDATE_FOR_THE_LOVE_OF_ALL_THATS_HOLY_WHAT_WERE_YOU_THINKING +#ifndef ONT_VALIDATE // Create a persistent debug messenger that stays alive for the application's // lifetime to capture all subsequent events. auto vkCreateDebugUtilsMessengerFuncPtr = @@ -475,7 +475,7 @@ inline VulkanContext createVulkanContext() { inline void cleanupVulkanContext(VulkanContext &ctx) { vkDestroyDevice(ctx.device, nullptr); -#ifndef ONT_VALIDATE_FOR_THE_LOVE_OF_ALL_THATS_HOLY_WHAT_WERE_YOU_THINKING +#ifndef ONT_VALIDATE auto vkDestroyDebugUtilsMessengerFuncPtr = (PFN_vkDestroyDebugUtilsMessengerEXT)vkGetInstanceProcAddr( ctx.instance, "vkDestroyDebugUtilsMessengerEXT"); @@ -1241,7 +1241,7 @@ inline void cleanupVulkan(VulkanContext &ctx, ImageResources &res) { vkDestroyImage(ctx.device, res.image, nullptr); vkFreeMemory(ctx.device, res.memory, nullptr); vkDestroyDevice(ctx.device, nullptr); -#ifndef ONT_VALIDATE_FOR_THE_LOVE_OF_ALL_THATS_HOLY_WHAT_WERE_YOU_THINKING +#ifndef ONT_VALIDATE auto vkDestroyDebugUtilsMessengerFuncPtr = (PFN_vkDestroyDebugUtilsMessengerEXT)vkGetInstanceProcAddr( ctx.instance, "vkDestroyDebugUtilsMessengerEXT"); From 3f0ce1ebbbb50d6d1aa5349b14444863e8ec8220 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Fri, 15 May 2026 16:29:27 -0700 Subject: [PATCH 06/17] DG2 testing going well. Reenabling lots of tests --- .../D3D12_sycl_buffer_timeline_semaphore.cpp | 4 -- .../D3D12_sycl_buffer_win32_name_native.cpp | 3 -- .../D3D12_sycl_interop_1D_read.cpp | 26 ++++++------- .../D3D12_sycl_interop_1D_write_unsampled.cpp | 28 +++++++------- .../D3D12_sycl_interop_2D_arithmetic.cpp | 27 +++++++------ .../D3D12_sycl_interop_2D_read.cpp | 27 +++++++------ .../D3D12_sycl_interop_2D_write_unsampled.cpp | 28 +++++++------- .../D3D12_sycl_interop_3D_read.cpp | 32 ++++++++-------- .../D3D12_sycl_interop_3D_write_unsampled.cpp | 38 ++++++++++++------- .../vulkan_sycl_2d_arithmetic.cpp | 18 ++++----- .../vulkan_sycl_buffer_binary_semaphore.cpp | 3 -- 11 files changed, 117 insertions(+), 117 deletions(-) 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 28f65653b5161..f54b298ea260f 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 @@ -2,10 +2,6 @@ // REQUIRES: aspect-ext_oneapi_external_semaphore_import // REQUIRES: windows -// UNSUPPORTED: gpu-intel-dg2 -// UNSUPPORTED-TRACKER: GSD-12428 -// semaphores-do-not-work-dg2 - // UNSUPPORTED: gpu-intel-gen12 // UNSUPPORTED-TRACKER: GSD-12427 // Gen12-semaphores-work-but-this-test-hangs. diff --git a/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_win32_name_native.cpp b/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_win32_name_native.cpp index 581e48d3188ec..75b0d5f767498 100644 --- a/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_win32_name_native.cpp +++ b/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_win32_name_native.cpp @@ -2,9 +2,6 @@ // REQUIRES: aspect-ext_oneapi_external_semaphore_import // REQUIRES: windows -// UNSUPPORTED: gpu-intel-dg2 -// UNSUPPORTED-TRACKER: GSD-12428 - // UNSUPPORTED: gpu-intel-gen12 // UNSUPPORTED-TRACKER: GSD-12427 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 e1e2ef6301a52..72b35e05c2292 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 @@ -74,19 +74,19 @@ // Semaphore coverage tests -// At this time, semaphores aren't working on DG2 (GSD-12428), and can hang on BMG if run in parallel (GSD-12436). -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 4 --semaphores 33x -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 2 --semaphores 33x -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type int32 --channels 1 --semaphores 33x -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 4 --semaphores 33x -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type int16 --channels 2 --semaphores 33x -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint16 --channels 1 --semaphores 33x -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint8 --channels 4 --semaphores 33x -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type int8 --channels 2 --semaphores 33x -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 4 --sampled --semaphores 33x -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 2 --sampled --semaphores 33x -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type int32 --channels 1 --sampled --semaphores 33x -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 4 --sampled --semaphores 33x +// At this time, semaphores can hang on BMG if run in parallel (GSD-12436). +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 4 --semaphores 33x +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 2 --semaphores 33x +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int32 --channels 1 --semaphores 33x +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 4 --semaphores 33x +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int16 --channels 2 --semaphores 33x +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint16 --channels 1 --semaphores 33x +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint8 --channels 4 --semaphores 33x +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int8 --channels 2 --semaphores 33x +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 4 --sampled --semaphores 33x +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 2 --sampled --semaphores 33x +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int32 --channels 1 --sampled --semaphores 33x +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 4 --sampled --semaphores 33x // clang-format on 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 c3ba154eeaa0c..96aecac90d79b 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 @@ -7,7 +7,7 @@ // clang-format off /* - clang++.exe -fsycl -o ds1w.exe D3D12_sycl_interop_1D_write.cpp -ld3d12 -ldxgi -ld3dcompiler + clang++.exe -fsycl -o ds1w.exe D3D12_sycl_interop_1D_write_unsampled.cpp -ld3d12 -ldxgi -ld3dcompiler FLAGS: --sampled ERROR: Sampled image writes are not supported @@ -47,19 +47,19 @@ // RUN: %{run} %t.exe --type unorm8 --channels 4 33x // Semaphore coverage tests -// At this time, semaphores aren't working on DG2 (GSD-12428), and can hang on BMG if run in parallel (GSD-12436). -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 4 --semaphores 33x -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 1 --semaphores 33x -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 2 --semaphores 33x -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type int32 --channels 4 --semaphores 33x -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 1 --semaphores 33x -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type int16 --channels 2 --semaphores 33x -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint16 --channels 4 --semaphores 33x -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint8 --channels 1 --semaphores 33x -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type int8 --channels 2 --semaphores 33x -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type unorm8 --channels 4 --semaphores 33x -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 4 --semaphores 33x -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 2 --semaphores 33x +// At this time, semaphores can hang on BMG if run in parallel (GSD-12436). +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 4 --semaphores 33x +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 1 --semaphores 33x +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 2 --semaphores 33x +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int32 --channels 4 --semaphores 33x +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 1 --semaphores 33x +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int16 --channels 2 --semaphores 33x +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint16 --channels 4 --semaphores 33x +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint8 --channels 1 --semaphores 33x +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int8 --channels 2 --semaphores 33x +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type unorm8 --channels 4 --semaphores 33x +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 4 --semaphores 33x +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 2 --semaphores 33x // clang-format on 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 49e185d282f35..b8c9a84880ea8 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 @@ -26,7 +26,6 @@ DG2: - WORKS, including --sampled - - semaphores segfault DG2 $ sycl-ls [level_zero:gpu][level_zero:0] Intel(R) oneAPI Unified Runtime over @@ -93,19 +92,19 @@ // RUN-IF: !gpu-intel-bmg, %{run} %t.exe --type unorm8 --channels 4 --sampled 32x33 // Semaphore coverage tests -// At this time, semaphores aren't working on DG2 (GSD-12428), and can hang on BMG if run in parallel (GSD-12436). -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 4 --semaphores 32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 2 --semaphores 32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type int32 --channels 1 --semaphores 32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 4 --semaphores 32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type int16 --channels 2 --semaphores 32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint16 --channels 1 --semaphores 32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint8 --channels 4 --semaphores 32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type int8 --channels 2 --semaphores 32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 4 --sampled --semaphores 32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 2 --sampled --semaphores 32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type int32 --channels 1 --sampled --semaphores 32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type unorm8 --channels 4 --sampled --semaphores 32x33 +// At this time, semaphores can hang on BMG if run in parallel (GSD-12436). +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 4 --semaphores 32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 2 --semaphores 32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int32 --channels 1 --semaphores 32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 4 --semaphores 32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int16 --channels 2 --semaphores 32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint16 --channels 1 --semaphores 32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint8 --channels 4 --semaphores 32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int8 --channels 2 --semaphores 32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 4 --sampled --semaphores 32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 2 --sampled --semaphores 32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int32 --channels 1 --sampled --semaphores 32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type unorm8 --channels 4 --sampled --semaphores 32x33 // clang-format on 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 e51b8172df4d8..b956c0c86a19b 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 @@ -35,7 +35,6 @@ DG2: - WORKS, including --sampled - - semaphores segfault DG2 $ sycl-ls [level_zero:gpu][level_zero:0] Intel(R) oneAPI Unified Runtime over @@ -98,19 +97,19 @@ // RUN: %{run} %t.exe --type int8 --channels 4 --sampled 32x33 // Semaphore coverage tests -// At this time, semaphores aren't working on DG2 (GSD-12428), and can hang on BMG if run in parallel (GSD-12436). -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 4 --semaphores 32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 2 --semaphores 32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type int32 --channels 1 --semaphores 32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 4 --semaphores 32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type int16 --channels 2 --semaphores 32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint16 --channels 1 --semaphores 32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint8 --channels 4 --semaphores 32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type int8 --channels 2 --semaphores 32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 4 --sampled --semaphores 32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 2 --sampled --semaphores 32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type int32 --channels 1 --sampled --semaphores 32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 4 --sampled --semaphores 32x33 +// At this time, semaphores can hang on BMG if run in parallel (GSD-12436). +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 4 --semaphores 32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 2 --semaphores 32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int32 --channels 1 --semaphores 32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 4 --semaphores 32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int16 --channels 2 --semaphores 32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint16 --channels 1 --semaphores 32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint8 --channels 4 --semaphores 32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int8 --channels 2 --semaphores 32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 4 --sampled --semaphores 32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 2 --sampled --semaphores 32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int32 --channels 1 --sampled --semaphores 32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 4 --sampled --semaphores 32x33 // clang-format on 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 fd59cdfcd0286..40b77c2f90f20 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 @@ -7,7 +7,7 @@ // clang-format off /* - clang++.exe -fsycl -o ds2w.exe D3D12_sycl_interop_2D_write.cpp -ld3d12 -ldxgi -ld3dcompiler + clang++.exe -fsycl -o ds2w.exe D3D12_sycl_interop_2D_write_unsampled.cpp -ld3d12 -ldxgi -ld3dcompiler FLAGS: --sampled ERROR: Sampled image writes are not supported @@ -47,19 +47,19 @@ // RUN: %{run} %t.exe --type unorm8 --channels 4 32x33 // Semaphore coverage tests -// At this time, semaphores aren't working on DG2 (GSD-12428), and can hang on BMG if run in parallel (GSD-12436). -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 4 --semaphores 32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 1 --semaphores 32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 2 --semaphores 32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type int32 --channels 4 --semaphores 32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 1 --semaphores 32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type int16 --channels 2 --semaphores 32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint16 --channels 4 --semaphores 32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint8 --channels 1 --semaphores 32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type int8 --channels 2 --semaphores 32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type unorm8 --channels 4 --semaphores 32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 4 --semaphores 32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 2 --semaphores 32x33 +// At this time, semaphores can hang on BMG if run in parallel (GSD-12436). +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 4 --semaphores 32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 1 --semaphores 32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 2 --semaphores 32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int32 --channels 4 --semaphores 32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 1 --semaphores 32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int16 --channels 2 --semaphores 32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint16 --channels 4 --semaphores 32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint8 --channels 1 --semaphores 32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int8 --channels 2 --semaphores 32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type unorm8 --channels 4 --semaphores 32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 4 --semaphores 32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 2 --semaphores 32x33 // clang-format on 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 5e2d44af8978f..7ccf60be5ae76 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 @@ -6,12 +6,12 @@ // UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/20384 // also GSD-12430 -// UNSUPPORTED: gpu-intel-dg2 -// UNSUPPORTED-TRACKER: GSD-12430 - // RUN: %{build} -o %t.exe %link-directx // RUN: %{run} %t.exe --type float --channels 4 8x8x8 + +// clang++.exe -fsycl -o ds3r.exe D3D12_sycl_interop_3D_read.cpp -ld3d12 -ldxgi -ld3dcompiler + // clang-format off // RUN: %{run} %t.exe --type float --channels 1 33x32x31 @@ -67,19 +67,19 @@ // RUN-IF: !gpu-intel-gen12, %{run} %t.exe --type unorm8 --channels 4 --sampled 7x9x8 // Semaphore coverage tests -// At this time, semaphores aren't working on DG2 (GSD-12428), and can hang on BMG if run in parallel (GSD-12436). -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 4 --semaphores 16x17x15 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 2 --semaphores 17x16x15 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type int32 --channels 1 --semaphores 9x8x7 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 4 --semaphores 33x31x32 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type int16 --channels 2 --semaphores 15x17x16 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint16 --channels 1 --semaphores 9x7x8 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint8 --channels 4 --semaphores 32x31x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type int8 --channels 2 --semaphores 16x15x17 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 4 --sampled --semaphores 31x32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 2 --sampled --semaphores 15x16x17 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type int32 --channels 1 --sampled --semaphores 7x8x9 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 4 --sampled --semaphores 32x31x33 +// At this time, semaphores can hang on BMG if run in parallel (GSD-12436). +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 4 --semaphores 16x17x15 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 2 --semaphores 17x16x15 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int32 --channels 1 --semaphores 9x8x7 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 4 --semaphores 33x31x32 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int16 --channels 2 --semaphores 15x17x16 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint16 --channels 1 --semaphores 9x7x8 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint8 --channels 4 --semaphores 32x31x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int8 --channels 2 --semaphores 16x15x17 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 4 --sampled --semaphores 31x32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 2 --sampled --semaphores 15x16x17 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int32 --channels 1 --sampled --semaphores 7x8x9 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 4 --sampled --semaphores 32x31x33 /* clang++.exe -fsycl -o ds3r.exe D3D12_sycl_interop_3D_read.cpp -ld3d12 -ldxgi -ld3dcompiler 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 220b201588a4e..78a55c9720e39 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 @@ -5,6 +5,18 @@ // RUN: %{build} -o %t.exe %link-directx // RUN: %{run} %t.exe --type float --channels 4 8x8x8 +/* + clang++.exe -fsycl -o ds3w.exe D3D12_sycl_interop_3D_write_unsampled.cpp -ld3d12 -ldxgi -ld3dcompiler + + FLAGS: + --sampled ERROR: Sampled image writes are not supported + --semaphores Use DX12 Fences for SYCL Interop Sync + --channels X Set number of channels (1, 2, or 4). Default is 4 (RGBA) + --type XXX Set data type (float, half, uint32, int32, uint16, int16, uint8, int8, unorm8). + Default is float + WxHxD Set custom Width x Height x Depth (e.g. 8x4x2) +*/ + // clang-format off // RUN: %{run} %t.exe --type float --channels 1 33x32x31 @@ -36,19 +48,19 @@ // RUN: %{run} %t.exe --type unorm8 --channels 4 7x9x8 // Semaphore coverage tests -// At this time, semaphores aren't working on DG2 (GSD-12428), and can hang on BMG if run in parallel (GSD-12436). -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 4 --semaphores 16x17x15 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 1 --semaphores 31x32x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 2 --semaphores 17x16x15 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type int32 --channels 4 --semaphores 9x8x7 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 1 --semaphores 33x31x32 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type int16 --channels 2 --semaphores 15x17x16 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint16 --channels 4 --semaphores 9x7x8 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint8 --channels 1 --semaphores 32x31x33 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type int8 --channels 2 --semaphores 16x15x17 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type unorm8 --channels 4 --semaphores 7x9x8 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 4 --semaphores 15x16x17 -// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 2 --semaphores 32x31x33 +// At this time, semaphores can hang on BMG if run in parallel (GSD-12436). +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 4 --semaphores 16x17x15 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 1 --semaphores 31x32x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 2 --semaphores 17x16x15 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int32 --channels 4 --semaphores 9x8x7 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 1 --semaphores 33x31x32 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int16 --channels 2 --semaphores 15x17x16 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint16 --channels 4 --semaphores 9x7x8 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint8 --channels 1 --semaphores 32x31x33 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int8 --channels 2 --semaphores 16x15x17 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type unorm8 --channels 4 --semaphores 7x9x8 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 4 --semaphores 15x16x17 +// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 2 --semaphores 32x31x33 /* clang++.exe -fsycl -o ds3w.exe D3D12_sycl_interop_3D_write.cpp -ld3d12 -ldxgi -ld3dcompiler 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 d9cd21b0bfb2f..79095298c6b10 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 @@ -7,8 +7,8 @@ // UNSUPPORTED: linux // UNSUPPORTED-TRACKER: GSD-12357 -// XFAIL: windows && gpu-intel-dg2 -// XFAIL-TRACKER: https://github.com/intel/llvm/issues/21985 + +// clang-format off /* Run ALL the vulkan formats through the gauntlet. sampled and unsampled. @@ -102,24 +102,24 @@ /* Vulkan/SYCL 2D Arithmetic (A + B = C) - clang++ -fsycl -o vs_2d_arith.bin vulkan_sycl_2d_arithmetic.cpp -lvulkan - -I$VULKAN_SDK/include -L$VULKAN_SDK/lib + clang++ -fsycl -o vs_2d_arith.bin vulkan_sycl_2d_arithmetic.cpp -lvulkan -I$VULKAN_SDK/include -L$VULKAN_SDK/lib - clang++ -fsycl -o vs_2d_arith.exe vulkan_sycl_2d_arithmetic.cpp - -Wno-ignored-attributes -lvulkan-1 -I$VULKAN_SDK/Include -L$VULKAN_SDK/Lib + clang++ -fsycl -o vs_2d_arith.exe vulkan_sycl_2d_arithmetic.cpp -Wno-ignored-attributes -lvulkan-1 -I$VULKAN_SDK/Include -L$VULKAN_SDK/Lib FLAGS --semaphores Use Vulkan Semaphores for SYCL Interop Sync --linear Use LINEAR tiling for the Vulkan Image (default is OPTIMAL) --channels X Set number of channels (1, 2, or 4). Default is 4 (RGBA) - --type XXX Set data type (float, half, uint32, int32, uint16, int16, - uint8, int8, unorm8). Default is float WxH Set custom Width x - Height (e.g. 8x4) + --type XXX Set data type (float, half, uint32, int32, uint16, int16, uint8, int8, unorm8). Default is float + WxH Set custom Width x Height (e.g. 8x4) --sampled // RUN: %{run} %t.out --type float --semaphores // RUN: %{run} %t.out --type unorm8 --sampled --semaphores */ + +// clang-format on + #include "vulkan_setup.hpp" #include 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 e4d4cb3e663c1..a1ebda0f09072 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 @@ -8,9 +8,6 @@ // UNSUPPORTED: windows && gpu-intel-gen12 // UNSUPPORTED-TRACKER: URLZA-723 -// XFAIL: windows && gpu-intel-dg2 -// XFAIL-TRACKER: https://github.com/intel/llvm/issues/21985 - // UNSUPPORTED: windows && arch-intel_gpu_bmg_g21 // UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/21986 From 16f8fa2c18fb0aecd2068d1a90018002cc8119ff Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Mon, 18 May 2026 12:00:14 -0700 Subject: [PATCH 07/17] exploratory Signed-off-by: Chris Perkins --- .../dx12_interop/D3D12_sycl_buffer_timeline_semaphore.cpp | 4 ++++ .../dx12_interop/D3D12_sycl_buffer_win32_name_native.cpp | 4 ++++ .../dx12_interop/D3D12_sycl_interop_3D_read.cpp | 6 ++++-- .../vulkan_sycl_buffer_timeline_semaphore.cpp | 3 --- .../vulkan_interop/vulkan_sycl_image_interop_read_1d.cpp | 3 --- .../vulkan_interop/vulkan_sycl_image_interop_read_2d.cpp | 3 --- .../vulkan_sycl_image_interop_write_2d_unsampled.cpp | 3 --- .../vulkan_sycl_image_unsampled_timeline_semaphore.cpp | 3 --- 8 files changed, 12 insertions(+), 17 deletions(-) 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 f54b298ea260f..28f65653b5161 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 @@ -2,6 +2,10 @@ // REQUIRES: aspect-ext_oneapi_external_semaphore_import // REQUIRES: windows +// UNSUPPORTED: gpu-intel-dg2 +// UNSUPPORTED-TRACKER: GSD-12428 +// semaphores-do-not-work-dg2 + // UNSUPPORTED: gpu-intel-gen12 // UNSUPPORTED-TRACKER: GSD-12427 // Gen12-semaphores-work-but-this-test-hangs. diff --git a/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_win32_name_native.cpp b/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_win32_name_native.cpp index 75b0d5f767498..2768cb386999f 100644 --- a/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_win32_name_native.cpp +++ b/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_win32_name_native.cpp @@ -2,6 +2,10 @@ // REQUIRES: aspect-ext_oneapi_external_semaphore_import // REQUIRES: windows +// UNSUPPORTED: gpu-intel-dg2 +// UNSUPPORTED-TRACKER: GSD-12428 +// semaphores-do-not-work-dg2 + // UNSUPPORTED: gpu-intel-gen12 // UNSUPPORTED-TRACKER: GSD-12427 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 7ccf60be5ae76..56483ca32d6b3 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 @@ -6,14 +6,16 @@ // UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/20384 // also GSD-12430 +// UNSUPPORTED: gpu-intel-dg2 +// UNSUPPORTED-TRACKER: GSD-12430 + // RUN: %{build} -o %t.exe %link-directx // RUN: %{run} %t.exe --type float --channels 4 8x8x8 +// clang-format off // clang++.exe -fsycl -o ds3r.exe D3D12_sycl_interop_3D_read.cpp -ld3d12 -ldxgi -ld3dcompiler -// clang-format off - // RUN: %{run} %t.exe --type float --channels 1 33x32x31 // RUN: %{run} %t.exe --type float --channels 2 32x33x31 // RUN: %{run} %t.exe --type float --channels 4 31x32x33 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 fcba2223ab136..f574416a51095 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 @@ -10,9 +10,6 @@ // UNSUPPORTED: windows && gpu-intel-gen12 // UNSUPPORTED-TRACKER: URLZA-723 -// XFAIL: windows && gpu-intel-dg2 -// XFAIL-TRACKER: https://github.com/intel/llvm/issues/21985 - // UNSUPPORTED: windows && arch-intel_gpu_bmg_g21 // UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/21986 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 d0f429456f575..762b1847bd4b6 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 @@ -2,9 +2,6 @@ // REQUIRES: aspect-ext_oneapi_external_memory_import || (windows && level_zero && aspect-ext_oneapi_bindless_images) // REQUIRES: vulkan -// XFAIL: windows && gpu-intel-dg2 -// XFAIL-TRACKER: https://github.com/intel/llvm/issues/21985 - // RUN: %{build} %link-vulkan -o %t.out %if target-spir %{ -Wno-ignored-attributes %} /* 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 ff9626d70a4a9..cc784bd8c269e 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 @@ -7,9 +7,6 @@ // UNSUPPORTED: linux // UNSUPPORTED-TRACKER: GSD-12357 -// XFAIL: windows && gpu-intel-dg2 -// XFAIL-TRACKER: https://github.com/intel/llvm/issues/21985 - // XFAIL: windows && arch-intel_gpu_bmg_g21 // XFAIL-TRACKER: https://github.com/intel/llvm/issues/21986 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 09a429424d541..1675e16f985c5 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 @@ -7,9 +7,6 @@ // UNSUPPORTED: linux // UNSUPPORTED-TRACKER: GSD-12357 -// XFAIL: windows && gpu-intel-dg2 -// XFAIL-TRACKER: https://github.com/intel/llvm/issues/21985 - // UNSUPPORTED: windows && arch-intel_gpu_bmg_g21 // UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/21986 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 228b825ec45c1..dd3ea3ee5dcd1 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 @@ -5,9 +5,6 @@ // UNSUPPORTED: linux && gpu-intel-dg2 // UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/21764 -// XFAIL: windows && gpu-intel-dg2 -// XFAIL-TRACKER: https://github.com/intel/llvm/issues/21985 - // RUN: %{build} %link-vulkan -o %t.out %if target-spir %{ -Wno-ignored-attributes %} // RUN: %{run} %t.out --type float --channels 1 From 77ad8f390ac3a33b6344b2f9bc7baf15e214001a Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Mon, 18 May 2026 13:56:26 -0700 Subject: [PATCH 08/17] clang-format never fails Signed-off-by: Chris Perkins --- .../D3D12_sycl_interop_3D_write_unsampled.cpp | 14 +++----------- 1 file changed, 3 insertions(+), 11 deletions(-) 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 78a55c9720e39..a6913ef3549d0 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 @@ -5,6 +5,8 @@ // RUN: %{build} -o %t.exe %link-directx // RUN: %{run} %t.exe --type float --channels 4 8x8x8 +// clang-format off + /* clang++.exe -fsycl -o ds3w.exe D3D12_sycl_interop_3D_write_unsampled.cpp -ld3d12 -ldxgi -ld3dcompiler @@ -17,7 +19,7 @@ WxHxD Set custom Width x Height x Depth (e.g. 8x4x2) */ -// clang-format off + // RUN: %{run} %t.exe --type float --channels 1 33x32x31 // RUN: %{run} %t.exe --type float --channels 2 32x33x31 @@ -62,17 +64,7 @@ // RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 4 --semaphores 15x16x17 // RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 2 --semaphores 32x31x33 -/* - clang++.exe -fsycl -o ds3w.exe D3D12_sycl_interop_3D_write.cpp -ld3d12 -ldxgi -ld3dcompiler - FLAGS: - --sampled ERROR: Sampled image writes are not supported - --semaphores Use DX12 Fences for SYCL Interop Sync - --channels X Set number of channels (1, 2, or 4). Default is 4 (RGBA) - --type XXX Set data type (float, half, uint32, int32, uint16, int16, uint8, int8, unorm8). - Default is float - WxHxD Set custom Width x Height x Depth (e.g. 8x4x41) -*/ // clang-format on #include "d3d12_setup.hpp" From cc223b3ec701eec8a818788d09dbccc4771cb02f Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Mon, 18 May 2026 14:18:57 -0700 Subject: [PATCH 09/17] clang-format is decidedly hostile. Signed-off-by: Chris Perkins --- .../dx12_interop/D3D12_sycl_interop_3D_write_unsampled.cpp | 1 - .../bindless_images/vulkan_interop/vulkan_sycl_2d_arithmetic.cpp | 1 - 2 files changed, 2 deletions(-) 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 a6913ef3549d0..ea6a83db0ed09 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 @@ -64,7 +64,6 @@ // RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 4 --semaphores 15x16x17 // RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 2 --semaphores 32x31x33 - // clang-format on #include "d3d12_setup.hpp" 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 79095298c6b10..fc4483b3d04b9 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 @@ -7,7 +7,6 @@ // UNSUPPORTED: linux // UNSUPPORTED-TRACKER: GSD-12357 - // clang-format off /* From 57c5d75069fef28a9725c61d067811b6f413b4f7 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Mon, 18 May 2026 16:49:32 -0700 Subject: [PATCH 10/17] BMG enablement for DX12 tests --- .../D3D12_sycl_buffer_timeline_semaphore.cpp | 5 --- .../D3D12_sycl_interop_1D_read.cpp | 25 ++++++----- .../D3D12_sycl_interop_1D_write_unsampled.cpp | 25 ++++++----- .../D3D12_sycl_interop_2D_arithmetic.cpp | 41 ++++++++----------- .../D3D12_sycl_interop_2D_read.cpp | 33 ++++++--------- .../D3D12_sycl_interop_2D_write_unsampled.cpp | 25 ++++++----- .../D3D12_sycl_interop_3D_write_unsampled.cpp | 25 ++++++----- 7 files changed, 79 insertions(+), 100 deletions(-) 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 28f65653b5161..fb1ff574d8c75 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 @@ -10,11 +10,6 @@ // UNSUPPORTED-TRACKER: GSD-12427 // Gen12-semaphores-work-but-this-test-hangs. -// UNSUPPORTED: arch-intel_gpu_bmg_g21 -// UNSUPPORTED-TRACKER: GSD-12436 -// this test works on BMG, but if run in parallel with itself, or with other -// semaphore tests, it can hang. - // RUN: %{build} %link-directx -o %t.exe %if target-spir %{ -Wno-ignored-attributes %} // RUN: %{run} %t.exe --no-sem // RUN: %{run} %t.exe 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 72b35e05c2292..9cf9a9e090b20 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 @@ -74,19 +74,18 @@ // Semaphore coverage tests -// At this time, semaphores can hang on BMG if run in parallel (GSD-12436). -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 4 --semaphores 33x -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 2 --semaphores 33x -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int32 --channels 1 --semaphores 33x -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 4 --semaphores 33x -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int16 --channels 2 --semaphores 33x -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint16 --channels 1 --semaphores 33x -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint8 --channels 4 --semaphores 33x -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int8 --channels 2 --semaphores 33x -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 4 --sampled --semaphores 33x -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 2 --sampled --semaphores 33x -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int32 --channels 1 --sampled --semaphores 33x -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 4 --sampled --semaphores 33x +// RUN: %{run} %t.exe --type float --channels 4 --semaphores 33x +// RUN: %{run} %t.exe --type half --channels 2 --semaphores 33x +// RUN: %{run} %t.exe --type int32 --channels 1 --semaphores 33x +// RUN: %{run} %t.exe --type uint32 --channels 4 --semaphores 33x +// RUN: %{run} %t.exe --type int16 --channels 2 --semaphores 33x +// RUN: %{run} %t.exe --type uint16 --channels 1 --semaphores 33x +// RUN: %{run} %t.exe --type uint8 --channels 4 --semaphores 33x +// RUN: %{run} %t.exe --type int8 --channels 2 --semaphores 33x +// RUN: %{run} %t.exe --type float --channels 4 --sampled --semaphores 33x +// RUN: %{run} %t.exe --type half --channels 2 --sampled --semaphores 33x +// RUN: %{run} %t.exe --type int32 --channels 1 --sampled --semaphores 33x +// RUN: %{run} %t.exe --type uint32 --channels 4 --sampled --semaphores 33x // clang-format on 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 96aecac90d79b..dbb172a650288 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 @@ -47,19 +47,18 @@ // RUN: %{run} %t.exe --type unorm8 --channels 4 33x // Semaphore coverage tests -// At this time, semaphores can hang on BMG if run in parallel (GSD-12436). -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 4 --semaphores 33x -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 1 --semaphores 33x -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 2 --semaphores 33x -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int32 --channels 4 --semaphores 33x -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 1 --semaphores 33x -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int16 --channels 2 --semaphores 33x -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint16 --channels 4 --semaphores 33x -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint8 --channels 1 --semaphores 33x -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int8 --channels 2 --semaphores 33x -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type unorm8 --channels 4 --semaphores 33x -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 4 --semaphores 33x -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 2 --semaphores 33x +// RUN: %{run} %t.exe --type float --channels 4 --semaphores 33x +// RUN: %{run} %t.exe --type float --channels 1 --semaphores 33x +// RUN: %{run} %t.exe --type half --channels 2 --semaphores 33x +// RUN: %{run} %t.exe --type int32 --channels 4 --semaphores 33x +// RUN: %{run} %t.exe --type uint32 --channels 1 --semaphores 33x +// RUN: %{run} %t.exe --type int16 --channels 2 --semaphores 33x +// RUN: %{run} %t.exe --type uint16 --channels 4 --semaphores 33x +// RUN: %{run} %t.exe --type uint8 --channels 1 --semaphores 33x +// RUN: %{run} %t.exe --type int8 --channels 2 --semaphores 33x +// RUN: %{run} %t.exe --type unorm8 --channels 4 --semaphores 33x +// RUN: %{run} %t.exe --type half --channels 4 --semaphores 33x +// RUN: %{run} %t.exe --type uint32 --channels 2 --semaphores 33x // clang-format on 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 b8c9a84880ea8..3d6470ae47b44 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 @@ -2,10 +2,6 @@ // REQUIRES: aspect-ext_oneapi_external_memory_import // REQUIRES: windows -// UNSUPPORTED: arch-intel_gpu_bmg_g21 -// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/20384 -// also GSD-12429 - // RUN: %{build} -o %t.exe %link-directx // RUN: %{run} %t.exe --type float --channels 4 8x8 @@ -60,9 +56,9 @@ // RUN: %{run} %t.exe --type int8 --channels 1 32x33 // RUN: %{run} %t.exe --type int8 --channels 2 32x33 // RUN: %{run} %t.exe --type int8 --channels 4 32x33 -// RUN-IF: !gpu-intel-bmg, %{run} %t.exe --type unorm8 --channels 1 32x33 -// RUN-IF: !gpu-intel-bmg, %{run} %t.exe --type unorm8 --channels 2 32x33 -// RUN-IF: !gpu-intel-bmg, %{run} %t.exe --type unorm8 --channels 4 32x33 +// RUN: %{run} %t.exe --type unorm8 --channels 1 32x33 +// RUN: %{run} %t.exe --type unorm8 --channels 2 32x33 +// RUN: %{run} %t.exe --type unorm8 --channels 4 32x33 // RUN: %{run} %t.exe --type float --channels 1 --sampled 32x33 // RUN: %{run} %t.exe --type float --channels 2 --sampled 32x33 // RUN: %{run} %t.exe --type float --channels 4 --sampled 32x33 @@ -87,24 +83,23 @@ // RUN: %{run} %t.exe --type int8 --channels 1 --sampled 32x33 // RUN: %{run} %t.exe --type int8 --channels 2 --sampled 32x33 // RUN: %{run} %t.exe --type int8 --channels 4 --sampled 32x33 -// RUN-IF: !gpu-intel-bmg, %{run} %t.exe --type unorm8 --channels 1 --sampled 32x33 -// RUN-IF: !gpu-intel-bmg, %{run} %t.exe --type unorm8 --channels 2 --sampled 32x33 -// RUN-IF: !gpu-intel-bmg, %{run} %t.exe --type unorm8 --channels 4 --sampled 32x33 +// RUN: %{run} %t.exe --type unorm8 --channels 1 --sampled 32x33 +// RUN: %{run} %t.exe --type unorm8 --channels 2 --sampled 32x33 +// RUN: %{run} %t.exe --type unorm8 --channels 4 --sampled 32x33 // Semaphore coverage tests -// At this time, semaphores can hang on BMG if run in parallel (GSD-12436). -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 4 --semaphores 32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 2 --semaphores 32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int32 --channels 1 --semaphores 32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 4 --semaphores 32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int16 --channels 2 --semaphores 32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint16 --channels 1 --semaphores 32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint8 --channels 4 --semaphores 32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int8 --channels 2 --semaphores 32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 4 --sampled --semaphores 32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 2 --sampled --semaphores 32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int32 --channels 1 --sampled --semaphores 32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type unorm8 --channels 4 --sampled --semaphores 32x33 +// RUN: %{run} %t.exe --type float --channels 4 --semaphores 32x33 +// RUN: %{run} %t.exe --type half --channels 2 --semaphores 32x33 +// RUN: %{run} %t.exe --type int32 --channels 1 --semaphores 32x33 +// RUN: %{run} %t.exe --type uint32 --channels 4 --semaphores 32x33 +// RUN: %{run} %t.exe --type int16 --channels 2 --semaphores 32x33 +// RUN: %{run} %t.exe --type uint16 --channels 1 --semaphores 32x33 +// RUN: %{run} %t.exe --type uint8 --channels 4 --semaphores 32x33 +// RUN: %{run} %t.exe --type int8 --channels 2 --semaphores 32x33 +// RUN: %{run} %t.exe --type float --channels 4 --sampled --semaphores 32x33 +// RUN: %{run} %t.exe --type half --channels 2 --sampled --semaphores 32x33 +// RUN: %{run} %t.exe --type int32 --channels 1 --sampled --semaphores 32x33 +// RUN: %{run} %t.exe --type unorm8 --channels 4 --sampled --semaphores 32x33 // clang-format on 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 b956c0c86a19b..737eda30047d4 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 @@ -2,10 +2,6 @@ // REQUIRES: aspect-ext_oneapi_external_memory_import // REQUIRES: windows -// UNSUPPORTED: arch-intel_gpu_bmg_g21 -// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/20384 -// also: GSD-12429 - // RUN: %{build} -o %t.exe %link-directx // RUN: %{run} %t.exe --type float --channels 4 32x33 @@ -23,9 +19,7 @@ WxH Set custom Width x Height (e.g. 8x4) - BMG: - - 1x1 works, nothing else. Suggesting offset/pitch issue. - - semaphores segfault. suggesting segfaulting semaphores. + BMG $ sycl-ls [level_zero:gpu][level_zero:0] Intel(R) oneAPI Unified Runtime over @@ -97,19 +91,18 @@ // RUN: %{run} %t.exe --type int8 --channels 4 --sampled 32x33 // Semaphore coverage tests -// At this time, semaphores can hang on BMG if run in parallel (GSD-12436). -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 4 --semaphores 32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 2 --semaphores 32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int32 --channels 1 --semaphores 32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 4 --semaphores 32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int16 --channels 2 --semaphores 32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint16 --channels 1 --semaphores 32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint8 --channels 4 --semaphores 32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int8 --channels 2 --semaphores 32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 4 --sampled --semaphores 32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 2 --sampled --semaphores 32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int32 --channels 1 --sampled --semaphores 32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 4 --sampled --semaphores 32x33 +// RUN: %{run} %t.exe --type float --channels 4 --semaphores 32x33 +// RUN: %{run} %t.exe --type half --channels 2 --semaphores 32x33 +// RUN: %{run} %t.exe --type int32 --channels 1 --semaphores 32x33 +// RUN: %{run} %t.exe --type uint32 --channels 4 --semaphores 32x33 +// RUN: %{run} %t.exe --type int16 --channels 2 --semaphores 32x33 +// RUN: %{run} %t.exe --type uint16 --channels 1 --semaphores 32x33 +// RUN: %{run} %t.exe --type uint8 --channels 4 --semaphores 32x33 +// RUN: %{run} %t.exe --type int8 --channels 2 --semaphores 32x33 +// RUN: %{run} %t.exe --type float --channels 4 --sampled --semaphores 32x33 +// RUN: %{run} %t.exe --type half --channels 2 --sampled --semaphores 32x33 +// RUN: %{run} %t.exe --type int32 --channels 1 --sampled --semaphores 32x33 +// RUN: %{run} %t.exe --type uint32 --channels 4 --sampled --semaphores 32x33 // clang-format on 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 40b77c2f90f20..40afca9508b14 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 @@ -47,19 +47,18 @@ // RUN: %{run} %t.exe --type unorm8 --channels 4 32x33 // Semaphore coverage tests -// At this time, semaphores can hang on BMG if run in parallel (GSD-12436). -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 4 --semaphores 32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 1 --semaphores 32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 2 --semaphores 32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int32 --channels 4 --semaphores 32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 1 --semaphores 32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int16 --channels 2 --semaphores 32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint16 --channels 4 --semaphores 32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint8 --channels 1 --semaphores 32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int8 --channels 2 --semaphores 32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type unorm8 --channels 4 --semaphores 32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 4 --semaphores 32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 2 --semaphores 32x33 +// RUN: %{run} %t.exe --type float --channels 4 --semaphores 32x33 +// RUN: %{run} %t.exe --type float --channels 1 --semaphores 32x33 +// RUN: %{run} %t.exe --type half --channels 2 --semaphores 32x33 +// RUN: %{run} %t.exe --type int32 --channels 4 --semaphores 32x33 +// RUN: %{run} %t.exe --type uint32 --channels 1 --semaphores 32x33 +// RUN: %{run} %t.exe --type int16 --channels 2 --semaphores 32x33 +// RUN: %{run} %t.exe --type uint16 --channels 4 --semaphores 32x33 +// RUN: %{run} %t.exe --type uint8 --channels 1 --semaphores 32x33 +// RUN: %{run} %t.exe --type int8 --channels 2 --semaphores 32x33 +// RUN: %{run} %t.exe --type unorm8 --channels 4 --semaphores 32x33 +// RUN: %{run} %t.exe --type half --channels 4 --semaphores 32x33 +// RUN: %{run} %t.exe --type uint32 --channels 2 --semaphores 32x33 // clang-format on 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 ea6a83db0ed09..bbf7207f2344f 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 @@ -50,19 +50,18 @@ // RUN: %{run} %t.exe --type unorm8 --channels 4 7x9x8 // Semaphore coverage tests -// At this time, semaphores can hang on BMG if run in parallel (GSD-12436). -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 4 --semaphores 16x17x15 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 1 --semaphores 31x32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 2 --semaphores 17x16x15 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int32 --channels 4 --semaphores 9x8x7 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 1 --semaphores 33x31x32 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int16 --channels 2 --semaphores 15x17x16 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint16 --channels 4 --semaphores 9x7x8 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint8 --channels 1 --semaphores 32x31x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int8 --channels 2 --semaphores 16x15x17 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type unorm8 --channels 4 --semaphores 7x9x8 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 4 --semaphores 15x16x17 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 2 --semaphores 32x31x33 +// RUN: %{run} %t.exe --type float --channels 4 --semaphores 16x17x15 +// RUN: %{run} %t.exe --type float --channels 1 --semaphores 31x32x33 +// RUN: %{run} %t.exe --type half --channels 2 --semaphores 17x16x15 +// RUN: %{run} %t.exe --type int32 --channels 4 --semaphores 9x8x7 +// RUN: %{run} %t.exe --type uint32 --channels 1 --semaphores 33x31x32 +// RUN: %{run} %t.exe --type int16 --channels 2 --semaphores 15x17x16 +// RUN: %{run} %t.exe --type uint16 --channels 4 --semaphores 9x7x8 +// RUN: %{run} %t.exe --type uint8 --channels 1 --semaphores 32x31x33 +// RUN: %{run} %t.exe --type int8 --channels 2 --semaphores 16x15x17 +// RUN: %{run} %t.exe --type unorm8 --channels 4 --semaphores 7x9x8 +// RUN: %{run} %t.exe --type half --channels 4 --semaphores 15x16x17 +// RUN: %{run} %t.exe --type uint32 --channels 2 --semaphores 32x31x33 // clang-format on From ced6462a93af97cfd486ec5fd3f78e35f1b9e292 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Mon, 18 May 2026 17:42:39 -0700 Subject: [PATCH 11/17] enable some moar vulkan tests --- .../vulkan_interop/vulkan_sycl_buffer_binary_semaphore.cpp | 3 --- .../vulkan_interop/vulkan_sycl_buffer_timeline_semaphore.cpp | 3 --- .../vulkan_interop/vulkan_sycl_image_interop_read_2d.cpp | 3 --- .../vulkan_sycl_image_interop_write_1d_unsampled.cpp | 2 +- .../vulkan_sycl_image_interop_write_2d_unsampled.cpp | 3 --- 5 files changed, 1 insertion(+), 13 deletions(-) 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 a1ebda0f09072..67b0146ae5693 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 @@ -8,9 +8,6 @@ // UNSUPPORTED: windows && gpu-intel-gen12 // UNSUPPORTED-TRACKER: URLZA-723 -// UNSUPPORTED: windows && arch-intel_gpu_bmg_g21 -// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/21986 - // RUN: %{build} %link-vulkan -o %t.out %if target-spir %{ -Wno-ignored-attributes %} // RUN: %{run} %t.out --no-sem // RUN: %{run} %t.out --dual-sem 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 f574416a51095..e0662b192a716 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 @@ -10,9 +10,6 @@ // UNSUPPORTED: windows && gpu-intel-gen12 // UNSUPPORTED-TRACKER: URLZA-723 -// UNSUPPORTED: windows && arch-intel_gpu_bmg_g21 -// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/21986 - // RUN: %{build} %link-vulkan -o %t.out %if target-spir %{ -Wno-ignored-attributes %} // RUN: %{run} %t.out --no-sem // RUN: %{run} %t.out 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 cc784bd8c269e..e296899d307a7 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 @@ -7,9 +7,6 @@ // UNSUPPORTED: linux // UNSUPPORTED-TRACKER: GSD-12357 -// XFAIL: windows && arch-intel_gpu_bmg_g21 -// XFAIL-TRACKER: https://github.com/intel/llvm/issues/21986 - /* Run ALL the vulkan formats through the gauntlet. sampled and unsampled. This entire test takes less than 30 seconds on a slow machine. MUCH faster 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 441f1670c310f..3af7df26e2b5a 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 @@ -3,7 +3,7 @@ // REQUIRES: aspect-ext_oneapi_external_memory_import || (windows && level_zero && aspect-ext_oneapi_bindless_images) // REQUIRES: vulkan -// UNSUPPORTED: windows +// maybe working on Windows. Test. // UNSUPPORTED-TRACKER: CMPLRLLVM-73525 // RUN: %{build} %link-vulkan -o %t.out %if target-spir %{ -Wno-ignored-attributes %} 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 1675e16f985c5..27fcbc5d6aa33 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 @@ -7,9 +7,6 @@ // UNSUPPORTED: linux // UNSUPPORTED-TRACKER: GSD-12357 -// UNSUPPORTED: windows && arch-intel_gpu_bmg_g21 -// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/21986 - /* Run all the vulkan formats through a write test. Note this is unsampled only, you can't "write" with the image sampler. From bed00b04bc29531d099b8fedad455386a3509437 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 19 May 2026 11:41:06 -0700 Subject: [PATCH 12/17] clang-format is hereby uninvited from all my future social gatherings and parties Signed-off-by: Chris Perkins --- .../vulkan_sycl_image_interop_write_1d_unsampled.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) 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 3af7df26e2b5a..41223a2bd34ad 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 @@ -3,7 +3,7 @@ // REQUIRES: aspect-ext_oneapi_external_memory_import || (windows && level_zero && aspect-ext_oneapi_bindless_images) // REQUIRES: vulkan -// maybe working on Windows. Test. +// maybe working on Windows. Test. // UNSUPPORTED-TRACKER: CMPLRLLVM-73525 // RUN: %{build} %link-vulkan -o %t.out %if target-spir %{ -Wno-ignored-attributes %} From 4a0b7f6319754df1cff608810594e92e3e792dff Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 19 May 2026 15:09:36 -0700 Subject: [PATCH 13/17] bump Signed-off-by: Chris Perkins --- .../vulkan_sycl_image_interop_read_2d.cpp | 10 +++++----- .../vulkan_sycl_image_interop_write_1d_unsampled.cpp | 2 +- 2 files changed, 6 insertions(+), 6 deletions(-) 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 e296899d307a7..164203d8eeb7c 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 @@ -21,6 +21,7 @@ will fail. This is being tracked as a separate issue. */ +// clang-format off // RUN: %{run} %t.out --type float --channels 1 32x33 // RUN: %{run} %t.out --type float --channels 2 32x33 @@ -73,9 +74,9 @@ // RUN: %{run} %t.out --type int8 --channels 1 --sampled 32x33 // RUN: %{run} %t.out --type int8 --channels 2 --sampled 32x33 // RUN: %{run} %t.out --type int8 --channels 4 --sampled 32x33 -// RUN: %{run} %t.out --type unorm8 --channels 1 --sampled 32x33 -// RUN: %{run} %t.out --type unorm8 --channels 2 --sampled 32x33 -// RUN: %{run} %t.out --type unorm8 --channels 4 --sampled 32x33 +// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.out --type unorm8 --channels 1 --sampled 32x33 +// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.out --type unorm8 --channels 2 --sampled 32x33 +// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.out --type unorm8 --channels 4 --sampled 32x33 // RUN: %{run} %t.out --type float --channels 1 32x33 --semaphores // RUN: %{run} %t.out --type float --channels 4 32x33 --semaphores @@ -87,9 +88,8 @@ // RUN: %{run} %t.out --type int32 --channels 4 --sampled 32x33 --semaphores // RUN: %{run} %t.out --type int16 --channels 4 --sampled 32x33 --semaphores // RUN: %{run} %t.out --type uint8 --channels 2 --sampled 32x33 --semaphores -// RUN: %{run} %t.out --type unorm8 --channels 4 --sampled 32x33 --semaphores +// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.out --type unorm8 --channels 4 --sampled 32x33 --semaphores -// clang-format off /* Vulkan/SYCL 2D Image Read Test (Sampled + Unsampled) 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 41223a2bd34ad..441f1670c310f 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 @@ -3,7 +3,7 @@ // REQUIRES: aspect-ext_oneapi_external_memory_import || (windows && level_zero && aspect-ext_oneapi_bindless_images) // REQUIRES: vulkan -// maybe working on Windows. Test. +// UNSUPPORTED: windows // UNSUPPORTED-TRACKER: CMPLRLLVM-73525 // RUN: %{build} %link-vulkan -o %t.out %if target-spir %{ -Wno-ignored-attributes %} From 212abd3e1dfcf763fee6e4353bc38c163fdc8628 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 19 May 2026 15:45:31 -0700 Subject: [PATCH 14/17] Eric's suggestion Signed-off-by: Chris Perkins --- .../D3D12_sycl_buffer_timeline_semaphore.cpp | 14 +++++++++----- .../D3D12_sycl_buffer_win32_name_native.cpp | 14 +++++++++----- .../dx12_interop/D3D12_sycl_interop_1D_read.cpp | 14 +++++++++----- .../D3D12_sycl_interop_1D_write_unsampled.cpp | 14 +++++++++----- .../D3D12_sycl_interop_2D_arithmetic.cpp | 14 +++++++++----- .../dx12_interop/D3D12_sycl_interop_2D_read.cpp | 14 +++++++++----- .../D3D12_sycl_interop_2D_write_unsampled.cpp | 14 +++++++++----- .../dx12_interop/D3D12_sycl_interop_3D_read.cpp | 14 +++++++++----- .../D3D12_sycl_interop_3D_write_unsampled.cpp | 14 +++++++++----- .../vulkan_interop/vulkan_sycl_2d_arithmetic.cpp | 14 +++++++++----- .../vulkan_interop/vulkan_sycl_buffer.cpp | 14 +++++++++----- .../vulkan_sycl_buffer_binary_semaphore.cpp | 14 +++++++++----- .../vulkan_sycl_buffer_timeline_semaphore.cpp | 14 +++++++++----- .../vulkan_sycl_image_interop_read_1d.cpp | 14 +++++++++----- .../vulkan_sycl_image_interop_read_2d.cpp | 14 +++++++++----- .../vulkan_sycl_image_interop_read_3d.cpp | 14 +++++++++----- ...ulkan_sycl_image_interop_write_1d_unsampled.cpp | 14 +++++++++----- ...ulkan_sycl_image_interop_write_2d_unsampled.cpp | 14 +++++++++----- ...ulkan_sycl_image_interop_write_3d_unsampled.cpp | 14 +++++++++----- 19 files changed, 171 insertions(+), 95 deletions(-) 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 fb1ff574d8c75..e697842d9a760 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 @@ -116,11 +116,15 @@ int main(int argc, char **argv) { // SYCL INTEROP try { - // External semaphore ops require an in-order queue backed by immediate - // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). - sycl::queue q{ - {sycl::property::queue::in_order{}, - sycl::ext::intel::property::queue::immediate_command_list{}}}; + // 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. + sycl::property_list qProps = + useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue:: + immediate_command_list{}} + : sycl::property_list{sycl::property::queue::in_order{}}; + sycl::queue q{qProps}; auto device = q.get_device(); auto context = q.get_context(); diff --git a/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_win32_name_native.cpp b/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_win32_name_native.cpp index 0278c50d0e2a3..4adfde483eb90 100644 --- a/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_win32_name_native.cpp +++ b/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_win32_name_native.cpp @@ -185,11 +185,15 @@ int main(int argc, char **argv) { // SYCL INTEROP - using resource_win32_name NATIVELY try { - // External semaphore ops require an in-order queue backed by immediate - // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). - sycl::queue q{ - {sycl::property::queue::in_order{}, - sycl::ext::intel::property::queue::immediate_command_list{}}}; + // 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. + sycl::property_list qProps = + useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue:: + immediate_command_list{}} + : sycl::property_list{sycl::property::queue::in_order{}}; + sycl::queue q{qProps}; auto device = q.get_device(); auto context = q.get_context(); 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 9cf9a9e090b20..bbb978f426fb4 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 @@ -187,11 +187,15 @@ int runTest( // SYCL Import and Verification try { - // External semaphore ops require an in-order queue backed by immediate - // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). - sycl::queue q{ - {sycl::property::queue::in_order{}, - sycl::ext::intel::property::queue::immediate_command_list{}}}; + // 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. + sycl::property_list qProps = + useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue:: + immediate_command_list{}} + : sycl::property_list{sycl::property::queue::in_order{}}; + sycl::queue q{qProps}; syclexp::external_mem_descriptor extMemDesc{ imgRes.sharedHandle, syclexp::external_mem_handle_type::win32_nt_handle, 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 dbb172a650288..569581ae071f4 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 @@ -148,11 +148,15 @@ int runTest( } try { - // External semaphore ops require an in-order queue backed by immediate - // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). - sycl::queue q{ - {sycl::property::queue::in_order{}, - sycl::ext::intel::property::queue::immediate_command_list{}}}; + // 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. + sycl::property_list qProps = + useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue:: + immediate_command_list{}} + : sycl::property_list{sycl::property::queue::in_order{}}; + sycl::queue q{qProps}; syclexp::external_mem_descriptor extMemDesc{ imgRes.sharedHandle, syclexp::external_mem_handle_type::win32_nt_handle, 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 3d6470ae47b44..55e6e5064b8d7 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 @@ -410,11 +410,15 @@ int runTest( signalExportableFence(ctx, extFenceB); try { - // External semaphore ops require an in-order queue backed by immediate - // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). - sycl::queue q{ - {sycl::property::queue::in_order{}, - sycl::ext::intel::property::queue::immediate_command_list{}}}; + // 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. + sycl::property_list qProps = + useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue:: + immediate_command_list{}} + : sycl::property_list{sycl::property::queue::in_order{}}; + sycl::queue q{qProps}; auto extMemA = syclexp::import_external_memory( syclexp::external_mem_descriptor{ 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 737eda30047d4..c1d69f8c9e59d 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 @@ -201,11 +201,15 @@ int runTest( // SYCL Import and Verification try { - // External semaphore ops require an in-order queue backed by immediate - // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). - sycl::queue q{ - {sycl::property::queue::in_order{}, - sycl::ext::intel::property::queue::immediate_command_list{}}}; + // 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. + sycl::property_list qProps = + useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue:: + immediate_command_list{}} + : sycl::property_list{sycl::property::queue::in_order{}}; + sycl::queue q{qProps}; syclexp::external_mem_descriptor extMemDesc{ imgRes.sharedHandle, syclexp::external_mem_handle_type::win32_nt_handle, 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 40afca9508b14..75f30256565ea 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 @@ -189,11 +189,15 @@ int runTest( } try { - // External semaphore ops require an in-order queue backed by immediate - // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). - sycl::queue q{ - {sycl::property::queue::in_order{}, - sycl::ext::intel::property::queue::immediate_command_list{}}}; + // 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. + sycl::property_list qProps = + useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue:: + immediate_command_list{}} + : sycl::property_list{sycl::property::queue::in_order{}}; + sycl::queue q{qProps}; syclexp::external_mem_descriptor extMemDesc{ imgRes.sharedHandle, syclexp::external_mem_handle_type::win32_nt_handle, 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 56483ca32d6b3..61b0b07732d18 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 @@ -204,11 +204,15 @@ int runTest( // SYCL Import and Verification try { - // External semaphore ops require an in-order queue backed by immediate - // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). - sycl::queue q{ - {sycl::property::queue::in_order{}, - sycl::ext::intel::property::queue::immediate_command_list{}}}; + // 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. + sycl::property_list qProps = + useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue:: + immediate_command_list{}} + : sycl::property_list{sycl::property::queue::in_order{}}; + sycl::queue q{qProps}; syclexp::external_mem_descriptor extMemDesc{ imgRes.sharedHandle, syclexp::external_mem_handle_type::win32_nt_handle, 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 bbf7207f2344f..f1fb08791ea45 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 @@ -151,11 +151,15 @@ int runTest( } try { - // External semaphore ops require an in-order queue backed by immediate - // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). - sycl::queue q{ - {sycl::property::queue::in_order{}, - sycl::ext::intel::property::queue::immediate_command_list{}}}; + // 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. + sycl::property_list qProps = + useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue:: + immediate_command_list{}} + : sycl::property_list{sycl::property::queue::in_order{}}; + sycl::queue q{qProps}; syclexp::external_mem_descriptor extMemDesc{ imgRes.sharedHandle, syclexp::external_mem_handle_type::win32_nt_handle, 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 f7fbbe4fc0eee..82c9c1b21f8bf 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 @@ -246,11 +246,15 @@ int runTest( }); try { - // External semaphore ops require an in-order queue backed by immediate - // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). - sycl::queue q{ - {sycl::property::queue::in_order{}, - sycl::ext::intel::property::queue::immediate_command_list{}}}; + // 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. + sycl::property_list qProps = + useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue:: + immediate_command_list{}} + : sycl::property_list{sycl::property::queue::in_order{}}; + sycl::queue q{qProps}; #ifdef _WIN32 auto extMemA = syclexp::import_external_memory( syclexp::external_mem_descriptor{ 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 903525509b9e8..68b1e14b55fb0 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 @@ -275,11 +275,15 @@ int main(int argc, char **argv) { // SYCL INTEROP try { - // External semaphore ops require an in-order queue backed by immediate - // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). - sycl::queue q{ - {sycl::property::queue::in_order{}, - sycl::ext::intel::property::queue::immediate_command_list{}}}; + // 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. + sycl::property_list qProps = + useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue:: + immediate_command_list{}} + : sycl::property_list{sycl::property::queue::in_order{}}; + sycl::queue q{qProps}; auto device = q.get_device(); auto context = q.get_context(); 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 67b0146ae5693..6514ff324ca2d 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 @@ -144,11 +144,15 @@ int main(int argc, char **argv) { // SYCL INTEROP try { - // External semaphore ops require an in-order queue backed by immediate - // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). - sycl::queue q{ - {sycl::property::queue::in_order{}, - sycl::ext::intel::property::queue::immediate_command_list{}}}; + // 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. + sycl::property_list qProps = + useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue:: + immediate_command_list{}} + : sycl::property_list{sycl::property::queue::in_order{}}; + sycl::queue q{qProps}; auto device = q.get_device(); auto context = q.get_context(); 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 e0662b192a716..4e55b18a2160f 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 @@ -135,11 +135,15 @@ int main(int argc, char **argv) { // SYCL INTEROP try { - // External semaphore ops require an in-order queue backed by immediate - // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). - sycl::queue q{ - {sycl::property::queue::in_order{}, - sycl::ext::intel::property::queue::immediate_command_list{}}}; + // 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. + sycl::property_list qProps = + useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue:: + immediate_command_list{}} + : sycl::property_list{sycl::property::queue::in_order{}}; + sycl::queue q{qProps}; auto device = q.get_device(); auto context = q.get_context(); 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 762b1847bd4b6..f3b497c881563 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 @@ -255,11 +255,15 @@ int runTest( // SYCL Import and Verification namespace syclexp = sycl::ext::oneapi::experimental; try { - // External semaphore ops require an in-order queue backed by immediate - // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). - sycl::queue q{ - {sycl::property::queue::in_order{}, - sycl::ext::intel::property::queue::immediate_command_list{}}}; + // 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. + sycl::property_list qProps = + useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue:: + immediate_command_list{}} + : sycl::property_list{sycl::property::queue::in_order{}}; + sycl::queue q{qProps}; // Import Memory (Platform Specific) #ifdef _WIN32 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 164203d8eeb7c..d4830b7198947 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 @@ -240,11 +240,15 @@ int runTest( // SYCL Import and Verification namespace syclexp = sycl::ext::oneapi::experimental; try { - // External semaphore ops require an in-order queue backed by immediate - // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). - sycl::queue q{ - {sycl::property::queue::in_order{}, - sycl::ext::intel::property::queue::immediate_command_list{}}}; + // 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. + sycl::property_list qProps = + useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue:: + immediate_command_list{}} + : sycl::property_list{sycl::property::queue::in_order{}}; + sycl::queue q{qProps}; // Import Memory (Platform Specific) #ifdef _WIN32 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 017e3a26a8b79..c92afa758ad34 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 @@ -139,11 +139,15 @@ int runTest( // SYCL Import and Verification namespace syclexp = sycl::ext::oneapi::experimental; try { - // External semaphore ops require an in-order queue backed by immediate - // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). - sycl::queue q{ - {sycl::property::queue::in_order{}, - sycl::ext::intel::property::queue::immediate_command_list{}}}; + // 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. + sycl::property_list qProps = + useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue:: + immediate_command_list{}} + : sycl::property_list{sycl::property::queue::in_order{}}; + sycl::queue q{qProps}; // Import Memory (Platform Specific) #ifdef _WIN32 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 441f1670c310f..b8660778d7f54 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 @@ -213,11 +213,15 @@ int runTest( } try { - // External semaphore ops require an in-order queue backed by immediate - // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). - sycl::queue q{ - {sycl::property::queue::in_order{}, - sycl::ext::intel::property::queue::immediate_command_list{}}}; + // 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. + sycl::property_list qProps = + useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue:: + immediate_command_list{}} + : sycl::property_list{sycl::property::queue::in_order{}}; + sycl::queue q{qProps}; // IMPORT MEMORY #ifdef _WIN32 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 27fcbc5d6aa33..a828d95c2df4d 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,11 +203,15 @@ int runTest( } try { - // External semaphore ops require an in-order queue backed by immediate - // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). - sycl::queue q{ - {sycl::property::queue::in_order{}, - sycl::ext::intel::property::queue::immediate_command_list{}}}; + // 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. + sycl::property_list qProps = + useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue:: + immediate_command_list{}} + : sycl::property_list{sycl::property::queue::in_order{}}; + sycl::queue q{qProps}; // IMPORT MEMORY #ifdef _WIN32 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 145eddb998043..0da9a5149c3ad 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 @@ -153,11 +153,15 @@ int runTest( } try { - // External semaphore ops require an in-order queue backed by immediate - // command lists (see sycl_ext_oneapi_bindless_images.asciidoc). - sycl::queue q{ - {sycl::property::queue::in_order{}, - sycl::ext::intel::property::queue::immediate_command_list{}}}; + // 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. + sycl::property_list qProps = + useSemaphores ? sycl::property_list{sycl::property::queue::in_order{}, + sycl::ext::intel::property::queue:: + immediate_command_list{}} + : sycl::property_list{sycl::property::queue::in_order{}}; + sycl::queue q{qProps}; // IMPORT MEMORY #ifdef _WIN32 From 3a0c3144dbada9a38846e5677059453ec4323a8d Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 19 May 2026 17:13:53 -0700 Subject: [PATCH 15/17] updating tests. Signed-off-by: Chris Perkins --- .../external_semaphore_regular_cl_fails.cpp | 2 +- .../external_semaphore_regular_cl_fails.cpp | 2 +- .../vulkan_sycl_image_interop_read_2d.cpp | 6 +- ..._sycl_image_interop_write_1d_unsampled.cpp | 168 ++++++++--------- ..._sycl_image_interop_write_2d_unsampled.cpp | 175 +++++++++--------- ..._sycl_image_interop_write_3d_unsampled.cpp | 170 ++++++++--------- 6 files changed, 266 insertions(+), 257 deletions(-) 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 index 31a3983d52b14..aa416c0b82f71 100644 --- 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 @@ -1,6 +1,6 @@ // -// REQUIRES: aspect-ext_oneapi_external_semaphore_import, windows +// 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 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 index d325d49defe5e..c7616be3060db 100644 --- 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 @@ -1,7 +1,7 @@ // // REQUIRES: aspect-ext_oneapi_external_semaphore_import -// REQUIRES: vulkan +// REQUIRES: vulkan && level_zero // // RUN: %{build} %link-vulkan -o %t.out %if target-spir %{ -Wno-ignored-attributes %} // RUN: %{run} %t.out 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 d4830b7198947..1c02ce0d683b5 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 @@ -47,9 +47,9 @@ // RUN: %{run} %t.out --type int8 --channels 1 32x33 // RUN: %{run} %t.out --type int8 --channels 2 32x33 // RUN: %{run} %t.out --type int8 --channels 4 32x33 -// RUN: %{run} %t.out --type unorm8 --channels 1 32x33 -// RUN: %{run} %t.out --type unorm8 --channels 2 32x33 -// RUN: %{run} %t.out --type unorm8 --channels 4 32x33 +// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.out --type unorm8 --channels 1 32x33 +// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.out --type unorm8 --channels 2 32x33 +// RUN-IF: (!gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.out --type unorm8 --channels 4 32x33 // RUN: %{run} %t.out --type float --channels 1 --sampled 32x33 // RUN: %{run} %t.out --type float --channels 2 --sampled 32x33 // RUN: %{run} %t.out --type float --channels 4 --sampled 32x33 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 b8660778d7f54..adba096a55b6c 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 @@ -322,97 +322,99 @@ int runTest( syclexp::destroy_image_handle(unsampledHandle, q.get_device(), q.get_context()); syclexp::release_external_memory(extMem, q.get_device(), q.get_context()); + + // Vulkan Verify + vkDeviceWaitIdle(vkCtx.device); + VkBuffer verifyBuffer; + VkDeviceMemory verifyMem; + size_t dataSize = width * channels * sizeof(T); + VkBufferCreateInfo bi = {VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO}; + bi.size = dataSize; + bi.usage = VK_BUFFER_USAGE_TRANSFER_DST_BIT; + vkCreateBuffer(vkCtx.device, &bi, nullptr, &verifyBuffer); + VkMemoryRequirements req; + vkGetBufferMemoryRequirements(vkCtx.device, verifyBuffer, &req); + VkMemoryAllocateInfo ai = {VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO}; + ai.allocationSize = req.size; + ai.memoryTypeIndex = + findMemoryType(vkCtx.physicalDevice, req.memoryTypeBits, + VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | + VK_MEMORY_PROPERTY_HOST_COHERENT_BIT); + vkAllocateMemory(vkCtx.device, &ai, nullptr, &verifyMem); + vkBindBufferMemory(vkCtx.device, verifyBuffer, verifyMem, 0); + + { + VkCommandPoolCreateInfo poolInfo = { + VK_STRUCTURE_TYPE_COMMAND_POOL_CREATE_INFO}; + poolInfo.queueFamilyIndex = vkCtx.queueFamilyIndex; + VkCommandPool pool; + vkCreateCommandPool(vkCtx.device, &poolInfo, nullptr, &pool); + VkCommandBuffer cmd; + VkCommandBufferAllocateInfo ca = { + VK_STRUCTURE_TYPE_COMMAND_BUFFER_ALLOCATE_INFO}; + ca.commandPool = pool; + ca.level = VK_COMMAND_BUFFER_LEVEL_PRIMARY; + ca.commandBufferCount = 1; + vkAllocateCommandBuffers(vkCtx.device, &ca, &cmd); + VkCommandBufferBeginInfo bi = { + VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO}; + vkBeginCommandBuffer(cmd, &bi); + VkBufferImageCopy reg = {}; + reg.imageSubresource = {VK_IMAGE_ASPECT_COLOR_BIT, 0, 0, 1}; + reg.imageExtent = extent; + vkCmdCopyImageToBuffer(cmd, imgRes.image, VK_IMAGE_LAYOUT_GENERAL, + verifyBuffer, 1, ®); + vkEndCommandBuffer(cmd); + VkSubmitInfo si = {VK_STRUCTURE_TYPE_SUBMIT_INFO}; + si.commandBufferCount = 1; + si.pCommandBuffers = &cmd; + std::vector waitStages = { + VK_PIPELINE_STAGE_TRANSFER_BIT}; + if (useSemaphores) { + si.waitSemaphoreCount = 1; + si.pWaitSemaphores = &vkSem; + si.pWaitDstStageMask = waitStages.data(); + } + vkQueueSubmit(vkCtx.queue, 1, &si, VK_NULL_HANDLE); + vkQueueWaitIdle(vkCtx.queue); + vkDestroyCommandPool(vkCtx.device, pool, nullptr); + } + + void *ptr; + vkMapMemory(vkCtx.device, verifyMem, 0, dataSize, 0, &ptr); + T *vData = (T *)ptr; + bool passed = true; + int errorCount = 0; + for (size_t i = 0; i < width * channels; ++i) { + T expected = generateTestValue(i / channels, i % channels, width); + if (!checkValue(vData[i], expected)) { + passed = false; + if (errorCount++ < 5) + std::cout << "Mismatch at " << i << " Got: " << (double)vData[i] + << " Exp: " << (double)expected << std::endl; + } + } + vkUnmapMemory(vkCtx.device, verifyMem); + if (passed) + std::cout << "SUCCESS!" << std::endl; + else + std::cout << "FAILURE! (" << errorCount << " errors)" << std::endl; + + vkDestroyBuffer(vkCtx.device, verifyBuffer, nullptr); + vkFreeMemory(vkCtx.device, verifyMem, nullptr); if (useSemaphores) { syclexp::release_external_semaphore(extSem, q.get_device(), q.get_context()); + vkDestroySemaphore(vkCtx.device, vkSem, nullptr); } + cleanupVulkan(vkCtx, imgRes); + return passed ? 0 : 1; + } catch (std::exception &e) { std::cerr << "SYCL Exception: " << e.what() << std::endl; + cleanupVulkan(vkCtx, imgRes); return 1; } - - // Vulkan Verify - vkDeviceWaitIdle(vkCtx.device); - VkBuffer verifyBuffer; - VkDeviceMemory verifyMem; - size_t dataSize = width * channels * sizeof(T); - VkBufferCreateInfo bi = {VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO}; - bi.size = dataSize; - bi.usage = VK_BUFFER_USAGE_TRANSFER_DST_BIT; - vkCreateBuffer(vkCtx.device, &bi, nullptr, &verifyBuffer); - VkMemoryRequirements req; - vkGetBufferMemoryRequirements(vkCtx.device, verifyBuffer, &req); - VkMemoryAllocateInfo ai = {VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO}; - ai.allocationSize = req.size; - ai.memoryTypeIndex = findMemoryType(vkCtx.physicalDevice, req.memoryTypeBits, - VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | - VK_MEMORY_PROPERTY_HOST_COHERENT_BIT); - vkAllocateMemory(vkCtx.device, &ai, nullptr, &verifyMem); - vkBindBufferMemory(vkCtx.device, verifyBuffer, verifyMem, 0); - - { - VkCommandPoolCreateInfo poolInfo = { - VK_STRUCTURE_TYPE_COMMAND_POOL_CREATE_INFO}; - poolInfo.queueFamilyIndex = vkCtx.queueFamilyIndex; - VkCommandPool pool; - vkCreateCommandPool(vkCtx.device, &poolInfo, nullptr, &pool); - VkCommandBuffer cmd; - VkCommandBufferAllocateInfo ca = { - VK_STRUCTURE_TYPE_COMMAND_BUFFER_ALLOCATE_INFO}; - ca.commandPool = pool; - ca.level = VK_COMMAND_BUFFER_LEVEL_PRIMARY; - ca.commandBufferCount = 1; - vkAllocateCommandBuffers(vkCtx.device, &ca, &cmd); - VkCommandBufferBeginInfo bi = {VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO}; - vkBeginCommandBuffer(cmd, &bi); - VkBufferImageCopy reg = {}; - reg.imageSubresource = {VK_IMAGE_ASPECT_COLOR_BIT, 0, 0, 1}; - reg.imageExtent = extent; - vkCmdCopyImageToBuffer(cmd, imgRes.image, VK_IMAGE_LAYOUT_GENERAL, - verifyBuffer, 1, ®); - vkEndCommandBuffer(cmd); - VkSubmitInfo si = {VK_STRUCTURE_TYPE_SUBMIT_INFO}; - si.commandBufferCount = 1; - si.pCommandBuffers = &cmd; - std::vector waitStages = { - VK_PIPELINE_STAGE_TRANSFER_BIT}; - if (useSemaphores) { - si.waitSemaphoreCount = 1; - si.pWaitSemaphores = &vkSem; - si.pWaitDstStageMask = waitStages.data(); - } - vkQueueSubmit(vkCtx.queue, 1, &si, VK_NULL_HANDLE); - vkQueueWaitIdle(vkCtx.queue); - vkDestroyCommandPool(vkCtx.device, pool, nullptr); - } - - void *ptr; - vkMapMemory(vkCtx.device, verifyMem, 0, dataSize, 0, &ptr); - T *vData = (T *)ptr; - bool passed = true; - int errorCount = 0; - for (size_t i = 0; i < width * channels; ++i) { - T expected = generateTestValue(i / channels, i % channels, width); - if (!checkValue(vData[i], expected)) { - passed = false; - if (errorCount++ < 5) - std::cout << "Mismatch at " << i << " Got: " << (double)vData[i] - << " Exp: " << (double)expected << std::endl; - } - } - vkUnmapMemory(vkCtx.device, verifyMem); - if (passed) - std::cout << "SUCCESS!" << std::endl; - else - std::cout << "FAILURE! (" << errorCount << " errors)" << std::endl; - - vkDestroyBuffer(vkCtx.device, verifyBuffer, nullptr); - vkFreeMemory(vkCtx.device, verifyMem, nullptr); - if (useSemaphores) - vkDestroySemaphore(vkCtx.device, vkSem, nullptr); - cleanupVulkan(vkCtx, imgRes); - - return passed ? 0 : 1; } int main(int argc, char **argv) { 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 a828d95c2df4d..14a8ec67c267a 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 @@ -327,100 +327,103 @@ int runTest( syclexp::destroy_image_handle(unsampledHandle, q.get_device(), q.get_context()); syclexp::release_external_memory(extMem, q.get_device(), q.get_context()); + + // Vulkan Verify + vkDeviceWaitIdle(vkCtx.device); + VkBuffer verifyBuffer; + VkDeviceMemory verifyMem; + size_t dataSize = width * height * channels * sizeof(T); + VkBufferCreateInfo bi = {VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO}; + bi.size = dataSize; + bi.usage = VK_BUFFER_USAGE_TRANSFER_DST_BIT; + vkCreateBuffer(vkCtx.device, &bi, nullptr, &verifyBuffer); + VkMemoryRequirements req; + vkGetBufferMemoryRequirements(vkCtx.device, verifyBuffer, &req); + VkMemoryAllocateInfo ai = {VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO}; + ai.allocationSize = req.size; + ai.memoryTypeIndex = + findMemoryType(vkCtx.physicalDevice, req.memoryTypeBits, + VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | + VK_MEMORY_PROPERTY_HOST_COHERENT_BIT); + vkAllocateMemory(vkCtx.device, &ai, nullptr, &verifyMem); + vkBindBufferMemory(vkCtx.device, verifyBuffer, verifyMem, 0); + + { + VkCommandPoolCreateInfo poolInfo = { + VK_STRUCTURE_TYPE_COMMAND_POOL_CREATE_INFO}; + poolInfo.queueFamilyIndex = vkCtx.queueFamilyIndex; + VkCommandPool pool; + vkCreateCommandPool(vkCtx.device, &poolInfo, nullptr, &pool); + VkCommandBuffer cmd; + VkCommandBufferAllocateInfo ca = { + VK_STRUCTURE_TYPE_COMMAND_BUFFER_ALLOCATE_INFO}; + ca.commandPool = pool; + ca.level = VK_COMMAND_BUFFER_LEVEL_PRIMARY; + ca.commandBufferCount = 1; + vkAllocateCommandBuffers(vkCtx.device, &ca, &cmd); + VkCommandBufferBeginInfo bi = { + VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO}; + vkBeginCommandBuffer(cmd, &bi); + VkBufferImageCopy reg = {}; + reg.imageSubresource = {VK_IMAGE_ASPECT_COLOR_BIT, 0, 0, 1}; + reg.imageExtent = extent; + vkCmdCopyImageToBuffer(cmd, imgRes.image, VK_IMAGE_LAYOUT_GENERAL, + verifyBuffer, 1, ®); + vkEndCommandBuffer(cmd); + VkSubmitInfo si = {VK_STRUCTURE_TYPE_SUBMIT_INFO}; + si.commandBufferCount = 1; + si.pCommandBuffers = &cmd; + std::vector waitStages = { + VK_PIPELINE_STAGE_TRANSFER_BIT}; + if (useSemaphores) { + si.waitSemaphoreCount = 1; + si.pWaitSemaphores = &vkSem; + si.pWaitDstStageMask = waitStages.data(); + } + vkQueueSubmit(vkCtx.queue, 1, &si, VK_NULL_HANDLE); + vkQueueWaitIdle(vkCtx.queue); + vkDestroyCommandPool(vkCtx.device, pool, nullptr); + } + + void *ptr; + vkMapMemory(vkCtx.device, verifyMem, 0, dataSize, 0, &ptr); + T *vData = (T *)ptr; + bool passed = true; + int errorCount = 0; + size_t totalPixels = width * height; + + for (size_t i = 0; i < totalPixels * channels; ++i) { + size_t pixelIdx = i / channels; + int ch = i % channels; + T expected = generateTestValue(pixelIdx, ch, totalPixels); + if (!checkValue(vData[i], expected)) { + passed = false; + if (errorCount++ < 5) + std::cout << "Mismatch at " << i << " Got: " << (double)vData[i] + << " Exp: " << (double)expected << std::endl; + } + } + vkUnmapMemory(vkCtx.device, verifyMem); + if (passed) + std::cout << "SUCCESS!" << std::endl; + else + std::cout << "FAILURE! (" << errorCount << " errors)" << std::endl; + + vkDestroyBuffer(vkCtx.device, verifyBuffer, nullptr); + vkFreeMemory(vkCtx.device, verifyMem, nullptr); if (useSemaphores) { syclexp::release_external_semaphore(extSem, q.get_device(), q.get_context()); + vkDestroySemaphore(vkCtx.device, vkSem, nullptr); } + cleanupVulkan(vkCtx, imgRes); + return passed ? 0 : 1; + } catch (std::exception &e) { std::cerr << "SYCL Exception: " << e.what() << std::endl; + cleanupVulkan(vkCtx, imgRes); return 1; } - - // Vulkan Verify - vkDeviceWaitIdle(vkCtx.device); - VkBuffer verifyBuffer; - VkDeviceMemory verifyMem; - size_t dataSize = width * height * channels * sizeof(T); - VkBufferCreateInfo bi = {VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO}; - bi.size = dataSize; - bi.usage = VK_BUFFER_USAGE_TRANSFER_DST_BIT; - vkCreateBuffer(vkCtx.device, &bi, nullptr, &verifyBuffer); - VkMemoryRequirements req; - vkGetBufferMemoryRequirements(vkCtx.device, verifyBuffer, &req); - VkMemoryAllocateInfo ai = {VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO}; - ai.allocationSize = req.size; - ai.memoryTypeIndex = findMemoryType(vkCtx.physicalDevice, req.memoryTypeBits, - VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | - VK_MEMORY_PROPERTY_HOST_COHERENT_BIT); - vkAllocateMemory(vkCtx.device, &ai, nullptr, &verifyMem); - vkBindBufferMemory(vkCtx.device, verifyBuffer, verifyMem, 0); - - { - VkCommandPoolCreateInfo poolInfo = { - VK_STRUCTURE_TYPE_COMMAND_POOL_CREATE_INFO}; - poolInfo.queueFamilyIndex = vkCtx.queueFamilyIndex; - VkCommandPool pool; - vkCreateCommandPool(vkCtx.device, &poolInfo, nullptr, &pool); - VkCommandBuffer cmd; - VkCommandBufferAllocateInfo ca = { - VK_STRUCTURE_TYPE_COMMAND_BUFFER_ALLOCATE_INFO}; - ca.commandPool = pool; - ca.level = VK_COMMAND_BUFFER_LEVEL_PRIMARY; - ca.commandBufferCount = 1; - vkAllocateCommandBuffers(vkCtx.device, &ca, &cmd); - VkCommandBufferBeginInfo bi = {VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO}; - vkBeginCommandBuffer(cmd, &bi); - VkBufferImageCopy reg = {}; - reg.imageSubresource = {VK_IMAGE_ASPECT_COLOR_BIT, 0, 0, 1}; - reg.imageExtent = extent; - vkCmdCopyImageToBuffer(cmd, imgRes.image, VK_IMAGE_LAYOUT_GENERAL, - verifyBuffer, 1, ®); - vkEndCommandBuffer(cmd); - VkSubmitInfo si = {VK_STRUCTURE_TYPE_SUBMIT_INFO}; - si.commandBufferCount = 1; - si.pCommandBuffers = &cmd; - std::vector waitStages = { - VK_PIPELINE_STAGE_TRANSFER_BIT}; - if (useSemaphores) { - si.waitSemaphoreCount = 1; - si.pWaitSemaphores = &vkSem; - si.pWaitDstStageMask = waitStages.data(); - } - vkQueueSubmit(vkCtx.queue, 1, &si, VK_NULL_HANDLE); - vkQueueWaitIdle(vkCtx.queue); - vkDestroyCommandPool(vkCtx.device, pool, nullptr); - } - - void *ptr; - vkMapMemory(vkCtx.device, verifyMem, 0, dataSize, 0, &ptr); - T *vData = (T *)ptr; - bool passed = true; - int errorCount = 0; - size_t totalPixels = width * height; - - for (size_t i = 0; i < totalPixels * channels; ++i) { - size_t pixelIdx = i / channels; - int ch = i % channels; - T expected = generateTestValue(pixelIdx, ch, totalPixels); - if (!checkValue(vData[i], expected)) { - passed = false; - if (errorCount++ < 5) - std::cout << "Mismatch at " << i << " Got: " << (double)vData[i] - << " Exp: " << (double)expected << std::endl; - } - } - vkUnmapMemory(vkCtx.device, verifyMem); - if (passed) - std::cout << "SUCCESS!" << std::endl; - else - std::cout << "FAILURE! (" << errorCount << " errors)" << std::endl; - - vkDestroyBuffer(vkCtx.device, verifyBuffer, nullptr); - vkFreeMemory(vkCtx.device, verifyMem, nullptr); - if (useSemaphores) - vkDestroySemaphore(vkCtx.device, vkSem, nullptr); - cleanupVulkan(vkCtx, imgRes); - return passed ? 0 : 1; } int main(int argc, char **argv) { 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 0da9a5149c3ad..17526a62bc7f3 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 @@ -276,97 +276,101 @@ int runTest( syclexp::destroy_image_handle(unsampledHandle, q.get_device(), q.get_context()); syclexp::release_external_memory(extMem, q.get_device(), q.get_context()); + + // Vulkan Verify + vkDeviceWaitIdle(vkCtx.device); + VkBuffer verifyBuffer; + VkDeviceMemory verifyMem; + size_t dataSize = width * height * depth * channels * sizeof(T); + VkBufferCreateInfo bi = {VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO}; + bi.size = dataSize; + bi.usage = VK_BUFFER_USAGE_TRANSFER_DST_BIT; + vkCreateBuffer(vkCtx.device, &bi, nullptr, &verifyBuffer); + VkMemoryRequirements req; + vkGetBufferMemoryRequirements(vkCtx.device, verifyBuffer, &req); + VkMemoryAllocateInfo ai = {VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO}; + ai.allocationSize = req.size; + ai.memoryTypeIndex = + findMemoryType(vkCtx.physicalDevice, req.memoryTypeBits, + VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | + VK_MEMORY_PROPERTY_HOST_COHERENT_BIT); + vkAllocateMemory(vkCtx.device, &ai, nullptr, &verifyMem); + vkBindBufferMemory(vkCtx.device, verifyBuffer, verifyMem, 0); + + { + VkCommandPoolCreateInfo poolInfo = { + VK_STRUCTURE_TYPE_COMMAND_POOL_CREATE_INFO}; + poolInfo.queueFamilyIndex = vkCtx.queueFamilyIndex; + VkCommandPool pool; + vkCreateCommandPool(vkCtx.device, &poolInfo, nullptr, &pool); + VkCommandBuffer cmd; + VkCommandBufferAllocateInfo ca = { + VK_STRUCTURE_TYPE_COMMAND_BUFFER_ALLOCATE_INFO}; + ca.commandPool = pool; + ca.level = VK_COMMAND_BUFFER_LEVEL_PRIMARY; + ca.commandBufferCount = 1; + vkAllocateCommandBuffers(vkCtx.device, &ca, &cmd); + VkCommandBufferBeginInfo bi = { + VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO}; + vkBeginCommandBuffer(cmd, &bi); + VkBufferImageCopy reg = {}; + reg.imageSubresource = {VK_IMAGE_ASPECT_COLOR_BIT, 0, 0, 1}; + reg.imageExtent = extent; + vkCmdCopyImageToBuffer(cmd, imgRes.image, VK_IMAGE_LAYOUT_GENERAL, + verifyBuffer, 1, ®); + vkEndCommandBuffer(cmd); + VkSubmitInfo si = {VK_STRUCTURE_TYPE_SUBMIT_INFO}; + si.commandBufferCount = 1; + si.pCommandBuffers = &cmd; + std::vector waitStages = { + VK_PIPELINE_STAGE_TRANSFER_BIT}; + if (useSemaphores) { + si.waitSemaphoreCount = 1; + si.pWaitSemaphores = &vkSem; + si.pWaitDstStageMask = waitStages.data(); + } + vkQueueSubmit(vkCtx.queue, 1, &si, VK_NULL_HANDLE); + vkQueueWaitIdle(vkCtx.queue); + vkDestroyCommandPool(vkCtx.device, pool, nullptr); + } + + void *ptr; + vkMapMemory(vkCtx.device, verifyMem, 0, dataSize, 0, &ptr); + T *vData = (T *)ptr; + bool passed = true; + int errorCount = 0; + size_t totalPixels = width * height * depth; + for (size_t i = 0; i < totalPixels * channels; ++i) { + T expected = + generateTestValue(i / channels, i % channels, totalPixels); + if (!checkValue(vData[i], expected)) { + passed = false; + if (errorCount++ < 5) + std::cout << "Mismatch at " << i << " Got: " << (double)vData[i] + << " Exp: " << (double)expected << std::endl; + } + } + vkUnmapMemory(vkCtx.device, verifyMem); + if (passed) + std::cout << "SUCCESS!" << std::endl; + else + std::cout << "FAILURE! (" << errorCount << " errors)" << std::endl; + + vkDestroyBuffer(vkCtx.device, verifyBuffer, nullptr); + vkFreeMemory(vkCtx.device, verifyMem, nullptr); if (useSemaphores) { syclexp::release_external_semaphore(extSem, q.get_device(), q.get_context()); + vkDestroySemaphore(vkCtx.device, vkSem, nullptr); } + cleanupVulkan(vkCtx, imgRes); + return passed ? 0 : 1; + } catch (std::exception &e) { std::cerr << "SYCL Exception: " << e.what() << std::endl; + cleanupVulkan(vkCtx, imgRes); return 1; } - - // Vulkan Verify - vkDeviceWaitIdle(vkCtx.device); - VkBuffer verifyBuffer; - VkDeviceMemory verifyMem; - size_t dataSize = width * height * depth * channels * sizeof(T); - VkBufferCreateInfo bi = {VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO}; - bi.size = dataSize; - bi.usage = VK_BUFFER_USAGE_TRANSFER_DST_BIT; - vkCreateBuffer(vkCtx.device, &bi, nullptr, &verifyBuffer); - VkMemoryRequirements req; - vkGetBufferMemoryRequirements(vkCtx.device, verifyBuffer, &req); - VkMemoryAllocateInfo ai = {VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO}; - ai.allocationSize = req.size; - ai.memoryTypeIndex = findMemoryType(vkCtx.physicalDevice, req.memoryTypeBits, - VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | - VK_MEMORY_PROPERTY_HOST_COHERENT_BIT); - vkAllocateMemory(vkCtx.device, &ai, nullptr, &verifyMem); - vkBindBufferMemory(vkCtx.device, verifyBuffer, verifyMem, 0); - - { - VkCommandPoolCreateInfo poolInfo = { - VK_STRUCTURE_TYPE_COMMAND_POOL_CREATE_INFO}; - poolInfo.queueFamilyIndex = vkCtx.queueFamilyIndex; - VkCommandPool pool; - vkCreateCommandPool(vkCtx.device, &poolInfo, nullptr, &pool); - VkCommandBuffer cmd; - VkCommandBufferAllocateInfo ca = { - VK_STRUCTURE_TYPE_COMMAND_BUFFER_ALLOCATE_INFO}; - ca.commandPool = pool; - ca.level = VK_COMMAND_BUFFER_LEVEL_PRIMARY; - ca.commandBufferCount = 1; - vkAllocateCommandBuffers(vkCtx.device, &ca, &cmd); - VkCommandBufferBeginInfo bi = {VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO}; - vkBeginCommandBuffer(cmd, &bi); - VkBufferImageCopy reg = {}; - reg.imageSubresource = {VK_IMAGE_ASPECT_COLOR_BIT, 0, 0, 1}; - reg.imageExtent = extent; - vkCmdCopyImageToBuffer(cmd, imgRes.image, VK_IMAGE_LAYOUT_GENERAL, - verifyBuffer, 1, ®); - vkEndCommandBuffer(cmd); - VkSubmitInfo si = {VK_STRUCTURE_TYPE_SUBMIT_INFO}; - si.commandBufferCount = 1; - si.pCommandBuffers = &cmd; - std::vector waitStages = { - VK_PIPELINE_STAGE_TRANSFER_BIT}; - if (useSemaphores) { - si.waitSemaphoreCount = 1; - si.pWaitSemaphores = &vkSem; - si.pWaitDstStageMask = waitStages.data(); - } - vkQueueSubmit(vkCtx.queue, 1, &si, VK_NULL_HANDLE); - vkQueueWaitIdle(vkCtx.queue); - vkDestroyCommandPool(vkCtx.device, pool, nullptr); - } - - void *ptr; - vkMapMemory(vkCtx.device, verifyMem, 0, dataSize, 0, &ptr); - T *vData = (T *)ptr; - bool passed = true; - int errorCount = 0; - size_t totalPixels = width * height * depth; - for (size_t i = 0; i < totalPixels * channels; ++i) { - T expected = generateTestValue(i / channels, i % channels, totalPixels); - if (!checkValue(vData[i], expected)) { - passed = false; - if (errorCount++ < 5) - std::cout << "Mismatch at " << i << " Got: " << (double)vData[i] - << " Exp: " << (double)expected << std::endl; - } - } - vkUnmapMemory(vkCtx.device, verifyMem); - if (passed) - std::cout << "SUCCESS!" << std::endl; - else - std::cout << "FAILURE! (" << errorCount << " errors)" << std::endl; - - vkDestroyBuffer(vkCtx.device, verifyBuffer, nullptr); - vkFreeMemory(vkCtx.device, verifyMem, nullptr); - if (useSemaphores) - vkDestroySemaphore(vkCtx.device, vkSem, nullptr); - cleanupVulkan(vkCtx, imgRes); - return passed ? 0 : 1; } int main(int argc, char **argv) { From 7a0e921fcf6ed6a5c607c5b5772db64b04031fda Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 20 May 2026 16:17:30 -0700 Subject: [PATCH 16/17] moar reenablement --- .../D3D12_sycl_buffer_win32_name_native.cpp | 2 - .../D3D12_sycl_interop_3D_read.cpp | 47 ++++++++----------- 2 files changed, 20 insertions(+), 29 deletions(-) diff --git a/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_win32_name_native.cpp b/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_win32_name_native.cpp index 4adfde483eb90..d517a872a0c99 100644 --- a/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_win32_name_native.cpp +++ b/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_win32_name_native.cpp @@ -9,8 +9,6 @@ // UNSUPPORTED: gpu-intel-gen12 // UNSUPPORTED-TRACKER: GSD-12427 -// UNSUPPORTED: arch-intel_gpu_bmg_g21 -// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/22028 // RUN: %{build} %link-directx -o %t.exe %if target-spir %{ -Wno-ignored-attributes %} // RUN: %{run} %t.exe --no-sem 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 61b0b07732d18..d6ea87d3c4ae5 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 @@ -2,13 +2,6 @@ // REQUIRES: aspect-ext_oneapi_external_memory_import // REQUIRES: windows -// UNSUPPORTED: arch-intel_gpu_bmg_g21 -// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/20384 -// also GSD-12430 - -// UNSUPPORTED: gpu-intel-dg2 -// UNSUPPORTED-TRACKER: GSD-12430 - // RUN: %{build} -o %t.exe %link-directx // RUN: %{run} %t.exe --type float --channels 4 8x8x8 @@ -64,24 +57,23 @@ // RUN: %{run} %t.exe --type int8 --channels 1 --sampled 17x16x15 // RUN: %{run} %t.exe --type int8 --channels 2 --sampled 16x15x17 // RUN: %{run} %t.exe --type int8 --channels 4 --sampled 15x17x16 -// RUN-IF: !gpu-intel-gen12, %{run} %t.exe --type unorm8 --channels 1 --sampled 9x8x7 -// RUN-IF: !gpu-intel-gen12, %{run} %t.exe --type unorm8 --channels 2 --sampled 8x7x9 -// RUN-IF: !gpu-intel-gen12, %{run} %t.exe --type unorm8 --channels 4 --sampled 7x9x8 +// RUN-IF: (!gpu-intel-gen12 && !gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type unorm8 --channels 1 --sampled 9x8x7 +// RUN-IF: (!gpu-intel-gen12 && !gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type unorm8 --channels 2 --sampled 8x7x9 +// RUN-IF: (!gpu-intel-gen12 && !gpu-intel-dg2 && !arch-intel_gpu_bmg_g21), %{run} %t.exe --type unorm8 --channels 4 --sampled 7x9x8 // Semaphore coverage tests -// At this time, semaphores can hang on BMG if run in parallel (GSD-12436). -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 4 --semaphores 16x17x15 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 2 --semaphores 17x16x15 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int32 --channels 1 --semaphores 9x8x7 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 4 --semaphores 33x31x32 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int16 --channels 2 --semaphores 15x17x16 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint16 --channels 1 --semaphores 9x7x8 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint8 --channels 4 --semaphores 32x31x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int8 --channels 2 --semaphores 16x15x17 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type float --channels 4 --sampled --semaphores 31x32x33 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type half --channels 2 --sampled --semaphores 15x16x17 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type int32 --channels 1 --sampled --semaphores 7x8x9 -// RUN-IF: (!arch-intel_gpu_bmg_g21), %{run} %t.exe --type uint32 --channels 4 --sampled --semaphores 32x31x33 +// RUN: %{run} %t.exe --type float --channels 4 --semaphores 16x17x15 +// RUN: %{run} %t.exe --type half --channels 2 --semaphores 17x16x15 +// RUN: %{run} %t.exe --type int32 --channels 1 --semaphores 9x8x7 +// RUN: %{run} %t.exe --type uint32 --channels 4 --semaphores 33x31x32 +// RUN: %{run} %t.exe --type int16 --channels 2 --semaphores 15x17x16 +// RUN: %{run} %t.exe --type uint16 --channels 1 --semaphores 9x7x8 +// RUN: %{run} %t.exe --type uint8 --channels 4 --semaphores 32x31x33 +// RUN: %{run} %t.exe --type int8 --channels 2 --semaphores 16x15x17 +// RUN: %{run} %t.exe --type float --channels 4 --sampled --semaphores 31x32x33 +// RUN: %{run} %t.exe --type half --channels 2 --sampled --semaphores 15x16x17 +// RUN: %{run} %t.exe --type int32 --channels 1 --sampled --semaphores 7x8x9 +// RUN: %{run} %t.exe --type uint32 --channels 4 --sampled --semaphores 32x31x33 /* clang++.exe -fsycl -o ds3r.exe D3D12_sycl_interop_3D_read.cpp -ld3d12 -ldxgi -ld3dcompiler @@ -96,12 +88,13 @@ WxHxD Set custom Width x Height x Depth (e.g. 16x15x14) - BMG: - + BMG: + - WORKS, including --sampled and semaphores + - unorm8 does NOT work DG2: - - WORKS, including --sampled - - semaphores segfault + - WORKS, including --sampled and semaphores + - unorm8 does NOT work. GEN12: - WORKS, including --sampled and semaphores From fdfc513c2b51a081259df21b5c2bebb3a058f4cc Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 20 May 2026 16:32:09 -0700 Subject: [PATCH 17/17] clang-for-nothing Signed-off-by: Chris Perkins --- .../dx12_interop/D3D12_sycl_buffer_win32_name_native.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_win32_name_native.cpp b/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_win32_name_native.cpp index d517a872a0c99..3d0eba09f493a 100644 --- a/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_win32_name_native.cpp +++ b/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_win32_name_native.cpp @@ -9,7 +9,6 @@ // UNSUPPORTED: gpu-intel-gen12 // UNSUPPORTED-TRACKER: GSD-12427 - // RUN: %{build} %link-directx -o %t.exe %if target-spir %{ -Wno-ignored-attributes %} // RUN: %{run} %t.exe --no-sem // RUN: %{run} %t.exe