fixes

cmake/features.cmake

@@ -44,7 +44,7 @@ endif()
 set(OV_GPU_DEFAULT_RT "L0")
 if (ENABLE_INTEL_GPU)
     ov_option_enum (GPU_RT_TYPE "Type of GPU runtime. Supported value: OCL and L0" ${OV_GPU_DEFAULT_RT} ALLOWED_VALUES L0 OCL)
-    if (GPU_RT_TYPE STREQUAL L0)
+    if (GPU_RT_TYPE STREQUAL "L0")
         # There's no interop with native L0 in onednn API. Temporary disable onednn when L0 runtime is selected
         set(ENABLE_ONEDNN_FOR_GPU_DEFAULT OFF)
     endif()

src/plugins/intel_gpu/include/intel_gpu/runtime/device_info.hpp

@@ -75,6 +75,8 @@ struct device_info {
     bool supports_imad;                         ///< Does engine support int8 mad.
     bool supports_immad;                        ///< Does engine support int8 multi mad.
 
+    bool supports_mutable_command_list;         ///< Does the target runtime/device support mutable command list feature
+
     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
 

src/plugins/intel_gpu/src/runtime/ocl/ocl_device.cpp

@@ -322,6 +322,8 @@ device_info init_device_info(const cl::Device& device, const cl::Context& contex
         info.num_ccs = std::max<uint32_t>(num_queues, info.num_ccs);
     }
 
+    info.supports_mutable_command_list = false;
+
     // Not supported
     info.timer_resolution = 0;
     info.kernel_timestamp_valid_bits = 0;

src/plugins/intel_gpu/src/runtime/ze/ze_device.cpp

@@ -28,14 +28,28 @@ namespace ze {
 
 namespace {
 
+bool supports_extension(const std::vector<ze_driver_extension_properties_t>& extensions, const std::string& ext_name, uint32_t ext_ver) {
+    return std::find_if(extensions.begin(), extensions.end(), [&ext_name, &ext_ver](const ze_driver_extension_properties_t& ep) {
+        return std::string(ep.name) == ext_name && ep.version == ext_ver;
+    }) != extensions.end();
+}
+
 device_info init_device_info(ze_driver_handle_t driver, ze_device_handle_t device) {
     device_info info;
 
+    uint32_t num_ext = 0;
+    ZE_CHECK(zeDriverGetExtensionProperties(driver, &num_ext, nullptr));
+
+    std::vector<ze_driver_extension_properties_t> extensions(num_ext);
+    ZE_CHECK(zeDriverGetExtensionProperties(driver, &num_ext, &extensions[0]));
+
     ze_driver_properties_t driver_properties;
     ZE_CHECK(zeDriverGetProperties(driver, &driver_properties));
 
+    bool supports_luid = supports_extension(extensions, ZE_DEVICE_LUID_EXT_NAME, ZE_DEVICE_LUID_EXT_VERSION_1_0);
+
     ze_device_ip_version_ext_t ip_version_properties = {ZE_STRUCTURE_TYPE_DEVICE_IP_VERSION_EXT, nullptr, 0};
-    ze_device_properties_t device_properties{ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES_1_2, &ip_version_properties};
+    ze_device_properties_t device_properties{ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES_1_2, supports_luid ? &ip_version_properties : nullptr};
     ZE_CHECK(zeDeviceGetProperties(device, &device_properties));
 
     ze_device_compute_properties_t device_compute_properties{ZE_STRUCTURE_TYPE_DEVICE_COMPUTE_PROPERTIES};
@@ -66,12 +80,6 @@ device_info init_device_info(ze_driver_handle_t driver, ze_device_handle_t devic
     }
     ZE_CHECK(zeDeviceGetMemoryProperties(device, &memory_properties_count, &device_memory_properties[0]));
 
-    uint32_t extension_properties_count = 0;
-    ZE_CHECK(zeDriverGetExtensionProperties(driver, &extension_properties_count, nullptr));
-
-    std::vector<ze_driver_extension_properties_t> driver_extension_properties(extension_properties_count);
-    ZE_CHECK(zeDriverGetExtensionProperties(driver, &extension_properties_count, &driver_extension_properties[0]));
-
     ze_device_memory_access_properties_t device_memory_access_properties{ZE_STRUCTURE_TYPE_DEVICE_MEMORY_ACCESS_PROPERTIES};
     ZE_CHECK(zeDeviceGetMemoryAccessProperties(device, &device_memory_access_properties));
 
@@ -149,13 +157,14 @@ device_info init_device_info(ze_driver_handle_t driver, ze_device_handle_t devic
     static_assert(ZE_MAX_DEVICE_LUID_SIZE_EXT == ov::device::LUID::MAX_LUID_SIZE, "");
     std::copy_n(&device_properties.uuid.id[0], ZE_MAX_DEVICE_UUID_SIZE, info.uuid.uuid.begin());
 
-    {
+    if (supports_luid) {
         ze_device_luid_ext_properties_t luid_props{ZE_STRUCTURE_TYPE_DEVICE_LUID_EXT_PROPERTIES, nullptr};
         ze_device_properties_t device_properties{ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES, &luid_props};
         if (zeDeviceGetProperties(device, &device_properties) == ZE_RESULT_SUCCESS)
             std::copy_n(&luid_props.luid.id[0], ZE_MAX_DEVICE_LUID_SIZE_EXT, info.luid.luid.begin());
     }
 
+    info.supports_mutable_command_list = supports_extension(extensions, ZE_MUTABLE_COMMAND_LIST_EXP_NAME, ZE_MUTABLE_COMMAND_LIST_EXP_VERSION_1_0);
     return info;
 }
 

src/plugins/intel_gpu/src/runtime/ze/ze_device_detector.cpp

@@ -15,13 +15,11 @@ namespace cldnn {
 namespace ze {
 
 static std::vector<ze_device_handle_t> get_sub_devices(ze_device_handle_t root_device) {
-
     uint32_t n_subdevices = 0;
     ZE_CHECK(zeDeviceGetSubDevices(root_device, &n_subdevices, nullptr));
     if (n_subdevices == 0)
         return {};
 
-
     std::vector<ze_device_handle_t> subdevices(n_subdevices);
 
     ZE_CHECK(zeDeviceGetSubDevices(root_device, &n_subdevices, &subdevices[0]));

src/plugins/intel_gpu/src/runtime/ze/ze_event.cpp

@@ -19,8 +19,8 @@ std::chrono::nanoseconds timestamp_to_duration(const device_info& device_info, c
     const double timestamp_freq = NS_IN_SEC / device_info.timer_resolution;
     const uint64_t timestamp_max_value = ~(-1L << device_info.kernel_timestamp_valid_bits);
 
-    auto d = (timestamp.kernelEnd >= timestamp.kernelStart) ? ( timestamp.kernelEnd - timestamp.kernelStart ) * timestamp_freq
-                                                            : (( timestamp_max_value - timestamp.kernelStart) + timestamp.kernelEnd + 1 ) * timestamp_freq;
+    auto d = (timestamp.kernelEnd >= timestamp.kernelStart) ? (timestamp.kernelEnd - timestamp.kernelStart) * timestamp_freq
+                                                            : ((timestamp_max_value - timestamp.kernelStart) + timestamp.kernelEnd + 1) * timestamp_freq;
 
     return std::chrono::nanoseconds(static_cast<uint64_t>(d));
 }

src/plugins/intel_gpu/src/runtime/ze/ze_stream.cpp

@@ -166,7 +166,6 @@ ze_stream::ze_stream(const ze_engine &engine, const ExecutionConfig& config)
     : stream(config.get_property(ov::intel_gpu::queue_type), stream::get_expected_sync_method(config))
     , _engine(engine)
     , m_pool(engine, config.get_property(ov::enable_profiling)) {
-
     ze_command_queue_desc_t command_queue_desc = {};
     command_queue_desc.stype = ZE_STRUCTURE_TYPE_COMMAND_QUEUE_DESC;
     command_queue_desc.pNext = nullptr;
@@ -221,11 +220,11 @@ event::ptr ze_stream::enqueue_kernel(kernel& kernel,
     ze_group_count_t args = { global.groupCountX / local.groupCountX, global.groupCountY / local.groupCountY, global.groupCountZ / local.groupCountZ };
     ZE_CHECK(zeKernelSetGroupSize(kern, local.groupCountX, local.groupCountY, local.groupCountZ));
     ZE_CHECK(zeCommandListAppendLaunchKernel(m_command_list,
-                                    kern,
-                                    &args,
-                                    set_output_event ? std::dynamic_pointer_cast<ze_base_event>(ev)->get() : nullptr,
-                                    dep_events_ptr == nullptr ? 0 : dep_events_ptr->size(),
-                                    dep_events_ptr == nullptr ? 0 : &dep_events_ptr->front()));
+                                             kern,
+                                             &args,
+                                             set_output_event ? std::dynamic_pointer_cast<ze_base_event>(ev)->get() : nullptr,
+                                             dep_events_ptr == nullptr ? 0 : static_cast<uint32_t>(dep_events_ptr->size()),
+                                             dep_events_ptr == nullptr ? 0 : &dep_events_ptr->front()));
 
     return ev;
 }
@@ -253,7 +252,10 @@ event::ptr ze_stream::enqueue_marker(std::vector<ze_event::ptr> const& deps, boo
             return create_user_event(true);
 
         auto ev = create_base_event();
-        ZE_CHECK(zeCommandListAppendBarrier(m_command_list, std::dynamic_pointer_cast<ze_base_event>(ev)->get(), dep_events.size(), &dep_events.front()));
+        ZE_CHECK(zeCommandListAppendBarrier(m_command_list,
+                                            std::dynamic_pointer_cast<ze_base_event>(ev)->get(),
+                                            static_cast<uint32_t>(dep_events.size()),
+                                            &dep_events.front()));
         return ev;
     } else if (m_sync_method == SyncMethods::barriers) {
         sync_events(deps, is_output);