[SYCL][ESIMD] Add capability to specify 64 bit offsets to esimd functions #7411

fineg74 · 2022-11-16T05:41:45Z

No description provided.

This reverts commit 918f2d0.

fineg74 · 2022-11-16T05:42:36Z

Complementary test PR : intel/llvm-test-suite#1385

fineg74 · 2022-11-16T17:45:49Z

Test failures in
SYCL :: ESIMD/Stencil.cpp
SYCL :: ESIMD/stencil2.cpp
are expected and are fixed in test PR

kbobrovs · 2022-11-16T19:28:02Z

sycl/include/sycl/ext/intel/esimd/memory.hpp

-scatter(Tx *p, simd<uint32_t, N> offsets, simd<Tx, N> vals,
+template <typename Tx, int N, class T = detail::__raw_t<Tx>, typename Toffset>
+__ESIMD_API std::enable_if_t<detail::isPowerOf2(N, 32) &&
+                             (std::is_same_v<Toffset, uint32_t> ||


I believe this can break existing code. E.g. int offsets will now cause compilation error, but should have been working in the past.

The only code breaking problem I saw so far is when offset is a simd_view i.e. something like:
scatter(p, x.select<1,8>(0), vals) and the reason is that simd_view converts to simd using operator() which can convert the simd_view to simd of any type. It wasn't the issue when functions received only a single type but becomes the issue when function can accept multiple types so the compiler can't choose which type to use.
Using multiple functions instead of template i.e. having scatter(Tx *p, simd<uint32_t, N> offsets, simd<Tx, N> vals) and scatter(Tx *p, simd<uint64_t, N> offsets, simd<Tx, N> vals) doesn't solve the problem as the compiler faces the same issue.
I am not sure that the code we have now is good one as it would accept even vector of floats without any issue (I have 2 tests where offsets and vals were flipped and did not cause any compilation issue

Backward compatibility of non-experimental code must not be broken. This is a hard requirement. The only exception is approved list of APIs changed in major releases.

The only code breaking problem I saw so far is when offset is a simd_view i.e. something like:
scatter(p, x.select<1,8>(0), vals)

that is a problem too.
I tried short as offset type for scatter - it works now. So it should continue to work. My suggestion is:

create template<class OffsetT> ... scatter_impl(..., simd<OffsetT, N> offsets,...) in namespace detail.

create ...scatter(Tx *p, simd<uint64_t, N> offsets,...

have ...scatter(Tx *p, simd<uint64_t, N> offsets,... and ...scatter(Tx *p, simd<uint32_t, N> offsets,... delegate to scatter_impl.

I am not sure that the code we have now is good one as it would accept even vector of floats without any issue (I have 2 tests where offsets and vals were flipped and did not cause any compilation issue

Implicit conversion between arithmetic types in most contexts is basic C++ behavior, esimd::simd tries to follow. The user logic error you mention is desirable to protect against, of course, but even the enable_if_t you added does not fully shield from this if element type is uint32/64.

Removed type checks for offset types, so it should be now compatible with existing code.
I believe the suggested approach would break the interface as there is ambiguity which function to use.
For example here is the error I got when I tried to use this approach with test where offsets were defined as vector of floats while it compiles perfectly with current version with removed offset type checks. I believe there will be similar issues when offset types other than uint32_t or uint64_t will be provided.

/home/gregory/src/dpc/sysl_workspace/work/tests/llvm-test-suite/SYCL/ESIMD/Stencil.cpp:179:17: error: call to 'scatter' is ambiguous scatter<float, WIDTH>(outputMatrix, sum, elm16_off, p); ^~~~~~~~~~~~~~~~~~~~~ /home/gregory/src/dpc/sysl_workspace/work/llvm/build/bin/../include/sycl/ext/intel/esimd/memory.hpp:218:1: note: candidate function [with Tx = float, N = 16, T = float] scatter(Tx *p, simd<uint32_t, N> offsets, simd<Tx, N> vals, ^ /home/gregory/src/dpc/sysl_workspace/work/llvm/build/bin/../include/sycl/ext/intel/esimd/memory.hpp:225:1: note: candidate function [with Tx = float, N = 16, T = float] scatter(Tx *p, simd<uint64_t, N> offsets, simd<Tx, N> vals,

kbobrovs · 2022-11-16T20:55:34Z

sycl/include/sycl/ext/intel/esimd/memory.hpp

+          int N, typename Toffset>
+__ESIMD_API std::enable_if_t<(N == 8 || N == 16 || N == 32) && sizeof(T) == 4 &&
+                                 (std::is_same_v<Toffset, uint32_t> ||
+                                  std::is_same_v<Toffset, uint64_t>),


Please make sure there is a test with 64-bit offset and N.

kbobrovs · 2022-11-16T20:59:22Z

sycl/include/sycl/ext/intel/esimd/memory.hpp

-scatter_rgba(T *p, simd<uint32_t, N> offsets,
+          int N, typename Toffset>
+__ESIMD_API std::enable_if_t<(N == 8 || N == 16 || N == 32) && sizeof(T) == 4 &&
+                             (std::is_same_v<Toffset, uint32_t> ||


I'd suggest to use the same mechanism for template parameter checking in all APIs. Here std::enable_if_t is used, but in lsc_gather below - static_assert.
static_assert seems better choice, as gives clearer idea of the problem to the user.

The issue is that non-lsc API uses std::enable_if_t approach while lsc API uses static_assert.
We probably need to decide on common approach between these 2 APIs

That is what I suggest - use static_assert everywhere

kbobrovs · 2022-11-16T21:00:33Z

sycl/include/sycl/ext/intel/esimd/memory.hpp

-__ESIMD_API simd<Tx, N> atomic_update(Tx *p, simd<unsigned, N> offset,
-                                      simd<Tx, N> src0, simd_mask<N> mask) {
+template <atomic_op Op, typename Tx, int N, typename Toffset>
+__ESIMD_API std::enable_if_t<std::is_same_v<Toffset, uint32_t> ||


same here an in other places - template parameter checking approach should be consistent in all APIs.

This reverts commit 70f9020.

kbobrovs · 2022-11-29T18:49:33Z

sycl/include/sycl/ext/intel/esimd/memory.hpp

-template <typename Tx, int N, class T = detail::__raw_t<Tx>>
-__ESIMD_API std::enable_if_t<detail::isPowerOf2(N, 32), simd<Tx, N>>
-gather(const Tx *p, simd<uint32_t, N> offsets, simd_mask<N> mask = 1) {
+template <typename Tx, int N, class T = detail::__raw_t<Tx>, typename Toffset>


Here and in other places:
Parameters with default values (T) should go after parameters w/o default values.

Actually, T calculation should be moved out of the parameter list and replaced with using T = detail::__raw_t<Tx>, it is never supposed to be set by the user.

kbobrovs · 2022-11-29T18:56:15Z

sycl/include/sycl/ext/intel/esimd/memory.hpp

+scatter_rgba(T *p, simd_view<Toffset, RegionTy> offsets,
+             simd<T, N * get_num_channels_enabled(RGBAMask)> vals,
+             simd_mask<N> mask = 1) {
+  using Ty = typename simd_view<Toffset, RegionTy>::element_type;


add assert that the number of offsets matches the number of vals

They are not expected to have the same number of elements. One element in offsets controls multiple elements in vals as specified by mask. It is enforced by template specialization

kbobrovs · 2022-11-29T18:57:43Z

sycl/include/sycl/ext/intel/esimd/memory.hpp

@@ -611,6 +647,17 @@ scatter_rgba(T *p, simd<uint32_t, N> offsets,
      addrs.data(), vals.data(), mask.data());
 }

+template <rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T,


Here and in other new overloads:
please add doxygen, and explain what is the difference with the other scatter_rgba overload.

kbobrovs · 2022-11-30T03:35:55Z

sycl/include/sycl/ext/intel/esimd/memory.hpp

+/// Loads ("gathers") elements from different memory locations and returns a
+/// vector of them. Each memory location is base address plus an offset - a
+/// value of the corresponding element in the input offset vector. Access to
+/// any element's memory location can be disabled via the input vector of
+/// predicates (mask).


Here and in other places:

Suggested change

/// Loads ("gathers") elements from different memory locations and returns a

/// vector of them. Each memory location is base address plus an offset - a

/// value of the corresponding element in the input offset vector. Access to

/// any element's memory location can be disabled via the input vector of

/// predicates (mask).

/// A variation of \c gather API with \c offsets represented as a \c simd_view object.

kbobrovs · 2022-11-30T03:37:26Z

sycl/include/sycl/ext/intel/esimd/memory.hpp

@@ -580,6 +613,39 @@ gather_rgba(const T *p, simd<Toffset, N> offsets, simd_mask<N> mask = 1) {
      addrs.data(), mask.data());
 }

+/// @anchor usm_gather_rgba


Here my comment above is especially relevant.

kbobrovs · 2022-11-30T20:04:41Z

sycl/include/sycl/ext/intel/esimd/memory.hpp

@@ -148,6 +152,19 @@ gather(const Tx *p, simd<uint32_t, N> offsets, simd_mask<N> mask = 1) {
                                                                 mask.data());
 }

+/// A variation of \c gather API with \c offsets represented as \c simd_view
+/// object
+///


My review request was actually to get rid of description duplication only. Please return parameter doxygen comments.

kbobrovs · 2022-12-01T08:18:39Z

Looks like more fixes are needed (might be a test problem):

******************** TEST 'SYCL :: ESIMD/Stencil.cpp' FAILED ********************
...
/__w/llvm/llvm/toolchain/bin/../include/sycl/ext/intel/esimd/memory.hpp:195:3: error: static assertion failed due to requirement 'std::is_integral_v<float>': Unsupported offset type
  static_assert(std::is_integral_v<Toffset>, "Unsupported offset type");
  ^             ~~~~~~~~~~~~~~~~~~~~~~~~~~~
/__w/llvm/llvm/llvm_test_suite/SYCL/ESIMD/Stencil.cpp:179:17: note: in instantiation of function template specialization 'sycl::_V1::ext::intel::esimd::scatter<float, 16, float>' requested here
                scatter<float, WIDTH>(outputMatrix, sum, elm16_off, p);
                ^
1 error generated.

fineg74 · 2022-12-01T17:32:57Z

Looks like more fixes are needed (might be a test problem):

******************** TEST 'SYCL :: ESIMD/Stencil.cpp' FAILED ********************
...
/__w/llvm/llvm/toolchain/bin/../include/sycl/ext/intel/esimd/memory.hpp:195:3: error: static assertion failed due to requirement 'std::is_integral_v<float>': Unsupported offset type
  static_assert(std::is_integral_v<Toffset>, "Unsupported offset type");
  ^             ~~~~~~~~~~~~~~~~~~~~~~~~~~~
/__w/llvm/llvm/llvm_test_suite/SYCL/ESIMD/Stencil.cpp:179:17: note: in instantiation of function template specialization 'sycl::_V1::ext::intel::esimd::scatter<float, 16, float>' requested here
                scatter<float, WIDTH>(outputMatrix, sum, elm16_off, p);
                ^
1 error generated.

It is known issue and it is fixed in test PR.

fineg74 · 2022-12-01T21:38:40Z

/verify with intel/llvm-test-suite#1385

kbobrovs · 2022-12-05T08:50:28Z

waiting for comments on failed tests

fineg74 · 2022-12-05T19:05:07Z

/verify with intel/llvm-test-suite#1385

fineg74 · 2022-12-07T04:14:49Z

/verify with intel/llvm-test-suite#1385

kbobrovs · 2022-12-08T18:51:30Z

Please provide analysis of the failed tests (or state which ones are unrelated)

fineg74 · 2022-12-09T17:56:55Z

Test failures in
SYCL :: ESIMD/Stencil.cpp
SYCL :: ESIMD/stencil2.cpp
are expected and are fixed in test PR

fineg74 · 2022-12-13T23:21:36Z

/verify with intel/llvm-test-suite#1385

fineg74 · 2022-12-14T00:46:28Z

Test failures in
SYCL :: ESIMD/Stencil.cpp
SYCL :: ESIMD/stencil2.cpp
are expected and are fixed in test PR

fineg74 · 2022-12-14T03:40:46Z

llvm-test-suite failures

HostInteropTask/host-task-failure.cpp
Basic/group_async_copy.cpp
DeviceLib/imf_fp16_trivial_test.cpp
DeviceLib/imf_fp32_test.cpp
DeviceLib/imf_half_type_cast.cpp
Reduction/reduction_big_data.cpp
Reduction/reduction_nd_N_vars.cpp
Reduction/reduction_nd_conditional.cpp
Reduction/reduction_nd_dw.cpp
Reduction/reduction_nd_ext_half.cpp
Reduction/reduction_nd_lambda.cpp
Reduction/reduction_nd_rw.cpp
Reduction/reduction_range_1d_dw.cpp
Reduction/reduction_range_1d_rw.cpp
Reduction/reduction_range_2d_dw.cpp
Reduction/reduction_range_2d_rw.cpp
Reduction/reduction_range_3d_dw.cpp
Reduction/reduction_range_3d_rw.cpp
Reduction/reduction_range_N_vars.cpp
Reduction/reduction_usm.cpp
Reduction/reduction_usm_dw.cpp
are not related to the change

fineg74 added 3 commits November 15, 2022 11:47

Add support for 64 bit offsets

7306408

Fix test failures

918f2d0

Revert "Fix test failures"

a6e2ea3

This reverts commit 918f2d0.

fineg74 requested a review from a team as a code owner November 16, 2022 05:41

fineg74 mentioned this pull request Nov 16, 2022

[SYCL][ESIMD] Add tests for ESIMD functions accepting 64 bit offsets intel/llvm-test-suite#1385

Merged

kbobrovs reviewed Nov 16, 2022

View reviewed changes

fineg74 added 3 commits November 16, 2022 14:20

Address PR comments

70f9020

Revert "Address PR comments"

eaef64f

This reverts commit 70f9020.

Address PR comments

7e5eb5c

kbobrovs reviewed Nov 29, 2022

View reviewed changes

Address PR comments

9a4a2b2

kbobrovs reviewed Nov 30, 2022

View reviewed changes

Address PR comments

30cd76e

kbobrovs reviewed Nov 30, 2022

View reviewed changes

Address PR comments

b912510

kbobrovs approved these changes Dec 1, 2022

View reviewed changes

Merge remote-tracking branch 'origin/sycl' into 64bitOffset

e9b90eb

Merge remote-tracking branch 'intel_llvm1/sycl' into 64bitOffset

9a12347

v-klochkov approved these changes Dec 14, 2022

View reviewed changes

v-klochkov merged commit c63f802 into intel:sycl Dec 14, 2022

fineg74 deleted the 64bitOffset branch December 14, 2022 18:35

fineg74 restored the 64bitOffset branch May 17, 2023 17:03

fineg74 deleted the 64bitOffset branch May 17, 2023 17:04

v-klochkov mentioned this pull request Nov 29, 2023

[SYCL][ESIMD] Implement unified memory API - block_store slm and local accessors #11921

Merged

[SYCL][ESIMD] Add capability to specify 64 bit offsets to esimd functions #7411

[SYCL][ESIMD] Add capability to specify 64 bit offsets to esimd functions #7411

Uh oh!

Conversation

fineg74 commented Nov 16, 2022

Uh oh!

fineg74 commented Nov 16, 2022

Uh oh!

fineg74 commented Nov 16, 2022 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

kbobrovs Nov 29, 2022 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

kbobrovs Nov 30, 2022 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

kbobrovs commented Dec 1, 2022

Uh oh!

fineg74 commented Dec 1, 2022

Uh oh!

fineg74 commented Dec 1, 2022

Uh oh!

kbobrovs commented Dec 5, 2022

Uh oh!

fineg74 commented Dec 5, 2022

Uh oh!

fineg74 commented Dec 7, 2022

Uh oh!

kbobrovs commented Dec 8, 2022

Uh oh!

fineg74 commented Dec 9, 2022

Uh oh!

fineg74 commented Dec 13, 2022

Uh oh!

fineg74 commented Dec 14, 2022

Uh oh!

fineg74 commented Dec 14, 2022

Uh oh!

Uh oh!

fineg74 commented Nov 16, 2022 •

edited

Loading

kbobrovs Nov 29, 2022 •

edited

Loading

kbobrovs Nov 30, 2022 •

edited

Loading