COPY OF #22641 [SYCL][UR][OpenCL] Fix profiling tag ordering on OpenCL GPU

sycl/include/sycl/ext/oneapi/experimental/profiling_tag.hpp

@@ -21,22 +21,24 @@ namespace ext::oneapi::experimental {
 inline event submit_profiling_tag(queue &Queue,
                                   const sycl::detail::code_location &CodeLoc =
                                       sycl::detail::code_location::current()) {
-  if (Queue.get_device().has(aspect::ext_oneapi_queue_profiling_tag)) {
-    return Queue.submit(
-        [=](handler &CGH) {
-          sycl::detail::HandlerAccess::internalProfilingTagImpl(CGH);
-        },
-        CodeLoc);
-  }
-
-  // If it is not supported natively on the device, we use another path if
-  // profiling is enabled.
-  if (!Queue.has_property<sycl::property::queue::enable_profiling>())
+  // The profiling tag can be serviced natively when the device advertises the
+  // ext_oneapi_queue_profiling_tag aspect. Otherwise, we can still service it
+  // as long as the queue has profiling enabled. In both cases we submit an
+  // internal profiling-tag command group: the runtime records the timestamp
+  // using a native device command where possible and otherwise falls back to a
+  // barrier (see CGType::ProfilingTag handling in the scheduler).
+  if (!Queue.get_device().has(aspect::ext_oneapi_queue_profiling_tag) &&
+      !Queue.has_property<sycl::property::queue::enable_profiling>())
     throw sycl::exception(
         make_error_code(errc::invalid),
         "Device must either have aspect::ext_oneapi_queue_profiling_tag or the "
         "queue must have profiling enabled.");
-  return Queue.ext_oneapi_submit_barrier();
+
+  return Queue.submit(
+      [=](handler &CGH) {
+        sycl::detail::HandlerAccess::internalProfilingTagImpl(CGH);
+      },
+      CodeLoc);
 }
 
 } // namespace ext::oneapi::experimental

sycl/source/detail/event_impl.cpp

@@ -427,8 +427,9 @@ uint64_t
 event_impl::get_profiling_info<info::event_profiling::command_submit>() {
   checkProfilingPreconditions();
   if (isProfilingTagEvent()) {
-    // Tag events report command_submit through the adapter.
-    return get_event_profiling_info<info::event_profiling::command_submit>(
+    // The empty tag command uses its completion timestamp for all three
+    // queries.
+    return get_event_profiling_info<info::event_profiling::command_end>(
         this->getHandle(), this->getAdapter());
   }
 
@@ -464,6 +465,9 @@ event_impl::get_profiling_info<info::event_profiling::command_start>() {
   if (!MIsHostEvent) {
     auto Handle = getHandle();
     if (Handle) {
+      if (isProfilingTagEvent())
+        return get_event_profiling_info<info::event_profiling::command_end>(
+            Handle, this->getAdapter());
       return get_event_profiling_info<info::event_profiling::command_start>(
           Handle, this->getAdapter());
     }

sycl/source/detail/scheduler/commands.cpp

@@ -22,6 +22,7 @@
 #include <detail/scheduler/commands.hpp>
 #include <detail/scheduler/scheduler.hpp>
 #include <detail/stream_impl.hpp>
+#include <detail/ur_utils.hpp>
 #include <detail/xpti_registry.hpp>
 #include <sycl/access/access.hpp>
 #include <sycl/backend_types.hpp>
@@ -3735,17 +3736,14 @@ ur_result_t ExecCGCommand::enqueueImpQueue() {
     assert(MQueue && "Profiling tag requires a valid queue");
     adapter_impl &Adapter = MQueue->getAdapter();
 
-    bool IsInOrderQueue = MQueue->isInOrder();
-    ur_event_handle_t *TimestampDeps = nullptr;
-    size_t NumTimestampDeps = 0;
-
-    // TO DO - once the following WA removed: to change call to call_nocheck and
-    // return operation result to Command::enqueue (see other CG types). Set
-    // UREvent to EventImpl only for successful case.
+    const bool IsInOrderQueue = MQueue->isInOrder();
 
     // If the queue is not in-order, the implementation will need to first
     // insert a marker event that the timestamp waits for.
     ur_event_handle_t PreTimestampMarkerEvent{};
+    std::optional<OwnedUrEvent> OwnedPreTimestampMarkerEvent;
+    ur_event_handle_t *TimestampDeps = nullptr;
+    size_t NumTimestampDeps = 0;
     if (!IsInOrderQueue) {
       // FIXME: urEnqueueEventsWait on the L0 adapter requires a double-release.
       //        Use that instead once it has been fixed.
@@ -3754,29 +3752,57 @@ ur_result_t ExecCGCommand::enqueueImpQueue() {
           MQueue->getHandleRef(),
           /*num_events_in_wait_list=*/0,
           /*event_wait_list=*/nullptr, &PreTimestampMarkerEvent);
+      OwnedPreTimestampMarkerEvent.emplace(PreTimestampMarkerEvent, Adapter,
+                                           /*TakeOwnership=*/true);
       TimestampDeps = &PreTimestampMarkerEvent;
       NumTimestampDeps = 1;
     }
 
-    Adapter.call<UrApiKind::urEnqueueTimestampRecordingExp>(
-        MQueue->getHandleRef(),
-        /*blocking=*/false, NumTimestampDeps, TimestampDeps, Event);
-
-    // If the queue is not in-order, we need to insert a barrier. This barrier
-    // does not need output events as it will implicitly enforce the following
-    // enqueue is blocked until it finishes.
-    if (!IsInOrderQueue) {
-      // We also need to release the timestamp event from the marker.
-      Adapter.call<UrApiKind::urEventRelease>(PreTimestampMarkerEvent);
-      // FIXME: Due to a bug in the L0 UR adapter, we will leak events if we do
-      //        not pass an output event to the UR call. Once that is fixed,
-      //        this immediately-deleted event can be removed.
-      ur_event_handle_t PostTimestampBarrierEvent{};
-      Adapter.call<UrApiKind::urEnqueueEventsWaitWithBarrier>(
-          MQueue->getHandleRef(),
-          /*num_events_in_wait_list=*/0,
-          /*event_wait_list=*/nullptr, &PostTimestampBarrierEvent);
-      Adapter.call<UrApiKind::urEventRelease>(PostTimestampBarrierEvent);
+    // Try to record a device timestamp natively. Not every backend can do so:
+    // the OpenCL backend can only record a reliable timestamp on a
+    // profiling-enabled queue (see intel/llvm#22229). When the recording is
+    // unsupported we fall back to a plain barrier, whose event still carries
+    // (best-effort) profiling information and provides the same ordering.
+    ur_result_t TimestampResult =
+        Adapter.call_nocheck<UrApiKind::urEnqueueTimestampRecordingExp>(
+            MQueue->getHandleRef(),
+            /*blocking=*/false, NumTimestampDeps, TimestampDeps, Event);
+
+    if (TimestampResult == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) {
+      if (!IsInOrderQueue) {
+        // The pre-timestamp barrier already provides the required ordering and
+        // profiling information, so reuse its event instead of submitting a
+        // second barrier.
+        if (Event)
+          *Event = OwnedPreTimestampMarkerEvent->TransferOwnership();
+      } else {
+        // An in-order queue has no pre-timestamp barrier to reuse.
+        if (auto Result =
+                Adapter.call_nocheck<UrApiKind::urEnqueueEventsWaitWithBarrier>(
+                    MQueue->getHandleRef(),
+                    /*num_events_in_wait_list=*/0,
+                    /*event_wait_list=*/nullptr, Event);
+            Result != UR_RESULT_SUCCESS)
+          return Result;
+      }
+    } else {
+      if (TimestampResult != UR_RESULT_SUCCESS)
+        return TimestampResult;
+
+      // If the queue is not in-order, we need to insert a barrier. This barrier
+      // does not need output events as it will implicitly enforce the following
+      // enqueue is blocked until it finishes.
+      if (!IsInOrderQueue) {
+        // FIXME: Due to a bug in the L0 UR adapter, we will leak events if we
+        //        do not pass an output event to the UR call. Once that is
+        //        fixed, this immediately-deleted event can be removed.
+        ur_event_handle_t PostTimestampBarrierEvent{};
+        Adapter.call<UrApiKind::urEnqueueEventsWaitWithBarrier>(
+            MQueue->getHandleRef(),
+            /*num_events_in_wait_list=*/0,
+            /*event_wait_list=*/nullptr, &PostTimestampBarrierEvent);
+        Adapter.call<UrApiKind::urEventRelease>(PostTimestampBarrierEvent);
+      }
     }
 
     SetEventHandleOrDiscard();

sycl/source/event.cpp

@@ -87,7 +87,9 @@ event::get_profiling_info() const {
                           "recording state.");
   }
 
-  if constexpr (!std::is_same_v<Param, info::event_profiling::command_submit>) {
+  // A profiling tag uses its device-recorded completion time for submit too.
+  if (!std::is_same_v<Param, info::event_profiling::command_submit> ||
+      impl->isProfilingTagEvent()) {
     impl->wait();
   }
   return impl->template get_profiling_info<Param>();

sycl/test-e2e/ProfilingTag/common.hpp

@@ -32,10 +32,10 @@ int run_test_on_queue(sycl::queue &Queue) {
   sycl::event EndTagE =
       sycl::ext::oneapi::experimental::submit_profiling_tag(Queue);
 
-  Queue.wait();
-
   int Failures = 0;
 
+  // Query submit without an explicit wait: the tag timestamp must be ready
+  // before the profiling query returns.
   uint64_t StartTagSubmit =
       StartTagE
           .get_profiling_info<sycl::info::event_profiling::command_submit>();
@@ -63,14 +63,14 @@ int run_test_on_queue(sycl::queue &Queue) {
   CHECK(Failures, StartTagEnd != 0)
   CHECK(Failures, EndTagSubmit != 0)
   CHECK(Failures, EndTagStart != 0)
-  CHECK(Failures, StartTagSubmit != 0)
+  CHECK(Failures, EndTagEnd != 0)
 
   CHECK(Failures, StartTagSubmit <= StartTagEnd)
   CHECK(Failures, StartTagSubmit <= StartTagStart)
-  CHECK(Failures, StartTagStart <= StartTagEnd)
+  CHECK(Failures, StartTagStart == StartTagEnd)
   CHECK(Failures, EndTagSubmit <= EndTagEnd)
   CHECK(Failures, EndTagSubmit <= EndTagStart)
-  CHECK(Failures, EndTagStart <= EndTagEnd)
+  CHECK(Failures, EndTagStart == EndTagEnd)
   CHECK(Failures, StartTagEnd <= EndTagEnd)
 
   if (Queue.has_property<sycl::property::queue::enable_profiling>()) {

sycl/test-e2e/ProfilingTag/in_order_profiling_queue.cpp

@@ -8,11 +8,6 @@
 // Note: Extension should work even on devices that do not support the
 //       ext_oneapi_queue_profiling_tag aspect.
 
-// Bug in OpenCL GPU driver causes fallback solution to return end time later
-// than the submission of the following work.
-// UNSUPPORTED: opencl && gpu
-// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/22229
-
 // HIP backend currently returns invalid values for submission time queries.
 // UNSUPPORTED: hip
 // UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/12904

sycl/test-e2e/ProfilingTag/profiling_queue.cpp

@@ -8,11 +8,6 @@
 // Note: Extension should work even on devices that do not support the
 //       ext_oneapi_queue_profiling_tag aspect.
 
-// Bug in OpenCL GPU driver causes fallback solution to return end time later
-// than the submission of the following work.
-// UNSUPPORTED: opencl && gpu
-// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/22229
-
 // HIP backend currently returns invalid values for submission time queries.
 // UNSUPPORTED: hip
 // UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/12904