[SYCL][ESIMD] Add support for align flags for simd::copy_from/to operations

[sndmitriev]

This patch adds new parameter to simd<>::copy_from/copy_to which allows
to specify memory alignment for the load/store address. Depending on the
provided alignment, operation is lowered to an appropriate low-level memory
operation with matching alignment constraints.

Signed-off-by: Sergey Dmitriev serguei.n.dmitriev@intel.com

[kbobrovs]

please also add a reference to an E2E test to this PR

[kbobrovs]

@Pennycook, @rolandschulz - you might be interested in this too.

[kbobrovs]

thanks, LGTM, but some tests fail in Jenkins.

[sndmitriev]

Yes, some tests are failing, I will check why.

[v-klochkov]

Should it be Align >= OperandSize::DWORD or Align >= OperandSize::OWORD here?

[sndmitriev]

Block load requires offset to be at least dword aligned, so it should be OperandSize::DWORD. block_load() then will check what alignment is and will use either aligned load if alignment is >= OperandSize::OWORD or unaligned otherwise.

[kbobrovs]

please also add a reference to an E2E test to this PR

@sndmitriev, friendly reminder. Is there E2E test yet?

[sndmitriev]

please also add a reference to an E2E test to this PR

@sndmitriev, friendly reminder. Is there E2E test yet?

Yes, it is here intel/llvm-test-suite#560

[sndmitriev]

I expect to see two failures on precommit testing

SYCL :: ESIMD/reduction.cpp

Test passes on GPU but fails on host. Looks like host implementation of gather<int16_t, …> works incorrectly in some cases. I will try to create a small reproducer for this problem.

SYCL :: ESIMD/kmeans/kmeans.cpp

Compiler cannot resolve an ambiguity between array and load constructors for simd object

/// Construct from an array. To allow e.g. simd_mask_type<N> m({1,0,0,1,...}).
template <int N1, class = std::enable_if_t<N1 == N>>
simd_obj_impl(const Ty(&&Arr)[N1]) noexcept

/// Load constructor.
template <typename Flags = element_aligned_tag,
typename = std::enable_if_t<is_simd_flag_type_v<Flags>>>
simd_obj_impl(const Ty *ptr, Flags = {}) noexcept {

I can update test to resolve this ambiguity, but I am not sure if this is the right way to fix this problem. Do you have any suggestions?

[sndmitriev]

The problem with int16_t gather on host can be reproduced on the following test

$ cat gather16.cpp 
#include <CL/sycl.hpp>
#include <array>
#include <iostream>
#include <numeric>

#include <sycl/ext/intel/experimental/esimd.hpp>

using namespace cl::sycl;
using namespace sycl::ext::intel::experimental;
using namespace sycl::ext::intel::experimental::esimd;

int main(void) {
  constexpr int N = 8;
  using T = int16_t;

  std::array<T, N> Ref;
  std::iota(Ref.begin(), Ref.end(), 1);

  simd<T, N> Tst = gather<T, N>(Ref.data(), simd<uint32_t, N>(0u, sizeof(T)));

  for (int I = 0; I < N; ++I)
    if (Ref[I] != Tst[I]) {
      std::cout << "failed at " << I << ": " << Tst[I] << " (Tst) != " << Ref[I] << " (Ref)\n";
      return 1;
    }

  std::cout << "passed\n";
  return 0;
}
$ clang++ -fsycl gather16.cpp 
$ ./a.out 
failed at 1: 0 (Tst) != 2 (Ref)
$ 

[kbobrovs]

I can update test to resolve this ambiguity, but I am not sure if this is the right way to fix this problem. Do you have any suggestions?

We might need to remove the array-based constructor until we can find a better solution :( I noticed lately that it does not get called in supposed context anyway (tests somehow did not catch that). Unless removal of initializer_list-based constructor makes compiler pick-up the array-based one in wanted cases. Can the ambiguity be resolved on the ESIMD library side?

[sndmitriev]

I am not sure how this can be resolved in the ESIMD library. Maybe it makes sense to remove load constructor from this patch until we find the right solution for the ambiguity problem?

[sndmitriev]

I just noticed that on kmeans ambiguity was actually reported for simd::copy_to/from rather than for the constructors, so please ignore my notes about the constructors. The ambiguity has to be resolved for copy_to/from overloads (pointer version vs array specialization).