[SYCL][L0 v2] Use counter-based events

sycl/test-e2e/Adapters/level_zero/event-leak.cpp

@@ -1,5 +1,8 @@
 // REQUIRES: level_zero, level_zero_dev_kit
 //
+// UNSUPPORTED: windows && level_zero_v2_adapter
+// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/20852
+//
 // RUN: %{build} %level_zero_options -o %t.out
 // RUN: %{l0_leak_check} %{run} %t.out wait 2>&1 | FileCheck %s
 // RUN: %{l0_leak_check} %{run} %t.out nowait 2>&1 | FileCheck %s

sycl/test-e2e/DeprecatedFeatures/DiscardEvents/discard_events_l0_leak.cpp

@@ -1,5 +1,8 @@
 // REQUIRES: level_zero
 //
+// UNSUPPORTED: windows && level_zero_v2_adapter
+// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/20852
+//
 // RUN: %{build} -o %t.out
 //
 // RUN: env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0 SYCL_PI_LEVEL_ZERO_BATCH_SIZE=4 ONEAPI_DEVICE_SELECTOR='level_zero:*' %{l0_leak_check} %{run} %t.out wait  2>&1 | FileCheck %s

sycl/test-e2e/ProfilingTag/profile_tag_leak.cpp

@@ -1,5 +1,8 @@
 // REQUIRES: level_zero
-
+//
+// UNSUPPORTED: windows && level_zero_v2_adapter
+// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/20852
+//
 // RUN: %{build} -o %t.out
 // RUN: %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK
 

sycl/test-e2e/Regression/reduction_resource_leak_dw.cpp

@@ -1,5 +1,8 @@
 // REQUIRES: level_zero
 //
+// UNSUPPORTED: windows && level_zero_v2_adapter
+// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/20852
+//
 // RUN: %{build} -o %t.out
 // RUN: %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s
 //

sycl/test-e2e/USM/usm_leak_check.cpp

@@ -1,4 +1,7 @@
 // REQUIRES: level_zero
+//
+// UNSUPPORTED: windows && level_zero_v2_adapter
+// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/20852
 
 // RUN: %{build} -Wno-error=deprecated-declarations -o %t.out
 
@@ -36,10 +39,8 @@ template <typename T, size_t N> void sycl_buffer(queue &Q) {
     accessor accB{bufferB, cgh, read_only};
     accessor accC{bufferC, cgh, write_only};
 
-    cgh.parallel_for<class K<T>>(numElems,
-		[=](id<1> wiID) {
-      accC[wiID] = accA[wiID] + accB[wiID];
-    });
+    cgh.parallel_for<class K<T>>(
+        numElems, [=](id<1> wiID) { accC[wiID] = accA[wiID] + accB[wiID]; });
   });
 }
 

unified-runtime/source/adapters/level_zero/v2/context.cpp

@@ -80,23 +80,26 @@ ur_context_handle_t_::ur_context_handle_t_(ze_context_handle_t hContext,
                         phDevices[0]->Platform->ZeMutableCmdListExt.Supported}),
       eventPoolCacheImmediate(
           this, phDevices[0]->Platform->getNumDevices(),
-          [context = this](DeviceId /* deviceId*/, v2::event_flags_t flags)
-              -> std::unique_ptr<v2::event_provider> {
+          [context = this, platform = phDevices[0]->Platform](
+              DeviceId deviceId,
+              v2::event_flags_t flags) -> std::unique_ptr<v2::event_provider> {
+            auto device = platform->getDeviceById(deviceId);
+
+            // TODO: just use per-context id?
+            return v2::createProvider(platform, context, v2::QUEUE_IMMEDIATE,
+                                      device, flags);
+          }),
+      eventPoolCacheRegular(
+          this, phDevices[0]->Platform->getNumDevices(),
+          [context = this, platform = phDevices[0]->Platform](
+              DeviceId deviceId,
+              v2::event_flags_t flags) -> std::unique_ptr<v2::event_provider> {
+            auto device = platform->getDeviceById(deviceId);
+
             // TODO: just use per-context id?
-            return std::make_unique<v2::provider_normal>(
-                context, v2::QUEUE_IMMEDIATE, flags);
+            return v2::createProvider(platform, context, v2::QUEUE_REGULAR,
+                                      device, flags);
           }),
-      eventPoolCacheRegular(this, phDevices[0]->Platform->getNumDevices(),
-                            [context = this, platform = phDevices[0]->Platform](
-                                DeviceId deviceId, v2::event_flags_t flags)
-                                -> std::unique_ptr<v2::event_provider> {
-                              std::ignore = deviceId;
-                              std::ignore = platform;
-
-                              // TODO: just use per-context id?
-                              return std::make_unique<v2::provider_normal>(
-                                  context, v2::QUEUE_REGULAR, flags);
-                            }),
       nativeEventsPool(this, std::make_unique<v2::provider_normal>(
                                  this, v2::QUEUE_IMMEDIATE,
                                  v2::EVENT_FLAGS_PROFILING_ENABLED)),

unified-runtime/source/adapters/level_zero/v2/event_provider.hpp

@@ -28,6 +28,11 @@ enum event_flag_t {
 };
 static constexpr size_t EVENT_FLAGS_USED_BITS = 2;
 
+enum queue_type {
+  QUEUE_REGULAR,
+  QUEUE_IMMEDIATE,
+};
+
 class event_provider;
 
 namespace raii {

unified-runtime/source/adapters/level_zero/v2/event_provider_counter.cpp

@@ -13,6 +13,7 @@
 #include "context.hpp"
 #include "event_provider.hpp"
 #include "event_provider_counter.hpp"
+#include "event_provider_normal.hpp"
 #include "loader/ze_loader.h"
 
 #include "../device.hpp"
@@ -22,10 +23,17 @@ namespace v2 {
 
 provider_counter::provider_counter(ur_platform_handle_t platform,
                                    ur_context_handle_t context,
-                                   ur_device_handle_t device) {
+                                   queue_type queueType,
+                                   ur_device_handle_t device,
+                                   event_flags_t flags)
+    : queueType(queueType), flags(flags) {
+  assert(flags & EVENT_FLAGS_COUNTER);
+
+  // Try to get the counter-based event extension function
   ZE2UR_CALL_THROWS(zeDriverGetExtensionFunctionAddress,
-                    (platform->ZeDriver, "zexCounterBasedEventCreate",
+                    (platform->ZeDriver, "zexCounterBasedEventCreate2",
                      (void **)&this->eventCreateFunc));
+
   ZE2UR_CALL_THROWS(zelLoaderTranslateHandle,
                     (ZEL_HANDLE_CONTEXT, context->getZeHandle(),
                      (void **)&translatedContext));
@@ -34,17 +42,41 @@ provider_counter::provider_counter(ur_platform_handle_t platform,
       (ZEL_HANDLE_DEVICE, device->ZeDevice, (void **)&translatedDevice));
 }
 
+static zex_counter_based_event_exp_flags_t createZeFlags(queue_type queueType,
+                                                         event_flags_t flags) {
+  zex_counter_based_event_exp_flags_t zeFlags =
+      ZEX_COUNTER_BASED_EVENT_FLAG_HOST_VISIBLE;
+  if (flags & EVENT_FLAGS_PROFILING_ENABLED) {
+    zeFlags |= ZEX_COUNTER_BASED_EVENT_FLAG_KERNEL_TIMESTAMP;
+  }
+
+  if (queueType == QUEUE_IMMEDIATE) {
+    zeFlags |= ZEX_COUNTER_BASED_EVENT_FLAG_IMMEDIATE;
+  } else {
+    zeFlags |= ZEX_COUNTER_BASED_EVENT_FLAG_NON_IMMEDIATE;
+  }
+
+  return zeFlags;
+}
+
 raii::cache_borrowed_event provider_counter::allocate() {
   if (freelist.empty()) {
-    ZeStruct<ze_event_desc_t> desc;
-    desc.index = 0;
-    desc.signal = ZE_EVENT_SCOPE_FLAG_HOST;
-    desc.wait = 0;
+    zex_counter_based_event_desc_t desc = {};
+    desc.stype = ZEX_STRUCTURE_COUNTER_BASED_EVENT_DESC;
+    desc.flags = createZeFlags(queueType, flags);
+    desc.signalScope = ZE_EVENT_SCOPE_FLAG_HOST;
+
+    uint32_t equivalentFlags = ZE_EVENT_POOL_FLAG_HOST_VISIBLE;
+    if (flags & EVENT_FLAGS_PROFILING_ENABLED) {
+      equivalentFlags |= ZE_EVENT_POOL_FLAG_KERNEL_TIMESTAMP;
+    }
+    UR_LOG(DEBUG, "ze_event_pool_desc_t flags set to: {}", equivalentFlags);
+
     ze_event_handle_t handle;
 
     // TODO: allocate host and device buffers to use here
-    ZE2UR_CALL_THROWS(eventCreateFunc, (translatedContext, translatedDevice,
-                                        nullptr, nullptr, 0, &desc, &handle));
+    ZE2UR_CALL_THROWS(eventCreateFunc,
+                      (translatedContext, translatedDevice, &desc, &handle));
 
     freelist.emplace_back(handle);
   }
@@ -57,8 +89,29 @@ raii::cache_borrowed_event provider_counter::allocate() {
       [this](ze_event_handle_t handle) { freelist.push_back(handle); });
 }
 
-event_flags_t provider_counter::eventFlags() const {
-  return EVENT_FLAGS_COUNTER;
+event_flags_t provider_counter::eventFlags() const { return flags; }
+
+std::unique_ptr<event_provider> createProvider(ur_platform_handle_t platform,
+                                               ur_context_handle_t context,
+                                               queue_type queueType,
+                                               ur_device_handle_t device,
+                                               event_flags_t flags) {
+  // Only try counter-based events if the flag is set
+  if (flags & EVENT_FLAGS_COUNTER) {
+    // Try to create a counter-based event provider first
+    try {
+      return std::make_unique<provider_counter>(platform, context, queueType,
+                                                device, flags);
+    } catch (...) {
+      // If the new counter-based API (zexCounterBasedEventCreate2) is not
+      // available, fall back to normal provider which support counter-based
+      // events using the old API
+      return std::make_unique<provider_normal>(context, queueType, flags);
+    }
+  }
+
+  // Counter-based events not requested, use normal events
+  return std::make_unique<provider_normal>(context, queueType, flags);
 }
 
 } // namespace v2

unified-runtime/source/adapters/level_zero/v2/event_provider_counter.hpp

@@ -25,22 +25,27 @@
 
 #include "../device.hpp"
 
+#include <level_zero/driver_experimental/zex_event.h>
+#include <level_zero/ze_intel_gpu.h>
+
 namespace v2 {
 
 typedef ze_result_t (*zexCounterBasedEventCreate)(
     ze_context_handle_t hContext, ze_device_handle_t hDevice,
-    uint64_t *deviceAddress, uint64_t *hostAddress, uint64_t completionValue,
-    const ze_event_desc_t *desc, ze_event_handle_t *phEvent);
+    const zex_counter_based_event_desc_t *desc, ze_event_handle_t *phEvent);
 
 class provider_counter : public event_provider {
 public:
   provider_counter(ur_platform_handle_t platform, ur_context_handle_t,
-                   ur_device_handle_t);
+                   queue_type, ur_device_handle_t, event_flags_t);
 
   raii::cache_borrowed_event allocate() override;
   event_flags_t eventFlags() const override;
 
 private:
+  queue_type queueType;
+  event_flags_t flags;
+
   ze_context_handle_t translatedContext;
   ze_device_handle_t translatedDevice;
 
@@ -49,4 +54,12 @@ class provider_counter : public event_provider {
   std::vector<raii::ze_event_handle_t> freelist;
 };
 
+// Factory function that creates a counter-based provider with fallback to
+// normal provider
+std::unique_ptr<event_provider> createProvider(ur_platform_handle_t platform,
+                                               ur_context_handle_t context,
+                                               queue_type queueType,
+                                               ur_device_handle_t device,
+                                               event_flags_t flags);
+
 } // namespace v2

unified-runtime/source/adapters/level_zero/v2/event_provider_normal.hpp

@@ -21,17 +21,13 @@
 
 #include "common.hpp"
 #include "event.hpp"
+#include "event_provider.hpp"
 
 #include "../device.hpp"
 #include "../ur_interface_loader.hpp"
 
 namespace v2 {
 
-enum queue_type {
-  QUEUE_REGULAR,
-  QUEUE_IMMEDIATE,
-};
-
 class provider_pool {
 public:
   provider_pool(ur_context_handle_t, queue_type, event_flags_t flags);

unified-runtime/test/adapters/level_zero/v2/event_pool_test.cpp

@@ -137,8 +137,8 @@ struct EventPoolTest : public uur::urQueueTestWithParam<ProviderParams> {
           // the provider
           switch (params.provider) {
           case TEST_PROVIDER_COUNTER:
-            return std::make_unique<provider_counter>(platform, context,
-                                                      device);
+            return std::make_unique<provider_counter>(
+                platform, context, params.queue, device, params.flags);
           case TEST_PROVIDER_NORMAL:
             return std::make_unique<provider_normal>(context, params.queue,
                                                      flags);