From 7508fc71786e744ba2315ba16a6b7994cb88ba62 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Thu, 27 Nov 2025 05:48:23 -0800 Subject: [PATCH 01/15] [SYCL][E2E] Enable ipc_memory test on Windows Signed-off-by: Larsen, Steffen --- sycl/test-e2e/Experimental/ipc_memory.cpp | 3 --- sycl/test-e2e/Experimental/ipc_put_after_free.cpp | 3 --- 2 files changed, 6 deletions(-) diff --git a/sycl/test-e2e/Experimental/ipc_memory.cpp b/sycl/test-e2e/Experimental/ipc_memory.cpp index 7366935a0cba4..8221144a001df 100644 --- a/sycl/test-e2e/Experimental/ipc_memory.cpp +++ b/sycl/test-e2e/Experimental/ipc_memory.cpp @@ -1,8 +1,5 @@ // REQUIRES: aspect-usm_device_allocations && aspect-ext_oneapi_ipc_memory -// UNSUPPORTED: level_zero && windows -// UNSUPPORTED-TRACKER: UMFW-348 - // DEFINE: %{cpp20} = %if cl_options %{/clang:-std=c++20%} %else %{-std=c++20%} // RUN: %{build} -o %t.out diff --git a/sycl/test-e2e/Experimental/ipc_put_after_free.cpp b/sycl/test-e2e/Experimental/ipc_put_after_free.cpp index 2851a2589eb7f..f376e4b0314fa 100644 --- a/sycl/test-e2e/Experimental/ipc_put_after_free.cpp +++ b/sycl/test-e2e/Experimental/ipc_put_after_free.cpp @@ -1,8 +1,5 @@ // REQUIRES: aspect-usm_device_allocations && aspect-ext_oneapi_ipc_memory -// UNSUPPORTED: level_zero && windows -// UNSUPPORTED-TRACKER: UMFW-348 - // RUN: %{build} -o %t.out // RUN: %{run} %t.out From 3e194c34d88db7b07df7a146d77393ca7bd3e822 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Fri, 28 Nov 2025 04:07:47 -0800 Subject: [PATCH 02/15] Bump UMF temporarily Signed-off-by: Larsen, Steffen --- unified-runtime/source/common/CMakeLists.txt | 6 +----- 1 file changed, 1 insertion(+), 5 deletions(-) diff --git a/unified-runtime/source/common/CMakeLists.txt b/unified-runtime/source/common/CMakeLists.txt index e21a0bce53841..18e83ae072690 100644 --- a/unified-runtime/source/common/CMakeLists.txt +++ b/unified-runtime/source/common/CMakeLists.txt @@ -81,11 +81,7 @@ if(umf_FOUND) else() set(UMF_REPO "https://github.com/oneapi-src/unified-memory-framework.git") - # commit 1209db2c5702b5de773ffc117b03e62f57f9554f - # Author: Łukasz Ślusarczyk - # Date: Mon Aug 25 13:35:07 2025 +0200 - # Add resident device change call - set(UMF_TAG v1.1.0-dev3) + set(UMF_TAG v1.1.0-dev4) if(NOT FETCHCONTENT_SOURCE_DIR_UNIFIED-MEMORY-FRAMEWORK) message(STATUS "Will fetch Unified Memory Framework from ${UMF_REPO}") From 73cffa0393db8e66e48bc11420e67423613fcbd5 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Tue, 16 Dec 2025 04:57:42 -0800 Subject: [PATCH 03/15] Use CreateProcess on Windows Signed-off-by: Larsen, Steffen --- sycl/test-e2e/Experimental/ipc_memory.cpp | 26 ++++++++++++++++++++--- 1 file changed, 23 insertions(+), 3 deletions(-) diff --git a/sycl/test-e2e/Experimental/ipc_memory.cpp b/sycl/test-e2e/Experimental/ipc_memory.cpp index 8221144a001df..935ca7960deff 100644 --- a/sycl/test-e2e/Experimental/ipc_memory.cpp +++ b/sycl/test-e2e/Experimental/ipc_memory.cpp @@ -19,6 +19,8 @@ #include #include #include +#elif defined(__WIN32__) || defined(_WIN32) +#include #endif // defined(__linux__) namespace syclexp = sycl::ext::oneapi::experimental; @@ -26,6 +28,26 @@ namespace syclexp = sycl::ext::oneapi::experimental; constexpr size_t N = 32; constexpr const char *CommsFile = "ipc_comms.txt"; +void spawn_and_sync(std::string Exe) { + std::cout << "Spawning: " << Exe << " 1" << std::endl; +#if defined(__WIN32__) || defined(_WIN32) + STARTUPINFO StartupInfo; + PROCESS_INFORMATION ProcInfo; + + std::memset(&ProcInfo, 0, sizeof(ProcInfo)); + std::memset(&StartupInfo, 0, sizeof(StartupInfo)); + StartupInfo.cb = sizeof(si); + CreateProcess(Exe.c_str(), "1", NULL, NULL, TRUE, 0, NULL, NULL, + &StartupInfo, &ProcInfo); + WaitForSingleObject(ProcInfo.hProcess, 30000); + CloseHandle(ProcInfo.hProcess); + CloseHandle(ProcInfo.hThread); +#else + std::string Cmd = Exe + " 1"; + std::system(Cmd.c_str()); +#endif +} + int spawner(int argc, char *argv[]) { assert(argc == 1); sycl::queue Q; @@ -64,9 +86,7 @@ int spawner(int argc, char *argv[]) { } // Spawn other process with an argument. - std::string Cmd = std::string{argv[0]} + " 1"; - std::cout << "Spawning: " << Cmd << std::endl; - std::system(Cmd.c_str()); + spawn_and_sync(std::string{argv[0]}); } int Failures = 0; From 5b3177326c02175e75621091033d7609c2cb5a31 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Tue, 16 Dec 2025 05:12:42 -0800 Subject: [PATCH 04/15] Fix formatting Signed-off-by: Larsen, Steffen --- sycl/test-e2e/Experimental/ipc_memory.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/test-e2e/Experimental/ipc_memory.cpp b/sycl/test-e2e/Experimental/ipc_memory.cpp index 935ca7960deff..b9f0cf3684c4d 100644 --- a/sycl/test-e2e/Experimental/ipc_memory.cpp +++ b/sycl/test-e2e/Experimental/ipc_memory.cpp @@ -37,8 +37,8 @@ void spawn_and_sync(std::string Exe) { std::memset(&ProcInfo, 0, sizeof(ProcInfo)); std::memset(&StartupInfo, 0, sizeof(StartupInfo)); StartupInfo.cb = sizeof(si); - CreateProcess(Exe.c_str(), "1", NULL, NULL, TRUE, 0, NULL, NULL, - &StartupInfo, &ProcInfo); + CreateProcess(Exe.c_str(), "1", NULL, NULL, TRUE, 0, NULL, NULL, &StartupInfo, + &ProcInfo); WaitForSingleObject(ProcInfo.hProcess, 30000); CloseHandle(ProcInfo.hProcess); CloseHandle(ProcInfo.hThread); From 295a9662ddf6f2fa6e10f9382358faa62b012ab5 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Tue, 16 Dec 2025 07:44:40 -0800 Subject: [PATCH 05/15] Change to CreateProcessA Signed-off-by: Larsen, Steffen --- sycl/test-e2e/Experimental/ipc_memory.cpp | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/sycl/test-e2e/Experimental/ipc_memory.cpp b/sycl/test-e2e/Experimental/ipc_memory.cpp index b9f0cf3684c4d..8417dc4bf0dd4 100644 --- a/sycl/test-e2e/Experimental/ipc_memory.cpp +++ b/sycl/test-e2e/Experimental/ipc_memory.cpp @@ -29,21 +29,21 @@ constexpr size_t N = 32; constexpr const char *CommsFile = "ipc_comms.txt"; void spawn_and_sync(std::string Exe) { - std::cout << "Spawning: " << Exe << " 1" << std::endl; + std::string Cmd = Exe + " 1"; + std::cout << "Spawning: " << Cmd << std::endl; #if defined(__WIN32__) || defined(_WIN32) STARTUPINFO StartupInfo; PROCESS_INFORMATION ProcInfo; std::memset(&ProcInfo, 0, sizeof(ProcInfo)); std::memset(&StartupInfo, 0, sizeof(StartupInfo)); - StartupInfo.cb = sizeof(si); - CreateProcess(Exe.c_str(), "1", NULL, NULL, TRUE, 0, NULL, NULL, &StartupInfo, - &ProcInfo); + StartupInfo.cb = sizeof(StartupInfo); + CreateProcessA(NULL, Cmd.c_str(), NULL, NULL, TRUE, 0, NULL, NULL, + &StartupInfo, &ProcInfo); WaitForSingleObject(ProcInfo.hProcess, 30000); CloseHandle(ProcInfo.hProcess); CloseHandle(ProcInfo.hThread); #else - std::string Cmd = Exe + " 1"; std::system(Cmd.c_str()); #endif } From 3a0907c9e3996c81eb56e5961cd9763a1865d5aa Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Tue, 16 Dec 2025 22:27:49 -0800 Subject: [PATCH 06/15] Const cast Signed-off-by: Larsen, Steffen --- sycl/test-e2e/Experimental/ipc_memory.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/test-e2e/Experimental/ipc_memory.cpp b/sycl/test-e2e/Experimental/ipc_memory.cpp index 8417dc4bf0dd4..4b2fd17b073b2 100644 --- a/sycl/test-e2e/Experimental/ipc_memory.cpp +++ b/sycl/test-e2e/Experimental/ipc_memory.cpp @@ -38,8 +38,8 @@ void spawn_and_sync(std::string Exe) { std::memset(&ProcInfo, 0, sizeof(ProcInfo)); std::memset(&StartupInfo, 0, sizeof(StartupInfo)); StartupInfo.cb = sizeof(StartupInfo); - CreateProcessA(NULL, Cmd.c_str(), NULL, NULL, TRUE, 0, NULL, NULL, - &StartupInfo, &ProcInfo); + CreateProcessA(NULL, const_cast(Cmd.c_str()), NULL, NULL, TRUE, 0, + NULL, NULL, &StartupInfo, &ProcInfo); WaitForSingleObject(ProcInfo.hProcess, 30000); CloseHandle(ProcInfo.hProcess); CloseHandle(ProcInfo.hThread); From 650fd712e9d51e25b38348271140e628e79f26fb Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Wed, 17 Dec 2025 00:25:48 -0800 Subject: [PATCH 07/15] Add more output Signed-off-by: Larsen, Steffen --- sycl/test-e2e/Experimental/ipc_memory.cpp | 12 ++++++++++-- 1 file changed, 10 insertions(+), 2 deletions(-) diff --git a/sycl/test-e2e/Experimental/ipc_memory.cpp b/sycl/test-e2e/Experimental/ipc_memory.cpp index 4b2fd17b073b2..c5db7aecad6c5 100644 --- a/sycl/test-e2e/Experimental/ipc_memory.cpp +++ b/sycl/test-e2e/Experimental/ipc_memory.cpp @@ -48,7 +48,8 @@ void spawn_and_sync(std::string Exe) { #endif } -int spawner(int argc, char *argv[]) { +int spawner(int argc, char *argv[]) try { + std::cout << "Running spanwer..." << std::endl; assert(argc == 1); sycl::queue Q; @@ -101,9 +102,13 @@ int spawner(int argc, char *argv[]) { } sycl::free(DataPtr, Q); return Failures; +} catch (sycl::exception &e) { + std::cout << "Spawner failed: " << e.what() << std::endl; + throw; } -int consumer() { +int consumer() try { + std::cout << "Running consumer..." << std::endl; sycl::queue Q; // Read the handle data. @@ -143,6 +148,9 @@ int consumer() { syclexp::ipc_memory::close(DataPtr, Q.get_context()); return Failures; +} catch (sycl::exception &e) { + std::cout << "Consumer failed: " << e.what() << std::endl; + throw; } int main(int argc, char *argv[]) { From aba6447da3a5954e8c3a966cf12a93cebc2f36ff Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Wed, 17 Dec 2025 03:23:41 -0800 Subject: [PATCH 08/15] Enable UMF windows workaround Signed-off-by: Larsen, Steffen --- .../source/adapters/level_zero/memory.cpp | 17 +++++++++++++++++ 1 file changed, 17 insertions(+) diff --git a/unified-runtime/source/adapters/level_zero/memory.cpp b/unified-runtime/source/adapters/level_zero/memory.cpp index 1a31be1e57595..f35f9b5fbdb45 100644 --- a/unified-runtime/source/adapters/level_zero/memory.cpp +++ b/unified-runtime/source/adapters/level_zero/memory.cpp @@ -13,6 +13,10 @@ #include #include +#ifdef _WIN32 +#include +#endif + #include "context.hpp" #include "event.hpp" #include "helpers/memory_helpers.hpp" @@ -1952,9 +1956,20 @@ ur_result_t urEnqueueWriteHostPipe( return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } +void enableWindowsUMFIPCWorkaround() { +#ifdef _WIN32 + // UMF on Windows currently requires a workaround for IPC to work. + int useImportExportForIPC = 1; + umfCtlSet("umf.provider.default.LEVEL_ZERO.params.use_import_export_for_IPC", + &useImportExportForIPC, sizeof(useImportExportForIPC)); +#endif +} + ur_result_t urIPCGetMemHandleExp(ur_context_handle_t, void *pMem, void **ppIPCMemHandleData, size_t *pIPCMemHandleDataSizeRet) { + enableWindowsUMFIPCWorkaround(); + umf_memory_pool_handle_t umfPool; auto urRet = umf::umf2urResult(umfPoolByPtr(pMem, &umfPool)); if (urRet) @@ -1983,6 +1998,8 @@ ur_result_t urIPCOpenMemHandleExp(ur_context_handle_t hContext, ur_device_handle_t hDevice, void *pIPCMemHandleData, size_t ipcMemHandleDataSize, void **ppMem) { + enableWindowsUMFIPCWorkaround(); + auto *pool = hContext->DefaultPool.getPool(usm::pool_descriptor{ &hContext->DefaultPool, hContext, hDevice, UR_USM_TYPE_DEVICE, false}); if (!pool) From 150e10a5fc8fd74700422ed647aadbca419bc913 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Wed, 17 Dec 2025 03:26:47 -0800 Subject: [PATCH 09/15] Add workaround to L0v2 and CUDA Signed-off-by: Larsen, Steffen --- unified-runtime/source/adapters/cuda/memory.cpp | 17 +++++++++++++++++ .../source/adapters/level_zero/v2/memory.cpp | 17 +++++++++++++++++ 2 files changed, 34 insertions(+) diff --git a/unified-runtime/source/adapters/cuda/memory.cpp b/unified-runtime/source/adapters/cuda/memory.cpp index 6eb0c8e12ad3a..7d7f499050770 100644 --- a/unified-runtime/source/adapters/cuda/memory.cpp +++ b/unified-runtime/source/adapters/cuda/memory.cpp @@ -10,6 +10,10 @@ #include +#ifdef _WIN32 +#include +#endif + #include "common.hpp" #include "context.hpp" #include "enqueue.hpp" @@ -591,9 +595,20 @@ CUsurfObject SurfaceMem::getSurface(const ur_device_handle_t Device) { return SurfObjs[OuterMemStruct->getContext()->getDeviceIndex(Device)]; } +void enableWindowsUMFIPCWorkaround() { +#ifdef _WIN32 + // UMF on Windows currently requires a workaround for IPC to work. + int useImportExportForIPC = 1; + umfCtlSet("umf.provider.default.LEVEL_ZERO.params.use_import_export_for_IPC", + &useImportExportForIPC, sizeof(useImportExportForIPC)); +#endif +} + UR_APIEXPORT ur_result_t UR_APICALL urIPCGetMemHandleExp(ur_context_handle_t, void *pMem, void **ppIPCMemHandleData, size_t *pIPCMemHandleDataSizeRet) { + enableWindowsUMFIPCWorkaround(); + umf_memory_pool_handle_t umfPool; auto urRet = umf::umf2urResult(umfPoolByPtr(pMem, &umfPool)); if (urRet) @@ -622,6 +637,8 @@ urIPCPutMemHandleExp(ur_context_handle_t, void *pIPCMemHandleData) { UR_APIEXPORT ur_result_t UR_APICALL urIPCOpenMemHandleExp( ur_context_handle_t, ur_device_handle_t hDevice, void *pIPCMemHandleData, size_t ipcMemHandleDataSize, void **ppMem) { + enableWindowsUMFIPCWorkaround(); + umf_memory_pool_handle_t umfPool = hDevice->MemoryPoolDevice; size_t umfHandleSize = 0; diff --git a/unified-runtime/source/adapters/level_zero/v2/memory.cpp b/unified-runtime/source/adapters/level_zero/v2/memory.cpp index 6473ebf69fdae..87ad2acad16e3 100644 --- a/unified-runtime/source/adapters/level_zero/v2/memory.cpp +++ b/unified-runtime/source/adapters/level_zero/v2/memory.cpp @@ -8,6 +8,10 @@ // //===----------------------------------------------------------------------===// +#ifdef _WIN32 +#include +#endif + #include "memory.hpp" #include "../ur_interface_loader.hpp" #include "context.hpp" @@ -781,9 +785,20 @@ ur_result_t urMemImageGetInfo(ur_mem_handle_t /*hMemory*/, return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } +void enableWindowsUMFIPCWorkaround() { +#ifdef _WIN32 + // UMF on Windows currently requires a workaround for IPC to work. + int useImportExportForIPC = 1; + umfCtlSet("umf.provider.default.LEVEL_ZERO.params.use_import_export_for_IPC", + &useImportExportForIPC, sizeof(useImportExportForIPC)); +#endif +} + ur_result_t urIPCGetMemHandleExp(ur_context_handle_t, void *pMem, void **ppIPCMemHandleData, size_t *pIPCMemHandleDataSizeRet) { + enableWindowsUMFIPCWorkaround(); + umf_memory_pool_handle_t umfPool; auto urRet = umf::umf2urResult(umfPoolByPtr(pMem, &umfPool)); if (urRet) @@ -812,6 +827,8 @@ ur_result_t urIPCOpenMemHandleExp(ur_context_handle_t hContext, ur_device_handle_t hDevice, void *pIPCMemHandleData, size_t ipcMemHandleDataSize, void **ppMem) { + enableWindowsUMFIPCWorkaround(); + auto *pool = hContext->getDefaultUSMPool()->getPool( usm::pool_descriptor{hContext->getDefaultUSMPool(), hContext, hDevice, UR_USM_TYPE_DEVICE, false}); From 4e5663fdfa4b341930cff30b14784531cb1cc822 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Wed, 17 Dec 2025 04:34:01 -0800 Subject: [PATCH 10/15] Provide specific provider Signed-off-by: Larsen, Steffen --- .../source/adapters/cuda/memory.cpp | 13 ------------ .../source/adapters/level_zero/memory.cpp | 20 ++++++++++++------ .../source/adapters/level_zero/v2/memory.cpp | 21 ++++++++++++------- 3 files changed, 28 insertions(+), 26 deletions(-) diff --git a/unified-runtime/source/adapters/cuda/memory.cpp b/unified-runtime/source/adapters/cuda/memory.cpp index 7d7f499050770..7172fabcab2ae 100644 --- a/unified-runtime/source/adapters/cuda/memory.cpp +++ b/unified-runtime/source/adapters/cuda/memory.cpp @@ -595,20 +595,9 @@ CUsurfObject SurfaceMem::getSurface(const ur_device_handle_t Device) { return SurfObjs[OuterMemStruct->getContext()->getDeviceIndex(Device)]; } -void enableWindowsUMFIPCWorkaround() { -#ifdef _WIN32 - // UMF on Windows currently requires a workaround for IPC to work. - int useImportExportForIPC = 1; - umfCtlSet("umf.provider.default.LEVEL_ZERO.params.use_import_export_for_IPC", - &useImportExportForIPC, sizeof(useImportExportForIPC)); -#endif -} - UR_APIEXPORT ur_result_t UR_APICALL urIPCGetMemHandleExp(ur_context_handle_t, void *pMem, void **ppIPCMemHandleData, size_t *pIPCMemHandleDataSizeRet) { - enableWindowsUMFIPCWorkaround(); - umf_memory_pool_handle_t umfPool; auto urRet = umf::umf2urResult(umfPoolByPtr(pMem, &umfPool)); if (urRet) @@ -637,8 +626,6 @@ urIPCPutMemHandleExp(ur_context_handle_t, void *pIPCMemHandleData) { UR_APIEXPORT ur_result_t UR_APICALL urIPCOpenMemHandleExp( ur_context_handle_t, ur_device_handle_t hDevice, void *pIPCMemHandleData, size_t ipcMemHandleDataSize, void **ppMem) { - enableWindowsUMFIPCWorkaround(); - umf_memory_pool_handle_t umfPool = hDevice->MemoryPoolDevice; size_t umfHandleSize = 0; diff --git a/unified-runtime/source/adapters/level_zero/memory.cpp b/unified-runtime/source/adapters/level_zero/memory.cpp index f35f9b5fbdb45..b0da772a7ee9f 100644 --- a/unified-runtime/source/adapters/level_zero/memory.cpp +++ b/unified-runtime/source/adapters/level_zero/memory.cpp @@ -1956,25 +1956,33 @@ ur_result_t urEnqueueWriteHostPipe( return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } -void enableWindowsUMFIPCWorkaround() { +inline ur_result_t enableWindowsUMFIPCWorkaround( + [[maybe_unused]] umf_memory_pool_handle_t umfPool) { #ifdef _WIN32 // UMF on Windows currently requires a workaround for IPC to work. + umf_memory_provider_handle_t umfProvider = nullptr; + UR_CALL(umf::umf2urResult(umfPoolGetMemoryProvider(umfPool, &umfProvider))); + if (urRet) + return urRet; int useImportExportForIPC = 1; - umfCtlSet("umf.provider.default.LEVEL_ZERO.params.use_import_export_for_IPC", - &useImportExportForIPC, sizeof(useImportExportForIPC)); + UR_CALL(umf::umf2urResult(umfCtlSet( + "umf.provider.by_handle.{}.LEVEL_ZERO.params.use_import_export_for_IPC", + &useImportExportForIPC, sizeof(useImportExportForIPC), umfProvider))); #endif + return UR_RESULT_SUCCESS; } ur_result_t urIPCGetMemHandleExp(ur_context_handle_t, void *pMem, void **ppIPCMemHandleData, size_t *pIPCMemHandleDataSizeRet) { - enableWindowsUMFIPCWorkaround(); umf_memory_pool_handle_t umfPool; auto urRet = umf::umf2urResult(umfPoolByPtr(pMem, &umfPool)); if (urRet) return urRet; + UR_CALL(enableWindowsUMFIPCWorkaround(umfPool)); + // Fast path for returning the size of the handle only. if (!ppIPCMemHandleData) return umf::umf2urResult( @@ -1998,14 +2006,14 @@ ur_result_t urIPCOpenMemHandleExp(ur_context_handle_t hContext, ur_device_handle_t hDevice, void *pIPCMemHandleData, size_t ipcMemHandleDataSize, void **ppMem) { - enableWindowsUMFIPCWorkaround(); - auto *pool = hContext->DefaultPool.getPool(usm::pool_descriptor{ &hContext->DefaultPool, hContext, hDevice, UR_USM_TYPE_DEVICE, false}); if (!pool) return UR_RESULT_ERROR_INVALID_CONTEXT; umf_memory_pool_handle_t umfPool = pool->UmfPool.get(); + UR_CALL(enableWindowsUMFIPCWorkaround(umfPool)); + size_t umfHandleSize = 0; auto urRet = umf::umf2urResult(umfPoolGetIPCHandleSize(umfPool, &umfHandleSize)); diff --git a/unified-runtime/source/adapters/level_zero/v2/memory.cpp b/unified-runtime/source/adapters/level_zero/v2/memory.cpp index 87ad2acad16e3..cb21871945ccf 100644 --- a/unified-runtime/source/adapters/level_zero/v2/memory.cpp +++ b/unified-runtime/source/adapters/level_zero/v2/memory.cpp @@ -785,25 +785,32 @@ ur_result_t urMemImageGetInfo(ur_mem_handle_t /*hMemory*/, return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } -void enableWindowsUMFIPCWorkaround() { +inline ur_result_t enableWindowsUMFIPCWorkaround( + [[maybe_unused]] umf_memory_pool_handle_t umfPool) { #ifdef _WIN32 // UMF on Windows currently requires a workaround for IPC to work. + umf_memory_provider_handle_t umfProvider = nullptr; + UR_CALL(umf::umf2urResult(umfPoolGetMemoryProvider(umfPool, &umfProvider))); + if (urRet) + return urRet; int useImportExportForIPC = 1; - umfCtlSet("umf.provider.default.LEVEL_ZERO.params.use_import_export_for_IPC", - &useImportExportForIPC, sizeof(useImportExportForIPC)); + UR_CALL(umf::umf2urResult(umfCtlSet( + "umf.provider.by_handle.{}.LEVEL_ZERO.params.use_import_export_for_IPC", + &useImportExportForIPC, sizeof(useImportExportForIPC), umfProvider))); #endif + return UR_RESULT_SUCCESS; } ur_result_t urIPCGetMemHandleExp(ur_context_handle_t, void *pMem, void **ppIPCMemHandleData, size_t *pIPCMemHandleDataSizeRet) { - enableWindowsUMFIPCWorkaround(); - umf_memory_pool_handle_t umfPool; auto urRet = umf::umf2urResult(umfPoolByPtr(pMem, &umfPool)); if (urRet) return urRet; + UR_CALL(enableWindowsUMFIPCWorkaround(umfPool)); + // Fast path for returning the size of the handle only. if (!ppIPCMemHandleData) return umf::umf2urResult( @@ -827,8 +834,6 @@ ur_result_t urIPCOpenMemHandleExp(ur_context_handle_t hContext, ur_device_handle_t hDevice, void *pIPCMemHandleData, size_t ipcMemHandleDataSize, void **ppMem) { - enableWindowsUMFIPCWorkaround(); - auto *pool = hContext->getDefaultUSMPool()->getPool( usm::pool_descriptor{hContext->getDefaultUSMPool(), hContext, hDevice, UR_USM_TYPE_DEVICE, false}); @@ -836,6 +841,8 @@ ur_result_t urIPCOpenMemHandleExp(ur_context_handle_t hContext, return UR_RESULT_ERROR_INVALID_CONTEXT; umf_memory_pool_handle_t umfPool = pool->umfPool.get(); + UR_CALL(enableWindowsUMFIPCWorkaround(umfPool)); + size_t umfHandleSize = 0; auto urRet = umf::umf2urResult(umfPoolGetIPCHandleSize(umfPool, &umfHandleSize)); From 52628b851511373d6aff2e223342a74299a15310 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Wed, 17 Dec 2025 05:16:23 -0800 Subject: [PATCH 11/15] Fix Signed-off-by: Larsen, Steffen --- unified-runtime/source/adapters/level_zero/memory.cpp | 2 -- unified-runtime/source/adapters/level_zero/v2/memory.cpp | 2 -- 2 files changed, 4 deletions(-) diff --git a/unified-runtime/source/adapters/level_zero/memory.cpp b/unified-runtime/source/adapters/level_zero/memory.cpp index b0da772a7ee9f..bd02e9a2b02d9 100644 --- a/unified-runtime/source/adapters/level_zero/memory.cpp +++ b/unified-runtime/source/adapters/level_zero/memory.cpp @@ -1962,8 +1962,6 @@ inline ur_result_t enableWindowsUMFIPCWorkaround( // UMF on Windows currently requires a workaround for IPC to work. umf_memory_provider_handle_t umfProvider = nullptr; UR_CALL(umf::umf2urResult(umfPoolGetMemoryProvider(umfPool, &umfProvider))); - if (urRet) - return urRet; int useImportExportForIPC = 1; UR_CALL(umf::umf2urResult(umfCtlSet( "umf.provider.by_handle.{}.LEVEL_ZERO.params.use_import_export_for_IPC", diff --git a/unified-runtime/source/adapters/level_zero/v2/memory.cpp b/unified-runtime/source/adapters/level_zero/v2/memory.cpp index cb21871945ccf..45cccfec76d7f 100644 --- a/unified-runtime/source/adapters/level_zero/v2/memory.cpp +++ b/unified-runtime/source/adapters/level_zero/v2/memory.cpp @@ -791,8 +791,6 @@ inline ur_result_t enableWindowsUMFIPCWorkaround( // UMF on Windows currently requires a workaround for IPC to work. umf_memory_provider_handle_t umfProvider = nullptr; UR_CALL(umf::umf2urResult(umfPoolGetMemoryProvider(umfPool, &umfProvider))); - if (urRet) - return urRet; int useImportExportForIPC = 1; UR_CALL(umf::umf2urResult(umfCtlSet( "umf.provider.by_handle.{}.LEVEL_ZERO.params.use_import_export_for_IPC", From b6a96697e83ad171da90f00b224458208392aeae Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Thu, 8 Jan 2026 07:23:01 -0800 Subject: [PATCH 12/15] Remove windows disablement Signed-off-by: Larsen, Steffen --- unified-runtime/source/adapters/level_zero/device.cpp | 5 ----- 1 file changed, 5 deletions(-) diff --git a/unified-runtime/source/adapters/level_zero/device.cpp b/unified-runtime/source/adapters/level_zero/device.cpp index ea9dd479daf02..a0ab6614bb7bc 100644 --- a/unified-runtime/source/adapters/level_zero/device.cpp +++ b/unified-runtime/source/adapters/level_zero/device.cpp @@ -1307,12 +1307,7 @@ ur_result_t urDeviceGetInfo( #endif } case UR_DEVICE_INFO_IPC_MEMORY_SUPPORT_EXP: -#ifdef _WIN32 - // TODO: Remove when IPC memory works in UMF on Windows. - return ReturnValue(false); -#else return ReturnValue(true); -#endif case UR_DEVICE_INFO_ASYNC_BARRIER: return ReturnValue(false); case UR_DEVICE_INFO_HOST_PIPE_READ_WRITE_SUPPORT: From edec3880c188df90700eef04a004e90c229c0c34 Mon Sep 17 00:00:00 2001 From: "Kornev, Nikita" Date: Wed, 22 Apr 2026 18:15:38 +0200 Subject: [PATCH 13/15] [SYCL][NFC] Enable IPC tests UMF issue was fixed. --- sycl/test-e2e/Experimental/ipc_memory.cpp | 3 --- sycl/test-e2e/Experimental/ipc_put_after_free.cpp | 3 --- 2 files changed, 6 deletions(-) diff --git a/sycl/test-e2e/Experimental/ipc_memory.cpp b/sycl/test-e2e/Experimental/ipc_memory.cpp index 7366935a0cba4..8221144a001df 100644 --- a/sycl/test-e2e/Experimental/ipc_memory.cpp +++ b/sycl/test-e2e/Experimental/ipc_memory.cpp @@ -1,8 +1,5 @@ // REQUIRES: aspect-usm_device_allocations && aspect-ext_oneapi_ipc_memory -// UNSUPPORTED: level_zero && windows -// UNSUPPORTED-TRACKER: UMFW-348 - // DEFINE: %{cpp20} = %if cl_options %{/clang:-std=c++20%} %else %{-std=c++20%} // RUN: %{build} -o %t.out diff --git a/sycl/test-e2e/Experimental/ipc_put_after_free.cpp b/sycl/test-e2e/Experimental/ipc_put_after_free.cpp index 2851a2589eb7f..f376e4b0314fa 100644 --- a/sycl/test-e2e/Experimental/ipc_put_after_free.cpp +++ b/sycl/test-e2e/Experimental/ipc_put_after_free.cpp @@ -1,8 +1,5 @@ // REQUIRES: aspect-usm_device_allocations && aspect-ext_oneapi_ipc_memory -// UNSUPPORTED: level_zero && windows -// UNSUPPORTED-TRACKER: UMFW-348 - // RUN: %{build} -o %t.out // RUN: %{run} %t.out From 151b5f59704243a3520928266dcc5341de177749 Mon Sep 17 00:00:00 2001 From: "Kornev, Nikita" Date: Wed, 22 Apr 2026 19:08:28 +0200 Subject: [PATCH 14/15] update-level-zero-adapter --- unified-runtime/source/adapters/level_zero/device.cpp | 5 ----- 1 file changed, 5 deletions(-) diff --git a/unified-runtime/source/adapters/level_zero/device.cpp b/unified-runtime/source/adapters/level_zero/device.cpp index ed3e9a8e88b9d..4dcf56970fd1b 100644 --- a/unified-runtime/source/adapters/level_zero/device.cpp +++ b/unified-runtime/source/adapters/level_zero/device.cpp @@ -1367,12 +1367,7 @@ ur_result_t urDeviceGetInfo( #endif } case UR_DEVICE_INFO_IPC_MEMORY_SUPPORT_EXP: -#ifdef _WIN32 - // TODO: Remove when IPC memory works in UMF on Windows. - return ReturnValue(false); -#else return ReturnValue(true); -#endif case UR_DEVICE_INFO_ASYNC_BARRIER: return ReturnValue(false); case UR_DEVICE_INFO_HOST_PIPE_READ_WRITE_SUPPORT: From a2504afed1fe4e297fc372a78a85084ebae468b7 Mon Sep 17 00:00:00 2001 From: "Kornev, Nikita" Date: Thu, 23 Apr 2026 12:03:33 +0200 Subject: [PATCH 15/15] log --- sycl/test-e2e/Experimental/ipc_memory.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/Experimental/ipc_memory.cpp b/sycl/test-e2e/Experimental/ipc_memory.cpp index 8221144a001df..86729585bfdbc 100644 --- a/sycl/test-e2e/Experimental/ipc_memory.cpp +++ b/sycl/test-e2e/Experimental/ipc_memory.cpp @@ -3,7 +3,7 @@ // DEFINE: %{cpp20} = %if cl_options %{/clang:-std=c++20%} %else %{-std=c++20%} // RUN: %{build} -o %t.out -// RUN: %{run} %t.out +// RUN: %{run} SYCL_UR_TRACE=-1 %t.out // RUN: %{build} -DUSE_VIEW %{cpp20} -o %t.view.out // RUN: %{run} %t.view.out