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

…14123)

Current implementation of profiling info for NOP barriers is
inconsistent
with other events from the same queue (e.g., if the previous event
started
after the barrier was submitted). To make them consistent while keeping
the optimization, we would need to duplicate the event on our side and
make the duplicate check and potentially use profiling info of its
previous event.

Instead, as the first step, disable the NOP optimization during
profiling
since profiling is known to incur a performance hit anyway. The proper
duplicate event approach can be implemented as a follow up if this
causes issues for users.

Partially reverts #12949

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);
+}