From 1261e0518812ca8f5aecbf9d2801d0dd45b60974 Mon Sep 17 00:00:00 2001
From: fineg74 <61437305+fineg74@users.noreply.github.com>
Date: Wed, 10 Jan 2024 11:45:31 -0800
Subject: [PATCH] [SYCL][ESIMD] Implement rdtsc() - ReaD TimeStamp Counter API
 (#12315)

---
 llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp     |  3 +-
 .../experimental/esimd/detail/math_intrin.hpp |  3 +
 .../ext/intel/experimental/esimd/math.hpp     |  7 ++
 sycl/test-e2e/ESIMD/rdtsc.cpp                 | 69 +++++++++++++++++++
 4 files changed, 81 insertions(+), 1 deletion(-)
 create mode 100644 sycl/test-e2e/ESIMD/rdtsc.cpp

diff --git a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp
index 0379987f8af40..0f767e434d918 100644
--- a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp
+++ b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp
@@ -661,7 +661,8 @@ class ESIMDIntrinDescTable {
         {"addc", {"addc", {l(0)}}},
         {"subb", {"subb", {l(0)}}},
         {"bfn", {"bfn", {a(0), a(1), a(2), t(0)}}},
-        {"srnd", {"srnd", {a(0), a(1)}}}};
+        {"srnd", {"srnd", {a(0), a(1)}}},
+        {"timestamp",{"timestamp",{}}}};
   }
   // clang-format on
 
diff --git a/sycl/include/sycl/ext/intel/experimental/esimd/detail/math_intrin.hpp b/sycl/include/sycl/ext/intel/experimental/esimd/detail/math_intrin.hpp
index 7be20c432ebfc..52aba4c817c96 100644
--- a/sycl/include/sycl/ext/intel/experimental/esimd/detail/math_intrin.hpp
+++ b/sycl/include/sycl/ext/intel/experimental/esimd/detail/math_intrin.hpp
@@ -152,6 +152,9 @@ __ESIMD_INTRIN __ESIMD_raw_vec_t(sycl::half, N)
                  __ESIMD_DNS::vector_type_t<uint16_t, N> src2)
         __ESIMD_INTRIN_END;
 
+__ESIMD_INTRIN __ESIMD_raw_vec_t(uint32_t, 4)
+    __esimd_timestamp() __ESIMD_INTRIN_END;
+
 #undef __ESIMD_raw_vec_t
 #undef __ESIMD_cpp_vec_t
 
diff --git a/sycl/include/sycl/ext/intel/experimental/esimd/math.hpp b/sycl/include/sycl/ext/intel/experimental/esimd/math.hpp
index e69f5686e2aa6..ca212c66e1129 100644
--- a/sycl/include/sycl/ext/intel/experimental/esimd/math.hpp
+++ b/sycl/include/sycl/ext/intel/experimental/esimd/math.hpp
@@ -1723,6 +1723,13 @@ __ESIMD_API std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T>::value &&
   return __ESIMD_NS::bfn<FuncControl>(src0, src1, src2);
 }
 
+/// rdtsc - get the value of timestamp counter.
+/// \return the current value of timestamp counter
+ESIMD_INLINE uint64_t rdtsc() {
+  __ESIMD_NS::simd<uint32_t, 4> retv = __esimd_timestamp();
+  return retv.template bit_cast_view<uint64_t>()[0];
+}
+
 /// @} sycl_esimd_logical
 
 } // namespace ext::intel::experimental::esimd
diff --git a/sycl/test-e2e/ESIMD/rdtsc.cpp b/sycl/test-e2e/ESIMD/rdtsc.cpp
new file mode 100644
index 0000000000000..a263348a884d3
--- /dev/null
+++ b/sycl/test-e2e/ESIMD/rdtsc.cpp
@@ -0,0 +1,69 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+//==- rdtsc.cpp - Test to verify rdtsc0 and sr0 functionlity----------------==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+// This is basic test to validate rdtsc function.
+
+#include <cmath>
+#include <iostream>
+#include <sycl/ext/intel/esimd.hpp>
+#include <sycl/ext/intel/esimd/simd.hpp>
+#include <sycl/sycl.hpp>
+#include <vector>
+
+int ErrCnt = 0;
+template <typename DataT>
+using shared_allocator = sycl::usm_allocator<DataT, sycl::usm::alloc::shared>;
+template <typename DataT>
+using shared_vector = std::vector<DataT, shared_allocator<DataT>>;
+
+int test_rdtsc() {
+  sycl::queue Queue;
+  shared_allocator<uint64_t> Allocator(Queue);
+  constexpr int32_t SIZE = 32;
+
+  shared_vector<uint64_t> VectorOutputRDTSC(SIZE, 0, Allocator);
+
+  auto GlobalRange = sycl::range<1>(SIZE);
+  sycl::range<1> LocalRange{1};
+  sycl::nd_range<1> Range(GlobalRange, LocalRange);
+
+  {
+    Queue.submit([&](sycl::handler &cgh) {
+      uint64_t *VectorOutputRDTSCPtr = VectorOutputRDTSC.data();
+
+      auto Kernel = ([=](sycl::nd_item<1> ndi) [[intel::sycl_explicit_simd]] {
+        using namespace sycl::ext::intel::esimd;
+        auto Idx = ndi.get_global_id(0);
+        uint64_t StartCounter = sycl::ext::intel::experimental::esimd::rdtsc();
+        simd<uint64_t, 1> VectorResultRDTSC(VectorOutputRDTSCPtr + Idx);
+        uint64_t EndCounter = sycl::ext::intel::experimental::esimd::rdtsc();
+        VectorResultRDTSC += EndCounter > StartCounter;
+
+        VectorResultRDTSC.copy_to(VectorOutputRDTSCPtr + Idx);
+      });
+
+      cgh.parallel_for(Range, Kernel);
+    });
+    Queue.wait();
+  }
+
+  return std::any_of(VectorOutputRDTSC.begin(), VectorOutputRDTSC.end(),
+                     [](uint64_t v) { return v == 0; });
+}
+
+int main() {
+
+  int TestResult = test_rdtsc();
+
+  if (!TestResult) {
+    std::cout << "Pass" << std::endl;
+  }
+  return TestResult;
+}