Skip to content

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

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 2 commits into from
Apr 17, 2024
Merged

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

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
6935dc5
Move rdtsc function out of experimental namespace
fineg74 Apr 16, 2024
eda046d
Address PR comments
fineg74 Apr 16, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 2 additions & 0 deletions sycl/include/sycl/ext/intel/esimd/detail/math_intrin.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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__

Expand Down
7 changes: 7 additions & 0 deletions sycl/include/sycl/ext/intel/esimd/math.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
3 changes: 0 additions & 3 deletions sycl/include/sycl/ext/intel/experimental/esimd/detail/math_intrin.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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

Expand Down
8 changes: 3 additions & 5 deletions sycl/include/sycl/ext/intel/experimental/esimd/math.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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

Expand Down
6 changes: 4 additions & 2 deletions sycl/test-e2e/ESIMD/rdtsc.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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();
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

nit: maybe we call both experimental and non-experimental based on Idx % 2 or something

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Done

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);
Expand Down