[SYCL][UR] Fix the accuracy of command submission timestamp (#18735)

Current "submission time" calculation is inaccurate because we don't use
both synchronized timestamps returned by zeDeviceGetGlobalTimestamps but
using only device timestamp from that call and use std::chrono "close"
to that call to record the host time. This estimation becomes inaccurate
pretty quickly. This PR fixes this problem using the known fact that L0
runtime implementation uses CLOCK_MONOTONIC_RAW on Linux and
QueryPerformanceCounter on Windows.
So, with this fix, at the first call, we use both device and host
timestamps from zeDeviceGetGlobalTimestamps, subsequent calls (when
device timestamp is not requested) will return corresponding host
timestamp only, without making the `zeDeviceGetGlobalTimestamps` call
which has high latency.

Even though this approach improves accuracy and submit time doesn't
become "invalid" (submit time > start_time) fast, it still doesn't
guarantee that it will not happen. So, there will be additional fix done
in #18717 to fix that. Test is also
updated there to check larger number of iterations.

Also apply the same fix as
138bef7
for cuda adapter (i.e. record host time after cuEventSynchronize for
more precise measurement)

Author: againull

sycl/source/detail/device_impl.cpp

@@ -324,19 +324,32 @@ ur_native_handle_t device_impl::getNative() const {
 // clock drift between host and device.
 //
 uint64_t device_impl::getCurrentDeviceTime() {
-  using namespace std::chrono;
-  uint64_t HostTime =
-      duration_cast<nanoseconds>(steady_clock::now().time_since_epoch())
-          .count();
+  auto GetGlobalTimestamps = [this](ur_device_handle_t Device,
+                                    uint64_t *DeviceTime, uint64_t *HostTime) {
+    auto Result =
+        getAdapter()->call_nocheck<UrApiKind::urDeviceGetGlobalTimestamps>(
+            Device, DeviceTime, HostTime);
+    if (Result == UR_RESULT_ERROR_INVALID_OPERATION) {
+      // NOTE(UR port): Removed the call to GetLastError because  we shouldn't
+      // be calling it after ERROR_INVALID_OPERATION: there is no
+      // adapter-specific error.
+      throw detail::set_ur_error(
+          sycl::exception(
+              make_error_code(errc::feature_not_supported),
+              "Device and/or backend does not support querying timestamp."),
+          UR_RESULT_ERROR_INVALID_OPERATION);
+    } else {
+      getAdapter()->checkUrResult<errc::feature_not_supported>(Result);
+    }
+  };
 
+  uint64_t HostTime = 0;
+  uint64_t Diff = 0;
   // To account for potential clock drift between host clock and device clock.
   // The value set is arbitrary: 200 seconds
-  std::shared_lock<std::shared_mutex> ReadLock(MDeviceHostBaseTimeMutex);
   constexpr uint64_t TimeTillRefresh = 200e9;
-  assert(HostTime >= MDeviceHostBaseTime.second);
-  uint64_t Diff = HostTime - MDeviceHostBaseTime.second;
-
   // If getCurrentDeviceTime is called for the first time or we have to refresh.
+  std::shared_lock<std::shared_mutex> ReadLock(MDeviceHostBaseTimeMutex);
   if (!MDeviceHostBaseTime.second || Diff > TimeTillRefresh) {
     ReadLock.unlock();
     std::unique_lock<std::shared_mutex> WriteLock(MDeviceHostBaseTimeMutex);
@@ -346,38 +359,12 @@ uint64_t device_impl::getCurrentDeviceTime() {
       // MDeviceHostBaseTime, so we can just return the current device time.
       return MDeviceHostBaseTime.first + Diff;
     }
-    const auto &Adapter = getAdapter();
-    auto Result = Adapter->call_nocheck<UrApiKind::urDeviceGetGlobalTimestamps>(
-        MDevice, &MDeviceHostBaseTime.first, &MDeviceHostBaseTime.second);
-    // We have to remember base host timestamp right after UR 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 getAdapter()
-    // and urDeviceGetGlobalTimestamps calls may take significant amount of
-    // time, for example on the first call to getAdapter adapters 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 urDeviceGetGlobalTimestamps.
-    HostTime =
-        duration_cast<nanoseconds>(steady_clock::now().time_since_epoch())
-            .count();
-    if (Result == UR_RESULT_ERROR_INVALID_OPERATION) {
-      // NOTE(UR port): Removed the call to GetLastError because  we shouldn't
-      // be calling it after ERROR_INVALID_OPERATION: there is no
-      // adapter-specific error.
-      throw detail::set_ur_error(
-          sycl::exception(
-              make_error_code(errc::feature_not_supported),
-              "Device and/or backend does not support querying timestamp."),
-          UR_RESULT_ERROR_INVALID_OPERATION);
-    } else {
-      Adapter->checkUrResult<errc::feature_not_supported>(Result);
-    }
-    // Until next sync we will compute device time based on the host time
-    // returned in HostTime, so make this our base host time.
-    MDeviceHostBaseTime.second = HostTime;
-    Diff = 0;
+    GetGlobalTimestamps(MDevice, &MDeviceHostBaseTime.first,
+                        &MDeviceHostBaseTime.second);
+  } else {
+    GetGlobalTimestamps(MDevice, nullptr, &HostTime);
+    assert(HostTime >= MDeviceHostBaseTime.second);
+    Diff = HostTime - MDeviceHostBaseTime.second;
   }
   return MDeviceHostBaseTime.first + Diff;
 }

unified-runtime/source/adapters/cuda/device.cpp

@@ -1320,6 +1320,14 @@ ur_result_t UR_APICALL urDeviceGetGlobalTimestamps(ur_device_handle_t hDevice,
     UR_CHECK_ERROR(cuEventCreate(&Event, CU_EVENT_DEFAULT));
     UR_CHECK_ERROR(cuEventRecord(Event, 0));
   }
+
+  if (pDeviceTimestamp) {
+    UR_CHECK_ERROR(cuEventSynchronize(Event));
+    *pDeviceTimestamp = hDevice->getElapsedTime(Event);
+  }
+
+  // Record the host timestamp after the cuEventSynchronize() call for more
+  // precise measurement.
   if (pHostTimestamp) {
 
     using namespace std::chrono;
@@ -1328,11 +1336,6 @@ ur_result_t UR_APICALL urDeviceGetGlobalTimestamps(ur_device_handle_t hDevice,
             .count();
   }
 
-  if (pDeviceTimestamp) {
-    UR_CHECK_ERROR(cuEventSynchronize(Event));
-    *pDeviceTimestamp = hDevice->getElapsedTime(Event);
-  }
-
   return UR_RESULT_SUCCESS;
 }
 

unified-runtime/source/adapters/level_zero/device.cpp

@@ -15,6 +15,11 @@
 #include "ur_util.hpp"
 #include <algorithm>
 #include <climits>
+#if defined(__linux__)
+#include <ctime>
+#elif defined(_WIN32)
+#include <windows.h>
+#endif
 #include <optional>
 #include <vector>
 
@@ -1569,6 +1574,40 @@ ur_result_t urDeviceGetGlobalTimestamps(
     /// [out][optional] pointer to the Host's global timestamp that correlates
     /// with the Device's global timestamp value
     uint64_t *HostTimestamp) {
+  if (!DeviceTimestamp && HostTimestamp) {
+    // If only HostTimestamp is requested, we need to avoid making a call to
+    // zeDeviceGetGlobalTimestamps which has higher latency. This is a
+    // workaround for the fact that Level Zero does not provide a way to get the
+    // host timestamp directly. It is known that current implementation of L0
+    // runtime uses CLOCK_MONOTONIC_RAW on Linux and QueryPerformanceCounter on
+    // Windows.
+#if defined(__linux__)
+    timespec Monotonic;
+    if (clock_gettime(CLOCK_MONOTONIC_RAW, &Monotonic) != 0) {
+      UR_LOG(ERR, "Failed to get CLOCK_MONOTONIC time");
+      return UR_RESULT_ERROR_UNINITIALIZED;
+    }
+    *HostTimestamp = static_cast<uint64_t>(Monotonic.tv_sec) * 1'000'000'000 +
+                     static_cast<uint64_t>(Monotonic.tv_nsec);
+    return UR_RESULT_SUCCESS;
+#elif defined(_WIN32)
+    // Use QueryPerformanceCounter on Windows
+    uint64_t Counter;
+    if (!QueryPerformanceCounter((LARGE_INTEGER *)&Counter)) {
+      UR_LOG(ERR, "Failed to get performance counter");
+      return UR_RESULT_ERROR_UNINITIALIZED;
+    }
+    LARGE_INTEGER Frequency;
+    if (!QueryPerformanceFrequency(&Frequency)) {
+      UR_LOG(ERR, "Failed to get performance frequency");
+      return UR_RESULT_ERROR_UNINITIALIZED;
+    }
+    *HostTimestamp = static_cast<uint64_t>(
+        (static_cast<double>(Counter) * 1'000'000'000 / Frequency.QuadPart));
+    return UR_RESULT_SUCCESS;
+#endif
+  }
+
   const uint64_t &ZeTimerResolution =
       Device->ZeDeviceProperties->timerResolution;
   const uint64_t TimestampMaxCount = Device->getTimestampMask();