[SYCL] Disable in-order queue barrier optimization while profiling

sycl/source/detail/event_impl.cpp

@@ -167,15 +167,11 @@ event_impl::event_impl(sycl::detail::pi::PiEvent Event,
   }
 }
 
-event_impl::event_impl(const QueueImplPtr &Queue) {
+event_impl::event_impl(const QueueImplPtr &Queue)
+    : MQueue{Queue},
+      MIsProfilingEnabled{Queue->is_host() || Queue->MIsProfilingEnabled},
+      MFallbackProfiling{MIsProfilingEnabled && Queue->isProfilingFallback()} {
   this->setContextImpl(Queue->getContextImplPtr());
-  this->associateWithQueue(Queue);
-}
-
-void event_impl::associateWithQueue(const QueueImplPtr &Queue) {
-  MQueue = Queue;
-  MIsProfilingEnabled = Queue->is_host() || Queue->MIsProfilingEnabled;
-  MFallbackProfiling = MIsProfilingEnabled && Queue->isProfilingFallback();
   if (Queue->is_host()) {
     MState.store(HES_NotComplete);
     if (Queue->has_property<property::queue::enable_profiling>()) {
@@ -337,11 +333,6 @@ template <>
 uint64_t
 event_impl::get_profiling_info<info::event_profiling::command_start>() {
   checkProfilingPreconditions();
-
-  // For nop command start time is equal to submission time.
-  if (isNOP() && MSubmitTime)
-    return MSubmitTime;
-
   if (!MHostEvent) {
     if (MEvent) {
       auto StartTime =
@@ -369,11 +360,6 @@ event_impl::get_profiling_info<info::event_profiling::command_start>() {
 template <>
 uint64_t event_impl::get_profiling_info<info::event_profiling::command_end>() {
   checkProfilingPreconditions();
-
-  // For nop command end time is equal to submission time.
-  if (isNOP() && MSubmitTime)
-    return MSubmitTime;
-
   if (!MHostEvent) {
     if (MEvent) {
       auto EndTime =

sycl/source/detail/event_impl.hpp

@@ -244,11 +244,6 @@ class event_impl {
     MSubmittedQueue = SubmittedQueue;
   };
 
-  /// Associate event with provided queue.
-  ///
-  /// @return
-  void associateWithQueue(const QueueImplPtr &Queue);
-
   /// Indicates if this event is not associated with any command and doesn't
   /// have native handle.
   ///

sycl/source/queue.cpp

@@ -214,22 +214,7 @@ getBarrierEventForInorderQueueHelper(const detail::QueueImplPtr QueueImpl) {
   assert(!QueueImpl->getCommandGraph() &&
          "Should not be called in on graph recording.");
 
-  auto LastEvent = QueueImpl->getLastEvent();
-  if (QueueImpl->MDiscardEvents) {
-    std::cout << "Discard event enabled" << std::endl;
-    return LastEvent;
-  }
-
-  auto LastEventImpl = detail::getSyclObjImpl(LastEvent);
-  // If last event is default constructed event then we want to associate it
-  // with the queue and record submission time if profiling is enabled. Such
-  // event corresponds to NOP and its submit time is same as start time and
-  // end time.
-  if (!LastEventImpl->isContextInitialized()) {
-    LastEventImpl->associateWithQueue(QueueImpl);
-    LastEventImpl->setSubmissionTime();
-  }
-  return detail::createSyclObjFromImpl<event>(LastEventImpl);
+  return QueueImpl->getLastEvent();
 }
 
 /// Prevents any commands submitted afterward to this queue from executing
@@ -240,7 +225,7 @@ getBarrierEventForInorderQueueHelper(const detail::QueueImplPtr QueueImpl) {
 /// \return a SYCL event object, which corresponds to the queue the command
 /// group is being enqueued on.
 event queue::ext_oneapi_submit_barrier(const detail::code_location &CodeLoc) {
-  if (is_in_order() && !impl->getCommandGraph())
+  if (is_in_order() && !impl->getCommandGraph() && !impl->MIsProfilingEnabled)
     return getBarrierEventForInorderQueueHelper(impl);
 
   return submit([=](handler &CGH) { CGH.ext_oneapi_barrier(); }, CodeLoc);
@@ -262,7 +247,8 @@ event queue::ext_oneapi_submit_barrier(const std::vector<event> &WaitList,
         auto EventImpl = detail::getSyclObjImpl(Event);
         return !EventImpl->isContextInitialized() || EventImpl->isNOP();
       });
-  if (is_in_order() && !impl->getCommandGraph() && AllEventsEmptyOrNop)
+  if (is_in_order() && !impl->getCommandGraph() && !impl->MIsProfilingEnabled &&
+      AllEventsEmptyOrNop)
     return getBarrierEventForInorderQueueHelper(impl);
 
   return submit([=](handler &CGH) { CGH.ext_oneapi_barrier(WaitList); },

sycl/test-e2e/Regression/in_order_barrier_profiling.cpp

@@ -0,0 +1,42 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+//==----------------- in_order_barrier_profiling.cpp -----------------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+// Level Zero adapter has a similar in-order queue barrier optimization that
+// leads to incorrect profiling values.
+// UNSUPPORTED: level_zero
+#include <sycl/detail/core.hpp>
+
+#include <sycl/properties/all_properties.hpp>
+
+using namespace sycl;
+
+// Checks that the barrier profiling info is consistent with the previous
+// command, despite the fact that the latter started after the barrier was
+// submitted.
+int main() {
+  queue Q({property::queue::in_order(), property::queue::enable_profiling()});
+
+  buffer<int, 1> Buf(range<1>(1));
+  event KernelEvent;
+  event BarrierEvent;
+  {
+    auto HostAcc = Buf.get_access();
+    KernelEvent = Q.submit([&](handler &cgh) {
+      auto Acc = Buf.get_access(cgh);
+      cgh.single_task([=]() {});
+    });
+    BarrierEvent = Q.ext_oneapi_submit_barrier();
+  }
+  uint64_t KernelEnd =
+      KernelEvent.get_profiling_info<info::event_profiling::command_end>();
+  uint64_t BarrierStart =
+      BarrierEvent.get_profiling_info<info::event_profiling::command_start>();
+  assert(KernelEnd <= BarrierStart);
+}