From 885d46643c87a82e3716112ef5d2e0f0979cfd95 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 28 May 2026 15:54:47 -0700 Subject: [PATCH 1/2] Revert "[SYCL] add support for bindless image resource_win32_name (named handles) (#21899)" This reverts commit e985fcba62f322b0f77c536dc1420369c15f0fa8. --- .../ext/oneapi/bindless_images_interop.hpp | 2 - sycl/source/detail/bindless_images.cpp | 213 +----------------- sycl/test/abi/sycl_symbols_windows.dump | 4 - 3 files changed, 5 insertions(+), 214 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/bindless_images_interop.hpp b/sycl/include/sycl/ext/oneapi/bindless_images_interop.hpp index 6bdd154ed9a8f..4a5891ba171b0 100644 --- a/sycl/include/sycl/ext/oneapi/bindless_images_interop.hpp +++ b/sycl/include/sycl/ext/oneapi/bindless_images_interop.hpp @@ -53,8 +53,6 @@ struct resource_win32_handle { // Windows external name type struct resource_win32_name { const void *name; - void *device; // ID3D12Device* (ID3D12Device1 or higher) for opening the - // named handle via OpenSharedHandleByName }; /// Opaque external memory descriptor type diff --git a/sycl/source/detail/bindless_images.cpp b/sycl/source/detail/bindless_images.cpp index bf5989db54b53..f5ffd0c4952fd 100644 --- a/sycl/source/detail/bindless_images.cpp +++ b/sycl/source/detail/bindless_images.cpp @@ -17,24 +17,6 @@ #include -#if defined(_WIN32) || defined(_WIN64) -#include -#include -#define WIN32_LEAN_AND_MEAN -#include - -namespace { -// Track handles opened via resource_win32_name so we can close them on release. -// Declared here so release_external_memory/semaphore can access them. Separate -// maps per UR handle type (memory vs. semaphore) — the two handle types are -// unrelated and must not be confused via reinterpret_cast. -std::mutex g_win32NameHandlesMutex; -std::unordered_map g_win32NameMemHandles; -std::unordered_map - g_win32NameSemHandles; -} // namespace -#endif - namespace sycl { inline namespace _V1 { namespace ext::oneapi::experimental { @@ -514,9 +496,11 @@ __SYCL_EXPORT external_mem import_external_memory( externalMemDesc, syclQueue.get_device(), syclQueue.get_context()); } -__SYCL_EXPORT image_mem_handle map_external_image_memory( - external_mem extMem, const image_descriptor &desc, - const sycl::device &syclDevice, const sycl::context &syclContext) { +__SYCL_EXPORT +image_mem_handle map_external_image_memory(external_mem extMem, + const image_descriptor &desc, + const sycl::device &syclDevice, + const sycl::context &syclContext) { desc.verify(); auto [urDevice, urCtx, Adapter] = get_ur_handles(syclDevice, syclContext); @@ -571,18 +555,6 @@ void *map_external_linear_memory(external_mem extMem, uint64_t offset, __SYCL_EXPORT void release_external_memory(external_mem extMem, const sycl::device &syclDevice, const sycl::context &syclContext) { -#if defined(_WIN32) || defined(_WIN64) - // Close handle if it was opened via resource_win32_name - { - std::lock_guard lock(g_win32NameHandlesMutex); - auto it = g_win32NameMemHandles.find(extMem.raw_handle); - if (it != g_win32NameMemHandles.end()) { - CloseHandle(it->second); - g_win32NameMemHandles.erase(it); - } - } -#endif - auto [urDevice, urCtx, Adapter] = get_ur_handles(syclDevice, syclContext); Adapter @@ -732,18 +704,6 @@ __SYCL_EXPORT void release_external_semaphore(external_semaphore externalSemaphore, const sycl::device &syclDevice, const sycl::context &syclContext) { -#if defined(_WIN32) || defined(_WIN64) - // Close handle if it was opened via resource_win32_name - { - std::lock_guard lock(g_win32NameHandlesMutex); - auto it = g_win32NameSemHandles.find(externalSemaphore.raw_handle); - if (it != g_win32NameSemHandles.end()) { - CloseHandle(it->second); - g_win32NameSemHandles.erase(it); - } - } -#endif - auto [urDevice, urCtx, Adapter] = get_ur_handles(syclDevice, syclContext); Adapter->call< @@ -1008,169 +968,6 @@ __SYCL_EXPORT bool is_image_handle_supported( syclQueue.get_context()); } -// ============================================================================ -// resource_win32_name support - Windows-specific code at end of file -// ============================================================================ -#if defined(_WIN32) || defined(_WIN64) - -// Include D3D12 for OpenSharedHandleByName -#ifdef __has_include -#if __has_include() -#include -#define SYCL_HAS_D3D12_INTEROP 1 -#endif -#endif - -namespace { - -// Closes a HANDLE on scope exit unless released. Used so that if the UR import -// call throws after we've opened a named handle, we don't leak it. -struct NamedHandleGuard { - HANDLE h; - explicit NamedHandleGuard(HANDLE handle) : h(handle) {} - ~NamedHandleGuard() { - if (h) - CloseHandle(h); - } - HANDLE release() { - HANDLE out = h; - h = nullptr; - return out; - } - NamedHandleGuard(const NamedHandleGuard &) = delete; - NamedHandleGuard &operator=(const NamedHandleGuard &) = delete; -}; - -HANDLE openNamedHandleImpl(void *device, const void *name) { -#ifdef SYCL_HAS_D3D12_INTEROP - // OpenSharedHandleByName requires ID3D12Device1 or higher - auto d3dDeviceBase = static_cast(device); - ID3D12Device1 *d3dDevice1 = nullptr; - - HRESULT hr = d3dDeviceBase->QueryInterface(IID_PPV_ARGS(&d3dDevice1)); - if (FAILED(hr)) { - return nullptr; - } - - HANDLE openedHandle = nullptr; - const wchar_t *wname = static_cast(name); - - hr = d3dDevice1->OpenSharedHandleByName(wname, GENERIC_ALL, &openedHandle); - d3dDevice1->Release(); - - return SUCCEEDED(hr) ? openedHandle : nullptr; -#else - (void)device; - (void)name; - throw sycl::exception( - sycl::make_error_code(sycl::errc::feature_not_supported), - "resource_win32_name requires D3D12 headers. " - "Use resource_win32_handle instead."); -#endif -} -} // anonymous namespace - -template <> -__SYCL_EXPORT external_mem import_external_memory( - external_mem_descriptor externalMemDesc, - const sycl::device &syclDevice, const sycl::context &syclContext) { - - if (!externalMemDesc.external_resource.device) { - throw sycl::exception(sycl::make_error_code(sycl::errc::invalid), - "D3D12 device required in resource_win32_name"); - } - - HANDLE openedHandle = - openNamedHandleImpl(externalMemDesc.external_resource.device, - externalMemDesc.external_resource.name); - - if (!openedHandle) { - throw sycl::exception(sycl::make_error_code(sycl::errc::invalid), - "Failed to open named Win32 handle"); - } - - // Close openedHandle if the delegated UR import throws before we've - // handed ownership to the tracking map. - NamedHandleGuard guard(openedHandle); - - // Use existing resource_win32_handle implementation - external_mem_descriptor handleDesc{ - {openedHandle}, - externalMemDesc.handle_type, - externalMemDesc.size_in_bytes}; - - // Delegate to existing resource_win32_handle implementation - external_mem result = import_external_memory( - handleDesc, syclDevice, syclContext); - - // Track the opened handle so we can close it on release - { - std::lock_guard lock(g_win32NameHandlesMutex); - g_win32NameMemHandles[result.raw_handle] = openedHandle; - } - guard.release(); - - return result; -} - -template <> -__SYCL_EXPORT external_mem import_external_memory( - external_mem_descriptor externalMemDesc, - const sycl::queue &syclQueue) { - return import_external_memory( - externalMemDesc, syclQueue.get_device(), syclQueue.get_context()); -} - -template <> -__SYCL_EXPORT external_semaphore import_external_semaphore( - external_semaphore_descriptor externalSemaphoreDesc, - const sycl::device &syclDevice, const sycl::context &syclContext) { - - if (!externalSemaphoreDesc.external_resource.device) { - throw sycl::exception(sycl::make_error_code(sycl::errc::invalid), - "D3D12 device required in resource_win32_name"); - } - - HANDLE openedHandle = - openNamedHandleImpl(externalSemaphoreDesc.external_resource.device, - externalSemaphoreDesc.external_resource.name); - - if (!openedHandle) { - throw sycl::exception(sycl::make_error_code(sycl::errc::invalid), - "Failed to open named Win32 semaphore"); - } - - // Close openedHandle if the delegated UR import throws before we've - // handed ownership to the tracking map. - NamedHandleGuard guard(openedHandle); - - // Delegate to existing resource_win32_handle implementation - external_semaphore_descriptor handleDesc{ - {openedHandle}, externalSemaphoreDesc.handle_type}; - - external_semaphore result = - import_external_semaphore(handleDesc, syclDevice, syclContext); - - // Track the opened handle so we can close it on release - { - std::lock_guard lock(g_win32NameHandlesMutex); - g_win32NameSemHandles[result.raw_handle] = openedHandle; - } - guard.release(); - - return result; -} - -template <> -__SYCL_EXPORT external_semaphore import_external_semaphore( - external_semaphore_descriptor externalSemaphoreDesc, - const sycl::queue &syclQueue) { - return import_external_semaphore( - externalSemaphoreDesc, syclQueue.get_device(), syclQueue.get_context()); -} - -#endif // _WIN32 || _WIN64 - } // namespace ext::oneapi::experimental } // namespace _V1 } // namespace sycl diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index f07fa3e2e50b1..1fccbf6808ab1 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -233,14 +233,10 @@ ??$import_external_memory@Uresource_fd@experimental@oneapi@ext@_V1@sycl@@@experimental@oneapi@ext@_V1@sycl@@YA?AUexternal_mem@01234@U?$external_mem_descriptor@Uresource_fd@experimental@oneapi@ext@_V1@sycl@@@01234@AEBVqueue@34@@Z ??$import_external_memory@Uresource_win32_handle@experimental@oneapi@ext@_V1@sycl@@@experimental@oneapi@ext@_V1@sycl@@YA?AUexternal_mem@01234@U?$external_mem_descriptor@Uresource_win32_handle@experimental@oneapi@ext@_V1@sycl@@@01234@AEBVdevice@34@AEBVcontext@34@@Z ??$import_external_memory@Uresource_win32_handle@experimental@oneapi@ext@_V1@sycl@@@experimental@oneapi@ext@_V1@sycl@@YA?AUexternal_mem@01234@U?$external_mem_descriptor@Uresource_win32_handle@experimental@oneapi@ext@_V1@sycl@@@01234@AEBVqueue@34@@Z -??$import_external_memory@Uresource_win32_name@experimental@oneapi@ext@_V1@sycl@@@experimental@oneapi@ext@_V1@sycl@@YA?AUexternal_mem@01234@U?$external_mem_descriptor@Uresource_win32_name@experimental@oneapi@ext@_V1@sycl@@@01234@AEBVdevice@34@AEBVcontext@34@@Z -??$import_external_memory@Uresource_win32_name@experimental@oneapi@ext@_V1@sycl@@@experimental@oneapi@ext@_V1@sycl@@YA?AUexternal_mem@01234@U?$external_mem_descriptor@Uresource_win32_name@experimental@oneapi@ext@_V1@sycl@@@01234@AEBVqueue@34@@Z ??$import_external_semaphore@Uresource_fd@experimental@oneapi@ext@_V1@sycl@@@experimental@oneapi@ext@_V1@sycl@@YA?AUexternal_semaphore@01234@U?$external_semaphore_descriptor@Uresource_fd@experimental@oneapi@ext@_V1@sycl@@@01234@AEBVdevice@34@AEBVcontext@34@@Z ??$import_external_semaphore@Uresource_fd@experimental@oneapi@ext@_V1@sycl@@@experimental@oneapi@ext@_V1@sycl@@YA?AUexternal_semaphore@01234@U?$external_semaphore_descriptor@Uresource_fd@experimental@oneapi@ext@_V1@sycl@@@01234@AEBVqueue@34@@Z ??$import_external_semaphore@Uresource_win32_handle@experimental@oneapi@ext@_V1@sycl@@@experimental@oneapi@ext@_V1@sycl@@YA?AUexternal_semaphore@01234@U?$external_semaphore_descriptor@Uresource_win32_handle@experimental@oneapi@ext@_V1@sycl@@@01234@AEBVdevice@34@AEBVcontext@34@@Z ??$import_external_semaphore@Uresource_win32_handle@experimental@oneapi@ext@_V1@sycl@@@experimental@oneapi@ext@_V1@sycl@@YA?AUexternal_semaphore@01234@U?$external_semaphore_descriptor@Uresource_win32_handle@experimental@oneapi@ext@_V1@sycl@@@01234@AEBVqueue@34@@Z -??$import_external_semaphore@Uresource_win32_name@experimental@oneapi@ext@_V1@sycl@@@experimental@oneapi@ext@_V1@sycl@@YA?AUexternal_semaphore@01234@U?$external_semaphore_descriptor@Uresource_win32_name@experimental@oneapi@ext@_V1@sycl@@@01234@AEBVdevice@34@AEBVcontext@34@@Z -??$import_external_semaphore@Uresource_win32_name@experimental@oneapi@ext@_V1@sycl@@@experimental@oneapi@ext@_V1@sycl@@YA?AUexternal_semaphore@01234@U?$external_semaphore_descriptor@Uresource_win32_name@experimental@oneapi@ext@_V1@sycl@@@01234@AEBVqueue@34@@Z ??$is_image_handle_supported@Usampled_image_handle@experimental@oneapi@ext@_V1@sycl@@@experimental@oneapi@ext@_V1@sycl@@YA_NAEBUimage_descriptor@01234@W4image_memory_handle_type@01234@AEBVdevice@34@AEBVcontext@34@@Z ??$is_image_handle_supported@Usampled_image_handle@experimental@oneapi@ext@_V1@sycl@@@experimental@oneapi@ext@_V1@sycl@@YA_NAEBUimage_descriptor@01234@W4image_memory_handle_type@01234@AEBVqueue@34@@Z ??$is_image_handle_supported@Uunsampled_image_handle@experimental@oneapi@ext@_V1@sycl@@@experimental@oneapi@ext@_V1@sycl@@YA_NAEBUimage_descriptor@01234@W4image_memory_handle_type@01234@AEBVdevice@34@AEBVcontext@34@@Z From 4c562848baab9d63fcd220aca830ef608187b00e Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Fri, 29 May 2026 10:42:17 -0700 Subject: [PATCH 2/2] remove test AGAIN. --- .../D3D12_sycl_buffer_win32_name_native.cpp | 398 ------------------ 1 file changed, 398 deletions(-) delete mode 100644 sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_win32_name_native.cpp 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 deleted file mode 100644 index 9a8b020e6ea6a..0000000000000 --- a/sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_buffer_win32_name_native.cpp +++ /dev/null @@ -1,398 +0,0 @@ -// REQUIRES: aspect-ext_oneapi_external_memory_import -// 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 - -// 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 -// RUN: %{run} %t.exe - -// clang-format off -/* - DirectX 12 / SYCL Buffer + Fence (Timeline) Interop Test - resource_win32_name - - clang++.exe -fsycl -o dsbwnn.exe D3D12_sycl_buffer_win32_name_native.cpp -ld3d12 -ldxgi -ld3dcompiler - - Tests native resource_win32_name support in SYCL. - SYCL runtime internally converts named handles to regular handles. - - FLAGS: --no-sem Don't use semaphores for SYCL/D3D12 synchronization - --iterations N Number of iterations to run (default: 10) - --size M Number of uint32_t elements in the buffer (default: 1024) -*/ -// clang-format on - -#include "d3d12_setup.hpp" -#include -#include -#include -#include -#include - -#define WIN32_LEAN_AND_MEAN -#include - -namespace syclexp = sycl::ext::oneapi::experimental; - -// Named resource structure -struct D3D12NamedBuffer { - Microsoft::WRL::ComPtr resource; - std::wstring name; - HANDLE keepAliveHandle; // Must keep at least one handle open for the name to - // persist -}; - -struct D3D12NamedFence { - Microsoft::WRL::ComPtr fence; - std::wstring name; - HANDLE keepAliveHandle; -}; - -// Create a buffer with a named shared handle -D3D12NamedBuffer createNamedExportableBuffer(D3D12Context &ctx, size_t size, - const wchar_t *name) { - D3D12NamedBuffer result; - result.name = name; - - D3D12_HEAP_PROPERTIES heapProps = {}; - heapProps.Type = D3D12_HEAP_TYPE_DEFAULT; - - D3D12_RESOURCE_DESC desc = {}; - desc.Dimension = D3D12_RESOURCE_DIMENSION_BUFFER; - desc.Width = size; - desc.Height = 1; - desc.DepthOrArraySize = 1; - desc.MipLevels = 1; - desc.SampleDesc.Count = 1; - desc.Layout = D3D12_TEXTURE_LAYOUT_ROW_MAJOR; - desc.Flags = D3D12_RESOURCE_FLAG_ALLOW_UNORDERED_ACCESS; - - HRESULT hr = ctx.device->CreateCommittedResource( - &heapProps, D3D12_HEAP_FLAG_SHARED, &desc, D3D12_RESOURCE_STATE_COMMON, - nullptr, IID_PPV_ARGS(&result.resource)); - if (FAILED(hr)) { - throw std::runtime_error("Failed to create named buffer"); - } - - // Create a NAMED shared handle and keep it open so the name persists - hr = ctx.device->CreateSharedHandle(result.resource.Get(), nullptr, - GENERIC_ALL, name, - &result.keepAliveHandle); - if (FAILED(hr)) { - throw std::runtime_error("Failed to create named shared handle"); - } - - std::wcout << L"[D3D12] Created named buffer: " << name << std::endl; - - return result; -} - -// Create a fence with a named shared handle -D3D12NamedFence createNamedExportableFence(D3D12Context &ctx, - const wchar_t *name) { - D3D12NamedFence result; - result.name = name; - - HRESULT hr = ctx.device->CreateFence(0, D3D12_FENCE_FLAG_SHARED, - IID_PPV_ARGS(&result.fence)); - if (FAILED(hr)) { - throw std::runtime_error("Failed to create fence"); - } - - // Create a NAMED shared handle and keep it open so the name persists - hr = ctx.device->CreateSharedHandle(result.fence.Get(), nullptr, GENERIC_ALL, - name, &result.keepAliveHandle); - if (FAILED(hr)) { - throw std::runtime_error("Failed to create named shared fence handle"); - } - - std::wcout << L"[D3D12] Created named fence: " << name << std::endl; - - return result; -} - -int main(int argc, char **argv) { - bool useSemaphores = true; - int iterations = 10; - size_t numElements = 1024; - - for (int i = 1; i < argc; ++i) { - std::string arg = argv[i]; - if (arg == "--no-sem") - useSemaphores = false; - if (arg == "--iterations" && i + 1 < argc) - iterations = std::stoi(argv[++i]); - if (arg == "--size" && i + 1 < argc) - numElements = std::stoi(argv[++i]); - } - - size_t bufferSize = numElements * sizeof(uint32_t); - - std::cout << "Running SYCL D3D12 resource_win32_name Native Test\n"; - std::cout << "Elements: " << numElements << " | Iterations: " << iterations - << " | Semaphores: " << (useSemaphores ? "ON" : "OFF") << "\n"; - - // D3D12 SETUP - D3D12Context d3dCtx = createD3D12Context(); - - // Create NAMED exportable buffers - D3D12NamedBuffer inBuf = createNamedExportableBuffer( - d3dCtx, bufferSize, L"Global\\SYCLTestInputBuffer3"); - D3D12NamedBuffer outBuf = createNamedExportableBuffer( - d3dCtx, bufferSize, L"Global\\SYCLTestOutputBuffer3"); - - // Host-visible staging buffers (not named, not exported) - D3D12BufferResources inStaging = createUploadBuffer(d3dCtx, bufferSize); - D3D12BufferResources outStaging = createReadbackBuffer(d3dCtx, bufferSize); - - // Interop Timeline Fence - NAMED - D3D12NamedFence extFence; - if (useSemaphores) { - extFence = createNamedExportableFence(d3dCtx, L"Global\\SYCLTestFence3"); - } - - // Set initial buffer states - d3dCtx.cmdAlloc->Reset(); - d3dCtx.cmdList->Reset(d3dCtx.cmdAlloc.Get(), nullptr); - D3D12_RESOURCE_BARRIER initialBarriers[2] = {}; - initialBarriers[0].Type = D3D12_RESOURCE_BARRIER_TYPE_TRANSITION; - initialBarriers[0].Transition.pResource = inBuf.resource.Get(); - initialBarriers[0].Transition.StateBefore = D3D12_RESOURCE_STATE_COMMON; - initialBarriers[0].Transition.StateAfter = D3D12_RESOURCE_STATE_COPY_DEST; - initialBarriers[0].Transition.Subresource = - D3D12_RESOURCE_BARRIER_ALL_SUBRESOURCES; - - initialBarriers[1].Type = D3D12_RESOURCE_BARRIER_TYPE_TRANSITION; - initialBarriers[1].Transition.pResource = outBuf.resource.Get(); - initialBarriers[1].Transition.StateBefore = D3D12_RESOURCE_STATE_COMMON; - initialBarriers[1].Transition.StateAfter = - D3D12_RESOURCE_STATE_UNORDERED_ACCESS; - initialBarriers[1].Transition.Subresource = - D3D12_RESOURCE_BARRIER_ALL_SUBRESOURCES; - d3dCtx.cmdList->ResourceBarrier(2, initialBarriers); - d3dCtx.cmdList->Close(); - executeAndWait(d3dCtx); - - // SYCL INTEROP - using resource_win32_name NATIVELY - try { - sycl::queue q; - auto device = q.get_device(); - auto context = q.get_context(); - - std::cout << "[SYCL] Device: " - << device.get_info() << std::endl; - - // Import buffers BY NAME using resource_win32_name - // Pass D3D12 device pointer so SYCL can open the named handle - std::cout << "[SYCL] Importing input buffer by name (native support)" - << std::endl; - syclexp::external_mem_descriptor inDesc{ - {(const void *)inBuf.name.c_str(), d3dCtx.device.Get()}, - syclexp::external_mem_handle_type::win32_nt_handle, - bufferSize}; - syclexp::external_mem inExtMem = - syclexp::import_external_memory(inDesc, device, context); - - std::cout << "[SYCL] Importing output buffer by name (native support)" - << std::endl; - syclexp::external_mem_descriptor outDesc{ - {(const void *)outBuf.name.c_str(), d3dCtx.device.Get()}, - syclexp::external_mem_handle_type::win32_nt_handle, - bufferSize}; - syclexp::external_mem outExtMem = - syclexp::import_external_memory(outDesc, device, context); - - // Import timeline fence BY NAME - syclexp::external_semaphore syclSem{}; - if (useSemaphores) { - std::cout << "[SYCL] Importing fence by name (native support)" - << std::endl; - auto semDesc = - syclexp::external_semaphore_descriptor{ - {(const void *)extFence.name.c_str(), d3dCtx.device.Get()}, - syclexp::external_semaphore_handle_type::win32_nt_dx12_fence}; - syclSem = syclexp::import_external_semaphore(semDesc, device, context); - } - - uint32_t *inPtr = static_cast( - syclexp::map_external_linear_memory(inExtMem, 0, bufferSize, q)); - uint32_t *outPtr = static_cast( - syclexp::map_external_linear_memory(outExtMem, 0, bufferSize, q)); - - std::cout << "[Test] Starting " << iterations << " iteration test..." - << std::endl; - - for (int i = 1; i <= iterations; ++i) { - uint64_t d3dSignalVal = (uint64_t)(2 * i - 1); - uint64_t syclSignalVal = (uint64_t)(2 * i); - - // D3D12: Upload and copy - void *mapped; - inStaging.resource->Map(0, nullptr, &mapped); - auto *data = static_cast(mapped); - for (size_t j = 0; j < numElements; ++j) - data[j] = (uint32_t)i; - inStaging.resource->Unmap(0, nullptr); - - d3dCtx.cmdAlloc->Reset(); - d3dCtx.cmdList->Reset(d3dCtx.cmdAlloc.Get(), nullptr); - - d3dCtx.cmdList->CopyBufferRegion(inBuf.resource.Get(), 0, - inStaging.resource.Get(), 0, bufferSize); - - // Barrier: CopyDest -> UAV - D3D12_RESOURCE_BARRIER barrier = {}; - barrier.Type = D3D12_RESOURCE_BARRIER_TYPE_TRANSITION; - barrier.Transition.pResource = inBuf.resource.Get(); - barrier.Transition.StateBefore = D3D12_RESOURCE_STATE_COPY_DEST; - barrier.Transition.StateAfter = D3D12_RESOURCE_STATE_UNORDERED_ACCESS; - barrier.Transition.Subresource = D3D12_RESOURCE_BARRIER_ALL_SUBRESOURCES; - d3dCtx.cmdList->ResourceBarrier(1, &barrier); - - d3dCtx.cmdList->Close(); - ID3D12CommandList *ppCommandLists[] = {d3dCtx.cmdList.Get()}; - d3dCtx.cmdQueue->ExecuteCommandLists(1, ppCommandLists); - - if (useSemaphores) { - d3dCtx.cmdQueue->Signal(extFence.fence.Get(), d3dSignalVal); - } - - // Host wait for upload to finish - d3dCtx.fenceValue++; - d3dCtx.cmdQueue->Signal(d3dCtx.fence.Get(), d3dCtx.fenceValue); - d3dCtx.fence->SetEventOnCompletion(d3dCtx.fenceValue, d3dCtx.fenceEvent); - WaitForSingleObject(d3dCtx.fenceEvent, INFINITE); - - std::cout << " [" << i << "] D3D12 upload done" << std::endl; - - // SYCL: Wait, execute, signal - if (useSemaphores) { - std::cout << ", SYCL sem-wait(" << d3dSignalVal << ")..." << std::endl; - q.ext_oneapi_wait_external_semaphore(syclSem, d3dSignalVal); - std::cout << "ok" << std::endl; - } - - q.submit([&](sycl::handler &h) { - h.parallel_for(sycl::range<1>(numElements), [=](sycl::item<1> item) { - size_t id = item.get_id(0); - outPtr[id] = inPtr[id] * 2; - }); - }); - - if (useSemaphores) { - std::cout << ", SYCL sem-signal(" << syclSignalVal << ")" << std::endl; - q.ext_oneapi_signal_external_semaphore(syclSem, syclSignalVal); - std::cout << "ok" << std::endl; - } - q.wait(); - std::cout << ", SYCL done" << std::endl; - - // D3D12: Readback and verify - if (useSemaphores) { - d3dCtx.cmdQueue->Wait(extFence.fence.Get(), syclSignalVal); - } - - d3dCtx.cmdAlloc->Reset(); - d3dCtx.cmdList->Reset(d3dCtx.cmdAlloc.Get(), nullptr); - - // Barrier: UAV -> CopySource - barrier.Transition.pResource = outBuf.resource.Get(); - barrier.Transition.StateBefore = D3D12_RESOURCE_STATE_UNORDERED_ACCESS; - barrier.Transition.StateAfter = D3D12_RESOURCE_STATE_COPY_SOURCE; - d3dCtx.cmdList->ResourceBarrier(1, &barrier); - - d3dCtx.cmdList->CopyBufferRegion(outStaging.resource.Get(), 0, - outBuf.resource.Get(), 0, bufferSize); - - // Barrier: revert outBuf back to UAV for next iteration - barrier.Transition.StateBefore = D3D12_RESOURCE_STATE_COPY_SOURCE; - barrier.Transition.StateAfter = D3D12_RESOURCE_STATE_UNORDERED_ACCESS; - d3dCtx.cmdList->ResourceBarrier(1, &barrier); - - d3dCtx.cmdList->Close(); - d3dCtx.cmdQueue->ExecuteCommandLists(1, ppCommandLists); - - // Host wait for readback - std::cout << ", d3d-fence..." << std::endl; - d3dCtx.fenceValue++; - d3dCtx.cmdQueue->Signal(d3dCtx.fence.Get(), d3dCtx.fenceValue); - d3dCtx.fence->SetEventOnCompletion(d3dCtx.fenceValue, d3dCtx.fenceEvent); - - if (WaitForSingleObject(d3dCtx.fenceEvent, 5000) == WAIT_TIMEOUT) { - std::cerr << "\nTIMEOUT on host wait!\n"; - return 1; - } - std::cout << "ok" << std::flush; - - // Verify data - outStaging.resource->Map(0, nullptr, &mapped); - auto *outData = static_cast(mapped); - uint32_t expected = (uint32_t)i * 2; - int errors = 0; - for (size_t j = 0; j < numElements; ++j) { - if (outData[j] != expected) { - if (errors++ < 5) - std::cerr << " [" << j << "]: got " << outData[j] << " expected " - << expected << std::endl; - } - } - outStaging.resource->Unmap(0, nullptr); - - if (errors > 0) { - std::cerr << "\nFAILURE at iteration " << i << ": " << errors - << " mismatches" << std::endl; - return 1; - } - - // Reset inBuf state to COPY_DEST for next iteration - d3dCtx.cmdAlloc->Reset(); - d3dCtx.cmdList->Reset(d3dCtx.cmdAlloc.Get(), nullptr); - barrier.Transition.pResource = inBuf.resource.Get(); - barrier.Transition.StateBefore = D3D12_RESOURCE_STATE_UNORDERED_ACCESS; - barrier.Transition.StateAfter = D3D12_RESOURCE_STATE_COPY_DEST; - d3dCtx.cmdList->ResourceBarrier(1, &barrier); - d3dCtx.cmdList->Close(); - executeAndWait(d3dCtx); - - if (i % 5 == 0 || i == 1) - std::cout << " PASS" << std::endl; - else - std::cout << " ok" << std::endl; - } - - std::cout << "SUCCESS! All " << iterations << " iterations passed." - << std::endl; - - // SYCL Cleanup - handles are automatically closed by release functions - syclexp::unmap_external_linear_memory(inPtr, q); - syclexp::unmap_external_linear_memory(outPtr, q); - if (useSemaphores) - syclexp::release_external_semaphore(syclSem, device, context); - syclexp::release_external_memory(inExtMem, device, context); - syclexp::release_external_memory(outExtMem, device, context); - - } catch (sycl::exception &e) { - std::cerr << "SYCL Exception: " << e.what() << std::endl; - return 1; - } - - // D3D12 Cleanup - CloseHandle(inBuf.keepAliveHandle); - CloseHandle(outBuf.keepAliveHandle); - if (useSemaphores) - CloseHandle(extFence.keepAliveHandle); - cleanupBuffer(inStaging); - cleanupBuffer(outStaging); - if (d3dCtx.fenceEvent) - CloseHandle(d3dCtx.fenceEvent); - - return 0; -}