[SYCL][ESIMD] Fix assignment to private globals (#9608)

The implicitly generated copy assignment operator doesn't do what we
need because we need vloads and vstores for globals.

Add an explicit copy assignment operator to do what we want.

Note this makes the `simd` and `simd_mask` classes not trivially move
assignable or copy assignable, but that is a fundamental requirement to
support globals, so it being trivially assignable was a bug.

Signed-off-by: Sarnie, Nick <[email protected]>

Author: sarnex

sycl/include/sycl/ext/intel/esimd/detail/simd_mask_impl.hpp

@@ -132,6 +132,11 @@ class simd_mask_impl
     return *this;
   }
 
+  /// Copy assignment operator.
+  simd_mask_impl &operator=(const simd_mask_impl &other) noexcept {
+    return base_type::operator=(other);
+  }
+
   /// Conversion to boolean. Available only when the number of elements is 1.
   /// @return true if the element is non-zero, false otherwise.
   template <class T1 = simd_mask_impl,

sycl/include/sycl/ext/intel/esimd/detail/simd_obj_impl.hpp

@@ -347,6 +347,12 @@ class simd_obj_impl {
     copy_from(acc, offset, Flags{});
   }
 
+  /// Copy assignment operator.
+  Derived &operator=(const simd_obj_impl &other) noexcept {
+    set(other.data());
+    return cast_this_to_derived();
+  }
+
   /// Type conversion into a scalar:
   /// <code><simd_obj_impl<RawTy, 1, simd<Ty,1>></code> to \c Ty.
   template <typename T = simd_obj_impl,

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

@@ -126,6 +126,11 @@ class simd : public detail::simd_obj_impl<
     return sycl::ext::oneapi::experimental::simd<Ty, N1>(base_type::data());
   }
 
+  /// Copy assignment operator.
+  simd &operator=(const simd &other) noexcept {
+    return base_type::operator=(other);
+  }
+
   /// Prefix increment, increments elements of this object.
   /// @return Reference to this object.
   simd &operator++() {

sycl/test-e2e/ESIMD/api/functional/operators/operator_assignment_glb.cpp

@@ -0,0 +1,94 @@
+//==--- operator_assignment_glb.cpp  - DPC++ ESIMD on-device test ---==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===---------------------------------------------------------------===//
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+#include "common.hpp"
+#include <iostream>
+#include <sycl/ext/intel/esimd.hpp>
+#include <sycl/sycl.hpp>
+
+using namespace sycl;
+using namespace sycl::ext::intel::esimd;
+
+constexpr unsigned VL = 16;
+
+ESIMD_PRIVATE ESIMD_REGISTER(0) simd<float, VL> va;
+ESIMD_PRIVATE ESIMD_REGISTER(0) simd<float, VL> vc;
+
+int main(void) {
+  constexpr unsigned Size = 1024 * 128;
+
+  std::vector<float> A(Size);
+  std::vector<float> B(Size);
+  std::vector<float> C(Size);
+
+  for (unsigned i = 0; i < Size; ++i) {
+    A[i] = B[i] = i;
+    C[i] = 0.0f;
+  }
+
+  buffer<float, 1> bufa(A.data(), A.size());
+  buffer<float, 1> bufb(B.data(), B.size());
+  buffer<float, 1> bufc(C.data(), C.size());
+
+  try {
+    // We need that many workgroups
+    range<1> GlobalRange{Size / VL};
+
+    // We need that many threads in each group
+    range<1> LocalRange{1};
+
+    queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());
+
+    auto dev = q.get_device();
+    std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
+
+    auto e = q.submit([&](handler &cgh) {
+      auto PA = bufa.get_access<access::mode::read>(cgh);
+      auto PB = bufb.get_access<access::mode::read>(cgh);
+      auto PC = bufc.get_access<access::mode::write>(cgh);
+      cgh.parallel_for<class Test>(
+          GlobalRange * LocalRange, [=](id<1> i) SYCL_ESIMD_KERNEL {
+            using namespace sycl::ext::intel::esimd;
+            unsigned int offset = i * VL * sizeof(float);
+            va.copy_from(PA, offset);
+            simd<float, VL> vb;
+            vb.copy_from(PB, offset);
+            vc = va + vb;
+            vc.copy_to(PC, offset);
+          });
+    });
+    e.wait();
+  } catch (sycl::exception const &e) {
+    std::cout << "SYCL exception caught: " << e.what() << '\n';
+
+    return 1;
+  }
+
+  sycl::host_accessor A_acc(bufa);
+  sycl::host_accessor B_acc(bufb);
+  sycl::host_accessor C_acc(bufc);
+  int err_cnt = 0;
+
+  for (unsigned i = 0; i < Size; ++i) {
+    if (A_acc[i] + B_acc[i] != C_acc[i]) {
+      if (++err_cnt < 10) {
+        std::cout << "failed at index " << i << ", " << C_acc[i]
+                  << " != " << A_acc[i] << " + " << B_acc[i] << "\n";
+      }
+    }
+  }
+  if (err_cnt > 0) {
+    std::cout << "  pass rate: "
+              << ((float)(Size - err_cnt) / (float)Size) * 100.0f << "% ("
+              << (Size - err_cnt) << "/" << Size << ")\n";
+  }
+
+  std::cout << (err_cnt > 0 ? "FAILED\n" : "Passed\n");
+  return err_cnt > 0 ? 1 : 0;
+}

sycl/test-e2e/ESIMD/api/functional/operators/operator_assignment_glb_mask.cpp

@@ -0,0 +1,90 @@
+//==- operator_assignment_glb_mask.cpp  - DPC++ ESIMD on-device test -==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===---------------------------------------------------------------===//
+// TODO: Remove when GPU driver is updated
+// REQUIRES: gpu-intel-pvc
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+#include "common.hpp"
+#include <iostream>
+#include <sycl/ext/intel/esimd.hpp>
+#include <sycl/sycl.hpp>
+
+using namespace sycl;
+using namespace sycl::ext::intel::esimd;
+
+constexpr unsigned VL = 16;
+
+ESIMD_PRIVATE ESIMD_REGISTER(0) simd_mask<VL> va;
+ESIMD_PRIVATE ESIMD_REGISTER(0) simd_mask<VL> vb;
+
+int main(void) {
+  constexpr unsigned Size = 1024;
+
+  std::vector<unsigned short> A(Size);
+  std::vector<unsigned short> B(Size);
+
+  for (unsigned i = 0; i < Size; ++i) {
+    A[i] = i % std::numeric_limits<unsigned short>::max();
+    B[i] = 0;
+  }
+
+  buffer<unsigned short, 1> bufa(A.data(), A.size());
+  buffer<unsigned short, 1> bufb(B.data(), B.size());
+
+  try {
+    // We need that many workgroups
+    range<1> GlobalRange{Size / VL};
+
+    // We need that many threads in each group
+    range<1> LocalRange{1};
+
+    queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());
+
+    auto dev = q.get_device();
+    std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
+
+    auto e = q.submit([&](handler &cgh) {
+      auto PA = bufa.get_access<access::mode::read>(cgh);
+      auto PB = bufb.get_access<access::mode::write>(cgh);
+      cgh.parallel_for<class Test>(
+          GlobalRange * LocalRange, [=](id<1> i) SYCL_ESIMD_KERNEL {
+            using namespace sycl::ext::intel::esimd;
+            unsigned int offset = i * VL * sizeof(unsigned short);
+            va.copy_from(PA, offset);
+            vb = va;
+            vb.copy_to(PB, offset);
+          });
+    });
+    e.wait();
+  } catch (sycl::exception const &e) {
+    std::cout << "SYCL exception caught: " << e.what() << '\n';
+
+    return 1;
+  }
+
+  sycl::host_accessor A_acc(bufa);
+  sycl::host_accessor B_acc(bufb);
+  int err_cnt = 0;
+
+  for (unsigned i = 0; i < Size; ++i) {
+    if (A_acc[i] != B_acc[i]) {
+      if (++err_cnt < 10) {
+        std::cout << "failed at index " << i << ", " << B_acc[i]
+                  << " != " << A_acc[i] << "\n";
+      }
+    }
+  }
+  if (err_cnt > 0) {
+    std::cout << "  pass rate: "
+              << ((float)(Size - err_cnt) / (float)Size) * 100.0f << "% ("
+              << (Size - err_cnt) << "/" << Size << ")\n";
+  }
+
+  std::cout << (err_cnt > 0 ? "FAILED\n" : "Passed\n");
+  return err_cnt > 0 ? 1 : 0;
+}

sycl/test/esimd/simd.cpp

@@ -33,11 +33,11 @@ template <class T> void test_simd_class_traits() SYCL_ESIMD_FUNCTION {
                 "type trait mismatch");
   static_assert(std::is_copy_assignable<simd<T, 4>>::value,
                 "type trait mismatch");
-  static_assert(std::is_trivially_copy_assignable<simd<T, 4>>::value,
+  static_assert(!std::is_trivially_copy_assignable<simd<T, 4>>::value,
                 "type trait mismatch");
   static_assert(std::is_move_assignable<simd<T, 4>>::value,
                 "type trait mismatch");
-  static_assert(std::is_trivially_move_assignable<simd<T, 4>>::value,
+  static_assert(!std::is_trivially_move_assignable<simd<T, 4>>::value,
                 "type trait mismatch");
 }