Cycles: Add Vulkan/oneAPI graphics interop

This PR adds Vulkan/oneAPI graphics interop to Cycles. Just like for
CUDA and HIP interop, persistent memory mapping is used, as there could
potentially be some overhead of continuously mapping/unmapping buffers.

Pull Request: https://projects.blender.org/blender/blender/pulls/144442

intern/cycles/device/CMakeLists.txt

@@ -77,6 +77,8 @@ set(SRC_ONEAPI
   oneapi/device_impl.h
   oneapi/device.cpp
   oneapi/device.h
+  oneapi/graphics_interop.cpp
+  oneapi/graphics_interop.h
   oneapi/queue.cpp
   oneapi/queue.h
 )
@@ -195,6 +197,33 @@ if(WITH_CYCLES_DEVICE_ONEAPI)
   list(APPEND INC_SYS
     ${SYCL_INCLUDE_DIR}
   )
+
+  # Test for the presence of sycl::ext::oneapi::experimental::unmap_external_linear_memory (necessary for interop).
+  # https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc
+  include(CheckCXXSourceCompiles)
+  set(CMAKE_REQUIRED_INCLUDES "${SYCL_INCLUDE_DIR}")
+  set(CMAKE_REQUIRED_LIBRARIES "${SYCL_LIBRARIES}")
+  check_cxx_source_compiles("
+      #include <sycl/ext/oneapi/bindless_images.hpp>
+      int main()
+      {
+          sycl::queue sycl_queue{sycl::gpu_selector_v};
+          sycl::ext::oneapi::experimental::unmap_external_linear_memory(nullptr, sycl_queue);
+      }
+  " SYCL_UNMAP_EXTERNAL_LINEAR_MEMORY_SUPPORTED)
+  if (SYCL_UNMAP_EXTERNAL_LINEAR_MEMORY_SUPPORTED)
+    foreach(FILE ${SRC_ONEAPI})
+      set_source_files_properties(
+          ${FILE} PROPERTIES COMPILE_DEFINITIONS "SYCL_LINEAR_MEMORY_INTEROP_AVAILABLE"
+      )
+    endforeach()
+  else()
+    message(WARNING
+      "The installed SYCL version does not support unmap_external_linear_memory. "
+      "Upgrade to oneAPI >= 2025.3 or DPC++ >= 6.2 to support Vulkan-oneAPI interoperability.")
+  endif()
+  unset(CMAKE_REQUIRED_INCLUDES)
+  unset(CMAKE_REQUIRED_LIBRARIES)
 endif()
 
 if(WITH_OPENIMAGEDENOISE)

intern/cycles/device/oneapi/device_impl.cpp

@@ -1,4 +1,4 @@
-/* SPDX-FileCopyrightText: 2021-2022 Intel Corporation
+/* SPDX-FileCopyrightText: 2021-2025 Intel Corporation
  *
  * SPDX-License-Identifier: Apache-2.0 */
 
@@ -26,6 +26,8 @@
 #  include "kernel/device/oneapi/globals.h"
 #  include "kernel/device/oneapi/kernel.h"
 
+#  include "session/display_driver.h"
+
 #  if defined(WITH_EMBREE_GPU) && defined(EMBREE_SYCL_SUPPORT) && !defined(SYCL_LANGUAGE_VERSION)
 /* These declarations are missing from embree headers when compiling from a compiler that doesn't
  * support SYCL. */
@@ -945,11 +947,46 @@ unique_ptr<DeviceQueue> OneapiDevice::gpu_queue_create()
   return make_unique<OneapiDeviceQueue>(this);
 }
 
-bool OneapiDevice::should_use_graphics_interop(const GraphicsInteropDevice & /*interop_device*/,
+bool OneapiDevice::should_use_graphics_interop(const GraphicsInteropDevice &interop_device,
-                                               const bool /*log*/)
+                                               const bool log)
 {
-  /* NOTE(@nsirgien): oneAPI doesn't yet support direct writing into graphics API objects, so
+#  ifdef SYCL_LINEAR_MEMORY_INTEROP_AVAILABLE
-   * return false. */
+  if (interop_device.type != GraphicsInteropDevice::VULKAN) {
+    /* SYCL only supports interop with Vulkan and D3D. */
+    return false;
+  }
+
+  try {
+    const sycl::device &device = reinterpret_cast<sycl::queue *>(device_queue_)->get_device();
+    if (!device.has(sycl::aspect::ext_oneapi_external_memory_import)) {
+      return false;
+    }
+
+    /* This extension is in the namespace "sycl::ext::intel",
+     * but also available on non-Intel GPUs. */
+    sycl::detail::uuid_type uuid = device.get_info<sycl::ext::intel::info::device::uuid>();
+    const bool found = (uuid.size() == interop_device.uuid.size() &&
+                        memcmp(uuid.data(), interop_device.uuid.data(), uuid.size()) == 0);
+
+    if (log) {
+      if (found) {
+        LOG_INFO << "Graphics interop: found matching Vulkan device for oneAPI";
+      }
+      else {
+        LOG_INFO << "Graphics interop: no matching Vulkan device for oneAPI";
+      }
+
+      LOG_INFO << "Graphics Interop: oneAPI UUID " << string_hex(uuid.data(), uuid.size())
+               << ", Vulkan UUID "
+               << string_hex(interop_device.uuid.data(), interop_device.uuid.size());
+    }
+
+    return found;
+  }
+  catch (sycl::exception &e) {
+    LOG_ERROR << "Could not release external Vulkan memory: " << e.what();
+  }
+#  endif
   return false;
 }
 

intern/cycles/device/oneapi/device_impl.h

@@ -1,4 +1,4 @@
-/* SPDX-FileCopyrightText: 2021-2022 Intel Corporation
+/* SPDX-FileCopyrightText: 2021-2025 Intel Corporation
  *
  * SPDX-License-Identifier: Apache-2.0 */
 

intern/cycles/device/oneapi/graphics_interop.cpp

@@ -0,0 +1,168 @@
+/* SPDX-FileCopyrightText: 2025 Intel Corporation
+ *
+ * SPDX-License-Identifier: Apache-2.0 */
+
+#if defined(WITH_ONEAPI) && defined(SYCL_LINEAR_MEMORY_INTEROP_AVAILABLE)
+
+#  include "device/oneapi/graphics_interop.h"
+
+#  include "device/oneapi/device.h"
+#  include "device/oneapi/device_impl.h"
+#  include "device/oneapi/queue.h"
+
+#  include "session/display_driver.h"
+
+#  ifdef _WIN32
+#    include "util/windows.h"
+#  else
+#    include <unistd.h>
+#  endif
+
+CCL_NAMESPACE_BEGIN
+
+OneapiDeviceGraphicsInterop::OneapiDeviceGraphicsInterop(OneapiDeviceQueue *queue)
+    : queue_(queue), device_(static_cast<OneapiDevice *>(queue->device))
+{
+}
+
+OneapiDeviceGraphicsInterop::~OneapiDeviceGraphicsInterop()
+{
+  free();
+}
+
+void OneapiDeviceGraphicsInterop::set_buffer(GraphicsInteropBuffer &interop_buffer)
+{
+  if (interop_buffer.is_empty()) {
+    free();
+    return;
+  }
+
+  need_zero_ |= interop_buffer.take_zero();
+
+  if (!interop_buffer.has_new_handle()) {
+    return;
+  }
+
+  free();
+
+  if (interop_buffer.get_type() != GraphicsInteropDevice::VULKAN) {
+    /* SYCL only supports interop with Vulkan and D3D. */
+    LOG_ERROR
+        << "oneAPI interop set_buffer called for invalid graphics API. Only Vulkan is supported.";
+    return;
+  }
+
+#  ifdef _WIN32
+  /* import_external_memory will not take ownership of the handle. */
+  vulkan_windows_handle_ = reinterpret_cast<void *>(interop_buffer.take_handle());
+  auto sycl_mem_handle_type =
+      sycl::ext::oneapi::experimental::external_mem_handle_type::win32_nt_handle;
+  sycl::ext::oneapi::experimental::external_mem_descriptor<
+      sycl::ext::oneapi::experimental::resource_win32_handle>
+      sycl_external_mem_descriptor{vulkan_windows_handle_, sycl_mem_handle_type};
+#  else
+  /* import_external_memory will take ownership of the file descriptor. */
+  auto sycl_mem_handle_type = sycl::ext::oneapi::experimental::external_mem_handle_type::opaque_fd;
+  sycl::ext::oneapi::experimental::external_mem_descriptor<
+      sycl::ext::oneapi::experimental::resource_fd>
+      sycl_external_mem_descriptor{static_cast<int>(interop_buffer.take_handle()),
+                                   sycl_mem_handle_type};
+#  endif
+
+  sycl::queue *sycl_queue = reinterpret_cast<sycl::queue *>(device_->sycl_queue());
+  try {
+    sycl_external_memory_ = sycl::ext::oneapi::experimental::import_external_memory(
+        sycl_external_mem_descriptor, *sycl_queue);
+  }
+  catch (sycl::exception &e) {
+#  ifdef _WIN32
+    CloseHandle(HANDLE(vulkan_windows_handle_));
+    vulkan_windows_handle_ = nullptr;
+#  else
+    close(sycl_external_mem_descriptor.external_resource.file_descriptor);
+#  endif
+    LOG_ERROR << "Error importing Vulkan memory: " << e.what();
+    return;
+  }
+
+  buffer_size_ = interop_buffer.get_size();
+
+  /* Like the CUDA/HIP backend, we map the buffer persistently. */
+  try {
+    sycl_memory_ptr_ = sycl::ext::oneapi::experimental::map_external_linear_memory(
+        sycl_external_memory_, 0, buffer_size_, *sycl_queue);
+  }
+  catch (sycl::exception &e) {
+    try {
+      sycl::ext::oneapi::experimental::release_external_memory(sycl_external_memory_, *sycl_queue);
+    }
+    catch (sycl::exception &e) {
+      LOG_ERROR << "Could not release external Vulkan memory: " << e.what();
+    }
+    sycl_external_memory_ = {};
+    buffer_size_ = 0;
+    /* Only need to close Windows handle, as file descriptor is owned by compute API. */
+#  ifdef _WIN32
+    CloseHandle(HANDLE(vulkan_windows_handle_));
+    vulkan_windows_handle_ = nullptr;
+#  endif
+    LOG_ERROR << "Error mapping external Vulkan memory: " << e.what();
+    return;
+  }
+}
+
+device_ptr OneapiDeviceGraphicsInterop::map()
+{
+  if (sycl_memory_ptr_ && need_zero_) {
+    try {
+      /* We do not wait on the returned event here, as CUDA also uses "cuMemsetD8Async". */
+      sycl::queue *sycl_queue = reinterpret_cast<sycl::queue *>(device_->sycl_queue());
+      sycl_queue->memset(sycl_memory_ptr_, 0, buffer_size_);
+    }
+    catch (sycl::exception &e) {
+      LOG_ERROR << "Error clearing external Vulkan memory: " << e.what();
+      return device_ptr(0);
+    }
+    need_zero_ = false;
+  }
+
+  return reinterpret_cast<device_ptr>(sycl_memory_ptr_);
+}
+
+void OneapiDeviceGraphicsInterop::unmap() {}
+
+void OneapiDeviceGraphicsInterop::free()
+{
+  if (sycl_external_memory_.raw_handle) {
+    sycl::queue *sycl_queue = reinterpret_cast<sycl::queue *>(device_->sycl_queue());
+    try {
+      sycl::ext::oneapi::experimental::unmap_external_linear_memory(sycl_memory_ptr_, *sycl_queue);
+    }
+    catch (sycl::exception &e) {
+      LOG_ERROR << "Could not unmap external Vulkan memory: " << e.what();
+    }
+    try {
+      sycl::ext::oneapi::experimental::release_external_memory(sycl_external_memory_, *sycl_queue);
+    }
+    catch (sycl::exception &e) {
+      LOG_ERROR << "Could not release external Vulkan memory: " << e.what();
+    }
+    sycl_memory_ptr_ = {};
+    sycl_external_memory_ = {};
+  }
+
+#  ifdef _WIN32
+  if (vulkan_windows_handle_) {
+    CloseHandle(HANDLE(vulkan_windows_handle_));
+    vulkan_windows_handle_ = nullptr;
+  }
+#  endif
+
+  buffer_size_ = 0;
+
+  need_zero_ = false;
+}
+
+CCL_NAMESPACE_END
+
+#endif

intern/cycles/device/oneapi/graphics_interop.h

@@ -0,0 +1,61 @@
+/* SPDX-FileCopyrightText: 2025 Intel Corporation
+ *
+ * SPDX-License-Identifier: Apache-2.0 */
+
+#if defined(WITH_ONEAPI) && defined(SYCL_LINEAR_MEMORY_INTEROP_AVAILABLE)
+
+#  include <sycl/sycl.hpp>
+
+#  include "device/graphics_interop.h"
+#  include "session/display_driver.h"
+
+#  include "device/oneapi/device.h"
+#  include "device/oneapi/queue.h"
+
+CCL_NAMESPACE_BEGIN
+
+class OneapiDevice;
+class OneapiDeviceQueue;
+
+class OneapiDeviceGraphicsInterop : public DeviceGraphicsInterop {
+ public:
+  explicit OneapiDeviceGraphicsInterop(OneapiDeviceQueue *queue);
+
+  OneapiDeviceGraphicsInterop(const OneapiDeviceGraphicsInterop &other) = delete;
+  OneapiDeviceGraphicsInterop(OneapiDeviceGraphicsInterop &&other) noexcept = delete;
+
+  ~OneapiDeviceGraphicsInterop() override;
+
+  OneapiDeviceGraphicsInterop &operator=(const OneapiDeviceGraphicsInterop &other) = delete;
+  OneapiDeviceGraphicsInterop &operator=(OneapiDeviceGraphicsInterop &&other) = delete;
+
+  void set_buffer(GraphicsInteropBuffer &interop_buffer) override;
+
+  device_ptr map() override;
+  void unmap() override;
+
+ protected:
+  OneapiDeviceQueue *queue_ = nullptr;
+  OneapiDevice *device_ = nullptr;
+
+  /* Size of the buffer in bytes. */
+  size_t buffer_size_ = 0;
+
+  /* The destination was requested to be cleared. */
+  bool need_zero_ = false;
+
+  /* Oneapi resources. */
+  sycl::ext::oneapi::experimental::external_mem sycl_external_memory_{};
+  void *sycl_memory_ptr_ = nullptr;
+
+  /* Vulkan handle to free. */
+#  ifdef _WIN32
+  void *vulkan_windows_handle_ = nullptr;
+#  endif
+
+  void free();
+};
+
+CCL_NAMESPACE_END
+
+#endif

intern/cycles/device/oneapi/queue.cpp

@@ -1,4 +1,4 @@
-/* SPDX-FileCopyrightText: 2021-2022 Intel Corporation
+/* SPDX-FileCopyrightText: 2021-2025 Intel Corporation
  *
  * SPDX-License-Identifier: Apache-2.0 */
 
@@ -6,6 +6,7 @@
 
 #  include "device/oneapi/queue.h"
 #  include "device/oneapi/device_impl.h"
+#  include "device/oneapi/graphics_interop.h"
 #  include "util/log.h"
 
 #  include "kernel/device/oneapi/kernel.h"
@@ -142,6 +143,13 @@ void OneapiDeviceQueue::copy_from_device(device_memory &mem)
   oneapi_device_->mem_copy_from(mem);
 }
 
+#  ifdef SYCL_LINEAR_MEMORY_INTEROP_AVAILABLE
+unique_ptr<DeviceGraphicsInterop> OneapiDeviceQueue::graphics_interop_create()
+{
+  return make_unique<OneapiDeviceGraphicsInterop>(this);
+}
+#  endif
+
 CCL_NAMESPACE_END
 
 #endif /* WITH_ONEAPI */

intern/cycles/device/oneapi/queue.h

@@ -1,4 +1,4 @@
-/* SPDX-FileCopyrightText: 2021-2022 Intel Corporation
+/* SPDX-FileCopyrightText: 2021-2025 Intel Corporation
  *
  * SPDX-License-Identifier: Apache-2.0 */
 
@@ -46,6 +46,10 @@ class OneapiDeviceQueue : public DeviceQueue {
     return true;
   }
 
+#  ifdef SYCL_LINEAR_MEMORY_INTEROP_AVAILABLE
+  unique_ptr<DeviceGraphicsInterop> graphics_interop_create() override;
+#  endif
+
  protected:
   OneapiDevice *oneapi_device_;
   unique_ptr<KernelContext> kernel_context_;