[SYCL] Implement sycl_ext_oneapi_clock

llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td

@@ -94,6 +94,7 @@ def AspectExt_oneapi_async_memory_alloc : Aspect<"ext_oneapi_async_memory_alloc"
 def AspectExt_intel_device_info_luid : Aspect<"ext_intel_device_info_luid">;
 def AspectExt_intel_device_info_node_mask : Aspect<"ext_intel_device_info_node_mask">;
 def Aspectext_oneapi_exportable_device_mem : Aspect<"ext_oneapi_exportable_device_mem">;
+def Aspectext_oneapi_clock : Aspect<"ext_oneapi_clock">;
 
 // Deprecated aspects
 def AspectInt64_base_atomics : Aspect<"int64_base_atomics">;
@@ -168,7 +169,8 @@ def : TargetInfo<"__TestAspectList",
     AspectExt_oneapi_async_memory_alloc,
     AspectExt_intel_device_info_luid,
     AspectExt_intel_device_info_node_mask,
-    Aspectext_oneapi_exportable_device_mem],
+    Aspectext_oneapi_exportable_device_mem,
+    Aspectext_oneapi_clock],
     []>;
 // This definition serves the only purpose of testing whether the deprecated aspect list defined in here and in SYCL RT
 // match.

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

@@ -0,0 +1,47 @@
+//==-------- clock.hpp --- SYCL extension for clock() free function --------==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <sycl/__spirv/spirv_ops.hpp>
+#include <sycl/aspects.hpp>
+#include <sycl/exception.hpp>
+
+namespace sycl {
+inline namespace _V1 {
+namespace ext::oneapi::experimental {
+
+enum class clock_scope : int {
+  // Aligned with SPIR-V Scope<id> values
+  device = 1,
+  work_group = 2,
+  sub_group = 3
+};
+
+#ifdef __SYCL_DEVICE_ONLY__
+[[__sycl_detail__::__uses_aspects__(sycl::aspect::ext_oneapi_clock)]]
+#endif // __SYCL_DEVICE_ONLY__
+inline uint64_t
+clock([[maybe_unused]] clock_scope scope = clock_scope::sub_group) {
+#ifdef __SYCL_DEVICE_ONLY__
+#if defined(__NVPTX__) || defined(__AMDGCN__)
+  // Currently clock() is not supported on NVPTX and AMDGCN.
+  return 0;
+#else
+  return __spirv_ReadClockKHR(static_cast<int>(scope));
+#endif // defined(__NVPTX__) || defined(__AMDGCN__)
+#else
+  throw sycl::exception(
+      make_error_code(errc::runtime),
+      "sycl::ext::oneapi::experimental::clock() is not supported on host.");
+#endif // __SYCL_DEVICE_ONLY__
+}
+
+} // namespace ext::oneapi::experimental
+} // namespace _V1
+} // namespace sycl

sycl/include/sycl/info/aspects.def

@@ -80,4 +80,5 @@ __SYCL_ASPECT(ext_oneapi_async_memory_alloc, 87)
 __SYCL_ASPECT(ext_intel_device_info_luid, 88)
 __SYCL_ASPECT(ext_intel_device_info_node_mask, 89)
 __SYCL_ASPECT(ext_oneapi_exportable_device_mem, 90)
+__SYCL_ASPECT(ext_oneapi_clock, 91)
 

sycl/include/sycl/sycl.hpp

@@ -113,6 +113,7 @@ can be disabled by setting SYCL_DISABLE_FSYCL_SYCLHPP_WARNING macro.")
 #include <sycl/ext/oneapi/experimental/bfloat16_math.hpp>
 #include <sycl/ext/oneapi/experimental/builtins.hpp>
 #include <sycl/ext/oneapi/experimental/chunk.hpp>
+#include <sycl/ext/oneapi/experimental/clock.hpp>
 #include <sycl/ext/oneapi/experimental/cluster_group_prop.hpp>
 #include <sycl/ext/oneapi/experimental/composite_device.hpp>
 #include <sycl/ext/oneapi/experimental/cuda/barrier.hpp>

sycl/source/detail/device_impl.hpp

@@ -1579,6 +1579,7 @@ class device_impl : public std::enable_shared_from_this<device_impl> {
                  UR_DEVICE_INFO_MEMORY_EXPORT_EXPORTABLE_DEVICE_MEM_EXP>()
           .value_or(0);
     }
+    CASE(ext_oneapi_clock) { return has_extension("cl_khr_kernel_clock"); }
     else {
       return false; // This device aspect has not been implemented yet.
     }

sycl/source/feature_test.hpp.in

@@ -120,6 +120,7 @@ inline namespace _V1 {
 #define SYCL_KHR_FREE_FUNCTION_COMMANDS 1
 #define SYCL_KHR_QUEUE_EMPTY_QUERY 1
 #define SYCL_EXT_ONEAPI_MEMORY_EXPORT 1
+#define SYCL_EXT_ONEAPI_CLOCK 1
 // In progress yet
 #define SYCL_EXT_ONEAPI_ATOMIC16 0
 #define SYCL_KHR_DEFAULT_CONTEXT 1

sycl/test-e2e/Experimental/clock.cpp

@@ -0,0 +1,43 @@
+// REQUIRES: aspect-ext_oneapi_clock, aspect-usm_shared_allocations
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/oneapi/experimental/clock.hpp>
+#include <sycl/usm.hpp>
+
+int main() {
+  sycl::queue q;
+  uint64_t *data = sycl::malloc_shared<uint64_t>(3, q);
+
+  q.single_task([=]() {
+    uint64_t sg_clock_start = sycl::ext::oneapi::experimental::clock(
+        sycl::ext::oneapi::experimental::clock_scope::sub_group);
+    uint64_t wg_clock_start = sycl::ext::oneapi::experimental::clock(
+        sycl::ext::oneapi::experimental::clock_scope::work_group);
+    uint64_t dev_clock_start = sycl::ext::oneapi::experimental::clock(
+        sycl::ext::oneapi::experimental::clock_scope::device);
+
+    int count = 0;
+    for (int i = 0; i < 1e6; ++i)
+      count++;
+
+    uint64_t sg_clock_end = sycl::ext::oneapi::experimental::clock(
+        sycl::ext::oneapi::experimental::clock_scope::sub_group);
+    uint64_t wg_clock_end = sycl::ext::oneapi::experimental::clock(
+        sycl::ext::oneapi::experimental::clock_scope::work_group);
+    uint64_t dev_clock_end = sycl::ext::oneapi::experimental::clock(
+        sycl::ext::oneapi::experimental::clock_scope::device);
+    data[0] = sg_clock_end - sg_clock_start;
+    data[1] = wg_clock_end - wg_clock_start;
+    data[2] = dev_clock_end - dev_clock_start;
+  });
+  q.wait();
+
+  assert(data[0] > 0);
+  assert(data[1] > 0);
+  assert(data[2] > 0);
+  sycl::free(data, q);
+
+  return 0;
+}