-
Notifications
You must be signed in to change notification settings - Fork 801
[ESIMD] Introduce atomic_update<native::lsc::fadd>(...) and similar ops. #6629
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
ec6ce71 to
c6e1ba2
Compare
Signed-off-by: Konstantin S Bobrovsky <konstantin.s.bobrovsky@intel.com>
c6e1ba2 to
92c00d5
Compare
v-klochkov
left a comment
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.
Most of my comments are minor and could be fixed in a separate change,
but there are 2 comments related to naming of the atomic ops that are created in non-experimental namespace and may need sooner resolution to not pollute non-experimental namespace with something to be changed later/soon.
sycl/include/sycl/ext/intel/experimental/esimd/detail/math_intrin.hpp
Outdated
Show resolved
Hide resolved
… detail/memory_intrin.hpp
| } | ||
|
|
||
| /// Base case for checking if a type U is one of the types. | ||
| template <typename U> constexpr bool is_type() { return false; } |
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.
BTW, we have type checker with exactly same definition, but a better name:
https://github.com/intel/llvm/blob/sycl/sycl/include/sycl/ext/intel/experimental/esimd/detail/util.hpp#L32
sycl::ext::intel::esimd::native::lsc::atomic_opesimd::atomic_updateis invoked withesimd::native::lsc::atomic_op(rather thanesimd::atomic_op), then implementation is redirected to Xe-specific LSC-based implementation (won't run until Gen12)esimd::atomic_opcause warning and are replaced with corresponding LSC operations.Complementary E2E test: intel/llvm-test-suite#1171
Signed-off-by: Konstantin S Bobrovsky konstantin.s.bobrovsky@intel.com