-
Notifications
You must be signed in to change notification settings - Fork 787
[ESIMD] Fix slm_gather/scatter. #4771
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
- remove limitations on element size being 4 bytes only - enable 1,8,16,32 elements as allowed by hardware - share implementation with surface-based gather/scatter Signed-off-by: Konstantin S Bobrovsky <[email protected]>
729582b
to
04dc860
Compare
sycl/include/sycl/ext/intel/experimental/esimd/detail/memory_intrin.hpp
Outdated
Show resolved
Hide resolved
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 have several comments requiring only NFC fixes, which can be fixed in this PR or some time later.
/// \ingroup sycl_esimd | ||
template <typename T, int N, typename AccessorTy, | ||
CacheHint L1H = CacheHint::None, CacheHint L3H = CacheHint::None> | ||
ESIMD_INLINE ESIMD_NODEBUG typename sycl::detail::enable_if_t< |
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.
ESIMD_INLINE ESIMD_NODEBUG typename sycl::detail::enable_if_t< | |
ESIMD_INLINE ESIMD_NODEBUG std::enable_if_t< |
The Jenkins failure is a failure of the single test:
and expects |
__esimd_gather_masked_scaled2
intrinsic which is the best choice to mapgather
to according to the VC BE teamComplementary new E2E tests and test fixes: intel/llvm-test-suite#517
Signed-off-by: Konstantin S Bobrovsky [email protected]