diff --git a/intern/cycles/device/CMakeLists.txt b/intern/cycles/device/CMakeLists.txt index 70461e029c1..9e3dfeb9792 100644 --- a/intern/cycles/device/CMakeLists.txt +++ b/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 + 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) diff --git a/intern/cycles/device/oneapi/device_impl.cpp b/intern/cycles/device/oneapi/device_impl.cpp index fe1dad306e1..98411a62c2e 100644 --- a/intern/cycles/device/oneapi/device_impl.cpp +++ b/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 OneapiDevice::gpu_queue_create() return make_unique(this); } -bool OneapiDevice::should_use_graphics_interop(const GraphicsInteropDevice & /*interop_device*/, - const bool /*log*/) +bool OneapiDevice::should_use_graphics_interop(const GraphicsInteropDevice &interop_device, + const bool log) { - /* NOTE(@nsirgien): oneAPI doesn't yet support direct writing into graphics API objects, so - * return false. */ +# ifdef SYCL_LINEAR_MEMORY_INTEROP_AVAILABLE + if (interop_device.type != GraphicsInteropDevice::VULKAN) { + /* SYCL only supports interop with Vulkan and D3D. */ + return false; + } + + try { + const sycl::device &device = reinterpret_cast(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(); + 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; } diff --git a/intern/cycles/device/oneapi/device_impl.h b/intern/cycles/device/oneapi/device_impl.h index 046eac68e32..d2aafb4ceca 100644 --- a/intern/cycles/device/oneapi/device_impl.h +++ b/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 */ diff --git a/intern/cycles/device/oneapi/graphics_interop.cpp b/intern/cycles/device/oneapi/graphics_interop.cpp new file mode 100644 index 00000000000..0132faa4041 --- /dev/null +++ b/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 +# endif + +CCL_NAMESPACE_BEGIN + +OneapiDeviceGraphicsInterop::OneapiDeviceGraphicsInterop(OneapiDeviceQueue *queue) + : queue_(queue), device_(static_cast(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(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(interop_buffer.take_handle()), + sycl_mem_handle_type}; +# endif + + sycl::queue *sycl_queue = reinterpret_cast(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(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(sycl_memory_ptr_); +} + +void OneapiDeviceGraphicsInterop::unmap() {} + +void OneapiDeviceGraphicsInterop::free() +{ + if (sycl_external_memory_.raw_handle) { + sycl::queue *sycl_queue = reinterpret_cast(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 diff --git a/intern/cycles/device/oneapi/graphics_interop.h b/intern/cycles/device/oneapi/graphics_interop.h new file mode 100644 index 00000000000..f3234bfee5d --- /dev/null +++ b/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 + +# 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 diff --git a/intern/cycles/device/oneapi/queue.cpp b/intern/cycles/device/oneapi/queue.cpp index 671d02645c8..5e2712866f1 100644 --- a/intern/cycles/device/oneapi/queue.cpp +++ b/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 OneapiDeviceQueue::graphics_interop_create() +{ + return make_unique(this); +} +# endif + CCL_NAMESPACE_END #endif /* WITH_ONEAPI */ diff --git a/intern/cycles/device/oneapi/queue.h b/intern/cycles/device/oneapi/queue.h index 949a3c8a86d..f4a9b1410f4 100644 --- a/intern/cycles/device/oneapi/queue.h +++ b/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 graphics_interop_create() override; +# endif + protected: OneapiDevice *oneapi_device_; unique_ptr kernel_context_;