Skip to content

[WIP][DO-NOT-MERGE] explore bindless image #22046

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Draft
cperkinsintel wants to merge 16 commits into
base: sycl
Choose a base branch
from
Draft

[WIP][DO-NOT-MERGE] explore bindless image #22046

Show file tree
Hide file tree
Changes from 9 commits
Commits
Show all changes
16 commits
Select commit Hold shift + click to select a range
f9743a9
bindless image semaphore usage requires queues that use immediate com…
cperkinsintel May 13, 2026
ed6c1d3
narrower include
cperkinsintel May 15, 2026
a58cfbc
revise coupla tests
cperkinsintel May 15, 2026
d49c119
validation optional
cperkinsintel May 15, 2026
eb3f2f1
spelling correction
cperkinsintel May 15, 2026
3f0ce1e
DG2 testing going well. Reenabling lots of tests
cperkinsintel May 15, 2026
16f8fa2
exploratory
cperkinsintel May 18, 2026
77ad8f3
clang-format never fails
cperkinsintel May 18, 2026
cc223b3
clang-format is decidedly hostile.
cperkinsintel May 18, 2026
57c5d75
BMG enablement for DX12 tests
cperkinsintel May 18, 2026
ced6462
enable some moar vulkan tests
cperkinsintel May 19, 2026
bed00b0
clang-format is hereby uninvited from all my future social gatherings…
cperkinsintel May 19, 2026
4a0b7f6
bump
cperkinsintel May 19, 2026
88ba550
merge
cperkinsintel May 19, 2026
212abd3
Eric's suggestion
cperkinsintel May 19, 2026
3a0c314
updating tests.
cperkinsintel May 20, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
19 changes: 12 additions & 7 deletions sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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.
Expand Down
11 changes: 10 additions & 1 deletion sycl/test-e2e/bindless_images/dx11_interop/read_write_unsampled.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17,6 +17,7 @@
#include "dx11_interop.h"

#include <sycl/ext/oneapi/bindless_images.hpp>
#include <sycl/properties/queue_properties.hpp>

#ifdef TEST_SEMAPHORE_IMPORT
#include <d3d11_4.h> // Used for ID3D11Device5 / ID3D11DeviceContext4 / ID3D11Fence
Expand Down Expand Up @@ -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
Expand Down
7 changes: 6 additions & 1 deletion sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_timeline_semaphore.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -50,6 +50,7 @@
#include <string>
#include <sycl/detail/core.hpp>
#include <sycl/ext/oneapi/bindless_images.hpp>
#include <sycl/properties/queue_properties.hpp>
#include <vector>

#define WIN32_LEAN_AND_MEAN
Expand Down Expand Up @@ -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{}}};
Copy link
Copy Markdown
Contributor

@enm-intel enm-intel May 18, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Since these properties are only needed for external semaphores, perhaps you could condition their inclusion on the useSemaphores variable:

    sycl::property_list props{
        useSemaphores
            ? {sycl::property::queue::in_order{},
               sycl::ext::intel::property::queue::immediate_command_list{}}
            : sycl::property_list{}};
    sycl::queue q{props};

auto device = q.get_device();
auto context = q.get_context();

Expand Down
8 changes: 7 additions & 1 deletion sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_win32_name_native.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,6 +4,7 @@

// UNSUPPORTED: gpu-intel-dg2
// UNSUPPORTED-TRACKER: GSD-12428
// semaphores-do-not-work-dg2

// UNSUPPORTED: gpu-intel-gen12
// UNSUPPORTED-TRACKER: GSD-12427
Expand Down Expand Up @@ -32,6 +33,7 @@
#include <string>
#include <sycl/detail/core.hpp>
#include <sycl/ext/oneapi/bindless_images.hpp>
#include <sycl/properties/queue_properties.hpp>
#include <vector>

#define WIN32_LEAN_AND_MEAN
Expand Down Expand Up @@ -180,7 +182,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();

Expand Down
33 changes: 19 additions & 14 deletions sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_1D_read.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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

Expand All @@ -98,6 +98,7 @@
#include <sycl/detail/core.hpp>
#include <sycl/ext/oneapi/bindless_images.hpp>
#include <sycl/ext/oneapi/bindless_images_interop.hpp>
#include <sycl/properties/queue_properties.hpp>

namespace syclexp = sycl::ext::oneapi::experimental;

Expand Down Expand Up @@ -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<syclexp::resource_win32_handle> extMemDesc{
imgRes.sharedHandle, syclexp::external_mem_handle_type::win32_nt_handle,
Expand Down
35 changes: 20 additions & 15 deletions sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_1D_write_unsampled.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -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

Expand All @@ -71,6 +71,7 @@
#include <sycl/detail/core.hpp>
#include <sycl/ext/oneapi/bindless_images.hpp>
#include <sycl/ext/oneapi/bindless_images_interop.hpp>
#include <sycl/properties/queue_properties.hpp>

namespace syclexp = sycl::ext::oneapi::experimental;

Expand Down Expand Up @@ -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<syclexp::resource_win32_handle> extMemDesc{
imgRes.sharedHandle, syclexp::external_mem_handle_type::win32_nt_handle,
Expand Down
34 changes: 19 additions & 15 deletions sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_2D_arithmetic.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -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

Expand All @@ -118,6 +117,7 @@
#include <sycl/detail/core.hpp>
#include <sycl/ext/oneapi/bindless_images.hpp>
#include <sycl/ext/oneapi/bindless_images_interop.hpp>
#include <sycl/properties/queue_properties.hpp>
#include <vector>

namespace syclexp = sycl::ext::oneapi::experimental;
Expand Down Expand Up @@ -415,7 +415,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<syclexp::resource_win32_handle>{
Expand Down
34 changes: 19 additions & 15 deletions sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_2D_read.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -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

Expand All @@ -122,6 +121,7 @@
#include <sycl/detail/core.hpp>
#include <sycl/ext/oneapi/bindless_images.hpp>
#include <sycl/ext/oneapi/bindless_images_interop.hpp>
#include <sycl/properties/queue_properties.hpp>

namespace syclexp = sycl::ext::oneapi::experimental;

Expand Down Expand Up @@ -208,7 +208,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<syclexp::resource_win32_handle> extMemDesc{
imgRes.sharedHandle, syclexp::external_mem_handle_type::win32_nt_handle,
Expand Down
Loading
Loading