intel · bratpiorka · Apr 22, 2026 · Apr 22, 2026 · Apr 23, 2026 · May 13, 2026
@@ -50,6 +50,11 @@ on:
         type: string
         required: false
 
+      e2e_extra_lit_opts:
+        type: string
+        required: false
+        default: ''
+
       e2e_binaries_new_offload_model_artifact:
         type: string
         required: false
@@ -251,7 +256,7 @@ jobs:
         testing_mode: build-only
         target_devices: all
         binaries_artifact: ${{ inputs.e2e_binaries_artifact }}
-        extra_lit_opts: --param sycl_build_targets="spir"
+        extra_lit_opts: --param sycl_build_targets="spir" ${{ inputs.e2e_extra_lit_opts }}
         cxx: ${{ inputs.cxx }}
 
     - name: Build E2E tests with New Offload Model

@@ -58,6 +58,7 @@ jobs:
     with:
       changes: ${{ needs.detect_changes.outputs.filters }}
       e2e_binaries_artifact: sycl_windows_e2ebin
+      e2e_extra_lit_opts: --filter=Experimental/ipc_memory\.cpp$
       e2e_binaries_new_offload_model_artifact: sycl_windows_e2ebin_with_new_offload_model
 
   run_prebuilt_e2e_tests:
@@ -84,6 +85,7 @@ jobs:
       toolchain_artifact_filename: ${{ needs.build.outputs.toolchain_artifact_filename }}
       testing_mode: run-only
       binaries_artifact: sycl_windows_e2ebin
+      extra_lit_opts: --filter=Experimental/ipc_memory\.cpp$
 
   run_prebuilt_e2e_with_new_offload_model_tests:
     needs: build

@@ -76,7 +76,7 @@ runs:
     continue-on-error: true
     shell: bash
     env:
-      LIT_OPTS: -v --no-progress-bar --show-unsupported --show-pass --show-xfail --max-time ${{ inputs.testing_mode == 'run-only' && 1200 || 3600 }} --time-tests --param print_features=True --param test-mode=${{ inputs.testing_mode }} --param sycl_devices=${{ inputs.target_devices }} ${{ inputs.extra_lit_opts }}
+      LIT_OPTS: -vv --no-progress-bar --show-unsupported --show-pass --show-xfail --max-time ${{ inputs.testing_mode == 'run-only' && 1200 || 3600 }} --time-tests --param print_features=True --param test-mode=${{ inputs.testing_mode }} --param sycl_devices=${{ inputs.target_devices }} ${{ inputs.extra_lit_opts }}
     run: |
       cmake --build build-e2e --target check-sycl-e2e > e2e.log 2>&1
   # Two steps below are duplicated between Lin/Win actions, updates must change both
@@ -87,6 +87,43 @@ runs:
       echo "::group::Show Full E2E Log"
       cat e2e.log
       echo "::endgroup::"
+  - name: E2E failure diagnostics
+    if: steps.run_e2e.outcome != 'success'
+    shell: bash
+    run: |
+      echo "::group::IPC Failure Diagnostics"
+      echo "working-directory=$(pwd)"
+      if [ -f build-e2e/CMakeCache.txt ]; then
+        echo "::group::Relevant CMakeCache entries"
+        grep -nE "CMAKE_CXX_COMPILER:|LLVM_LIT:|SYCL_E2E_TESTS_LIT_FLAGS:|CMAKE_GENERATOR:|LEVEL_ZERO_" build-e2e/CMakeCache.txt || true
+        echo "::endgroup::"
+      fi
+      if [ -f build-e2e/build.ninja ]; then
+        echo "::group::check-sycl-e2e target excerpt"
+        grep -nE "check-sycl-e2e|ipc_memory\.cpp" build-e2e/build.ninja || true
+        echo "::endgroup::"
+      fi
+      if [ -d build-e2e/Experimental/Output ]; then
+        echo "::group::Experimental Output listing"
+        find build-e2e/Experimental/Output -maxdepth 1 -type f | sort || true
+        echo "::endgroup::"
+        for path in build-e2e/Experimental/Output/ipc_memory.cpp*; do
+          [ -e "$path" ] || continue
+          echo "::group::Artifact $path"
+          ls -l "$path"
+          echo "::endgroup::"
+        done
+        for exe in build-e2e/Experimental/Output/ipc_memory.cpp*.out; do
+          [ -f "$exe" ] || continue
+          echo "::group::Direct rerun $exe"
+          "$(realpath "$exe")" 2>&1 || true
+          echo "::endgroup::"
+        done
+      fi
+      echo "::group::ipc_memory excerpt from e2e.log"
+      grep -nE "ipc_memory\.cpp|FAILED:|executed command|command stdout|command stderr|error:" e2e.log || true
+      echo "::endgroup::"
+      echo "::endgroup::"
   - name: Report E2E Failures
     if: steps.run_e2e.outcome != 'success'
     shell: bash

@@ -1,37 +1,138 @@
 // 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-unfiltered-devices} %t.out
 // RUN: %{build} -DUSE_VIEW %{cpp20} -o %t.view.out
-// RUN: %{run} %t.view.out
+// RUN: %{run-unfiltered-devices} %t.view.out
 
 #include <sycl/detail/core.hpp>
 #include <sycl/ext/oneapi/experimental/ipc_memory.hpp>
 #include <sycl/usm.hpp>
 
+#include <cassert>
 #include <cstdio>
 #include <cstdlib>
+#include <cstring>
 #include <fstream>
+#include <iostream>
+#include <memory>
+#include <stdexcept>
+#include <string>
 
 #if defined(__linux__)
 #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[]) {
+static void print_env(const char *Name) {
+  const char *Value = std::getenv(Name);
+  std::cout << Name << '=' << (Value ? Value : "<unset>") << std::endl;
+}
+
+static void set_env(const char *Name, const char *Value) {
+#if defined(__WIN32__) || defined(_WIN32)
+  if (!SetEnvironmentVariableA(Name, Value)) {
+    std::cout << "SetEnvironmentVariableA failed for " << Name << ": "
+              << GetLastError() << std::endl;
+    throw std::runtime_error("SetEnvironmentVariableA failed");
+  }
+#else
+  if (setenv(Name, Value, 1) != 0) {
+    throw std::runtime_error("setenv failed");
+  }
+#endif
+}
+
+static void configure_runtime_diagnostics_env() {
+  set_env("ONEAPI_DEVICE_SELECTOR", "level_zero:gpu");
+  set_env("UR_L0_V2_ENABLE_WINDOWS_IPC_WA", "1");
+  set_env("SYCL_UR_TRACE", "-1");
+  set_env("UR_LOG_LOADER", "level:debug;output:stdout;flush:debug");
+  set_env("UR_LOG_LEVEL_ZERO", "level:debug;output:stdout;flush:debug");
+  set_env("UMF_LOG", "level:debug;flush:debug;output:stdout;pid:yes");
+}
+
+static void print_runtime_diagnostics(const char *Role, sycl::queue &Q) {
+  std::cout << '[' << Role << "] backend=" << static_cast<int>(Q.get_backend())
+            << " device=" << Q.get_device().get_info<sycl::info::device::name>()
+            << std::endl;
+  print_env("ONEAPI_DEVICE_SELECTOR");
+  print_env("UR_L0_V2_ENABLE_WINDOWS_IPC_WA");
+  print_env("SYCL_UR_TRACE");
+  print_env("UR_LOG_LOADER");
+  print_env("UR_LOG_LEVEL_ZERO");
+  print_env("UMF_LOG");
+}
+
+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);
+  BOOL Created =
+      CreateProcessA(NULL, const_cast<char *>(Cmd.c_str()), NULL, NULL, TRUE, 0,
+                     NULL, NULL, &StartupInfo, &ProcInfo);
+  std::cout << "CreateProcessA result: " << Created << std::endl;
+  if (!Created) {
+    std::cout << "CreateProcessA GetLastError: " << GetLastError() << std::endl;
+    throw std::runtime_error("CreateProcessA failed");
+  }
+
+  DWORD WaitStatus = WaitForSingleObject(ProcInfo.hProcess, 30000);
+  std::cout << "WaitForSingleObject result: " << WaitStatus << std::endl;
+  if (WaitStatus == WAIT_FAILED) {
+    std::cout << "WaitForSingleObject GetLastError: " << GetLastError()
+              << std::endl;
+    CloseHandle(ProcInfo.hProcess);
+    CloseHandle(ProcInfo.hThread);
+    throw std::runtime_error("WaitForSingleObject failed");
+  }
+
+  if (WaitStatus == WAIT_TIMEOUT) {
+    CloseHandle(ProcInfo.hProcess);
+    CloseHandle(ProcInfo.hThread);
+    throw std::runtime_error("Child process timed out");
+  }
+
+  DWORD ExitCode = 0;
+  if (!GetExitCodeProcess(ProcInfo.hProcess, &ExitCode)) {
+    std::cout << "GetExitCodeProcess GetLastError: " << GetLastError()
+              << std::endl;
+    CloseHandle(ProcInfo.hProcess);
+    CloseHandle(ProcInfo.hThread);
+    throw std::runtime_error("GetExitCodeProcess failed");
+  }
+  std::cout << "Child exit code: " << ExitCode << std::endl;
+  CloseHandle(ProcInfo.hProcess);
+  CloseHandle(ProcInfo.hThread);
+  if (ExitCode != 0)
+    throw std::runtime_error("Child process returned non-zero exit code");
+#else
+  std::system(Cmd.c_str());
+#endif
+}
+
+int spawner(int argc, char *argv[]) try {
+  std::cout << "Running spanwer..." << std::endl;
   assert(argc == 1);
+  configure_runtime_diagnostics_env();
   sycl::queue Q;
+  print_runtime_diagnostics("spawner", Q);
 
 #if defined(__linux__)
   // UMF currently requires ptrace permissions to be set for the spawner. As
@@ -60,16 +161,15 @@ int spawner(int argc, char *argv[]) {
       syclexp::ipc_memory::handle_data_t HandleData = Handle.data();
 #endif
       size_t HandleDataSize = HandleData.size();
+      std::cout << "Spawner handle size: " << HandleDataSize << std::endl;
       std::fstream FS(CommsFile, std::ios_base::out | std::ios_base::binary);
       FS.write(reinterpret_cast<const char *>(&HandleDataSize), sizeof(size_t));
       FS.write(reinterpret_cast<const char *>(HandleData.data()),
                HandleDataSize);
     }
 
     // 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,15 +184,22 @@ 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;
+  configure_runtime_diagnostics_env();
   sycl::queue Q;
+  print_runtime_diagnostics("consumer", Q);
 
   // Read the handle data.
   std::fstream FS(CommsFile, std::ios_base::in | std::ios_base::binary);
   size_t HandleSize = 0;
   FS.read(reinterpret_cast<char *>(&HandleSize), sizeof(size_t));
+  std::cout << "Consumer handle size: " << HandleSize << std::endl;
   std::unique_ptr<std::byte[]> HandleData{new std::byte[HandleSize]};
   FS.read(reinterpret_cast<char *>(HandleData.get()), HandleSize);
 
@@ -105,6 +212,8 @@ int consumer() {
 #endif
   int *DataPtr = reinterpret_cast<int *>(
       syclexp::ipc_memory::open(Handle, Q.get_context(), Q.get_device()));
+  std::cout << "Consumer open succeeded: " << static_cast<void *>(DataPtr)
+            << std::endl;
 
   // Test the data already in the USM pointer.
   int Failures = 0;
@@ -124,8 +233,12 @@ int consumer() {
 
   // Close the IPC pointer.
   syclexp::ipc_memory::close(DataPtr, Q.get_context());
+  std::cout << "Consumer close succeeded" << std::endl;
 
   return Failures;
+} catch (sycl::exception &e) {
+  std::cout << "Consumer failed: " << e.what() << std::endl;
+  throw;
 }
 
 int main(int argc, char *argv[]) {

@@ -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
 

@@ -145,6 +145,10 @@ Environment Variables
 | UR_L0_DISABLE_USM_ALLOCATOR                 | Controls the use of the USM allocator.                       | "0": USM allocator is enabled.                               | "0"              |
 |                                             |                                                              | Any other value: USM allocator is disabled.                  |                  |
 +---------------------------------------------+--------------------------------------------------------------+--------------------------------------------------------------+------------------+
+| UR_L0_V2_ENABLE_WINDOWS_IPC_WA              | Enables the Windows UMF IPC workaround during Level Zero v2  | "0" or unset: The workaround is disabled.                  | "0"              |
+|                                             | provider creation. This must be enabled before provider/pool | "1": The workaround is enabled.                            |                  |
+|                                             | creation so UMF import/export IPC support is configured.     |                                                              |                  |
++---------------------------------------------+--------------------------------------------------------------+--------------------------------------------------------------+------------------+
 | UR_L0_CMD_BUFFER_USE_IMMEDIATE_APPEND_PATH  | Controls which command-buffer implementation path is used.   | "1": the immediate append path will always be enabled as     | Unset            |
 |                                             | The paths rely on different APIs to enqueue command-buffers. | long as the pre-requisites are met.                          |                  |
 |                                             | The immediate append path relies on                          | "0": the immediate append path will always be disabled.      |                  |

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

@@ -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:

@@ -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));