[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/scheduler/commands.cpp

@@ -3651,17 +3651,13 @@ 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{};
+    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.
@@ -3674,25 +3670,51 @@ ur_result_t ExecCGCommand::enqueueImpQueue() {
       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) {
+      // Release the marker the (unused) timestamp recording would have waited
+      // for.
+      if (!IsInOrderQueue)
+        Adapter.call<UrApiKind::urEventRelease>(PreTimestampMarkerEvent);
+      // A barrier with an empty wait list waits for all previously-submitted
+      // work and blocks subsequent work, providing the same ordering semantics
+      // as a native profiling tag.
+      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) {
+        // 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);
+      }
     }
 
     SetEventHandleOrDiscard();

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

sycl/unittests/Extensions/ProfilingTag.cpp

@@ -30,6 +30,14 @@ inline ur_result_t after_urEnqueueTimestampRecordingExp(void *) {
   return UR_RESULT_SUCCESS;
 }
 
+// Simulates a backend that cannot record a device timestamp (e.g. the OpenCL
+// backend on a queue without profiling enabled), so that the scheduler falls
+// back to a barrier.
+inline ur_result_t replace_urEnqueueTimestampRecordingExpUnsupported(void *) {
+  ++counter_urEnqueueTimestampRecordingExp;
+  return UR_RESULT_ERROR_UNSUPPORTED_FEATURE;
+}
+
 thread_local std::optional<ur_profiling_info_t> LatestProfilingQuery;
 inline ur_result_t after_urEventGetProfilingInfo(void *pParams) {
   auto &Params =
@@ -205,24 +213,57 @@ TEST_F(ProfilingTagTest, ProfilingTagFallbackDefaultQueue) {
   }
 }
 
-TEST_F(ProfilingTagTest, ProfilingTagFallbackProfilingQueue) {
+// Without the aspect but with profiling enabled, the tag is still serviced via
+// the profiling-tag command group. If the backend can record a device
+// timestamp (as the OpenCL backend now does on a profiling-enabled queue), the
+// native recording path is used rather than a bare barrier.
+TEST_F(ProfilingTagTest, ProfilingTagFallbackProfilingQueueTimestamp) {
   mock::getCallbacks().set_after_callback("urDeviceGetInfo",
                                           &after_urDeviceGetInfo<false>);
   mock::getCallbacks().set_after_callback(
       "urEnqueueTimestampRecordingExp", &after_urEnqueueTimestampRecordingExp);
   mock::getCallbacks().set_after_callback(
-      "urEnqueueEventsWaitWithBarrierExt",
+      "urEnqueueEventsWaitWithBarrier",
       &after_urEnqueueEventsWaitWithBarrierExt);
 
   sycl::context Ctx{sycl::platform()};
   sycl::queue Queue{Ctx,
                     sycl::default_selector_v,
-                    {sycl::property::queue::enable_profiling()}};
+                    {sycl::property::queue::enable_profiling(),
+                     sycl::property::queue::in_order()}};
+  sycl::device Dev = Queue.get_device();
+
+  ASSERT_FALSE(Dev.has(sycl::aspect::ext_oneapi_queue_profiling_tag));
+
+  sycl::event E = sycl::ext::oneapi::experimental::submit_profiling_tag(Queue);
+  ASSERT_EQ(size_t{1}, counter_urEnqueueTimestampRecordingExp);
+  ASSERT_EQ(size_t{0}, counter_urEnqueueEventsWaitWithBarrierExt);
+}
+
+// If the backend reports that it cannot record a device timestamp, the
+// scheduler must gracefully fall back to a barrier.
+TEST_F(ProfilingTagTest, ProfilingTagFallbackProfilingQueueBarrier) {
+  mock::getCallbacks().set_after_callback("urDeviceGetInfo",
+                                          &after_urDeviceGetInfo<false>);
+  mock::getCallbacks().set_replace_callback(
+      "urEnqueueTimestampRecordingExp",
+      &replace_urEnqueueTimestampRecordingExpUnsupported);
+  mock::getCallbacks().set_after_callback(
+      "urEnqueueEventsWaitWithBarrier",
+      &after_urEnqueueEventsWaitWithBarrierExt);
+
+  sycl::context Ctx{sycl::platform()};
+  sycl::queue Queue{Ctx,
+                    sycl::default_selector_v,
+                    {sycl::property::queue::enable_profiling(),
+                     sycl::property::queue::in_order()}};
   sycl::device Dev = Queue.get_device();
 
   ASSERT_FALSE(Dev.has(sycl::aspect::ext_oneapi_queue_profiling_tag));
 
   sycl::event E = sycl::ext::oneapi::experimental::submit_profiling_tag(Queue);
-  ASSERT_EQ(size_t{0}, counter_urEnqueueTimestampRecordingExp);
+  // The timestamp recording is attempted and reports unsupported, so a single
+  // barrier is used as the fallback on this in-order queue.
+  ASSERT_EQ(size_t{1}, counter_urEnqueueTimestampRecordingExp);
   ASSERT_EQ(size_t{1}, counter_urEnqueueEventsWaitWithBarrierExt);
 }

unified-runtime/source/adapters/opencl/context.hpp

@@ -13,6 +13,7 @@
 #include "common/ur_ref_count.hpp"
 #include "device.hpp"
 
+#include <mutex>
 #include <vector>
 
 namespace ur::opencl {
@@ -25,6 +26,15 @@ struct ur_context_handle_t_ : handle_base {
   bool IsNativeHandleOwned = true;
   ur::RefCount RefCount;
 
+  // Small device buffer used to emulate a timestamp recording via a real,
+  // profilable command (see urEnqueueTimestampRecordingExp). Some OpenCL
+  // drivers timestamp synchronization-only commands (barriers/markers) at the
+  // point they are inserted into the pipeline rather than when the preceding
+  // work actually completes, which breaks profiling-tag orderings
+  // (see https://github.com/intel/llvm/issues/22229). Created lazily.
+  std::mutex TimestampRecordingBufferMutex;
+  cl_mem TimestampRecordingBuffer = nullptr;
+
   ur_context_handle_t_(const ur_context_handle_t_ &) = delete;
   ur_context_handle_t_ &operator=(const ur_context_handle_t_ &) = delete;
 
@@ -40,6 +50,21 @@ struct ur_context_handle_t_ : handle_base {
   static ur_result_t makeWithNative(native_type Ctx, uint32_t DevCount,
                                     const ur_device_handle_t *phDevices,
                                     ur_context_handle_t &Context);
+
+  // Returns the small internal buffer used by urEnqueueTimestampRecordingExp,
+  // creating it on first use. Thread-safe.
+  ur_result_t getTimestampRecordingBuffer(cl_mem *OutBuffer) {
+    std::lock_guard<std::mutex> Lock(TimestampRecordingBufferMutex);
+    if (!TimestampRecordingBuffer) {
+      cl_int CLErr = CL_SUCCESS;
+      TimestampRecordingBuffer = clCreateBuffer(
+          CLContext, CL_MEM_READ_WRITE, sizeof(cl_uint), nullptr, &CLErr);
+      CL_RETURN_ON_FAILURE(CLErr);
+    }
+    *OutBuffer = TimestampRecordingBuffer;
+    return UR_RESULT_SUCCESS;
+  }
+
   ~ur_context_handle_t_() noexcept {
     // If we're reasonably sure this context is about to be destroyed we should
     // clear the ext function pointer cache. This isn't foolproof sadly but it
@@ -50,6 +75,9 @@ struct ur_context_handle_t_ : handle_base {
     for (uint32_t i = 0; i < DeviceCount; i++) {
       ur::opencl::urDeviceRelease(cast(Devices[i]));
     }
+    if (TimestampRecordingBuffer) {
+      clReleaseMemObject(TimestampRecordingBuffer);
+    }
     if (IsNativeHandleOwned) {
       clReleaseContext(CLContext);
     }

unified-runtime/source/adapters/opencl/event.cpp

@@ -309,10 +309,52 @@ urEventSetCallback(ur_event_handle_t hEvent, ur_execution_info_t execStatus,
   return UR_RESULT_SUCCESS;
 }
 
-UR_APIEXPORT ur_result_t UR_APICALL
-urEnqueueTimestampRecordingExp(ur_queue_handle_t, bool, uint32_t,
-                               const ur_event_handle_t *, ur_event_handle_t *) {
-  return UR_RESULT_ERROR_UNSUPPORTED_FEATURE;
+UR_APIEXPORT ur_result_t UR_APICALL urEnqueueTimestampRecordingExp(
+    ur_queue_handle_t hQueue, bool Blocking, uint32_t numEventsInWaitList,
+    const ur_event_handle_t *phEventWaitList, ur_event_handle_t *phEvent) {
+  // OpenCL has no native "record device timestamp" primitive. Some drivers
+  // (notably the Intel GPU driver) timestamp synchronization-only commands
+  // (barriers/markers) at the point they are inserted into the pipeline rather
+  // than when the preceding work actually completes, which breaks profiling-tag
+  // orderings (see https://github.com/intel/llvm/issues/22229). We therefore
+  // emulate a timestamp by enqueuing a lightweight *real* device command (a
+  // small buffer fill), whose profiling timestamps do reflect completion of the
+  // preceding work.
+  //
+  // This relies on OpenCL command profiling, which is only available when the
+  // queue was created with CL_QUEUE_PROFILING_ENABLE. If profiling is not
+  // enabled we cannot honor the request and report it as unsupported so the
+  // caller can fall back.
+  auto Queue = cast(hQueue);
+
+  cl_command_queue_properties QueueProperties = 0;
+  CL_RETURN_ON_FAILURE(clGetCommandQueueInfo(
+      Queue->CLQueue, CL_QUEUE_PROPERTIES, sizeof(QueueProperties),
+      &QueueProperties, nullptr));
+  if (!(QueueProperties & CL_QUEUE_PROFILING_ENABLE))
+    return UR_RESULT_ERROR_UNSUPPORTED_FEATURE;
+
+  cl_mem Buffer = nullptr;
+  UR_RETURN_ON_FAILURE(Queue->Context->getTimestampRecordingBuffer(&Buffer));
+
+  std::vector<cl_event> CLWaitEvents(numEventsInWaitList);
+  for (uint32_t I = 0; I < numEventsInWaitList; I++)
+    CLWaitEvents[I] = cast(phEventWaitList[I])->CLEvent;
+
+  const cl_uint Pattern = 0;
+  cl_event Event = nullptr;
+  CL_RETURN_ON_FAILURE(clEnqueueFillBuffer(
+      Queue->CLQueue, Buffer, &Pattern, sizeof(Pattern), /*offset=*/0,
+      /*size=*/sizeof(Pattern), numEventsInWaitList, CLWaitEvents.data(),
+      ifUrEvent(phEvent, Event)));
+
+  UR_RETURN_ON_FAILURE(
+      createUREvent(Event, cast(Queue->Context), cast(Queue), phEvent));
+
+  if (Blocking)
+    CL_RETURN_ON_FAILURE(clFinish(Queue->CLQueue));
+
+  return UR_RESULT_SUCCESS;
 }
 
 UR_APIEXPORT ur_result_t UR_APICALL