diff --git a/sycl/test-e2e/Experimental/ipc_memory.cpp b/sycl/test-e2e/Experimental/ipc_memory.cpp index 7366935a0cba4..c5db7aecad6c5 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 @@ -22,6 +19,8 @@ #include #include #include +#elif defined(__WIN32__) || defined(_WIN32) +#include #endif // defined(__linux__) namespace syclexp = sycl::ext::oneapi::experimental; @@ -29,7 +28,28 @@ namespace syclexp = sycl::ext::oneapi::experimental; constexpr size_t N = 32; constexpr const char *CommsFile = "ipc_comms.txt"; -int spawner(int argc, char *argv[]) { +void spawn_and_sync(std::string Exe) { + 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(StartupInfo); + 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); +#else + std::system(Cmd.c_str()); +#endif +} + +int spawner(int argc, char *argv[]) try { + std::cout << "Running spanwer..." << std::endl; assert(argc == 1); sycl::queue Q; @@ -67,9 +87,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; @@ -84,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. @@ -126,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[]) { 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 diff --git a/unified-runtime/source/adapters/cuda/memory.cpp b/unified-runtime/source/adapters/cuda/memory.cpp index 6eb0c8e12ad3a..7172fabcab2ae 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" 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: diff --git a/unified-runtime/source/adapters/level_zero/memory.cpp b/unified-runtime/source/adapters/level_zero/memory.cpp index 1a31be1e57595..bd02e9a2b02d9 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,14 +1956,31 @@ ur_result_t urEnqueueWriteHostPipe( return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } +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))); + int useImportExportForIPC = 1; + 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) { + 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( @@ -1989,6 +2010,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)); diff --git a/unified-runtime/source/adapters/level_zero/v2/memory.cpp b/unified-runtime/source/adapters/level_zero/v2/memory.cpp index 6473ebf69fdae..45cccfec76d7f 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,6 +785,20 @@ ur_result_t urMemImageGetInfo(ur_mem_handle_t /*hMemory*/, return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } +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))); + int useImportExportForIPC = 1; + 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) { @@ -789,6 +807,8 @@ ur_result_t urIPCGetMemHandleExp(ur_context_handle_t, void *pMem, if (urRet) return urRet; + UR_CALL(enableWindowsUMFIPCWorkaround(umfPool)); + // Fast path for returning the size of the handle only. if (!ppIPCMemHandleData) return umf::umf2urResult( @@ -819,6 +839,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));