[SYCL][Reduction] Make reducer uncopyable and immovable

sycl/include/sycl/reduction.hpp

@@ -422,6 +422,13 @@ class ReductionIdentityContainer<
   static constexpr bool has_identity = false;
 };
 
+// Token class to help with the in-place construction of reducers.
+template <class BinaryOperation, typename IdentityContainerT>
+struct ReducerToken {
+  const IdentityContainerT &IdentityContainer;
+  const BinaryOperation BOp;
+};
+
 } // namespace detail
 
 /// Specialization of the generic class 'reducer'. It is used for reductions
@@ -458,6 +465,14 @@ class reducer<
   reducer(const IdentityContainerT &IdentityContainer, BinaryOperation BOp)
       : MValue(GetInitialValue(IdentityContainer)),
         MIdentity(IdentityContainer), MBinaryOp(BOp) {}
+  reducer(
+      const detail::ReducerToken<BinaryOperation, IdentityContainerT> &Token)
+      : reducer(Token.IdentityContainer, Token.BOp) {}
+
+  reducer(const reducer &) = delete;
+  reducer(reducer &&) = delete;
+  reducer &operator=(const reducer &) = delete;
+  reducer &operator=(reducer &&) = delete;
 
   reducer &combine(const T &Partial) {
     if constexpr (has_identity)
@@ -515,6 +530,14 @@ class reducer<
   reducer() : MValue(getIdentity()) {}
   reducer(const IdentityContainerT & /* Identity */, BinaryOperation)
       : MValue(getIdentity()) {}
+  reducer(
+      const detail::ReducerToken<BinaryOperation, IdentityContainerT> &Token)
+      : reducer(Token.IdentityContainer, Token.BOp) {}
+
+  reducer(const reducer &) = delete;
+  reducer(reducer &&) = delete;
+  reducer &operator=(const reducer &) = delete;
+  reducer &operator=(reducer &&) = delete;
 
   reducer &combine(const T &Partial) {
     BinaryOperation BOp;
@@ -553,6 +576,14 @@ class reducer<T, BinaryOperation, Dims, Extent, IdentityContainerT, View,
 public:
   reducer(internal_value_type &Ref, BinaryOperation BOp)
       : MElement(Ref), MBinaryOp(BOp) {}
+  reducer(
+      const detail::ReducerToken<BinaryOperation, IdentityContainerT> &Token)
+      : reducer(Token.IdentityContainer, Token.BOp) {}
+
+  reducer(const reducer &) = delete;
+  reducer(reducer &&) = delete;
+  reducer &operator=(const reducer &) = delete;
+  reducer &operator=(reducer &&) = delete;
 
   reducer &combine(const T &Partial) {
     if constexpr (has_identity)
@@ -599,6 +630,14 @@ class reducer<
   reducer(const IdentityContainerT &IdentityContainer, BinaryOperation BOp)
       : MValue(GetInitialValue(IdentityContainer)),
         MIdentity(IdentityContainer), MBinaryOp(BOp) {}
+  reducer(
+      const detail::ReducerToken<BinaryOperation, IdentityContainerT> &Token)
+      : reducer(Token.IdentityContainer, Token.BOp) {}
+
+  reducer(const reducer &) = delete;
+  reducer(reducer &&) = delete;
+  reducer &operator=(const reducer &) = delete;
+  reducer &operator=(reducer &&) = delete;
 
   reducer<T, BinaryOperation, Dims - 1, Extent, IdentityContainerT, true>
   operator[](size_t Index) {
@@ -650,6 +689,14 @@ class reducer<
   reducer() : MValue(getIdentity()) {}
   reducer(const IdentityContainerT & /* Identity */, BinaryOperation)
       : MValue(getIdentity()) {}
+  reducer(
+      const detail::ReducerToken<BinaryOperation, IdentityContainerT> &Token)
+      : reducer(Token.IdentityContainer, Token.BOp) {}
+
+  reducer(const reducer &) = delete;
+  reducer(reducer &&) = delete;
+  reducer &operator=(const reducer &) = delete;
+  reducer &operator=(reducer &&) = delete;
 
   // SYCL 2020 revision 4 says this should be const, but this is a bug
   // see https://github.com/KhronosGroup/SYCL-Docs/pull/252
@@ -746,6 +793,8 @@ class reduction_impl_algo {
 
   using identity_container_type =
       ReductionIdentityContainer<T, BinaryOperation, ExplicitIdentity>;
+  using reducer_token_type =
+      detail::ReducerToken<BinaryOperation, identity_container_type>;
   using reducer_type =
       reducer<T, BinaryOperation, Dims, Extent, identity_container_type>;
   using result_type = T;
@@ -2061,8 +2110,11 @@ void reduCGFuncMulti(handler &CGH, KernelType KernelFunc,
       // Pass all reductions to user's lambda in the same order as supplied
       // Each reducer initializes its own storage
       auto ReduIndices = std::index_sequence_for<Reductions...>();
-      auto ReducersTuple = std::tuple{typename Reductions::reducer_type{
-          std::get<Is>(IdentitiesTuple), std::get<Is>(BOPsTuple)}...};
+      auto ReducerTokensTuple =
+          std::tuple{typename Reductions::reducer_token_type{
+              std::get<Is>(IdentitiesTuple), std::get<Is>(BOPsTuple)}...};
+      auto ReducersTuple = std::tuple<typename Reductions::reducer_type...>{
+          std::get<Is>(ReducerTokensTuple)...};
       std::apply([&](auto &...Reducers) { KernelFunc(NDIt, Reducers...); },
                  ReducersTuple);
 

sycl/test/basic_tests/reduction/reducer_copy_move.cpp

@@ -0,0 +1,79 @@
+// RUN: %clangxx -fsycl -fsyntax-only %s
+
+// Tests that the reducer class is neither movable nor copyable.
+
+#include <sycl/sycl.hpp>
+
+#include <type_traits>
+
+template <class T> struct PlusWithoutIdentity {
+  T operator()(const T &A, const T &B) const { return A + B; }
+};
+
+template <typename ReducerT> static constexpr void checkReducer() {
+  static_assert(!std::is_copy_constructible_v<ReducerT>);
+  static_assert(!std::is_move_constructible_v<ReducerT>);
+  static_assert(!std::is_copy_assignable_v<ReducerT>);
+  static_assert(!std::is_move_assignable_v<ReducerT>);
+}
+
+int main() {
+  sycl::queue Q;
+
+  int *ScalarMem = sycl::malloc_shared<int>(1, Q);
+  int *SpanMem = sycl::malloc_shared<int>(8, Q);
+  auto ScalarRed1 = sycl::reduction(ScalarMem, std::plus<int>{});
+  auto ScalarRed2 = sycl::reduction(ScalarMem, PlusWithoutIdentity<int>{});
+  auto SpanRed1 =
+      sycl::reduction(sycl::span<int, 8>{SpanMem, 8}, std::plus<int>{});
+  auto SpanRed2 = sycl::reduction(sycl::span<int, 8>{SpanMem, 8},
+                                  PlusWithoutIdentity<int>{});
+
+  Q.parallel_for(sycl::range<1>{1024}, ScalarRed1,
+                 [=](sycl::item<1>, auto &Reducer) {
+                   checkReducer<std::remove_reference_t<decltype(Reducer)>>();
+                 });
+
+  Q.parallel_for(sycl::nd_range<1>{1024, 1024}, ScalarRed1,
+                 [=](sycl::nd_item<1>, auto &Reducer) {
+                   checkReducer<std::remove_reference_t<decltype(Reducer)>>();
+                 });
+
+  Q.parallel_for(sycl::range<1>{1024}, ScalarRed2,
+                 [=](sycl::item<1>, auto &Reducer) {
+                   checkReducer<std::remove_reference_t<decltype(Reducer)>>();
+                 });
+
+  Q.parallel_for(sycl::nd_range<1>{1024, 1024}, ScalarRed2,
+                 [=](sycl::nd_item<1>, auto &Reducer) {
+                   checkReducer<std::remove_reference_t<decltype(Reducer)>>();
+                 });
+
+  Q.parallel_for(
+      sycl::range<1>{1024}, SpanRed1, [=](sycl::item<1>, auto &Reducer) {
+        checkReducer<std::remove_reference_t<decltype(Reducer)>>();
+        checkReducer<std::remove_reference_t<decltype(Reducer[0])>>();
+      });
+
+  Q.parallel_for(
+      sycl::nd_range<1>{1024, 1024}, SpanRed1,
+      [=](sycl::nd_item<1>, auto &Reducer) {
+        checkReducer<std::remove_reference_t<decltype(Reducer)>>();
+        checkReducer<std::remove_reference_t<decltype(Reducer[0])>>();
+      });
+
+  Q.parallel_for(
+      sycl::range<1>{1024}, SpanRed2, [=](sycl::item<1>, auto &Reducer) {
+        checkReducer<std::remove_reference_t<decltype(Reducer)>>();
+        checkReducer<std::remove_reference_t<decltype(Reducer[0])>>();
+      });
+
+  Q.parallel_for(
+      sycl::nd_range<1>{1024, 1024}, SpanRed2,
+      [=](sycl::nd_item<1>, auto &Reducer) {
+        checkReducer<std::remove_reference_t<decltype(Reducer)>>();
+        checkReducer<std::remove_reference_t<decltype(Reducer[0])>>();
+      });
+
+  return 0;
+}