[SYCL] Fix computation of the submit time based on host timestamps (#12104)

getCurrentDeviceTime uses the following formula to calculate device
time:
[base device timestamp] + ([current host timestamp] - [base host
timestamp]).
Host time stamps are queried using std::chrono. And base host timestamp
is updated in two cases:
1. The first call to getCurrentDeviceTime.
2. When refresh is needed.

Problem is that currently we remember base host timestamp at the wrong
moment: there is a large gap between the point when we query and
remember base host timestamp and the point when we get device timestamp
from plugin, so ([current host timestamp] - [base host timestamp])
includes execution time of things like getPlugin() which may be
significant, especially on the first call when plugins initialization
happen.

As a result we add incorrect difference to the base device time and
calculated submission time is incorrect, it is sometimes greater than
command_start time.

This patch fixes that problem by querying base host time properly after
piGetDeviceAndHostTimer call.

Author: againull

sycl/source/detail/device_impl.cpp

@@ -599,8 +599,8 @@ ext::oneapi::experimental::architecture device_impl::getDeviceArch() const {
   return MDeviceArch;
 }
 
-// On first call this function queries for device timestamp
-// along with host synchronized timestamp and stores it in memeber varaible
+// On the first call this function queries for device timestamp
+// along with host synchronized timestamp and stores it in member variable
 // MDeviceHostBaseTime. Subsequent calls to this function would just retrieve
 // the host timestamp, compute difference against the host timestamp in
 // MDeviceHostBaseTime and calculate the device timestamp based on the
@@ -622,14 +622,28 @@ uint64_t device_impl::getCurrentDeviceTime() {
   // To account for potential clock drift between host clock and device clock.
   // The value set is arbitrary: 200 seconds
   constexpr uint64_t TimeTillRefresh = 200e9;
+  assert(HostTime >= MDeviceHostBaseTime.second);
   uint64_t Diff = HostTime - MDeviceHostBaseTime.second;
 
-  if (Diff > TimeTillRefresh || Diff <= 0) {
+  // If getCurrentDeviceTime is called for the first time or we have to refresh.
+  if (!MDeviceHostBaseTime.second || Diff > TimeTillRefresh) {
     const auto &Plugin = getPlugin();
     auto Result =
         Plugin->call_nocheck<detail::PiApiKind::piGetDeviceAndHostTimer>(
             MDevice, &MDeviceHostBaseTime.first, &MDeviceHostBaseTime.second);
-
+    // We have to remember base host timestamp right after PI call and it is
+    // going to be used for calculation of the device timestamp at the next
+    // getCurrentDeviceTime() call. We need to do it here because getPlugin()
+    // and piGetDeviceAndHostTimer calls may take significant amount of time,
+    // for example on the first call to getPlugin plugins may need to be
+    // initialized. If we use timestamp from the beginning of the function then
+    // the difference between host timestamps of the current
+    // getCurrentDeviceTime and the next getCurrentDeviceTime will be incorrect
+    // because it will include execution time of the code before we get device
+    // timestamp from piGetDeviceAndHostTimer.
+    HostTime =
+        duration_cast<nanoseconds>(steady_clock::now().time_since_epoch())
+            .count();
     if (Result == PI_ERROR_INVALID_OPERATION) {
       char *p = nullptr;
       Plugin->call_nocheck<detail::PiApiKind::piPluginGetLastError>(&p);

sycl/source/detail/device_impl.hpp

@@ -274,7 +274,7 @@ class device_impl {
   mutable std::once_flag MDeviceNameFlag;
   mutable ext::oneapi::experimental::architecture MDeviceArch{};
   mutable std::once_flag MDeviceArchFlag;
-  std::pair<uint64_t, uint64_t> MDeviceHostBaseTime;
+  std::pair<uint64_t, uint64_t> MDeviceHostBaseTime{0, 0};
 }; // class device_impl
 
 } // namespace detail

sycl/test-e2e/Basic/submit_time.cpp

@@ -1,34 +1,32 @@
 // RUN: %{build} -o %t.out
-// RUN: %{run} %t.out
+// There is an issue with reported device time for the L0 backend, works only on
+// pvc for now. No such problems for other backends.
+// RUN: %if (!ext_oneapi_level_zero || gpu-intel-pvc) %{ %{run} %t.out %}
 
-// Check that device_impl::isGetDeviceAndHostTimerSupported() is not spoiling
-// device_impl::MDeviceHostBaseTime values used for submit timestamp
-// calculation.
+// Check that submission time is calculated properly.
 
 #include <sycl/sycl.hpp>
 
-using namespace sycl;
-
 int main(void) {
-  sycl::queue queue(
-      sycl::property_list{sycl::property::queue::enable_profiling()});
-  sycl::event event = queue.submit([&](sycl::handler &cgh) {
-    cgh.parallel_for<class set_value>(sycl::range<1>{1024},
-                                      [=](sycl::id<1> idx) {});
-  });
+  sycl::queue q({sycl::property::queue::enable_profiling{}});
+  int *data = malloc_host<int>(1024, q);
 
-  // SYCL RT internally calls device_impl::isGetDeviceAndHostTimerSupported()
-  // to decide how to calculate "submit" timestamp - either using backend API
-  // call or using fallback implementation.
-  auto submit =
-      event.get_profiling_info<sycl::info::event_profiling::command_submit>();
-  auto start =
-      event.get_profiling_info<sycl::info::event_profiling::command_start>();
-  auto end =
-      event.get_profiling_info<sycl::info::event_profiling::command_end>();
+  for (int i = 0; i < 20; i++) {
+    auto event = q.submit([&](sycl::handler &cgh) {
+      cgh.parallel_for<class KernelTime>(
+          sycl::range<1>(1024), [=](sycl::id<1> idx) { data[idx] = idx; });
+    });
 
-  if (!(submit <= start) || !(start <= end))
-    return -1;
+    event.wait();
+    auto submit_time =
+        event.get_profiling_info<sycl::info::event_profiling::command_submit>();
+    auto start_time =
+        event.get_profiling_info<sycl::info::event_profiling::command_start>();
+    auto end_time =
+        event.get_profiling_info<sycl::info::event_profiling::command_end>();
 
+    assert((submit_time <= start_time) && (start_time <= end_time));
+  }
+  sycl::free(data, q);
   return 0;
 }