[SYCL][NFC] Enable IPC tests

sycl/test-e2e/Experimental/ipc_memory.cpp

@@ -1,12 +1,9 @@
 // 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
-// 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
 
@@ -22,14 +19,37 @@
 #include <linux/prctl.h>
 #include <sys/prctl.h>
 #include <unistd.h>
+#elif defined(__WIN32__) || defined(_WIN32)
+#include <windows.h>
 #endif // defined(__linux__)
 
 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<char *>(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[]) {

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
 

unified-runtime/source/adapters/cuda/memory.cpp

@@ -10,6 +10,10 @@
 
 #include <cuda.h>
 
+#ifdef _WIN32
+#include <umf/experimental/ctl.h>
+#endif
+
 #include "common.hpp"
 #include "context.hpp"
 #include "enqueue.hpp"

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:

unified-runtime/source/adapters/level_zero/memory.cpp

@@ -13,6 +13,10 @@
 #include <string.h>
 #include <ur/ur.hpp>
 
+#ifdef _WIN32
+#include <umf/experimental/ctl.h>
+#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));

unified-runtime/source/adapters/level_zero/v2/memory.cpp

@@ -8,6 +8,10 @@
 //
 //===----------------------------------------------------------------------===//
 
+#ifdef _WIN32
+#include <umf/experimental/ctl.h>
+#endif
+
 #include "memory.hpp"
 
 #include "../ur_interface_loader.hpp"
@@ -860,6 +864,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) {
@@ -868,6 +886,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(
@@ -898,6 +918,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));