rest

Author: KornevNikita

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/source/detail/device_impl.hpp

@@ -1579,6 +1579,9 @@ 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,40 @@
+// REQUIRES: ext_oneapi_clock
+
+#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);
+
+  return 0;
+}