[SYCL][Reduction] Fix return type of reduction combine (#8125)

The combine member of reducers and the operator shortcuts should return
a reference to the reducer. This commit fixes this return type.

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

Author: steffenlarsen

sycl/include/sycl/reduction.hpp

@@ -179,46 +179,48 @@ template <class Reducer> class combiner {
 public:
   template <typename _T = Ty, int _Dims = Dims>
   enable_if_t<(_Dims == 0) && IsPlus<_T, BinaryOp>::value &&
-              is_geninteger<_T>::value>
+                  is_geninteger<_T>::value,
+              Reducer &>
   operator++() {
-    static_cast<Reducer *>(this)->combine(static_cast<_T>(1));
+    return static_cast<Reducer *>(this)->combine(static_cast<_T>(1));
   }
 
   template <typename _T = Ty, int _Dims = Dims>
   enable_if_t<(_Dims == 0) && IsPlus<_T, BinaryOp>::value &&
-              is_geninteger<_T>::value>
+                  is_geninteger<_T>::value,
+              Reducer &>
   operator++(int) {
-    static_cast<Reducer *>(this)->combine(static_cast<_T>(1));
+    return static_cast<Reducer *>(this)->combine(static_cast<_T>(1));
   }
 
   template <typename _T = Ty, int _Dims = Dims>
-  enable_if_t<(_Dims == 0) && IsPlus<_T, BinaryOp>::value>
+  enable_if_t<(_Dims == 0) && IsPlus<_T, BinaryOp>::value, Reducer &>
   operator+=(const _T &Partial) {
-    static_cast<Reducer *>(this)->combine(Partial);
+    return static_cast<Reducer *>(this)->combine(Partial);
   }
 
   template <typename _T = Ty, int _Dims = Dims>
-  enable_if_t<(_Dims == 0) && IsMultiplies<_T, BinaryOp>::value>
+  enable_if_t<(_Dims == 0) && IsMultiplies<_T, BinaryOp>::value, Reducer &>
   operator*=(const _T &Partial) {
-    static_cast<Reducer *>(this)->combine(Partial);
+    return static_cast<Reducer *>(this)->combine(Partial);
   }
 
   template <typename _T = Ty, int _Dims = Dims>
-  enable_if_t<(_Dims == 0) && IsBitOR<_T, BinaryOp>::value>
+  enable_if_t<(_Dims == 0) && IsBitOR<_T, BinaryOp>::value, Reducer &>
   operator|=(const _T &Partial) {
-    static_cast<Reducer *>(this)->combine(Partial);
+    return static_cast<Reducer *>(this)->combine(Partial);
   }
 
   template <typename _T = Ty, int _Dims = Dims>
-  enable_if_t<(_Dims == 0) && IsBitXOR<_T, BinaryOp>::value>
+  enable_if_t<(_Dims == 0) && IsBitXOR<_T, BinaryOp>::value, Reducer &>
   operator^=(const _T &Partial) {
-    static_cast<Reducer *>(this)->combine(Partial);
+    return static_cast<Reducer *>(this)->combine(Partial);
   }
 
   template <typename _T = Ty, int _Dims = Dims>
-  enable_if_t<(_Dims == 0) && IsBitAND<_T, BinaryOp>::value>
+  enable_if_t<(_Dims == 0) && IsBitAND<_T, BinaryOp>::value, Reducer &>
   operator&=(const _T &Partial) {
-    static_cast<Reducer *>(this)->combine(Partial);
+    return static_cast<Reducer *>(this)->combine(Partial);
   }
 
 private:
@@ -339,7 +341,10 @@ class reducer<
   reducer(const T &Identity, BinaryOperation BOp)
       : MValue(Identity), MIdentity(Identity), MBinaryOp(BOp) {}
 
-  void combine(const T &Partial) { MValue = MBinaryOp(MValue, Partial); }
+  reducer &combine(const T &Partial) {
+    MValue = MBinaryOp(MValue, Partial);
+    return *this;
+  }
 
   T getIdentity() const { return MIdentity; }
 
@@ -371,9 +376,10 @@ class reducer<
   reducer() : MValue(getIdentity()) {}
   reducer(const T & /* Identity */, BinaryOperation) : MValue(getIdentity()) {}
 
-  void combine(const T &Partial) {
+  reducer &combine(const T &Partial) {
     BinaryOperation BOp;
     MValue = BOp(MValue, Partial);
+    return *this;
   }
 
   static T getIdentity() {
@@ -396,7 +402,10 @@ class reducer<T, BinaryOperation, Dims, Extent, View,
 public:
   reducer(T &Ref, BinaryOperation BOp) : MElement(Ref), MBinaryOp(BOp) {}
 
-  void combine(const T &Partial) { MElement = MBinaryOp(MElement, Partial); }
+  reducer &combine(const T &Partial) {
+    MElement = MBinaryOp(MElement, Partial);
+    return *this;
+  }
 
 private:
   T &MElement;

sycl/test/basic_tests/reduction/reduction_combine_return_type.cpp

@@ -0,0 +1,42 @@
+// RUN: %clangxx -fsycl -fsyntax-only -sycl-std=2020 %s
+
+// Tests the return type of combiner operations on reducers.
+
+#include <sycl/sycl.hpp>
+
+#include <type_traits>
+
+int main() {
+  sycl::queue Q;
+
+  int *PlusMem = sycl::malloc_device<int>(1, Q);
+  int *MultMem = sycl::malloc_device<int>(1, Q);
+  int *BitAndMem = sycl::malloc_device<int>(1, Q);
+  int *BitOrMem = sycl::malloc_device<int>(1, Q);
+  int *BitXorMem = sycl::malloc_device<int>(1, Q);
+  Q.submit([&](sycl::handler &CGH) {
+    auto PlusReduction = sycl::reduction(PlusMem, sycl::plus<int>());
+    auto MultReduction = sycl::reduction(MultMem, sycl::multiplies<int>());
+    auto BitAndReduction = sycl::reduction(BitAndMem, sycl::bit_and<int>());
+    auto BitOrReduction = sycl::reduction(BitOrMem, sycl::bit_or<int>());
+    auto BitXorReduction = sycl::reduction(PlusMem, sycl::bit_xor<int>());
+    CGH.parallel_for(sycl::range<1>(10), PlusReduction, MultReduction,
+                     BitAndReduction, BitOrReduction, BitXorReduction,
+                     [=](sycl::id<1>, auto &Plus, auto &Mult, auto &BitAnd,
+                         auto &BitOr, auto &BitXor) {
+                       (Plus.combine(1) += 1).combine(1);
+                       (Plus.combine(1)++).combine(1);
+                       (++Plus.combine(1)).combine(1);
+                       (Mult.combine(1) *= 1).combine(1);
+                       (BitAnd.combine(1) &= 1).combine(1);
+                       (BitOr.combine(1) |= 1).combine(1);
+                       (BitXor.combine(1) ^= 1).combine(1);
+                     });
+  });
+  sycl::free(PlusMem, Q);
+  sycl::free(MultMem, Q);
+  sycl::free(BitAndMem, Q);
+  sycl::free(BitOrMem, Q);
+  sycl::free(BitXorMem, Q);
+  return 0;
+}