-
Notifications
You must be signed in to change notification settings - Fork 788
[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.
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
Conversation
detail::is_esimd_scalar<TArg>::value, | ||
TRes> | ||
ESIMD_NODEBUG | ||
ESIMD_INLINE std::enable_if_t<detail::is_esimd_scalar<TRes>::value && |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
nit: there are some whitespace changes of this pattern that in my opinion decrease readability, i would expect we only need to touch the rdtsc function and not these
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Fixed
@@ -1302,6 +1301,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 |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I think it's @return
based on other functions
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Fixed
@@ -41,9 +41,9 @@ 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(); |
There was a problem hiding this comment.
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
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Done
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Looks good to me.
Can you please check if this extension is ready for usage, maybe create a copy/draft of this PR with OpReadClockKHR
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
The change is correct as written now, and the experiment with using SPIRV instead of GenX intrinsic can be done in a separate PR.
No description provided.