From 00b6389ac6d7ca6630b350b6911f5fd7107d274c Mon Sep 17 00:00:00 2001
From: Artur Gainullin <artur.gainullin@intel.com>
Date: Wed, 29 Jul 2026 20:50:19 -0700
Subject: [PATCH 1/5] [UR][L0v2][SYCL] Avoid redundant global timestamp query
 for profiling tag events

submit_profiling_tag and enqueue_signal_event already record a GPU-written
global timestamp via zeCommandListAppendWriteGlobalTimestamp. The additional
urDeviceGetGlobalTimestamps call used to capture the submission and start times
adds no value, since users only rely on the end timestamp.

Change the specification and implementation so that a tag event carries a
single timestamp, reported alike for command_submit, command_start and
command_end.

This drops the call in both the UR L0 v2 adapter and the SYCL RT, removes the
now-unused start-timestamp machinery.
Tests updated accordingly (the three fields are now equal for tag operations).

Assisted-By: Claude
---
 .../sycl_ext_oneapi_profiling_tag.asciidoc    | 31 +++------
 .../sycl_ext_oneapi_reusable_events.asciidoc  | 33 ++++------
 sycl/source/detail/event_impl.cpp             |  7 +-
 sycl/source/detail/queue_impl.cpp             |  2 +
 sycl/test-e2e/ProfilingTag/common.hpp         | 18 ++---
 .../level_zero/v2/command_list_manager.cpp    |  2 +-
 .../source/adapters/level_zero/v2/event.cpp   | 65 +++++++------------
 .../source/adapters/level_zero/v2/event.hpp   | 16 ++---
 .../adapters/level_zero/v2/queue_batched.cpp  | 13 ++--
 9 files changed, 68 insertions(+), 119 deletions(-)

diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_profiling_tag.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_profiling_tag.asciidoc
index d61a493e576c5..76ba1e1e2caae 100644
--- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_profiling_tag.asciidoc
+++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_profiling_tag.asciidoc
@@ -145,36 +145,25 @@ the "tag" command.
 In either case, the event's status becomes
 `info::event_command_status::complete` when all commands submitted to the queue
 prior to the call to `submit_profiling_tag` have completed.
-The event's `info::event_profiling::command_submit` timestamp reflects the
-time at which `submit_profiling_tag` is called.
-The event's `info::event_profiling::command_end` timestamp reflects the time
-at which the event enters the "complete" state.
-The event's `info::event_profiling::command_start` timestamp reflects the time
-that the profiling tag command starts executing.
-This timestamp is between the `info::event_profiling::command_submit` and
-`info::event_profiling::command_end` timestamps.
+The event's `info::event_profiling::command_submit`, `command_start`, and
+`command_end` timestamps are all equal to the device time at which the tag
+command executes.
+
+[_Note:_ A tag captures a single device timestamp, so all three fields report
+the same value.
+This value shares the same time base as other commands on the queue, so the
+elapsed device time between two tags is the difference of their timestamps.
+_{endnote}_]
 
 It is unspecified whether the event ever has the
 `info::event_command_status::running` status.
 Implementations are encouraged to transition the event directly from the
-"submitted" status to the "complete" status and are encouraged to set the
-"command_start" timestamp to the same value as the "command_end" timestamp.
+"submitted" status to the "complete" status.
 
 _Throws:_ A synchronous `exception` with the `errc::invalid` error code if
 the queue's device does not have the aspect `ext_oneapi_queue_profiling_tag`
 and the queue was not constructed with the `property::queue::enable_profiling`
 property.
-
-[_Note:_ In order to understand why the "command_start" and "command_end"
-timestamps are encouraged to be the same, think of the barrier as an empty
-kernel with an implicit set of dependencies on all previous commands in the
-same queue.
-This theoretical kernel starts executing when the dependencies are resolved.
-Since the kernel is empty, the end time is the same as the start time.
-The "command_start" and "command_end" timestamps are not required to be the
-same, though, in order to accommodate an implementation where the barrier is
-implemented by submitting an actual kernel, which has non-zero execution time.
-_{endnote}_]
 |====
 
 
diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_reusable_events.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_reusable_events.asciidoc
index 4f01a25864ac4..ff47f14fa9edf 100644
--- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_reusable_events.asciidoc
+++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_reusable_events.asciidoc
@@ -274,20 +274,21 @@ After the barrier completes, a "tag" operation sets the status of `evt` to
 The event's timestamp information is also set if the queue `q` was created with
 the `property::queue::enable_profiling` property or if the event `e` was created
 with the `ext::oneapi::experimental::enable_profiling` property.
-The event's `info::event_profiling::command_submit` timestamp reflects the time
-at which `enqueue_signal_event` is called.
-The event's `info::event_profiling::command_end` timestamp reflects the time at
-which the event enters the "complete" state.
-The event's `info::event_profiling::command_start` timestamp reflects the time
-that the tag operation starts executing.
-This timestamp is between the `info::event_profiling::command_submit` and
-`info::event_profiling::command_end` timestamps.
+When set, the event's `info::event_profiling::command_submit`,
+`command_start`, and `command_end` timestamps are all equal to the device time
+at which the tag operation executes.
+
+[_Note:_ A tag captures a single device timestamp, so all three fields report
+the same value.
+This value shares the same time base as other commands on the queue, so the
+elapsed device time between two tag operations is the difference of their
+timestamps.
+_{endnote}_]
 
 It is unspecified whether the event ever has the
 `info::event_command_status::running` status.
 Implementations are encouraged to transition the event directly from the
-"submitted" status to the "complete" status and are encouraged to set the
-"command_start" timestamp to the same value as the "command_end" timestamp.
+"submitted" status to the "complete" status.
 
 _Throws:_
 
@@ -298,18 +299,6 @@ _Throws:_
    timestamps and if the device associated with `q` does not have
    `aspect::ext_oneapi_per_event_profiling`.
 
-[_Note:_ In order to understand why the "command_start" and "command_end"
-timestamps are encouraged to be the same, think of the tag operation as an empty
-kernel with an implicit set of dependencies on all previous commands in the
-same queue.
-This theoretical kernel starts executing when the dependencies are resolved.
-Since the kernel is empty, the end time is the same as the start time.
-The "command_start" and "command_end" timestamps are not required to be the
-same, though, in order to accommodate an implementation where the tag operation
-is implemented by submitting an actual kernel, which has non-zero execution
-time.
-_{endnote}_]
-
 === Interaction with other event APIs
 
 An event _E_ created via `make_event` can be used as a command dependency (e.g.
diff --git a/sycl/source/detail/event_impl.cpp b/sycl/source/detail/event_impl.cpp
index 4084486139af3..b3dac0b6e6720 100644
--- a/sycl/source/detail/event_impl.cpp
+++ b/sycl/source/detail/event_impl.cpp
@@ -410,8 +410,7 @@ uint64_t
 event_impl::get_profiling_info<info::event_profiling::command_submit>() {
   checkProfilingPreconditions();
   if (isProfilingTagEvent()) {
-    // For profiling tag events we rely on the submission time reported as
-    // the start time has undefined behavior.
+    // A tag event has a single device timestamp; all fields return it.
     return get_event_profiling_info<info::event_profiling::command_submit>(
         this->getHandle(), this->getAdapter());
   }
@@ -608,7 +607,9 @@ void event_impl::cleanDepEventsThroughOneLevel() {
 }
 
 void event_impl::setSubmissionTime() {
-  if (!MIsProfilingEnabled && !MProfilingTagEvent)
+  // Tag events use their single device timestamp for all fields, so there is
+  // no host-side submission time to record.
+  if (!MIsProfilingEnabled || MProfilingTagEvent)
     return;
 
   if (std::shared_ptr<queue_impl> Queue =
diff --git a/sycl/source/detail/queue_impl.cpp b/sycl/source/detail/queue_impl.cpp
index b8bd5c3bc620f..4c92a53a4030c 100644
--- a/sycl/source/detail/queue_impl.cpp
+++ b/sycl/source/detail/queue_impl.cpp
@@ -518,6 +518,8 @@ EventImplPtr queue_impl::submit_barrier_scheduler_bypass(
     ResEvent->setWorkerQueue(weak_from_this());
     ResEvent->setPotentiallyNativeRecorded(
         getContextImpl().isNativeRecordingActive());
+    if (EventForReuse)
+      ResEvent->markAsProfilingTagEvent();
     ResEvent->setSubmissionTime();
     ResEvent->setEnqueued();
     ResEvent->setStateIncomplete();
diff --git a/sycl/test-e2e/ProfilingTag/common.hpp b/sycl/test-e2e/ProfilingTag/common.hpp
index 9aab2e8876eb1..97bf89e8cb481 100644
--- a/sycl/test-e2e/ProfilingTag/common.hpp
+++ b/sycl/test-e2e/ProfilingTag/common.hpp
@@ -58,19 +58,13 @@ int run_test_on_queue(sycl::queue &Queue) {
   std::cout << "EndTagStart: " << EndTagStart << std::endl;
   std::cout << "EndTagEnd: " << EndTagEnd << std::endl;
 
-  CHECK(Failures, StartTagSubmit != 0)
-  CHECK(Failures, StartTagStart != 0)
+  // All three timestamps are equal to the GPU-written completion time.
   CHECK(Failures, StartTagEnd != 0)
-  CHECK(Failures, EndTagSubmit != 0)
-  CHECK(Failures, EndTagStart != 0)
-  CHECK(Failures, StartTagSubmit != 0)
-
-  CHECK(Failures, StartTagSubmit <= StartTagEnd)
-  CHECK(Failures, StartTagSubmit <= StartTagStart)
-  CHECK(Failures, StartTagStart <= StartTagEnd)
-  CHECK(Failures, EndTagSubmit <= EndTagEnd)
-  CHECK(Failures, EndTagSubmit <= EndTagStart)
-  CHECK(Failures, EndTagStart <= EndTagEnd)
+  CHECK(Failures, EndTagEnd != 0)
+  CHECK(Failures, StartTagSubmit == StartTagStart)
+  CHECK(Failures, StartTagStart == StartTagEnd)
+  CHECK(Failures, EndTagSubmit == EndTagStart)
+  CHECK(Failures, EndTagStart == EndTagEnd)
   CHECK(Failures, StartTagEnd <= EndTagEnd)
 
   if (Queue.has_property<sycl::property::queue::enable_profiling>()) {
diff --git a/unified-runtime/source/adapters/level_zero/v2/command_list_manager.cpp b/unified-runtime/source/adapters/level_zero/v2/command_list_manager.cpp
index de4defaad57cd..16ef6f96d8bbc 100644
--- a/unified-runtime/source/adapters/level_zero/v2/command_list_manager.cpp
+++ b/unified-runtime/source/adapters/level_zero/v2/command_list_manager.cpp
@@ -602,7 +602,7 @@ ur_result_t ur_command_list_manager::appendTimestampRecordingExp(
 
   auto [pWaitEvents, numWaitEvents, _] = waitListView;
 
-  phEvent->recordStartTimestamp();
+  phEvent->initTimestampRecording();
 
   auto [timestampPtr, zeSignalEvent] = phEvent->getEventEndTimestampAndHandle();
 
diff --git a/unified-runtime/source/adapters/level_zero/v2/event.cpp b/unified-runtime/source/adapters/level_zero/v2/event.cpp
index ffef28393d829..c56436055004b 100644
--- a/unified-runtime/source/adapters/level_zero/v2/event.cpp
+++ b/unified-runtime/source/adapters/level_zero/v2/event.cpp
@@ -56,44 +56,38 @@ uint64_t event_profiling_data_t::getEventEndTimestamp() {
   assert(zeTimerResolution);
   assert(timestampMaxValue);
 
-  adjustedEventEndTimestamp = adjustEndEventTimestamp(
-      adjustedEventStartTimestamp, recordEventEndTimestamp, timestampMaxValue,
-      zeTimerResolution);
+  // A timestamp-recording event holds a single GPU-written global timestamp,
+  // so there is no separate start value to detect a wrap-around against.
+  adjustedEventEndTimestamp =
+      (recordEventEndTimestamp & timestampMaxValue) * zeTimerResolution;
 
   return adjustedEventEndTimestamp;
 }
 
 void event_profiling_data_t::reset() {
-  // This ensures that the event is consider as not timestamped.
-  // We can't touch the recordEventEndTimestamp
-  // as it may still be overwritten by the driver.
-  // In case event is resued and recordStartTimestamp
-  // is called again, adjustedEventEndTimestamp will always be updated correctly
-  // to the new value as we wait for the event to be signaled.
-  // If the event is reused on another queue, this means that the original
-  // queue must have been destroyed (and the even pool released back to the
-  // context) and the timstamp is already wrriten, so there's no race-condition
-  // possible.
-  adjustedEventStartTimestamp = 0;
+  // This ensures that the event is considered as not timestamped.
+  // We can't touch recordEventEndTimestamp as it may still be overwritten by
+  // the driver. When the event is reused and initTimestampRecording is called
+  // again, adjustedEventEndTimestamp will be recomputed correctly as we wait
+  // for the event to be signaled before reading it.
+  // If the event is reused on another queue, the original queue must have been
+  // destroyed (and the event pool released back to the context) and the
+  // timestamp is already written, so no race-condition is possible.
   adjustedEventEndTimestamp = 0;
   timestampRecorded = false;
 }
 
-void event_profiling_data_t::recordStartTimestamp(ur_device_handle_t hDevice) {
+void event_profiling_data_t::initTimestampRecording(
+    ur_device_handle_t hDevice) {
   zeTimerResolution = hDevice->getTimerResolution();
   timestampMaxValue = hDevice->getTimestampMask();
-
-  uint64_t deviceStartTimestamp = 0;
-  UR_CALL_THROWS(ur::level_zero::urDeviceGetGlobalTimestamps(
-      common_cast(hDevice), &deviceStartTimestamp, nullptr));
-
-  assert(adjustedEventStartTimestamp == 0);
-  adjustedEventStartTimestamp = deviceStartTimestamp;
   timestampRecorded = true;
 }
 
-uint64_t event_profiling_data_t::getEventStartTimestmap() const {
-  return adjustedEventStartTimestamp;
+void ur_event_handle_t_::initTimestampRecording() {
+  assert(hQueue);
+  assert(hDevice);
+  profilingData.initTimestampRecording(hDevice);
 }
 
 bool event_profiling_data_t::recordingStarted() const {
@@ -145,18 +139,6 @@ void ur_event_handle_t_::onWaitListUse() {
   }
 }
 
-void ur_event_handle_t_::recordStartTimestamp() {
-  // queue and device must be set before calling this
-  assert(hQueue);
-  assert(hDevice);
-
-  profilingData.recordStartTimestamp(hDevice);
-}
-
-uint64_t ur_event_handle_t_::getEventStartTimestmap() const {
-  return profilingData.getEventStartTimestmap();
-}
-
 uint64_t ur_event_handle_t_::getEventEndTimestamp() {
   return profilingData.getEventEndTimestamp();
 }
@@ -342,19 +324,16 @@ ur_result_t urEventGetProfilingInfo(
 
   UrReturnHelper returnValue(propValueSize, pPropValue, pPropValueSizeRet);
 
-  // For timestamped events we have the timestamps ready directly on the event
-  // handle, so we short-circuit the return.
+  // For timestamp-recording events the GPU wrote a single global timestamp.
+  // All profiling fields return that value: submit == start == end.
   if (isTimestampedEvent) {
-    uint64_t contextStartTime = event->getEventStartTimestmap();
     switch (propName) {
     case UR_PROFILING_INFO_COMMAND_QUEUED:
     case UR_PROFILING_INFO_COMMAND_SUBMIT:
-      return returnValue(contextStartTime);
-    case UR_PROFILING_INFO_COMMAND_END:
     case UR_PROFILING_INFO_COMMAND_START:
-    case UR_PROFILING_INFO_COMMAND_COMPLETE: {
+    case UR_PROFILING_INFO_COMMAND_END:
+    case UR_PROFILING_INFO_COMMAND_COMPLETE:
       return returnValue(event->getEventEndTimestamp());
-    }
     default:
       UR_LOG(ERR, "urEventGetProfilingInfo: not supported ParamName");
       return UR_RESULT_ERROR_INVALID_VALUE;
diff --git a/unified-runtime/source/adapters/level_zero/v2/event.hpp b/unified-runtime/source/adapters/level_zero/v2/event.hpp
index f7e6fe0a83c64..0c30ca81ed43a 100644
--- a/unified-runtime/source/adapters/level_zero/v2/event.hpp
+++ b/unified-runtime/source/adapters/level_zero/v2/event.hpp
@@ -28,8 +28,9 @@ class event_pool;
 struct event_profiling_data_t {
   event_profiling_data_t(ze_event_handle_t hZeEvent) : hZeEvent(hZeEvent) {}
 
-  void recordStartTimestamp(ur_device_handle_t hDevice);
-  uint64_t getEventStartTimestmap() const;
+  // Prepare the event to hold a single GPU-written global timestamp: cache the
+  // device timer resolution/mask and mark the event as timestamp-recording.
+  void initTimestampRecording(ur_device_handle_t hDevice);
 
   uint64_t getEventEndTimestamp();
   uint64_t *eventEndTimestampAddr();
@@ -43,7 +44,6 @@ struct event_profiling_data_t {
 private:
   ze_event_handle_t hZeEvent;
 
-  uint64_t adjustedEventStartTimestamp = 0;
   uint64_t recordEventEndTimestamp = 0;
   uint64_t adjustedEventEndTimestamp = 0;
 
@@ -127,16 +127,14 @@ struct ur_event_handle_t_ : v2::ur_object_t {
   // Get the device associated with this event
   ur_device_handle_t getDevice() const;
 
-  // Record the start timestamp of the event, to be obtained by
-  // urEventGetProfilingInfo. setQueue should be
-  // called before this.
-  void recordStartTimestamp();
+  // Mark this event as recording a single GPU-written global timestamp,
+  // obtainable via urEventGetProfilingInfo. setQueue must be called first.
+  void initTimestampRecording();
 
-  // Get pointer to the end timestamp, and ze event handle.
+  // Get pointer to the timestamp storage, and ze event handle.
   // Caller is responsible for signaling the event once the timestamp is ready.
   std::pair<uint64_t *, ze_event_handle_t> getEventEndTimestampAndHandle();
 
-  uint64_t getEventStartTimestmap() const;
   uint64_t getEventEndTimestamp();
 
   ur::RefCount RefCount;
diff --git a/unified-runtime/source/adapters/level_zero/v2/queue_batched.cpp b/unified-runtime/source/adapters/level_zero/v2/queue_batched.cpp
index 88bebe21f10f5..10670dfc8f008 100644
--- a/unified-runtime/source/adapters/level_zero/v2/queue_batched.cpp
+++ b/unified-runtime/source/adapters/level_zero/v2/queue_batched.cpp
@@ -863,14 +863,11 @@ ur_result_t ur_queue_batched_t::bindlessImagesSignalExternalSemaphoreExp(
       getEvent(lockedBatch, phEvent));
 }
 
-// In case of queues with batched submissions, which use regular command lists
-// (similarly to command buffers), the start timestamp would be recorded as the
-// operation is submitted (event.recordStartTimestamp() in
-// appendTimestampRecordingExp does not use the queue but directly the device),
-// but the end timestamp would wait for the submission of the given regular
-// command list. The difference between the start and end timestamps would
-// reflect the delay in the batch submission, the difference between end
-// timestamps would reflect the actual time of execution.
+// Queues with batched submissions use regular command lists (similarly to
+// command buffers). The tag event records a single GPU-written global
+// timestamp via zeCommandListAppendWriteGlobalTimestamp, which becomes
+// available once the batch's regular command list is submitted. The difference
+// between two tag timestamps reflects the actual time of execution.
 //
 
 ur_result_t ur_queue_batched_t::enqueueTimestampRecordingExp(

From b84e6e64bdec8aec5d463863f522da0350947473 Mon Sep 17 00:00:00 2001
From: Artur Gainullin <artur.gainullin@intel.com>
Date: Fri, 31 Jul 2026 10:23:33 -0700
Subject: [PATCH 2/5] [SYCL] Revert profiling tag spec and test changes

The follow-up commits keep command_submit <= command_start <= command_end for
tag events (immediate queues report a single timestamp for all three; batched
queues still record a distinct submission time). The observable semantics thus
match the original specification, so the spec wording and ProfilingTag test
assertions are restored to their pre-change form.

Assisted-By: Claude
---
 .../sycl_ext_oneapi_profiling_tag.asciidoc    | 31 +++++++++++------
 .../sycl_ext_oneapi_reusable_events.asciidoc  | 33 ++++++++++++-------
 sycl/test-e2e/ProfilingTag/common.hpp         | 18 ++++++----
 3 files changed, 55 insertions(+), 27 deletions(-)

diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_profiling_tag.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_profiling_tag.asciidoc
index 76ba1e1e2caae..d61a493e576c5 100644
--- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_profiling_tag.asciidoc
+++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_profiling_tag.asciidoc
@@ -145,25 +145,36 @@ the "tag" command.
 In either case, the event's status becomes
 `info::event_command_status::complete` when all commands submitted to the queue
 prior to the call to `submit_profiling_tag` have completed.
-The event's `info::event_profiling::command_submit`, `command_start`, and
-`command_end` timestamps are all equal to the device time at which the tag
-command executes.
-
-[_Note:_ A tag captures a single device timestamp, so all three fields report
-the same value.
-This value shares the same time base as other commands on the queue, so the
-elapsed device time between two tags is the difference of their timestamps.
-_{endnote}_]
+The event's `info::event_profiling::command_submit` timestamp reflects the
+time at which `submit_profiling_tag` is called.
+The event's `info::event_profiling::command_end` timestamp reflects the time
+at which the event enters the "complete" state.
+The event's `info::event_profiling::command_start` timestamp reflects the time
+that the profiling tag command starts executing.
+This timestamp is between the `info::event_profiling::command_submit` and
+`info::event_profiling::command_end` timestamps.
 
 It is unspecified whether the event ever has the
 `info::event_command_status::running` status.
 Implementations are encouraged to transition the event directly from the
-"submitted" status to the "complete" status.
+"submitted" status to the "complete" status and are encouraged to set the
+"command_start" timestamp to the same value as the "command_end" timestamp.
 
 _Throws:_ A synchronous `exception` with the `errc::invalid` error code if
 the queue's device does not have the aspect `ext_oneapi_queue_profiling_tag`
 and the queue was not constructed with the `property::queue::enable_profiling`
 property.
+
+[_Note:_ In order to understand why the "command_start" and "command_end"
+timestamps are encouraged to be the same, think of the barrier as an empty
+kernel with an implicit set of dependencies on all previous commands in the
+same queue.
+This theoretical kernel starts executing when the dependencies are resolved.
+Since the kernel is empty, the end time is the same as the start time.
+The "command_start" and "command_end" timestamps are not required to be the
+same, though, in order to accommodate an implementation where the barrier is
+implemented by submitting an actual kernel, which has non-zero execution time.
+_{endnote}_]
 |====
 
 
diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_reusable_events.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_reusable_events.asciidoc
index ff47f14fa9edf..4f01a25864ac4 100644
--- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_reusable_events.asciidoc
+++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_reusable_events.asciidoc
@@ -274,21 +274,20 @@ After the barrier completes, a "tag" operation sets the status of `evt` to
 The event's timestamp information is also set if the queue `q` was created with
 the `property::queue::enable_profiling` property or if the event `e` was created
 with the `ext::oneapi::experimental::enable_profiling` property.
-When set, the event's `info::event_profiling::command_submit`,
-`command_start`, and `command_end` timestamps are all equal to the device time
-at which the tag operation executes.
-
-[_Note:_ A tag captures a single device timestamp, so all three fields report
-the same value.
-This value shares the same time base as other commands on the queue, so the
-elapsed device time between two tag operations is the difference of their
-timestamps.
-_{endnote}_]
+The event's `info::event_profiling::command_submit` timestamp reflects the time
+at which `enqueue_signal_event` is called.
+The event's `info::event_profiling::command_end` timestamp reflects the time at
+which the event enters the "complete" state.
+The event's `info::event_profiling::command_start` timestamp reflects the time
+that the tag operation starts executing.
+This timestamp is between the `info::event_profiling::command_submit` and
+`info::event_profiling::command_end` timestamps.
 
 It is unspecified whether the event ever has the
 `info::event_command_status::running` status.
 Implementations are encouraged to transition the event directly from the
-"submitted" status to the "complete" status.
+"submitted" status to the "complete" status and are encouraged to set the
+"command_start" timestamp to the same value as the "command_end" timestamp.
 
 _Throws:_
 
@@ -299,6 +298,18 @@ _Throws:_
    timestamps and if the device associated with `q` does not have
    `aspect::ext_oneapi_per_event_profiling`.
 
+[_Note:_ In order to understand why the "command_start" and "command_end"
+timestamps are encouraged to be the same, think of the tag operation as an empty
+kernel with an implicit set of dependencies on all previous commands in the
+same queue.
+This theoretical kernel starts executing when the dependencies are resolved.
+Since the kernel is empty, the end time is the same as the start time.
+The "command_start" and "command_end" timestamps are not required to be the
+same, though, in order to accommodate an implementation where the tag operation
+is implemented by submitting an actual kernel, which has non-zero execution
+time.
+_{endnote}_]
+
 === Interaction with other event APIs
 
 An event _E_ created via `make_event` can be used as a command dependency (e.g.
diff --git a/sycl/test-e2e/ProfilingTag/common.hpp b/sycl/test-e2e/ProfilingTag/common.hpp
index 97bf89e8cb481..9aab2e8876eb1 100644
--- a/sycl/test-e2e/ProfilingTag/common.hpp
+++ b/sycl/test-e2e/ProfilingTag/common.hpp
@@ -58,13 +58,19 @@ int run_test_on_queue(sycl::queue &Queue) {
   std::cout << "EndTagStart: " << EndTagStart << std::endl;
   std::cout << "EndTagEnd: " << EndTagEnd << std::endl;
 
-  // All three timestamps are equal to the GPU-written completion time.
+  CHECK(Failures, StartTagSubmit != 0)
+  CHECK(Failures, StartTagStart != 0)
   CHECK(Failures, StartTagEnd != 0)
-  CHECK(Failures, EndTagEnd != 0)
-  CHECK(Failures, StartTagSubmit == StartTagStart)
-  CHECK(Failures, StartTagStart == StartTagEnd)
-  CHECK(Failures, EndTagSubmit == EndTagStart)
-  CHECK(Failures, EndTagStart == EndTagEnd)
+  CHECK(Failures, EndTagSubmit != 0)
+  CHECK(Failures, EndTagStart != 0)
+  CHECK(Failures, StartTagSubmit != 0)
+
+  CHECK(Failures, StartTagSubmit <= StartTagEnd)
+  CHECK(Failures, StartTagSubmit <= StartTagStart)
+  CHECK(Failures, StartTagStart <= StartTagEnd)
+  CHECK(Failures, EndTagSubmit <= EndTagEnd)
+  CHECK(Failures, EndTagSubmit <= EndTagStart)
+  CHECK(Failures, EndTagStart <= EndTagEnd)
   CHECK(Failures, StartTagEnd <= EndTagEnd)
 
   if (Queue.has_property<sycl::property::queue::enable_profiling>()) {

From c5da4ec1ae2dc20b42734e4cfa968ae0c62caa6c Mon Sep 17 00:00:00 2001
From: Artur Gainullin <artur.gainullin@intel.com>
Date: Fri, 31 Jul 2026 10:23:44 -0700
Subject: [PATCH 3/5] [UR][L0v2][SYCL] Skip submission timestamp query for
 immediate-queue tags

The initial version reported command_submit == command_start == command_end for
all tag events and dropped the urDeviceGetGlobalTimestamps query. That query is
only redundant for immediate queues, where the tag is submitted right away.

Reintroduce the query for batched queues (ur_queue_batched_t), where submission
is deferred and command_submit meaningfully precedes execution, while keeping
immediate queues on the fast path (no query; the GPU-written timestamp is
reported for all fields). appendTimestampRecordingExp gains a recordSubmit flag
set by the queue. The recorded submission timestamp is also the reference for
end-timestamp wrap-around correction.

The SYCL RT no longer records a host submission time for tag events, relying on
the adapter for command_submit.

Assisted-By: Claude
---
 sycl/source/detail/event_impl.cpp             |  5 +-
 .../level_zero/v2/command_list_manager.cpp    |  5 +-
 .../level_zero/v2/command_list_manager.hpp    |  3 +-
 .../source/adapters/level_zero/v2/event.cpp   | 59 +++++++++++++------
 .../source/adapters/level_zero/v2/event.hpp   | 19 ++++--
 .../adapters/level_zero/v2/queue_batched.cpp  | 12 ++--
 .../v2/queue_immediate_in_order.hpp           |  3 +-
 .../v2/queue_immediate_out_of_order.hpp       |  3 +-
 8 files changed, 72 insertions(+), 37 deletions(-)

diff --git a/sycl/source/detail/event_impl.cpp b/sycl/source/detail/event_impl.cpp
index b3dac0b6e6720..786a86f877131 100644
--- a/sycl/source/detail/event_impl.cpp
+++ b/sycl/source/detail/event_impl.cpp
@@ -410,7 +410,7 @@ uint64_t
 event_impl::get_profiling_info<info::event_profiling::command_submit>() {
   checkProfilingPreconditions();
   if (isProfilingTagEvent()) {
-    // A tag event has a single device timestamp; all fields return it.
+    // Tag events report command_submit through the adapter.
     return get_event_profiling_info<info::event_profiling::command_submit>(
         this->getHandle(), this->getAdapter());
   }
@@ -607,8 +607,7 @@ void event_impl::cleanDepEventsThroughOneLevel() {
 }
 
 void event_impl::setSubmissionTime() {
-  // Tag events use their single device timestamp for all fields, so there is
-  // no host-side submission time to record.
+  // Tag events obtain command_submit from the adapter.
   if (!MIsProfilingEnabled || MProfilingTagEvent)
     return;
 
diff --git a/unified-runtime/source/adapters/level_zero/v2/command_list_manager.cpp b/unified-runtime/source/adapters/level_zero/v2/command_list_manager.cpp
index 16ef6f96d8bbc..6a239a145a0a4 100644
--- a/unified-runtime/source/adapters/level_zero/v2/command_list_manager.cpp
+++ b/unified-runtime/source/adapters/level_zero/v2/command_list_manager.cpp
@@ -593,7 +593,8 @@ ur_result_t ur_command_list_manager::appendUSMMemcpy2D(
 }
 
 ur_result_t ur_command_list_manager::appendTimestampRecordingExp(
-    bool blocking, wait_list_view &waitListView, ur_event_handle_t phEvent) {
+    bool blocking, wait_list_view &waitListView, ur_event_handle_t phEvent,
+    bool recordSubmit) {
   TRACK_SCOPE_LATENCY("ur_command_list_manager::appendTimestampRecordingExp");
 
   if (!phEvent) {
@@ -602,7 +603,7 @@ ur_result_t ur_command_list_manager::appendTimestampRecordingExp(
 
   auto [pWaitEvents, numWaitEvents, _] = waitListView;
 
-  phEvent->initTimestampRecording();
+  phEvent->initTimestampRecording(recordSubmit);
 
   auto [timestampPtr, zeSignalEvent] = phEvent->getEventEndTimestampAndHandle();
 
diff --git a/unified-runtime/source/adapters/level_zero/v2/command_list_manager.hpp b/unified-runtime/source/adapters/level_zero/v2/command_list_manager.hpp
index 8ba0bd797256e..f0fae46ca2ee1 100644
--- a/unified-runtime/source/adapters/level_zero/v2/command_list_manager.hpp
+++ b/unified-runtime/source/adapters/level_zero/v2/command_list_manager.hpp
@@ -212,7 +212,8 @@ struct ur_command_list_manager {
       const ur_event_handle_t *phEventWaitList, ur_event_handle_t phEvent);
   ur_result_t appendTimestampRecordingExp(bool blocking,
                                           wait_list_view &waitListView,
-                                          ur_event_handle_t phEvent);
+                                          ur_event_handle_t phEvent,
+                                          bool recordSubmit);
   ur_result_t
   appendCommandBufferExp(ur_exp_command_buffer_handle_t hCommandBuffer,
                          wait_list_view &waitListView,
diff --git a/unified-runtime/source/adapters/level_zero/v2/event.cpp b/unified-runtime/source/adapters/level_zero/v2/event.cpp
index c56436055004b..11d53ce393314 100644
--- a/unified-runtime/source/adapters/level_zero/v2/event.cpp
+++ b/unified-runtime/source/adapters/level_zero/v2/event.cpp
@@ -56,38 +56,58 @@ uint64_t event_profiling_data_t::getEventEndTimestamp() {
   assert(zeTimerResolution);
   assert(timestampMaxValue);
 
-  // A timestamp-recording event holds a single GPU-written global timestamp,
-  // so there is no separate start value to detect a wrap-around against.
+  // recordedSubmitTimestamp is the wrap-around reference; it is 0 when no
+  // submission timestamp was recorded, in which case no adjustment is applied.
   adjustedEventEndTimestamp =
-      (recordEventEndTimestamp & timestampMaxValue) * zeTimerResolution;
+      adjustEndEventTimestamp(recordedSubmitTimestamp, recordEventEndTimestamp,
+                              timestampMaxValue, zeTimerResolution);
 
   return adjustedEventEndTimestamp;
 }
 
 void event_profiling_data_t::reset() {
-  // This ensures that the event is considered as not timestamped.
-  // We can't touch recordEventEndTimestamp as it may still be overwritten by
-  // the driver. When the event is reused and initTimestampRecording is called
-  // again, adjustedEventEndTimestamp will be recomputed correctly as we wait
-  // for the event to be signaled before reading it.
-  // If the event is reused on another queue, the original queue must have been
-  // destroyed (and the event pool released back to the context) and the
-  // timestamp is already written, so no race-condition is possible.
+  // This ensures that the event is consider as not timestamped.
+  // We can't touch the recordEventEndTimestamp
+  // as it may still be overwritten by the driver.
+  // In case event is resued and initTimestampRecording
+  // is called again, adjustedEventEndTimestamp will always be updated correctly
+  // to the new value as we wait for the event to be signaled.
+  // If the event is reused on another queue, this means that the original
+  // queue must have been destroyed (and the even pool released back to the
+  // context) and the timstamp is already wrriten, so there's no race-condition
+  // possible.
   adjustedEventEndTimestamp = 0;
+  recordedSubmitTimestamp = 0;
   timestampRecorded = false;
 }
 
-void event_profiling_data_t::initTimestampRecording(
-    ur_device_handle_t hDevice) {
+void event_profiling_data_t::initTimestampRecording(ur_device_handle_t hDevice,
+                                                    bool recordSubmit) {
   zeTimerResolution = hDevice->getTimerResolution();
   timestampMaxValue = hDevice->getTimestampMask();
+
+  // Recording a submission timestamp requires an extra device query whose cost
+  // dominates the tag latency, so only do it when requested (batched queues).
+  if (recordSubmit) {
+    UR_CALL_THROWS(ur::level_zero::urDeviceGetGlobalTimestamps(
+        common_cast(hDevice), &recordedSubmitTimestamp, nullptr));
+  }
+
   timestampRecorded = true;
 }
 
-void ur_event_handle_t_::initTimestampRecording() {
+void ur_event_handle_t_::initTimestampRecording(bool recordSubmit) {
+  // queue and device must be set before calling this
   assert(hQueue);
   assert(hDevice);
-  profilingData.initTimestampRecording(hDevice);
+
+  profilingData.initTimestampRecording(hDevice, recordSubmit);
+}
+
+uint64_t event_profiling_data_t::getEventSubmitTimestamp() {
+  // Fall back to the completion timestamp when no submission time was recorded.
+  return recordedSubmitTimestamp ? recordedSubmitTimestamp
+                                 : getEventEndTimestamp();
 }
 
 bool event_profiling_data_t::recordingStarted() const {
@@ -139,6 +159,10 @@ void ur_event_handle_t_::onWaitListUse() {
   }
 }
 
+uint64_t ur_event_handle_t_::getEventSubmitTimestamp() {
+  return profilingData.getEventSubmitTimestamp();
+}
+
 uint64_t ur_event_handle_t_::getEventEndTimestamp() {
   return profilingData.getEventEndTimestamp();
 }
@@ -324,12 +348,13 @@ ur_result_t urEventGetProfilingInfo(
 
   UrReturnHelper returnValue(propValueSize, pPropValue, pPropValueSizeRet);
 
-  // For timestamp-recording events the GPU wrote a single global timestamp.
-  // All profiling fields return that value: submit == start == end.
+  // For timestamped events we have the timestamps ready directly on the event
+  // handle, so we short-circuit the return.
   if (isTimestampedEvent) {
     switch (propName) {
     case UR_PROFILING_INFO_COMMAND_QUEUED:
     case UR_PROFILING_INFO_COMMAND_SUBMIT:
+      return returnValue(event->getEventSubmitTimestamp());
     case UR_PROFILING_INFO_COMMAND_START:
     case UR_PROFILING_INFO_COMMAND_END:
     case UR_PROFILING_INFO_COMMAND_COMPLETE:
diff --git a/unified-runtime/source/adapters/level_zero/v2/event.hpp b/unified-runtime/source/adapters/level_zero/v2/event.hpp
index 0c30ca81ed43a..6f3cacc304663 100644
--- a/unified-runtime/source/adapters/level_zero/v2/event.hpp
+++ b/unified-runtime/source/adapters/level_zero/v2/event.hpp
@@ -28,10 +28,13 @@ class event_pool;
 struct event_profiling_data_t {
   event_profiling_data_t(ze_event_handle_t hZeEvent) : hZeEvent(hZeEvent) {}
 
-  // Prepare the event to hold a single GPU-written global timestamp: cache the
-  // device timer resolution/mask and mark the event as timestamp-recording.
-  void initTimestampRecording(ur_device_handle_t hDevice);
+  // Cache the device timer resolution/mask and mark the event as
+  // timestamp-recording. When recordSubmit is set, also capture a submission
+  // timestamp for command_submit; otherwise command_submit reports the
+  // GPU-written timestamp like command_start/command_end.
+  void initTimestampRecording(ur_device_handle_t hDevice, bool recordSubmit);
 
+  uint64_t getEventSubmitTimestamp();
   uint64_t getEventEndTimestamp();
   uint64_t *eventEndTimestampAddr();
 
@@ -44,6 +47,8 @@ struct event_profiling_data_t {
 private:
   ze_event_handle_t hZeEvent;
 
+  // Submission timestamp captured at enqueue; 0 if none was recorded.
+  uint64_t recordedSubmitTimestamp = 0;
   uint64_t recordEventEndTimestamp = 0;
   uint64_t adjustedEventEndTimestamp = 0;
 
@@ -127,14 +132,16 @@ struct ur_event_handle_t_ : v2::ur_object_t {
   // Get the device associated with this event
   ur_device_handle_t getDevice() const;
 
-  // Mark this event as recording a single GPU-written global timestamp,
-  // obtainable via urEventGetProfilingInfo. setQueue must be called first.
-  void initTimestampRecording();
+  // Mark this event as recording a GPU-written global timestamp, obtainable via
+  // urEventGetProfilingInfo. setQueue must be called first. recordSubmit also
+  // captures a separate submission timestamp for command_submit.
+  void initTimestampRecording(bool recordSubmit);
 
   // Get pointer to the timestamp storage, and ze event handle.
   // Caller is responsible for signaling the event once the timestamp is ready.
   std::pair<uint64_t *, ze_event_handle_t> getEventEndTimestampAndHandle();
 
+  uint64_t getEventSubmitTimestamp();
   uint64_t getEventEndTimestamp();
 
   ur::RefCount RefCount;
diff --git a/unified-runtime/source/adapters/level_zero/v2/queue_batched.cpp b/unified-runtime/source/adapters/level_zero/v2/queue_batched.cpp
index 10670dfc8f008..e2a200c02af0b 100644
--- a/unified-runtime/source/adapters/level_zero/v2/queue_batched.cpp
+++ b/unified-runtime/source/adapters/level_zero/v2/queue_batched.cpp
@@ -863,11 +863,10 @@ ur_result_t ur_queue_batched_t::bindlessImagesSignalExternalSemaphoreExp(
       getEvent(lockedBatch, phEvent));
 }
 
-// Queues with batched submissions use regular command lists (similarly to
-// command buffers). The tag event records a single GPU-written global
-// timestamp via zeCommandListAppendWriteGlobalTimestamp, which becomes
-// available once the batch's regular command list is submitted. The difference
-// between two tag timestamps reflects the actual time of execution.
+// Batched submission defers the tag to the device until the batch is flushed,
+// so command_submit meaningfully precedes execution: record a submission
+// timestamp (recordSubmit=true). command_start/command_end still come from the
+// GPU-written timestamp.
 //
 
 ur_result_t ur_queue_batched_t::enqueueTimestampRecordingExp(
@@ -881,7 +880,8 @@ ur_result_t ur_queue_batched_t::enqueueTimestampRecordingExp(
   markIssuedCommandInBatch(lockedBatch);
 
   UR_CALL(lockedBatch->getListManager().appendTimestampRecordingExp(
-      false, waitListView, getEvent(lockedBatch, phEvent)));
+      false, waitListView, getEvent(lockedBatch, phEvent),
+      /*recordSubmit=*/true));
 
   if (blocking) {
     UR_CALL(queueFinishUnlocked(lockedBatch));
diff --git a/unified-runtime/source/adapters/level_zero/v2/queue_immediate_in_order.hpp b/unified-runtime/source/adapters/level_zero/v2/queue_immediate_in_order.hpp
index 0b1bdeeb9d693..10a4b1d7a122a 100644
--- a/unified-runtime/source/adapters/level_zero/v2/queue_immediate_in_order.hpp
+++ b/unified-runtime/source/adapters/level_zero/v2/queue_immediate_in_order.hpp
@@ -484,7 +484,8 @@ struct ur_queue_immediate_in_order_t : ur_object_t, ur_queue_t_ {
 
     return commandListManager.lock()->appendTimestampRecordingExp(
         blocking, waitListView,
-        createEventIfRequested(eventPool.get(), phEvent, this));
+        createEventIfRequested(eventPool.get(), phEvent, this),
+        /*recordSubmit=*/false);
   }
 
   ur_result_t
diff --git a/unified-runtime/source/adapters/level_zero/v2/queue_immediate_out_of_order.hpp b/unified-runtime/source/adapters/level_zero/v2/queue_immediate_out_of_order.hpp
index 525a3b6802523..0ebadbd8722e3 100644
--- a/unified-runtime/source/adapters/level_zero/v2/queue_immediate_out_of_order.hpp
+++ b/unified-runtime/source/adapters/level_zero/v2/queue_immediate_out_of_order.hpp
@@ -544,7 +544,8 @@ struct ur_queue_immediate_out_of_order_t : ur_object_t, ur_queue_t_ {
     return commandListManagers.lock()[commandListId]
         .appendTimestampRecordingExp(
             blocking, waitListView,
-            createEventIfRequested(eventPool.get(), phEvent, this));
+            createEventIfRequested(eventPool.get(), phEvent, this),
+            /*recordSubmit=*/false);
   }
 
   ur_result_t

From 9690996297a91f2d79fd4a2b2318e4c387551eb6 Mon Sep 17 00:00:00 2001
From: Artur Gainullin <artur.gainullin@intel.com>
Date: Fri, 31 Jul 2026 10:23:51 -0700
Subject: [PATCH 4/5] [UR][L0v1] Skip submission timestamp query for
 immediate-queue tags

Apply the same optimization to the L0 v1 adapter: in
urEnqueueTimestampRecordingExp, query urDeviceGetGlobalTimestamps for
command_submit only for batched (non-immediate) queues, where submission is
deferred. Immediate queues report the GPU-written timestamp for command_submit
as well. The recorded submission timestamp, when present, is the reference for
end-timestamp wrap-around correction.

Assisted-By: Claude
---
 .../source/adapters/level_zero/event.cpp      | 41 +++++++++++--------
 .../source/adapters/level_zero/event.hpp      | 15 ++++---
 .../source/adapters/level_zero/queue.cpp      |  6 ++-
 3 files changed, 38 insertions(+), 24 deletions(-)

diff --git a/unified-runtime/source/adapters/level_zero/event.cpp b/unified-runtime/source/adapters/level_zero/event.cpp
index 51bb588bbb9ec..065e9a19de046 100644
--- a/unified-runtime/source/adapters/level_zero/event.cpp
+++ b/unified-runtime/source/adapters/level_zero/event.cpp
@@ -563,13 +563,17 @@ ur_result_t urEventGetProfilingInfo(
   // handle, so we short-circuit the return.
   // We don't support user events with timestamps due to requiring the UrQueue.
   if (isTimestampedEvent && Event->UrQueue) {
-    uint64_t ContextStartTime = Event->RecordEventStartTimestamp;
     switch (PropName) {
     case UR_PROFILING_INFO_COMMAND_QUEUED:
     case UR_PROFILING_INFO_COMMAND_SUBMIT:
-      return ReturnValue(ContextStartTime);
+      // When no submission timestamp was recorded, fall through to report the
+      // completion time for command_submit as well.
+      if (Event->RecordEventSubmitTimestamp)
+        return ReturnValue(Event->RecordEventSubmitTimestamp);
+      [[fallthrough]];
+    case UR_PROFILING_INFO_COMMAND_START:
     case UR_PROFILING_INFO_COMMAND_END:
-    case UR_PROFILING_INFO_COMMAND_START: {
+    case UR_PROFILING_INFO_COMMAND_COMPLETE: {
       // If RecordEventEndTimestamp on the event is non-zero it means it has
       // collected the result of the queue already. In that case it has been
       // adjusted and is ready for immediate return.
@@ -596,8 +600,10 @@ ur_result_t urEventGetProfilingInfo(
       // Handle a possible wrap-around (the underlying HW counter is < 64-bit).
       // Note, it will not report correct time if there were multiple wrap
       // arounds, and the longer term plan is to enlarge the capacity of the
-      // HW timestamps.
-      if (ContextEndTime < ContextStartTime)
+      // HW timestamps. Only batched queues record a submission timestamp to
+      // compare against.
+      if (Event->RecordEventSubmitTimestamp &&
+          ContextEndTime < Event->RecordEventSubmitTimestamp)
         ContextEndTime += TimestampMaxValue * ZeTimerResolution;
 
       // Now that we have the result, there is no need to keep it in the queue
@@ -608,10 +614,6 @@ ur_result_t urEventGetProfilingInfo(
 
       return ReturnValue(ContextEndTime);
     }
-    case UR_PROFILING_INFO_COMMAND_COMPLETE:
-      UR_LOG(ERR, "urEventGetProfilingInfo: "
-                  "UR_PROFILING_INFO_COMMAND_COMPLETE not supported");
-      return UR_RESULT_ERROR_UNSUPPORTED_ENUMERATION;
     default:
       UR_LOG(ERR, "urEventGetProfilingInfo: not supported ParamName");
       return UR_RESULT_ERROR_INVALID_VALUE;
@@ -777,17 +779,24 @@ ur_result_t urEnqueueTimestampRecordingExp(
   ze_event_handle_t ZeEvent = OutEventInternal->ZeEvent;
   OutEventInternal->WaitList = TmpWaitList;
 
-  // Reset the end timestamp, in case it has been previously used.
+  // Reset the timestamps, in case they have been previously used.
+  OutEventInternal->RecordEventSubmitTimestamp = 0;
   OutEventInternal->RecordEventEndTimestamp = 0;
 
-  uint64_t DeviceStartTimestamp = 0;
-  UR_CALL(ur::level_zero::urDeviceGetGlobalTimestamps(
-      common_cast(Device), &DeviceStartTimestamp, nullptr));
-  OutEventInternal->RecordEventStartTimestamp = DeviceStartTimestamp;
-
   // Mark this event as timestamped
   OutEventInternal->IsTimestamped = true;
 
+  // Batched submission defers the tag to the device until the batch is flushed,
+  // so command_submit meaningfully precedes execution: record a submission
+  // timestamp. Immediate queues skip this query (its cost dominates the tag
+  // latency) and report the GPU-written timestamp for command_submit too.
+  if (!Queue->UsingImmCmdLists) {
+    uint64_t DeviceSubmitTimestamp = 0;
+    UR_CALL(ur::level_zero::urDeviceGetGlobalTimestamps(
+        common_cast(Device), &DeviceSubmitTimestamp, nullptr));
+    OutEventInternal->RecordEventSubmitTimestamp = DeviceSubmitTimestamp;
+  }
+
   // Create a new entry in the queue's recordings.
   Queue->EndTimeRecordings[*OutEventInternalPtr] = 0;
 
@@ -1503,7 +1512,7 @@ ur_result_t ur::level_zero::v1::ur_event_handle_t_::reset() {
   completionBatch = std::nullopt;
   OriginAllocEvent = nullptr;
   IsTimestamped = false;
-  RecordEventStartTimestamp = 0;
+  RecordEventSubmitTimestamp = 0;
   RecordEventEndTimestamp = 0;
 
   if (!isHostVisible())
diff --git a/unified-runtime/source/adapters/level_zero/event.hpp b/unified-runtime/source/adapters/level_zero/event.hpp
index 9a977b8efa5e3..da31357fd04b0 100644
--- a/unified-runtime/source/adapters/level_zero/event.hpp
+++ b/unified-runtime/source/adapters/level_zero/event.hpp
@@ -216,12 +216,15 @@ struct ur_event_handle_t_ : ur_object_t {
   // Indicates within creation of proxy event.
   bool IsCreatingHostProxyEvent = {false};
 
-  // Indicates the recorded start and end timestamps for the event. These are
-  // only set for events returned by timestamp recording enqueue functions.
-  // A non-zero value for RecordEventStartTimestamp indicates the event was the
-  // result of a timestamp recording. If RecordEventEndTimestamp is non-zero, it
-  // means the event has fetched the end-timestamp from the queue.
-  uint64_t RecordEventStartTimestamp = 0;
+  // Submission timestamp for a tag event, captured at enqueue via
+  // urDeviceGetGlobalTimestamps. Recorded only for batched queues, where
+  // submission is deferred; 0 otherwise (command_submit then equals the
+  // GPU-written timestamp).
+  uint64_t RecordEventSubmitTimestamp = 0;
+
+  // The GPU-written global timestamp for a timestamp-recording event. Set to a
+  // non-zero adjusted value once the end-timestamp has been fetched from the
+  // queue; IsTimestamped tells whether the event is timestamp-recording.
   uint64_t RecordEventEndTimestamp = 0;
 
   // Besides each PI object keeping a total reference count in
diff --git a/unified-runtime/source/adapters/level_zero/queue.cpp b/unified-runtime/source/adapters/level_zero/queue.cpp
index 1844140ef9c1c..91fbe0136ab1b 100644
--- a/unified-runtime/source/adapters/level_zero/queue.cpp
+++ b/unified-runtime/source/adapters/level_zero/queue.cpp
@@ -1667,8 +1667,10 @@ void ur::level_zero::v1::ur_queue_handle_t_::clearEndTimeRecordings() {
     // Handle a possible wrap-around (the underlying HW counter is < 64-bit).
     // Note, it will not report correct time if there were multiple wrap
     // arounds, and the longer term plan is to enlarge the capacity of the
-    // HW timestamps.
-    if (ContextEndTime < Event->RecordEventStartTimestamp)
+    // HW timestamps. Only batched queues record a submission timestamp to
+    // compare against.
+    if (Event->RecordEventSubmitTimestamp &&
+        ContextEndTime < Event->RecordEventSubmitTimestamp)
       ContextEndTime += TimestampMaxValue * ZeTimerResolution;
 
     // Store it in the event.

From 3114b5825b87d245083a7a8399896ff8327c8dab Mon Sep 17 00:00:00 2001
From: Artur Gainullin <artur.gainullin@intel.com>
Date: Fri, 31 Jul 2026 13:58:22 -0700
Subject: [PATCH 5/5] [UR][L0][SYCL] Always report equal timestamps for
 profiling tag events

Following the spec agreement, a tag operation (submit_profiling_tag /
enqueue_signal_event) is an empty command, so command_submit, command_start and
command_end all report the single GPU-written completion timestamp.
This removes the batched-queue special case added earlier (the extra
urDeviceGetGlobalTimestamps query for command_submit and the associated
wrap-around handling) from both the L0 v1 and v2 adapters.

Assisted-By: Claude
---
 .../source/adapters/level_zero/event.cpp      | 36 +++---------------
 .../source/adapters/level_zero/event.hpp      |  6 ---
 .../source/adapters/level_zero/queue.cpp      | 18 ++-------
 .../level_zero/v2/command_list_manager.cpp    |  5 +--
 .../level_zero/v2/command_list_manager.hpp    |  3 +-
 .../source/adapters/level_zero/v2/event.cpp   | 37 +++++--------------
 .../source/adapters/level_zero/v2/event.hpp   | 15 ++------
 .../adapters/level_zero/v2/queue_batched.cpp  |  9 +----
 .../v2/queue_immediate_in_order.hpp           |  3 +-
 .../v2/queue_immediate_out_of_order.hpp       |  3 +-
 10 files changed, 28 insertions(+), 107 deletions(-)

diff --git a/unified-runtime/source/adapters/level_zero/event.cpp b/unified-runtime/source/adapters/level_zero/event.cpp
index 065e9a19de046..b264618564fc9 100644
--- a/unified-runtime/source/adapters/level_zero/event.cpp
+++ b/unified-runtime/source/adapters/level_zero/event.cpp
@@ -564,13 +564,10 @@ ur_result_t urEventGetProfilingInfo(
   // We don't support user events with timestamps due to requiring the UrQueue.
   if (isTimestampedEvent && Event->UrQueue) {
     switch (PropName) {
+    // The tag is an empty command, so all timestamps are the single
+    // GPU-written completion time.
     case UR_PROFILING_INFO_COMMAND_QUEUED:
     case UR_PROFILING_INFO_COMMAND_SUBMIT:
-      // When no submission timestamp was recorded, fall through to report the
-      // completion time for command_submit as well.
-      if (Event->RecordEventSubmitTimestamp)
-        return ReturnValue(Event->RecordEventSubmitTimestamp);
-      [[fallthrough]];
     case UR_PROFILING_INFO_COMMAND_START:
     case UR_PROFILING_INFO_COMMAND_END:
     case UR_PROFILING_INFO_COMMAND_COMPLETE: {
@@ -588,7 +585,8 @@ ur_result_t urEventGetProfilingInfo(
         return UR_RESULT_ERROR_UNKNOWN;
       auto &EndTimeRecording = Entry->second;
 
-      // End time needs to be adjusted for resolution and valid bits.
+      // End time needs to be adjusted for resolution and valid bits. A single
+      // timestamp has no separate start value to detect wrap-around against.
       uint64_t ContextEndTime =
           (EndTimeRecording & TimestampMaxValue) * ZeTimerResolution;
 
@@ -597,15 +595,6 @@ ur_result_t urEventGetProfilingInfo(
       if (ContextEndTime == 0)
         return ReturnValue(ContextEndTime);
 
-      // Handle a possible wrap-around (the underlying HW counter is < 64-bit).
-      // Note, it will not report correct time if there were multiple wrap
-      // arounds, and the longer term plan is to enlarge the capacity of the
-      // HW timestamps. Only batched queues record a submission timestamp to
-      // compare against.
-      if (Event->RecordEventSubmitTimestamp &&
-          ContextEndTime < Event->RecordEventSubmitTimestamp)
-        ContextEndTime += TimestampMaxValue * ZeTimerResolution;
-
       // Now that we have the result, there is no need to keep it in the queue
       // anymore, so we cache it on the event and evict the record from the
       // queue.
@@ -757,8 +746,6 @@ ur_result_t urEnqueueTimestampRecordingExp(
   // Lock automatically releases when this goes out of scope.
   std::scoped_lock<ur_shared_mutex> lock(Queue->Mutex);
 
-  ur_device_handle_t Device = Queue->Device;
-
   bool UseCopyEngine = false;
   ur_ze_event_list_t TmpWaitList;
   UR_CALL(TmpWaitList.createAndRetainUrZeEventList(
@@ -779,24 +766,12 @@ ur_result_t urEnqueueTimestampRecordingExp(
   ze_event_handle_t ZeEvent = OutEventInternal->ZeEvent;
   OutEventInternal->WaitList = TmpWaitList;
 
-  // Reset the timestamps, in case they have been previously used.
-  OutEventInternal->RecordEventSubmitTimestamp = 0;
+  // Reset the end timestamp, in case it has been previously used.
   OutEventInternal->RecordEventEndTimestamp = 0;
 
   // Mark this event as timestamped
   OutEventInternal->IsTimestamped = true;
 
-  // Batched submission defers the tag to the device until the batch is flushed,
-  // so command_submit meaningfully precedes execution: record a submission
-  // timestamp. Immediate queues skip this query (its cost dominates the tag
-  // latency) and report the GPU-written timestamp for command_submit too.
-  if (!Queue->UsingImmCmdLists) {
-    uint64_t DeviceSubmitTimestamp = 0;
-    UR_CALL(ur::level_zero::urDeviceGetGlobalTimestamps(
-        common_cast(Device), &DeviceSubmitTimestamp, nullptr));
-    OutEventInternal->RecordEventSubmitTimestamp = DeviceSubmitTimestamp;
-  }
-
   // Create a new entry in the queue's recordings.
   Queue->EndTimeRecordings[*OutEventInternalPtr] = 0;
 
@@ -1512,7 +1487,6 @@ ur_result_t ur::level_zero::v1::ur_event_handle_t_::reset() {
   completionBatch = std::nullopt;
   OriginAllocEvent = nullptr;
   IsTimestamped = false;
-  RecordEventSubmitTimestamp = 0;
   RecordEventEndTimestamp = 0;
 
   if (!isHostVisible())
diff --git a/unified-runtime/source/adapters/level_zero/event.hpp b/unified-runtime/source/adapters/level_zero/event.hpp
index da31357fd04b0..c971396f9df26 100644
--- a/unified-runtime/source/adapters/level_zero/event.hpp
+++ b/unified-runtime/source/adapters/level_zero/event.hpp
@@ -216,12 +216,6 @@ struct ur_event_handle_t_ : ur_object_t {
   // Indicates within creation of proxy event.
   bool IsCreatingHostProxyEvent = {false};
 
-  // Submission timestamp for a tag event, captured at enqueue via
-  // urDeviceGetGlobalTimestamps. Recorded only for batched queues, where
-  // submission is deferred; 0 otherwise (command_submit then equals the
-  // GPU-written timestamp).
-  uint64_t RecordEventSubmitTimestamp = 0;
-
   // The GPU-written global timestamp for a timestamp-recording event. Set to a
   // non-zero adjusted value once the end-timestamp has been fetched from the
   // queue; IsTimestamped tells whether the event is timestamp-recording.
diff --git a/unified-runtime/source/adapters/level_zero/queue.cpp b/unified-runtime/source/adapters/level_zero/queue.cpp
index 91fbe0136ab1b..45342df193a34 100644
--- a/unified-runtime/source/adapters/level_zero/queue.cpp
+++ b/unified-runtime/source/adapters/level_zero/queue.cpp
@@ -1660,21 +1660,11 @@ void ur::level_zero::v1::ur_queue_handle_t_::clearEndTimeRecordings() {
     auto &Event = Entry.first;
     auto &EndTimeRecording = Entry.second;
 
-    // Write the result back to the event if it is not dead.
-    uint64_t ContextEndTime =
+    // Write the result back to the event if it is not dead. A single
+    // GPU-written timestamp has no separate start value to detect wrap-around
+    // against.
+    Event->RecordEventEndTimestamp =
         (EndTimeRecording & TimestampMaxValue) * ZeTimerResolution;
-
-    // Handle a possible wrap-around (the underlying HW counter is < 64-bit).
-    // Note, it will not report correct time if there were multiple wrap
-    // arounds, and the longer term plan is to enlarge the capacity of the
-    // HW timestamps. Only batched queues record a submission timestamp to
-    // compare against.
-    if (Event->RecordEventSubmitTimestamp &&
-        ContextEndTime < Event->RecordEventSubmitTimestamp)
-      ContextEndTime += TimestampMaxValue * ZeTimerResolution;
-
-    // Store it in the event.
-    Event->RecordEventEndTimestamp = ContextEndTime;
   }
   EndTimeRecordings.clear();
   EvictedEndTimeRecordings.clear();
diff --git a/unified-runtime/source/adapters/level_zero/v2/command_list_manager.cpp b/unified-runtime/source/adapters/level_zero/v2/command_list_manager.cpp
index 6a239a145a0a4..16ef6f96d8bbc 100644
--- a/unified-runtime/source/adapters/level_zero/v2/command_list_manager.cpp
+++ b/unified-runtime/source/adapters/level_zero/v2/command_list_manager.cpp
@@ -593,8 +593,7 @@ ur_result_t ur_command_list_manager::appendUSMMemcpy2D(
 }
 
 ur_result_t ur_command_list_manager::appendTimestampRecordingExp(
-    bool blocking, wait_list_view &waitListView, ur_event_handle_t phEvent,
-    bool recordSubmit) {
+    bool blocking, wait_list_view &waitListView, ur_event_handle_t phEvent) {
   TRACK_SCOPE_LATENCY("ur_command_list_manager::appendTimestampRecordingExp");
 
   if (!phEvent) {
@@ -603,7 +602,7 @@ ur_result_t ur_command_list_manager::appendTimestampRecordingExp(
 
   auto [pWaitEvents, numWaitEvents, _] = waitListView;
 
-  phEvent->initTimestampRecording(recordSubmit);
+  phEvent->initTimestampRecording();
 
   auto [timestampPtr, zeSignalEvent] = phEvent->getEventEndTimestampAndHandle();
 
diff --git a/unified-runtime/source/adapters/level_zero/v2/command_list_manager.hpp b/unified-runtime/source/adapters/level_zero/v2/command_list_manager.hpp
index f0fae46ca2ee1..8ba0bd797256e 100644
--- a/unified-runtime/source/adapters/level_zero/v2/command_list_manager.hpp
+++ b/unified-runtime/source/adapters/level_zero/v2/command_list_manager.hpp
@@ -212,8 +212,7 @@ struct ur_command_list_manager {
       const ur_event_handle_t *phEventWaitList, ur_event_handle_t phEvent);
   ur_result_t appendTimestampRecordingExp(bool blocking,
                                           wait_list_view &waitListView,
-                                          ur_event_handle_t phEvent,
-                                          bool recordSubmit);
+                                          ur_event_handle_t phEvent);
   ur_result_t
   appendCommandBufferExp(ur_exp_command_buffer_handle_t hCommandBuffer,
                          wait_list_view &waitListView,
diff --git a/unified-runtime/source/adapters/level_zero/v2/event.cpp b/unified-runtime/source/adapters/level_zero/v2/event.cpp
index 11d53ce393314..babef838be86d 100644
--- a/unified-runtime/source/adapters/level_zero/v2/event.cpp
+++ b/unified-runtime/source/adapters/level_zero/v2/event.cpp
@@ -56,11 +56,10 @@ uint64_t event_profiling_data_t::getEventEndTimestamp() {
   assert(zeTimerResolution);
   assert(timestampMaxValue);
 
-  // recordedSubmitTimestamp is the wrap-around reference; it is 0 when no
-  // submission timestamp was recorded, in which case no adjustment is applied.
+  // A timestamp-recording event holds a single GPU-written global timestamp,
+  // so there is no separate start value to detect a wrap-around against.
   adjustedEventEndTimestamp =
-      adjustEndEventTimestamp(recordedSubmitTimestamp, recordEventEndTimestamp,
-                              timestampMaxValue, zeTimerResolution);
+      (recordEventEndTimestamp & timestampMaxValue) * zeTimerResolution;
 
   return adjustedEventEndTimestamp;
 }
@@ -77,37 +76,22 @@ void event_profiling_data_t::reset() {
   // context) and the timstamp is already wrriten, so there's no race-condition
   // possible.
   adjustedEventEndTimestamp = 0;
-  recordedSubmitTimestamp = 0;
   timestampRecorded = false;
 }
 
-void event_profiling_data_t::initTimestampRecording(ur_device_handle_t hDevice,
-                                                    bool recordSubmit) {
+void event_profiling_data_t::initTimestampRecording(
+    ur_device_handle_t hDevice) {
   zeTimerResolution = hDevice->getTimerResolution();
   timestampMaxValue = hDevice->getTimestampMask();
-
-  // Recording a submission timestamp requires an extra device query whose cost
-  // dominates the tag latency, so only do it when requested (batched queues).
-  if (recordSubmit) {
-    UR_CALL_THROWS(ur::level_zero::urDeviceGetGlobalTimestamps(
-        common_cast(hDevice), &recordedSubmitTimestamp, nullptr));
-  }
-
   timestampRecorded = true;
 }
 
-void ur_event_handle_t_::initTimestampRecording(bool recordSubmit) {
+void ur_event_handle_t_::initTimestampRecording() {
   // queue and device must be set before calling this
   assert(hQueue);
   assert(hDevice);
 
-  profilingData.initTimestampRecording(hDevice, recordSubmit);
-}
-
-uint64_t event_profiling_data_t::getEventSubmitTimestamp() {
-  // Fall back to the completion timestamp when no submission time was recorded.
-  return recordedSubmitTimestamp ? recordedSubmitTimestamp
-                                 : getEventEndTimestamp();
+  profilingData.initTimestampRecording(hDevice);
 }
 
 bool event_profiling_data_t::recordingStarted() const {
@@ -159,10 +143,6 @@ void ur_event_handle_t_::onWaitListUse() {
   }
 }
 
-uint64_t ur_event_handle_t_::getEventSubmitTimestamp() {
-  return profilingData.getEventSubmitTimestamp();
-}
-
 uint64_t ur_event_handle_t_::getEventEndTimestamp() {
   return profilingData.getEventEndTimestamp();
 }
@@ -354,10 +334,11 @@ ur_result_t urEventGetProfilingInfo(
     switch (propName) {
     case UR_PROFILING_INFO_COMMAND_QUEUED:
     case UR_PROFILING_INFO_COMMAND_SUBMIT:
-      return returnValue(event->getEventSubmitTimestamp());
     case UR_PROFILING_INFO_COMMAND_START:
     case UR_PROFILING_INFO_COMMAND_END:
     case UR_PROFILING_INFO_COMMAND_COMPLETE:
+      // The tag is an empty command, so all timestamps are the single
+      // GPU-written completion time.
       return returnValue(event->getEventEndTimestamp());
     default:
       UR_LOG(ERR, "urEventGetProfilingInfo: not supported ParamName");
diff --git a/unified-runtime/source/adapters/level_zero/v2/event.hpp b/unified-runtime/source/adapters/level_zero/v2/event.hpp
index 6f3cacc304663..208371062fda7 100644
--- a/unified-runtime/source/adapters/level_zero/v2/event.hpp
+++ b/unified-runtime/source/adapters/level_zero/v2/event.hpp
@@ -29,12 +29,9 @@ struct event_profiling_data_t {
   event_profiling_data_t(ze_event_handle_t hZeEvent) : hZeEvent(hZeEvent) {}
 
   // Cache the device timer resolution/mask and mark the event as
-  // timestamp-recording. When recordSubmit is set, also capture a submission
-  // timestamp for command_submit; otherwise command_submit reports the
-  // GPU-written timestamp like command_start/command_end.
-  void initTimestampRecording(ur_device_handle_t hDevice, bool recordSubmit);
+  // timestamp-recording.
+  void initTimestampRecording(ur_device_handle_t hDevice);
 
-  uint64_t getEventSubmitTimestamp();
   uint64_t getEventEndTimestamp();
   uint64_t *eventEndTimestampAddr();
 
@@ -47,8 +44,6 @@ struct event_profiling_data_t {
 private:
   ze_event_handle_t hZeEvent;
 
-  // Submission timestamp captured at enqueue; 0 if none was recorded.
-  uint64_t recordedSubmitTimestamp = 0;
   uint64_t recordEventEndTimestamp = 0;
   uint64_t adjustedEventEndTimestamp = 0;
 
@@ -133,15 +128,13 @@ struct ur_event_handle_t_ : v2::ur_object_t {
   ur_device_handle_t getDevice() const;
 
   // Mark this event as recording a GPU-written global timestamp, obtainable via
-  // urEventGetProfilingInfo. setQueue must be called first. recordSubmit also
-  // captures a separate submission timestamp for command_submit.
-  void initTimestampRecording(bool recordSubmit);
+  // urEventGetProfilingInfo. setQueue must be called first.
+  void initTimestampRecording();
 
   // Get pointer to the timestamp storage, and ze event handle.
   // Caller is responsible for signaling the event once the timestamp is ready.
   std::pair<uint64_t *, ze_event_handle_t> getEventEndTimestampAndHandle();
 
-  uint64_t getEventSubmitTimestamp();
   uint64_t getEventEndTimestamp();
 
   ur::RefCount RefCount;
diff --git a/unified-runtime/source/adapters/level_zero/v2/queue_batched.cpp b/unified-runtime/source/adapters/level_zero/v2/queue_batched.cpp
index e2a200c02af0b..c39a2a4ce8d5c 100644
--- a/unified-runtime/source/adapters/level_zero/v2/queue_batched.cpp
+++ b/unified-runtime/source/adapters/level_zero/v2/queue_batched.cpp
@@ -863,12 +863,6 @@ ur_result_t ur_queue_batched_t::bindlessImagesSignalExternalSemaphoreExp(
       getEvent(lockedBatch, phEvent));
 }
 
-// Batched submission defers the tag to the device until the batch is flushed,
-// so command_submit meaningfully precedes execution: record a submission
-// timestamp (recordSubmit=true). command_start/command_end still come from the
-// GPU-written timestamp.
-//
-
 ur_result_t ur_queue_batched_t::enqueueTimestampRecordingExp(
     bool blocking, uint32_t numEventsInWaitList,
     const ur_event_handle_t *phEventWaitList, ur_event_handle_t *phEvent) {
@@ -880,8 +874,7 @@ ur_result_t ur_queue_batched_t::enqueueTimestampRecordingExp(
   markIssuedCommandInBatch(lockedBatch);
 
   UR_CALL(lockedBatch->getListManager().appendTimestampRecordingExp(
-      false, waitListView, getEvent(lockedBatch, phEvent),
-      /*recordSubmit=*/true));
+      false, waitListView, getEvent(lockedBatch, phEvent)));
 
   if (blocking) {
     UR_CALL(queueFinishUnlocked(lockedBatch));
diff --git a/unified-runtime/source/adapters/level_zero/v2/queue_immediate_in_order.hpp b/unified-runtime/source/adapters/level_zero/v2/queue_immediate_in_order.hpp
index 10a4b1d7a122a..0b1bdeeb9d693 100644
--- a/unified-runtime/source/adapters/level_zero/v2/queue_immediate_in_order.hpp
+++ b/unified-runtime/source/adapters/level_zero/v2/queue_immediate_in_order.hpp
@@ -484,8 +484,7 @@ struct ur_queue_immediate_in_order_t : ur_object_t, ur_queue_t_ {
 
     return commandListManager.lock()->appendTimestampRecordingExp(
         blocking, waitListView,
-        createEventIfRequested(eventPool.get(), phEvent, this),
-        /*recordSubmit=*/false);
+        createEventIfRequested(eventPool.get(), phEvent, this));
   }
 
   ur_result_t
diff --git a/unified-runtime/source/adapters/level_zero/v2/queue_immediate_out_of_order.hpp b/unified-runtime/source/adapters/level_zero/v2/queue_immediate_out_of_order.hpp
index 0ebadbd8722e3..525a3b6802523 100644
--- a/unified-runtime/source/adapters/level_zero/v2/queue_immediate_out_of_order.hpp
+++ b/unified-runtime/source/adapters/level_zero/v2/queue_immediate_out_of_order.hpp
@@ -544,8 +544,7 @@ struct ur_queue_immediate_out_of_order_t : ur_object_t, ur_queue_t_ {
     return commandListManagers.lock()[commandListId]
         .appendTimestampRecordingExp(
             blocking, waitListView,
-            createEventIfRequested(eventPool.get(), phEvent, this),
-            /*recordSubmit=*/false);
+            createEventIfRequested(eventPool.get(), phEvent, this));
   }
 
   ur_result_t