Skip to content

[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

Merged
merged 3 commits into from
Oct 18, 2021
Merged

Conversation

kbobrovs
Copy link
Contributor

@kbobrovs kbobrovs commented Oct 16, 2021

  • add __esimd_gather_masked_scaled2 intrinsic which is the best choice to map gather to according to the VC BE team
  • 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

Complementary new E2E tests and test fixes: intel/llvm-test-suite#517

Signed-off-by: Konstantin S Bobrovsky [email protected]

- 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]>
@kbobrovs kbobrovs force-pushed the fix_slm_gather_scatter branch from 729582b to 04dc860 Compare October 17, 2021 22:42
@kbobrovs kbobrovs changed the title [WIP] Fix slm_gather/scatter. [ESIMD] Fix slm_gather/scatter. Oct 17, 2021
v-klochkov
v-klochkov previously approved these changes Oct 18, 2021
Copy link
Contributor

@v-klochkov v-klochkov left a 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<
Copy link
Contributor

Choose a reason for hiding this comment

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

Suggested change
ESIMD_INLINE ESIMD_NODEBUG typename sycl::detail::enable_if_t<
ESIMD_INLINE ESIMD_NODEBUG std::enable_if_t<

@kbobrovs
Copy link
Contributor Author

The Jenkins failure is a failure of the single test:
[2021-10-18T01:23:58.551Z] Failed Tests (1):
[2021-10-18T01:23:58.551Z] SYCL :: ESIMD/api/slm_gather_scatter.cpp
which is invalid and fixed in intel/llvm-test-suite#517
The test basically does

simd_mask<N> pred = 1;
pred[MASKED_LANE] = 0;
simd<T,N> val = slm_gather<...>(..., pred);

and expects val[MASKED_LANE] to contain '0', which is wrong, as its contents is undefined.

@kbobrovs kbobrovs requested a review from v-klochkov October 18, 2021 18:43
@kbobrovs kbobrovs merged commit 701e480 into intel:sycl Oct 18, 2021
@kbobrovs kbobrovs deleted the fix_slm_gather_scatter branch January 19, 2022 22:52
@kbobrovs kbobrovs restored the fix_slm_gather_scatter branch January 19, 2022 22:52
@kbobrovs kbobrovs deleted the fix_slm_gather_scatter branch January 19, 2022 22:53
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

2 participants