[SYCL][ESIMD] Move rdtsc function out of experimental namespace (#13417)

Author: fineg74

sycl/include/sycl/ext/intel/esimd/detail/math_intrin.hpp

@@ -193,6 +193,8 @@ __ESIMD_INTRIN __ESIMD_raw_vec_t(T1, N)
     __esimd_ssdp4a_sat(__ESIMD_raw_vec_t(T2, N) src0,
                        __ESIMD_raw_vec_t(T3, N) src1,
                        __ESIMD_raw_vec_t(T4, N) src2) __ESIMD_INTRIN_END;
+__ESIMD_INTRIN __ESIMD_raw_vec_t(uint32_t, 4)
+    __esimd_timestamp() __ESIMD_INTRIN_END;
 
 #ifdef __SYCL_DEVICE_ONLY__
 

sycl/include/sycl/ext/intel/esimd/math.hpp

@@ -1302,6 +1302,13 @@ __ESIMD_API uint32_t subb(uint32_t &borrow, uint32_t src0, uint32_t src1) {
   return Res[0];
 }
 
+/// rdtsc - get the value of timestamp counter.
+/// @return the current value of timestamp counter
+__ESIMD_API uint64_t rdtsc() {
+  __ESIMD_NS::simd<uint32_t, 4> retv = __esimd_timestamp();
+  return retv.template bit_cast_view<uint64_t>()[0];
+}
+
 /// @} sycl_esimd_math
 
 } // namespace ext::intel::esimd

sycl/include/sycl/ext/intel/experimental/esimd/detail/math_intrin.hpp

@@ -151,9 +151,6 @@ __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
 

sycl/include/sycl/ext/intel/experimental/esimd/math.hpp

@@ -1724,11 +1724,9 @@ __ESIMD_API std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T>::value &&
 }
 
 /// 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];
-}
+/// @return the current value of timestamp counter
+__SYCL_DEPRECATED("Please use sycl::ext::intel::esimd::rdtsc();")
+ESIMD_INLINE uint64_t rdtsc() { return __ESIMD_NS::rdtsc(); }
 
 /// @} sycl_esimd_logical
 

sycl/test-e2e/ESIMD/rdtsc.cpp

@@ -41,9 +41,11 @@ int test_rdtsc() {
       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();
+        uint64_t StartCounter =
+            Idx % 2 == 0 ? __ESIMD_NS::rdtsc() : __ESIMD_ENS::rdtsc();
         simd<uint64_t, 1> VectorResultRDTSC(VectorOutputRDTSCPtr + Idx);
-        uint64_t EndCounter = sycl::ext::intel::experimental::esimd::rdtsc();
+        uint64_t EndCounter =
+            Idx % 2 == 0 ? __ESIMD_NS::rdtsc() : __ESIMD_ENS::rdtsc();
         VectorResultRDTSC += EndCounter > StartCounter;
 
         VectorResultRDTSC.copy_to(VectorOutputRDTSCPtr + Idx);