diff --git a/src/inference/include/openvino/runtime/intel_gpu/remote_properties.hpp b/src/inference/include/openvino/runtime/intel_gpu/remote_properties.hpp index 977580aef39ad2..1c35990cfefbaa 100644 --- a/src/inference/include/openvino/runtime/intel_gpu/remote_properties.hpp +++ b/src/inference/include/openvino/runtime/intel_gpu/remote_properties.hpp @@ -30,7 +30,7 @@ using os_handle_param = void*; enum class ContextType { OCL = 0, //!< Pure OpenCL context VA_SHARED = 1, //!< Context shared with a video decoding device - ZE = 2, //!< Pure Level0 context + ZE = 2, //!< Pure Level Zero context }; /** @cond INTERNAL */ @@ -52,10 +52,10 @@ inline std::istream& operator>>(std::istream& is, ContextType& context_type) { is >> str; if (str == "OCL") { context_type = ContextType::OCL; - } else if (str == "ZE") { - context_type = ContextType::ZE; } else if (str == "VA_SHARED") { context_type = ContextType::VA_SHARED; + } else if (str == "ZE") { + context_type = ContextType::ZE; } else { OPENVINO_THROW("Unsupported context type: ", str); } diff --git a/src/plugins/intel_gpu/docs/blog/ze_ocl_interop_2026_3.md b/src/plugins/intel_gpu/docs/blog/ze_ocl_interop_2026_3.md new file mode 100644 index 00000000000000..14f2a1a529e600 --- /dev/null +++ b/src/plugins/intel_gpu/docs/blog/ze_ocl_interop_2026_3.md @@ -0,0 +1,50 @@ +# Support for Level Zero - OpenCL interoperability: New in OpenVINO 2026.3 + +*By Jakub Kasprzak | June 23, 2026* + +Some projects maintain their own OpenCL GPU pipeline and require features such as memory, context or queue sharing. For such users OpenVINO GPU plugin implements Remote Context and Remote Tensor interface. However, this OpenCL interoperability works only if OpenVINO GPU runtime is set to OpenCL and prevents OpenVINO from switching to Level Zero runtime. +OpenVINO 2026.3 introduces Level Zero - OpenCL interoperability enabling users that depend on OpenVINO Remote Context/Tensor OpenCL API to swap OpenVINO GPU runtime to Level Zero without changing their OpenCL based integration logic. This post explains: how to enable this feature and provides OpenVINO maintainers with design description. + +## How to enable Level Zero - OpenCL interoperability + +OpenVINO GPU plugin with Level Zero runtime detects if interoperability is supported on current system and enables it automatically. User must ensure that currently installed Intel GPU driver supports and enables "Level Zero executing OpenCL" (LEO) feature. + +## Level Zero - OpenCL interoperability design + +*This section describes internal design of the feature and is intended for OpenVINO maintainers* + +### Object lifetimes + +This subsection describes part of the design responsible for managing resources. Main goal was to provide modular API that will minimize chances of memory leaks and other memory bugs. + +Logic for managing all (Level Zero and OpenCL) GPU resources is provided by `resource_owner` template class. It contains `is_borrowed` flag to distinguish between resources that need to be released and resources that should not be released (borrowed resources, for example: external memory buffers provided by the user). Once `resource_owner` is destroyed it will release the resource or do nothing if resource was borrowed. Basically `resource_owner` with `is_borrowed=false` is a generic RAII wrapper which means that only one such object can exist for a given resource and that it can not be copied (only move semantics are allowed). + +In order to work correctly `resource_owner` template must be provided with information about the type of resource it manages and how to release it. This information is defined for Level Zero and OpenCL resources in `ze_resource_info` and `ocl_resource_info` class template specializations respectively. For example `ocl_resource_info` defines `handle_t` as `cl_command_queue` and `deleter_t` as a functor that calls `clReleaseCommandQueue`. Finally `ze_owner` and `ocl_owner` aliases are defined by combining `resource_owner` with `ze_resource_info` and `ocl_resource_info` respectively. + +### Binding corresponding Level Zero - OpenCL handles + +It is important to recognize that with interoperability some Level Zero handles will have corresponding OpenCL handles and vice versa. For example, after converting Level Zero immediate command list to OpenCL command queue we end up with two different handles referring to the same entity. `ze_ocl_owner_impl` is a template that is responsible for combining Level Zero handle with OpenCL handle and managing them as a single resource. Internally it contains `ze_owner` and tuple of optional `ocl_owner`s. By design `ze_ocl_owner_impl` always contains Level Zero resource and OpenCL resources can be attached to it. Once OpenCL resource of some given type is attached to `ze_ocl_owner_impl` attempt to attach another OpenCL resource of the same type will result in **exception** (this design is to prevent exporting single Level Zero resource multiple times). + +Not all handle combinations are valid. For example in current context it does not make sense to attach/bind `cl_command_queue` handle to `ze_ocl_owner_impl` that holds `ze_image_handle_t`. To prevent this `ze_ocl_owner` template inherits only valid `ze_ocl_owner_impl` specializations. For example `ze_ocl_owner` is `ze_ocl_owner_impl` meaning it holds `ze_command_list_handle_t` and optionally `cl_command_queue` handles. If specialization is not defined explicitly it fall backs to default implementation where only Level Zero handle is allowed, for example `ze_ocl_owner` holds only `ze_kernel_handle_t` and attempt to attach or access any type of OpenCL handle will result in **compilation error**. + +### Shared resource ownership + +Finally `ze_resource` template combines `std::shared_ptr` and `ze_ocl_owner` to implement shared ownership model for GPU resources. This is especially useful for handling lifetime of Level Zero events and event pools where `ze_resource` can be copied to all event objects created from given pool ensuring that the pool resource will be destroyed only after last event from the pool was destroyed (same scheme is used to manage lifetime of Level Zero modules and kernels). + +Additionaly aliases in form of `ze_*_resource` such as `ze_event_pool_resource` are defined for developer convenience. + +### Interoperability integration + +Basic interoperability is implemented by a `ze_ocl_interop` singleton that detects interoperability support and provides direct (no abstraction) Level Zero - OpenCL handle to handle conversion interface. When `ze_ocl_interop` is accessed for the first time it will initialize GPU runtimes and map corresponding Level Zero and OpenCL devices. + +Conversion interface provided by `ze_ocl_interop` is wrapped by `ze_import_*` and `ze_export_ocl` functions from `ze_resource_interop.hpp` in order to simplify conversion logic when used with `ze_resource` objects. For example calling `ze_export_ocl_command_queue` for given `ze_command_list_resource` will: +1. Return early if `ze_command_list_resource` already has attached OpenCL command queue handle +2. Create new OpenCL command queue based on present Level Zero handle +3. Attach new OpenCL handle to `ze_command_list_resource` +4. Guarantee strong exception safety + +### Remote context + +Type of remote context decides what kind of objects are exposed by the GPU plugin. During model compilation OpenVINO creates remote context and internal engine that by default match currently selected GPU runtime. For example, when GPU runtime is set to Level Zero this causes the creation of Level Zero remote context and `ze_engine`. Remote context will query internal GPU plugin objects for handles based on the remote context type. So if Level Zero remote context is used only Level Zero handles can be requested from internal objects and OpenCL interoperability is not required. + +If user provides OpenCL context or queue or OpenVINO detects that current system supports Level Zero - OpenCL interoperability then remote context type will be changed to OpenCL but engine will still match currently selected GPU runtime. In result OpenCL remote context will request OpenCL handles from internal Level Zero GPU plugin objects like `ze_engine` forcing usage of interoperability features. diff --git a/src/plugins/intel_gpu/include/intel_gpu/plugin/remote_context.hpp b/src/plugins/intel_gpu/include/intel_gpu/plugin/remote_context.hpp index 15956e73683659..f0b7cd2e40e965 100644 --- a/src/plugins/intel_gpu/include/intel_gpu/plugin/remote_context.hpp +++ b/src/plugins/intel_gpu/include/intel_gpu/plugin/remote_context.hpp @@ -95,11 +95,8 @@ class RemoteContextImpl : public ov::IRemoteContext { ov::intel_gpu::gpu_handle_param m_va_display = nullptr; ov::intel_gpu::gpu_handle_param m_external_queue = nullptr; -#ifdef OV_GPU_WITH_ZE_RT - ContextType m_type = ContextType::ZE; -#else - ContextType m_type = ContextType::OCL; -#endif + + ContextType m_type; std::string m_device_name = ""; static const size_t cache_capacity = 100; cldnn::LruCache m_memory_cache = cldnn::LruCache(cache_capacity); diff --git a/src/plugins/intel_gpu/include/intel_gpu/runtime/device_info.hpp b/src/plugins/intel_gpu/include/intel_gpu/runtime/device_info.hpp index bef507037b0243..1a0b6014c65e68 100644 --- a/src/plugins/intel_gpu/include/intel_gpu/runtime/device_info.hpp +++ b/src/plugins/intel_gpu/include/intel_gpu/runtime/device_info.hpp @@ -122,8 +122,9 @@ struct device_info { bool supports_usm; ///< Does engine support unified shared memory. bool has_separate_cache; ///< Does the target hardware has separate cache for usm_device and usm_host - bool supports_cp_offload; ///< [ZE] Does the command queue support copy offload - bool supports_counter_based_events; ///< [ZE] Does the target runtime support counter based events + bool supports_cp_offload; ///< [L0] Does the command queue support copy offload + bool supports_counter_based_events; ///< [L0] Does the target runtime support counter based events + bool supports_leo; ///< [L0] Does the device support Level Zero - OpenCL interoperability (LEO) std::vector supported_simd_sizes; ///< List of SIMD sizes supported by current device and compiler diff --git a/src/plugins/intel_gpu/include/intel_gpu/runtime/device_query.hpp b/src/plugins/intel_gpu/include/intel_gpu/runtime/device_query.hpp index 1ff25f67065354..760b7daf3171ad 100644 --- a/src/plugins/intel_gpu/include/intel_gpu/runtime/device_query.hpp +++ b/src/plugins/intel_gpu/include/intel_gpu/runtime/device_query.hpp @@ -17,11 +17,6 @@ namespace cldnn { struct device_query { public: static int device_id; - /// @brief Get default engine type - static engine_types get_default_engine_type(); - - /// @brief Get default runtime type - static runtime_types get_default_runtime_type(); explicit device_query(engine_types engine_type, runtime_types runtime_type, diff --git a/src/plugins/intel_gpu/include/intel_gpu/runtime/engine.hpp b/src/plugins/intel_gpu/include/intel_gpu/runtime/engine.hpp index 374586c1d222b5..ebd5c51ebbdf95 100644 --- a/src/plugins/intel_gpu/include/intel_gpu/runtime/engine.hpp +++ b/src/plugins/intel_gpu/include/intel_gpu/runtime/engine.hpp @@ -113,8 +113,8 @@ class engine { /// Returns device object associated with the engine const device::ptr get_device() const; - /// Returns user context handle which was used to create the engine - virtual void* get_user_context() const = 0; + /// Returns L0 or OpenCL user context handle which was used to create the engine. + virtual void* get_user_context(runtime_types rt_type) const = 0; /// Returns the total maximum amount of GPU memory allocated by engine in current process for all allocation types uint64_t get_max_used_device_memory() const; diff --git a/src/plugins/intel_gpu/include/intel_gpu/runtime/engine_configuration.hpp b/src/plugins/intel_gpu/include/intel_gpu/runtime/engine_configuration.hpp index 31c28145e5783d..8854c9f4061f32 100644 --- a/src/plugins/intel_gpu/include/intel_gpu/runtime/engine_configuration.hpp +++ b/src/plugins/intel_gpu/include/intel_gpu/runtime/engine_configuration.hpp @@ -64,4 +64,10 @@ inline std::ostream& operator<<(std::ostream& os, const backend_types& type) { return os; } +/// @brief Get default engine type +engine_types get_default_engine_type(); + +/// @brief Get default runtime type +runtime_types get_default_runtime_type(); + } // namespace cldnn diff --git a/src/plugins/intel_gpu/include/intel_gpu/runtime/memory.hpp b/src/plugins/intel_gpu/include/intel_gpu/runtime/memory.hpp index d3bd2a6f035aca..44ec4773c15d37 100644 --- a/src/plugins/intel_gpu/include/intel_gpu/runtime/memory.hpp +++ b/src/plugins/intel_gpu/include/intel_gpu/runtime/memory.hpp @@ -58,7 +58,11 @@ struct memory { size_t size() const { return _bytes_count; } size_t count() const { return _layout.count(); } - virtual shared_mem_params get_internal_params() const = 0; + /// @brief Return internal params for specified runtime type + virtual shared_mem_params get_internal_params(runtime_types rt_type) const = 0; + shared_mem_params get_internal_params() const { + return get_internal_params(get_default_runtime_type()); + } virtual bool is_allocated_by(const engine& engine) const { return &engine == _engine && _type != allocation_type::unknown; } engine* get_engine() const { return _engine; } const layout& get_layout() const { return _layout; } @@ -156,7 +160,7 @@ struct simple_attached_memory : memory { void unlock(const stream& /* stream */) override {} event::ptr fill(stream& /* stream */, unsigned char, const std::vector&, bool) override { return nullptr; } event::ptr fill(stream& /* stream */, const std::vector&, bool) override { return nullptr; } - shared_mem_params get_internal_params() const override { return { shared_mem_type::shared_mem_empty, nullptr, nullptr, nullptr, + shared_mem_params get_internal_params(runtime_types rt_type) const override { return { shared_mem_type::shared_mem_empty, nullptr, nullptr, nullptr, #ifdef _WIN32 nullptr, #else diff --git a/src/plugins/intel_gpu/include/intel_gpu/runtime/memory_caps.hpp b/src/plugins/intel_gpu/include/intel_gpu/runtime/memory_caps.hpp index c41bdd383fc7f7..047c9a66d5f949 100644 --- a/src/plugins/intel_gpu/include/intel_gpu/runtime/memory_caps.hpp +++ b/src/plugins/intel_gpu/include/intel_gpu/runtime/memory_caps.hpp @@ -110,7 +110,7 @@ using shared_surface = uint32_t; /// @brief Low-level API handles required for using cldnn memory objects in external API calls. struct shared_mem_params { shared_mem_type mem_type; ///< shared buffer type - shared_handle context; ///< OpenCL context for external operations + shared_handle context; ///< OpenCL or Level Zero context for external operations shared_handle user_device; ///< DX/VA device for external operations shared_handle mem; ///< memory object handle #ifdef _WIN32 diff --git a/src/plugins/intel_gpu/src/plugin/remote_context.cpp b/src/plugins/intel_gpu/src/plugin/remote_context.cpp index f932abe14eb2ce..86a121335b5539 100644 --- a/src/plugins/intel_gpu/src/plugin/remote_context.cpp +++ b/src/plugins/intel_gpu/src/plugin/remote_context.cpp @@ -23,12 +23,29 @@ Type extract_object(const ov::AnyMap& params, const ov::Property& p) { return res.as(); } +ContextType get_default_context_type() { + #ifdef OV_GPU_WITH_ZE_RT + return ContextType::ZE; + #elif defined(OV_GPU_WITH_OCL_RT) + return ContextType::OCL; + #else + #error "Expected OpenVINO GPU runtime macros to be defined" + #endif +} + } // namespace RemoteContextImpl::RemoteContextImpl(const std::string& device_name, std::vector devices, bool initialize_ctx) : m_device_name(device_name) { OPENVINO_ASSERT(devices.size() == 1, "[GPU] Currently context can be created for single device only"); m_device = devices.front(); + m_type = get_default_context_type(); + const auto &info = m_device->get_info(); + if (m_type == ContextType::ZE && info.supports_leo) { + m_type = ContextType::OCL; + GPU_DEBUG_INFO << "Enabled Level Zero - OpenCL interoperability for " + << m_device_name << " (" << info.dev_name << ")" << std::endl; + } if (initialize_ctx) { initialize(); @@ -39,6 +56,7 @@ RemoteContextImpl::RemoteContextImpl(const std::mapsecond; @@ -89,19 +110,20 @@ const cldnn::engine& RemoteContextImpl::get_engine() const { } void RemoteContextImpl::init_properties() { - properties = { ov::intel_gpu::ocl_context(m_engine->get_user_context()) }; - switch (m_type) { case ContextType::OCL: properties.insert(ov::intel_gpu::context_type(ov::intel_gpu::ContextType::OCL)); + properties.insert(ov::intel_gpu::ocl_context(m_engine->get_user_context(cldnn::runtime_types::ocl))); properties.insert(ov::intel_gpu::ocl_queue(m_external_queue)); break; case ContextType::VA_SHARED: properties.insert(ov::intel_gpu::context_type(ov::intel_gpu::ContextType::VA_SHARED)); + properties.insert(ov::intel_gpu::ocl_context(m_engine->get_user_context(cldnn::runtime_types::ocl))); properties.insert(ov::intel_gpu::va_device(m_va_display)); break; case ContextType::ZE: properties.insert(ov::intel_gpu::context_type(ov::intel_gpu::ContextType::ZE)); + properties.insert(ov::intel_gpu::ocl_context(m_engine->get_user_context(cldnn::runtime_types::ze))); break; default: OPENVINO_THROW("[GPU] Unsupported shared context type ", m_type); @@ -256,7 +278,7 @@ void RemoteContextImpl::initialize() { m_device->initialize(); // Initialize associated device before use m_engine = cldnn::engine::create( - cldnn::device_query::get_default_engine_type(), cldnn::device_query::get_default_runtime_type(), m_device); + cldnn::get_default_engine_type(), cldnn::get_default_runtime_type(), m_device); init_properties(); diff --git a/src/plugins/intel_gpu/src/plugin/remote_tensor.cpp b/src/plugins/intel_gpu/src/plugin/remote_tensor.cpp index 74ff175b250bad..f9efbd47b5c57c 100644 --- a/src/plugins/intel_gpu/src/plugin/remote_tensor.cpp +++ b/src/plugins/intel_gpu/src/plugin/remote_tensor.cpp @@ -319,7 +319,6 @@ void RemoteTensorImpl::allocate() { switch (m_mem_type) { case TensorType::BT_BUF_INTERNAL: { - // BT_BUF_INTERNAL should map to cl_mem however ZE engine can not allocate cl_mem if (engine.supports_allocation(cldnn::allocation_type::cl_mem)) { m_memory_object = engine.allocate_memory(m_layout, cldnn::allocation_type::cl_mem, reset); } else if (engine.supports_allocation(cldnn::allocation_type::sycl_buffer)) { @@ -450,7 +449,19 @@ std::shared_ptr RemoteTensorImpl::get_context() const { void RemoteTensorImpl::update_properties() { OPENVINO_ASSERT(is_allocated(), "[GPU] Can't initialize RemoteTensorImpl parameters as memory was not allocated"); - auto params = m_memory_object->get_internal_params(); + const auto &ctx_props = m_context->get_property(); + const auto it = ctx_props.find(ov::intel_gpu::context_type.name()); + OPENVINO_ASSERT(it != ctx_props.end(), "[GPU] Could not find context type in RemoteContext properties"); + const auto ctx_type = it->second.as(); + + cldnn::shared_mem_params params; + if (ctx_type == ContextType::OCL || ctx_type == ContextType::VA_SHARED) { + params = m_memory_object->get_internal_params(cldnn::runtime_types::ocl); + } else if (ctx_type == ContextType::ZE) { + params = m_memory_object->get_internal_params(cldnn::runtime_types::ze); + } else { + OPENVINO_THROW("[GPU] Can't update RemoteTensorImpl properties for unsupported context type (", ctx_type, ")"); + } switch (m_mem_type) { case TensorType::BT_BUF_INTERNAL: diff --git a/src/plugins/intel_gpu/src/runtime/device_query.cpp b/src/plugins/intel_gpu/src/runtime/device_query.cpp index 5340d86a26c288..edae3c27f20c32 100644 --- a/src/plugins/intel_gpu/src/runtime/device_query.cpp +++ b/src/plugins/intel_gpu/src/runtime/device_query.cpp @@ -15,29 +15,6 @@ namespace cldnn { int device_query::device_id = -1; -engine_types device_query::get_default_engine_type() { - auto engine_type = engine_types::ocl; -#ifdef OV_GPU_WITH_ZE_RT - engine_type = engine_types::ze; -#elif defined(OV_GPU_WITH_OCL_RT) - engine_type = engine_types::ocl; -#elif defined(OV_GPU_WITH_SYCL_RT) - engine_type = engine_types::sycl; -#endif - return engine_type; -} -runtime_types device_query::get_default_runtime_type() { - auto rt_type = runtime_types::ocl; -#ifdef OV_GPU_WITH_ZE_RT - rt_type = runtime_types::ze; -#elif defined(OV_GPU_WITH_OCL_RT) - rt_type = runtime_types::ocl; -#elif defined(OV_GPU_WITH_SYCL_RT) - rt_type = runtime_types::sycl; -#endif - return rt_type; -} - device_query::device_query(void* user_context, void* user_device, int ctx_device_id, @@ -59,14 +36,21 @@ device_query::device_query(engine_types engine_type, int target_tile_id, bool initialize_devices) { switch (runtime_type) { -#ifdef OV_GPU_WITH_OCL_RT case runtime_types::ocl: { OPENVINO_ASSERT(engine_type == engine_types::ocl || engine_type == engine_types::sycl); ocl::ocl_device_detector ocl_detector; _available_devices = ocl_detector.get_available_devices(user_context, user_device, ctx_device_id, target_tile_id, initialize_devices); +#ifdef OV_GPU_WITH_ZE_RT + // If running with ZE runtime, convert found OCL devices to ZE devices + std::map ze_devices; + for (auto& device : _available_devices) { + auto ze_device = ze::create_ze_device_from_ocl_device(device.second, initialize_devices); + ze_devices[device.first] = ze_device; + } + _available_devices = std::move(ze_devices); +#endif break; } -#endif #ifdef OV_GPU_WITH_ZE_RT case runtime_types::ze: { OPENVINO_ASSERT(engine_type == engine_types::ze); diff --git a/src/plugins/intel_gpu/src/runtime/engine_configuration.cpp b/src/plugins/intel_gpu/src/runtime/engine_configuration.cpp new file mode 100644 index 00000000000000..1c84b66495ba55 --- /dev/null +++ b/src/plugins/intel_gpu/src/runtime/engine_configuration.cpp @@ -0,0 +1,32 @@ +// Copyright (C) 2018-2026 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +#include "intel_gpu/runtime/engine_configuration.hpp" + +namespace cldnn { + +engine_types get_default_engine_type() { +#ifdef OV_GPU_WITH_ZE_RT + return engine_types::ze; +#elif defined(OV_GPU_WITH_OCL_RT) + return engine_types::ocl; +#elif defined(OV_GPU_WITH_SYCL_RT) + return engine_types::sycl; +#else + #error "Expected OpenVINO GPU runtime macros to be defined" +#endif +} + +runtime_types get_default_runtime_type() { +#ifdef OV_GPU_WITH_ZE_RT + return runtime_types::ze; +#elif defined(OV_GPU_WITH_OCL_RT) + return runtime_types::ocl; +#elif defined(OV_GPU_WITH_SYCL_RT) + return runtime_types::sycl; +#else + #error "Expected OpenVINO GPU runtime macros to be defined" +#endif +} +} // namespace cldnn diff --git a/src/plugins/intel_gpu/src/runtime/ocl/ocl_device.cpp b/src/plugins/intel_gpu/src/runtime/ocl/ocl_device.cpp index ff37b3ab54c373..f95df0fbf50d3f 100644 --- a/src/plugins/intel_gpu/src/runtime/ocl/ocl_device.cpp +++ b/src/plugins/intel_gpu/src/runtime/ocl/ocl_device.cpp @@ -355,6 +355,7 @@ device_info init_device_info(const cl::Device& device, const cl::Context& contex info.device_memory_ordinal = 0; info.supports_cp_offload = false; info.supports_counter_based_events = false; + info.supports_leo = false; #if defined(ENABLE_ONEDNN_FOR_GPU) && defined(OV_GPU_WITH_OCL_RT) using namespace dnnl::impl::gpu::intel::jit; diff --git a/src/plugins/intel_gpu/src/runtime/ocl/ocl_engine.cpp b/src/plugins/intel_gpu/src/runtime/ocl/ocl_engine.cpp index d90436c6c8be6a..f56d1f021c86b4 100644 --- a/src/plugins/intel_gpu/src/runtime/ocl/ocl_engine.cpp +++ b/src/plugins/intel_gpu/src/runtime/ocl/ocl_engine.cpp @@ -140,7 +140,7 @@ memory::ptr ocl_engine::import_buffer(const layout& layout, ov::intel_gpu::os_ha }; cl_int errcode = CL_SUCCESS; - auto cl_ctx = static_cast(get_user_context()); + auto cl_ctx = static_cast(get_user_context(runtime_types::ocl)); OPENVINO_ASSERT(cl_ctx != nullptr, "[GPU] OpenCL context is null while importing external buffer"); const auto byte_size = layout.bytes_count(); cl_mem imported = clCreateBufferWithProperties(cl_ctx, props, CL_MEM_READ_WRITE, byte_size, nullptr, &errcode); @@ -350,7 +350,9 @@ bool ocl_engine::is_the_same_buffer(const memory& mem1, const memory& mem2) { reinterpret_cast(mem2).get_buffer()); } -void* ocl_engine::get_user_context() const { +void* ocl_engine::get_user_context(runtime_types rt_type) const { + OPENVINO_ASSERT(rt_type == runtime_types::ocl, + "[GPU] OCL engine can only provide OCL context but requested context for ", rt_type); auto& cl_device = downcast(*_device); return static_cast(cl_device.get_context().get()); } diff --git a/src/plugins/intel_gpu/src/runtime/ocl/ocl_engine.hpp b/src/plugins/intel_gpu/src/runtime/ocl/ocl_engine.hpp index 9315eb65d0f618..586a7fd02ff421 100644 --- a/src/plugins/intel_gpu/src/runtime/ocl/ocl_engine.hpp +++ b/src/plugins/intel_gpu/src/runtime/ocl/ocl_engine.hpp @@ -30,10 +30,9 @@ class ocl_engine : public engine { memory_ptr reinterpret_buffer(const memory& memory, const layout& new_layout) override; memory_ptr import_buffer(const layout&, ov::intel_gpu::os_handle_param external_handle) override; bool is_the_same_buffer(const memory& mem1, const memory& mem2) override; - void release_external_memory(cl_mem) const; - void* get_user_context() const override; + void* get_user_context(runtime_types rt_type) const override; allocation_type get_default_allocation_type() const override { return allocation_type::cl_mem; } allocation_type detect_usm_allocation_type(const void* memory) const override; diff --git a/src/plugins/intel_gpu/src/runtime/ocl/ocl_memory.cpp b/src/plugins/intel_gpu/src/runtime/ocl/ocl_memory.cpp index 9e2818c4e14383..c0cf6288bd5d3b 100644 --- a/src/plugins/intel_gpu/src/runtime/ocl/ocl_memory.cpp +++ b/src/plugins/intel_gpu/src/runtime/ocl/ocl_memory.cpp @@ -122,7 +122,8 @@ event::ptr gpu_buffer::fill(stream& stream, unsigned char pattern, const std::ve return ev; } -shared_mem_params gpu_buffer::get_internal_params() const { +shared_mem_params gpu_buffer::get_internal_params(runtime_types rt_type) const { + OPENVINO_ASSERT(rt_type == runtime_types::ocl, "[GPU] Can not provide internal params for non-OCL runtime"); auto cl_engine = downcast(_engine); return {shared_mem_type::shared_mem_buffer, static_cast(cl_engine->get_cl_context().get()), nullptr, static_cast(_buffer.get()), @@ -366,7 +367,8 @@ void gpu_image2d::unlock(const stream& stream) { } -shared_mem_params gpu_image2d::get_internal_params() const { +shared_mem_params gpu_image2d::get_internal_params(runtime_types rt_type) const { + OPENVINO_ASSERT(rt_type == runtime_types::ocl, "[GPU] gpu_image2d can not provide internal params for non-OCL runtime"); auto cl_engine = downcast(_engine); return {shared_mem_type::shared_mem_image, static_cast(cl_engine->get_cl_context().get()), nullptr, static_cast(_buffer.get()), @@ -445,7 +447,8 @@ gpu_media_buffer::gpu_media_buffer(ocl_engine* engine, surface(params.surface), plane(params.plane) { } -shared_mem_params gpu_media_buffer::get_internal_params() const { +shared_mem_params gpu_media_buffer::get_internal_params(runtime_types rt_type) const { + OPENVINO_ASSERT(rt_type == runtime_types::ocl, "[GPU] gpu_media_buffer can not provide internal params for non-OCL runtime"); auto cl_engine = downcast(_engine); return {shared_mem_type::shared_mem_vasurface, static_cast(cl_engine->get_cl_context().get()), device, static_cast(_buffer.get()), surface, plane }; @@ -460,7 +463,8 @@ gpu_dx_buffer::gpu_dx_buffer(ocl_engine* engine, device(params.user_device), resource(params.mem) { } -shared_mem_params gpu_dx_buffer::get_internal_params() const { +shared_mem_params gpu_dx_buffer::get_internal_params(runtime_types rt_type) const { + OPENVINO_ASSERT(rt_type == runtime_types::ocl, "[GPU] gpu_dx_buffer can not provide internal params for non-OCL runtime"); auto cl_engine = downcast(_engine); return {shared_mem_type::shared_mem_dxbuffer, static_cast(cl_engine->get_cl_context().get()), device, static_cast(_buffer.get()), resource, 0 }; @@ -687,7 +691,8 @@ dnnl::memory gpu_usm::get_onednn_grouped_memory(dnnl::memory::desc desc, const m } #endif -shared_mem_params gpu_usm::get_internal_params() const { +shared_mem_params gpu_usm::get_internal_params(runtime_types rt_type) const { + OPENVINO_ASSERT(rt_type == runtime_types::ocl, "[GPU] gpu_usm can not provide internal params for non-OCL runtime"); auto cl_engine = downcast(_engine); return { shared_mem_type::shared_mem_usm, // shared_mem_type diff --git a/src/plugins/intel_gpu/src/runtime/ocl/ocl_memory.hpp b/src/plugins/intel_gpu/src/runtime/ocl/ocl_memory.hpp index a2cc3db172c294..4eb41ad49948b0 100644 --- a/src/plugins/intel_gpu/src/runtime/ocl/ocl_memory.hpp +++ b/src/plugins/intel_gpu/src/runtime/ocl/ocl_memory.hpp @@ -39,7 +39,7 @@ struct gpu_buffer : public lockable_gpu_mem, public memory { void unlock(const stream& stream) override; event::ptr fill(stream& stream, unsigned char pattern, const std::vector& dep_events = {}, bool blocking = true) override; event::ptr fill(stream& stream, const std::vector& dep_events = {}, bool blocking = true) override; - shared_mem_params get_internal_params() const override; + shared_mem_params get_internal_params(runtime_types rt_type) const override; const cl::Buffer& get_buffer() const { assert(0 == _lock_count); return _buffer; @@ -74,7 +74,7 @@ struct gpu_image2d : public lockable_gpu_mem, public memory { void unlock(const stream& stream) override; event::ptr fill(stream& stream, unsigned char pattern, const std::vector& dep_events = {}, bool blocking = true) override; event::ptr fill(stream& stream, const std::vector& dep_events = {}, bool blocking = true) override; - shared_mem_params get_internal_params() const override; + shared_mem_params get_internal_params(runtime_types rt_type) const override; const cl::Image2D& get_buffer() const { assert(0 == _lock_count); return _buffer; @@ -94,7 +94,7 @@ struct gpu_image2d : public lockable_gpu_mem, public memory { struct gpu_media_buffer : public gpu_image2d { gpu_media_buffer(ocl_engine* engine, const layout& new_layout, shared_mem_params params); - shared_mem_params get_internal_params() const override; + shared_mem_params get_internal_params(runtime_types rt_type) const override; private: void* device; #ifdef _WIN32 @@ -108,7 +108,7 @@ struct gpu_media_buffer : public gpu_image2d { #ifdef _WIN32 struct gpu_dx_buffer : public gpu_buffer { gpu_dx_buffer(ocl_engine* engine, const layout& new_layout, shared_mem_params VAEncMiscParameterTypeSubMbPartPel); - shared_mem_params get_internal_params() const override; + shared_mem_params get_internal_params(runtime_types rt_type) const override; private: void* device; void* resource; @@ -128,7 +128,7 @@ struct gpu_usm : public lockable_gpu_mem, public memory { event::ptr fill(stream& stream, unsigned char pattern, const std::vector& dep_events = {}, bool blocking = true) override; event::ptr fill(stream& stream, const std::vector& dep_events = {}, bool blocking = true) override; - shared_mem_params get_internal_params() const override; + shared_mem_params get_internal_params(runtime_types rt_type) const override; event::ptr copy_from(stream& stream, const void* data_ptr, size_t src_offset, size_t dst_offset, size_t size, bool blocking) override; event::ptr copy_from(stream& stream, const memory& src_mem, size_t src_offset, size_t dst_offset, size_t size, bool blocking) override; diff --git a/src/plugins/intel_gpu/src/runtime/ze/ze_common.hpp b/src/plugins/intel_gpu/src/runtime/ze/ze_common.hpp index ffdfc9580ea80e..98544a4f2b3046 100644 --- a/src/plugins/intel_gpu/src/runtime/ze/ze_common.hpp +++ b/src/plugins/intel_gpu/src/runtime/ze/ze_common.hpp @@ -7,6 +7,7 @@ #include "openvino/core/except.hpp" #define ZERO_API_KEEP_SYMBOLS_LIST_MACRO #include "openvino/zero_api.hpp" +#undef ZERO_API_KEEP_SYMBOLS_LIST_MACRO #include #include @@ -57,7 +58,7 @@ inline const ::ov::ZeroApi& get_ze_api_instance() { inline typename std::invoke_result::type wrapped_##symbol(Args... args) { \ const auto& ze_api = get_ze_api_instance(); \ if (ze_api.symbol == nullptr) { \ - OPENVINO_THROW("Unsupported symbol " #symbol); \ + return ZE_RESULT_ERROR_UNSUPPORTED_FEATURE; \ } \ return ze_api.symbol(std::forward(args)...); \ } diff --git a/src/plugins/intel_gpu/src/runtime/ze/ze_counter_based_event.cpp b/src/plugins/intel_gpu/src/runtime/ze/ze_counter_based_event.cpp index bc6cd1e075334f..c43884aa959ee8 100644 --- a/src/plugins/intel_gpu/src/runtime/ze/ze_counter_based_event.cpp +++ b/src/plugins/intel_gpu/src/runtime/ze/ze_counter_based_event.cpp @@ -13,7 +13,7 @@ using namespace cldnn; using namespace ze; void ze_counter_based_event::wait_impl() { - OV_ZE_EXPECT(ze::zeEventHostSynchronize(m_event, endless_wait)); + OV_ZE_EXPECT(ze::zeEventHostSynchronize(m_event.handle(), endless_wait)); } void ze_counter_based_event::set_impl() { @@ -21,7 +21,7 @@ void ze_counter_based_event::set_impl() { } bool ze_counter_based_event::is_set_impl() { - auto ret = ze::zeEventQueryStatus(m_event); + auto ret = ze::zeEventQueryStatus(m_event.handle()); switch (ret) { case ZE_RESULT_SUCCESS: return true; @@ -36,7 +36,7 @@ bool ze_counter_based_event::is_set_impl() { } ze_event_handle_t ze_counter_based_event::get_handle() const { - return m_event; + return m_event.handle(); } std::optional ze_counter_based_event::query_timestamp() { @@ -44,7 +44,7 @@ std::optional ze_counter_based_event::query_timest return std::nullopt; } ze_kernel_timestamp_result_t timestamp{}; - OV_ZE_EXPECT(ze::zeEventQueryKernelTimestamp(m_event, ×tamp)); + OV_ZE_EXPECT(ze::zeEventQueryKernelTimestamp(m_event.handle(), ×tamp)); return timestamp; } @@ -66,7 +66,3 @@ bool ze_counter_based_event::get_profiling_info_impl(std::list(queue_stamp, *this, event); + OV_ZE_EXPECT(func_zexCounterBasedEventCreate2(m_engine.get_context().handle(), m_engine.get_device().handle(), &desc, &event)); + auto ev_holder = ze_event_resource(event); + auto cb_event = std::make_shared(queue_stamp, *this, ev_holder); return cb_event; } diff --git a/src/plugins/intel_gpu/src/runtime/ze/ze_device.cpp b/src/plugins/intel_gpu/src/runtime/ze/ze_device.cpp index c400ac8f254335..384016409aaf22 100644 --- a/src/plugins/intel_gpu/src/runtime/ze/ze_device.cpp +++ b/src/plugins/intel_gpu/src/runtime/ze/ze_device.cpp @@ -4,6 +4,7 @@ #include "ze_device.hpp" #include "ze_common.hpp" +#include "ze_ocl_interop.hpp" #include "compute_runtime/zex_common.h" #include "compute_runtime/ze_intel_gpu.h" @@ -219,6 +220,7 @@ device_info init_device_info(ze_driver_handle_t driver, ze_device_handle_t devic info.supports_immad = supports_dp_properties && (dp_properties.flags & ZE_INTEL_DEVICE_MODULE_EXP_FLAG_DPAS) != 0; info.supports_usm = device_memory_access_properties.hostAllocCapabilities && device_memory_access_properties.deviceAllocCapabilities; + info.supports_leo = ze_ocl_interop::get_instance().check_support(device); // FIXME: Could not find how to retrieve those from ZE info.supports_work_group_collective_functions = false; @@ -321,6 +323,9 @@ memory_capabilities init_memory_caps(ze_device_handle_t device, const device_inf } if (device_memory_access_properties.deviceAllocCapabilities) { memory_caps.push_back(allocation_type::usm_device); + if (info.supports_leo) { + memory_caps.push_back(allocation_type::cl_mem); + } } } if (info.supports_image) { @@ -333,49 +338,45 @@ memory_capabilities init_memory_caps(ze_device_handle_t device, const device_inf } // namespace -ze_device::ze_device(ze_driver_handle_t driver, ze_device_handle_t device, bool initialize) -: _driver(driver) -, _device(device) -, _info(init_device_info(driver, device)) -, _mem_caps(init_memory_caps(device, _info)) { - if (initialize) { - this->initialize(); +ze_device::ze_device(ze_driver_resource driver, ze_device_resource device, ze_context_resource context, bool initialize_device) +: _driver(std::move(driver)) +, _device(std::move(device)) +, _context(std::move(context)) +, _info(init_device_info(_driver.handle(), _device.handle())) +, _mem_caps(init_memory_caps(_device.handle(), _info)) { + OPENVINO_ASSERT(!_driver.is_empty(), "[GPU] Expected non-empty driver when creating ze_device"); + OPENVINO_ASSERT(!_device.is_empty(), "[GPU] Expected non-empty device when creating ze_device"); + if (initialize_device) { + initialize(); } } void ze_device::initialize() { - if (_is_initialized) + if (is_initialized()) return; ze_context_desc_t context_desc = { ZE_STRUCTURE_TYPE_CONTEXT_DESC, nullptr, 0 }; - OV_ZE_EXPECT(ze::zeContextCreate(_driver, &context_desc, &_context)); - _is_initialized = true; + ze_context_handle_t ctx = nullptr; + OV_ZE_EXPECT(ze::zeContextCreate(_driver.handle(), &context_desc, &ctx)); + _context = ze_context_resource{ctx}; } bool ze_device::is_initialized() const { - return _is_initialized; + return _context.is_empty() == false; } bool ze_device::is_same(const device::ptr other) { - auto casted = downcast(other.get()); + auto casted = std::dynamic_pointer_cast(other); if (!casted) return false; - if (is_initialized() && casted->is_initialized()) { - // Do not compare contexts as one driver can have many different contexts - return _device == casted->get_device() && _driver == casted->get_driver(); - } - return _info.is_same_device(casted->_info); + return _device.handle() == casted->get_device().handle() + && _driver.handle() == casted->get_driver().handle(); } void ze_device::set_mem_caps(const memory_capabilities& memory_capabilities) { _mem_caps = memory_capabilities; } -ze_device::~ze_device() { - if (_is_initialized) - OV_ZE_WARN(ze::zeContextDestroy(_context)); -} - } // namespace ze } // namespace cldnn diff --git a/src/plugins/intel_gpu/src/runtime/ze/ze_device.hpp b/src/plugins/intel_gpu/src/runtime/ze/ze_device.hpp index bbd65c03250b85..70c2637c493724 100644 --- a/src/plugins/intel_gpu/src/runtime/ze/ze_device.hpp +++ b/src/plugins/intel_gpu/src/runtime/ze/ze_device.hpp @@ -6,13 +6,14 @@ #include "intel_gpu/runtime/device.hpp" #include "ze_common.hpp" +#include "ze_resource.hpp" namespace cldnn { namespace ze { struct ze_device : public device { public: - ze_device(ze_driver_handle_t driver, ze_device_handle_t device, bool initialize = true); + ze_device(ze_driver_resource driver, ze_device_resource device, ze_context_resource context, bool initialize_device = true); ze_device(const ze_device &other) = delete; ze_device& operator=(const ze_device &other) = delete; @@ -22,20 +23,19 @@ struct ze_device : public device { void initialize() override; bool is_initialized() const override; - ze_driver_handle_t get_driver() const { return _driver; } - ze_device_handle_t get_device() const { return _device; } - ze_context_handle_t get_context() const { return _context; } + ze_driver_resource get_driver() const { return _driver; } + ze_device_resource get_device() const { return _device; } + ze_context_resource get_context() const { return _context; } bool is_same(const device::ptr other) override; void set_mem_caps(const memory_capabilities& memory_capabilities) override; - ~ze_device(); + ~ze_device() = default; private: - ze_driver_handle_t _driver = nullptr; - ze_device_handle_t _device = nullptr; - ze_context_handle_t _context = nullptr; - bool _is_initialized = false; + ze_driver_resource _driver; + ze_device_resource _device; + ze_context_resource _context; device_info _info; memory_capabilities _mem_caps; diff --git a/src/plugins/intel_gpu/src/runtime/ze/ze_device_detector.cpp b/src/plugins/intel_gpu/src/runtime/ze/ze_device_detector.cpp index e92b5379be69cf..5b2a990a7351ce 100644 --- a/src/plugins/intel_gpu/src/runtime/ze/ze_device_detector.cpp +++ b/src/plugins/intel_gpu/src/runtime/ze/ze_device_detector.cpp @@ -5,6 +5,8 @@ #include "ze_device_detector.hpp" #include "ze_device.hpp" #include "ze_common.hpp" +#include "ze_resource_interop.hpp" +#include "ocl/ocl_device.hpp" #include @@ -49,7 +51,7 @@ std::map ze_device_detector::get_available_devices(voi auto root_device = std::dynamic_pointer_cast(dptr); OPENVINO_ASSERT(root_device != nullptr, "[GPU] Invalid device type created in ocl_device_detector"); - auto sub_devices = get_sub_devices(root_device->get_device()); + auto sub_devices = get_sub_devices(root_device->get_device().handle()); if (!sub_devices.empty()) { uint32_t sub_idx = 0; for (auto& sub_device : sub_devices) { @@ -57,7 +59,8 @@ std::map ze_device_detector::get_available_devices(voi sub_idx++; continue; } - auto sub_device_ptr = std::make_shared(root_device->get_driver(), sub_device, initialize_devices); + ze_device_resource sub_device_res(sub_device); + auto sub_device_ptr = std::make_shared(root_device->get_driver(), sub_device_res, ze_context_resource{}, initialize_devices); ret[map_id + "." + std::to_string(sub_idx++)] = sub_device_ptr; } } @@ -90,7 +93,9 @@ std::vector ze_device_detector::create_device_list(bool initialize_ OV_ZE_EXPECT(ze::zeDeviceGetProperties(all_devices[d], &device_properties)); if (ZE_DEVICE_TYPE_GPU == device_properties.type) { - ret.emplace_back(std::make_shared(all_drivers[i], all_devices[d], initialize_devices)); + ze_driver_resource driver_res(all_drivers[i]); + ze_device_resource device_res(all_devices[d]); + ret.emplace_back(std::make_shared(driver_res, device_res, ze_context_resource{}, initialize_devices)); } } catch (std::exception& ex) { GPU_DEBUG_LOG << "Devices query/creation failed for driver " << i << ex.what() << std::endl; @@ -111,5 +116,19 @@ std::vector ze_device_detector::create_device_list_from_user_device OPENVINO_NOT_IMPLEMENTED; } +device::ptr create_ze_device_from_ocl_device(device::ptr ocl_device, bool initialize) { + auto ocl_dev = std::dynamic_pointer_cast(ocl_device); + OPENVINO_ASSERT(ocl_dev != nullptr, "[GPU] Invalid device type for OpenCL to ZE device conversion"); + auto device = ocl_dev->get_device().get(); + auto ze_driver_res = ze_import_driver(device); + auto ze_device_res = ze_import_device(device); + auto context = ocl_dev->get_context().get(); + if (context == nullptr) { + return std::make_shared(ze_driver_res, ze_device_res, ze_context_resource{}, initialize); + } + auto ze_context_res = ze_import_context(context); + return std::make_shared(ze_driver_res, ze_device_res, ze_context_res, initialize); +} + } // namespace ze } // namespace cldnn diff --git a/src/plugins/intel_gpu/src/runtime/ze/ze_device_detector.hpp b/src/plugins/intel_gpu/src/runtime/ze/ze_device_detector.hpp index ce3969cc79b0a6..44929c4d5aa6f6 100644 --- a/src/plugins/intel_gpu/src/runtime/ze/ze_device_detector.hpp +++ b/src/plugins/intel_gpu/src/runtime/ze/ze_device_detector.hpp @@ -28,5 +28,7 @@ class ze_device_detector { std::vector create_device_list_from_user_device(void* user_device) const; }; +device::ptr create_ze_device_from_ocl_device(device::ptr ocl_device, bool initialize_devices = false); + } // namespace ze } // namespace cldnn diff --git a/src/plugins/intel_gpu/src/runtime/ze/ze_engine.cpp b/src/plugins/intel_gpu/src/runtime/ze/ze_engine.cpp index 832ebdd31be007..08590b0ddb59d7 100644 --- a/src/plugins/intel_gpu/src/runtime/ze/ze_engine.cpp +++ b/src/plugins/intel_gpu/src/runtime/ze/ze_engine.cpp @@ -13,8 +13,7 @@ #include "ze_stream.hpp" #include "ze_device.hpp" #include "ze_kernel.hpp" -#include "ze_module_holder.hpp" -#include "ze_kernel_holder.hpp" +#include "ze_resource_interop.hpp" #include #include #include @@ -41,25 +40,27 @@ void ze_engine::create_onednn_engine(const ExecutionConfig& config) { const std::lock_guard lock(onednn_mutex); OPENVINO_ASSERT(_device->get_info().vendor_id == INTEL_VENDOR_ID, "[GPU] OneDNN engine can be used for Intel GPUs only"); if (!_onednn_engine) { - auto casted = std::dynamic_pointer_cast(_device); - _onednn_engine = std::make_shared(dnnl::ze_interop::make_engine(casted->get_driver(), casted->get_device(), casted->get_context())); + _onednn_engine = std::make_shared(dnnl::ze_interop::make_engine( + get_driver().handle(), + get_device().handle(), + get_context().handle() + )); } } #endif - -ze_driver_handle_t ze_engine::get_driver() const { +ze_driver_resource ze_engine::get_driver() const { auto casted = std::dynamic_pointer_cast(_device); OPENVINO_ASSERT(casted, "[GPU] Invalid device type for ze_engine"); return casted->get_driver(); } -ze_context_handle_t ze_engine::get_context() const { +ze_context_resource ze_engine::get_context() const { auto casted = std::dynamic_pointer_cast(_device); OPENVINO_ASSERT(casted, "[GPU] Invalid device type for ze_engine"); return casted->get_context(); } -ze_device_handle_t ze_engine::get_device() const { +ze_device_resource ze_engine::get_device() const { auto casted = std::dynamic_pointer_cast(_device); OPENVINO_ASSERT(casted, "[GPU] Invalid device type for ze_engine"); return casted->get_device(); @@ -78,7 +79,7 @@ memory::ptr ze_engine::allocate_memory(const layout& layout, allocation_type typ memory::ptr res; if (layout.format.is_image_2d()) { res = std::make_shared(this, layout); - } else if (memory_capabilities::is_usm_type(type)){ + } else if (memory_capabilities::is_usm_type(type) || type == allocation_type::cl_mem){ res = std::make_shared(this, layout, type); } else { OPENVINO_THROW("[GPU] Unsupported allocation type: ", type); @@ -109,16 +110,17 @@ memory::ptr ze_engine::reinterpret_buffer(const memory& memory, const layout& ne bool from_memory_pool = memory.from_memory_pool; memory::ptr reinterpret_memory = nullptr; - if (memory_capabilities::is_usm_type(memory.get_allocation_type())) { + if (memory_capabilities::is_usm_type(memory.get_allocation_type()) + || memory.get_allocation_type() == allocation_type::cl_mem) { reinterpret_memory = std::make_shared(this, new_layout, - reinterpret_cast(memory).get_buffer(), + reinterpret_cast(memory).get_resource(), memory.get_allocation_type(), memory.get_mem_tracker()); } else if (new_layout.format.is_image_2d()) { reinterpret_memory = std::make_shared(this, new_layout, - reinterpret_cast(memory).get_handle(), + reinterpret_cast(memory).get_resource(), memory.get_mem_tracker()); } else { OPENVINO_THROW("[GPU] Unexpected memory type for reinterpret_buffer"); @@ -128,18 +130,23 @@ memory::ptr ze_engine::reinterpret_buffer(const memory& memory, const layout& ne } memory::ptr ze_engine::reinterpret_handle(const layout& new_layout, shared_mem_params params) { + // Convert OCL handles from `params` to L0 and create memory objects if (params.mem_type == shared_mem_type::shared_mem_usm) { - ze::UsmMemory usm_buffer(get_context(), get_device(), params.mem); - size_t actual_mem_size = 0; - OV_ZE_EXPECT(ze::zeMemGetAddressRange(get_context(), params.mem, nullptr, &actual_mem_size)); - auto requested_mem_size = new_layout.bytes_count(); - OPENVINO_ASSERT(actual_mem_size >= requested_mem_size, - "[GPU] shared USM buffer has smaller size (", actual_mem_size, - ") than specified layout (", requested_mem_size, ")"); - return std::make_shared(this, new_layout, usm_buffer, nullptr); + // USM memory does not need to be converted + auto ctx = get_context(); + ov_ze_usm_handle usm_handle{ctx.handle(), params.mem}; + const bool is_borrowed = true; + ze_usm_resource usm_res(usm_handle, is_borrowed); + return std::make_shared(this, new_layout, usm_res, nullptr); + } else if (params.mem_type == shared_mem_type::shared_mem_buffer) { + auto ctx = get_context(); + auto ocl_buffer = static_cast(params.mem); + auto imported_buffer = ze_import_usm(ocl_buffer, ctx); + return std::make_shared(this, new_layout, imported_buffer, allocation_type::cl_mem, nullptr); } else if (params.mem_type == shared_mem_type::shared_mem_image) { - auto image = reinterpret_cast(params.mem); - return std::make_shared(this, new_layout, image, nullptr); + auto ocl_image = static_cast(params.mem); + auto imported_image = ze_import_image(ocl_image); + return std::make_shared(this, new_layout, imported_image, nullptr); } else { OPENVINO_THROW("[GPU] Unsupported shared memory type: ", params.mem_type); } @@ -152,11 +159,14 @@ memory_ptr ze_engine::create_subbuffer(const memory& memory, const layout& new_l } OPENVINO_ASSERT(memory_capabilities::is_usm_type(memory.get_allocation_type()), "[GPU] Trying to create subbuffer for non usm memory"); auto& new_buf = reinterpret_cast(memory); - auto ptr = new_buf.get_buffer().get(); - auto sub_buffer = ze::UsmMemory(get_context(), get_device(), ptr, byte_offset); + auto ptr = new_buf.buffer_ptr(); + auto ctx = get_context(); + ov_ze_usm_handle usm_handle{ctx.handle(), reinterpret_cast(ptr) + byte_offset}; + const bool is_borrowed = true; + ze_usm_resource usm_res(usm_handle, is_borrowed); return std::make_shared(this, new_layout, - sub_buffer, + usm_res, memory.get_allocation_type(), memory.get_mem_tracker()); } @@ -172,8 +182,17 @@ bool ze_engine::is_the_same_buffer(const memory& mem1, const memory& mem2) { return false; if (&mem1 == &mem2) return true; - - return (reinterpret_cast(mem1).get_buffer().get() == reinterpret_cast(mem2).get_buffer().get()); + + auto alloc_type = mem1.get_allocation_type(); + if (memory_capabilities::is_usm_type(mem1.get_allocation_type()) || alloc_type == allocation_type::cl_mem) { + const auto &usm1 = downcast(mem1); + const auto &usm2 = downcast(mem2); + return usm1.buffer_ptr() == usm2.buffer_ptr(); + } else { + const auto &img1 = downcast(mem1); + const auto &img2 = downcast(mem2); + return img1.get_resource().handle() == img2.get_resource().handle(); + } } std::shared_ptr ze_engine::create_kernel_builder() const { @@ -182,9 +201,16 @@ std::shared_ptr ze_engine::create_kernel_builder() const { return std::make_shared(*casted); } -void* ze_engine::get_user_context() const { - auto& casted = downcast(*_device); - return static_cast(casted.get_context()); +void* ze_engine::get_user_context(runtime_types rt_type) const { + auto ctx = get_context(); + if (rt_type == runtime_types::ze) { + return ctx.handle(); + } else if (rt_type == runtime_types::ocl) { + ze_export_ocl_context(ctx, get_device()); + return ctx.ocl_handle(); + } else { + OPENVINO_THROW("[GPU] ZE engine cannot provide context for ", rt_type); + } } stream::ptr ze_engine::create_stream(const ExecutionConfig& config) const { @@ -192,7 +218,9 @@ stream::ptr ze_engine::create_stream(const ExecutionConfig& config) const { } stream::ptr ze_engine::create_stream(const ExecutionConfig& config, void* handle) const { - OPENVINO_NOT_IMPLEMENTED; + cl_command_queue ocl_handle = static_cast(handle); + auto ze_cmd_list = ze_import_command_list(ocl_handle); + return std::make_shared(*this, config, ze_cmd_list); } std::shared_ptr ze_engine::create(const device::ptr device, runtime_types runtime_type) { diff --git a/src/plugins/intel_gpu/src/runtime/ze/ze_engine.hpp b/src/plugins/intel_gpu/src/runtime/ze/ze_engine.hpp index 0082bf6996a1bf..ebacf2a8719217 100644 --- a/src/plugins/intel_gpu/src/runtime/ze/ze_engine.hpp +++ b/src/plugins/intel_gpu/src/runtime/ze/ze_engine.hpp @@ -5,6 +5,7 @@ #pragma once #include "ze_common.hpp" +#include "ze_resource.hpp" #include "intel_gpu/runtime/memory.hpp" #include "intel_gpu/runtime/engine.hpp" #include "intel_gpu/runtime/stream.hpp" @@ -30,14 +31,14 @@ class ze_engine : public engine { memory_ptr reinterpret_buffer(const memory& memory, const layout& new_layout) override; bool is_the_same_buffer(const memory& mem1, const memory& mem2) override; - void* get_user_context() const override; + void* get_user_context(runtime_types rt_type) const override; allocation_type get_default_allocation_type() const override { return allocation_type::usm_device; } allocation_type detect_usm_allocation_type(const void* memory) const override; - ze_context_handle_t get_context() const; - ze_driver_handle_t get_driver() const; - ze_device_handle_t get_device() const; + ze_context_resource get_context() const; + ze_driver_resource get_driver() const; + ze_device_resource get_device() const; stream_ptr create_stream(const ExecutionConfig& config) const override; stream_ptr create_stream(const ExecutionConfig& config, void *handle) const override; diff --git a/src/plugins/intel_gpu/src/runtime/ze/ze_event.cpp b/src/plugins/intel_gpu/src/runtime/ze/ze_event.cpp index 316e14da287ab1..5af712d07e50c5 100644 --- a/src/plugins/intel_gpu/src/runtime/ze/ze_event.cpp +++ b/src/plugins/intel_gpu/src/runtime/ze/ze_event.cpp @@ -14,19 +14,19 @@ using namespace ze; void ze_event::reset() { event::reset(); - OV_ZE_EXPECT(ze::zeEventHostReset(m_event)); + OV_ZE_EXPECT(ze::zeEventHostReset(m_event.handle())); } void ze_event::wait_impl() { - OV_ZE_EXPECT(ze::zeEventHostSynchronize(m_event, endless_wait)); + OV_ZE_EXPECT(ze::zeEventHostSynchronize(m_event.handle(), endless_wait)); } void ze_event::set_impl() { - OV_ZE_EXPECT(ze::zeEventHostSignal(m_event)); + OV_ZE_EXPECT(ze::zeEventHostSignal(m_event.handle())); } bool ze_event::is_set_impl() { - auto ret = ze::zeEventQueryStatus(m_event); + auto ret = ze::zeEventQueryStatus(m_event.handle()); switch (ret) { case ZE_RESULT_SUCCESS: return true; @@ -45,12 +45,12 @@ std::optional ze_event::query_timestamp() { return std::nullopt; } ze_kernel_timestamp_result_t timestamp{}; - OV_ZE_EXPECT(ze::zeEventQueryKernelTimestamp(m_event, ×tamp)); + OV_ZE_EXPECT(ze::zeEventQueryKernelTimestamp(m_event.handle(), ×tamp)); return timestamp; } ze_event_handle_t ze_event::get_handle() const { - return m_event; + return m_event.handle(); } bool ze_event::get_profiling_info_impl(std::list& info) { @@ -71,7 +71,3 @@ bool ze_event::get_profiling_info_impl(std::list event_pool) + ze_event(uint64_t queue_stamp, const ze_base_event_factory& factory, ze_event_resource ev, ze_event_pool_resource ev_pool) : ze_base_event(queue_stamp) - , m_event_pool(event_pool) , m_factory(factory) + , m_event_pool(ev_pool) , m_event(ev) { - // Ensure event handle is not null - OPENVINO_ASSERT(ev != nullptr, "[GPU] Trying to create event with null handle"); + OPENVINO_ASSERT(!m_event.is_empty(), "[GPU] Attempt to create event with empty resource"); + OPENVINO_ASSERT(!m_event_pool.is_empty(), "[GPU] Attempt to create event with empty event pool resource"); } - ze_event(const ze_event &) = delete; - ze_event& operator=(const ze_event &) = delete; - ~ze_event(); void reset() override; std::optional query_timestamp() override; @@ -38,9 +34,10 @@ struct ze_event : public ze_base_event { // TODO: Implement add_event_handler_impl // bool add_event_handler_impl(event_handler, void*) override; - std::shared_ptr m_event_pool; const ze_base_event_factory& m_factory; - ze_event_handle_t m_event; + // To ensure correct lifetime, drop event pool after dropping event + ze_event_pool_resource m_event_pool; + ze_event_resource m_event; }; } // namespace ze diff --git a/src/plugins/intel_gpu/src/runtime/ze/ze_event_factory.cpp b/src/plugins/intel_gpu/src/runtime/ze/ze_event_factory.cpp index e1874a509fbb83..10cd41d3083dca 100644 --- a/src/plugins/intel_gpu/src/runtime/ze/ze_event_factory.cpp +++ b/src/plugins/intel_gpu/src/runtime/ze/ze_event_factory.cpp @@ -13,18 +13,27 @@ using namespace ze; ze_event_factory::ze_event_factory(const ze_engine &engine, bool enable_profiling, uint32_t capacity) : ze_base_event_factory(engine, enable_profiling) -, m_current_pool(nullptr) , m_capacity(capacity) , m_num_used(0) { } event::ptr ze_event_factory::create_event(uint64_t queue_stamp) { std::lock_guard lock(_mutex); - if (m_num_used >= m_capacity || !m_current_pool) { + if (m_num_used >= m_capacity || m_current_pool.is_empty()) { m_num_used = 0; ze_event_pool_flags_t flags = is_profiling_enabled() ? ZE_EVENT_POOL_FLAG_KERNEL_TIMESTAMP : 0; flags |= ZE_EVENT_POOL_FLAG_HOST_VISIBLE; - m_current_pool = std::make_shared(m_engine, m_capacity, flags); + ze_event_pool_desc_t event_pool_desc = { + ZE_STRUCTURE_TYPE_EVENT_POOL_DESC, + nullptr, + flags, + m_capacity + }; + auto ctx_handle = m_engine.get_context().handle(); + auto device_handle = m_engine.get_device().handle(); + ze_event_pool_handle_t event_pool; + OV_ZE_EXPECT(ze::zeEventPoolCreate(ctx_handle, &event_pool_desc, 1, &device_handle, &event_pool)); + m_current_pool = ze_event_pool_resource{event_pool}; } ze_event_handle_t event; @@ -35,7 +44,8 @@ event::ptr ze_event_factory::create_event(uint64_t queue_stamp) { ZE_EVENT_SCOPE_FLAG_HOST, 0 }; - OV_ZE_EXPECT(ze::zeEventCreate(m_current_pool->m_handle, &event_desc, &event)); + OV_ZE_EXPECT(ze::zeEventCreate(m_current_pool.handle(), &event_desc, &event)); + auto event_holder = ze_event_resource(event); - return std::make_shared(queue_stamp, *this, event, m_current_pool); + return std::make_shared(queue_stamp, *this, event_holder, m_current_pool); } diff --git a/src/plugins/intel_gpu/src/runtime/ze/ze_event_factory.hpp b/src/plugins/intel_gpu/src/runtime/ze/ze_event_factory.hpp index bf57db27858f49..51dd197a862d80 100644 --- a/src/plugins/intel_gpu/src/runtime/ze/ze_event_factory.hpp +++ b/src/plugins/intel_gpu/src/runtime/ze/ze_event_factory.hpp @@ -5,9 +5,9 @@ #pragma once #include "ze_base_event_factory.hpp" -#include "ze_event_pool.hpp" +#include "ze_resource.hpp" -#include "mutex" +#include namespace cldnn { namespace ze { @@ -19,7 +19,7 @@ struct ze_event_factory : public ze_base_event_factory { event::ptr create_event(uint64_t queue_stamp) override; protected: std::mutex _mutex; - std::shared_ptr m_current_pool; + ze_event_pool_resource m_current_pool; const uint32_t m_capacity; uint32_t m_num_used; }; diff --git a/src/plugins/intel_gpu/src/runtime/ze/ze_event_pool.cpp b/src/plugins/intel_gpu/src/runtime/ze/ze_event_pool.cpp deleted file mode 100644 index 45196578e8c836..00000000000000 --- a/src/plugins/intel_gpu/src/runtime/ze/ze_event_pool.cpp +++ /dev/null @@ -1,28 +0,0 @@ -// Copyright (C) 2018-2026 Intel Corporation -// SPDX-License-Identifier: Apache-2.0 -// - -#include "ze_event_pool.hpp" -#include "ze_event.hpp" -#include "ze_common.hpp" - -namespace cldnn { -namespace ze { - -ze_event_pool::ze_event_pool(const ze_engine& engine, uint32_t capacity, ze_event_pool_flags_t flags) - : m_engine(engine) { - ze_event_pool_desc_t event_pool_desc = { - ZE_STRUCTURE_TYPE_EVENT_POOL_DESC, - nullptr, - flags, - capacity - }; - auto device = engine.get_device(); - OV_ZE_EXPECT(ze::zeEventPoolCreate(engine.get_context(), &event_pool_desc, 1, &device, &m_handle)); -} - -ze_event_pool::~ze_event_pool() { - OV_ZE_WARN(ze::zeEventPoolDestroy(m_handle)); -} -} // namespace ze -} // namespace cldnn diff --git a/src/plugins/intel_gpu/src/runtime/ze/ze_event_pool.hpp b/src/plugins/intel_gpu/src/runtime/ze/ze_event_pool.hpp deleted file mode 100644 index 3ecd0d28892afb..00000000000000 --- a/src/plugins/intel_gpu/src/runtime/ze/ze_event_pool.hpp +++ /dev/null @@ -1,27 +0,0 @@ -// Copyright (C) 2018-2026 Intel Corporation -// SPDX-License-Identifier: Apache-2.0 -// - -#pragma once - -#include "ze_engine.hpp" - -namespace cldnn { -namespace ze { - -// RAII wrapper for Level Zero event pool -struct ze_event_pool { - ze_event_pool(const ze_engine& engine, uint32_t capacity, ze_event_pool_flags_t flags); - ~ze_event_pool(); - ze_event_pool(const ze_event_pool&) = delete; - ze_event_pool& operator=(const ze_event_pool&) = delete; - - using ptr = std::shared_ptr; - - ze_event_pool_handle_t m_handle; - const ze_engine& m_engine; -}; - -} // namespace ze -} // namespace cldnn - diff --git a/src/plugins/intel_gpu/src/runtime/ze/ze_kernel.hpp b/src/plugins/intel_gpu/src/runtime/ze/ze_kernel.hpp index ea8de8d275fe96..b0f197ab99ca53 100644 --- a/src/plugins/intel_gpu/src/runtime/ze/ze_kernel.hpp +++ b/src/plugins/intel_gpu/src/runtime/ze/ze_kernel.hpp @@ -7,7 +7,7 @@ #include "intel_gpu/runtime/kernel.hpp" #include "openvino/core/except.hpp" #include "ze_common.hpp" -#include "ze_kernel_holder.hpp" +#include "ze_resource.hpp" #include @@ -16,8 +16,8 @@ namespace ze { class ze_kernel : public kernel { public: - static void create_kernels_from_module(std::shared_ptr module, std::vector &out) { - ze_module_handle_t module_handle = module->get_module_handle(); + static void create_kernels_from_module(ze_module_resource module_holder, ze_module_build_log_resource build_log_holder, std::vector &out) { + ze_module_handle_t module_handle = module_holder.handle(); uint32_t kernel_count = 0; OV_ZE_EXPECT(ze::zeModuleGetKernelNames(module_handle, &kernel_count, nullptr)); std::vector kernel_names(kernel_count); @@ -37,22 +37,28 @@ class ze_kernel : public kernel { kernel_desc.pKernelName = name_cstr; ze_kernel_handle_t kernel_handle; OV_ZE_EXPECT(ze::zeKernelCreate(module_handle, &kernel_desc, &kernel_handle)); - auto kernel_holder = std::make_shared(kernel_handle, module); - out.push_back(std::make_shared(kernel_holder, name)); + ze_kernel_resource kernel_holder(kernel_handle); + out.push_back(std::make_shared(module_holder, kernel_holder, name, build_log_holder)); } } - ze_kernel(std::shared_ptr kernel, const std::string& kernel_id) - : m_kernel(kernel) - , m_kernel_id(kernel_id) { } + ze_kernel(ze_module_resource module_res, ze_kernel_resource kernel_res, const std::string& kernel_id, ze_module_build_log_resource build_log) + : m_module(std::move(module_res)) + , m_build_log(std::move(build_log)) + , m_kernel(std::move(kernel_res)) + , m_kernel_id(kernel_id) { + // Allow empty build log holder + OPENVINO_ASSERT(!m_kernel.is_empty(), "[GPU] Attempt to create kernel with empty kernel resource"); + OPENVINO_ASSERT(!m_module.is_empty(), "[GPU] Attempt to create kernel with empty module resource"); + } - ze_kernel_handle_t get_kernel_handle() const { return m_kernel->get_kernel_handle(); } - ze_module_handle_t get_module_handle() const { return m_kernel->get_module()->get_module_handle(); } + ze_kernel_handle_t get_kernel_handle() const { return m_kernel.handle(); } + ze_module_handle_t get_module_handle() const { return m_module.handle(); } std::string get_id() const override { return m_kernel_id; } std::shared_ptr clone(bool reuse_kernel_handle = false) const override { if (reuse_kernel_handle) { - return std::make_shared(m_kernel, m_kernel_id); + return std::make_shared(m_module, m_kernel, m_kernel_id, m_build_log); } else { ze_kernel_handle_t cloned_handle; ze_module_handle_t module_handle = get_module_handle(); @@ -62,8 +68,8 @@ class ze_kernel : public kernel { descriptor.flags = 0; descriptor.pKernelName = m_kernel_id.c_str(); OV_ZE_EXPECT(ze::zeKernelCreate(module_handle, &descriptor, &cloned_handle)); - auto kernel_holder = std::make_shared(cloned_handle, m_kernel->get_module()); - return std::make_shared(kernel_holder, m_kernel_id); + ze_kernel_resource kernel_holder(cloned_handle); + return std::make_shared(m_module, kernel_holder, m_kernel_id, m_build_log); } } @@ -87,7 +93,10 @@ class ze_kernel : public kernel { } std::string get_build_log() const override { - ze_module_build_log_handle_t build_log_handle = m_kernel->get_module()->get_build_log_handle(); + if (m_build_log.is_empty()) { + return {}; + } + ze_module_build_log_handle_t build_log_handle = m_build_log.handle(); size_t log_size = 0; OV_ZE_EXPECT(ze::zeModuleBuildLogGetString(build_log_handle, &log_size, nullptr)); @@ -97,7 +106,10 @@ class ze_kernel : public kernel { } private: - std::shared_ptr m_kernel; + // To ensure correct lifetime, drop kernel first then build log and module + ze_module_resource m_module; + ze_module_build_log_resource m_build_log; + ze_kernel_resource m_kernel; std::string m_kernel_id; }; diff --git a/src/plugins/intel_gpu/src/runtime/ze/ze_kernel_builder.cpp b/src/plugins/intel_gpu/src/runtime/ze/ze_kernel_builder.cpp index 9e339a0a29fb76..f8115951267d36 100644 --- a/src/plugins/intel_gpu/src/runtime/ze/ze_kernel_builder.cpp +++ b/src/plugins/intel_gpu/src/runtime/ze/ze_kernel_builder.cpp @@ -8,7 +8,7 @@ #include "ocl/ocl_kernel_builder.hpp" #include "ocl/ocl_device_detector.hpp" -#include +#include using namespace cldnn; using namespace ze; @@ -33,18 +33,19 @@ void ze_kernel_builder::init_ocl_builder() const { } bool ze_kernel_builder::check_ze_build_support() const { - static std::unordered_map cache; + static std::map cache; static std::mutex m; // Prevent multiple threads from checking support at the same time std::lock_guard lock(m); const char src[] = R"(__kernel void k(){})"; auto src_bytes = sizeof(src); - auto dev_handle = m_device.get_device(); + auto dev_handle = m_device.get_device().handle(); if (cache.find(dev_handle) != cache.end()) { return cache.at(dev_handle); } try { - build_module_ze(src, src_bytes, KernelFormat::SOURCE, ""); + std::vector tmp_out; + build_kernels_ze(src, src_bytes, KernelFormat::SOURCE, "", tmp_out); cache[dev_handle] = true; } catch (std::exception&) { GPU_DEBUG_INFO << "[GPU] Device(" << dev_handle << ") does not support kernel compilation from source through Level Zero" << std::endl; @@ -53,7 +54,7 @@ bool ze_kernel_builder::check_ze_build_support() const { return cache.at(dev_handle); } -std::shared_ptr ze_kernel_builder::build_module_ze(const void *src, size_t src_bytes, KernelFormat src_format, const std::string &options) const { +void ze_kernel_builder::build_kernels_ze(const void *src, size_t src_bytes, KernelFormat src_format, const std::string &options, std::vector &out) const { ze_module_desc_t module_desc = { ZE_STRUCTURE_TYPE_MODULE_DESC, nullptr, @@ -80,33 +81,36 @@ std::shared_ptr ze_kernel_builder::build_module_ze(const void } ze_module_handle_t module_handle; ze_module_build_log_handle_t log_handle; - ze_result_t build_result = ze::zeModuleCreate(m_device.get_context(), m_device.get_device(), &module_desc, &module_handle, &log_handle); + auto ctx_handle = m_device.get_context().handle(); + auto device_handle = m_device.get_device().handle(); + ze_result_t build_result = ze::zeModuleCreate(ctx_handle, device_handle, &module_desc, &module_handle, &log_handle); + ze_module_build_log_resource build_log_holder(log_handle); if (build_result != ZE_RESULT_SUCCESS) { size_t log_size = 0; OV_ZE_EXPECT(ze::zeModuleBuildLogGetString(log_handle, &log_size, nullptr)); std::string log(log_size, ' '); OV_ZE_EXPECT(ze::zeModuleBuildLogGetString(log_handle, &log_size, log.data())); - OV_ZE_EXPECT(ze::zeModuleBuildLogDestroy(log_handle)); GPU_DEBUG_INFO << "-------- Kernel build error" << std::endl; GPU_DEBUG_INFO << log << std::endl; GPU_DEBUG_INFO << "-------- End of Kernel build error" << std::endl; OPENVINO_THROW("[GPU] Failed to build module"); } - return std::make_shared(module_handle, log_handle); + ze_module_resource module_holder(module_handle); + ze_kernel::create_kernels_from_module(module_holder, build_log_holder, out); } -std::shared_ptr ze_kernel_builder::build_module_ocl(const void *src, size_t src_bytes, KernelFormat src_format, const std::string &options) const { +void ze_kernel_builder::build_kernels_ocl(const void *src, size_t src_bytes, KernelFormat src_format, const std::string &options, std::vector &out) const { OPENVINO_ASSERT(src_format == KernelFormat::SOURCE, "[GPU] Level Zero kernel builder should only fallback to OCL when building kernels from source"); OPENVINO_ASSERT(m_ocl_builder != nullptr, "[GPU] Level Zero kernel builder expected initialized OCL builder"); std::vector tmp; m_ocl_builder->build_kernels(src, src_bytes, src_format, options, tmp); OPENVINO_ASSERT(tmp.size() > 0, "[GPU] Level Zero kernel builder expected non-empty module"); auto binary = tmp[0]->get_binary(); - return build_module_ze(binary.data(), binary.size(), KernelFormat::NATIVE_BIN, options); + + build_kernels_ze(binary.data(), binary.size(), KernelFormat::NATIVE_BIN, options, out); } void ze_kernel_builder::build_kernels(const void *src, size_t src_bytes, KernelFormat src_format, const std::string &options, std::vector &out) const { - std::shared_ptr module_holder; if (src_format == KernelFormat::SOURCE && !check_ze_build_support()) { { std::lock_guard lock(this->m_mutex); @@ -115,9 +119,8 @@ void ze_kernel_builder::build_kernels(const void *src, size_t src_bytes, KernelF init_ocl_builder(); } } - module_holder = build_module_ocl(src, src_bytes, src_format, options); + build_kernels_ocl(src, src_bytes, src_format, options, out); } else { - module_holder = build_module_ze(src, src_bytes, src_format, options); + build_kernels_ze(src, src_bytes, src_format, options, out); } - ze_kernel::create_kernels_from_module(module_holder, out); } diff --git a/src/plugins/intel_gpu/src/runtime/ze/ze_kernel_builder.hpp b/src/plugins/intel_gpu/src/runtime/ze/ze_kernel_builder.hpp index 435e5b1517315e..5836b520807d85 100644 --- a/src/plugins/intel_gpu/src/runtime/ze/ze_kernel_builder.hpp +++ b/src/plugins/intel_gpu/src/runtime/ze/ze_kernel_builder.hpp @@ -27,10 +27,10 @@ class ze_kernel_builder : public kernel_builder{ private: /// @brief Check if ZE can build kernels from source bool check_ze_build_support() const; - /// @brief Build module through ZE API - std::shared_ptr build_module_ze(const void *src, size_t src_bytes, KernelFormat src_format, const std::string &options) const; - /// @brief Build module through OCL API and repackage to ZE module - std::shared_ptr build_module_ocl(const void *src, size_t src_bytes, KernelFormat src_format, const std::string &options) const; + /// @brief Build kernels through ZE API + void build_kernels_ze(const void *src, size_t src_bytes, KernelFormat src_format, const std::string &options, std::vector &out) const; + /// @brief Build kernels through OCL API and repackage to ZE module + void build_kernels_ocl(const void *src, size_t src_bytes, KernelFormat src_format, const std::string &options, std::vector &out) const; void init_ocl_builder() const; const ze_device &m_device; // OCL workaround for legacy devices that does not support ZE compilation diff --git a/src/plugins/intel_gpu/src/runtime/ze/ze_kernel_holder.hpp b/src/plugins/intel_gpu/src/runtime/ze/ze_kernel_holder.hpp deleted file mode 100644 index 3f449ca728bea0..00000000000000 --- a/src/plugins/intel_gpu/src/runtime/ze/ze_kernel_holder.hpp +++ /dev/null @@ -1,32 +0,0 @@ -// Copyright (C) 2018-2026 Intel Corporation -// SPDX-License-Identifier: Apache-2.0 -// - -#pragma once - -#include "ze_common.hpp" -#include "ze_module_holder.hpp" - -#include - -namespace cldnn { -namespace ze { - -// RAII wrapper for Level Zero kernel -class ze_kernel_holder { -public: - // Take ownership of existing kernel handle - explicit ze_kernel_holder(ze_kernel_handle_t kernel, std::shared_ptr module) : m_kernel(kernel), m_module(module) {} - ze_kernel_holder(const ze_kernel_holder& other) = delete; - ze_kernel_holder& operator=(const ze_kernel_holder& other) = delete; - ~ze_kernel_holder() { - OV_ZE_WARN(ze::zeKernelDestroy(m_kernel)); - } - ze_kernel_handle_t get_kernel_handle() { return m_kernel; } - std::shared_ptr get_module() { return m_module; } -private: - ze_kernel_handle_t m_kernel; - std::shared_ptr m_module; -}; -} // namespace ze -} // namespace cldnn diff --git a/src/plugins/intel_gpu/src/runtime/ze/ze_memory.cpp b/src/plugins/intel_gpu/src/runtime/ze/ze_memory.cpp index 32a1bc00b1c29b..d916d43bc00694 100644 --- a/src/plugins/intel_gpu/src/runtime/ze/ze_memory.cpp +++ b/src/plugins/intel_gpu/src/runtime/ze/ze_memory.cpp @@ -8,6 +8,7 @@ #include "ze_engine.hpp" #include "ze_stream.hpp" #include "ze_event.hpp" +#include "ze_resource_interop.hpp" #include "runtime_common.hpp" #include @@ -29,6 +30,17 @@ static inline cldnn::event::ptr create_event(stream& stream, size_t bytes_count) return stream.create_base_event(); } +bool check_allocation_range(ze_context_handle_t ctx, void *ptr, size_t expected_size) { + void *base = nullptr; + size_t allocation_size = 0; + OV_ZE_EXPECT(ze::zeMemGetAddressRange(ctx, ptr, &base, &allocation_size)); + const auto base_u = reinterpret_cast(base); + const auto ptr_u = reinterpret_cast(ptr); + const auto alloc_end_u = base_u + allocation_size; + const auto ptr_end_u = ptr_u + expected_size; + return (ptr_u >= base_u) && (ptr_end_u >= ptr_u) && (ptr_end_u <= alloc_end_u); +} + std::vector get_ze_events(const std::vector& events) { std::vector ze_events; ze_events.reserve(events.size()); @@ -41,6 +53,45 @@ std::vector get_ze_events(const std::vector& even return ze_events; } +ze_usm_resource allocate_usm_host(ze_context_resource context, size_t size) { + ze_host_mem_alloc_desc_t host_desc = {}; + host_desc.stype = ZE_STRUCTURE_TYPE_HOST_MEM_ALLOC_DESC; + host_desc.flags = 0; + host_desc.pNext = nullptr; + ov_ze_usm_handle usm_handle; + usm_handle.context = context.handle(); + OV_ZE_EXPECT(ze::zeMemAllocHost(usm_handle.context, &host_desc, size, 0, &usm_handle.ptr)); + return ze_usm_resource(usm_handle); +} + +ze_usm_resource allocate_usm_shared(ze_context_resource context, ze_device_resource device, size_t size, uint32_t ordinal) { + ze_device_mem_alloc_desc_t device_desc = {}; + device_desc.stype = ZE_STRUCTURE_TYPE_DEVICE_MEM_ALLOC_DESC; + device_desc.flags = 0; + device_desc.ordinal = ordinal; + device_desc.pNext = nullptr; + ze_host_mem_alloc_desc_t host_desc = {}; + host_desc.stype = ZE_STRUCTURE_TYPE_HOST_MEM_ALLOC_DESC; + host_desc.flags = 0; + host_desc.pNext = nullptr; + ov_ze_usm_handle usm_handle; + usm_handle.context = context.handle(); + OV_ZE_EXPECT(ze::zeMemAllocShared(usm_handle.context, &device_desc, &host_desc, size, 0, device.handle(), &usm_handle.ptr)); + return ze_usm_resource(usm_handle); +} + +ze_usm_resource allocate_usm_device(ze_context_resource context, ze_device_resource device, size_t size, uint32_t ordinal) { + ze_device_mem_alloc_desc_t device_desc = {}; + device_desc.stype = ZE_STRUCTURE_TYPE_DEVICE_MEM_ALLOC_DESC; + device_desc.flags = 0; + device_desc.ordinal = ordinal; + device_desc.pNext = nullptr; + ov_ze_usm_handle usm_handle; + usm_handle.context = context.handle(); + OV_ZE_EXPECT(ze::zeMemAllocDevice(usm_handle.context, &device_desc, size, 0, device.handle(), &usm_handle.ptr)); + return ze_usm_resource(usm_handle); +} + size_t get_element_size(const ze_image_format_layout_t& ze_layout) { switch(ze_layout) { case ZE_IMAGE_FORMAT_LAYOUT_8: @@ -95,54 +146,134 @@ std::pair get_width_height(const layout& layout) { break; } default: - OPENVINO_THROW("[GPU] 2D image allocation", "unsupported image type!"); + OPENVINO_THROW("[GPU] Unsupported image type!"); } return {width, height}; } + +cl_image_format get_cl_image_format(const layout &layout) { + cl_image_format image_format{}; + cl_channel_type &type = image_format.image_channel_data_type; + cl_channel_order &order = image_format.image_channel_order; + type = layout.data_type == data_types::f16 ? CL_HALF_FLOAT : CL_FLOAT; + order = CL_R; + switch (layout.format) { + case format::image_2d_weights_c4_fyx_b: + order = CL_RGBA; + break; + case format::image_2d_rgba: + order = CL_RGBA; + if (layout.feature() != 3 && layout.feature() != 4) { + OPENVINO_THROW("[GPU] Invalid number of channels in image_2d_rgba input image (should be 3 or 4)!"); + } + type = CL_UNORM_INT8; + break; + case format::nv12: + { + // [NHWC] dimensions order + auto shape = layout.get_shape(); + if (shape[3] == 2) { + order = CL_RG; + } else if (shape[3] > 2) { + OPENVINO_THROW("[GPU] Invalid number of channels in NV12 input image!"); + } + type = CL_UNORM_INT8; + break; + } + default: + OPENVINO_THROW("[GPU] Unexpected layout format"); + } + return image_format; +} + +cl_image_desc get_cl_image_desc(const layout &layout) { + cl_image_desc image_desc{CL_MEM_OBJECT_IMAGE2D, 0, 0, 0, 0, 0, 0, 0, 0, nullptr}; + auto &width = image_desc.image_width; + auto &height = image_desc.image_height; + switch (layout.format) { + case format::image_2d_weights_c1_b_fyx: + width = layout.batch(); + height = layout.spatial(0) * layout.feature() * layout.spatial(1); + break; + case format::image_2d_weights_winograd_6x3_s1_fbxyb: + height = layout.feature(); + width = layout.spatial(0) * layout.batch() * layout.spatial(1) * 8 / 3; + break; + case format::image_2d_weights_winograd_6x3_s1_xfbyb: + height = layout.feature() * layout.spatial(0) * 8 / 3; + width = layout.batch() * layout.spatial(1); + break; + case format::image_2d_weights_c4_fyx_b: + width = layout.batch(); + height = layout.spatial(0) * layout.feature() * layout.spatial(1); + break; + case format::image_2d_rgba: + width = layout.spatial(0); + height = layout.spatial(1); + break; + case format::nv12: + { + // [NHWC] dimensions order + auto shape = layout.get_shape(); + width = shape[2]; + height = shape[1]; + break; + } + default: + OPENVINO_THROW("[GPU] Unexpected layout format"); + } + return image_desc; +} + } // namespace allocation_type gpu_usm::detect_allocation_type(const ze_engine* engine, const void* mem_ptr) { ze_memory_allocation_properties_t props{ZE_STRUCTURE_TYPE_MEMORY_ALLOCATION_PROPERTIES}; ze_device_handle_t device = nullptr; - OV_ZE_EXPECT(ze::zeMemGetAllocProperties(engine->get_context(), mem_ptr, &props, &device)); + OV_ZE_EXPECT(ze::zeMemGetAllocProperties(engine->get_context().handle(), mem_ptr, &props, &device)); + allocation_type alloc_type = allocation_type::unknown; switch (props.type) { - case ZE_MEMORY_TYPE_DEVICE: return allocation_type::usm_device; - case ZE_MEMORY_TYPE_HOST: return allocation_type::usm_host; - case ZE_MEMORY_TYPE_SHARED: return allocation_type::usm_shared; - default: return allocation_type::unknown; + case ZE_MEMORY_TYPE_DEVICE: + alloc_type = allocation_type::usm_device; + break; + case ZE_MEMORY_TYPE_HOST: + alloc_type = allocation_type::usm_host; + break; + case ZE_MEMORY_TYPE_SHARED: + alloc_type = allocation_type::usm_shared; + break; + default: + alloc_type = allocation_type::unknown; + break; } - - return allocation_type::unknown; + return alloc_type; } -allocation_type gpu_usm::detect_allocation_type(const ze_engine* engine, const ze::UsmMemory& buffer) { - auto alloc_type = detect_allocation_type(engine, buffer.get()); - OPENVINO_ASSERT(alloc_type == allocation_type::usm_device || - alloc_type == allocation_type::usm_host || - alloc_type == allocation_type::usm_shared, "[GPU] Unsupported USM alloc type: " + to_string(alloc_type)); - return alloc_type; +allocation_type gpu_usm::detect_allocation_type(const ze_engine* engine, ze_usm_resource buffer) { + if (buffer.has_ocl_handle()) { + return allocation_type::cl_mem; + } + return detect_allocation_type(engine, buffer.handle().ptr); } -gpu_usm::gpu_usm(ze_engine* engine, const layout& new_layout, const ze::UsmMemory& buffer, allocation_type type, std::shared_ptr mem_tracker) +gpu_usm::gpu_usm(ze_engine* engine, const layout& new_layout, ze_usm_resource buffer, allocation_type type, std::shared_ptr mem_tracker) : lockable_gpu_mem() , memory(engine, new_layout, type, mem_tracker) - , _buffer(buffer) - , _host_buffer(engine->get_context(), engine->get_device()) { + , _buffer(std::move(buffer)) { + auto ctx_handle = engine->get_context().handle(); + auto ptr = _buffer.handle().ptr; + auto expected_size = new_layout.bytes_count(); + OPENVINO_ASSERT(check_allocation_range(ctx_handle, ptr, expected_size), + "[GPU] Allocation is smaller than the size required by the layout"); } -gpu_usm::gpu_usm(ze_engine* engine, const layout& new_layout, const ze::UsmMemory& buffer, std::shared_ptr mem_tracker) - : lockable_gpu_mem() - , memory(engine, new_layout, detect_allocation_type(engine, buffer), mem_tracker) - , _buffer(buffer) - , _host_buffer(engine->get_context(), engine->get_device()) { -} +gpu_usm::gpu_usm(ze_engine* engine, const layout& new_layout, ze_usm_resource buffer, std::shared_ptr mem_tracker) + : gpu_usm(engine, new_layout, std::move(buffer), detect_allocation_type(engine, buffer), mem_tracker) {} gpu_usm::gpu_usm(ze_engine* engine, const layout& layout, allocation_type type) : lockable_gpu_mem() - , memory(engine, layout, type, nullptr) - , _buffer(engine->get_context(), engine->get_device()) - , _host_buffer(engine->get_context(), engine->get_device()) { + , memory(engine, layout, type, nullptr) { auto actual_bytes_count = _bytes_count; if (actual_bytes_count == 0) actual_bytes_count = 1; @@ -150,49 +281,50 @@ gpu_usm::gpu_usm(ze_engine* engine, const layout& layout, allocation_type type) auto mem_ordinal = engine->get_device_info().device_memory_ordinal; switch (get_allocation_type()) { case allocation_type::usm_host: - _buffer.allocateHost(actual_bytes_count); + _buffer = allocate_usm_host(engine->get_context(), actual_bytes_count); break; case allocation_type::usm_shared: - _buffer.allocateShared(actual_bytes_count, mem_ordinal); + _buffer = allocate_usm_shared(engine->get_context(), engine->get_device(), actual_bytes_count, mem_ordinal); break; + case allocation_type::cl_mem: case allocation_type::usm_device: - _buffer.allocateDevice(actual_bytes_count, mem_ordinal); + _buffer = allocate_usm_device(engine->get_context(), engine->get_device(), actual_bytes_count, mem_ordinal); break; default: - OPENVINO_THROW("[GPU] Unknown unified shared memory type!"); + OPENVINO_THROW("[GPU] Requested unknown unified shared memory type!"); } - m_mem_tracker = std::make_shared(engine, _buffer.get(), actual_bytes_count, type); + m_mem_tracker = std::make_shared(engine, _buffer.handle().ptr, actual_bytes_count, type); } void* gpu_usm::lock(const stream& stream, mem_lock_type type) { std::lock_guard locker(_mutex); if (0 == _lock_count) { auto& _ze_stream = downcast(stream); - if (get_allocation_type() == allocation_type::usm_device) { + auto alloc_type = get_allocation_type(); + if (alloc_type == allocation_type::usm_device || alloc_type == allocation_type::cl_mem) { GPU_DEBUG_LOG << "Copy usm_device buffer to host buffer." << std::endl; - _host_buffer.allocateHost(_bytes_count); + auto *zero_engine = downcast(_engine); + _host_buffer = allocate_usm_host(zero_engine->get_context(), _bytes_count); // Always copy device data to host buffer (treat write as read_write internally). // This ensures the host buffer always has valid data, making nested locks safe. OV_ZE_EXPECT(zeCommandListAppendMemoryCopy(_ze_stream.get_queue(), - _host_buffer.get(), - _buffer.get(), + _host_buffer.handle().ptr, + _buffer.handle().ptr, _bytes_count, nullptr, 0, nullptr)); - OV_ZE_EXPECT(zeCommandListHostSynchronize(_ze_stream.get_queue(), endless_wait)); - _host_buffer_has_device_data = true; - _copy_back_to_device = (type != mem_lock_type::read); - _mapped_ptr = _host_buffer.get(); + OV_ZE_EXPECT(ze::zeCommandListHostSynchronize(_ze_stream.get_queue(), endless_wait)); + _mapped_ptr = _host_buffer.handle().ptr; } else { - _mapped_ptr = _buffer.get(); - } - } else if (get_allocation_type() == allocation_type::usm_device) { - if (type != mem_lock_type::read) { - _copy_back_to_device = true; + _mapped_ptr = _buffer.handle().ptr; } } + if (!_host_buffer.is_empty()) { + // Update write back flag on every lock call + _copy_back_to_device = _copy_back_to_device || (type != mem_lock_type::read); + } _lock_count++; return _mapped_ptr; } @@ -202,26 +334,30 @@ void gpu_usm::unlock(const stream& stream) { OPENVINO_ASSERT(_lock_count != 0, "[GPU] Trying to unlock an already unlocked buffer"); _lock_count--; if (0 == _lock_count) { - if (get_allocation_type() == allocation_type::usm_device) { - if (_copy_back_to_device) { + if (_copy_back_to_device) { auto& _ze_stream = downcast(stream); OV_ZE_EXPECT(zeCommandListAppendMemoryCopy(_ze_stream.get_queue(), - _buffer.get(), - _host_buffer.get(), + _buffer.handle().ptr, + _host_buffer.handle().ptr, _bytes_count, nullptr, 0, nullptr)); OV_ZE_EXPECT(zeCommandListHostSynchronize(_ze_stream.get_queue(), endless_wait)); - } - _host_buffer.freeMem(); - _copy_back_to_device = false; - _host_buffer_has_device_data = false; } + _copy_back_to_device = false; + _host_buffer.drop(); _mapped_ptr = nullptr; } } +void* gpu_usm::buffer_ptr() const { + if (_buffer.is_empty()) { + return nullptr; + } + return _buffer.handle().ptr; +} + event::ptr gpu_usm::fill(stream& stream, unsigned char pattern, const std::vector& dep_events, bool blocking) { auto& _ze_stream = downcast(stream); auto ev = _ze_stream.create_base_event(); @@ -229,7 +365,7 @@ event::ptr gpu_usm::fill(stream& stream, unsigned char pattern, const std::vecto auto ze_dep_events = get_ze_events(dep_events); const auto num_ze_dep_events = static_cast(ze_dep_events.size()); OV_ZE_EXPECT(ze::zeCommandListAppendMemoryFill(_ze_stream.get_queue(), - _buffer.get(), + _buffer.handle().ptr, &pattern, sizeof(unsigned char), _bytes_count, @@ -255,7 +391,7 @@ event::ptr gpu_usm::copy_from(stream& stream, const void* data_ptr, size_t src_o auto _ze_stream = downcast(&stream); auto _ze_event = downcast(result_event.get())->get_handle(); auto src_ptr = reinterpret_cast(data_ptr) + src_offset; - auto dst_ptr = reinterpret_cast(buffer_ptr()) + dst_offset; + auto dst_ptr = reinterpret_cast(_buffer.handle().ptr) + dst_offset; OV_ZE_EXPECT(ze::zeCommandListAppendMemoryCopy(_ze_stream->get_queue(), dst_ptr, @@ -280,11 +416,12 @@ event::ptr gpu_usm::copy_from(stream& stream, const memory& src_mem, size_t src_ auto _ze_stream = downcast(&stream); auto _ze_event = downcast(result_event.get())->get_handle(); - OPENVINO_ASSERT(memory_capabilities::is_usm_type(src_mem.get_allocation_type())); + auto alloc_type = src_mem.get_allocation_type(); + OPENVINO_ASSERT(memory_capabilities::is_usm_type(alloc_type) || alloc_type == allocation_type::cl_mem, "[GPU] Source memory for gpu_usm::copy_from(memory&) should be USM or OpenCL buffer"); auto usm_mem = downcast(&src_mem); auto src_ptr = reinterpret_cast(usm_mem->buffer_ptr()) + src_offset; - auto dst_ptr = reinterpret_cast(buffer_ptr()) + dst_offset; + auto dst_ptr = reinterpret_cast(_buffer.handle().ptr) + dst_offset; OV_ZE_EXPECT(ze::zeCommandListAppendMemoryCopy(_ze_stream->get_queue(), dst_ptr, @@ -329,7 +466,7 @@ event::ptr gpu_usm::copy_to(stream& stream, void* data_ptr, size_t src_offset, s dnnl::memory gpu_usm::get_onednn_memory(dnnl::memory::desc desc, int64_t offset) const { auto onednn_engine = _engine->get_onednn_engine(); dnnl::memory dnnl_mem = dnnl::ze_interop::make_memory(desc, onednn_engine, - reinterpret_cast(_buffer.get()) + offset); + reinterpret_cast(_buffer.handle().ptr) + offset); return dnnl_mem; } @@ -338,31 +475,53 @@ dnnl::memory gpu_usm::get_onednn_grouped_memory(dnnl::memory::desc desc, const m OPENVINO_ASSERT(memory_capabilities::is_usm_type(offsets.get_allocation_type())); OPENVINO_ASSERT(offsets.get_engine() == this->_engine); dnnl::memory dnnl_mem = dnnl::ze_interop::make_memory(desc, onednn_engine, - {reinterpret_cast(_buffer.get()), reinterpret_cast(offsets.buffer_ptr())}); + {reinterpret_cast(_buffer.handle().ptr), reinterpret_cast(offsets.buffer_ptr())}); return dnnl_mem; } #endif -shared_mem_params gpu_usm::get_internal_params() const { - auto casted = downcast(_engine); - return { - shared_mem_type::shared_mem_usm, // shared_mem_type - static_cast(casted->get_context()), // context handle - static_cast(casted->get_device()), // user_device handle - static_cast(_buffer.get()), // mem handle +shared_mem_params gpu_usm::get_internal_params(runtime_types rt_type) const { + auto params = shared_mem_params { + shared_mem_type::shared_mem_usm, + nullptr, + nullptr, + nullptr, #ifdef _WIN32 - nullptr, // surface handle + nullptr, #else - 0, // surface handle + 0, #endif - 0 // plane + 0 }; + auto zero_engine = downcast(_engine); + auto ctx_res = zero_engine->get_context(); + if (rt_type == runtime_types::ze) { + params.context = ctx_res.handle(); + params.mem = _buffer.handle().ptr; + return params; + } else if (rt_type == runtime_types::ocl) { + auto device_res = zero_engine->get_device(); + ze_export_ocl_context(ctx_res, device_res); + params.context = ctx_res.ocl_handle(); + if (get_allocation_type() == allocation_type::cl_mem) { + cl_mem_flags flags = 0; + size_t buffer_size = _bytes_count; + ze_export_ocl_mem(_buffer, ctx_res, device_res, flags, buffer_size); + params.mem_type = shared_mem_type::shared_mem_buffer; + params.mem = _buffer.ocl_handle(); + } else { + // No need to convert when exporting USM pointer + params.mem = _buffer.handle().ptr; + } + } else { + OPENVINO_THROW("[GPU] Unsupported runtime type for gpu_usm internal params"); + } + return params; } gpu_image2d::gpu_image2d(ze_engine* engine, const layout& layout) : lockable_gpu_mem() , memory(engine, layout, allocation_type::ze_image, nullptr) - , _host_buffer(engine->get_context(), engine->get_device()) , _width(0) , _height(0) { ze_image_desc_t image_desc = {}; @@ -450,19 +609,20 @@ gpu_image2d::gpu_image2d(ze_engine* engine, const layout& layout) image_desc.depth = 1; image_desc.arraylevels = 1; image_desc.miplevels = 0; + auto ctx_handle = engine->get_context().handle(); + auto device_handle = engine->get_device().handle(); ze_image_handle_t image_handle; - OV_ZE_EXPECT(ze::zeImageCreate(engine->get_context(), engine->get_device(), &image_desc, &image_handle)); - _image = std::make_shared(image_handle, false); + OV_ZE_EXPECT(ze::zeImageCreate(ctx_handle, device_handle, &image_desc, &image_handle)); + _image_holder = ze_image_resource(image_handle); size_t elem_size = get_element_size(image_desc.format.layout); - _bytes_count = elem_size * _width * _height; + OPENVINO_ASSERT(elem_size * _width * _height == layout.bytes_count(), "[GPU] Image size does not match layout bytes count"); m_mem_tracker = std::make_shared(engine, image_handle, layout.bytes_count(), allocation_type::ze_image); } -gpu_image2d::gpu_image2d(ze_engine* engine, const layout& new_layout, ze_image_handle_t image, std::shared_ptr mem_tracker) +gpu_image2d::gpu_image2d(ze_engine* engine, const layout& new_layout, ze_image_resource image, std::shared_ptr mem_tracker) : lockable_gpu_mem() , memory(engine, new_layout, allocation_type::ze_image, mem_tracker) - , _image(std::make_shared(image, true)) - , _host_buffer(engine->get_context(), engine->get_device()) { + , _image_holder(std::move(image)) { // No way to get width and height from Level Zero so we have to assume layout is correct std::tie(_width, _height) = get_width_height(new_layout); } @@ -476,12 +636,12 @@ void* gpu_image2d::lock(const stream& stream, mem_lock_type type) { } else { _needs_write_back = false; } - - _host_buffer.allocateHost(_bytes_count); + auto *zero_engine = downcast(_engine); + _host_buffer = allocate_usm_host(zero_engine->get_context(), _bytes_count); if (type != mem_lock_type::write) { OV_ZE_EXPECT(ze::zeCommandListAppendImageCopyToMemory(zero_stream.get_queue(), - _host_buffer.get(), - _image->get_handle(), + _host_buffer.handle().ptr, + _image_holder.handle(), nullptr, nullptr, 0, @@ -489,7 +649,7 @@ void* gpu_image2d::lock(const stream& stream, mem_lock_type type) { // Block thread and wait for copy and previous operations to finish OV_ZE_EXPECT(ze::zeCommandListHostSynchronize(zero_stream.get_queue(), endless_wait)); } - _mapped_ptr = _host_buffer.get(); + _mapped_ptr = _host_buffer.handle().ptr; } _lock_count++; return _mapped_ptr; @@ -504,8 +664,8 @@ void gpu_image2d::unlock(const stream& stream) { if (_needs_write_back) { auto& zero_stream = downcast(stream); OV_ZE_EXPECT(ze::zeCommandListAppendImageCopyFromMemory(zero_stream.get_queue(), - _image->get_handle(), - _host_buffer.get(), + _image_holder.handle(), + _host_buffer.handle().ptr, nullptr, nullptr, 0, @@ -513,7 +673,7 @@ void gpu_image2d::unlock(const stream& stream) { // Insert barrier to ensure that following commands have correct image data OV_ZE_EXPECT(ze::zeCommandListAppendBarrier(zero_stream.get_queue(), nullptr, 0, nullptr)); } - _host_buffer.freeMem(); + _host_buffer.drop(); _mapped_ptr = nullptr; } } @@ -530,13 +690,13 @@ event::ptr gpu_image2d::fill(stream& stream, unsigned char pattern, const std::v ze_dep_events.push_back(ev_fill_handle); // Level Zero does not have API to fill image directly // Workaround is to fill usm buffer and then copy it to image - auto context = zero_stream.get_engine().get_context(); auto device = zero_stream.get_engine().get_device(); - ze::UsmMemory fill_buffer(context, device); - fill_buffer.allocateDevice(_bytes_count, zero_stream.get_engine().get_device_info().device_memory_ordinal); + auto context = zero_stream.get_engine().get_context(); + auto mem_ordinal = zero_stream.get_engine().get_device_info().device_memory_ordinal; + ze_usm_resource fill_buffer = allocate_usm_device(context, device, _bytes_count, mem_ordinal); OV_ZE_EXPECT(ze::zeCommandListAppendMemoryFill(zero_stream.get_queue(), - fill_buffer.get(), + fill_buffer.handle().ptr, &pattern, sizeof(unsigned char), _bytes_count, @@ -545,8 +705,8 @@ event::ptr gpu_image2d::fill(stream& stream, unsigned char pattern, const std::v nullptr)); auto ev_result_handle = downcast(result_event.get())->get_handle(); OV_ZE_EXPECT(ze::zeCommandListAppendImageCopyFromMemory(zero_stream.get_queue(), - _image->get_handle(), - fill_buffer.get(), + _image_holder.handle(), + fill_buffer.handle().ptr, nullptr, ev_result_handle, static_cast(ze_dep_events.size()), @@ -563,19 +723,39 @@ event::ptr gpu_image2d::fill(stream& stream, unsigned char pattern, const std::v return result_event; } -shared_mem_params gpu_image2d::get_internal_params() const { - auto zero_engine = downcast(_engine); - return {shared_mem_type::shared_mem_image, static_cast(zero_engine->get_context()), nullptr, - static_cast(_image->get_handle()), +shared_mem_params gpu_image2d::get_internal_params(runtime_types rt_type) const { + auto params = shared_mem_params { + shared_mem_type::shared_mem_image, + nullptr, + nullptr, + nullptr, #ifdef _WIN32 nullptr, #else 0, #endif - 0}; + 0 + }; + auto zero_engine = downcast(_engine); + auto ctx_res = zero_engine->get_context(); + if (rt_type == runtime_types::ze) { + params.context = ctx_res.handle(); + params.mem = _image_holder.handle(); + return params; + } else if (rt_type == runtime_types::ocl) { + auto device_res = zero_engine->get_device(); + cl_mem_flags flags = 0; + cl_image_format img_fmt = get_cl_image_format(get_layout()); + cl_image_desc img_desc = get_cl_image_desc(get_layout()); + ze_export_ocl_image(_image_holder, ctx_res, device_res, flags, img_fmt, img_desc); + params.context = ctx_res.ocl_handle(); + params.mem = _image_holder.ocl_handle(); + } else { + OPENVINO_THROW("[GPU] Unsupported runtime type for gpu_image2d internal params"); + } + return params; } - event::ptr gpu_image2d::copy_from(stream& stream, const void* data_ptr, size_t src_offset, size_t dst_offset, size_t size, bool blocking) { auto result_event = create_event(stream, size); if (size == 0) @@ -589,7 +769,7 @@ event::ptr gpu_image2d::copy_from(stream& stream, const void* data_ptr, size_t s auto ev_result_handle = downcast(result_event.get())->get_handle(); OV_ZE_EXPECT(ze::zeCommandListAppendImageCopyFromMemory(zero_stream->get_queue(), - _image->get_handle(), + _image_holder.handle(), src_ptr, nullptr, ev_result_handle, @@ -614,8 +794,8 @@ event::ptr gpu_image2d::copy_from(stream& stream, const memory& src_mem, size_t auto src_image = downcast(&src_mem); auto ev_result_handle = downcast(result_event.get())->get_handle(); OV_ZE_EXPECT(ze::zeCommandListAppendImageCopy(zero_stream->get_queue(), - _image->get_handle(), - src_image->_image->get_handle(), + _image_holder.handle(), + src_image->_image_holder.handle(), ev_result_handle, 0, nullptr)); @@ -637,7 +817,7 @@ event::ptr gpu_image2d::copy_to(stream& stream, void* data_ptr, size_t src_offse auto ev_result_handle = downcast(result_event.get())->get_handle(); OV_ZE_EXPECT(ze::zeCommandListAppendImageCopyToMemory(zero_stream->get_queue(), dst_ptr, - _image->get_handle(), + _image_holder.handle(), nullptr, ev_result_handle, 0, diff --git a/src/plugins/intel_gpu/src/runtime/ze/ze_memory.hpp b/src/plugins/intel_gpu/src/runtime/ze/ze_memory.hpp index 5f2695df36c3f3..d0437b97c1f4c8 100644 --- a/src/plugins/intel_gpu/src/runtime/ze/ze_memory.hpp +++ b/src/plugins/intel_gpu/src/runtime/ze/ze_memory.hpp @@ -5,6 +5,7 @@ #pragma once #include "ze_common.hpp" +#include "ze_resource.hpp" #include "ze_engine.hpp" #include "ze_base_event.hpp" #include "intel_gpu/runtime/memory.hpp" @@ -19,124 +20,27 @@ struct lockable_gpu_mem { lockable_gpu_mem() : _lock_count(0), _mapped_ptr(nullptr), - _copy_back_to_device(false), - _host_buffer_has_device_data(false) {} + _copy_back_to_device(false) {} std::mutex _mutex; unsigned _lock_count; void* _mapped_ptr; bool _copy_back_to_device; - bool _host_buffer_has_device_data; -}; - -class UsmHolder { -public: - UsmHolder(ze_context_handle_t context, void* ptr, bool shared_memory = false) : _context(context), _ptr(ptr), _shared_memory(shared_memory) { - if (ptr == nullptr) - OPENVINO_THROW("[GPU] Can not create UsmHolder with nullptr"); - } - UsmHolder(const UsmHolder&) = delete; - UsmHolder& operator=(const UsmHolder&) = delete; - - void* ptr() { return _ptr; } - - ~UsmHolder() { - if (!_shared_memory) { - OV_ZE_WARN(ze::zeMemFree(_context, _ptr)); - } - } -private: - ze_context_handle_t _context; - void* _ptr; - bool _shared_memory = false; -}; - -class UsmMemory { -public: - explicit UsmMemory(ze_context_handle_t context, ze_device_handle_t device) - : _context(context) - , _device(device) {} - - UsmMemory(ze_context_handle_t context, ze_device_handle_t device, void* usm_ptr, size_t offset = 0) - : _context(context) - , _device(device) - , _usm_pointer(std::make_shared(_context, reinterpret_cast(usm_ptr) + offset, true)) {} - - void* get() const { - if (is_empty()) { - return nullptr; - } - return _usm_pointer->ptr(); - } - - bool is_empty() const { return _usm_pointer.get() == nullptr; } - - void allocateHost(size_t size) { - ze_host_mem_alloc_desc_t host_desc = {}; - host_desc.stype = ZE_STRUCTURE_TYPE_HOST_MEM_ALLOC_DESC; - host_desc.flags = 0; - host_desc.pNext = nullptr; - - void* memory = nullptr; - OV_ZE_EXPECT(ze::zeMemAllocHost(_context, &host_desc, size, 0, &memory)); - _usm_pointer = std::make_shared(_context, memory); - } - - void allocateShared(size_t size, uint32_t ordinal) { - ze_device_mem_alloc_desc_t device_desc = {}; - device_desc.stype = ZE_STRUCTURE_TYPE_DEVICE_MEM_ALLOC_DESC; - device_desc.flags = 0; - device_desc.ordinal = ordinal; - device_desc.pNext = nullptr; - - ze_host_mem_alloc_desc_t host_desc = {}; - host_desc.stype = ZE_STRUCTURE_TYPE_HOST_MEM_ALLOC_DESC; - host_desc.flags = 0; - host_desc.pNext = nullptr; - - void* memory = nullptr; - OV_ZE_EXPECT(ze::zeMemAllocShared(_context, &device_desc, &host_desc, size, 0, _device, &memory)); - _usm_pointer = std::make_shared(_context, memory); - } - - void allocateDevice(size_t size, uint32_t ordinal) { - ze_device_mem_alloc_desc_t device_desc = {}; - device_desc.stype = ZE_STRUCTURE_TYPE_DEVICE_MEM_ALLOC_DESC; - device_desc.flags = 0; - device_desc.ordinal = ordinal; - device_desc.pNext = nullptr; - - void* memory = nullptr; - OV_ZE_EXPECT(ze::zeMemAllocDevice(_context, &device_desc, size, 0, _device, &memory)); - _usm_pointer = std::make_shared(_context, memory); - } - - void freeMem() { - _usm_pointer.reset(); - } - - virtual ~UsmMemory() = default; - -protected: - ze_context_handle_t _context; - ze_device_handle_t _device; - std::shared_ptr _usm_pointer = nullptr; }; struct gpu_usm : public lockable_gpu_mem, public memory { - gpu_usm(ze_engine* engine, const layout& new_layout, const ze::UsmMemory& usm_buffer, allocation_type type, std::shared_ptr mem_tracker); - gpu_usm(ze_engine* engine, const layout& new_layout, const ze::UsmMemory& usm_buffer, std::shared_ptr mem_tracker); + gpu_usm(ze_engine* engine, const layout& new_layout, ze_usm_resource usm_buffer, allocation_type type, std::shared_ptr mem_tracker); + gpu_usm(ze_engine* engine, const layout& new_layout, ze_usm_resource usm_buffer, std::shared_ptr mem_tracker); gpu_usm(ze_engine* engine, const layout& layout, allocation_type type); void* lock(const stream& stream, mem_lock_type type) override; void unlock(const stream& stream) override; - const ze::UsmMemory& get_buffer() const { return _buffer; } - ze::UsmMemory& get_buffer() { return _buffer; } event::ptr fill(stream& stream, unsigned char pattern, const std::vector& dep_events = {}, bool blocking = true) override; event::ptr fill(stream& stream, const std::vector& dep_events = {}, bool blocking = true) override; - shared_mem_params get_internal_params() const override; - void* buffer_ptr() const override { return _buffer.get(); } + shared_mem_params get_internal_params(runtime_types rt_type) const override; + void* buffer_ptr() const override; + ze_usm_resource get_resource() const { return _buffer; } event::ptr copy_from(stream& stream, const void* data_ptr, size_t src_offset, size_t dst_offset, size_t size, bool blocking) override; event::ptr copy_from(stream& stream, const memory& src_mem, size_t src_offset, size_t dst_offset, size_t size, bool blocking) override; @@ -147,52 +51,34 @@ struct gpu_usm : public lockable_gpu_mem, public memory { #endif static allocation_type detect_allocation_type(const ze_engine* engine, const void* mem_ptr); - static allocation_type detect_allocation_type(const ze_engine* engine, const ze::UsmMemory& buffer); + static allocation_type detect_allocation_type(const ze_engine* engine, ze_usm_resource buffer); protected: - ze::UsmMemory _buffer; - ze::UsmMemory _host_buffer; -}; - -struct image_holder { -public: - image_holder(ze_image_handle_t buffer, bool is_shared = false) : _buffer(buffer), _is_shared(is_shared) { - OPENVINO_ASSERT(buffer != nullptr, "[GPU] Can not create image_holder with nullptr"); - } - image_holder(const image_holder&) = delete; - image_holder& operator=(const image_holder&) = delete; - ~image_holder() { - if (!_is_shared) { - OV_ZE_WARN(ze::zeImageDestroy(_buffer)); - } - } - - ze_image_handle_t get_handle() const { return _buffer; } -private: - ze_image_handle_t _buffer; - bool _is_shared; + mutable ze_usm_resource _buffer; + ze_usm_resource _host_buffer; }; struct gpu_image2d : public lockable_gpu_mem, public memory { - gpu_image2d(ze_engine* engine, const layout& new_layout, ze_image_handle_t image, std::shared_ptr mem_tracker); + gpu_image2d(ze_engine* engine, const layout& new_layout, ze_image_resource image, std::shared_ptr mem_tracker); gpu_image2d(ze_engine* engine, const layout& layout); void* lock(const stream& stream, mem_lock_type type = mem_lock_type::read_write) override; void unlock(const stream& stream) override; event::ptr fill(stream& stream, unsigned char pattern, const std::vector& dep_events = {}, bool blocking = true) override; - shared_mem_params get_internal_params() const override; + shared_mem_params get_internal_params(runtime_types rt_type) const override; ze_image_handle_t get_handle() const { OPENVINO_ASSERT(0 == _lock_count, "[GPU] Cannot get image handle when memory is locked"); - return _image->get_handle(); + return _image_holder.handle(); } + ze_image_resource get_resource() const { return _image_holder; } event::ptr copy_from(stream& stream, const void* data_ptr, size_t src_offset = 0, size_t dst_offset = 0, size_t size = 0, bool blocking = true) override; event::ptr copy_from(stream& stream, const memory& src_mem, size_t src_offset = 0, size_t dst_offset = 0, size_t size = 0, bool blocking = true) override; event::ptr copy_to(stream& stream, void* data_ptr, size_t src_offset = 0, size_t dst_offset = 0, size_t size = 0, bool blocking = true) const override; protected: - std::shared_ptr _image; - ze::UsmMemory _host_buffer; + mutable ze_image_resource _image_holder; + ze_usm_resource _host_buffer; size_t _width; size_t _height; bool _needs_write_back; diff --git a/src/plugins/intel_gpu/src/runtime/ze/ze_module_holder.hpp b/src/plugins/intel_gpu/src/runtime/ze/ze_module_holder.hpp deleted file mode 100644 index f640e66795e58d..00000000000000 --- a/src/plugins/intel_gpu/src/runtime/ze/ze_module_holder.hpp +++ /dev/null @@ -1,34 +0,0 @@ -// Copyright (C) 2018-2026 Intel Corporation -// SPDX-License-Identifier: Apache-2.0 -// - -#pragma once - -#include "ze_common.hpp" -#include -#include - -namespace cldnn { -namespace ze { - -// RAII wrapper for Level Zero module -class ze_module_holder { -public: - // Take ownership of existing module and build log handles - explicit ze_module_holder(ze_module_handle_t module, ze_module_build_log_handle_t build_log) : m_module(module), m_build_log(build_log) {} - - ze_module_holder(const ze_module_holder& other) = delete; - ze_module_holder& operator=(const ze_module_holder& other) = delete; - ~ze_module_holder() { - OV_ZE_WARN(ze::zeModuleBuildLogDestroy(m_build_log)); - OV_ZE_WARN(ze::zeModuleDestroy(m_module)); - } - ze_module_handle_t get_module_handle() const { return m_module; } - ze_module_build_log_handle_t get_build_log_handle() const { return m_build_log; } - -private: - ze_module_handle_t m_module; - ze_module_build_log_handle_t m_build_log; -}; -} // namespace ze -} // namespace cldnn diff --git a/src/plugins/intel_gpu/src/runtime/ze/ze_ocl_common.hpp b/src/plugins/intel_gpu/src/runtime/ze/ze_ocl_common.hpp new file mode 100644 index 00000000000000..605e68b2525b21 --- /dev/null +++ b/src/plugins/intel_gpu/src/runtime/ze/ze_ocl_common.hpp @@ -0,0 +1,21 @@ +// Copyright (C) 2018-2026 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +#pragma once + +#include "openvino/core/except.hpp" +#include "ocl/ocl_wrapper.hpp" + +#include +#include + +// Prints warning if OpenCL command does not return success result +#define OV_OCL_WARN(f) \ + do { \ + cl_int res_ = (f); \ + if (res_ != CL_SUCCESS) { \ + GPU_DEBUG_INFO << "[Warning] [GPU] " #f " command failed with code " \ + << std::to_string(res_); \ + } \ + } while (false) diff --git a/src/plugins/intel_gpu/src/runtime/ze/ze_ocl_interop.cpp b/src/plugins/intel_gpu/src/runtime/ze/ze_ocl_interop.cpp new file mode 100644 index 00000000000000..bb1e7ddd732b9c --- /dev/null +++ b/src/plugins/intel_gpu/src/runtime/ze/ze_ocl_interop.cpp @@ -0,0 +1,216 @@ +// Copyright (C) 2018-2026 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +#include "ze_ocl_interop.hpp" + +#include +#include + +namespace cldnn { +namespace ze { +namespace { + +static constexpr cl_uint CL_L0_CONTEXT_HANDLE = 0x42B0; +static constexpr cl_uint CL_L0_IMMEDIATE_CMD_LIST_HANDLE = 0x42B1; +static constexpr cl_uint CL_L0_MEM_OBJ_HANDLE = 0x42B3; +static constexpr cl_uint CL_L0_DEVICE_HANDLE = 0x42B4; + +inline void expect_success(cl_int error, void *handle, const std::string& message) { + if (error != CL_SUCCESS || handle == nullptr) { + OPENVINO_THROW(message, + " (Error code: ", std::to_string(error), + ", returned handle=", std::to_string((uintptr_t)handle), ")"); + } +} + +std::vector find_ocl_devices() { + cl_int error; + cl_uint numPlatforms = 0; + error = clGetPlatformIDs(0, nullptr, &numPlatforms); + if (error != CL_SUCCESS || numPlatforms == 0) { + return {}; + } + std::vector platform_ids(numPlatforms); + error = clGetPlatformIDs(numPlatforms, platform_ids.data(), nullptr); + if (error != CL_SUCCESS) { + return {}; + } + std::vector device_ids; + for (auto &platform : platform_ids) { + cl_uint num_devices = 0; + error = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 0, nullptr, &num_devices); + if (error != CL_SUCCESS || num_devices == 0) { + continue; + } + size_t old_size = device_ids.size(); + size_t new_size = old_size + num_devices; + device_ids.resize(new_size); + error = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, num_devices, device_ids.data() + old_size, nullptr); + if (error != CL_SUCCESS) { + continue; + } + } + return device_ids; +} + +std::vector> find_ze_devices() { + ze_init_driver_type_desc_t desc = { ZE_STRUCTURE_TYPE_INIT_DRIVER_TYPE_DESC, nullptr, ZE_INIT_DRIVER_TYPE_FLAG_GPU }; + ze_result_t error; + uint32_t driver_count = 0; + error = zeInitDrivers(&driver_count, nullptr, &desc); + if (error != ZE_RESULT_SUCCESS || driver_count == 0) { + return {}; + } + std::vector drivers(driver_count); + error = zeInitDrivers(&driver_count, drivers.data(), &desc); + if (error != ZE_RESULT_SUCCESS) { + return {}; + } + std::vector> all_devices; + for (auto &driver : drivers) { + uint32_t device_count = 0; + error = zeDeviceGet(driver, &device_count, nullptr); + if (error != ZE_RESULT_SUCCESS || device_count == 0) { + continue; + } + std::vector devices(device_count); + error = zeDeviceGet(driver, &device_count, devices.data()); + if (error != ZE_RESULT_SUCCESS) { + continue; + } + for (auto &device : devices) { + all_devices.emplace_back(driver, device); + } + } + return all_devices; +} + +} // namespace + +bool ze_ocl_interop::check_support(ze_device_handle_t ze_device) const { + try { + auto ocl_dev = find_ocl_device(ze_device); + return (ocl_dev != nullptr); + } + catch(...) { + return false; + } +}; + +void ze_ocl_interop::init() { + // Initialize Level Zero and find devices + auto ze_devices = find_ze_devices(); + // Initialize OpenCL and find devices + auto ocl_devices = find_ocl_devices(); + for (auto &ze_device : ze_devices) { + _device_map[ze_device.second] = nullptr; + } + for (auto &ocl_device : ocl_devices) { + ze_device_handle_t ze_dev = nullptr; + try { + ze_dev = get_ze_device(ocl_device); + } catch (...) { + continue; + } + _device_map[ze_dev] = ocl_device; + } +} + +cl_device_id ze_ocl_interop::find_ocl_device(ze_device_handle_t ze_device) const { + auto it = _device_map.find(ze_device); + if (it == _device_map.end() || it->second == nullptr) { + OPENVINO_THROW("[GPU] Failed to find matching OCL device for given ZE device"); + } + return it->second; +} + +ze_driver_handle_t ze_ocl_interop::find_ze_driver(cl_device_id ocl_device) const { + auto ze_device = get_ze_device(ocl_device); + auto all_devices = find_ze_devices(); + for (auto &device : all_devices) { + if (device.second == ze_device) { + return device.first; + } + } + OPENVINO_THROW("[GPU] Failed to find matching ZE driver for given ZE device"); +} + +ze_context_handle_t ze_ocl_interop::get_ze_context(cl_context context) const { + cl_int error; + ze_context_handle_t ze_context = nullptr; + error = clGetContextInfo(context, CL_L0_CONTEXT_HANDLE, sizeof(ze_context_handle_t), &ze_context, nullptr); + expect_success(error, ze_context, "[GPU] Attempt to extract ZE context from OCL context failed"); + return ze_context; +} + +ze_command_list_handle_t ze_ocl_interop::get_ze_cmd_list(cl_command_queue queue) const { + cl_int error; + ze_command_list_handle_t ze_cmd_list; + error = clGetCommandQueueInfo(queue, CL_L0_IMMEDIATE_CMD_LIST_HANDLE, sizeof(ze_command_list_handle_t), &ze_cmd_list, nullptr); + expect_success(error, ze_cmd_list, "[GPU] Attempt to extract ZE command list from OCL command queue failed"); + return ze_cmd_list; +} + +void* ze_ocl_interop::get_ze_usm(cl_mem ocl_mem) const { + cl_mem mem = reinterpret_cast(ocl_mem); + cl_int error; + void *ze_mem; + error = clGetMemObjectInfo(mem, CL_L0_MEM_OBJ_HANDLE, sizeof(void*), &ze_mem, nullptr); + expect_success(error, ze_mem, "[GPU] Attempt to extract ZE usm pointer from OCL buffer failed"); + return ze_mem; +} + +ze_image_handle_t ze_ocl_interop::get_ze_image(cl_mem ocl_mem) const { + cl_mem mem = reinterpret_cast(ocl_mem); + cl_int error; + ze_image_handle_t ze_image; + error = clGetMemObjectInfo(mem, CL_L0_MEM_OBJ_HANDLE, sizeof(ze_image_handle_t), &ze_image, nullptr); + expect_success(error, ze_image, "[GPU] Attempt to extract ZE image from OCL buffer failed"); + return ze_image; +} + +ze_device_handle_t ze_ocl_interop::get_ze_device(cl_device_id device) const { + ze_device_handle_t ze_device; + cl_int error = clGetDeviceInfo(device, CL_L0_DEVICE_HANDLE, sizeof(ze_device_handle_t), &ze_device, nullptr); + expect_success(error, ze_device, "[GPU] Attempt to extract ZE device from OCL device failed"); + return ze_device; +} + +cl_context ze_ocl_interop::create_cl_context(ze_context_handle_t context, const ocl_context_args& args) const { + cl_context_properties properties[] = {CL_L0_CONTEXT_HANDLE, reinterpret_cast(context), 0}; + constexpr cl_uint num_devices = 1; + cl_int error; + auto converted_context = clCreateContext(properties, num_devices, &args.device, nullptr, nullptr, &error); + expect_success(error, converted_context, "[GPU] Attempt to create CL context from ZE context failed"); + return converted_context; + +} + +cl_command_queue ze_ocl_interop::create_cl_queue(ze_command_list_handle_t cmd_list, const ocl_queue_args& args) const { + // Works only for immediate command lists + cl_queue_properties properties[] = {CL_L0_IMMEDIATE_CMD_LIST_HANDLE, reinterpret_cast(cmd_list), 0}; + cl_int error; + auto converted_queue = clCreateCommandQueueWithProperties(args.context, args.device, properties, &error); + expect_success(error, converted_queue, "[GPU] Attempt to create CL command queue from ZE command list failed"); + return converted_queue; +} + +cl_mem ze_ocl_interop::create_cl_buffer(void* usm_ptr, const ocl_buffer_args& args) const { + cl_mem_properties properties[] = {CL_L0_MEM_OBJ_HANDLE, reinterpret_cast(usm_ptr), 0}; + cl_int error; + cl_mem converted_mem = clCreateBufferWithProperties(args.context, properties, args.flags, args.size, nullptr, &error); + expect_success(error, converted_mem, "[GPU] Attempt to create CL buffer from ZE USM failed"); + return converted_mem; +} + +cl_mem ze_ocl_interop::create_cl_image(ze_image_handle_t image, const ocl_image_args &args) const { + cl_mem_properties properties[] = {CL_L0_MEM_OBJ_HANDLE, reinterpret_cast(image), 0}; + cl_int error; + cl_mem converted_mem = clCreateImageWithProperties(args.context, properties, args.flags, &args.format, &args.desc, nullptr, &error); + expect_success(error, converted_mem, "[GPU] Attempt to create CL image from ZE image failed"); + return converted_mem; +} + +} // namespace ze +} // namespace cldnn diff --git a/src/plugins/intel_gpu/src/runtime/ze/ze_ocl_interop.hpp b/src/plugins/intel_gpu/src/runtime/ze/ze_ocl_interop.hpp new file mode 100644 index 00000000000000..7f9ccb98f355bd --- /dev/null +++ b/src/plugins/intel_gpu/src/runtime/ze/ze_ocl_interop.hpp @@ -0,0 +1,92 @@ +// Copyright (C) 2018-2026 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +#pragma once + +#include "ze_common.hpp" +#include "ze_ocl_common.hpp" + +#include +#include + +namespace cldnn { +namespace ze { +struct ocl_context_args { + cl_device_id device; +}; +struct ocl_queue_args { + cl_context context; + cl_device_id device; +}; + +struct ocl_buffer_args { + cl_context context; + cl_mem_flags flags; + size_t size; +}; + +struct ocl_image_args { + cl_context context; + cl_mem_flags flags; + cl_image_format format; + cl_image_desc desc; +}; + +/// @brief Provides functions to convert between Level Zero and OpenCL handles. +/// +/// Converting from OpenCL to Level Zero generally works by extracting internal Level Zero handles from OpenCL objects. +/// Converting from Level Zero to OpenCL generally works by creating new OpenCL objects and passing to them Level Zero handles. +/// Both approaches enable user to use the same entity through Level Zero and OpenCL APIs. +/// +/// Functions named "create_*" work by creating new objects. +/// Handles returned by these functions must be released by the caller using the appropriate release functions. +/// +/// Functions named "get_*" work by extracting data from already existing objects. +/// Handles returned by these functions are borrowed and should not be released. +/// +/// Functions named "find_*" work by matching already existing objects. +/// Handles returned by these functions should not be released. +/// +/// Conversion functions require both Level Zero and OpenCL to be initialized beforehand. +/// All functions will either succeed and return a valid resource handle or throw an exception. +class ze_ocl_interop { +public: + static ze_ocl_interop& get_instance() { + static ze_ocl_interop instance; + static std::once_flag init_flag; + std::call_once(init_flag, []() { + instance.init(); + }); + + return instance; + } + ze_ocl_interop(const ze_ocl_interop&) = delete; + ze_ocl_interop& operator=(const ze_ocl_interop&) = delete; + + bool check_support(ze_device_handle_t ze_device) const; + + cl_device_id find_ocl_device(ze_device_handle_t ze_device) const; + ze_driver_handle_t find_ze_driver(cl_device_id ocl_device) const; + + ze_context_handle_t get_ze_context(cl_context context) const; + ze_command_list_handle_t get_ze_cmd_list(cl_command_queue queue) const; + void* get_ze_usm(cl_mem ocl_mem) const; + ze_image_handle_t get_ze_image(cl_mem ocl_mem) const; + ze_device_handle_t get_ze_device(cl_device_id device) const; + + cl_context create_cl_context(ze_context_handle_t context, const ocl_context_args& args) const; + cl_command_queue create_cl_queue(ze_command_list_handle_t cmd_list, const ocl_queue_args& args) const; + cl_mem create_cl_buffer(void* usm_ptr, const ocl_buffer_args& args) const; + cl_mem create_cl_image(ze_image_handle_t image, const ocl_image_args& args) const; +protected: + ze_ocl_interop() = default; + /// @brief Initialize OpenCL and L0 drivers and fill device_map + void init(); + + // L0 driver/device handles and OCL platform/device ids are global read only objects so we can cache them + std::map _device_map; +}; + +} // namespace ze +} // namespace cldnn diff --git a/src/plugins/intel_gpu/src/runtime/ze/ze_ocl_owner.hpp b/src/plugins/intel_gpu/src/runtime/ze/ze_ocl_owner.hpp new file mode 100644 index 00000000000000..eef51446ebdb5d --- /dev/null +++ b/src/plugins/intel_gpu/src/runtime/ze/ze_ocl_owner.hpp @@ -0,0 +1,180 @@ +// Copyright (C) 2018-2026 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +#pragma once + +#include "ze_owner.hpp" +#include "ze_ocl_common.hpp" + +#include +#include +#include +#include +#include + +namespace cldnn { +namespace ze { + +/// @brief Defines supported OpenCL resources for Level Zero interoperability +enum class ocl_resource_type : uint8_t { + platform, + device, + context, + command_queue, + mem_object, +}; + +/// @brief Provides information about specific OpenCL resource +template +struct ocl_resource_info { + static_assert(false, "Specialization for given resource type is not implemented"); +}; + +template <> +struct ocl_resource_info { + static constexpr ocl_resource_type resource = ocl_resource_type::platform; + using handle_t = cl_platform_id; + // Platform resource is not managed + struct deleter_t { + void operator()(handle_t handle) const noexcept {} + }; +}; + +template <> +struct ocl_resource_info { + static constexpr ocl_resource_type resource = ocl_resource_type::device; + using handle_t = cl_device_id; + // Device resource is not managed + struct deleter_t { + void operator()(handle_t handle) const noexcept {} + }; +}; + +template <> +struct ocl_resource_info { + static constexpr ocl_resource_type resource = ocl_resource_type::context; + using handle_t = cl_context; + struct deleter_t { + void operator()(handle_t handle) const noexcept { + OV_OCL_WARN(clReleaseContext(handle)); + } + }; +}; + +template <> +struct ocl_resource_info { + static constexpr ocl_resource_type resource = ocl_resource_type::command_queue; + using handle_t = cl_command_queue; + struct deleter_t { + void operator()(handle_t handle) const noexcept { + OV_OCL_WARN(clReleaseCommandQueue(handle)); + } + }; +}; + +template <> +struct ocl_resource_info { + static constexpr ocl_resource_type resource = ocl_resource_type::mem_object; + using handle_t = cl_mem; + struct deleter_t { + void operator()(handle_t handle) const noexcept { + OV_OCL_WARN(clReleaseMemObject(handle)); + } + }; +}; + +/// @brief Resource owner for OpenCL resources. +/// @tparam _resource_type OpenCL resource type managed by this owner. + +template +using ocl_owner = resource_owner>; + +/// @brief Resource owner for Level Zero resources and corresponding OpenCL resources. +/// +/// OpenCL resources are optional and can be attached after owner is created. +/// When destroyed owner will release non-shared OpenCL resources in declaration order before destroying Level Zero resource. +/// @tparam _resource_type Level Zero resource type managed by this owner. +/// @tparam ...ocl_resource_types Types of OpenCL resources that can be attached. +template +struct ze_ocl_owner_impl { + using ptr = std::shared_ptr; + using ze_handle_t = typename ze_resource_info<_resource_type>::handle_t; + ze_ocl_owner_impl(ze_handle_t ze_handle, bool is_borrowed = false) : _ze_owner(ze_handle, is_borrowed) {} + + ze_handle_t handle() const { + return _ze_owner.get_handle(); + } + + template + typename ocl_resource_info::handle_t ocl_handle() const { + static_assert((((ocl_resource_type == ocl_resource_types) || ...)), "Specified OCL resource type can not be obtained from this owner"); + auto &owner = std::get>>(_ocl_owners); + OPENVINO_ASSERT(owner.has_value(), "[GPU] Attempted to get ocl handle that is not attached to resource owner"); + return owner->get_handle(); + } + + template + void attach_ocl_handle(typename ocl_resource_info::handle_t handle, bool is_borrowed = false) { + static_assert((((ocl_resource_type == ocl_resource_types) || ...)), "Specified OCL resource type can not be attached to this owner"); + auto &owner = std::get>>(_ocl_owners); + OPENVINO_ASSERT(!owner.has_value(), "[GPU] Attempted to overwrite existing ocl handle attached to resource owner"); + owner.emplace(handle, is_borrowed); + } + + template + void attach_ocl_handle(ocl_owner &&moved_owner) { + static_assert((((ocl_resource_type == ocl_resource_types) || ...)), "Specified OCL resource type can not be attached to this owner"); + auto &owner = std::get>>(_ocl_owners); + OPENVINO_ASSERT(!owner.has_value(), "[GPU] Attempted to overwrite existing ocl handle attached to resource owner"); + owner.emplace(std::move(moved_owner)); + } + + template + bool has_ocl_handle() const { + static_assert((((ocl_resource_type == ocl_resource_types) || ...)), "Specified OCL resource type can not be obtained from this owner"); + auto &owner = std::get>>(_ocl_owners); + return owner.has_value(); + } + +private: + ze_owner<_resource_type> _ze_owner; + std::tuple>...> _ocl_owners; +}; + +/// @brief Resource owner for Level Zero resources and corresponding OpenCL resources. +/// +/// Only specific combinations of Level Zero and OpenCL resources are allowed, see specializations below. +/// @tparam _resource_type Level Zero resource type managed by this owner. +template +struct ze_ocl_owner : public ze_ocl_owner_impl<_resource_type> { + using ze_ocl_owner_impl<_resource_type>::ze_ocl_owner_impl; +}; + +template <> +struct ze_ocl_owner : public ze_ocl_owner_impl { + using ze_ocl_owner_impl::ze_ocl_owner_impl; +}; +template <> +struct ze_ocl_owner : public ze_ocl_owner_impl { + using ze_ocl_owner_impl::ze_ocl_owner_impl; +}; +template <> +struct ze_ocl_owner : public ze_ocl_owner_impl { + using ze_ocl_owner_impl::ze_ocl_owner_impl; +}; +template <> +struct ze_ocl_owner : public ze_ocl_owner_impl { + using ze_ocl_owner_impl::ze_ocl_owner_impl; +}; +template <> +struct ze_ocl_owner : public ze_ocl_owner_impl { + using ze_ocl_owner_impl::ze_ocl_owner_impl; +}; +template <> +struct ze_ocl_owner : public ze_ocl_owner_impl { + using ze_ocl_owner_impl::ze_ocl_owner_impl; +}; + +} // namespace ze +} // namespace cldnn diff --git a/src/plugins/intel_gpu/src/runtime/ze/ze_owner.hpp b/src/plugins/intel_gpu/src/runtime/ze/ze_owner.hpp new file mode 100644 index 00000000000000..339ade2edf0189 --- /dev/null +++ b/src/plugins/intel_gpu/src/runtime/ze/ze_owner.hpp @@ -0,0 +1,232 @@ +// Copyright (C) 2018-2026 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// +#pragma once + +#include "ze_common.hpp" + +#include +#include + +namespace cldnn { +namespace ze { + +/// @brief Defines supported Level Zero resources +enum class ze_resource_type : uint8_t { + driver, + device, + context, + command_queue, + command_list, + module, + kernel, + event_pool, + event, + image, + fence, + module_build_log, + usm_memory, +}; + +/// @brief Provides information about specific Level Zero resource +template +struct ze_resource_info { + static_assert(false, "Specialization for given resource type is not implemented"); +}; + +template <> +struct ze_resource_info { + static constexpr ze_resource_type resource = ze_resource_type::driver; + using handle_t = ze_driver_handle_t; + // Driver is a read only global construct that is not released + struct deleter_t { + void operator()(handle_t handle) const {} + }; +}; +template <> +struct ze_resource_info { + static constexpr ze_resource_type resource = ze_resource_type::device; + using handle_t = ze_device_handle_t; + // Device is a read only global construct that is not released + struct deleter_t { + void operator()(handle_t handle) const {} + }; +}; +template <> +struct ze_resource_info { + static constexpr ze_resource_type resource = ze_resource_type::context; + using handle_t = ze_context_handle_t; + struct deleter_t { + void operator()(handle_t handle) const { + OV_ZE_WARN(zeContextDestroy(handle)); + } + }; +}; + +template <> +struct ze_resource_info { + static constexpr ze_resource_type resource = ze_resource_type::command_queue; + using handle_t = ze_command_queue_handle_t; + struct deleter_t { + void operator()(handle_t handle) const { + OV_ZE_WARN(zeCommandQueueDestroy(handle)); + } + }; +}; + +template <> +struct ze_resource_info { + static constexpr ze_resource_type resource = ze_resource_type::command_list; + using handle_t = ze_command_list_handle_t; + struct deleter_t { + void operator()(handle_t handle) const { + OV_ZE_WARN(zeCommandListDestroy(handle)); + } + }; +}; + +template <> +struct ze_resource_info { + static constexpr ze_resource_type resource = ze_resource_type::module; + using handle_t = ze_module_handle_t; + struct deleter_t { + void operator()(handle_t handle) const { + OV_ZE_WARN(zeModuleDestroy(handle)); + } + }; +}; + +template <> +struct ze_resource_info { + static constexpr ze_resource_type resource = ze_resource_type::kernel; + using handle_t = ze_kernel_handle_t; + struct deleter_t { + void operator()(handle_t handle) const { + OV_ZE_WARN(zeKernelDestroy(handle)); + } + }; +}; + +template <> +struct ze_resource_info { + static constexpr ze_resource_type resource = ze_resource_type::event_pool; + using handle_t = ze_event_pool_handle_t; + struct deleter_t { + void operator()(handle_t handle) const { + OV_ZE_WARN(zeEventPoolDestroy(handle)); + } + }; +}; + +template <> +struct ze_resource_info { + static constexpr ze_resource_type resource = ze_resource_type::event; + using handle_t = ze_event_handle_t; + struct deleter_t { + void operator()(handle_t handle) const { + OV_ZE_WARN(zeEventDestroy(handle)); + } + }; +}; + +template <> +struct ze_resource_info { + static constexpr ze_resource_type resource = ze_resource_type::image; + using handle_t = ze_image_handle_t; + struct deleter_t { + void operator()(handle_t handle) const { + OV_ZE_WARN(zeImageDestroy(handle)); + } + }; +}; + +template <> +struct ze_resource_info { + static constexpr ze_resource_type resource = ze_resource_type::fence; + using handle_t = ze_fence_handle_t; + struct deleter_t { + void operator()(handle_t handle) const { + OV_ZE_WARN(zeFenceDestroy(handle)); + } + }; +}; + +template <> +struct ze_resource_info { + static constexpr ze_resource_type resource = ze_resource_type::module_build_log; + using handle_t = ze_module_build_log_handle_t; + struct deleter_t { + void operator()(handle_t handle) const { + OV_ZE_WARN(zeModuleBuildLogDestroy(handle)); + } + }; +}; + +struct ov_ze_usm_handle { + ze_context_handle_t context; + void* ptr; +}; + +template <> +struct ze_resource_info { + static constexpr ze_resource_type resource = ze_resource_type::usm_memory; + using handle_t = ov_ze_usm_handle; + struct deleter_t { + void operator()(handle_t handle) const { + OV_ZE_WARN(zeMemFree(handle.context, handle.ptr)); + } + }; +}; + +/// @brief Generic RAII owner that releases resources when destroyed. All resources should have exactly one owner. +/// @tparam _resource_info_t Provides information on how to handle resource. Must define appropriate handle_t and deleter_t types. +template +struct resource_owner { +public: + using ptr = std::shared_ptr; + using handle_t = typename _resource_info_t::handle_t; + using deleter_t = typename _resource_info_t::deleter_t; + + explicit resource_owner(handle_t handle, bool is_borrowed = false) : _handle(handle), _is_borrowed(is_borrowed) {} + resource_owner(const resource_owner& other) = delete; + resource_owner& operator=(const resource_owner& other) = delete; + resource_owner(resource_owner&& other) { + _handle = other._handle; + _is_borrowed = other._is_borrowed; + other._is_borrowed = true; // Mark as borrowed to prevent double free + } + resource_owner& operator=(resource_owner&& other) { + if (this != &other) { + release(); + _handle = other._handle; + _is_borrowed = other._is_borrowed; + other._is_borrowed = true; // Mark as borrowed to prevent double free + } + return *this; + } + ~resource_owner() { + release(); + } + + handle_t get_handle() const { + return _handle; + } +private: + void release() { + if (!_is_borrowed) { + deleter_t{}(_handle); + } + } + + handle_t _handle; + bool _is_borrowed; +}; + + +/// @brief Resource owner for Level Zero resources. +/// @tparam _resource_type Level Zero resource type managed by this owner. +template +using ze_owner = resource_owner>; + +} // namespace ze +} // namespace cldnn diff --git a/src/plugins/intel_gpu/src/runtime/ze/ze_resource.hpp b/src/plugins/intel_gpu/src/runtime/ze/ze_resource.hpp new file mode 100644 index 00000000000000..ee40b2be2e4e3e --- /dev/null +++ b/src/plugins/intel_gpu/src/runtime/ze/ze_resource.hpp @@ -0,0 +1,115 @@ +// Copyright (C) 2018-2026 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +#pragma once + +#include "ze_ocl_owner.hpp" + +namespace cldnn { +namespace ze { + +/// @brief Level Zero resource that can optionally have attached OpenCL resources. +/// +/// Copies of this class objects share the same resources. Resource are released when the last copy is destroyed. +/// @tparam _resource_type Level Zero resource type managed by this class. +template +struct ze_resource { + static constexpr ze_resource_type resource_type = _resource_type; + using ze_handle_t = typename ze_resource_info::handle_t; + using ze_ocl_owner_t = ze_ocl_owner; + + ze_resource() = default; + ze_resource(const ze_resource &) = default; + ze_resource(ze_resource&&) = default; + ze_resource& operator=(const ze_resource&) = default; + ze_resource& operator=(ze_resource&&) = default; + + /// @brief Create resource from existing holder. + /// @param holder Resource holder. + explicit ze_resource(typename ze_ocl_owner_t::ptr holder) + : _holder(holder) {} + + /// @brief Create ze_resource from Level Zero handle. Assumes passed handle is valid. + /// @param ze_handle Valid Level Zero object handle. + /// @param is_borrowed if false, takes ownership of the handle. + explicit ze_resource(ze_handle_t ze_handle, bool is_borrowed = false) + : _holder(std::make_shared(ze_handle, is_borrowed)) {} + + /// @brief Get Level Zero handle or throw if resource is empty. + ze_handle_t handle() const { + OPENVINO_ASSERT(_holder != nullptr, "[GPU] Attempted to get Level Zero handle from empty resource"); + return _holder->handle(); + } + + /// @brief Get OpenCL handle or throw if resource is empty. + template + typename ocl_resource_info::handle_t ocl_handle() const { + OPENVINO_ASSERT(_holder != nullptr, "[GPU] Attempted to get OpenCL handle from empty resource"); + return _holder->template ocl_handle(); + } + + /// @brief Attach OpenCL handle to the resource or throw if resource is empty. Assumes passed handle is valid. + /// + /// This function won't release passed handle in case exception is thrown. + template + void attach_ocl_handle(typename ocl_resource_info::handle_t handle, bool is_borrowed = false) { + OPENVINO_ASSERT(_holder != nullptr, "[GPU] Attempted to attach OpenCL handle to empty resource"); + return _holder->template attach_ocl_handle(handle, is_borrowed); + } + + /// @brief Attach OpenCL handle to the resource or throw if resource is empty. Assumes passed handle is valid. + /// + /// This function will release passed handle in case exception is thrown. + template + void attach_ocl_handle(ocl_owner &&owner) { + OPENVINO_ASSERT(_holder != nullptr, "[GPU] Attempted to attach OpenCL handle to empty resource"); + return _holder->template attach_ocl_handle(std::move(owner)); + } + + /// @brief Check if resource has specific OpenCL handle. + template + bool has_ocl_handle() const { + if (_holder == nullptr) { + return false; + } + return _holder->template has_ocl_handle(); + } + + /// @brief Get resource holder. Note that holder might be nullptr + typename ze_ocl_owner_t::ptr get_holder() const { + return _holder; + } + + /// @brief Drop resources and reset to empty state + void drop() { + _holder.reset(); + } + + /// @brief Returns true if object is not managing any Level Zero resource, false otherwise. + bool is_empty() const { + return _holder == nullptr; + } + +private: + typename ze_ocl_owner_t::ptr _holder; +}; + +// Aliases for convenience + +using ze_driver_resource = ze_resource; +using ze_device_resource = ze_resource; +using ze_context_resource = ze_resource; +using ze_command_queue_resource = ze_resource; +using ze_command_list_resource = ze_resource; +using ze_module_resource = ze_resource; +using ze_kernel_resource = ze_resource; +using ze_event_pool_resource = ze_resource; +using ze_event_resource = ze_resource; +using ze_image_resource = ze_resource; +using ze_fence_resource = ze_resource; +using ze_module_build_log_resource = ze_resource; +using ze_usm_resource = ze_resource; + +} // namespace ze +} // namespace cldnn diff --git a/src/plugins/intel_gpu/src/runtime/ze/ze_resource_interop.cpp b/src/plugins/intel_gpu/src/runtime/ze/ze_resource_interop.cpp new file mode 100644 index 00000000000000..fcd12939cbe260 --- /dev/null +++ b/src/plugins/intel_gpu/src/runtime/ze/ze_resource_interop.cpp @@ -0,0 +1,159 @@ +// Copyright (C) 2018-2026 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +#include "ze_resource_interop.hpp" + +namespace cldnn::ze { +ze_driver_resource ze_import_driver(cl_device_id ocl_device) { + ze_ocl_interop& interop = ze_ocl_interop::get_instance(); + auto ze_handle = interop.find_ze_driver(ocl_device); + const bool is_borrowed = true; + ze_driver_resource resource(ze_handle, is_borrowed); + + cl_platform_id platform; + cl_int error = clGetDeviceInfo(ocl_device, CL_DEVICE_PLATFORM, sizeof(cl_platform_id), &platform, nullptr); + OPENVINO_ASSERT(error == CL_SUCCESS, + "[GPU] Failed to get OpenCL platform from device handle (Error code: ", + std::to_string(error), + ")"); + resource.attach_ocl_handle(platform, is_borrowed); + return resource; +} + +ze_device_resource ze_import_device(cl_device_id ocl_device) { + ze_ocl_interop& interop = ze_ocl_interop::get_instance(); + auto ze_handle = interop.get_ze_device(ocl_device); + const bool is_borrowed = true; + ze_device_resource resource(ze_handle, is_borrowed); + resource.attach_ocl_handle(ocl_device, is_borrowed); + return resource; +} + +ze_context_resource ze_import_context(cl_context ocl_context) { + ze_ocl_interop& interop = ze_ocl_interop::get_instance(); + auto ze_handle = interop.get_ze_context(ocl_context); + const bool is_borrowed = true; + ze_context_resource resource(ze_handle, is_borrowed); + resource.attach_ocl_handle(ocl_context, is_borrowed); + return resource; +} + +ze_command_list_resource ze_import_command_list(cl_command_queue ocl_command_queue) { + ze_ocl_interop& interop = ze_ocl_interop::get_instance(); + auto ze_handle = interop.get_ze_cmd_list(ocl_command_queue); + const bool is_borrowed = true; + ze_command_list_resource resource(ze_handle, is_borrowed); + resource.attach_ocl_handle(ocl_command_queue, is_borrowed); + return resource; +} + +ze_usm_resource ze_import_usm(cl_mem ocl_buffer, ze_context_resource context) { + ze_ocl_interop& interop = ze_ocl_interop::get_instance(); + auto ze_usm_ptr = interop.get_ze_usm(ocl_buffer); + const bool is_borrowed = true; + ze_usm_resource resource({context.handle(), ze_usm_ptr}, is_borrowed); + resource.attach_ocl_handle(ocl_buffer, is_borrowed); + return resource; +} + +ze_image_resource ze_import_image(cl_mem ocl_image) { + auto& interop = ze_ocl_interop::get_instance(); + auto ze_handle = interop.get_ze_image(ocl_image); + ze_image_resource resource(ze_handle, true); + resource.attach_ocl_handle(ocl_image, true); + return resource; +} + +void ze_export_ocl_device(ze_device_resource device) { + if (device.has_ocl_handle()) { + return; + } + auto& interop = ze_ocl_interop::get_instance(); + auto ze_handle = device.handle(); + auto ocl_handle = interop.find_ocl_device(ze_handle); + ocl_owner ocl_owner(ocl_handle); + device.attach_ocl_handle(std::move(ocl_owner)); +} + +void ze_export_ocl_context(ze_context_resource context, ze_device_resource device) { + if (context.has_ocl_handle()) { + return; + } + + ze_export_ocl_device(device); + ocl_context_args context_args; + context_args.device = device.ocl_handle(); + + auto& interop = ze_ocl_interop::get_instance(); + auto ze_handle = context.handle(); + auto ocl_handle = interop.create_cl_context(ze_handle, context_args); + ocl_owner ocl_owner(ocl_handle); + context.attach_ocl_handle(std::move(ocl_owner)); +} + +void ze_export_ocl_command_queue(ze_command_list_resource cmd_list, + ze_context_resource context, + ze_device_resource device) { + if (cmd_list.has_ocl_handle()) { + return; + } + + ze_export_ocl_context(context, device); + ocl_queue_args queue_args; + queue_args.device = device.ocl_handle(); + queue_args.context = context.ocl_handle(); + + auto& interop = ze_ocl_interop::get_instance(); + auto ze_handle = cmd_list.handle(); + auto ocl_handle = interop.create_cl_queue(ze_handle, queue_args); + ocl_owner ocl_owner(ocl_handle); + cmd_list.attach_ocl_handle(std::move(ocl_owner)); +} + +void ze_export_ocl_mem(ze_usm_resource usm, + ze_context_resource context, + ze_device_resource device, + cl_mem_flags flags, + size_t size) { + if (usm.has_ocl_handle()) { + return; + } + + ze_export_ocl_context(context, device); + ocl_buffer_args buffer_args; + buffer_args.context = context.ocl_handle(); + buffer_args.flags = flags; + buffer_args.size = size; + + auto& interop = ze_ocl_interop::get_instance(); + auto ze_handle = usm.handle(); + auto ocl_handle = interop.create_cl_buffer(ze_handle.ptr, buffer_args); + ocl_owner ocl_owner(ocl_handle); + usm.attach_ocl_handle(std::move(ocl_owner)); +} + +void ze_export_ocl_image(ze_image_resource image, + ze_context_resource context, + ze_device_resource device, + cl_mem_flags flags, + cl_image_format format, + cl_image_desc desc) { + if (image.has_ocl_handle()) { + return; + } + + ze_export_ocl_context(context, device); + ocl_image_args image_args; + image_args.context = context.ocl_handle(); + image_args.flags = flags; + image_args.format = format; + image_args.desc = desc; + + auto& interop = ze_ocl_interop::get_instance(); + auto ze_handle = image.handle(); + auto ocl_handle = interop.create_cl_image(ze_handle, image_args); + ocl_owner ocl_owner(ocl_handle); + image.attach_ocl_handle(std::move(ocl_owner)); +} +} // namespace cldnn::ze diff --git a/src/plugins/intel_gpu/src/runtime/ze/ze_resource_interop.hpp b/src/plugins/intel_gpu/src/runtime/ze/ze_resource_interop.hpp new file mode 100644 index 00000000000000..16a8031230ed19 --- /dev/null +++ b/src/plugins/intel_gpu/src/runtime/ze/ze_resource_interop.hpp @@ -0,0 +1,100 @@ +// Copyright (C) 2018-2026 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +/** + * @brief Provide import/export functions for selected ze_resource objects. + * + * Wraps basic interop operations from ze_ocl_interop.hpp so that it works with ze_resource. + * @file ze_resource_interop.hpp + */ + +#pragma once + +#include "ze_ocl_interop.hpp" +#include "ze_resource.hpp" + +namespace cldnn::ze { +/// @brief Import Level Zero driver resource from OpenCL device. +/// @param ocl_device OpenCL device handle. +/// @return Level Zero driver resource extracted from OpenCL device handle. +ze_driver_resource ze_import_driver(cl_device_id ocl_device); + +/// @brief Import Level Zero device resource from OpenCL device. +/// @param ocl_device OpenCL device handle. +/// @return Level Zero device resource extracted from OpenCL device handle. +ze_device_resource ze_import_device(cl_device_id ocl_device); + +/// @brief Import Level Zero context resource from OpenCL context. +/// @param ocl_context OpenCL context handle. +/// @return Level Zero context resource extracted from OpenCL context handle. +ze_context_resource ze_import_context(cl_context ocl_context); + +/// @brief Import Level Zero immediate command list resource from OpenCL command queue. +/// @param ocl_command_queue OpenCL command queue handle. +/// @return Level Zero immediate command list resource extracted from OpenCL command queue handle. +ze_command_list_resource ze_import_command_list(cl_command_queue ocl_command_queue); + +/// @brief Import Level Zero USM (Unified Shared Memory) resource from OpenCL memory object. +/// @param ocl_buffer OpenCL memory object handle. +/// @param context Level Zero context resource. +/// @return Level Zero USM resource extracted from OpenCL memory object handle. +ze_usm_resource ze_import_usm(cl_mem ocl_buffer, ze_context_resource context); + +/// @brief Import Level Zero image resource from OpenCL image object. +/// @param ocl_image OpenCL image object handle. +/// @return Level Zero image resource extracted from OpenCL image object handle. +ze_image_resource ze_import_image(cl_mem ocl_image); + +/// @brief Exports Level Zero device to OpenCL device. Does nothing if resource already has OpenCL handle. +/// @param device Device resource to export and attach OpenCL handle to. +void ze_export_ocl_device(ze_device_resource device); + +/// @brief Exports Level Zero context to OpenCL context. Does nothing if resource already has OpenCL handle. +/// @param context Context resource to export and attach OpenCL handle to. +/// @param device Device resource to use for OpenCL context creation. If device does not have OpenCL handle, it will be +/// exported first. +void ze_export_ocl_context(ze_context_resource context, ze_device_resource device); + +/// @brief Exports Level Zero command list to OpenCL command queue. Does nothing if resource already has OpenCL handle. +/// @param cmd_list Command list resource to export and attach OpenCL handle to. +/// @param context Context resource to use for OpenCL command queue creation. If context does not have OpenCL handle, it +/// will be exported first. +/// @param device Device resource to use for OpenCL command queue creation. If device does not have OpenCL handle, it +/// will be exported first. +void ze_export_ocl_command_queue(ze_command_list_resource cmd_list, + ze_context_resource context, + ze_device_resource device); + +/// @brief Exports Level Zero USM (Unified Shared Memory) to OpenCL memory object. Does nothing if resource already has +/// OpenCL handle. +/// @param usm USM resource to export and attach OpenCL handle to. +/// @param context Context resource to use for OpenCL memory object creation. If context does not have OpenCL handle, it +/// will be exported first. +/// @param device Device resource to use for OpenCL memory object creation. If device does not have OpenCL handle, it +/// will be exported first. +/// @param flags OpenCL memory flags for the memory object. +/// @param size Size of the memory object. +void ze_export_ocl_mem(ze_usm_resource usm, + ze_context_resource context, + ze_device_resource device, + cl_mem_flags flags, + size_t size); + +/// @brief Exports Level Zero image to OpenCL image object. Does nothing if resource already has OpenCL handle. +/// @param image Image resource to export and attach OpenCL handle to. +/// @param context Context resource to use for OpenCL image object creation. If context does not have OpenCL handle, it +/// will be exported first. +/// @param device Device resource to use for OpenCL image object creation. If device does not have OpenCL handle, it +/// will be exported first. +/// @param flags OpenCL memory flags for the image object. +/// @param format OpenCL image format for the image object. +/// @param desc OpenCL image descriptor for the image object. +void ze_export_ocl_image(ze_image_resource image, + ze_context_resource context, + ze_device_resource device, + cl_mem_flags flags, + cl_image_format format, + cl_image_desc desc); + +} // namespace cldnn::ze diff --git a/src/plugins/intel_gpu/src/runtime/ze/ze_stream.cpp b/src/plugins/intel_gpu/src/runtime/ze/ze_stream.cpp index bdc2253ea327a3..4907aaf3ed76dc 100644 --- a/src/plugins/intel_gpu/src/runtime/ze/ze_stream.cpp +++ b/src/plugins/intel_gpu/src/runtime/ze/ze_stream.cpp @@ -31,6 +31,9 @@ #include #endif +// Required for queue type detection +#include "ze_ocl_common.hpp" + namespace cldnn { namespace ze { @@ -71,13 +74,12 @@ ze_result_t set_kernel_arg(ze_kernel_handle_t& kernel, uint32_t idx, cldnn::memo auto handle = image.get_handle(); GPU_DEBUG_TRACE_DETAIL << "kernel: " << kernel << " set arg (image) " << idx << " mem: " << handle << " size: " << mem->size() << std::endl; return ze::zeKernelSetArgumentValue(kernel, idx, sizeof(handle), &handle); - } else if (memory_capabilities::is_usm_type(mem->get_allocation_type())) { + } else if (memory_capabilities::is_usm_type(mem->get_allocation_type()) || mem->get_allocation_type() == allocation_type::cl_mem) { auto &usm = downcast(*mem); - const auto& buf = usm.get_buffer(); + auto ptr = usm.buffer_ptr(); auto mem_type = usm.get_allocation_type(); GPU_DEBUG_TRACE_DETAIL << "kernel: " << kernel << " set arg (" << mem_type << ") " << idx - << " mem: " << buf.get() << " size: " << mem->size() << std::endl; - auto ptr = buf.get(); + << " mem: " << ptr << " size: " << mem->size() << std::endl; return ze::zeKernelSetArgumentValue(kernel, idx, sizeof(ptr), &ptr); } else { return ZE_RESULT_ERROR_INVALID_ARGUMENT; @@ -193,6 +195,17 @@ void set_arguments_impl(ze_kernel_handle_t kernel, } } +QueueTypes detect_queue_type(ze_command_list_resource cmd_list) { + OPENVINO_ASSERT(cmd_list.has_ocl_handle(), "[GPU] Queue type detection requires OpenCL handle"); + auto queue = cmd_list.ocl_handle(); + cl_command_queue_properties properties; + + auto status = clGetCommandQueueInfo(queue, CL_QUEUE_PROPERTIES, sizeof(cl_command_queue_properties), &properties, nullptr); + OPENVINO_ASSERT(status == CL_SUCCESS, "[GPU] clGetCommandQueueInfo failed with error: ", status); + + return (properties & CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE) ? QueueTypes::out_of_order : QueueTypes::in_order; +} + } // namespace ze_stream::ze_stream(const ze_engine &engine, const ExecutionConfig& config) @@ -220,7 +233,12 @@ ze_stream::ze_stream(const ze_engine &engine, const ExecutionConfig& config) command_queue_desc.pNext = &cp_offload_desc; } - OV_ZE_EXPECT(ze::zeCommandListCreateImmediate(_engine.get_context(), _engine.get_device(), &command_queue_desc, &m_command_list)); + auto ctx_handle = engine.get_context().handle(); + auto device_handle = engine.get_device().handle(); + ze_command_list_handle_t cmd_list = nullptr; + OV_ZE_EXPECT(ze::zeCommandListCreateImmediate(ctx_handle, device_handle, &command_queue_desc, &cmd_list)); + m_cmd_list = ze_command_list_resource(cmd_list); + bool use_counter_based_events = m_queue_type == QueueTypes::in_order && info.supports_counter_based_events; m_user_ev_factory = std::make_shared(engine, config.get_enable_profiling()); if (use_counter_based_events) { @@ -236,13 +254,31 @@ ze_stream::ze_stream(const ze_engine &engine, const ExecutionConfig& config) << ")" << std::endl; } +ze_stream::ze_stream(const ze_engine& engine, const ExecutionConfig& config, ze_command_list_resource cmd_list) + : stream(detect_queue_type(cmd_list), stream::get_expected_sync_method(config)) + , _engine(engine) + , m_cmd_list(std::move(cmd_list)) { + const auto &info = engine.get_device_info(); + bool use_counter_based_events = m_queue_type == QueueTypes::in_order && info.supports_counter_based_events; + + m_user_ev_factory = std::make_shared(engine, config.get_enable_profiling()); + if (use_counter_based_events) { + m_ev_factory = std::make_shared(engine, config.get_enable_profiling()); + } else { + m_ev_factory = m_user_ev_factory; + } + GPU_DEBUG_INFO << "[GPU] Created L0 stream from existing command list (" + << "use_counter_based_events=" << use_counter_based_events + << ")" << std::endl; + } + + ze_stream::~ze_stream() { #ifdef ENABLE_ONEDNN_FOR_GPU - // Destroy OneDNN stream before destroying command list + // Destroy OneDNN stream before dropping command list _onednn_stream.reset(); #endif - if (m_command_list != nullptr) - OV_ZE_WARN(ze::zeCommandListDestroy(m_command_list)); + m_cmd_list.drop(); } void ze_stream::set_arguments(kernel& kernel, const kernel_arguments_desc& args_desc, const kernel_arguments_data& args) { @@ -283,7 +319,7 @@ event::ptr ze_stream::enqueue_kernel(kernel& kernel, auto local = to_group_count(args_desc.workGroups.local); ze_group_count_t args = { global.groupCountX / local.groupCountX, global.groupCountY / local.groupCountY, global.groupCountZ / local.groupCountZ }; OV_ZE_EXPECT(ze::zeKernelSetGroupSize(kern, local.groupCountX, local.groupCountY, local.groupCountZ)); - OV_ZE_EXPECT(ze::zeCommandListAppendLaunchKernel(m_command_list, + OV_ZE_EXPECT(ze::zeCommandListAppendLaunchKernel(m_cmd_list.handle(), kern, &args, set_output_event ? std::dynamic_pointer_cast(ev)->get_handle() : nullptr, @@ -294,13 +330,13 @@ event::ptr ze_stream::enqueue_kernel(kernel& kernel, } void ze_stream::enqueue_barrier() { - OV_ZE_EXPECT(ze::zeCommandListAppendBarrier(m_command_list, nullptr, 0, nullptr)); + OV_ZE_EXPECT(ze::zeCommandListAppendBarrier(m_cmd_list.handle(), nullptr, 0, nullptr)); } event::ptr ze_stream::enqueue_marker(std::vector const& deps, bool is_output) { if (deps.empty()) { auto ev = create_base_event(); - OV_ZE_EXPECT(ze::zeCommandListAppendBarrier(m_command_list, std::dynamic_pointer_cast(ev)->get_handle(), 0, nullptr)); + OV_ZE_EXPECT(ze::zeCommandListAppendBarrier(m_cmd_list.handle(), std::dynamic_pointer_cast(ev)->get_handle(), 0, nullptr)); return ev; } @@ -316,7 +352,7 @@ event::ptr ze_stream::enqueue_marker(std::vector const& deps, boo return create_user_event(true); auto ev = create_base_event(); - OV_ZE_EXPECT(ze::zeCommandListAppendBarrier(m_command_list, + OV_ZE_EXPECT(ze::zeCommandListAppendBarrier(m_cmd_list.handle(), std::dynamic_pointer_cast(ev)->get_handle(), static_cast(dep_events.size()), &dep_events.front())); @@ -360,7 +396,7 @@ void ze_stream::flush() const { } void ze_stream::finish() const { - OV_ZE_EXPECT(ze::zeCommandListHostSynchronize(m_command_list, endless_wait)); + OV_ZE_EXPECT(ze::zeCommandListHostSynchronize(m_cmd_list.handle(), endless_wait)); } void ze_stream::wait_for_events(const std::vector& events) { @@ -395,9 +431,9 @@ void ze_stream::sync_events(std::vector const& deps, bool is_output) if (is_output) { m_last_barrier_ev = std::dynamic_pointer_cast(create_base_event()); m_last_barrier_ev->set_queue_stamp(m_queue_counter.load()); - OV_ZE_EXPECT(ze::zeCommandListAppendBarrier(m_command_list, m_last_barrier_ev->get_handle(), 0, nullptr)); + OV_ZE_EXPECT(ze::zeCommandListAppendBarrier(m_cmd_list.handle(), m_last_barrier_ev->get_handle(), 0, nullptr)); } else { - OV_ZE_EXPECT(ze::zeCommandListAppendBarrier(m_command_list, nullptr, 0, nullptr)); + OV_ZE_EXPECT(ze::zeCommandListAppendBarrier(m_cmd_list.handle(), nullptr, 0, nullptr)); } m_last_barrier = ++m_queue_counter; } @@ -408,7 +444,7 @@ void ze_stream::sync_events(std::vector const& deps, bool is_output) } } -ze_context_handle_t ze_stream::get_context() const { +ze_context_resource ze_stream::get_context() const { return _engine.get_context(); } @@ -417,7 +453,7 @@ dnnl::stream& ze_stream::get_onednn_stream() { OPENVINO_ASSERT(m_queue_type == QueueTypes::in_order, "[GPU] Can't create onednn stream handle as onednn doesn't support out-of-order queue"); OPENVINO_ASSERT(_engine.get_device_info().vendor_id == INTEL_VENDOR_ID, "[GPU] Can't create onednn stream handle as for non-Intel devices"); if (!_onednn_stream) { - _onednn_stream = std::make_shared(dnnl::ze_interop::make_stream(_engine.get_onednn_engine(), m_command_list, m_ev_factory->is_profiling_enabled())); + _onednn_stream = std::make_shared(dnnl::ze_interop::make_stream(_engine.get_onednn_engine(), m_cmd_list.handle(), m_ev_factory->is_profiling_enabled())); } return *_onednn_stream; diff --git a/src/plugins/intel_gpu/src/runtime/ze/ze_stream.hpp b/src/plugins/intel_gpu/src/runtime/ze/ze_stream.hpp index ff13337d14e39b..f402aa0c610704 100644 --- a/src/plugins/intel_gpu/src/runtime/ze/ze_stream.hpp +++ b/src/plugins/intel_gpu/src/runtime/ze/ze_stream.hpp @@ -7,6 +7,7 @@ #include "intel_gpu/runtime/event.hpp" #include "intel_gpu/runtime/stream.hpp" #include "ze_common.hpp" +#include "ze_resource.hpp" #include "ze_engine.hpp" #include "ze_event.hpp" #include "ze_base_event_factory.hpp" @@ -16,20 +17,20 @@ namespace ze { class ze_stream : public stream { public: - ze_command_list_handle_t get_queue() const { return m_command_list; } + ze_command_list_handle_t get_queue() const { return m_cmd_list.handle(); } const ze_engine& get_engine() const { return _engine; } ze_stream(const ze_engine& engine, const ExecutionConfig& config); + ze_stream(const ze_engine& engine, const ExecutionConfig& config, ze_command_list_resource cmd_list); ze_stream(ze_stream&& other) : stream(other.m_queue_type, other.m_sync_method) , _engine(other._engine) - , m_command_list(other.m_command_list) + , m_cmd_list(std::move(other.m_cmd_list)) , m_queue_counter(other.m_queue_counter.load()) , m_last_barrier(other.m_last_barrier.load()) , m_last_barrier_ev(other.m_last_barrier_ev) , m_ev_factory(std::move(other.m_ev_factory)) , m_user_ev_factory(std::move(other.m_user_ev_factory)) { - other.m_command_list = nullptr; } ~ze_stream(); @@ -51,7 +52,7 @@ class ze_stream : public stream { event::ptr create_user_event(bool set) override; event::ptr create_base_event() override; std::unique_ptr create_surfaces_lock(const std::vector &mem) const override; - ze_context_handle_t get_context() const; + ze_context_resource get_context() const; #ifdef ENABLE_ONEDNN_FOR_GPU dnnl::stream& get_onednn_stream() override; @@ -61,7 +62,7 @@ class ze_stream : public stream { void sync_events(std::vector const& deps, bool is_output = false); const ze_engine& _engine; - mutable ze_command_list_handle_t m_command_list = 0; + ze_command_list_resource m_cmd_list; mutable std::atomic m_queue_counter{0}; std::atomic m_last_barrier{0}; std::shared_ptr m_last_barrier_ev = nullptr; diff --git a/src/plugins/intel_gpu/tests/functional/remote_tensor_tests/ocl_remote_tensor_tests.cpp b/src/plugins/intel_gpu/tests/functional/remote_tensor_tests/ocl_remote_tensor_tests.cpp index 122b7e604135ff..71fa38993f0bce 100644 --- a/src/plugins/intel_gpu/tests/functional/remote_tensor_tests/ocl_remote_tensor_tests.cpp +++ b/src/plugins/intel_gpu/tests/functional/remote_tensor_tests/ocl_remote_tensor_tests.cpp @@ -2,8 +2,6 @@ // SPDX-License-Identifier: Apache-2.0 // -#ifdef OV_GPU_WITH_OCL_RT - #include "openvino/core/preprocess/pre_post_process.hpp" #include "openvino/op/add.hpp" #include "openvino/op/constant.hpp" @@ -2947,4 +2945,3 @@ INSTANTIATE_TEST_SUITE_P(smoke_RemoteTensorDataType, OVRemoteTensorDataType_Test ov::element::Type_t::u16, ov::element::Type_t::u32)), OVRemoteTensorDataType_Test::getTestCaseName); -#endif // OV_GPU_WITH_OCL_RT diff --git a/src/plugins/intel_gpu/tests/unit/module_tests/ze/ze_device_test.cpp b/src/plugins/intel_gpu/tests/unit/module_tests/ze/ze_device_test.cpp index c693022c4dee08..6d79c29485c454 100644 --- a/src/plugins/intel_gpu/tests/unit/module_tests/ze/ze_device_test.cpp +++ b/src/plugins/intel_gpu/tests/unit/module_tests/ze/ze_device_test.cpp @@ -155,9 +155,9 @@ TEST(ze_devices_test, on_demand_initialization) { auto ze_device = std::dynamic_pointer_cast(device.second); ASSERT_NE(ze_device, nullptr); ASSERT_FALSE(ze_device->is_initialized()); - ASSERT_EQ(ze_device->get_context(), nullptr); - ASSERT_TRUE(ze_device->get_device() != nullptr); - ASSERT_TRUE(ze_device->get_driver() != nullptr); + ASSERT_EQ(ze_device->get_context().is_empty(), true); + ASSERT_TRUE(ze_device->get_device().handle() != nullptr); + ASSERT_TRUE(ze_device->get_driver().handle() != nullptr); ASSERT_FALSE(ze_device->get_info().execution_units_count == 0); ASSERT_FALSE(ze_device->get_info().vendor_id == 0); } @@ -170,8 +170,8 @@ TEST(ze_devices_test, on_demand_initialization) { auto ze_device = std::dynamic_pointer_cast(device.second); ASSERT_NE(ze_device, nullptr); ASSERT_TRUE(ze_device->is_initialized()); - ASSERT_TRUE(ze_device->get_device() != nullptr); - ASSERT_TRUE(ze_device->get_context() != nullptr); + ASSERT_TRUE(ze_device->get_device().handle() != nullptr); + ASSERT_TRUE(ze_device->get_context().handle() != nullptr); } } @@ -187,9 +187,9 @@ TEST(ze_devices_test, user_context_initialization_not_implemented) { auto initialized_device = std::dynamic_pointer_cast(devices.begin()->second); ASSERT_NE(initialized_device, nullptr); auto user_context = initialized_device->get_context(); - ASSERT_TRUE(user_context != nullptr); + ASSERT_TRUE(user_context.handle() != nullptr); - ASSERT_ANY_THROW(device_detector.get_available_devices(user_context, nullptr, 0, std::numeric_limits::max() /* ignore sub-devices */)); + ASSERT_ANY_THROW(device_detector.get_available_devices(user_context.handle(), nullptr, 0, std::numeric_limits::max() /* ignore sub-devices */)); } TEST(ze_devices_test, user_device_initialization_not_implemented) { @@ -204,9 +204,9 @@ TEST(ze_devices_test, user_device_initialization_not_implemented) { auto initialized_device = std::dynamic_pointer_cast(devices.begin()->second); ASSERT_NE(initialized_device, nullptr); auto user_device = initialized_device->get_device(); - ASSERT_TRUE(user_device != nullptr); + ASSERT_TRUE(user_device.handle() != nullptr); - ASSERT_ANY_THROW(device_detector.get_available_devices(nullptr, user_device, 0, std::numeric_limits::max() /* ignore sub-devices */)); + ASSERT_ANY_THROW(device_detector.get_available_devices(nullptr, user_device.handle(), 0, std::numeric_limits::max() /* ignore sub-devices */)); } #endif // OV_GPU_WITH_ZE_RT diff --git a/src/plugins/intel_gpu/tests/unit/module_tests/ze/ze_resource_test.cpp b/src/plugins/intel_gpu/tests/unit/module_tests/ze/ze_resource_test.cpp new file mode 100644 index 00000000000000..aea4ffdad94303 --- /dev/null +++ b/src/plugins/intel_gpu/tests/unit/module_tests/ze/ze_resource_test.cpp @@ -0,0 +1,214 @@ +// Copyright (C) 2018-2026 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +#include "ze/ze_resource.hpp" + +#include + +#include +#include +#include + +using namespace cldnn::ze; + +template +struct ze_resource_tag { + static constexpr ze_resource_type value = resource_type; +}; + +template +typename ocl_resource_info::handle_t make_ocl_handle(std::uintptr_t value) { + using handle_t = typename ocl_resource_info::handle_t; + return reinterpret_cast(value); +} + +template +typename ze_resource_info::handle_t make_ze_handle(std::uintptr_t value) { + using handle_t = typename ze_resource_info::handle_t; + if constexpr (resource_type == ze_resource_type::usm_memory) { + return ov_ze_usm_handle{reinterpret_cast(value), + reinterpret_cast(value)}; + } else { + return reinterpret_cast(value); + } +} + +template +void assert_ze_handle_eq(const typename ze_resource_info::handle_t& expected, + const typename ze_resource_info::handle_t& actual) { + if constexpr (resource_type == ze_resource_type::usm_memory) { + ASSERT_EQ(expected.context, actual.context); + ASSERT_EQ(expected.ptr, actual.ptr); + } else { + ASSERT_EQ(expected, actual); + } +} + +template +struct ocl_attachment_info; + +template <> +struct ocl_attachment_info { + static constexpr ocl_resource_type ocl_type = ocl_resource_type::platform; +}; + +template <> +struct ocl_attachment_info { + static constexpr ocl_resource_type ocl_type = ocl_resource_type::device; +}; + +template <> +struct ocl_attachment_info { + static constexpr ocl_resource_type ocl_type = ocl_resource_type::context; +}; + +template <> +struct ocl_attachment_info { + static constexpr ocl_resource_type ocl_type = ocl_resource_type::command_queue; +}; + +template <> +struct ocl_attachment_info { + static constexpr ocl_resource_type ocl_type = ocl_resource_type::mem_object; +}; + +template <> +struct ocl_attachment_info { + static constexpr ocl_resource_type ocl_type = ocl_resource_type::mem_object; +}; + +using all_ze_resource_types = ::testing::Types, + ze_resource_tag, + ze_resource_tag, + ze_resource_tag, + ze_resource_tag, + ze_resource_tag, + ze_resource_tag, + ze_resource_tag, + ze_resource_tag, + ze_resource_tag, + ze_resource_tag, + ze_resource_tag, + ze_resource_tag>; + +using ze_resource_types_with_ocl = ::testing::Types, + ze_resource_tag, + ze_resource_tag, + ze_resource_tag, + ze_resource_tag, + ze_resource_tag>; + +template +class ze_resource_all_types_test : public ::testing::Test {}; + +template +class ze_resource_with_ocl_types_test : public ::testing::Test {}; + +TYPED_TEST_SUITE(ze_resource_all_types_test, all_ze_resource_types); +TYPED_TEST_SUITE(ze_resource_with_ocl_types_test, ze_resource_types_with_ocl); + +TYPED_TEST(ze_resource_all_types_test, can_be_empty) { + constexpr auto resource_type = TypeParam::value; + using resource_t = ze_resource; + + resource_t resource; + + ASSERT_TRUE(resource.is_empty()); + ASSERT_EQ(resource.get_holder(), nullptr); + ASSERT_ANY_THROW(resource.handle()); +} + +TYPED_TEST(ze_resource_all_types_test, can_hold_ze_handle) { + constexpr auto resource_type = TypeParam::value; + using resource_t = ze_resource; + + const auto ze_handle = make_ze_handle(0x100); + resource_t resource(ze_handle, true); + + ASSERT_FALSE(resource.is_empty()); + ASSERT_NE(resource.get_holder(), nullptr); + assert_ze_handle_eq(ze_handle, resource.handle()); +} + +TYPED_TEST(ze_resource_all_types_test, can_share_ze_handle) { + constexpr auto resource_type = TypeParam::value; + using resource_t = ze_resource; + + const auto ze_handle = make_ze_handle(0x200); + std::weak_ptr weak_holder; + resource_t resource; + { + auto holder = std::make_shared(ze_handle, true); + resource = resource_t(holder); + weak_holder = holder; + } + ASSERT_FALSE(weak_holder.expired()); + ASSERT_FALSE(resource.is_empty()); + ASSERT_EQ(resource.get_holder(), weak_holder.lock()); + + auto copied_resource = resource; + ASSERT_EQ(copied_resource.get_holder(), weak_holder.lock()); + + resource.drop(); + ASSERT_TRUE(resource.is_empty()); + ASSERT_FALSE(weak_holder.expired()); + ASSERT_FALSE(copied_resource.is_empty()); + ASSERT_EQ(copied_resource.get_holder(), weak_holder.lock()); + + copied_resource.drop(); + ASSERT_TRUE(copied_resource.is_empty()); + ASSERT_TRUE(weak_holder.expired()); +} + +TYPED_TEST(ze_resource_with_ocl_types_test, can_not_attach_ocl_handle_to_empty_resource) { + constexpr auto resource_type = TypeParam::value; + constexpr auto ocl_type = ocl_attachment_info::ocl_type; + using resource_t = ze_resource; + + resource_t resource; + + ASSERT_FALSE(resource.template has_ocl_handle()); + ASSERT_ANY_THROW(resource.template ocl_handle()); + ASSERT_ANY_THROW(resource.template attach_ocl_handle(make_ocl_handle(0x300), true)); +} + +TYPED_TEST(ze_resource_with_ocl_types_test, can_attach_and_access_ocl_handle) { + constexpr auto resource_type = TypeParam::value; + constexpr auto ocl_type = ocl_attachment_info::ocl_type; + using resource_t = ze_resource; + + const auto ze_handle = make_ze_handle(0x400); + const auto ocl_handle = make_ocl_handle(0x500); + + resource_t resource(ze_handle, true); + + ASSERT_FALSE(resource.template has_ocl_handle()); + resource.template attach_ocl_handle(ocl_handle, true); + + ASSERT_TRUE(resource.template has_ocl_handle()); + ASSERT_EQ(resource.template ocl_handle(), ocl_handle); + ASSERT_ANY_THROW(resource.template attach_ocl_handle(ocl_handle, true)); +} + +TYPED_TEST(ze_resource_with_ocl_types_test, can_share_ocl_handle) { + constexpr auto resource_type = TypeParam::value; + constexpr auto ocl_type = ocl_attachment_info::ocl_type; + using resource_t = ze_resource; + using ocl_owner_t = ocl_owner; + + const auto ze_handle = make_ze_handle(0x600); + const auto ocl_handle = make_ocl_handle(0x700); + + resource_t resource(ze_handle, true); + ASSERT_FALSE(resource.template has_ocl_handle()); + auto copied_resource = resource; + { + ocl_owner_t owner(ocl_handle, true); + copied_resource.template attach_ocl_handle(std::move(owner)); + } + ASSERT_TRUE(resource.template has_ocl_handle()); + ASSERT_EQ(resource.template ocl_handle(), ocl_handle); + ASSERT_TRUE(copied_resource.template has_ocl_handle()); + ASSERT_EQ(copied_resource.template ocl_handle(), ocl_handle); +} diff --git a/src/plugins/intel_gpu/tests/unit/test_utils/test_utils.cpp b/src/plugins/intel_gpu/tests/unit/test_utils/test_utils.cpp index 8d6058aab047f4..504b693246aeb2 100644 --- a/src/plugins/intel_gpu/tests/unit/test_utils/test_utils.cpp +++ b/src/plugins/intel_gpu/tests/unit/test_utils/test_utils.cpp @@ -307,7 +307,7 @@ cldnn::ExecutionConfig get_test_default_config(const cldnn::engine& engine, std::shared_ptr create_test_engine() { auto ret = cldnn::engine::create( - cldnn::device_query::get_default_engine_type(), cldnn::device_query::get_default_runtime_type()); + cldnn::get_default_engine_type(), cldnn::get_default_runtime_type()); #ifdef ENABLE_ONEDNN_FOR_GPU if (ret->get_device_info().supports_immad) ret->create_onednn_engine({});