Defer barrier optimization path due to timestamp bugs

Signed-off-by: Larsen, Steffen <[email protected]>

Author: steffenlarsen

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

@@ -22,14 +22,6 @@ 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)) {
-    // If the queue is out-of-order and profiling is enabled, the implementation
-    // can save some operations by just using the required barrier event
-    // directly.
-    if (!Queue.is_in_order() &&
-        Queue.has_property<sycl::property::queue::enable_profiling>())
-      return Queue.ext_oneapi_submit_barrier();
-
-    // Otherwise, we use the internal implementation of the profiling tag.
     return Queue.submit(
         [=](handler &CGH) {
           sycl::detail::HandlerAccess::internalProfilingTagImpl(CGH);

sycl/source/detail/scheduler/commands.cpp

@@ -3448,9 +3448,6 @@ ur_result_t ExecCGCommand::enqueueImpQueue() {
 
     // If the queue is not in-order, the implementation will need to first
     // insert a marker event that the timestamp waits for.
-    // Note that with newer versions, MQueue should never have event profiling
-    // enabled here, as an optimization path in the headers will simply return
-    // the event from a barrier.
     ur_event_handle_t PreTimestampMarkerEvent{};
     if (!IsInOrderQueue) {
       // FIXME: urEnqueueEventsWait on the L0 adapter requires a double-release.