From 8c68366b30d5fea5d775353b5f8a14b22a5c6fe0 Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Wed, 13 Jul 2022 16:49:17 -0700
Subject: [PATCH 01/24] [NFC][SYCL] Remove Algorithm from reducer's template
 params

It isn't used for anything and the reducer is really lower level entity than the
reduction algorithm.
---
 sycl/include/sycl/ext/oneapi/reduction.hpp | 70 ++++++++++------------
 1 file changed, 30 insertions(+), 40 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/reduction.hpp b/sycl/include/sycl/ext/oneapi/reduction.hpp
index 3005dc3417944..a30db6672ae61 100644
--- a/sycl/include/sycl/ext/oneapi/reduction.hpp
+++ b/sycl/include/sycl/ext/oneapi/reduction.hpp
@@ -124,17 +124,17 @@ __SYCL_EXPORT size_t reduComputeWGSize(size_t NWorkItems, size_t MaxWGSize,
 /// functions and representing users' reduction variable.
 /// The generic version of the class represents those reductions of those
 /// types and operations for which the identity value is not known.
-/// The Algorithm template can be used to specialize reducers for different
-/// reduction algorithms. The View template describes whether the reducer
-/// owns its data or not: if View is 'true', then the reducer does not own
-/// its data and instead provides a view of data allocated elsewhere (i.e.
-/// via a reference or pointer member); if View is 'false', then the reducer
-/// owns its data. With the current default reduction algorithm, the top-level
-/// reducers that are passed to the user's lambda contain a private copy of
-/// the reduction variable, whereas any reducer created by a subscript operator
-/// contains a reference to a reduction variable allocated elsewhere.
+/// The View template describes whether the reducer owns its data or not: if
+/// View is 'true', then the reducer does not own its data and instead provides
+/// a view of data allocated elsewhere (i.e. via a reference or pointer member);
+/// if View is 'false', then the reducer owns its data. With the current default
+/// reduction algorithm, the top-level reducers that are passed to the user's
+/// lambda contain a private copy of the reduction variable, whereas any reducer
+/// created by a subscript operator contains a reference to a reduction variable
+/// allocated elsewhere. The Subst parameter is an implementation detail and is
+/// used to spell out restrictions using 'enable_if'.
 template <typename T, class BinaryOperation, int Dims, size_t Extent,
-          class Algorithm, bool View = false, typename Subst = void>
+          bool View = false, typename Subst = void>
 class reducer;
 
 /// Helper class for accessing reducer-defined types in CRTP
@@ -142,9 +142,8 @@ class reducer;
 template <typename Reducer> struct ReducerTraits;
 
 template <typename T, class BinaryOperation, int Dims, std::size_t Extent,
-          class Algorithm, bool View, typename Subst>
-struct ReducerTraits<
-    reducer<T, BinaryOperation, Dims, Extent, Algorithm, View, Subst>> {
+          bool View, typename Subst>
+struct ReducerTraits<reducer<T, BinaryOperation, Dims, Extent, View, Subst>> {
   using type = T;
   using op = BinaryOperation;
   static constexpr int dims = Dims;
@@ -319,14 +318,13 @@ template <class Reducer> class combiner {
 ///
 /// It stores a copy of the identity and binary operation associated with the
 /// reduction.
-template <typename T, class BinaryOperation, int Dims, size_t Extent,
-          class Algorithm, bool View>
+template <typename T, class BinaryOperation, int Dims, size_t Extent, bool View>
 class reducer<
-    T, BinaryOperation, Dims, Extent, Algorithm, View,
+    T, BinaryOperation, Dims, Extent, View,
     enable_if_t<Dims == 0 && Extent == 1 && View == false &&
                 !sycl::detail::IsKnownIdentityOp<T, BinaryOperation>::value>>
     : public combiner<
-          reducer<T, BinaryOperation, Dims, Extent, Algorithm, View,
+          reducer<T, BinaryOperation, Dims, Extent, View,
                   enable_if_t<Dims == 0 && Extent == 1 && View == false &&
                               !sycl::detail::IsKnownIdentityOp<
                                   T, BinaryOperation>::value>>> {
@@ -352,14 +350,13 @@ class reducer<
 ///
 /// It allows to reduce the size of the 'reducer' object by not holding
 /// the identity field inside it and allows to add a default constructor.
-template <typename T, class BinaryOperation, int Dims, size_t Extent,
-          class Algorithm, bool View>
+template <typename T, class BinaryOperation, int Dims, size_t Extent, bool View>
 class reducer<
-    T, BinaryOperation, Dims, Extent, Algorithm, View,
+    T, BinaryOperation, Dims, Extent, View,
     enable_if_t<Dims == 0 && Extent == 1 && View == false &&
                 sycl::detail::IsKnownIdentityOp<T, BinaryOperation>::value>>
     : public combiner<
-          reducer<T, BinaryOperation, Dims, Extent, Algorithm, View,
+          reducer<T, BinaryOperation, Dims, Extent, View,
                   enable_if_t<Dims == 0 && Extent == 1 && View == false &&
                               sycl::detail::IsKnownIdentityOp<
                                   T, BinaryOperation>::value>>> {
@@ -383,11 +380,10 @@ class reducer<
 
 /// Component of 'reducer' class for array reductions, representing a single
 /// element of the span (as returned by the subscript operator).
-template <typename T, class BinaryOperation, int Dims, size_t Extent,
-          class Algorithm, bool View>
-class reducer<T, BinaryOperation, Dims, Extent, Algorithm, View,
+template <typename T, class BinaryOperation, int Dims, size_t Extent, bool View>
+class reducer<T, BinaryOperation, Dims, Extent, View,
               enable_if_t<Dims == 0 && View == true>>
-    : public combiner<reducer<T, BinaryOperation, Dims, Extent, Algorithm, View,
+    : public combiner<reducer<T, BinaryOperation, Dims, Extent, View,
                               enable_if_t<Dims == 0 && View == true>>> {
 public:
   reducer(T &Ref, BinaryOperation BOp) : MElement(Ref), MBinaryOp(BOp) {}
@@ -401,13 +397,12 @@ class reducer<T, BinaryOperation, Dims, Extent, Algorithm, View,
 
 /// Specialization of 'reducer' class for array reductions exposing the
 /// subscript operator.
-template <typename T, class BinaryOperation, int Dims, size_t Extent,
-          class Algorithm, bool View>
+template <typename T, class BinaryOperation, int Dims, size_t Extent, bool View>
 class reducer<
-    T, BinaryOperation, Dims, Extent, Algorithm, View,
+    T, BinaryOperation, Dims, Extent, View,
     enable_if_t<Dims == 1 && View == false &&
                 !sycl::detail::IsKnownIdentityOp<T, BinaryOperation>::value>>
-    : public combiner<reducer<T, BinaryOperation, Dims, Extent, Algorithm, View,
+    : public combiner<reducer<T, BinaryOperation, Dims, Extent, View,
                               enable_if_t<Dims == 1 && View == false &&
                                           !sycl::detail::IsKnownIdentityOp<
                                               T, BinaryOperation>::value>>> {
@@ -415,8 +410,7 @@ class reducer<
   reducer(const T &Identity, BinaryOperation BOp)
       : MValue(Identity), MIdentity(Identity), MBinaryOp(BOp) {}
 
-  reducer<T, BinaryOperation, Dims - 1, Extent, Algorithm, true>
-  operator[](size_t Index) {
+  reducer<T, BinaryOperation, Dims - 1, Extent, true> operator[](size_t Index) {
     return {MValue[Index], MBinaryOp};
   }
 
@@ -432,13 +426,12 @@ class reducer<
 
 /// Specialization of 'reducer' class for array reductions accepting a span
 /// in cases where the identity value is known.
-template <typename T, class BinaryOperation, int Dims, size_t Extent,
-          class Algorithm, bool View>
+template <typename T, class BinaryOperation, int Dims, size_t Extent, bool View>
 class reducer<
-    T, BinaryOperation, Dims, Extent, Algorithm, View,
+    T, BinaryOperation, Dims, Extent, View,
     enable_if_t<Dims == 1 && View == false &&
                 sycl::detail::IsKnownIdentityOp<T, BinaryOperation>::value>>
-    : public combiner<reducer<T, BinaryOperation, Dims, Extent, Algorithm, View,
+    : public combiner<reducer<T, BinaryOperation, Dims, Extent, View,
                               enable_if_t<Dims == 1 && View == false &&
                                           sycl::detail::IsKnownIdentityOp<
                                               T, BinaryOperation>::value>>> {
@@ -448,8 +441,7 @@ class reducer<
 
   // SYCL 2020 revision 4 says this should be const, but this is a bug
   // see https://github.com/KhronosGroup/SYCL-Docs/pull/252
-  reducer<T, BinaryOperation, Dims - 1, Extent, Algorithm, true>
-  operator[](size_t Index) {
+  reducer<T, BinaryOperation, Dims - 1, Extent, true> operator[](size_t Index) {
     return {MValue[Index], BinaryOperation()};
   }
 
@@ -525,9 +517,7 @@ class reduction_impl_algo<
   using base = reduction_impl_common<T, BinaryOperation>;
 
 public:
-  using reducer_type =
-      reducer<T, BinaryOperation, Dims, Extent,
-              default_reduction_algorithm<IsUSM, IsPlaceholder, AccessorDims>>;
+  using reducer_type = reducer<T, BinaryOperation, Dims, Extent>;
   using result_type = T;
   using binary_operation = BinaryOperation;
 

From 417df4112391f9cd7eadb7f6220df53bd9a07c38 Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Thu, 14 Jul 2022 12:51:37 -0700
Subject: [PATCH 02/24] [WIP] SFINAE-out reduction_impl ctors

---
 sycl/include/sycl/ext/oneapi/reduction.hpp | 114 +++++++++++++++------
 sycl/include/sycl/handler.hpp              |   4 +-
 sycl/include/sycl/reduction.hpp            |  43 +++++---
 3 files changed, 109 insertions(+), 52 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/reduction.hpp b/sycl/include/sycl/ext/oneapi/reduction.hpp
index a30db6672ae61..74c62f8219f62 100644
--- a/sycl/include/sycl/ext/oneapi/reduction.hpp
+++ b/sycl/include/sycl/ext/oneapi/reduction.hpp
@@ -503,15 +503,16 @@ class default_reduction_algorithm {};
 
 /// Templated class for implementations of specific reduction algorithms
 template <typename T, class BinaryOperation, int Dims, size_t Extent,
-          class Algorithm>
+          typename RedOutVar, class Algorithm>
 class reduction_impl_algo;
 
 /// Original reduction algorithm is the default. It supports both USM and
 /// accessors via a single class
 template <typename T, class BinaryOperation, int Dims, size_t Extent,
-          bool IsUSM, access::placeholder IsPlaceholder, int AccessorDims>
+          typename RedOutVar, bool IsUSM, access::placeholder IsPlaceholder,
+          int AccessorDims>
 class reduction_impl_algo<
-    T, BinaryOperation, Dims, Extent,
+    T, BinaryOperation, Dims, Extent, RedOutVar,
     default_reduction_algorithm<IsUSM, IsPlaceholder, AccessorDims>>
     : public reduction_impl_common<T, BinaryOperation> {
   using base = reduction_impl_common<T, BinaryOperation>;
@@ -534,6 +535,7 @@ class reduction_impl_algo<
                access::target::device, IsPlaceholder,
                ext::oneapi::accessor_property_list<>>;
 
+
   static constexpr bool has_atomic_add_float64 =
       IsReduOptForAtomic64Add<T, BinaryOperation>::value;
   static constexpr bool has_fast_atomics =
@@ -705,6 +707,7 @@ class reduction_impl_algo<
     return Acc;
   }
 
+  RedOutVar *MRedOut;
   /// User's accessor to where the reduction must be written.
   std::shared_ptr<rw_accessor_type> MRWAcc;
   std::shared_ptr<dw_accessor_type> MDWAcc;
@@ -731,16 +734,38 @@ template <typename T> struct AreAllButLastReductions<T> {
   static constexpr bool value =
       !std::is_base_of<reduction_impl_base, std::remove_reference_t<T>>::value;
 };
+template <class T>
+struct is_rw_acc_t : public std::false_type {};
+
+template <class T, int AccessorDims, access::placeholder IsPlaceholder>
+struct is_rw_acc_t<
+    accessor<T, AccessorDims, access::mode::read_write, access::target::device,
+             IsPlaceholder, ext::oneapi::accessor_property_list<>>>
+    : public std::true_type {};
+
+template <class T>
+struct is_dw_acc_t : public std::false_type {};
+
+template <class T, int AccessorDims, access::placeholder IsPlaceholder>
+struct is_dw_acc_t<accessor<T, AccessorDims, access::mode::discard_write,
+                            access::target::device, IsPlaceholder,
+                            ext::oneapi::accessor_property_list<>>>
+    : public std::true_type {};
 
 /// This class encapsulates the reduction variable/accessor,
 /// the reduction operator and an optional operator identity.
 template <typename T, class BinaryOperation, int Dims, size_t Extent,
-          class Algorithm>
+          typename RedOutVar, class Algorithm>
 class reduction_impl
     : private reduction_impl_base,
-      public reduction_impl_algo<T, BinaryOperation, Dims, Extent, Algorithm> {
+      public reduction_impl_algo<T, BinaryOperation, Dims, Extent, RedOutVar,
+                                 Algorithm> {
 private:
-  using algo = reduction_impl_algo<T, BinaryOperation, Dims, Extent, Algorithm>;
+  using algo = reduction_impl_algo<T, BinaryOperation, Dims, Extent, RedOutVar,
+                                   Algorithm>;
+  static constexpr bool my_is_usm = std::is_same_v<RedOutVar, T *>;
+  static constexpr bool my_is_rw_acc = is_rw_acc_t<RedOutVar>::value;
+  static constexpr bool my_is_dw_acc = is_dw_acc_t<RedOutVar>::value;
 
 public:
   using reducer_type = typename algo::reducer_type;
@@ -753,8 +778,9 @@ class reduction_impl
   /// SYCL-2020.
   /// Constructs reduction_impl when the identity value is statically known.
   template <typename _T, typename AllocatorT,
-            std::enable_if_t<sycl::detail::IsKnownIdentityOp<
-                _T, BinaryOperation>::value> * = nullptr>
+            std::enable_if_t<
+                sycl::detail::IsKnownIdentityOp<_T, BinaryOperation>::value &&
+                my_is_rw_acc> * = nullptr>
   reduction_impl(buffer<_T, 1, AllocatorT> Buffer, handler &CGH,
                  bool InitializeToIdentity)
       : algo(reducer_type::getIdentity(), BinaryOperation(),
@@ -767,8 +793,10 @@ class reduction_impl
   }
 
   /// Constructs reduction_impl when the identity value is statically known.
-  template <typename _T = T, enable_if_t<sycl::detail::IsKnownIdentityOp<
-                                 _T, BinaryOperation>::value> * = nullptr>
+  template <
+      typename _T = T,
+      enable_if_t<sycl::detail::IsKnownIdentityOp<_T, BinaryOperation>::value &&
+                  my_is_rw_acc> * = nullptr>
   reduction_impl(rw_accessor_type &Acc)
       : algo(reducer_type::getIdentity(), BinaryOperation(), false,
              std::make_shared<rw_accessor_type>(Acc)) {
@@ -779,8 +807,10 @@ class reduction_impl
   }
 
   /// Constructs reduction_impl when the identity value is statically known.
-  template <typename _T = T, enable_if_t<sycl::detail::IsKnownIdentityOp<
-                                 _T, BinaryOperation>::value> * = nullptr>
+  template <
+      typename _T = T,
+      enable_if_t<sycl::detail::IsKnownIdentityOp<_T, BinaryOperation>::value &&
+                  my_is_dw_acc> * = nullptr>
   reduction_impl(dw_accessor_type &Acc)
       : algo(reducer_type::getIdentity(), BinaryOperation(), true,
              std::make_shared<dw_accessor_type>(Acc)) {
@@ -795,8 +825,8 @@ class reduction_impl
   /// and user still passed the identity value.
   template <
       typename _T, typename AllocatorT,
-      enable_if_t<sycl::detail::IsKnownIdentityOp<_T, BinaryOperation>::value>
-          * = nullptr>
+      enable_if_t<sycl::detail::IsKnownIdentityOp<_T, BinaryOperation>::value &&
+                  my_is_rw_acc> * = nullptr>
   reduction_impl(buffer<_T, 1, AllocatorT> Buffer, handler &CGH,
                  const T & /*Identity*/, BinaryOperation,
                  bool InitializeToIdentity)
@@ -822,8 +852,10 @@ class reduction_impl
 
   /// Constructs reduction_impl when the identity value is statically known,
   /// and user still passed the identity value.
-  template <typename _T = T, enable_if_t<sycl::detail::IsKnownIdentityOp<
-                                 _T, BinaryOperation>::value> * = nullptr>
+  template <
+      typename _T = T,
+      enable_if_t<sycl::detail::IsKnownIdentityOp<_T, BinaryOperation>::value &&
+                  my_is_rw_acc> * = nullptr>
   reduction_impl(rw_accessor_type &Acc, const T & /*Identity*/, BinaryOperation)
       : algo(reducer_type::getIdentity(), BinaryOperation(), false,
              std::make_shared<rw_accessor_type>(Acc)) {
@@ -846,8 +878,10 @@ class reduction_impl
 
   /// Constructs reduction_impl when the identity value is statically known,
   /// and user still passed the identity value.
-  template <typename _T = T, enable_if_t<sycl::detail::IsKnownIdentityOp<
-                                 _T, BinaryOperation>::value> * = nullptr>
+  template <
+      typename _T = T,
+      enable_if_t<sycl::detail::IsKnownIdentityOp<_T, BinaryOperation>::value &&
+                  my_is_dw_acc> * = nullptr>
   reduction_impl(dw_accessor_type &Acc, const T & /*Identity*/, BinaryOperation)
       : algo(reducer_type::getIdentity(), BinaryOperation(), true,
              std::make_shared<dw_accessor_type>(Acc)) {
@@ -870,10 +904,10 @@ class reduction_impl
 
   /// SYCL-2020.
   /// Constructs reduction_impl when the identity value is NOT known statically.
-  template <
-      typename _T, typename AllocatorT,
-      enable_if_t<!sycl::detail::IsKnownIdentityOp<_T, BinaryOperation>::value>
-          * = nullptr>
+  template <typename _T, typename AllocatorT,
+            enable_if_t<
+                !sycl::detail::IsKnownIdentityOp<_T, BinaryOperation>::value &&
+                my_is_rw_acc> * = nullptr>
   reduction_impl(buffer<_T, 1, AllocatorT> Buffer, handler &CGH,
                  const T &Identity, BinaryOperation BOp,
                  bool InitializeToIdentity)
@@ -888,7 +922,8 @@ class reduction_impl
 
   /// Constructs reduction_impl when the identity value is unknown.
   template <typename _T = T, enable_if_t<!sycl::detail::IsKnownIdentityOp<
-                                 _T, BinaryOperation>::value> * = nullptr>
+                                             _T, BinaryOperation>::value &&
+                                         my_is_rw_acc> * = nullptr>
   reduction_impl(rw_accessor_type &Acc, const T &Identity, BinaryOperation BOp)
       : algo(Identity, BOp, false, std::make_shared<rw_accessor_type>(Acc)) {
     if (Acc.size() != 1)
@@ -899,7 +934,8 @@ class reduction_impl
 
   /// Constructs reduction_impl when the identity value is unknown.
   template <typename _T = T, enable_if_t<!sycl::detail::IsKnownIdentityOp<
-                                 _T, BinaryOperation>::value> * = nullptr>
+                                             _T, BinaryOperation>::value &&
+                                         my_is_dw_acc> * = nullptr>
   reduction_impl(dw_accessor_type &Acc, const T &Identity, BinaryOperation BOp)
       : algo(Identity, BOp, true, std::make_shared<dw_accessor_type>(Acc)) {
     if (Acc.size() != 1)
@@ -912,8 +948,10 @@ class reduction_impl
   /// The \param VarPtr is a USM pointer to memory, to where the computed
   /// reduction value is added using BinaryOperation, i.e. it is expected that
   /// the memory is pre-initialized with some meaningful value.
-  template <typename _T = T, enable_if_t<sycl::detail::IsKnownIdentityOp<
-                                 _T, BinaryOperation>::value> * = nullptr>
+  template <
+      typename _T = T,
+      enable_if_t<sycl::detail::IsKnownIdentityOp<_T, BinaryOperation>::value &&
+                  my_is_usm> * = nullptr>
   reduction_impl(T *VarPtr, bool InitializeToIdentity = false)
       : algo(reducer_type::getIdentity(), BinaryOperation(),
              InitializeToIdentity, VarPtr) {}
@@ -923,8 +961,10 @@ class reduction_impl
   /// The \param VarPtr is a USM pointer to memory, to where the computed
   /// reduction value is added using BinaryOperation, i.e. it is expected that
   /// the memory is pre-initialized with some meaningful value.
-  template <typename _T = T, enable_if_t<sycl::detail::IsKnownIdentityOp<
-                                 _T, BinaryOperation>::value> * = nullptr>
+  template <
+      typename _T = T,
+      enable_if_t<sycl::detail::IsKnownIdentityOp<_T, BinaryOperation>::value &&
+                  my_is_usm> * = nullptr>
   reduction_impl(T *VarPtr, const T &Identity, BinaryOperation,
                  bool InitializeToIdentity = false)
       : algo(Identity, BinaryOperation(), InitializeToIdentity, VarPtr) {
@@ -946,7 +986,8 @@ class reduction_impl
   /// reduction value is added using BinaryOperation, i.e. it is expected that
   /// the memory is pre-initialized with some meaningful value.
   template <typename _T = T, enable_if_t<!sycl::detail::IsKnownIdentityOp<
-                                 _T, BinaryOperation>::value> * = nullptr>
+                                             _T, BinaryOperation>::value &&
+                                         my_is_usm> * = nullptr>
   reduction_impl(T *VarPtr, const T &Identity, BinaryOperation BOp,
                  bool InitializeToIdentity = false)
       : algo(Identity, BOp, InitializeToIdentity, VarPtr) {}
@@ -960,8 +1001,10 @@ class reduction_impl
 
   /// Constructs reduction_impl when the identity value is statically known
   /// and user passed an identity value anyway
-  template <typename _T = T, enable_if_t<sycl::detail::IsKnownIdentityOp<
-                                 _T, BinaryOperation>::value> * = nullptr>
+  template <
+      typename _T = T,
+      enable_if_t<sycl::detail::IsKnownIdentityOp<_T, BinaryOperation>::value &&
+                  my_is_usm> * = nullptr>
   reduction_impl(span<_T, Extent> Span, const T & /* Identity */,
                  BinaryOperation BOp, bool InitializeToIdentity = false)
       : algo(reducer_type::getIdentity(), BOp, InitializeToIdentity,
@@ -969,7 +1012,8 @@ class reduction_impl
 
   /// Constructs reduction_impl when the identity value is not statically known
   template <typename _T = T, enable_if_t<!sycl::detail::IsKnownIdentityOp<
-                                 _T, BinaryOperation>::value> * = nullptr>
+                                             _T, BinaryOperation>::value &&
+                                         my_is_usm> * = nullptr>
   reduction_impl(span<T, Extent> Span, const T &Identity, BinaryOperation BOp,
                  bool InitializeToIdentity = false)
       : algo(Identity, BOp, InitializeToIdentity, Span.data()) {}
@@ -2566,6 +2610,7 @@ tuple_select_elements(TupleT Tuple, std::index_sequence<Is...>) {
 template <typename T, class BinaryOperation, int Dims, access::mode AccMode,
           access::placeholder IsPH>
 detail::reduction_impl<T, BinaryOperation, 0, 1,
+                       accessor<T, Dims, AccMode, access::target::device, IsPH>,
                        detail::default_reduction_algorithm<false, IsPH, Dims>>
 reduction(accessor<T, Dims, AccMode, access::target::device, IsPH> &Acc,
           const T &Identity, BinaryOperation BOp) {
@@ -2581,6 +2626,7 @@ template <typename T, class BinaryOperation, int Dims, access::mode AccMode,
 std::enable_if_t<sycl::detail::IsKnownIdentityOp<T, BinaryOperation>::value,
                  detail::reduction_impl<
                      T, BinaryOperation, 0, 1,
+                     accessor<T, Dims, AccMode, access::target::device, IsPH>,
                      detail::default_reduction_algorithm<false, IsPH, Dims>>>
 reduction(accessor<T, Dims, AccMode, access::target::device, IsPH> &Acc,
           BinaryOperation) {
@@ -2593,7 +2639,7 @@ reduction(accessor<T, Dims, AccMode, access::target::device, IsPH> &Acc,
 /// \param Identity, and the binary operation used in the reduction.
 template <typename T, class BinaryOperation>
 detail::reduction_impl<
-    T, BinaryOperation, 0, 1,
+    T, BinaryOperation, 0, 1, T *,
     detail::default_reduction_algorithm<true, access::placeholder::false_t, 1>>
 reduction(T *VarPtr, const T &Identity, BinaryOperation BOp) {
   return {VarPtr, Identity, BOp};
@@ -2607,7 +2653,7 @@ reduction(T *VarPtr, const T &Identity, BinaryOperation BOp) {
 template <typename T, class BinaryOperation>
 std::enable_if_t<
     sycl::detail::IsKnownIdentityOp<T, BinaryOperation>::value,
-    detail::reduction_impl<T, BinaryOperation, 0, 1,
+    detail::reduction_impl<T, BinaryOperation, 0, 1, T *,
                            detail::default_reduction_algorithm<
                                true, access::placeholder::false_t, 1>>>
 reduction(T *VarPtr, BinaryOperation) {
diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp
index c6ee4bc17c675..20cf4b1f8c683 100644
--- a/sycl/include/sycl/handler.hpp
+++ b/sycl/include/sycl/handler.hpp
@@ -242,7 +242,7 @@ namespace ext {
 namespace oneapi {
 namespace detail {
 template <typename T, class BinaryOperation, int Dims, size_t Extent,
-          class Algorithm>
+          typename RedOutVar, class Algorithm>
 class reduction_impl_algo;
 
 using cl::sycl::detail::enable_if_t;
@@ -2667,7 +2667,7 @@ class __SYCL_EXPORT handler {
   // Make reduction friends to store buffers and arrays created for it
   // in handler from reduction methods.
   template <typename T, class BinaryOperation, int Dims, size_t Extent,
-            class Algorithm>
+            typename RedOutVar, class Algorithm>
   friend class ext::oneapi::detail::reduction_impl_algo;
 
 #ifndef __SYCL_DEVICE_ONLY__
diff --git a/sycl/include/sycl/reduction.hpp b/sycl/include/sycl/reduction.hpp
index 7416f38f31c6a..5e293ff298770 100644
--- a/sycl/include/sycl/reduction.hpp
+++ b/sycl/include/sycl/reduction.hpp
@@ -21,11 +21,15 @@ namespace sycl {
 /// Constructs a reduction object using the given buffer \p Var, handler \p CGH,
 /// reduction operation \p Combiner, and optional reduction properties.
 template <typename T, typename AllocatorT, typename BinaryOperation>
-std::enable_if_t<has_known_identity<BinaryOperation, T>::value,
-                 ext::oneapi::detail::reduction_impl<
-                     T, BinaryOperation, 0, 1,
-                     ext::oneapi::detail::default_reduction_algorithm<
-                         false, access::placeholder::true_t, 1>>>
+std::enable_if_t<
+    has_known_identity<BinaryOperation, T>::value,
+    ext::oneapi::detail::reduction_impl<
+        T, BinaryOperation, 0, 1,
+        accessor<T, 0, access::mode::read_write, access::target::device,
+                 access::placeholder::true_t,
+                 ext::oneapi::accessor_property_list<>>,
+        ext::oneapi::detail::default_reduction_algorithm<
+            false, access::placeholder::true_t, 1>>>
 reduction(buffer<T, 1, AllocatorT> Var, handler &CGH, BinaryOperation,
           const property_list &PropList = {}) {
   bool InitializeToIdentity =
@@ -38,11 +42,15 @@ reduction(buffer<T, 1, AllocatorT> Var, handler &CGH, BinaryOperation,
 /// The reduction algorithm may be less efficient for this variant as the
 /// reduction identity is not known statically and it is not provided by user.
 template <typename T, typename AllocatorT, typename BinaryOperation>
-std::enable_if_t<!has_known_identity<BinaryOperation, T>::value,
-                 ext::oneapi::detail::reduction_impl<
-                     T, BinaryOperation, 0, 1,
-                     ext::oneapi::detail::default_reduction_algorithm<
-                         false, access::placeholder::true_t, 1>>>
+std::enable_if_t<
+    !has_known_identity<BinaryOperation, T>::value,
+    ext::oneapi::detail::reduction_impl<
+        T, BinaryOperation, 0, 1,
+        accessor<T, 0, access::mode::read_write, access::target::device,
+                 access::placeholder::true_t,
+                 ext::oneapi::accessor_property_list<>>,
+        ext::oneapi::detail::default_reduction_algorithm<
+            false, access::placeholder::true_t, 1>>>
 reduction(buffer<T, 1, AllocatorT>, handler &, BinaryOperation,
           const property_list &PropList = {}) {
   // TODO: implement reduction that works even when identity is not known.
@@ -58,7 +66,7 @@ reduction(buffer<T, 1, AllocatorT>, handler &, BinaryOperation,
 template <typename T, typename BinaryOperation>
 std::enable_if_t<has_known_identity<BinaryOperation, T>::value,
                  ext::oneapi::detail::reduction_impl<
-                     T, BinaryOperation, 0, 1,
+                     T, BinaryOperation, 0, 1, T *,
                      ext::oneapi::detail::default_reduction_algorithm<
                          true, access::placeholder::false_t, 1>>>
 reduction(T *Var, BinaryOperation, const property_list &PropList = {}) {
@@ -75,7 +83,7 @@ reduction(T *Var, BinaryOperation, const property_list &PropList = {}) {
 template <typename T, typename BinaryOperation>
 std::enable_if_t<!has_known_identity<BinaryOperation, T>::value,
                  ext::oneapi::detail::reduction_impl<
-                     T, BinaryOperation, 0, 1,
+                     T, BinaryOperation, 0, 1, T *,
                      ext::oneapi::detail::default_reduction_algorithm<
                          true, access::placeholder::false_t, 1>>>
 reduction(T *, BinaryOperation, const property_list &PropList = {}) {
@@ -92,6 +100,9 @@ reduction(T *, BinaryOperation, const property_list &PropList = {}) {
 template <typename T, typename AllocatorT, typename BinaryOperation>
 ext::oneapi::detail::reduction_impl<
     T, BinaryOperation, 0, 1,
+    accessor<T, 0, access::mode::read_write, access::target::device,
+             access::placeholder::true_t,
+             ext::oneapi::accessor_property_list<>>,
     ext::oneapi::detail::default_reduction_algorithm<
         false, access::placeholder::true_t, 1>>
 reduction(buffer<T, 1, AllocatorT> Var, handler &CGH, const T &Identity,
@@ -106,7 +117,7 @@ reduction(buffer<T, 1, AllocatorT> Var, handler &CGH, const T &Identity,
 /// binary operation \p Combiner, and optional reduction properties.
 template <typename T, typename BinaryOperation>
 ext::oneapi::detail::reduction_impl<
-    T, BinaryOperation, 0, 1,
+    T, BinaryOperation, 0, 1, T *,
     ext::oneapi::detail::default_reduction_algorithm<
         true, access::placeholder::false_t, 1>>
 reduction(T *Var, const T &Identity, BinaryOperation Combiner,
@@ -124,7 +135,7 @@ template <typename T, size_t Extent, typename BinaryOperation>
 std::enable_if_t<Extent != dynamic_extent &&
                      has_known_identity<BinaryOperation, T>::value,
                  ext::oneapi::detail::reduction_impl<
-                     T, BinaryOperation, 1, Extent,
+                     T, BinaryOperation, 1, Extent, T *,
                      ext::oneapi::detail::default_reduction_algorithm<
                          true, access::placeholder::false_t, 1>>>
 reduction(span<T, Extent> Span, BinaryOperation,
@@ -143,7 +154,7 @@ template <typename T, size_t Extent, typename BinaryOperation>
 std::enable_if_t<Extent != dynamic_extent &&
                      !has_known_identity<BinaryOperation, T>::value,
                  ext::oneapi::detail::reduction_impl<
-                     T, BinaryOperation, 1, Extent,
+                     T, BinaryOperation, 1, Extent, T *,
                      ext::oneapi::detail::default_reduction_algorithm<
                          true, access::placeholder::false_t, 1>>>
 reduction(span<T, Extent>, BinaryOperation,
@@ -161,7 +172,7 @@ reduction(span<T, Extent>, BinaryOperation,
 template <typename T, size_t Extent, typename BinaryOperation>
 std::enable_if_t<Extent != dynamic_extent,
                  ext::oneapi::detail::reduction_impl<
-                     T, BinaryOperation, 1, Extent,
+                     T, BinaryOperation, 1, Extent, T *,
                      ext::oneapi::detail::default_reduction_algorithm<
                          true, access::placeholder::false_t, 1>>>
 reduction(span<T, Extent> Span, const T &Identity, BinaryOperation Combiner,

From a2718ef562d1c9e2895e1b6ce3038fc6319530fd Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Thu, 14 Jul 2022 13:06:45 -0700
Subject: [PATCH 03/24] Move my_is_* into base class

---
 sycl/include/sycl/ext/oneapi/reduction.hpp | 48 ++++++++++++----------
 1 file changed, 27 insertions(+), 21 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/reduction.hpp b/sycl/include/sycl/ext/oneapi/reduction.hpp
index 74c62f8219f62..5f6563136d649 100644
--- a/sycl/include/sycl/ext/oneapi/reduction.hpp
+++ b/sycl/include/sycl/ext/oneapi/reduction.hpp
@@ -496,6 +496,24 @@ template <typename T, class BinaryOperation> class reduction_impl_common {
   bool InitializeToIdentity;
 };
 
+template <class T>
+struct is_rw_acc_t : public std::false_type {};
+
+template <class T, int AccessorDims, access::placeholder IsPlaceholder>
+struct is_rw_acc_t<
+    accessor<T, AccessorDims, access::mode::read_write, access::target::device,
+             IsPlaceholder, ext::oneapi::accessor_property_list<>>>
+    : public std::true_type {};
+
+template <class T>
+struct is_dw_acc_t : public std::false_type {};
+
+template <class T, int AccessorDims, access::placeholder IsPlaceholder>
+struct is_dw_acc_t<accessor<T, AccessorDims, access::mode::discard_write,
+                            access::target::device, IsPlaceholder,
+                            ext::oneapi::accessor_property_list<>>>
+    : public std::true_type {};
+
 /// Types representing specific reduction algorithms
 /// Enables reduction_impl_algo to take additional algorithm-specific templates
 template <bool IsUSM, access::placeholder IsPlaceholder, int AccessorDims>
@@ -517,6 +535,11 @@ class reduction_impl_algo<
     : public reduction_impl_common<T, BinaryOperation> {
   using base = reduction_impl_common<T, BinaryOperation>;
 
+protected:
+  static constexpr bool my_is_usm = std::is_same_v<RedOutVar, T *>;
+  static constexpr bool my_is_rw_acc = is_rw_acc_t<RedOutVar>::value;
+  static constexpr bool my_is_dw_acc = is_dw_acc_t<RedOutVar>::value;
+
 public:
   using reducer_type = reducer<T, BinaryOperation, Dims, Extent>;
   using result_type = T;
@@ -734,24 +757,6 @@ template <typename T> struct AreAllButLastReductions<T> {
   static constexpr bool value =
       !std::is_base_of<reduction_impl_base, std::remove_reference_t<T>>::value;
 };
-template <class T>
-struct is_rw_acc_t : public std::false_type {};
-
-template <class T, int AccessorDims, access::placeholder IsPlaceholder>
-struct is_rw_acc_t<
-    accessor<T, AccessorDims, access::mode::read_write, access::target::device,
-             IsPlaceholder, ext::oneapi::accessor_property_list<>>>
-    : public std::true_type {};
-
-template <class T>
-struct is_dw_acc_t : public std::false_type {};
-
-template <class T, int AccessorDims, access::placeholder IsPlaceholder>
-struct is_dw_acc_t<accessor<T, AccessorDims, access::mode::discard_write,
-                            access::target::device, IsPlaceholder,
-                            ext::oneapi::accessor_property_list<>>>
-    : public std::true_type {};
-
 /// This class encapsulates the reduction variable/accessor,
 /// the reduction operator and an optional operator identity.
 template <typename T, class BinaryOperation, int Dims, size_t Extent,
@@ -763,9 +768,10 @@ class reduction_impl
 private:
   using algo = reduction_impl_algo<T, BinaryOperation, Dims, Extent, RedOutVar,
                                    Algorithm>;
-  static constexpr bool my_is_usm = std::is_same_v<RedOutVar, T *>;
-  static constexpr bool my_is_rw_acc = is_rw_acc_t<RedOutVar>::value;
-  static constexpr bool my_is_dw_acc = is_dw_acc_t<RedOutVar>::value;
+
+  using algo::my_is_usm;
+  using algo::my_is_rw_acc;
+  using algo::my_is_dw_acc;
 
 public:
   using reducer_type = typename algo::reducer_type;

From 7c686fced54a34211e3d1bb209563e8d07c948f8 Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Thu, 14 Jul 2022 13:21:39 -0700
Subject: [PATCH 04/24] SFINAE-out base class ctors as well

---
 sycl/include/sycl/ext/oneapi/reduction.hpp | 8 ++++++++
 1 file changed, 8 insertions(+)

diff --git a/sycl/include/sycl/ext/oneapi/reduction.hpp b/sycl/include/sycl/ext/oneapi/reduction.hpp
index 5f6563136d649..da8e661b96486 100644
--- a/sycl/include/sycl/ext/oneapi/reduction.hpp
+++ b/sycl/include/sycl/ext/oneapi/reduction.hpp
@@ -534,6 +534,9 @@ class reduction_impl_algo<
     default_reduction_algorithm<IsUSM, IsPlaceholder, AccessorDims>>
     : public reduction_impl_common<T, BinaryOperation> {
   using base = reduction_impl_common<T, BinaryOperation>;
+  using self = reduction_impl_algo<
+      T, BinaryOperation, Dims, Extent, RedOutVar,
+      default_reduction_algorithm<IsUSM, IsPlaceholder, AccessorDims>>;
 
 protected:
   static constexpr bool my_is_usm = std::is_same_v<RedOutVar, T *>;
@@ -572,12 +575,17 @@ class reduction_impl_algo<
   static constexpr size_t dims = Dims;
   static constexpr size_t num_elements = Extent;
 
+  template <class _self = self,
+            std::enable_if_t<_self::my_is_rw_acc> * = nullptr>
   reduction_impl_algo(const T &Identity, BinaryOperation BinaryOp, bool Init,
                       std::shared_ptr<rw_accessor_type> AccPointer)
       : base(Identity, BinaryOp, Init), MRWAcc(AccPointer){};
+  template <class _self = self,
+            std::enable_if_t<_self::my_is_dw_acc> * = nullptr>
   reduction_impl_algo(const T &Identity, BinaryOperation BinaryOp, bool Init,
                       std::shared_ptr<dw_accessor_type> AccPointer)
       : base(Identity, BinaryOp, Init), MDWAcc(AccPointer){};
+  template <class _self = self, std::enable_if_t<_self::my_is_usm> * = nullptr>
   reduction_impl_algo(const T &Identity, BinaryOperation BinaryOp, bool Init,
                       T *USMPointer)
       : base(Identity, BinaryOp, Init), MUSMPointer(USMPointer){};

From 3b1fd02a94f1073b753148149ae128260fedb0cc Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Thu, 14 Jul 2022 13:49:02 -0700
Subject: [PATCH 05/24] Init single reference to user reduction var

---
 sycl/include/sycl/ext/oneapi/reduction.hpp | 12 ++++++++----
 sycl/include/sycl/reduction.hpp            |  6 +++---
 2 files changed, 11 insertions(+), 7 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/reduction.hpp b/sycl/include/sycl/ext/oneapi/reduction.hpp
index da8e661b96486..0cd8333d6ccdd 100644
--- a/sycl/include/sycl/ext/oneapi/reduction.hpp
+++ b/sycl/include/sycl/ext/oneapi/reduction.hpp
@@ -579,16 +579,19 @@ class reduction_impl_algo<
             std::enable_if_t<_self::my_is_rw_acc> * = nullptr>
   reduction_impl_algo(const T &Identity, BinaryOperation BinaryOp, bool Init,
                       std::shared_ptr<rw_accessor_type> AccPointer)
-      : base(Identity, BinaryOp, Init), MRWAcc(AccPointer){};
+      : base(Identity, BinaryOp, Init), MRWAcc(AccPointer),
+        MRedOut(*AccPointer){};
   template <class _self = self,
             std::enable_if_t<_self::my_is_dw_acc> * = nullptr>
   reduction_impl_algo(const T &Identity, BinaryOperation BinaryOp, bool Init,
                       std::shared_ptr<dw_accessor_type> AccPointer)
-      : base(Identity, BinaryOp, Init), MDWAcc(AccPointer){};
+      : base(Identity, BinaryOp, Init), MDWAcc(AccPointer),
+        MRedOut(*AccPointer){};
   template <class _self = self, std::enable_if_t<_self::my_is_usm> * = nullptr>
   reduction_impl_algo(const T &Identity, BinaryOperation BinaryOp, bool Init,
                       T *USMPointer)
-      : base(Identity, BinaryOp, Init), MUSMPointer(USMPointer){};
+      : base(Identity, BinaryOp, Init), MUSMPointer(USMPointer),
+        MRedOut(MUSMPointer){};
 
   /// Associates the reduction accessor to user's memory with \p CGH handler
   /// to keep the accessor alive until the command group finishes the work.
@@ -738,7 +741,6 @@ class reduction_impl_algo<
     return Acc;
   }
 
-  RedOutVar *MRedOut;
   /// User's accessor to where the reduction must be written.
   std::shared_ptr<rw_accessor_type> MRWAcc;
   std::shared_ptr<dw_accessor_type> MDWAcc;
@@ -748,6 +750,8 @@ class reduction_impl_algo<
   /// USM pointer referencing the memory to where the result of the reduction
   /// must be written. Applicable/used only for USM reductions.
   T *MUSMPointer = nullptr;
+
+  RedOutVar &MRedOut;
 };
 
 /// Predicate returning true if all template type parameters except the last one
diff --git a/sycl/include/sycl/reduction.hpp b/sycl/include/sycl/reduction.hpp
index 5e293ff298770..bb8cb269980c6 100644
--- a/sycl/include/sycl/reduction.hpp
+++ b/sycl/include/sycl/reduction.hpp
@@ -25,7 +25,7 @@ std::enable_if_t<
     has_known_identity<BinaryOperation, T>::value,
     ext::oneapi::detail::reduction_impl<
         T, BinaryOperation, 0, 1,
-        accessor<T, 0, access::mode::read_write, access::target::device,
+        accessor<T, 1, access::mode::read_write, access::target::device,
                  access::placeholder::true_t,
                  ext::oneapi::accessor_property_list<>>,
         ext::oneapi::detail::default_reduction_algorithm<
@@ -46,7 +46,7 @@ std::enable_if_t<
     !has_known_identity<BinaryOperation, T>::value,
     ext::oneapi::detail::reduction_impl<
         T, BinaryOperation, 0, 1,
-        accessor<T, 0, access::mode::read_write, access::target::device,
+        accessor<T, 1, access::mode::read_write, access::target::device,
                  access::placeholder::true_t,
                  ext::oneapi::accessor_property_list<>>,
         ext::oneapi::detail::default_reduction_algorithm<
@@ -100,7 +100,7 @@ reduction(T *, BinaryOperation, const property_list &PropList = {}) {
 template <typename T, typename AllocatorT, typename BinaryOperation>
 ext::oneapi::detail::reduction_impl<
     T, BinaryOperation, 0, 1,
-    accessor<T, 0, access::mode::read_write, access::target::device,
+    accessor<T, 1, access::mode::read_write, access::target::device,
              access::placeholder::true_t,
              ext::oneapi::accessor_property_list<>>,
     ext::oneapi::detail::default_reduction_algorithm<

From edc843a784d9c97a73c15828e31e07786d77e98c Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Thu, 14 Jul 2022 14:14:33 -0700
Subject: [PATCH 06/24] Switch uses to generic/template MRedOut

---
 sycl/include/sycl/ext/oneapi/reduction.hpp | 93 +++++++++-------------
 sycl/include/sycl/handler.hpp              |  4 +-
 2 files changed, 41 insertions(+), 56 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/reduction.hpp b/sycl/include/sycl/ext/oneapi/reduction.hpp
index 0cd8333d6ccdd..a8a8bd2dd493f 100644
--- a/sycl/include/sycl/ext/oneapi/reduction.hpp
+++ b/sycl/include/sycl/ext/oneapi/reduction.hpp
@@ -538,11 +538,6 @@ class reduction_impl_algo<
       T, BinaryOperation, Dims, Extent, RedOutVar,
       default_reduction_algorithm<IsUSM, IsPlaceholder, AccessorDims>>;
 
-protected:
-  static constexpr bool my_is_usm = std::is_same_v<RedOutVar, T *>;
-  static constexpr bool my_is_rw_acc = is_rw_acc_t<RedOutVar>::value;
-  static constexpr bool my_is_dw_acc = is_dw_acc_t<RedOutVar>::value;
-
 public:
   using reducer_type = reducer<T, BinaryOperation, Dims, Extent>;
   using result_type = T;
@@ -568,7 +563,14 @@ class reduction_impl_algo<
       IsReduOptForFastAtomicFetch<T, BinaryOperation>::value;
   static constexpr bool has_fast_reduce =
       IsReduOptForFastReduce<T, BinaryOperation>::value;
+
   static constexpr bool is_usm = IsUSM;
+  static constexpr bool my_is_usm = std::is_same_v<RedOutVar, T *>;
+  static_assert(is_usm == my_is_usm);
+
+  static constexpr bool my_is_rw_acc = is_rw_acc_t<RedOutVar>::value;
+  static constexpr bool my_is_dw_acc = is_dw_acc_t<RedOutVar>::value;
+
   static constexpr bool is_placeholder =
       (IsPlaceholder == access::placeholder::true_t);
 
@@ -597,10 +599,10 @@ class reduction_impl_algo<
   /// to keep the accessor alive until the command group finishes the work.
   /// This function does not do anything for USM reductions.
   void associateWithHandler(handler &CGH) {
-    if (MRWAcc)
-      CGH.associateWithHandler(MRWAcc.get(), access::target::device);
-    else if (MDWAcc)
-      CGH.associateWithHandler(MDWAcc.get(), access::target::device);
+    if constexpr (!my_is_usm) {
+      static_assert(my_is_rw_acc || my_is_dw_acc);
+      CGH.associateWithHandler(&MRedOut, access::target::device);
+    }
   }
 
   /// Creates and returns a local accessor with the \p Size elements.
@@ -624,7 +626,7 @@ class reduction_impl_algo<
   template <bool IsOneWG, bool _IsUSM = is_usm>
   std::enable_if_t<IsOneWG && _IsUSM, result_type *>
   getWriteMemForPartialReds(size_t, handler &) {
-    return getUSMPointer();
+    return getUserRedVar();
   }
 
   /// Returns user's accessor passed to reduction for editing if that is
@@ -633,8 +635,8 @@ class reduction_impl_algo<
   template <bool IsOneWG, bool _IsUSM = is_usm>
   std::enable_if_t<IsOneWG && !_IsUSM, rw_accessor_type>
   getWriteMemForPartialReds(size_t, handler &CGH) {
-    if (MRWAcc)
-      return *MRWAcc;
+    if constexpr (my_is_rw_acc)
+      return MRedOut;
     return getWriteMemForPartialReds<false>(1, CGH);
   }
 
@@ -655,9 +657,11 @@ class reduction_impl_algo<
   /// Otherwise, a new buffer is created and accessor to that buffer is
   /// returned.
   rw_accessor_type getWriteAccForPartialReds(size_t Size, handler &CGH) {
-    if (Size == 1 && MRWAcc != nullptr) {
-      associateWithHandler(CGH);
-      return *MRWAcc;
+    if constexpr (my_is_rw_acc) {
+      if (Size == 1) {
+        associateWithHandler(CGH);
+        return MRedOut;
+      }
     }
 
     // Create a new output buffer and return an accessor to it.
@@ -673,8 +677,12 @@ class reduction_impl_algo<
   template <bool HasFastAtomics = (has_fast_atomics || has_atomic_add_float64)>
   std::enable_if_t<HasFastAtomics, rw_accessor_type>
   getReadWriteAccessorToInitializedMem(handler &CGH) {
-    if (!is_usm && !base::initializeToIdentity())
-      return *MRWAcc;
+    if constexpr (my_is_rw_acc) {
+      if (!base::initializeToIdentity())
+        return MRedOut;
+    }
+    assert(!(my_is_dw_acc && !base::initializeToIdentity()) &&
+           "Unexpected condition!");
 
     // TODO: Move to T[] in C++20 to simplify handling here
     // auto RWReduVal = std::make_shared<T[num_elements]>();
@@ -701,22 +709,7 @@ class reduction_impl_algo<
     return {*CounterBuf, CGH};
   }
 
-  bool hasUserDiscardWriteAccessor() { return MDWAcc != nullptr; }
-
-  template <bool _IsUSM = IsUSM>
-  std::enable_if_t<!_IsUSM, rw_accessor_type &> getUserReadWriteAccessor() {
-    return *MRWAcc;
-  }
-
-  template <bool _IsUSM = IsUSM>
-  std::enable_if_t<!_IsUSM, dw_accessor_type &> getUserDiscardWriteAccessor() {
-    return *MDWAcc;
-  }
-
-  result_type *getUSMPointer() {
-    assert(is_usm && "Unexpected call of getUSMPointer().");
-    return MUSMPointer;
-  }
+  RedOutVar &getUserRedVar() { return MRedOut; }
 
   static inline result_type *getOutPointer(const rw_accessor_type &OutAcc) {
     return OutAcc.get_pointer().get();
@@ -781,11 +774,11 @@ class reduction_impl
   using algo = reduction_impl_algo<T, BinaryOperation, Dims, Extent, RedOutVar,
                                    Algorithm>;
 
+public:
   using algo::my_is_usm;
   using algo::my_is_rw_acc;
   using algo::my_is_dw_acc;
 
-public:
   using reducer_type = typename algo::reducer_type;
   using rw_accessor_type = typename algo::rw_accessor_type;
   using dw_accessor_type = typename algo::dw_accessor_type;
@@ -1746,10 +1739,7 @@ std::enable_if_t<!Reduction::is_usm>
 reduSaveFinalResultToUserMem(handler &CGH, Reduction &Redu) {
   auto InAcc = Redu.getReadAccToPreviousPartialReds(CGH);
   Redu.associateWithHandler(CGH);
-  if (Redu.hasUserDiscardWriteAccessor())
-    CGH.copy(InAcc, Redu.getUserDiscardWriteAccessor());
-  else
-    CGH.copy(InAcc, Redu.getUserReadWriteAccessor());
+  CGH.copy(InAcc, Redu.getUserRedVar());
 }
 
 // This method is used for implementation of parallel_for accepting 1 reduction.
@@ -1762,7 +1752,7 @@ std::enable_if_t<Reduction::is_usm>
 reduSaveFinalResultToUserMem(handler &CGH, Reduction &Redu) {
   constexpr size_t NElements = Reduction::num_elements;
   auto InAcc = Redu.getReadAccToPreviousPartialReds(CGH);
-  auto UserVarPtr = Redu.getUSMPointer();
+  auto UserVarPtr = Redu.getUserRedVar();
   bool IsUpdateOfUserVar = !Redu.initializeToIdentity();
   auto BOp = Redu.getBinaryOperation();
   CGH.single_task<KernelName>([=] {
@@ -2559,21 +2549,16 @@ template <typename Reduction, typename... RestT>
 void reduSaveFinalResultToUserMemHelper(
     std::vector<event> &Events, std::shared_ptr<detail::queue_impl> Queue,
     bool IsHost, Reduction &Redu, RestT... Rest) {
-  // Reductions initialized with USM pointer currently do not require copying
-  // because the last kernel writes directly to the USM memory.
-  if constexpr (!Reduction::is_usm) {
-    if (Redu.hasUserDiscardWriteAccessor()) {
-      event CopyEvent =
-          withAuxHandler(Queue, IsHost, [&](handler &CopyHandler) {
-            auto InAcc = Redu.getReadAccToPreviousPartialReds(CopyHandler);
-            auto OutAcc = Redu.getUserDiscardWriteAccessor();
-            Redu.associateWithHandler(CopyHandler);
-            if (!Events.empty())
-              CopyHandler.depends_on(Events.back());
-            CopyHandler.copy(InAcc, OutAcc);
-          });
-      Events.push_back(CopyEvent);
-    }
+  if constexpr (Reduction::my_is_dw_acc) {
+    event CopyEvent = withAuxHandler(Queue, IsHost, [&](handler &CopyHandler) {
+      auto InAcc = Redu.getReadAccToPreviousPartialReds(CopyHandler);
+      auto OutAcc = Redu.getUserRedVar();
+      Redu.associateWithHandler(CopyHandler);
+      if (!Events.empty())
+        CopyHandler.depends_on(Events.back());
+      CopyHandler.copy(InAcc, OutAcc);
+    });
+    Events.push_back(CopyEvent);
   }
   reduSaveFinalResultToUserMemHelper(Events, Queue, IsHost, Rest...);
 }
diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp
index 20cf4b1f8c683..8d99605b699df 100644
--- a/sycl/include/sycl/handler.hpp
+++ b/sycl/include/sycl/handler.hpp
@@ -1655,7 +1655,7 @@ class __SYCL_EXPORT handler {
         *this, KernelFunc, Range, MaxWGSize, NumConcurrentWorkGroups, Redu);
     if (Reduction::is_usm ||
         (Reduction::has_fast_atomics && Redu.initializeToIdentity()) ||
-        (!Reduction::has_fast_atomics && Redu.hasUserDiscardWriteAccessor())) {
+        (!Reduction::has_fast_atomics && Reduction::my_is_dw_acc)) {
       this->finalize();
       MLastEvent = withAuxHandler(QueueCopy, [&](handler &CopyHandler) {
         ext::oneapi::detail::reduSaveFinalResultToUserMem<KernelName>(
@@ -1782,7 +1782,7 @@ class __SYCL_EXPORT handler {
       });
     } // end while (NWorkItems > 1)
 
-    if (Reduction::is_usm || Redu.hasUserDiscardWriteAccessor()) {
+    if (Reduction::is_usm || Reduction::my_is_dw_acc) {
       MLastEvent = withAuxHandler(QueueCopy, [&](handler &CopyHandler) {
         ext::oneapi::detail::reduSaveFinalResultToUserMem<KernelName>(
             CopyHandler, Redu);

From 445e8719047b57522b126cb86aef9c086a6000b6 Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Thu, 14 Jul 2022 14:26:32 -0700
Subject: [PATCH 07/24] Make MRedOut the *only* version of the user's reduction
 var

In other words, remove separate MRWAcc/MDWAcc/MUSMPointer.
---
 sycl/include/sycl/ext/oneapi/reduction.hpp | 108 ++++-----------------
 1 file changed, 17 insertions(+), 91 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/reduction.hpp b/sycl/include/sycl/ext/oneapi/reduction.hpp
index a8a8bd2dd493f..8e07bd5506ead 100644
--- a/sycl/include/sycl/ext/oneapi/reduction.hpp
+++ b/sycl/include/sycl/ext/oneapi/reduction.hpp
@@ -577,23 +577,9 @@ class reduction_impl_algo<
   static constexpr size_t dims = Dims;
   static constexpr size_t num_elements = Extent;
 
-  template <class _self = self,
-            std::enable_if_t<_self::my_is_rw_acc> * = nullptr>
   reduction_impl_algo(const T &Identity, BinaryOperation BinaryOp, bool Init,
-                      std::shared_ptr<rw_accessor_type> AccPointer)
-      : base(Identity, BinaryOp, Init), MRWAcc(AccPointer),
-        MRedOut(*AccPointer){};
-  template <class _self = self,
-            std::enable_if_t<_self::my_is_dw_acc> * = nullptr>
-  reduction_impl_algo(const T &Identity, BinaryOperation BinaryOp, bool Init,
-                      std::shared_ptr<dw_accessor_type> AccPointer)
-      : base(Identity, BinaryOp, Init), MDWAcc(AccPointer),
-        MRedOut(*AccPointer){};
-  template <class _self = self, std::enable_if_t<_self::my_is_usm> * = nullptr>
-  reduction_impl_algo(const T &Identity, BinaryOperation BinaryOp, bool Init,
-                      T *USMPointer)
-      : base(Identity, BinaryOp, Init), MUSMPointer(USMPointer),
-        MRedOut(MUSMPointer){};
+                      RedOutVar RedOut)
+      : base(Identity, BinaryOp, Init), MRedOut(std::move(RedOut)){};
 
   /// Associates the reduction accessor to user's memory with \p CGH handler
   /// to keep the accessor alive until the command group finishes the work.
@@ -734,17 +720,10 @@ class reduction_impl_algo<
     return Acc;
   }
 
-  /// User's accessor to where the reduction must be written.
-  std::shared_ptr<rw_accessor_type> MRWAcc;
-  std::shared_ptr<dw_accessor_type> MDWAcc;
-
   std::shared_ptr<buffer<T, buffer_dim>> MOutBufPtr;
 
-  /// USM pointer referencing the memory to where the result of the reduction
-  /// must be written. Applicable/used only for USM reductions.
-  T *MUSMPointer = nullptr;
-
-  RedOutVar &MRedOut;
+  /// User's accessor/USM pointer to where the reduction must be written.
+  RedOutVar MRedOut;
 };
 
 /// Predicate returning true if all template type parameters except the last one
@@ -795,7 +774,7 @@ class reduction_impl
   reduction_impl(buffer<_T, 1, AllocatorT> Buffer, handler &CGH,
                  bool InitializeToIdentity)
       : algo(reducer_type::getIdentity(), BinaryOperation(),
-             InitializeToIdentity, std::make_shared<rw_accessor_type>(Buffer)) {
+             InitializeToIdentity, rw_accessor_type{Buffer}) {
     algo::associateWithHandler(CGH);
     if (Buffer.size() != 1)
       throw sycl::runtime_error(errc::invalid,
@@ -807,24 +786,9 @@ class reduction_impl
   template <
       typename _T = T,
       enable_if_t<sycl::detail::IsKnownIdentityOp<_T, BinaryOperation>::value &&
-                  my_is_rw_acc> * = nullptr>
-  reduction_impl(rw_accessor_type &Acc)
-      : algo(reducer_type::getIdentity(), BinaryOperation(), false,
-             std::make_shared<rw_accessor_type>(Acc)) {
-    if (Acc.size() != 1)
-      throw sycl::runtime_error(errc::invalid,
-                                "Reduction variable must be a scalar.",
-                                PI_ERROR_INVALID_VALUE);
-  }
-
-  /// Constructs reduction_impl when the identity value is statically known.
-  template <
-      typename _T = T,
-      enable_if_t<sycl::detail::IsKnownIdentityOp<_T, BinaryOperation>::value &&
-                  my_is_dw_acc> * = nullptr>
-  reduction_impl(dw_accessor_type &Acc)
-      : algo(reducer_type::getIdentity(), BinaryOperation(), true,
-             std::make_shared<dw_accessor_type>(Acc)) {
+                  (my_is_rw_acc || my_is_dw_acc)> * = nullptr>
+  reduction_impl(RedOutVar &Acc)
+      : algo(reducer_type::getIdentity(), BinaryOperation(), false, Acc) {
     if (Acc.size() != 1)
       throw sycl::runtime_error(errc::invalid,
                                 "Reduction variable must be a scalar.",
@@ -842,7 +806,7 @@ class reduction_impl
                  const T & /*Identity*/, BinaryOperation,
                  bool InitializeToIdentity)
       : algo(reducer_type::getIdentity(), BinaryOperation(),
-             InitializeToIdentity, std::make_shared<rw_accessor_type>(Buffer)) {
+             InitializeToIdentity, rw_accessor_type{Buffer}) {
     algo::associateWithHandler(CGH);
     if (Buffer.size() != 1)
       throw sycl::runtime_error(errc::invalid,
@@ -866,36 +830,10 @@ class reduction_impl
   template <
       typename _T = T,
       enable_if_t<sycl::detail::IsKnownIdentityOp<_T, BinaryOperation>::value &&
-                  my_is_rw_acc> * = nullptr>
-  reduction_impl(rw_accessor_type &Acc, const T & /*Identity*/, BinaryOperation)
-      : algo(reducer_type::getIdentity(), BinaryOperation(), false,
-             std::make_shared<rw_accessor_type>(Acc)) {
-    if (Acc.size() != 1)
-      throw sycl::runtime_error(errc::invalid,
-                                "Reduction variable must be a scalar.",
-                                PI_ERROR_INVALID_VALUE);
-    // For now the implementation ignores the identity value given by user
-    // when the implementation knows the identity.
-    // The SPEC could prohibit passing identity parameter to operations with
-    // known identity, but that could have some bad consequences too.
-    // For example, at some moment the implementation may NOT know the identity
-    // for COMPLEX-PLUS reduction. User may create a program that would pass
-    // COMPLEX value (0,0) as identity for PLUS reduction. At some later moment
-    // when the implementation starts handling COMPLEX-PLUS as known operation
-    // the existing user's program remains compilable and working correctly.
-    // I.e. with this constructor here, adding more reduction operations to the
-    // list of known operations does not break the existing programs.
-  }
-
-  /// Constructs reduction_impl when the identity value is statically known,
-  /// and user still passed the identity value.
-  template <
-      typename _T = T,
-      enable_if_t<sycl::detail::IsKnownIdentityOp<_T, BinaryOperation>::value &&
-                  my_is_dw_acc> * = nullptr>
-  reduction_impl(dw_accessor_type &Acc, const T & /*Identity*/, BinaryOperation)
-      : algo(reducer_type::getIdentity(), BinaryOperation(), true,
-             std::make_shared<dw_accessor_type>(Acc)) {
+                  (my_is_rw_acc || my_is_dw_acc)> * = nullptr>
+  reduction_impl(RedOutVar &Acc, const T & /*Identity*/, BinaryOperation)
+      : algo(reducer_type::getIdentity(), BinaryOperation(), my_is_dw_acc,
+             Acc) {
     if (Acc.size() != 1)
       throw sycl::runtime_error(errc::invalid,
                                 "Reduction variable must be a scalar.",
@@ -923,7 +861,7 @@ class reduction_impl
                  const T &Identity, BinaryOperation BOp,
                  bool InitializeToIdentity)
       : algo(Identity, BOp, InitializeToIdentity,
-             std::make_shared<rw_accessor_type>(Buffer)) {
+             rw_accessor_type{Buffer}) {
     algo::associateWithHandler(CGH);
     if (Buffer.size() != 1)
       throw sycl::runtime_error(errc::invalid,
@@ -934,21 +872,9 @@ class reduction_impl
   /// Constructs reduction_impl when the identity value is unknown.
   template <typename _T = T, enable_if_t<!sycl::detail::IsKnownIdentityOp<
                                              _T, BinaryOperation>::value &&
-                                         my_is_rw_acc> * = nullptr>
-  reduction_impl(rw_accessor_type &Acc, const T &Identity, BinaryOperation BOp)
-      : algo(Identity, BOp, false, std::make_shared<rw_accessor_type>(Acc)) {
-    if (Acc.size() != 1)
-      throw sycl::runtime_error(errc::invalid,
-                                "Reduction variable must be a scalar.",
-                                PI_ERROR_INVALID_VALUE);
-  }
-
-  /// Constructs reduction_impl when the identity value is unknown.
-  template <typename _T = T, enable_if_t<!sycl::detail::IsKnownIdentityOp<
-                                             _T, BinaryOperation>::value &&
-                                         my_is_dw_acc> * = nullptr>
-  reduction_impl(dw_accessor_type &Acc, const T &Identity, BinaryOperation BOp)
-      : algo(Identity, BOp, true, std::make_shared<dw_accessor_type>(Acc)) {
+                                         (my_is_rw_acc || my_is_dw_acc)> * = nullptr>
+  reduction_impl(RedOutVar &Acc, const T &Identity, BinaryOperation BOp)
+      : algo(Identity, BOp, my_is_dw_acc, Acc) {
     if (Acc.size() != 1)
       throw sycl::runtime_error(errc::invalid,
                                 "Reduction variable must be a scalar.",

From 3a4d91d62cef0a4230b2d1932b39238f65ed54a9 Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Thu, 14 Jul 2022 15:34:55 -0700
Subject: [PATCH 08/24] Simplification of overloads with/without known identity

---
 sycl/include/sycl/ext/oneapi/reduction.hpp | 175 ++++++---------------
 1 file changed, 50 insertions(+), 125 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/reduction.hpp b/sycl/include/sycl/ext/oneapi/reduction.hpp
index 8e07bd5506ead..8c99aefc2bc16 100644
--- a/sycl/include/sycl/ext/oneapi/reduction.hpp
+++ b/sycl/include/sycl/ext/oneapi/reduction.hpp
@@ -752,6 +752,32 @@ class reduction_impl
 private:
   using algo = reduction_impl_algo<T, BinaryOperation, Dims, Extent, RedOutVar,
                                    Algorithm>;
+  using self = reduction_impl<T, BinaryOperation, Dims, Extent, RedOutVar, Algorithm>;
+
+  static constexpr bool is_known_identity =
+      sycl::detail::IsKnownIdentityOp<T, BinaryOperation>::value;
+
+  // TODO: Do we also need chooseBinOp?
+  static constexpr T chooseIdentity(const T &Identity) {
+    // For now the implementation ignores the identity value given by user
+    // when the implementation knows the identity.
+    // The SPEC could prohibit passing identity parameter to operations with
+    // known identity, but that could have some bad consequences too.
+    // For example, at some moment the implementation may NOT know the identity
+    // for COMPLEX-PLUS reduction. User may create a program that would pass
+    // COMPLEX value (0,0) as identity for PLUS reduction. At some later moment
+    // when the implementation starts handling COMPLEX-PLUS as known operation
+    // the existing user's program remains compilable and working correctly.
+    // I.e. with this constructor here, adding more reduction operations to the
+    // list of known operations does not break the existing programs.
+    if constexpr (is_known_identity) {
+      (void)Identity;
+      return reducer_type::getIdentity();
+
+    } else {
+      return Identity;
+    }
+  }
 
 public:
   using algo::my_is_usm;
@@ -767,11 +793,10 @@ class reduction_impl
 
   /// SYCL-2020.
   /// Constructs reduction_impl when the identity value is statically known.
-  template <typename _T, typename AllocatorT,
-            std::enable_if_t<
-                sycl::detail::IsKnownIdentityOp<_T, BinaryOperation>::value &&
-                my_is_rw_acc> * = nullptr>
-  reduction_impl(buffer<_T, 1, AllocatorT> Buffer, handler &CGH,
+  template <typename AllocatorT, typename _self = self,
+            std::enable_if_t<_self::is_known_identity && _self::my_is_rw_acc>
+                * = nullptr>
+  reduction_impl(buffer<T, 1, AllocatorT> Buffer, handler &CGH,
                  bool InitializeToIdentity)
       : algo(reducer_type::getIdentity(), BinaryOperation(),
              InitializeToIdentity, rw_accessor_type{Buffer}) {
@@ -783,10 +808,9 @@ class reduction_impl
   }
 
   /// Constructs reduction_impl when the identity value is statically known.
-  template <
-      typename _T = T,
-      enable_if_t<sycl::detail::IsKnownIdentityOp<_T, BinaryOperation>::value &&
-                  (my_is_rw_acc || my_is_dw_acc)> * = nullptr>
+  template <typename _self = self,
+            enable_if_t<_self::is_known_identity &&
+                        (my_is_rw_acc || my_is_dw_acc)> * = nullptr>
   reduction_impl(RedOutVar &Acc)
       : algo(reducer_type::getIdentity(), BinaryOperation(), false, Acc) {
     if (Acc.size() != 1)
@@ -798,69 +822,14 @@ class reduction_impl
   /// SYCL-2020.
   /// Constructs reduction_impl when the identity value is statically known,
   /// and user still passed the identity value.
-  template <
-      typename _T, typename AllocatorT,
-      enable_if_t<sycl::detail::IsKnownIdentityOp<_T, BinaryOperation>::value &&
-                  my_is_rw_acc> * = nullptr>
-  reduction_impl(buffer<_T, 1, AllocatorT> Buffer, handler &CGH,
-                 const T & /*Identity*/, BinaryOperation,
-                 bool InitializeToIdentity)
-      : algo(reducer_type::getIdentity(), BinaryOperation(),
-             InitializeToIdentity, rw_accessor_type{Buffer}) {
-    algo::associateWithHandler(CGH);
-    if (Buffer.size() != 1)
-      throw sycl::runtime_error(errc::invalid,
-                                "Reduction variable must be a scalar.",
-                                PI_ERROR_INVALID_VALUE);
-    // For now the implementation ignores the identity value given by user
-    // when the implementation knows the identity.
-    // The SPEC could prohibit passing identity parameter to operations with
-    // known identity, but that could have some bad consequences too.
-    // For example, at some moment the implementation may NOT know the identity
-    // for COMPLEX-PLUS reduction. User may create a program that would pass
-    // COMPLEX value (0,0) as identity for PLUS reduction. At some later moment
-    // when the implementation starts handling COMPLEX-PLUS as known operation
-    // the existing user's program remains compilable and working correctly.
-    // I.e. with this constructor here, adding more reduction operations to the
-    // list of known operations does not break the existing programs.
-  }
-
-  /// Constructs reduction_impl when the identity value is statically known,
-  /// and user still passed the identity value.
-  template <
-      typename _T = T,
-      enable_if_t<sycl::detail::IsKnownIdentityOp<_T, BinaryOperation>::value &&
-                  (my_is_rw_acc || my_is_dw_acc)> * = nullptr>
-  reduction_impl(RedOutVar &Acc, const T & /*Identity*/, BinaryOperation)
-      : algo(reducer_type::getIdentity(), BinaryOperation(), my_is_dw_acc,
-             Acc) {
-    if (Acc.size() != 1)
-      throw sycl::runtime_error(errc::invalid,
-                                "Reduction variable must be a scalar.",
-                                PI_ERROR_INVALID_VALUE);
-    // For now the implementation ignores the identity value given by user
-    // when the implementation knows the identity.
-    // The SPEC could prohibit passing identity parameter to operations with
-    // known identity, but that could have some bad consequences too.
-    // For example, at some moment the implementation may NOT know the identity
-    // for COMPLEX-PLUS reduction. User may create a program that would pass
-    // COMPLEX value (0,0) as identity for PLUS reduction. At some later moment
-    // when the implementation starts handling COMPLEX-PLUS as known operation
-    // the existing user's program remains compilable and working correctly.
-    // I.e. with this constructor here, adding more reduction operations to the
-    // list of known operations does not break the existing programs.
-  }
-
   /// SYCL-2020.
   /// Constructs reduction_impl when the identity value is NOT known statically.
-  template <typename _T, typename AllocatorT,
-            enable_if_t<
-                !sycl::detail::IsKnownIdentityOp<_T, BinaryOperation>::value &&
-                my_is_rw_acc> * = nullptr>
-  reduction_impl(buffer<_T, 1, AllocatorT> Buffer, handler &CGH,
+  template <typename AllocatorT, typename _self = self,
+            enable_if_t<_self::my_is_rw_acc> * = nullptr>
+  reduction_impl(buffer<T, 1, AllocatorT> Buffer, handler &CGH,
                  const T &Identity, BinaryOperation BOp,
                  bool InitializeToIdentity)
-      : algo(Identity, BOp, InitializeToIdentity,
+      : algo(chooseIdentity(Identity), BOp, InitializeToIdentity,
              rw_accessor_type{Buffer}) {
     algo::associateWithHandler(CGH);
     if (Buffer.size() != 1)
@@ -870,11 +839,10 @@ class reduction_impl
   }
 
   /// Constructs reduction_impl when the identity value is unknown.
-  template <typename _T = T, enable_if_t<!sycl::detail::IsKnownIdentityOp<
-                                             _T, BinaryOperation>::value &&
-                                         (my_is_rw_acc || my_is_dw_acc)> * = nullptr>
+  template <typename _self = self,
+            enable_if_t<_self::my_is_rw_acc || _self::my_is_dw_acc> * = nullptr>
   reduction_impl(RedOutVar &Acc, const T &Identity, BinaryOperation BOp)
-      : algo(Identity, BOp, my_is_dw_acc, Acc) {
+      : algo(chooseIdentity(Identity), BOp, my_is_dw_acc, Acc) {
     if (Acc.size() != 1)
       throw sycl::runtime_error(errc::invalid,
                                 "Reduction variable must be a scalar.",
@@ -885,75 +853,32 @@ class reduction_impl
   /// The \param VarPtr is a USM pointer to memory, to where the computed
   /// reduction value is added using BinaryOperation, i.e. it is expected that
   /// the memory is pre-initialized with some meaningful value.
-  template <
-      typename _T = T,
-      enable_if_t<sycl::detail::IsKnownIdentityOp<_T, BinaryOperation>::value &&
-                  my_is_usm> * = nullptr>
+  template <typename _self = self, enable_if_t<_self::is_known_identity &&
+                                               _self::my_is_usm> * = nullptr>
   reduction_impl(T *VarPtr, bool InitializeToIdentity = false)
       : algo(reducer_type::getIdentity(), BinaryOperation(),
              InitializeToIdentity, VarPtr) {}
 
-  /// Constructs reduction_impl when the identity value is statically known,
-  /// and user still passed the identity value.
   /// The \param VarPtr is a USM pointer to memory, to where the computed
   /// reduction value is added using BinaryOperation, i.e. it is expected that
   /// the memory is pre-initialized with some meaningful value.
-  template <
-      typename _T = T,
-      enable_if_t<sycl::detail::IsKnownIdentityOp<_T, BinaryOperation>::value &&
-                  my_is_usm> * = nullptr>
-  reduction_impl(T *VarPtr, const T &Identity, BinaryOperation,
-                 bool InitializeToIdentity = false)
-      : algo(Identity, BinaryOperation(), InitializeToIdentity, VarPtr) {
-    // For now the implementation ignores the identity value given by user
-    // when the implementation knows the identity.
-    // The SPEC could prohibit passing identity parameter to operations with
-    // known identity, but that could have some bad consequences too.
-    // For example, at some moment the implementation may NOT know the identity
-    // for COMPLEX-PLUS reduction. User may create a program that would pass
-    // COMPLEX value (0,0) as identity for PLUS reduction. At some later moment
-    // when the implementation starts handling COMPLEX-PLUS as known operation
-    // the existing user's program remains compilable and working correctly.
-    // I.e. with this constructor here, adding more reduction operations to the
-    // list of known operations does not break the existing programs.
-  }
-
-  /// Constructs reduction_impl when the identity value is unknown.
-  /// The \param VarPtr is a USM pointer to memory, to where the computed
-  /// reduction value is added using BinaryOperation, i.e. it is expected that
-  /// the memory is pre-initialized with some meaningful value.
-  template <typename _T = T, enable_if_t<!sycl::detail::IsKnownIdentityOp<
-                                             _T, BinaryOperation>::value &&
-                                         my_is_usm> * = nullptr>
+  template <typename _self = self, enable_if_t<_self::my_is_usm> * = nullptr>
   reduction_impl(T *VarPtr, const T &Identity, BinaryOperation BOp,
                  bool InitializeToIdentity = false)
-      : algo(Identity, BOp, InitializeToIdentity, VarPtr) {}
+      : algo(chooseIdentity(Identity), BOp, InitializeToIdentity, VarPtr) {}
 
   /// Constructs reduction_impl when the identity value is statically known
-  template <typename _T = T, enable_if_t<sycl::detail::IsKnownIdentityOp<
-                                 _T, BinaryOperation>::value> * = nullptr>
-  reduction_impl(span<_T, Extent> Span, bool InitializeToIdentity = false)
+  template <typename _self = self, enable_if_t<_self::is_known_identity &&
+                                               _self::my_is_usm> * = nullptr>
+  reduction_impl(span<T, Extent> Span, bool InitializeToIdentity = false)
       : algo(reducer_type::getIdentity(), BinaryOperation(),
              InitializeToIdentity, Span.data()) {}
 
-  /// Constructs reduction_impl when the identity value is statically known
-  /// and user passed an identity value anyway
-  template <
-      typename _T = T,
-      enable_if_t<sycl::detail::IsKnownIdentityOp<_T, BinaryOperation>::value &&
-                  my_is_usm> * = nullptr>
-  reduction_impl(span<_T, Extent> Span, const T & /* Identity */,
-                 BinaryOperation BOp, bool InitializeToIdentity = false)
-      : algo(reducer_type::getIdentity(), BOp, InitializeToIdentity,
-             Span.data()) {}
-
-  /// Constructs reduction_impl when the identity value is not statically known
-  template <typename _T = T, enable_if_t<!sycl::detail::IsKnownIdentityOp<
-                                             _T, BinaryOperation>::value &&
-                                         my_is_usm> * = nullptr>
+  template <typename _self = self, enable_if_t<_self::my_is_usm> * = nullptr>
   reduction_impl(span<T, Extent> Span, const T &Identity, BinaryOperation BOp,
                  bool InitializeToIdentity = false)
-      : algo(Identity, BOp, InitializeToIdentity, Span.data()) {}
+      : algo(chooseIdentity(Identity), BOp, InitializeToIdentity, Span.data()) {
+  }
 };
 
 /// A helper to pass undefined (sycl::detail::auto_name) names unmodified. We

From 862a4249bae6d3f3071cf900ed81d03f626ccc7d Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Thu, 14 Jul 2022 16:07:31 -0700
Subject: [PATCH 09/24] Move RedOutVar to the end of the template params list

I hoped that I'll be able to deduce it from arguments, but partial CTAD
doesn't exist, unfortunately.
---
 sycl/include/sycl/ext/oneapi/reduction.hpp | 46 +++++++++++----------
 sycl/include/sycl/handler.hpp              |  4 +-
 sycl/include/sycl/reduction.hpp            | 47 ++++++++++++----------
 3 files changed, 53 insertions(+), 44 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/reduction.hpp b/sycl/include/sycl/ext/oneapi/reduction.hpp
index 8c99aefc2bc16..61a7dbf78906c 100644
--- a/sycl/include/sycl/ext/oneapi/reduction.hpp
+++ b/sycl/include/sycl/ext/oneapi/reduction.hpp
@@ -521,22 +521,23 @@ class default_reduction_algorithm {};
 
 /// Templated class for implementations of specific reduction algorithms
 template <typename T, class BinaryOperation, int Dims, size_t Extent,
-          typename RedOutVar, class Algorithm>
+          class Algorithm, typename RedOutVar>
 class reduction_impl_algo;
 
 /// Original reduction algorithm is the default. It supports both USM and
 /// accessors via a single class
 template <typename T, class BinaryOperation, int Dims, size_t Extent,
-          typename RedOutVar, bool IsUSM, access::placeholder IsPlaceholder,
-          int AccessorDims>
+          bool IsUSM, access::placeholder IsPlaceholder, int AccessorDims,
+          typename RedOutVar>
 class reduction_impl_algo<
-    T, BinaryOperation, Dims, Extent, RedOutVar,
-    default_reduction_algorithm<IsUSM, IsPlaceholder, AccessorDims>>
+    T, BinaryOperation, Dims, Extent,
+    default_reduction_algorithm<IsUSM, IsPlaceholder, AccessorDims>, RedOutVar>
     : public reduction_impl_common<T, BinaryOperation> {
   using base = reduction_impl_common<T, BinaryOperation>;
   using self = reduction_impl_algo<
-      T, BinaryOperation, Dims, Extent, RedOutVar,
-      default_reduction_algorithm<IsUSM, IsPlaceholder, AccessorDims>>;
+      T, BinaryOperation, Dims, Extent,
+      default_reduction_algorithm<IsUSM, IsPlaceholder, AccessorDims>,
+      RedOutVar>;
 
 public:
   using reducer_type = reducer<T, BinaryOperation, Dims, Extent>;
@@ -744,15 +745,16 @@ template <typename T> struct AreAllButLastReductions<T> {
 /// This class encapsulates the reduction variable/accessor,
 /// the reduction operator and an optional operator identity.
 template <typename T, class BinaryOperation, int Dims, size_t Extent,
-          typename RedOutVar, class Algorithm>
+          class Algorithm, typename RedOutVar>
 class reduction_impl
     : private reduction_impl_base,
-      public reduction_impl_algo<T, BinaryOperation, Dims, Extent, RedOutVar,
-                                 Algorithm> {
+      public reduction_impl_algo<T, BinaryOperation, Dims, Extent, Algorithm,
+                                 RedOutVar> {
 private:
-  using algo = reduction_impl_algo<T, BinaryOperation, Dims, Extent, RedOutVar,
-                                   Algorithm>;
-  using self = reduction_impl<T, BinaryOperation, Dims, Extent, RedOutVar, Algorithm>;
+  using algo = reduction_impl_algo<T, BinaryOperation, Dims, Extent,
+                                   Algorithm, RedOutVar>;
+  using self =
+      reduction_impl<T, BinaryOperation, Dims, Extent, Algorithm, RedOutVar>;
 
   static constexpr bool is_known_identity =
       sycl::detail::IsKnownIdentityOp<T, BinaryOperation>::value;
@@ -2464,8 +2466,8 @@ tuple_select_elements(TupleT Tuple, std::index_sequence<Is...>) {
 template <typename T, class BinaryOperation, int Dims, access::mode AccMode,
           access::placeholder IsPH>
 detail::reduction_impl<T, BinaryOperation, 0, 1,
-                       accessor<T, Dims, AccMode, access::target::device, IsPH>,
-                       detail::default_reduction_algorithm<false, IsPH, Dims>>
+                       detail::default_reduction_algorithm<false, IsPH, Dims>,
+                       accessor<T, Dims, AccMode, access::target::device, IsPH>>
 reduction(accessor<T, Dims, AccMode, access::target::device, IsPH> &Acc,
           const T &Identity, BinaryOperation BOp) {
   return {Acc, Identity, BOp};
@@ -2480,8 +2482,8 @@ template <typename T, class BinaryOperation, int Dims, access::mode AccMode,
 std::enable_if_t<sycl::detail::IsKnownIdentityOp<T, BinaryOperation>::value,
                  detail::reduction_impl<
                      T, BinaryOperation, 0, 1,
-                     accessor<T, Dims, AccMode, access::target::device, IsPH>,
-                     detail::default_reduction_algorithm<false, IsPH, Dims>>>
+                     detail::default_reduction_algorithm<false, IsPH, Dims>,
+                     accessor<T, Dims, AccMode, access::target::device, IsPH>>>
 reduction(accessor<T, Dims, AccMode, access::target::device, IsPH> &Acc,
           BinaryOperation) {
   return {Acc};
@@ -2493,8 +2495,9 @@ reduction(accessor<T, Dims, AccMode, access::target::device, IsPH> &Acc,
 /// \param Identity, and the binary operation used in the reduction.
 template <typename T, class BinaryOperation>
 detail::reduction_impl<
-    T, BinaryOperation, 0, 1, T *,
-    detail::default_reduction_algorithm<true, access::placeholder::false_t, 1>>
+    T, BinaryOperation, 0, 1,
+    detail::default_reduction_algorithm<true, access::placeholder::false_t, 1>,
+    T *>
 reduction(T *VarPtr, const T &Identity, BinaryOperation BOp) {
   return {VarPtr, Identity, BOp};
 }
@@ -2507,9 +2510,10 @@ reduction(T *VarPtr, const T &Identity, BinaryOperation BOp) {
 template <typename T, class BinaryOperation>
 std::enable_if_t<
     sycl::detail::IsKnownIdentityOp<T, BinaryOperation>::value,
-    detail::reduction_impl<T, BinaryOperation, 0, 1, T *,
+    detail::reduction_impl<T, BinaryOperation, 0, 1,
                            detail::default_reduction_algorithm<
-                               true, access::placeholder::false_t, 1>>>
+                               true, access::placeholder::false_t, 1>,
+                           T *>>
 reduction(T *VarPtr, BinaryOperation) {
   return {VarPtr};
 }
diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp
index 8d99605b699df..220c2ea1122f3 100644
--- a/sycl/include/sycl/handler.hpp
+++ b/sycl/include/sycl/handler.hpp
@@ -242,7 +242,7 @@ namespace ext {
 namespace oneapi {
 namespace detail {
 template <typename T, class BinaryOperation, int Dims, size_t Extent,
-          typename RedOutVar, class Algorithm>
+          class Algorithm, typename RedOutVar>
 class reduction_impl_algo;
 
 using cl::sycl::detail::enable_if_t;
@@ -2667,7 +2667,7 @@ class __SYCL_EXPORT handler {
   // Make reduction friends to store buffers and arrays created for it
   // in handler from reduction methods.
   template <typename T, class BinaryOperation, int Dims, size_t Extent,
-            typename RedOutVar, class Algorithm>
+            class Algorithm, typename RedOutVar>
   friend class ext::oneapi::detail::reduction_impl_algo;
 
 #ifndef __SYCL_DEVICE_ONLY__
diff --git a/sycl/include/sycl/reduction.hpp b/sycl/include/sycl/reduction.hpp
index bb8cb269980c6..de9635e164356 100644
--- a/sycl/include/sycl/reduction.hpp
+++ b/sycl/include/sycl/reduction.hpp
@@ -25,11 +25,11 @@ std::enable_if_t<
     has_known_identity<BinaryOperation, T>::value,
     ext::oneapi::detail::reduction_impl<
         T, BinaryOperation, 0, 1,
+        ext::oneapi::detail::default_reduction_algorithm<
+            false, access::placeholder::true_t, 1>,
         accessor<T, 1, access::mode::read_write, access::target::device,
                  access::placeholder::true_t,
-                 ext::oneapi::accessor_property_list<>>,
-        ext::oneapi::detail::default_reduction_algorithm<
-            false, access::placeholder::true_t, 1>>>
+                 ext::oneapi::accessor_property_list<>>>>
 reduction(buffer<T, 1, AllocatorT> Var, handler &CGH, BinaryOperation,
           const property_list &PropList = {}) {
   bool InitializeToIdentity =
@@ -46,11 +46,11 @@ std::enable_if_t<
     !has_known_identity<BinaryOperation, T>::value,
     ext::oneapi::detail::reduction_impl<
         T, BinaryOperation, 0, 1,
+        ext::oneapi::detail::default_reduction_algorithm<
+            false, access::placeholder::true_t, 1>,
         accessor<T, 1, access::mode::read_write, access::target::device,
                  access::placeholder::true_t,
-                 ext::oneapi::accessor_property_list<>>,
-        ext::oneapi::detail::default_reduction_algorithm<
-            false, access::placeholder::true_t, 1>>>
+                 ext::oneapi::accessor_property_list<>>>>
 reduction(buffer<T, 1, AllocatorT>, handler &, BinaryOperation,
           const property_list &PropList = {}) {
   // TODO: implement reduction that works even when identity is not known.
@@ -66,9 +66,10 @@ reduction(buffer<T, 1, AllocatorT>, handler &, BinaryOperation,
 template <typename T, typename BinaryOperation>
 std::enable_if_t<has_known_identity<BinaryOperation, T>::value,
                  ext::oneapi::detail::reduction_impl<
-                     T, BinaryOperation, 0, 1, T *,
+                     T, BinaryOperation, 0, 1,
                      ext::oneapi::detail::default_reduction_algorithm<
-                         true, access::placeholder::false_t, 1>>>
+                         true, access::placeholder::false_t, 1>,
+                     T *>>
 reduction(T *Var, BinaryOperation, const property_list &PropList = {}) {
   bool InitializeToIdentity =
       PropList.has_property<property::reduction::initialize_to_identity>();
@@ -83,9 +84,10 @@ reduction(T *Var, BinaryOperation, const property_list &PropList = {}) {
 template <typename T, typename BinaryOperation>
 std::enable_if_t<!has_known_identity<BinaryOperation, T>::value,
                  ext::oneapi::detail::reduction_impl<
-                     T, BinaryOperation, 0, 1, T *,
+                     T, BinaryOperation, 0, 1,
                      ext::oneapi::detail::default_reduction_algorithm<
-                         true, access::placeholder::false_t, 1>>>
+                         true, access::placeholder::false_t, 1>,
+                     T *>>
 reduction(T *, BinaryOperation, const property_list &PropList = {}) {
   // TODO: implement reduction that works even when identity is not known.
   (void)PropList;
@@ -100,11 +102,11 @@ reduction(T *, BinaryOperation, const property_list &PropList = {}) {
 template <typename T, typename AllocatorT, typename BinaryOperation>
 ext::oneapi::detail::reduction_impl<
     T, BinaryOperation, 0, 1,
+    ext::oneapi::detail::default_reduction_algorithm<
+        false, access::placeholder::true_t, 1>,
     accessor<T, 1, access::mode::read_write, access::target::device,
              access::placeholder::true_t,
-             ext::oneapi::accessor_property_list<>>,
-    ext::oneapi::detail::default_reduction_algorithm<
-        false, access::placeholder::true_t, 1>>
+             ext::oneapi::accessor_property_list<>>>
 reduction(buffer<T, 1, AllocatorT> Var, handler &CGH, const T &Identity,
           BinaryOperation Combiner, const property_list &PropList = {}) {
   bool InitializeToIdentity =
@@ -117,9 +119,10 @@ reduction(buffer<T, 1, AllocatorT> Var, handler &CGH, const T &Identity,
 /// binary operation \p Combiner, and optional reduction properties.
 template <typename T, typename BinaryOperation>
 ext::oneapi::detail::reduction_impl<
-    T, BinaryOperation, 0, 1, T *,
+    T, BinaryOperation, 0, 1,
     ext::oneapi::detail::default_reduction_algorithm<
-        true, access::placeholder::false_t, 1>>
+        true, access::placeholder::false_t, 1>,
+    T *>
 reduction(T *Var, const T &Identity, BinaryOperation Combiner,
           const property_list &PropList = {}) {
   bool InitializeToIdentity =
@@ -135,9 +138,9 @@ template <typename T, size_t Extent, typename BinaryOperation>
 std::enable_if_t<Extent != dynamic_extent &&
                      has_known_identity<BinaryOperation, T>::value,
                  ext::oneapi::detail::reduction_impl<
-                     T, BinaryOperation, 1, Extent, T *,
+                     T, BinaryOperation, 1, Extent,
                      ext::oneapi::detail::default_reduction_algorithm<
-                         true, access::placeholder::false_t, 1>>>
+                         true, access::placeholder::false_t, 1>, T *>>
 reduction(span<T, Extent> Span, BinaryOperation,
           const property_list &PropList = {}) {
   bool InitializeToIdentity =
@@ -154,9 +157,10 @@ template <typename T, size_t Extent, typename BinaryOperation>
 std::enable_if_t<Extent != dynamic_extent &&
                      !has_known_identity<BinaryOperation, T>::value,
                  ext::oneapi::detail::reduction_impl<
-                     T, BinaryOperation, 1, Extent, T *,
+                     T, BinaryOperation, 1, Extent,
                      ext::oneapi::detail::default_reduction_algorithm<
-                         true, access::placeholder::false_t, 1>>>
+                         true, access::placeholder::false_t, 1>,
+                     T *>>
 reduction(span<T, Extent>, BinaryOperation,
           const property_list &PropList = {}) {
   // TODO: implement reduction that works even when identity is not known.
@@ -172,9 +176,10 @@ reduction(span<T, Extent>, BinaryOperation,
 template <typename T, size_t Extent, typename BinaryOperation>
 std::enable_if_t<Extent != dynamic_extent,
                  ext::oneapi::detail::reduction_impl<
-                     T, BinaryOperation, 1, Extent, T *,
+                     T, BinaryOperation, 1, Extent,
                      ext::oneapi::detail::default_reduction_algorithm<
-                         true, access::placeholder::false_t, 1>>>
+                         true, access::placeholder::false_t, 1>,
+                     T *>>
 reduction(span<T, Extent> Span, const T &Identity, BinaryOperation Combiner,
           const property_list &PropList = {}) {
   bool InitializeToIdentity =

From 96502df04f29467ab890f21a7bf50c718261fa44 Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Thu, 14 Jul 2022 16:43:16 -0700
Subject: [PATCH 10/24] Remove IsUSM from default_reduction_algorithm

---
 sycl/include/sycl/ext/oneapi/reduction.hpp | 30 ++++++++++------------
 sycl/include/sycl/reduction.hpp            | 19 +++++++-------
 2 files changed, 23 insertions(+), 26 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/reduction.hpp b/sycl/include/sycl/ext/oneapi/reduction.hpp
index 61a7dbf78906c..f7e8f598d8476 100644
--- a/sycl/include/sycl/ext/oneapi/reduction.hpp
+++ b/sycl/include/sycl/ext/oneapi/reduction.hpp
@@ -516,7 +516,7 @@ struct is_dw_acc_t<accessor<T, AccessorDims, access::mode::discard_write,
 
 /// Types representing specific reduction algorithms
 /// Enables reduction_impl_algo to take additional algorithm-specific templates
-template <bool IsUSM, access::placeholder IsPlaceholder, int AccessorDims>
+template <access::placeholder IsPlaceholder, int AccessorDims>
 class default_reduction_algorithm {};
 
 /// Templated class for implementations of specific reduction algorithms
@@ -527,17 +527,16 @@ class reduction_impl_algo;
 /// Original reduction algorithm is the default. It supports both USM and
 /// accessors via a single class
 template <typename T, class BinaryOperation, int Dims, size_t Extent,
-          bool IsUSM, access::placeholder IsPlaceholder, int AccessorDims,
+          access::placeholder IsPlaceholder, int AccessorDims,
           typename RedOutVar>
 class reduction_impl_algo<
     T, BinaryOperation, Dims, Extent,
-    default_reduction_algorithm<IsUSM, IsPlaceholder, AccessorDims>, RedOutVar>
+    default_reduction_algorithm<IsPlaceholder, AccessorDims>, RedOutVar>
     : public reduction_impl_common<T, BinaryOperation> {
   using base = reduction_impl_common<T, BinaryOperation>;
   using self = reduction_impl_algo<
       T, BinaryOperation, Dims, Extent,
-      default_reduction_algorithm<IsUSM, IsPlaceholder, AccessorDims>,
-      RedOutVar>;
+      default_reduction_algorithm<IsPlaceholder, AccessorDims>, RedOutVar>;
 
 public:
   using reducer_type = reducer<T, BinaryOperation, Dims, Extent>;
@@ -565,9 +564,8 @@ class reduction_impl_algo<
   static constexpr bool has_fast_reduce =
       IsReduOptForFastReduce<T, BinaryOperation>::value;
 
-  static constexpr bool is_usm = IsUSM;
   static constexpr bool my_is_usm = std::is_same_v<RedOutVar, T *>;
-  static_assert(is_usm == my_is_usm);
+  static constexpr bool is_usm = my_is_usm;
 
   static constexpr bool my_is_rw_acc = is_rw_acc_t<RedOutVar>::value;
   static constexpr bool my_is_dw_acc = is_dw_acc_t<RedOutVar>::value;
@@ -2466,7 +2464,7 @@ tuple_select_elements(TupleT Tuple, std::index_sequence<Is...>) {
 template <typename T, class BinaryOperation, int Dims, access::mode AccMode,
           access::placeholder IsPH>
 detail::reduction_impl<T, BinaryOperation, 0, 1,
-                       detail::default_reduction_algorithm<false, IsPH, Dims>,
+                       detail::default_reduction_algorithm<IsPH, Dims>,
                        accessor<T, Dims, AccMode, access::target::device, IsPH>>
 reduction(accessor<T, Dims, AccMode, access::target::device, IsPH> &Acc,
           const T &Identity, BinaryOperation BOp) {
@@ -2482,7 +2480,7 @@ template <typename T, class BinaryOperation, int Dims, access::mode AccMode,
 std::enable_if_t<sycl::detail::IsKnownIdentityOp<T, BinaryOperation>::value,
                  detail::reduction_impl<
                      T, BinaryOperation, 0, 1,
-                     detail::default_reduction_algorithm<false, IsPH, Dims>,
+                     detail::default_reduction_algorithm<IsPH, Dims>,
                      accessor<T, Dims, AccMode, access::target::device, IsPH>>>
 reduction(accessor<T, Dims, AccMode, access::target::device, IsPH> &Acc,
           BinaryOperation) {
@@ -2496,8 +2494,7 @@ reduction(accessor<T, Dims, AccMode, access::target::device, IsPH> &Acc,
 template <typename T, class BinaryOperation>
 detail::reduction_impl<
     T, BinaryOperation, 0, 1,
-    detail::default_reduction_algorithm<true, access::placeholder::false_t, 1>,
-    T *>
+    detail::default_reduction_algorithm<access::placeholder::false_t, 1>, T *>
 reduction(T *VarPtr, const T &Identity, BinaryOperation BOp) {
   return {VarPtr, Identity, BOp};
 }
@@ -2508,12 +2505,11 @@ reduction(T *VarPtr, const T &Identity, BinaryOperation BOp) {
 /// operation used in the reduction.
 /// The identity value is not passed to this version as it is statically known.
 template <typename T, class BinaryOperation>
-std::enable_if_t<
-    sycl::detail::IsKnownIdentityOp<T, BinaryOperation>::value,
-    detail::reduction_impl<T, BinaryOperation, 0, 1,
-                           detail::default_reduction_algorithm<
-                               true, access::placeholder::false_t, 1>,
-                           T *>>
+std::enable_if_t<sycl::detail::IsKnownIdentityOp<T, BinaryOperation>::value,
+                 detail::reduction_impl<T, BinaryOperation, 0, 1,
+                                        detail::default_reduction_algorithm<
+                                            access::placeholder::false_t, 1>,
+                                        T *>>
 reduction(T *VarPtr, BinaryOperation) {
   return {VarPtr};
 }
diff --git a/sycl/include/sycl/reduction.hpp b/sycl/include/sycl/reduction.hpp
index de9635e164356..340947e049e36 100644
--- a/sycl/include/sycl/reduction.hpp
+++ b/sycl/include/sycl/reduction.hpp
@@ -26,7 +26,7 @@ std::enable_if_t<
     ext::oneapi::detail::reduction_impl<
         T, BinaryOperation, 0, 1,
         ext::oneapi::detail::default_reduction_algorithm<
-            false, access::placeholder::true_t, 1>,
+            access::placeholder::true_t, 1>,
         accessor<T, 1, access::mode::read_write, access::target::device,
                  access::placeholder::true_t,
                  ext::oneapi::accessor_property_list<>>>>
@@ -47,7 +47,7 @@ std::enable_if_t<
     ext::oneapi::detail::reduction_impl<
         T, BinaryOperation, 0, 1,
         ext::oneapi::detail::default_reduction_algorithm<
-            false, access::placeholder::true_t, 1>,
+            access::placeholder::true_t, 1>,
         accessor<T, 1, access::mode::read_write, access::target::device,
                  access::placeholder::true_t,
                  ext::oneapi::accessor_property_list<>>>>
@@ -68,7 +68,7 @@ std::enable_if_t<has_known_identity<BinaryOperation, T>::value,
                  ext::oneapi::detail::reduction_impl<
                      T, BinaryOperation, 0, 1,
                      ext::oneapi::detail::default_reduction_algorithm<
-                         true, access::placeholder::false_t, 1>,
+                         access::placeholder::false_t, 1>,
                      T *>>
 reduction(T *Var, BinaryOperation, const property_list &PropList = {}) {
   bool InitializeToIdentity =
@@ -86,7 +86,7 @@ std::enable_if_t<!has_known_identity<BinaryOperation, T>::value,
                  ext::oneapi::detail::reduction_impl<
                      T, BinaryOperation, 0, 1,
                      ext::oneapi::detail::default_reduction_algorithm<
-                         true, access::placeholder::false_t, 1>,
+                         access::placeholder::false_t, 1>,
                      T *>>
 reduction(T *, BinaryOperation, const property_list &PropList = {}) {
   // TODO: implement reduction that works even when identity is not known.
@@ -103,7 +103,7 @@ template <typename T, typename AllocatorT, typename BinaryOperation>
 ext::oneapi::detail::reduction_impl<
     T, BinaryOperation, 0, 1,
     ext::oneapi::detail::default_reduction_algorithm<
-        false, access::placeholder::true_t, 1>,
+        access::placeholder::true_t, 1>,
     accessor<T, 1, access::mode::read_write, access::target::device,
              access::placeholder::true_t,
              ext::oneapi::accessor_property_list<>>>
@@ -121,7 +121,7 @@ template <typename T, typename BinaryOperation>
 ext::oneapi::detail::reduction_impl<
     T, BinaryOperation, 0, 1,
     ext::oneapi::detail::default_reduction_algorithm<
-        true, access::placeholder::false_t, 1>,
+        access::placeholder::false_t, 1>,
     T *>
 reduction(T *Var, const T &Identity, BinaryOperation Combiner,
           const property_list &PropList = {}) {
@@ -140,7 +140,8 @@ std::enable_if_t<Extent != dynamic_extent &&
                  ext::oneapi::detail::reduction_impl<
                      T, BinaryOperation, 1, Extent,
                      ext::oneapi::detail::default_reduction_algorithm<
-                         true, access::placeholder::false_t, 1>, T *>>
+                         access::placeholder::false_t, 1>,
+                     T *>>
 reduction(span<T, Extent> Span, BinaryOperation,
           const property_list &PropList = {}) {
   bool InitializeToIdentity =
@@ -159,7 +160,7 @@ std::enable_if_t<Extent != dynamic_extent &&
                  ext::oneapi::detail::reduction_impl<
                      T, BinaryOperation, 1, Extent,
                      ext::oneapi::detail::default_reduction_algorithm<
-                         true, access::placeholder::false_t, 1>,
+                         access::placeholder::false_t, 1>,
                      T *>>
 reduction(span<T, Extent>, BinaryOperation,
           const property_list &PropList = {}) {
@@ -178,7 +179,7 @@ std::enable_if_t<Extent != dynamic_extent,
                  ext::oneapi::detail::reduction_impl<
                      T, BinaryOperation, 1, Extent,
                      ext::oneapi::detail::default_reduction_algorithm<
-                         true, access::placeholder::false_t, 1>,
+                         access::placeholder::false_t, 1>,
                      T *>>
 reduction(span<T, Extent> Span, const T &Identity, BinaryOperation Combiner,
           const property_list &PropList = {}) {

From 252a71c262a08b529e37517cfe25077354dd5238 Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Thu, 14 Jul 2022 17:21:11 -0700
Subject: [PATCH 11/24] Remove IsPlaceholder from default_reduction_algorithm

---
 sycl/include/sycl/ext/oneapi/reduction.hpp | 78 +++++++++++-----------
 sycl/include/sycl/reduction.hpp            | 33 +++------
 2 files changed, 48 insertions(+), 63 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/reduction.hpp b/sycl/include/sycl/ext/oneapi/reduction.hpp
index f7e8f598d8476..15cc1be9684ad 100644
--- a/sycl/include/sycl/ext/oneapi/reduction.hpp
+++ b/sycl/include/sycl/ext/oneapi/reduction.hpp
@@ -499,24 +499,32 @@ template <typename T, class BinaryOperation> class reduction_impl_common {
 template <class T>
 struct is_rw_acc_t : public std::false_type {};
 
-template <class T, int AccessorDims, access::placeholder IsPlaceholder>
-struct is_rw_acc_t<
-    accessor<T, AccessorDims, access::mode::read_write, access::target::device,
-             IsPlaceholder, ext::oneapi::accessor_property_list<>>>
+template <class T, int AccessorDims, access::placeholder IsPlaceholder,
+          typename PropList>
+struct is_rw_acc_t<accessor<T, AccessorDims, access::mode::read_write,
+                            access::target::device, IsPlaceholder, PropList>>
     : public std::true_type {};
 
 template <class T>
 struct is_dw_acc_t : public std::false_type {};
 
-template <class T, int AccessorDims, access::placeholder IsPlaceholder>
+template <class T, int AccessorDims, access::placeholder IsPlaceholder,
+          typename PropList>
 struct is_dw_acc_t<accessor<T, AccessorDims, access::mode::discard_write,
-                            access::target::device, IsPlaceholder,
-                            ext::oneapi::accessor_property_list<>>>
+                            access::target::device, IsPlaceholder, PropList>>
+    : public std::true_type {};
+
+template <class T>
+struct is_placeholder_t : public std::false_type {};
+
+template <class T, int AccessorDims, access::mode Mode, typename PropList>
+struct is_placeholder_t<accessor<T, AccessorDims, Mode, access::target::device,
+                                 access::placeholder::true_t, PropList>>
     : public std::true_type {};
 
 /// Types representing specific reduction algorithms
 /// Enables reduction_impl_algo to take additional algorithm-specific templates
-template <access::placeholder IsPlaceholder, int AccessorDims>
+template <int AccessorDims>
 class default_reduction_algorithm {};
 
 /// Templated class for implementations of specific reduction algorithms
@@ -527,16 +535,15 @@ class reduction_impl_algo;
 /// Original reduction algorithm is the default. It supports both USM and
 /// accessors via a single class
 template <typename T, class BinaryOperation, int Dims, size_t Extent,
-          access::placeholder IsPlaceholder, int AccessorDims,
-          typename RedOutVar>
+          int AccessorDims, typename RedOutVar>
 class reduction_impl_algo<
     T, BinaryOperation, Dims, Extent,
-    default_reduction_algorithm<IsPlaceholder, AccessorDims>, RedOutVar>
+    default_reduction_algorithm<AccessorDims>, RedOutVar>
     : public reduction_impl_common<T, BinaryOperation> {
   using base = reduction_impl_common<T, BinaryOperation>;
-  using self = reduction_impl_algo<
-      T, BinaryOperation, Dims, Extent,
-      default_reduction_algorithm<IsPlaceholder, AccessorDims>, RedOutVar>;
+  using self =
+      reduction_impl_algo<T, BinaryOperation, Dims, Extent,
+                          default_reduction_algorithm<AccessorDims>, RedOutVar>;
 
 public:
   using reducer_type = reducer<T, BinaryOperation, Dims, Extent>;
@@ -548,15 +555,12 @@ class reduction_impl_algo<
   // AccessorDims also determines the dimensionality of some temp storage
   static constexpr int accessor_dim = AccessorDims;
   static constexpr int buffer_dim = (AccessorDims == 0) ? 1 : AccessorDims;
+  static constexpr access::placeholder is_placeholder =
+      is_placeholder_t<RedOutVar>::value ? access::placeholder::true_t
+                                         : access::placeholder::false_t;
   using rw_accessor_type = accessor<T, AccessorDims, access::mode::read_write,
-                                    access::target::device, IsPlaceholder,
+                                    access::target::device, is_placeholder,
                                     ext::oneapi::accessor_property_list<>>;
-  using dw_accessor_type =
-      accessor<T, AccessorDims, access::mode::discard_write,
-               access::target::device, IsPlaceholder,
-               ext::oneapi::accessor_property_list<>>;
-
-
   static constexpr bool has_atomic_add_float64 =
       IsReduOptForAtomic64Add<T, BinaryOperation>::value;
   static constexpr bool has_fast_atomics =
@@ -570,9 +574,6 @@ class reduction_impl_algo<
   static constexpr bool my_is_rw_acc = is_rw_acc_t<RedOutVar>::value;
   static constexpr bool my_is_dw_acc = is_dw_acc_t<RedOutVar>::value;
 
-  static constexpr bool is_placeholder =
-      (IsPlaceholder == access::placeholder::true_t);
-
   static constexpr size_t dims = Dims;
   static constexpr size_t num_elements = Extent;
 
@@ -705,14 +706,16 @@ class reduction_impl_algo<
   }
 
 private:
-  template <typename BufferT, access::placeholder IsPH = IsPlaceholder>
-  std::enable_if_t<IsPH == access::placeholder::false_t, rw_accessor_type>
+  template <typename BufferT, typename _self = self>
+  std::enable_if_t<_self::is_placeholder == access::placeholder::false_t,
+                   rw_accessor_type>
   createHandlerWiredReadWriteAccessor(handler &CGH, BufferT Buffer) {
     return {Buffer, CGH};
   }
 
-  template <typename BufferT, access::placeholder IsPH = IsPlaceholder>
-  std::enable_if_t<IsPH == access::placeholder::true_t, rw_accessor_type>
+  template <typename BufferT, typename _self = self>
+  std::enable_if_t<_self::is_placeholder == access::placeholder::true_t,
+                   rw_accessor_type>
   createHandlerWiredReadWriteAccessor(handler &CGH, BufferT Buffer) {
     rw_accessor_type Acc(Buffer);
     CGH.require(Acc);
@@ -786,7 +789,6 @@ class reduction_impl
 
   using reducer_type = typename algo::reducer_type;
   using rw_accessor_type = typename algo::rw_accessor_type;
-  using dw_accessor_type = typename algo::dw_accessor_type;
 
   // Only scalar and 1D array reductions are supported by SYCL 2020.
   static_assert(Dims <= 1, "Multi-dimensional reductions are not supported.");
@@ -2464,7 +2466,7 @@ tuple_select_elements(TupleT Tuple, std::index_sequence<Is...>) {
 template <typename T, class BinaryOperation, int Dims, access::mode AccMode,
           access::placeholder IsPH>
 detail::reduction_impl<T, BinaryOperation, 0, 1,
-                       detail::default_reduction_algorithm<IsPH, Dims>,
+                       detail::default_reduction_algorithm<Dims>,
                        accessor<T, Dims, AccMode, access::target::device, IsPH>>
 reduction(accessor<T, Dims, AccMode, access::target::device, IsPH> &Acc,
           const T &Identity, BinaryOperation BOp) {
@@ -2480,7 +2482,7 @@ template <typename T, class BinaryOperation, int Dims, access::mode AccMode,
 std::enable_if_t<sycl::detail::IsKnownIdentityOp<T, BinaryOperation>::value,
                  detail::reduction_impl<
                      T, BinaryOperation, 0, 1,
-                     detail::default_reduction_algorithm<IsPH, Dims>,
+                     detail::default_reduction_algorithm<Dims>,
                      accessor<T, Dims, AccMode, access::target::device, IsPH>>>
 reduction(accessor<T, Dims, AccMode, access::target::device, IsPH> &Acc,
           BinaryOperation) {
@@ -2492,9 +2494,8 @@ reduction(accessor<T, Dims, AccMode, access::target::device, IsPH> &Acc,
 /// the computed reduction must be stored \param VarPtr, identity value
 /// \param Identity, and the binary operation used in the reduction.
 template <typename T, class BinaryOperation>
-detail::reduction_impl<
-    T, BinaryOperation, 0, 1,
-    detail::default_reduction_algorithm<access::placeholder::false_t, 1>, T *>
+detail::reduction_impl<T, BinaryOperation, 0, 1,
+                       detail::default_reduction_algorithm<1>, T *>
 reduction(T *VarPtr, const T &Identity, BinaryOperation BOp) {
   return {VarPtr, Identity, BOp};
 }
@@ -2505,11 +2506,10 @@ reduction(T *VarPtr, const T &Identity, BinaryOperation BOp) {
 /// operation used in the reduction.
 /// The identity value is not passed to this version as it is statically known.
 template <typename T, class BinaryOperation>
-std::enable_if_t<sycl::detail::IsKnownIdentityOp<T, BinaryOperation>::value,
-                 detail::reduction_impl<T, BinaryOperation, 0, 1,
-                                        detail::default_reduction_algorithm<
-                                            access::placeholder::false_t, 1>,
-                                        T *>>
+std::enable_if_t<
+    sycl::detail::IsKnownIdentityOp<T, BinaryOperation>::value,
+    detail::reduction_impl<T, BinaryOperation, 0, 1,
+                           detail::default_reduction_algorithm<1>, T *>>
 reduction(T *VarPtr, BinaryOperation) {
   return {VarPtr};
 }
diff --git a/sycl/include/sycl/reduction.hpp b/sycl/include/sycl/reduction.hpp
index 340947e049e36..0057c12cd31e9 100644
--- a/sycl/include/sycl/reduction.hpp
+++ b/sycl/include/sycl/reduction.hpp
@@ -25,8 +25,7 @@ std::enable_if_t<
     has_known_identity<BinaryOperation, T>::value,
     ext::oneapi::detail::reduction_impl<
         T, BinaryOperation, 0, 1,
-        ext::oneapi::detail::default_reduction_algorithm<
-            access::placeholder::true_t, 1>,
+        ext::oneapi::detail::default_reduction_algorithm<1>,
         accessor<T, 1, access::mode::read_write, access::target::device,
                  access::placeholder::true_t,
                  ext::oneapi::accessor_property_list<>>>>
@@ -46,8 +45,7 @@ std::enable_if_t<
     !has_known_identity<BinaryOperation, T>::value,
     ext::oneapi::detail::reduction_impl<
         T, BinaryOperation, 0, 1,
-        ext::oneapi::detail::default_reduction_algorithm<
-            access::placeholder::true_t, 1>,
+        ext::oneapi::detail::default_reduction_algorithm<1>,
         accessor<T, 1, access::mode::read_write, access::target::device,
                  access::placeholder::true_t,
                  ext::oneapi::accessor_property_list<>>>>
@@ -67,9 +65,7 @@ template <typename T, typename BinaryOperation>
 std::enable_if_t<has_known_identity<BinaryOperation, T>::value,
                  ext::oneapi::detail::reduction_impl<
                      T, BinaryOperation, 0, 1,
-                     ext::oneapi::detail::default_reduction_algorithm<
-                         access::placeholder::false_t, 1>,
-                     T *>>
+                     ext::oneapi::detail::default_reduction_algorithm<1>, T *>>
 reduction(T *Var, BinaryOperation, const property_list &PropList = {}) {
   bool InitializeToIdentity =
       PropList.has_property<property::reduction::initialize_to_identity>();
@@ -85,9 +81,7 @@ template <typename T, typename BinaryOperation>
 std::enable_if_t<!has_known_identity<BinaryOperation, T>::value,
                  ext::oneapi::detail::reduction_impl<
                      T, BinaryOperation, 0, 1,
-                     ext::oneapi::detail::default_reduction_algorithm<
-                         access::placeholder::false_t, 1>,
-                     T *>>
+                     ext::oneapi::detail::default_reduction_algorithm<1>, T *>>
 reduction(T *, BinaryOperation, const property_list &PropList = {}) {
   // TODO: implement reduction that works even when identity is not known.
   (void)PropList;
@@ -102,8 +96,7 @@ reduction(T *, BinaryOperation, const property_list &PropList = {}) {
 template <typename T, typename AllocatorT, typename BinaryOperation>
 ext::oneapi::detail::reduction_impl<
     T, BinaryOperation, 0, 1,
-    ext::oneapi::detail::default_reduction_algorithm<
-        access::placeholder::true_t, 1>,
+    ext::oneapi::detail::default_reduction_algorithm<1>,
     accessor<T, 1, access::mode::read_write, access::target::device,
              access::placeholder::true_t,
              ext::oneapi::accessor_property_list<>>>
@@ -120,9 +113,7 @@ reduction(buffer<T, 1, AllocatorT> Var, handler &CGH, const T &Identity,
 template <typename T, typename BinaryOperation>
 ext::oneapi::detail::reduction_impl<
     T, BinaryOperation, 0, 1,
-    ext::oneapi::detail::default_reduction_algorithm<
-        access::placeholder::false_t, 1>,
-    T *>
+    ext::oneapi::detail::default_reduction_algorithm<1>, T *>
 reduction(T *Var, const T &Identity, BinaryOperation Combiner,
           const property_list &PropList = {}) {
   bool InitializeToIdentity =
@@ -139,9 +130,7 @@ std::enable_if_t<Extent != dynamic_extent &&
                      has_known_identity<BinaryOperation, T>::value,
                  ext::oneapi::detail::reduction_impl<
                      T, BinaryOperation, 1, Extent,
-                     ext::oneapi::detail::default_reduction_algorithm<
-                         access::placeholder::false_t, 1>,
-                     T *>>
+                     ext::oneapi::detail::default_reduction_algorithm<1>, T *>>
 reduction(span<T, Extent> Span, BinaryOperation,
           const property_list &PropList = {}) {
   bool InitializeToIdentity =
@@ -159,9 +148,7 @@ std::enable_if_t<Extent != dynamic_extent &&
                      !has_known_identity<BinaryOperation, T>::value,
                  ext::oneapi::detail::reduction_impl<
                      T, BinaryOperation, 1, Extent,
-                     ext::oneapi::detail::default_reduction_algorithm<
-                         access::placeholder::false_t, 1>,
-                     T *>>
+                     ext::oneapi::detail::default_reduction_algorithm<1>, T *>>
 reduction(span<T, Extent>, BinaryOperation,
           const property_list &PropList = {}) {
   // TODO: implement reduction that works even when identity is not known.
@@ -178,9 +165,7 @@ template <typename T, size_t Extent, typename BinaryOperation>
 std::enable_if_t<Extent != dynamic_extent,
                  ext::oneapi::detail::reduction_impl<
                      T, BinaryOperation, 1, Extent,
-                     ext::oneapi::detail::default_reduction_algorithm<
-                         access::placeholder::false_t, 1>,
-                     T *>>
+                     ext::oneapi::detail::default_reduction_algorithm<1>, T *>>
 reduction(span<T, Extent> Span, const T &Identity, BinaryOperation Combiner,
           const property_list &PropList = {}) {
   bool InitializeToIdentity =

From c15825bb900cd4aa154b77f9975cf6ed611d8b46 Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Thu, 14 Jul 2022 17:41:14 -0700
Subject: [PATCH 12/24] Remove default_reduction_algorithm

---
 sycl/include/sycl/ext/oneapi/reduction.hpp | 60 +++++++++-------------
 sycl/include/sycl/handler.hpp              |  4 +-
 sycl/include/sycl/reduction.hpp            | 44 ++++++----------
 3 files changed, 42 insertions(+), 66 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/reduction.hpp b/sycl/include/sycl/ext/oneapi/reduction.hpp
index 15cc1be9684ad..f926969a1508f 100644
--- a/sycl/include/sycl/ext/oneapi/reduction.hpp
+++ b/sycl/include/sycl/ext/oneapi/reduction.hpp
@@ -522,28 +522,24 @@ struct is_placeholder_t<accessor<T, AccessorDims, Mode, access::target::device,
                                  access::placeholder::true_t, PropList>>
     : public std::true_type {};
 
-/// Types representing specific reduction algorithms
-/// Enables reduction_impl_algo to take additional algorithm-specific templates
-template <int AccessorDims>
-class default_reduction_algorithm {};
 
-/// Templated class for implementations of specific reduction algorithms
-template <typename T, class BinaryOperation, int Dims, size_t Extent,
-          class Algorithm, typename RedOutVar>
-class reduction_impl_algo;
+template <class T>
+struct accessor_dim_t {
+  static constexpr int value = 1;
+};
+
+template <class T, int AccessorDims, access::mode Mode,
+          access::placeholder IsPH, typename PropList>
+struct accessor_dim_t<
+    accessor<T, AccessorDims, Mode, access::target::device, IsPH, PropList>> {
+  static constexpr int value = AccessorDims;
+};
 
-/// Original reduction algorithm is the default. It supports both USM and
-/// accessors via a single class
 template <typename T, class BinaryOperation, int Dims, size_t Extent,
-          int AccessorDims, typename RedOutVar>
-class reduction_impl_algo<
-    T, BinaryOperation, Dims, Extent,
-    default_reduction_algorithm<AccessorDims>, RedOutVar>
-    : public reduction_impl_common<T, BinaryOperation> {
+          typename RedOutVar>
+class reduction_impl_algo : public reduction_impl_common<T, BinaryOperation> {
   using base = reduction_impl_common<T, BinaryOperation>;
-  using self =
-      reduction_impl_algo<T, BinaryOperation, Dims, Extent,
-                          default_reduction_algorithm<AccessorDims>, RedOutVar>;
+  using self = reduction_impl_algo<T, BinaryOperation, Dims, Extent, RedOutVar>;
 
 public:
   using reducer_type = reducer<T, BinaryOperation, Dims, Extent>;
@@ -553,12 +549,12 @@ class reduction_impl_algo<
   // Buffers and accessors always describe scalar reductions (i.e. Dims == 0)
   // The input buffer/accessor is allowed to have different dimensionality
   // AccessorDims also determines the dimensionality of some temp storage
-  static constexpr int accessor_dim = AccessorDims;
-  static constexpr int buffer_dim = (AccessorDims == 0) ? 1 : AccessorDims;
+  static constexpr int accessor_dim = accessor_dim_t<RedOutVar>::value;
+  static constexpr int buffer_dim = (accessor_dim == 0) ? 1 : accessor_dim;
   static constexpr access::placeholder is_placeholder =
       is_placeholder_t<RedOutVar>::value ? access::placeholder::true_t
                                          : access::placeholder::false_t;
-  using rw_accessor_type = accessor<T, AccessorDims, access::mode::read_write,
+  using rw_accessor_type = accessor<T, accessor_dim, access::mode::read_write,
                                     access::target::device, is_placeholder,
                                     ext::oneapi::accessor_property_list<>>;
   static constexpr bool has_atomic_add_float64 =
@@ -746,16 +742,13 @@ template <typename T> struct AreAllButLastReductions<T> {
 /// This class encapsulates the reduction variable/accessor,
 /// the reduction operator and an optional operator identity.
 template <typename T, class BinaryOperation, int Dims, size_t Extent,
-          class Algorithm, typename RedOutVar>
+          typename RedOutVar>
 class reduction_impl
     : private reduction_impl_base,
-      public reduction_impl_algo<T, BinaryOperation, Dims, Extent, Algorithm,
-                                 RedOutVar> {
+      public reduction_impl_algo<T, BinaryOperation, Dims, Extent, RedOutVar> {
 private:
-  using algo = reduction_impl_algo<T, BinaryOperation, Dims, Extent,
-                                   Algorithm, RedOutVar>;
-  using self =
-      reduction_impl<T, BinaryOperation, Dims, Extent, Algorithm, RedOutVar>;
+  using algo = reduction_impl_algo<T, BinaryOperation, Dims, Extent, RedOutVar>;
+  using self = reduction_impl<T, BinaryOperation, Dims, Extent, RedOutVar>;
 
   static constexpr bool is_known_identity =
       sycl::detail::IsKnownIdentityOp<T, BinaryOperation>::value;
@@ -2466,7 +2459,6 @@ tuple_select_elements(TupleT Tuple, std::index_sequence<Is...>) {
 template <typename T, class BinaryOperation, int Dims, access::mode AccMode,
           access::placeholder IsPH>
 detail::reduction_impl<T, BinaryOperation, 0, 1,
-                       detail::default_reduction_algorithm<Dims>,
                        accessor<T, Dims, AccMode, access::target::device, IsPH>>
 reduction(accessor<T, Dims, AccMode, access::target::device, IsPH> &Acc,
           const T &Identity, BinaryOperation BOp) {
@@ -2482,7 +2474,6 @@ template <typename T, class BinaryOperation, int Dims, access::mode AccMode,
 std::enable_if_t<sycl::detail::IsKnownIdentityOp<T, BinaryOperation>::value,
                  detail::reduction_impl<
                      T, BinaryOperation, 0, 1,
-                     detail::default_reduction_algorithm<Dims>,
                      accessor<T, Dims, AccMode, access::target::device, IsPH>>>
 reduction(accessor<T, Dims, AccMode, access::target::device, IsPH> &Acc,
           BinaryOperation) {
@@ -2494,8 +2485,7 @@ reduction(accessor<T, Dims, AccMode, access::target::device, IsPH> &Acc,
 /// the computed reduction must be stored \param VarPtr, identity value
 /// \param Identity, and the binary operation used in the reduction.
 template <typename T, class BinaryOperation>
-detail::reduction_impl<T, BinaryOperation, 0, 1,
-                       detail::default_reduction_algorithm<1>, T *>
+detail::reduction_impl<T, BinaryOperation, 0, 1, T *>
 reduction(T *VarPtr, const T &Identity, BinaryOperation BOp) {
   return {VarPtr, Identity, BOp};
 }
@@ -2506,10 +2496,8 @@ reduction(T *VarPtr, const T &Identity, BinaryOperation BOp) {
 /// operation used in the reduction.
 /// The identity value is not passed to this version as it is statically known.
 template <typename T, class BinaryOperation>
-std::enable_if_t<
-    sycl::detail::IsKnownIdentityOp<T, BinaryOperation>::value,
-    detail::reduction_impl<T, BinaryOperation, 0, 1,
-                           detail::default_reduction_algorithm<1>, T *>>
+std::enable_if_t<sycl::detail::IsKnownIdentityOp<T, BinaryOperation>::value,
+                 detail::reduction_impl<T, BinaryOperation, 0, 1, T *>>
 reduction(T *VarPtr, BinaryOperation) {
   return {VarPtr};
 }
diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp
index 220c2ea1122f3..789a6c94aba84 100644
--- a/sycl/include/sycl/handler.hpp
+++ b/sycl/include/sycl/handler.hpp
@@ -242,7 +242,7 @@ namespace ext {
 namespace oneapi {
 namespace detail {
 template <typename T, class BinaryOperation, int Dims, size_t Extent,
-          class Algorithm, typename RedOutVar>
+          typename RedOutVar>
 class reduction_impl_algo;
 
 using cl::sycl::detail::enable_if_t;
@@ -2667,7 +2667,7 @@ class __SYCL_EXPORT handler {
   // Make reduction friends to store buffers and arrays created for it
   // in handler from reduction methods.
   template <typename T, class BinaryOperation, int Dims, size_t Extent,
-            class Algorithm, typename RedOutVar>
+            typename RedOutVar>
   friend class ext::oneapi::detail::reduction_impl_algo;
 
 #ifndef __SYCL_DEVICE_ONLY__
diff --git a/sycl/include/sycl/reduction.hpp b/sycl/include/sycl/reduction.hpp
index 0057c12cd31e9..e2b4facd83142 100644
--- a/sycl/include/sycl/reduction.hpp
+++ b/sycl/include/sycl/reduction.hpp
@@ -25,7 +25,6 @@ std::enable_if_t<
     has_known_identity<BinaryOperation, T>::value,
     ext::oneapi::detail::reduction_impl<
         T, BinaryOperation, 0, 1,
-        ext::oneapi::detail::default_reduction_algorithm<1>,
         accessor<T, 1, access::mode::read_write, access::target::device,
                  access::placeholder::true_t,
                  ext::oneapi::accessor_property_list<>>>>
@@ -45,7 +44,6 @@ std::enable_if_t<
     !has_known_identity<BinaryOperation, T>::value,
     ext::oneapi::detail::reduction_impl<
         T, BinaryOperation, 0, 1,
-        ext::oneapi::detail::default_reduction_algorithm<1>,
         accessor<T, 1, access::mode::read_write, access::target::device,
                  access::placeholder::true_t,
                  ext::oneapi::accessor_property_list<>>>>
@@ -62,10 +60,9 @@ reduction(buffer<T, 1, AllocatorT>, handler &, BinaryOperation,
 /// the given USM pointer \p Var, handler \p CGH, reduction operation
 /// \p Combiner, and optional reduction properties.
 template <typename T, typename BinaryOperation>
-std::enable_if_t<has_known_identity<BinaryOperation, T>::value,
-                 ext::oneapi::detail::reduction_impl<
-                     T, BinaryOperation, 0, 1,
-                     ext::oneapi::detail::default_reduction_algorithm<1>, T *>>
+std::enable_if_t<
+    has_known_identity<BinaryOperation, T>::value,
+    ext::oneapi::detail::reduction_impl<T, BinaryOperation, 0, 1, T *>>
 reduction(T *Var, BinaryOperation, const property_list &PropList = {}) {
   bool InitializeToIdentity =
       PropList.has_property<property::reduction::initialize_to_identity>();
@@ -78,10 +75,9 @@ reduction(T *Var, BinaryOperation, const property_list &PropList = {}) {
 /// The reduction algorithm may be less efficient for this variant as the
 /// reduction identity is not known statically and it is not provided by user.
 template <typename T, typename BinaryOperation>
-std::enable_if_t<!has_known_identity<BinaryOperation, T>::value,
-                 ext::oneapi::detail::reduction_impl<
-                     T, BinaryOperation, 0, 1,
-                     ext::oneapi::detail::default_reduction_algorithm<1>, T *>>
+std::enable_if_t<
+    !has_known_identity<BinaryOperation, T>::value,
+    ext::oneapi::detail::reduction_impl<T, BinaryOperation, 0, 1, T *>>
 reduction(T *, BinaryOperation, const property_list &PropList = {}) {
   // TODO: implement reduction that works even when identity is not known.
   (void)PropList;
@@ -96,7 +92,6 @@ reduction(T *, BinaryOperation, const property_list &PropList = {}) {
 template <typename T, typename AllocatorT, typename BinaryOperation>
 ext::oneapi::detail::reduction_impl<
     T, BinaryOperation, 0, 1,
-    ext::oneapi::detail::default_reduction_algorithm<1>,
     accessor<T, 1, access::mode::read_write, access::target::device,
              access::placeholder::true_t,
              ext::oneapi::accessor_property_list<>>>
@@ -111,9 +106,7 @@ reduction(buffer<T, 1, AllocatorT> Var, handler &CGH, const T &Identity,
 /// the given USM pointer \p Var, reduction identity value \p Identity,
 /// binary operation \p Combiner, and optional reduction properties.
 template <typename T, typename BinaryOperation>
-ext::oneapi::detail::reduction_impl<
-    T, BinaryOperation, 0, 1,
-    ext::oneapi::detail::default_reduction_algorithm<1>, T *>
+ext::oneapi::detail::reduction_impl<T, BinaryOperation, 0, 1, T *>
 reduction(T *Var, const T &Identity, BinaryOperation Combiner,
           const property_list &PropList = {}) {
   bool InitializeToIdentity =
@@ -126,11 +119,9 @@ reduction(T *Var, const T &Identity, BinaryOperation Combiner,
 /// the given sycl::span \p Span, reduction operation \p Combiner, and
 /// optional reduction properties.
 template <typename T, size_t Extent, typename BinaryOperation>
-std::enable_if_t<Extent != dynamic_extent &&
-                     has_known_identity<BinaryOperation, T>::value,
-                 ext::oneapi::detail::reduction_impl<
-                     T, BinaryOperation, 1, Extent,
-                     ext::oneapi::detail::default_reduction_algorithm<1>, T *>>
+std::enable_if_t<
+    Extent != dynamic_extent && has_known_identity<BinaryOperation, T>::value,
+    ext::oneapi::detail::reduction_impl<T, BinaryOperation, 1, Extent, T *>>
 reduction(span<T, Extent> Span, BinaryOperation,
           const property_list &PropList = {}) {
   bool InitializeToIdentity =
@@ -144,11 +135,9 @@ reduction(span<T, Extent> Span, BinaryOperation,
 /// The reduction algorithm may be less efficient for this variant as the
 /// reduction identity is not known statically and it is not provided by user.
 template <typename T, size_t Extent, typename BinaryOperation>
-std::enable_if_t<Extent != dynamic_extent &&
-                     !has_known_identity<BinaryOperation, T>::value,
-                 ext::oneapi::detail::reduction_impl<
-                     T, BinaryOperation, 1, Extent,
-                     ext::oneapi::detail::default_reduction_algorithm<1>, T *>>
+std::enable_if_t<
+    Extent != dynamic_extent && !has_known_identity<BinaryOperation, T>::value,
+    ext::oneapi::detail::reduction_impl<T, BinaryOperation, 1, Extent, T *>>
 reduction(span<T, Extent>, BinaryOperation,
           const property_list &PropList = {}) {
   // TODO: implement reduction that works even when identity is not known.
@@ -162,10 +151,9 @@ reduction(span<T, Extent>, BinaryOperation,
 /// the given sycl::span \p Span, reduction identity value \p Identity,
 /// reduction operation \p Combiner, and optional reduction properties.
 template <typename T, size_t Extent, typename BinaryOperation>
-std::enable_if_t<Extent != dynamic_extent,
-                 ext::oneapi::detail::reduction_impl<
-                     T, BinaryOperation, 1, Extent,
-                     ext::oneapi::detail::default_reduction_algorithm<1>, T *>>
+std::enable_if_t<
+    Extent != dynamic_extent,
+    ext::oneapi::detail::reduction_impl<T, BinaryOperation, 1, Extent, T *>>
 reduction(span<T, Extent> Span, const T &Identity, BinaryOperation Combiner,
           const property_list &PropList = {}) {
   bool InitializeToIdentity =

From add7ff24a8a9bde467b91040aa77816a6e83cee9 Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Thu, 14 Jul 2022 21:56:34 -0700
Subject: [PATCH 13/24] make_reduction for partial types deduction + process
 buffer/span on the caller's side

---
 sycl/include/sycl/ext/oneapi/reduction.hpp |  91 ++++++++--------
 sycl/include/sycl/reduction.hpp            | 114 ++++++++++-----------
 2 files changed, 101 insertions(+), 104 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/reduction.hpp b/sycl/include/sycl/ext/oneapi/reduction.hpp
index f926969a1508f..7d72432d77ad3 100644
--- a/sycl/include/sycl/ext/oneapi/reduction.hpp
+++ b/sycl/include/sycl/ext/oneapi/reduction.hpp
@@ -535,6 +535,18 @@ struct accessor_dim_t<
   static constexpr int value = AccessorDims;
 };
 
+template <class T> struct get_red_t;
+template <class T> struct get_red_t<T*> {
+  using type = T;
+};
+
+template <class T, int AccessorDims, access::mode Mode,
+          access::placeholder IsPH, typename PropList>
+struct get_red_t<
+    accessor<T, AccessorDims, Mode, access::target::device, IsPH, PropList>> {
+  using type = T;
+};
+
 template <typename T, class BinaryOperation, int Dims, size_t Extent,
           typename RedOutVar>
 class reduction_impl_algo : public reduction_impl_common<T, BinaryOperation> {
@@ -788,15 +800,14 @@ class reduction_impl
 
   /// SYCL-2020.
   /// Constructs reduction_impl when the identity value is statically known.
-  template <typename AllocatorT, typename _self = self,
+  template <typename _self = self,
             std::enable_if_t<_self::is_known_identity && _self::my_is_rw_acc>
                 * = nullptr>
-  reduction_impl(buffer<T, 1, AllocatorT> Buffer, handler &CGH,
-                 bool InitializeToIdentity)
+  reduction_impl(RedOutVar &Acc, handler &CGH, bool InitializeToIdentity)
       : algo(reducer_type::getIdentity(), BinaryOperation(),
-             InitializeToIdentity, rw_accessor_type{Buffer}) {
+             InitializeToIdentity, Acc) {
     algo::associateWithHandler(CGH);
-    if (Buffer.size() != 1)
+    if (Acc.size() != 1)
       throw sycl::runtime_error(errc::invalid,
                                 "Reduction variable must be a scalar.",
                                 PI_ERROR_INVALID_VALUE);
@@ -819,15 +830,12 @@ class reduction_impl
   /// and user still passed the identity value.
   /// SYCL-2020.
   /// Constructs reduction_impl when the identity value is NOT known statically.
-  template <typename AllocatorT, typename _self = self,
-            enable_if_t<_self::my_is_rw_acc> * = nullptr>
-  reduction_impl(buffer<T, 1, AllocatorT> Buffer, handler &CGH,
-                 const T &Identity, BinaryOperation BOp,
-                 bool InitializeToIdentity)
-      : algo(chooseIdentity(Identity), BOp, InitializeToIdentity,
-             rw_accessor_type{Buffer}) {
+  template <typename _self = self, enable_if_t<_self::my_is_rw_acc> * = nullptr>
+  reduction_impl(RedOutVar &Acc, handler &CGH, const T &Identity,
+                 BinaryOperation BOp, bool InitializeToIdentity)
+      : algo(chooseIdentity(Identity), BOp, InitializeToIdentity, Acc) {
     algo::associateWithHandler(CGH);
-    if (Buffer.size() != 1)
+    if (Acc.size() != 1)
       throw sycl::runtime_error(errc::invalid,
                                 "Reduction variable must be a scalar.",
                                 PI_ERROR_INVALID_VALUE);
@@ -861,21 +869,15 @@ class reduction_impl
   reduction_impl(T *VarPtr, const T &Identity, BinaryOperation BOp,
                  bool InitializeToIdentity = false)
       : algo(chooseIdentity(Identity), BOp, InitializeToIdentity, VarPtr) {}
-
-  /// Constructs reduction_impl when the identity value is statically known
-  template <typename _self = self, enable_if_t<_self::is_known_identity &&
-                                               _self::my_is_usm> * = nullptr>
-  reduction_impl(span<T, Extent> Span, bool InitializeToIdentity = false)
-      : algo(reducer_type::getIdentity(), BinaryOperation(),
-             InitializeToIdentity, Span.data()) {}
-
-  template <typename _self = self, enable_if_t<_self::my_is_usm> * = nullptr>
-  reduction_impl(span<T, Extent> Span, const T &Identity, BinaryOperation BOp,
-                 bool InitializeToIdentity = false)
-      : algo(chooseIdentity(Identity), BOp, InitializeToIdentity, Span.data()) {
-  }
 };
 
+template <class BinaryOp, int Dims, size_t Extent, typename RedOutVar,
+          typename... RestTy>
+auto make_reduction(RedOutVar RedVar, RestTy &&... Rest) {
+  return reduction_impl<typename get_red_t<RedOutVar>::type, BinaryOp, Dims,
+                        Extent, RedOutVar>{RedVar, std::forward<RestTy>(Rest)...};
+}
+
 /// A helper to pass undefined (sycl::detail::auto_name) names unmodified. We
 /// must do that to avoid name collisions.
 template <template <typename...> class Namer, class KernelName, class... Ts>
@@ -2458,11 +2460,9 @@ tuple_select_elements(TupleT Tuple, std::index_sequence<Is...>) {
 /// operation used in the reduction.
 template <typename T, class BinaryOperation, int Dims, access::mode AccMode,
           access::placeholder IsPH>
-detail::reduction_impl<T, BinaryOperation, 0, 1,
-                       accessor<T, Dims, AccMode, access::target::device, IsPH>>
-reduction(accessor<T, Dims, AccMode, access::target::device, IsPH> &Acc,
-          const T &Identity, BinaryOperation BOp) {
-  return {Acc, Identity, BOp};
+auto reduction(accessor<T, Dims, AccMode, access::target::device, IsPH> &Acc,
+               const T &Identity, BinaryOperation BOp) {
+  return detail::make_reduction<BinaryOperation, 0, 1>(Acc, Identity, BOp);
 }
 
 /// Creates and returns an object implementing the reduction functionality.
@@ -2470,14 +2470,12 @@ reduction(accessor<T, Dims, AccMode, access::target::device, IsPH> &Acc,
 /// must be stored \param Acc and the binary operation used in the reduction.
 /// The identity value is not passed to this version as it is statically known.
 template <typename T, class BinaryOperation, int Dims, access::mode AccMode,
-          access::placeholder IsPH>
-std::enable_if_t<sycl::detail::IsKnownIdentityOp<T, BinaryOperation>::value,
-                 detail::reduction_impl<
-                     T, BinaryOperation, 0, 1,
-                     accessor<T, Dims, AccMode, access::target::device, IsPH>>>
-reduction(accessor<T, Dims, AccMode, access::target::device, IsPH> &Acc,
-          BinaryOperation) {
-  return {Acc};
+          access::placeholder IsPH,
+          typename = std::enable_if_t<
+              sycl::detail::IsKnownIdentityOp<T, BinaryOperation>::value>>
+auto reduction(accessor<T, Dims, AccMode, access::target::device, IsPH> &Acc,
+               BinaryOperation) {
+  return detail::make_reduction<BinaryOperation, 0, 1>(Acc);
 }
 
 /// Creates and returns an object implementing the reduction functionality.
@@ -2485,9 +2483,8 @@ reduction(accessor<T, Dims, AccMode, access::target::device, IsPH> &Acc,
 /// the computed reduction must be stored \param VarPtr, identity value
 /// \param Identity, and the binary operation used in the reduction.
 template <typename T, class BinaryOperation>
-detail::reduction_impl<T, BinaryOperation, 0, 1, T *>
-reduction(T *VarPtr, const T &Identity, BinaryOperation BOp) {
-  return {VarPtr, Identity, BOp};
+auto reduction(T *VarPtr, const T &Identity, BinaryOperation BOp) {
+  return detail::make_reduction<BinaryOperation, 0, 1>(VarPtr, Identity, BOp);
 }
 
 /// Creates and returns an object implementing the reduction functionality.
@@ -2495,11 +2492,11 @@ reduction(T *VarPtr, const T &Identity, BinaryOperation BOp) {
 /// the computed reduction must be stored \param VarPtr, and the binary
 /// operation used in the reduction.
 /// The identity value is not passed to this version as it is statically known.
-template <typename T, class BinaryOperation>
-std::enable_if_t<sycl::detail::IsKnownIdentityOp<T, BinaryOperation>::value,
-                 detail::reduction_impl<T, BinaryOperation, 0, 1, T *>>
-reduction(T *VarPtr, BinaryOperation) {
-  return {VarPtr};
+template <typename T, class BinaryOperation,
+          typename = std::enable_if_t<
+              sycl::detail::IsKnownIdentityOp<T, BinaryOperation>::value>>
+auto reduction(T *VarPtr, BinaryOperation) {
+  return detail::make_reduction<BinaryOperation, 0, 1>(VarPtr);
 }
 
 // ---- has_known_identity
diff --git a/sycl/include/sycl/reduction.hpp b/sycl/include/sycl/reduction.hpp
index e2b4facd83142..64c054a9660a6 100644
--- a/sycl/include/sycl/reduction.hpp
+++ b/sycl/include/sycl/reduction.hpp
@@ -20,33 +20,32 @@ namespace sycl {
 
 /// Constructs a reduction object using the given buffer \p Var, handler \p CGH,
 /// reduction operation \p Combiner, and optional reduction properties.
-template <typename T, typename AllocatorT, typename BinaryOperation>
-std::enable_if_t<
-    has_known_identity<BinaryOperation, T>::value,
-    ext::oneapi::detail::reduction_impl<
-        T, BinaryOperation, 0, 1,
-        accessor<T, 1, access::mode::read_write, access::target::device,
-                 access::placeholder::true_t,
-                 ext::oneapi::accessor_property_list<>>>>
-reduction(buffer<T, 1, AllocatorT> Var, handler &CGH, BinaryOperation,
-          const property_list &PropList = {}) {
+template <
+    typename T, typename AllocatorT, typename BinaryOperation,
+    typename = std::enable_if_t<has_known_identity<BinaryOperation, T>::value>>
+auto reduction(buffer<T, 1, AllocatorT> Var, handler &CGH, BinaryOperation,
+               const property_list &PropList = {}) {
   bool InitializeToIdentity =
       PropList.has_property<property::reduction::initialize_to_identity>();
-  return {Var, CGH, InitializeToIdentity};
+  return ext::oneapi::detail::make_reduction<BinaryOperation, 0, 1>(
+      accessor<T, 1, access::mode::read_write, access::target::device,
+               access::placeholder::true_t,
+               ext::oneapi::accessor_property_list<>>{Var},
+      CGH, InitializeToIdentity);
 }
 
 /// Constructs a reduction object using the given buffer \p Var, handler \p CGH,
 /// reduction operation \p Combiner, and optional reduction properties.
 /// The reduction algorithm may be less efficient for this variant as the
 /// reduction identity is not known statically and it is not provided by user.
-template <typename T, typename AllocatorT, typename BinaryOperation>
-std::enable_if_t<
-    !has_known_identity<BinaryOperation, T>::value,
-    ext::oneapi::detail::reduction_impl<
-        T, BinaryOperation, 0, 1,
-        accessor<T, 1, access::mode::read_write, access::target::device,
-                 access::placeholder::true_t,
-                 ext::oneapi::accessor_property_list<>>>>
+template <
+    typename T, typename AllocatorT, typename BinaryOperation,
+    typename = std::enable_if_t<!has_known_identity<BinaryOperation, T>::value>>
+ext::oneapi::detail::reduction_impl<
+    T, BinaryOperation, 0, 1,
+    accessor<T, 1, access::mode::read_write, access::target::device,
+             access::placeholder::true_t,
+             ext::oneapi::accessor_property_list<>>>
 reduction(buffer<T, 1, AllocatorT>, handler &, BinaryOperation,
           const property_list &PropList = {}) {
   // TODO: implement reduction that works even when identity is not known.
@@ -59,14 +58,14 @@ reduction(buffer<T, 1, AllocatorT>, handler &, BinaryOperation,
 /// Constructs a reduction object using the reduction variable referenced by
 /// the given USM pointer \p Var, handler \p CGH, reduction operation
 /// \p Combiner, and optional reduction properties.
-template <typename T, typename BinaryOperation>
-std::enable_if_t<
-    has_known_identity<BinaryOperation, T>::value,
-    ext::oneapi::detail::reduction_impl<T, BinaryOperation, 0, 1, T *>>
-reduction(T *Var, BinaryOperation, const property_list &PropList = {}) {
+template <
+    typename T, typename BinaryOperation,
+    typename = std::enable_if_t<has_known_identity<BinaryOperation, T>::value>>
+auto reduction(T *Var, BinaryOperation, const property_list &PropList = {}) {
   bool InitializeToIdentity =
       PropList.has_property<property::reduction::initialize_to_identity>();
-  return {Var, InitializeToIdentity};
+  return ext::oneapi::detail::make_reduction<BinaryOperation, 0, 1>(
+      Var, InitializeToIdentity);
 }
 
 /// Constructs a reduction object using the reduction variable referenced by
@@ -74,10 +73,10 @@ reduction(T *Var, BinaryOperation, const property_list &PropList = {}) {
 /// \p Combiner, and optional reduction properties.
 /// The reduction algorithm may be less efficient for this variant as the
 /// reduction identity is not known statically and it is not provided by user.
-template <typename T, typename BinaryOperation>
-std::enable_if_t<
-    !has_known_identity<BinaryOperation, T>::value,
-    ext::oneapi::detail::reduction_impl<T, BinaryOperation, 0, 1, T *>>
+template <
+    typename T, typename BinaryOperation,
+    typename = std::enable_if_t<!has_known_identity<BinaryOperation, T>::value>>
+ext::oneapi::detail::reduction_impl<T, BinaryOperation, 0, 1, T *>
 reduction(T *, BinaryOperation, const property_list &PropList = {}) {
   // TODO: implement reduction that works even when identity is not known.
   (void)PropList;
@@ -90,43 +89,44 @@ reduction(T *, BinaryOperation, const property_list &PropList = {}) {
 /// reduction identity value \p Identity, reduction operation \p Combiner,
 /// and optional reduction properties.
 template <typename T, typename AllocatorT, typename BinaryOperation>
-ext::oneapi::detail::reduction_impl<
-    T, BinaryOperation, 0, 1,
-    accessor<T, 1, access::mode::read_write, access::target::device,
-             access::placeholder::true_t,
-             ext::oneapi::accessor_property_list<>>>
+auto
 reduction(buffer<T, 1, AllocatorT> Var, handler &CGH, const T &Identity,
           BinaryOperation Combiner, const property_list &PropList = {}) {
   bool InitializeToIdentity =
       PropList.has_property<property::reduction::initialize_to_identity>();
-  return {Var, CGH, Identity, Combiner, InitializeToIdentity};
+  return ext::oneapi::detail::make_reduction<BinaryOperation, 0, 1>(
+      accessor<T, 1, access::mode::read_write, access::target::device,
+               access::placeholder::true_t,
+               ext::oneapi::accessor_property_list<>>{Var},
+      CGH, Identity, Combiner, InitializeToIdentity);
 }
 
 /// Constructs a reduction object using the reduction variable referenced by
 /// the given USM pointer \p Var, reduction identity value \p Identity,
 /// binary operation \p Combiner, and optional reduction properties.
 template <typename T, typename BinaryOperation>
-ext::oneapi::detail::reduction_impl<T, BinaryOperation, 0, 1, T *>
-reduction(T *Var, const T &Identity, BinaryOperation Combiner,
-          const property_list &PropList = {}) {
+auto reduction(T *Var, const T &Identity, BinaryOperation Combiner,
+               const property_list &PropList = {}) {
   bool InitializeToIdentity =
       PropList.has_property<property::reduction::initialize_to_identity>();
-  return {Var, Identity, Combiner, InitializeToIdentity};
+  return ext::oneapi::detail::make_reduction<BinaryOperation, 0, 1>(
+      Var, Identity, Combiner, InitializeToIdentity);
 }
 
 #if __cplusplus >= 201703L
 /// Constructs a reduction object using the reduction variable referenced by
 /// the given sycl::span \p Span, reduction operation \p Combiner, and
 /// optional reduction properties.
-template <typename T, size_t Extent, typename BinaryOperation>
-std::enable_if_t<
-    Extent != dynamic_extent && has_known_identity<BinaryOperation, T>::value,
-    ext::oneapi::detail::reduction_impl<T, BinaryOperation, 1, Extent, T *>>
-reduction(span<T, Extent> Span, BinaryOperation,
-          const property_list &PropList = {}) {
+template <
+    typename T, size_t Extent, typename BinaryOperation,
+    typename = std::enable_if_t<Extent != dynamic_extent &&
+                                has_known_identity<BinaryOperation, T>::value>>
+auto reduction(span<T, Extent> Span, BinaryOperation,
+               const property_list &PropList = {}) {
   bool InitializeToIdentity =
       PropList.has_property<property::reduction::initialize_to_identity>();
-  return {Span, InitializeToIdentity};
+  return ext::oneapi::detail::make_reduction<BinaryOperation, 1, Extent>(
+      Span.data(), InitializeToIdentity);
 }
 
 /// Constructs a reduction object using the reduction variable referenced by
@@ -134,10 +134,11 @@ reduction(span<T, Extent> Span, BinaryOperation,
 /// optional reduction properties.
 /// The reduction algorithm may be less efficient for this variant as the
 /// reduction identity is not known statically and it is not provided by user.
-template <typename T, size_t Extent, typename BinaryOperation>
-std::enable_if_t<
-    Extent != dynamic_extent && !has_known_identity<BinaryOperation, T>::value,
-    ext::oneapi::detail::reduction_impl<T, BinaryOperation, 1, Extent, T *>>
+template <
+    typename T, size_t Extent, typename BinaryOperation,
+    typename = std::enable_if_t<Extent != dynamic_extent &&
+                                !has_known_identity<BinaryOperation, T>::value>>
+ext::oneapi::detail::reduction_impl<T, BinaryOperation, 1, Extent, T *>
 reduction(span<T, Extent>, BinaryOperation,
           const property_list &PropList = {}) {
   // TODO: implement reduction that works even when identity is not known.
@@ -150,15 +151,14 @@ reduction(span<T, Extent>, BinaryOperation,
 /// Constructs a reduction object using the reduction variable referenced by
 /// the given sycl::span \p Span, reduction identity value \p Identity,
 /// reduction operation \p Combiner, and optional reduction properties.
-template <typename T, size_t Extent, typename BinaryOperation>
-std::enable_if_t<
-    Extent != dynamic_extent,
-    ext::oneapi::detail::reduction_impl<T, BinaryOperation, 1, Extent, T *>>
-reduction(span<T, Extent> Span, const T &Identity, BinaryOperation Combiner,
-          const property_list &PropList = {}) {
+template <typename T, size_t Extent, typename BinaryOperation,
+          typename = std::enable_if_t<Extent != dynamic_extent>>
+auto reduction(span<T, Extent> Span, const T &Identity,
+               BinaryOperation Combiner, const property_list &PropList = {}) {
   bool InitializeToIdentity =
       PropList.has_property<property::reduction::initialize_to_identity>();
-  return {Span, Identity, Combiner, InitializeToIdentity};
+  return ext::oneapi::detail::make_reduction<BinaryOperation, 1, Extent>(
+      Span.data(), Identity, Combiner, InitializeToIdentity);
 }
 #endif
 

From 4eff2bb5b01d30dbdde40090a7192544234ac293 Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Fri, 15 Jul 2022 09:50:00 -0700
Subject: [PATCH 14/24] Sort ctors, minor change

---
 sycl/include/sycl/ext/oneapi/reduction.hpp | 64 ++++++++++------------
 1 file changed, 30 insertions(+), 34 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/reduction.hpp b/sycl/include/sycl/ext/oneapi/reduction.hpp
index 7d72432d77ad3..8298b357dbae6 100644
--- a/sycl/include/sycl/ext/oneapi/reduction.hpp
+++ b/sycl/include/sycl/ext/oneapi/reduction.hpp
@@ -798,21 +798,6 @@ class reduction_impl
   // Only scalar and 1D array reductions are supported by SYCL 2020.
   static_assert(Dims <= 1, "Multi-dimensional reductions are not supported.");
 
-  /// SYCL-2020.
-  /// Constructs reduction_impl when the identity value is statically known.
-  template <typename _self = self,
-            std::enable_if_t<_self::is_known_identity && _self::my_is_rw_acc>
-                * = nullptr>
-  reduction_impl(RedOutVar &Acc, handler &CGH, bool InitializeToIdentity)
-      : algo(reducer_type::getIdentity(), BinaryOperation(),
-             InitializeToIdentity, Acc) {
-    algo::associateWithHandler(CGH);
-    if (Acc.size() != 1)
-      throw sycl::runtime_error(errc::invalid,
-                                "Reduction variable must be a scalar.",
-                                PI_ERROR_INVALID_VALUE);
-  }
-
   /// Constructs reduction_impl when the identity value is statically known.
   template <typename _self = self,
             enable_if_t<_self::is_known_identity &&
@@ -825,15 +810,24 @@ class reduction_impl
                                 PI_ERROR_INVALID_VALUE);
   }
 
+  /// Constructs reduction_impl when the identity value is statically known.
+  /// The \param VarPtr is a USM pointer to memory, to where the computed
+  /// reduction value is added using BinaryOperation, i.e. it is expected that
+  /// the memory is pre-initialized with some meaningful value.
+  template <typename _self = self, enable_if_t<_self::is_known_identity &&
+                                               _self::my_is_usm> * = nullptr>
+  reduction_impl(RedOutVar VarPtr, bool InitializeToIdentity = false)
+      : algo(reducer_type::getIdentity(), BinaryOperation(),
+             InitializeToIdentity, VarPtr) {}
+
   /// SYCL-2020.
-  /// Constructs reduction_impl when the identity value is statically known,
-  /// and user still passed the identity value.
-  /// SYCL-2020.
-  /// Constructs reduction_impl when the identity value is NOT known statically.
-  template <typename _self = self, enable_if_t<_self::my_is_rw_acc> * = nullptr>
-  reduction_impl(RedOutVar &Acc, handler &CGH, const T &Identity,
-                 BinaryOperation BOp, bool InitializeToIdentity)
-      : algo(chooseIdentity(Identity), BOp, InitializeToIdentity, Acc) {
+  /// Constructs reduction_impl when the identity value is statically known.
+  template <typename _self = self,
+            std::enable_if_t<_self::is_known_identity && _self::my_is_rw_acc>
+                * = nullptr>
+  reduction_impl(RedOutVar &Acc, handler &CGH, bool InitializeToIdentity)
+      : algo(reducer_type::getIdentity(), BinaryOperation(),
+             InitializeToIdentity, Acc) {
     algo::associateWithHandler(CGH);
     if (Acc.size() != 1)
       throw sycl::runtime_error(errc::invalid,
@@ -852,23 +846,25 @@ class reduction_impl
                                 PI_ERROR_INVALID_VALUE);
   }
 
-  /// Constructs reduction_impl when the identity value is statically known.
-  /// The \param VarPtr is a USM pointer to memory, to where the computed
-  /// reduction value is added using BinaryOperation, i.e. it is expected that
-  /// the memory is pre-initialized with some meaningful value.
-  template <typename _self = self, enable_if_t<_self::is_known_identity &&
-                                               _self::my_is_usm> * = nullptr>
-  reduction_impl(T *VarPtr, bool InitializeToIdentity = false)
-      : algo(reducer_type::getIdentity(), BinaryOperation(),
-             InitializeToIdentity, VarPtr) {}
-
   /// The \param VarPtr is a USM pointer to memory, to where the computed
   /// reduction value is added using BinaryOperation, i.e. it is expected that
   /// the memory is pre-initialized with some meaningful value.
   template <typename _self = self, enable_if_t<_self::my_is_usm> * = nullptr>
-  reduction_impl(T *VarPtr, const T &Identity, BinaryOperation BOp,
+  reduction_impl(RedOutVar VarPtr, const T &Identity, BinaryOperation BOp,
                  bool InitializeToIdentity = false)
       : algo(chooseIdentity(Identity), BOp, InitializeToIdentity, VarPtr) {}
+
+  /// For placeholder accessor.
+  template <typename _self = self, enable_if_t<_self::my_is_rw_acc> * = nullptr>
+  reduction_impl(RedOutVar &Acc, handler &CGH, const T &Identity,
+                 BinaryOperation BOp, bool InitializeToIdentity)
+      : algo(chooseIdentity(Identity), BOp, InitializeToIdentity, Acc) {
+    algo::associateWithHandler(CGH);
+    if (Acc.size() != 1)
+      throw sycl::runtime_error(errc::invalid,
+                                "Reduction variable must be a scalar.",
+                                PI_ERROR_INVALID_VALUE);
+  }
 };
 
 template <class BinaryOp, int Dims, size_t Extent, typename RedOutVar,

From 630b58c02813a57ae4128676274b751600cfc07e Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Fri, 15 Jul 2022 09:50:52 -0700
Subject: [PATCH 15/24] Minor improvements

---
 sycl/include/sycl/ext/oneapi/reduction.hpp | 77 +++++++++++-----------
 sycl/include/sycl/handler.hpp              |  4 +-
 2 files changed, 41 insertions(+), 40 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/reduction.hpp b/sycl/include/sycl/ext/oneapi/reduction.hpp
index 8298b357dbae6..51d37709ca97c 100644
--- a/sycl/include/sycl/ext/oneapi/reduction.hpp
+++ b/sycl/include/sycl/ext/oneapi/reduction.hpp
@@ -576,11 +576,13 @@ class reduction_impl_algo : public reduction_impl_common<T, BinaryOperation> {
   static constexpr bool has_fast_reduce =
       IsReduOptForFastReduce<T, BinaryOperation>::value;
 
-  static constexpr bool my_is_usm = std::is_same_v<RedOutVar, T *>;
-  static constexpr bool is_usm = my_is_usm;
+  static constexpr bool is_usm = std::is_same_v<RedOutVar, T *>;
 
-  static constexpr bool my_is_rw_acc = is_rw_acc_t<RedOutVar>::value;
-  static constexpr bool my_is_dw_acc = is_dw_acc_t<RedOutVar>::value;
+  static constexpr bool is_rw_acc = is_rw_acc_t<RedOutVar>::value;
+  static constexpr bool is_dw_acc = is_dw_acc_t<RedOutVar>::value;
+  static constexpr bool is_acc = is_rw_acc | is_dw_acc;
+  static_assert(!is_rw_acc || !is_dw_acc, "Can be only one at once!");
+  static_assert(!is_usm || !is_acc, "Ca be only one at once!");
 
   static constexpr size_t dims = Dims;
   static constexpr size_t num_elements = Extent;
@@ -593,8 +595,7 @@ class reduction_impl_algo : public reduction_impl_common<T, BinaryOperation> {
   /// to keep the accessor alive until the command group finishes the work.
   /// This function does not do anything for USM reductions.
   void associateWithHandler(handler &CGH) {
-    if constexpr (!my_is_usm) {
-      static_assert(my_is_rw_acc || my_is_dw_acc);
+    if constexpr (is_acc) {
       CGH.associateWithHandler(&MRedOut, access::target::device);
     }
   }
@@ -629,7 +630,7 @@ class reduction_impl_algo : public reduction_impl_common<T, BinaryOperation> {
   template <bool IsOneWG, bool _IsUSM = is_usm>
   std::enable_if_t<IsOneWG && !_IsUSM, rw_accessor_type>
   getWriteMemForPartialReds(size_t, handler &CGH) {
-    if constexpr (my_is_rw_acc)
+    if constexpr (is_rw_acc)
       return MRedOut;
     return getWriteMemForPartialReds<false>(1, CGH);
   }
@@ -651,7 +652,7 @@ class reduction_impl_algo : public reduction_impl_common<T, BinaryOperation> {
   /// Otherwise, a new buffer is created and accessor to that buffer is
   /// returned.
   rw_accessor_type getWriteAccForPartialReds(size_t Size, handler &CGH) {
-    if constexpr (my_is_rw_acc) {
+    if constexpr (is_rw_acc) {
       if (Size == 1) {
         associateWithHandler(CGH);
         return MRedOut;
@@ -671,11 +672,11 @@ class reduction_impl_algo : public reduction_impl_common<T, BinaryOperation> {
   template <bool HasFastAtomics = (has_fast_atomics || has_atomic_add_float64)>
   std::enable_if_t<HasFastAtomics, rw_accessor_type>
   getReadWriteAccessorToInitializedMem(handler &CGH) {
-    if constexpr (my_is_rw_acc) {
+    if constexpr (is_rw_acc) {
       if (!base::initializeToIdentity())
         return MRedOut;
     }
-    assert(!(my_is_dw_acc && !base::initializeToIdentity()) &&
+    assert(!(is_dw_acc && !base::initializeToIdentity()) &&
            "Unexpected condition!");
 
     // TODO: Move to T[] in C++20 to simplify handling here
@@ -714,20 +715,21 @@ class reduction_impl_algo : public reduction_impl_common<T, BinaryOperation> {
   }
 
 private:
-  template <typename BufferT, typename _self = self>
-  std::enable_if_t<_self::is_placeholder == access::placeholder::false_t,
-                   rw_accessor_type>
-  createHandlerWiredReadWriteAccessor(handler &CGH, BufferT Buffer) {
-    return {Buffer, CGH};
-  }
-
-  template <typename BufferT, typename _self = self>
-  std::enable_if_t<_self::is_placeholder == access::placeholder::true_t,
-                   rw_accessor_type>
-  createHandlerWiredReadWriteAccessor(handler &CGH, BufferT Buffer) {
-    rw_accessor_type Acc(Buffer);
-    CGH.require(Acc);
-    return Acc;
+  template <typename BufferT>
+  rw_accessor_type createHandlerWiredReadWriteAccessor(handler &CGH,
+                                                       BufferT Buffer) {
+    // TODO:
+    // SYCL 2020: The accessor template parameter IsPlaceholder is allowed to be
+    // specified, but it has no bearing on whether the accessor instance is a
+    // placeholder. This is determined solely by the constructor used to create
+    // the instance. The associated type access::placeholder is also deprecated.
+    if constexpr (is_placeholder == access::placeholder::true_t) {
+      rw_accessor_type Acc(Buffer);
+      CGH.require(Acc);
+      return Acc;
+    } else {
+      return {Buffer, CGH};
+    }
   }
 
   std::shared_ptr<buffer<T, buffer_dim>> MOutBufPtr;
@@ -788,9 +790,10 @@ class reduction_impl
   }
 
 public:
-  using algo::my_is_usm;
-  using algo::my_is_rw_acc;
-  using algo::my_is_dw_acc;
+  using algo::is_usm;
+  using algo::is_rw_acc;
+  using algo::is_dw_acc;
+  using algo::is_acc;
 
   using reducer_type = typename algo::reducer_type;
   using rw_accessor_type = typename algo::rw_accessor_type;
@@ -800,8 +803,7 @@ class reduction_impl
 
   /// Constructs reduction_impl when the identity value is statically known.
   template <typename _self = self,
-            enable_if_t<_self::is_known_identity &&
-                        (my_is_rw_acc || my_is_dw_acc)> * = nullptr>
+            enable_if_t<_self::is_known_identity && _self::is_acc> * = nullptr>
   reduction_impl(RedOutVar &Acc)
       : algo(reducer_type::getIdentity(), BinaryOperation(), false, Acc) {
     if (Acc.size() != 1)
@@ -815,7 +817,7 @@ class reduction_impl
   /// reduction value is added using BinaryOperation, i.e. it is expected that
   /// the memory is pre-initialized with some meaningful value.
   template <typename _self = self, enable_if_t<_self::is_known_identity &&
-                                               _self::my_is_usm> * = nullptr>
+                                               _self::is_usm> * = nullptr>
   reduction_impl(RedOutVar VarPtr, bool InitializeToIdentity = false)
       : algo(reducer_type::getIdentity(), BinaryOperation(),
              InitializeToIdentity, VarPtr) {}
@@ -823,8 +825,8 @@ class reduction_impl
   /// SYCL-2020.
   /// Constructs reduction_impl when the identity value is statically known.
   template <typename _self = self,
-            std::enable_if_t<_self::is_known_identity && _self::my_is_rw_acc>
-                * = nullptr>
+            std::enable_if_t<_self::is_known_identity && _self::is_rw_acc> * =
+                nullptr>
   reduction_impl(RedOutVar &Acc, handler &CGH, bool InitializeToIdentity)
       : algo(reducer_type::getIdentity(), BinaryOperation(),
              InitializeToIdentity, Acc) {
@@ -836,10 +838,9 @@ class reduction_impl
   }
 
   /// Constructs reduction_impl when the identity value is unknown.
-  template <typename _self = self,
-            enable_if_t<_self::my_is_rw_acc || _self::my_is_dw_acc> * = nullptr>
+  template <typename _self = self, enable_if_t<_self::is_acc> * = nullptr>
   reduction_impl(RedOutVar &Acc, const T &Identity, BinaryOperation BOp)
-      : algo(chooseIdentity(Identity), BOp, my_is_dw_acc, Acc) {
+      : algo(chooseIdentity(Identity), BOp, is_dw_acc, Acc) {
     if (Acc.size() != 1)
       throw sycl::runtime_error(errc::invalid,
                                 "Reduction variable must be a scalar.",
@@ -849,13 +850,13 @@ class reduction_impl
   /// The \param VarPtr is a USM pointer to memory, to where the computed
   /// reduction value is added using BinaryOperation, i.e. it is expected that
   /// the memory is pre-initialized with some meaningful value.
-  template <typename _self = self, enable_if_t<_self::my_is_usm> * = nullptr>
+  template <typename _self = self, enable_if_t<_self::is_usm> * = nullptr>
   reduction_impl(RedOutVar VarPtr, const T &Identity, BinaryOperation BOp,
                  bool InitializeToIdentity = false)
       : algo(chooseIdentity(Identity), BOp, InitializeToIdentity, VarPtr) {}
 
   /// For placeholder accessor.
-  template <typename _self = self, enable_if_t<_self::my_is_rw_acc> * = nullptr>
+  template <typename _self = self, enable_if_t<_self::is_rw_acc> * = nullptr>
   reduction_impl(RedOutVar &Acc, handler &CGH, const T &Identity,
                  BinaryOperation BOp, bool InitializeToIdentity)
       : algo(chooseIdentity(Identity), BOp, InitializeToIdentity, Acc) {
@@ -2393,7 +2394,7 @@ template <typename Reduction, typename... RestT>
 void reduSaveFinalResultToUserMemHelper(
     std::vector<event> &Events, std::shared_ptr<detail::queue_impl> Queue,
     bool IsHost, Reduction &Redu, RestT... Rest) {
-  if constexpr (Reduction::my_is_dw_acc) {
+  if constexpr (Reduction::is_dw_acc) {
     event CopyEvent = withAuxHandler(Queue, IsHost, [&](handler &CopyHandler) {
       auto InAcc = Redu.getReadAccToPreviousPartialReds(CopyHandler);
       auto OutAcc = Redu.getUserRedVar();
diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp
index 789a6c94aba84..6425d59a9b5fe 100644
--- a/sycl/include/sycl/handler.hpp
+++ b/sycl/include/sycl/handler.hpp
@@ -1655,7 +1655,7 @@ class __SYCL_EXPORT handler {
         *this, KernelFunc, Range, MaxWGSize, NumConcurrentWorkGroups, Redu);
     if (Reduction::is_usm ||
         (Reduction::has_fast_atomics && Redu.initializeToIdentity()) ||
-        (!Reduction::has_fast_atomics && Reduction::my_is_dw_acc)) {
+        (!Reduction::has_fast_atomics && Reduction::is_dw_acc)) {
       this->finalize();
       MLastEvent = withAuxHandler(QueueCopy, [&](handler &CopyHandler) {
         ext::oneapi::detail::reduSaveFinalResultToUserMem<KernelName>(
@@ -1782,7 +1782,7 @@ class __SYCL_EXPORT handler {
       });
     } // end while (NWorkItems > 1)
 
-    if (Reduction::is_usm || Reduction::my_is_dw_acc) {
+    if (Reduction::is_usm || Reduction::is_dw_acc) {
       MLastEvent = withAuxHandler(QueueCopy, [&](handler &CopyHandler) {
         ext::oneapi::detail::reduSaveFinalResultToUserMem<KernelName>(
             CopyHandler, Redu);

From 5c20d6008283af0250525bc72b50c2fac5f45ea3 Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Fri, 15 Jul 2022 09:51:41 -0700
Subject: [PATCH 16/24] clang-format

---
 sycl/include/sycl/ext/oneapi/reduction.hpp | 55 ++++++++++------------
 sycl/include/sycl/reduction.hpp            |  5 +-
 2 files changed, 28 insertions(+), 32 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/reduction.hpp b/sycl/include/sycl/ext/oneapi/reduction.hpp
index 51d37709ca97c..0e50b3f761c0b 100644
--- a/sycl/include/sycl/ext/oneapi/reduction.hpp
+++ b/sycl/include/sycl/ext/oneapi/reduction.hpp
@@ -496,8 +496,7 @@ template <typename T, class BinaryOperation> class reduction_impl_common {
   bool InitializeToIdentity;
 };
 
-template <class T>
-struct is_rw_acc_t : public std::false_type {};
+template <class T> struct is_rw_acc_t : public std::false_type {};
 
 template <class T, int AccessorDims, access::placeholder IsPlaceholder,
           typename PropList>
@@ -505,8 +504,7 @@ struct is_rw_acc_t<accessor<T, AccessorDims, access::mode::read_write,
                             access::target::device, IsPlaceholder, PropList>>
     : public std::true_type {};
 
-template <class T>
-struct is_dw_acc_t : public std::false_type {};
+template <class T> struct is_dw_acc_t : public std::false_type {};
 
 template <class T, int AccessorDims, access::placeholder IsPlaceholder,
           typename PropList>
@@ -514,17 +512,14 @@ struct is_dw_acc_t<accessor<T, AccessorDims, access::mode::discard_write,
                             access::target::device, IsPlaceholder, PropList>>
     : public std::true_type {};
 
-template <class T>
-struct is_placeholder_t : public std::false_type {};
+template <class T> struct is_placeholder_t : public std::false_type {};
 
 template <class T, int AccessorDims, access::mode Mode, typename PropList>
 struct is_placeholder_t<accessor<T, AccessorDims, Mode, access::target::device,
                                  access::placeholder::true_t, PropList>>
     : public std::true_type {};
 
-
-template <class T>
-struct accessor_dim_t {
+template <class T> struct accessor_dim_t {
   static constexpr int value = 1;
 };
 
@@ -536,7 +531,7 @@ struct accessor_dim_t<
 };
 
 template <class T> struct get_red_t;
-template <class T> struct get_red_t<T*> {
+template <class T> struct get_red_t<T *> {
   using type = T;
 };
 
@@ -790,10 +785,10 @@ class reduction_impl
   }
 
 public:
-  using algo::is_usm;
-  using algo::is_rw_acc;
-  using algo::is_dw_acc;
   using algo::is_acc;
+  using algo::is_dw_acc;
+  using algo::is_rw_acc;
+  using algo::is_usm;
 
   using reducer_type = typename algo::reducer_type;
   using rw_accessor_type = typename algo::rw_accessor_type;
@@ -816,8 +811,8 @@ class reduction_impl
   /// The \param VarPtr is a USM pointer to memory, to where the computed
   /// reduction value is added using BinaryOperation, i.e. it is expected that
   /// the memory is pre-initialized with some meaningful value.
-  template <typename _self = self, enable_if_t<_self::is_known_identity &&
-                                               _self::is_usm> * = nullptr>
+  template <typename _self = self,
+            enable_if_t<_self::is_known_identity && _self::is_usm> * = nullptr>
   reduction_impl(RedOutVar VarPtr, bool InitializeToIdentity = false)
       : algo(reducer_type::getIdentity(), BinaryOperation(),
              InitializeToIdentity, VarPtr) {}
@@ -870,9 +865,10 @@ class reduction_impl
 
 template <class BinaryOp, int Dims, size_t Extent, typename RedOutVar,
           typename... RestTy>
-auto make_reduction(RedOutVar RedVar, RestTy &&... Rest) {
+auto make_reduction(RedOutVar RedVar, RestTy &&...Rest) {
   return reduction_impl<typename get_red_t<RedOutVar>::type, BinaryOp, Dims,
-                        Extent, RedOutVar>{RedVar, std::forward<RestTy>(Rest)...};
+                        Extent, RedOutVar>{RedVar,
+                                           std::forward<RestTy>(Rest)...};
 }
 
 /// A helper to pass undefined (sycl::detail::auto_name) names unmodified. We
@@ -1277,7 +1273,6 @@ void reduCGFuncForNDRangeFastAtomicsOnly(
     if (LID == 0) {
       Reducer.atomic_combine(Reduction::getOutPointer(Out));
     }
-
   });
 }
 
@@ -1805,7 +1800,9 @@ struct IsNonUsmReductionPredicate {
 };
 
 struct EmptyReductionPredicate {
-  template <typename T> struct Func { static constexpr bool value = false; };
+  template <typename T> struct Func {
+    static constexpr bool value = false;
+  };
 };
 
 template <bool Cond, size_t I> struct FilterElement {
@@ -2119,7 +2116,7 @@ void associateReduAccsWithHandlerHelper(handler &CGH, ReductionT &Redu) {
 template <typename ReductionT, typename... RestT,
           enable_if_t<(sizeof...(RestT) > 0), int> Z = 0>
 void associateReduAccsWithHandlerHelper(handler &CGH, ReductionT &Redu,
-                                        RestT &... Rest) {
+                                        RestT &...Rest) {
   Redu.associateWithHandler(CGH);
   associateReduAccsWithHandlerHelper(CGH, Rest...);
 }
@@ -2518,15 +2515,15 @@ __SYCL_INLINE_CONSTEXPR AccumulatorT known_identity_v =
 
 #ifdef __SYCL_INTERNAL_API
 namespace __SYCL2020_DEPRECATED("use 'ext::oneapi' instead") ONEAPI {
-  using namespace ext::oneapi;
-  namespace detail {
-  using cl::sycl::detail::queue_impl;
-  __SYCL_EXPORT size_t reduGetMaxWGSize(std::shared_ptr<queue_impl> Queue,
-                                        size_t LocalMemBytesPerWorkItem);
-  __SYCL_EXPORT size_t reduComputeWGSize(size_t NWorkItems, size_t MaxWGSize,
-                                         size_t &NWorkGroups);
-  } // namespace detail
-} // namespace ONEAPI
+using namespace ext::oneapi;
+namespace detail {
+using cl::sycl::detail::queue_impl;
+__SYCL_EXPORT size_t reduGetMaxWGSize(std::shared_ptr<queue_impl> Queue,
+                                      size_t LocalMemBytesPerWorkItem);
+__SYCL_EXPORT size_t reduComputeWGSize(size_t NWorkItems, size_t MaxWGSize,
+                                       size_t &NWorkGroups);
+} // namespace detail
+} // namespace __SYCL2020_DEPRECATED("use 'ext::oneapi' instead")ONEAPI
 #endif // __SYCL_INTERNAL_API
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)
diff --git a/sycl/include/sycl/reduction.hpp b/sycl/include/sycl/reduction.hpp
index 64c054a9660a6..44e37a3db254d 100644
--- a/sycl/include/sycl/reduction.hpp
+++ b/sycl/include/sycl/reduction.hpp
@@ -89,9 +89,8 @@ reduction(T *, BinaryOperation, const property_list &PropList = {}) {
 /// reduction identity value \p Identity, reduction operation \p Combiner,
 /// and optional reduction properties.
 template <typename T, typename AllocatorT, typename BinaryOperation>
-auto
-reduction(buffer<T, 1, AllocatorT> Var, handler &CGH, const T &Identity,
-          BinaryOperation Combiner, const property_list &PropList = {}) {
+auto reduction(buffer<T, 1, AllocatorT> Var, handler &CGH, const T &Identity,
+               BinaryOperation Combiner, const property_list &PropList = {}) {
   bool InitializeToIdentity =
       PropList.has_property<property::reduction::initialize_to_identity>();
   return ext::oneapi::detail::make_reduction<BinaryOperation, 0, 1>(

From 4aa74393b2bb760caa290ab89f220ddc5626771b Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Fri, 15 Jul 2022 11:13:24 -0700
Subject: [PATCH 17/24] Fix a typo in a comment

---
 sycl/include/sycl/ext/oneapi/reduction.hpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/sycl/include/sycl/ext/oneapi/reduction.hpp b/sycl/include/sycl/ext/oneapi/reduction.hpp
index 0e50b3f761c0b..2b41e151da3f1 100644
--- a/sycl/include/sycl/ext/oneapi/reduction.hpp
+++ b/sycl/include/sycl/ext/oneapi/reduction.hpp
@@ -577,7 +577,7 @@ class reduction_impl_algo : public reduction_impl_common<T, BinaryOperation> {
   static constexpr bool is_dw_acc = is_dw_acc_t<RedOutVar>::value;
   static constexpr bool is_acc = is_rw_acc | is_dw_acc;
   static_assert(!is_rw_acc || !is_dw_acc, "Can be only one at once!");
-  static_assert(!is_usm || !is_acc, "Ca be only one at once!");
+  static_assert(!is_usm || !is_acc, "Can be only one at once!");
 
   static constexpr size_t dims = Dims;
   static constexpr size_t num_elements = Extent;

From 5e305dc0fbc5c8cbad7c1a555511365e86a17489 Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Fri, 15 Jul 2022 11:13:37 -0700
Subject: [PATCH 18/24] Benefit from accessor's CTAD

---
 sycl/include/sycl/reduction.hpp | 10 ++--------
 1 file changed, 2 insertions(+), 8 deletions(-)

diff --git a/sycl/include/sycl/reduction.hpp b/sycl/include/sycl/reduction.hpp
index 44e37a3db254d..1ff6682db335e 100644
--- a/sycl/include/sycl/reduction.hpp
+++ b/sycl/include/sycl/reduction.hpp
@@ -28,10 +28,7 @@ auto reduction(buffer<T, 1, AllocatorT> Var, handler &CGH, BinaryOperation,
   bool InitializeToIdentity =
       PropList.has_property<property::reduction::initialize_to_identity>();
   return ext::oneapi::detail::make_reduction<BinaryOperation, 0, 1>(
-      accessor<T, 1, access::mode::read_write, access::target::device,
-               access::placeholder::true_t,
-               ext::oneapi::accessor_property_list<>>{Var},
-      CGH, InitializeToIdentity);
+      accessor{Var}, CGH, InitializeToIdentity);
 }
 
 /// Constructs a reduction object using the given buffer \p Var, handler \p CGH,
@@ -94,10 +91,7 @@ auto reduction(buffer<T, 1, AllocatorT> Var, handler &CGH, const T &Identity,
   bool InitializeToIdentity =
       PropList.has_property<property::reduction::initialize_to_identity>();
   return ext::oneapi::detail::make_reduction<BinaryOperation, 0, 1>(
-      accessor<T, 1, access::mode::read_write, access::target::device,
-               access::placeholder::true_t,
-               ext::oneapi::accessor_property_list<>>{Var},
-      CGH, Identity, Combiner, InitializeToIdentity);
+      accessor{Var}, CGH, Identity, Combiner, InitializeToIdentity);
 }
 
 /// Constructs a reduction object using the reduction variable referenced by

From cf039fe80e28d6ecc6e39ac18c7d29d7c5b5d59c Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Fri, 15 Jul 2022 15:50:03 -0700
Subject: [PATCH 19/24] WIP Simplify getWriteMemForPartialReds

---
 sycl/include/sycl/ext/oneapi/reduction.hpp | 33 ++++++----------------
 1 file changed, 8 insertions(+), 25 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/reduction.hpp b/sycl/include/sycl/ext/oneapi/reduction.hpp
index 2b41e151da3f1..7997735fa87ab 100644
--- a/sycl/include/sycl/ext/oneapi/reduction.hpp
+++ b/sycl/include/sycl/ext/oneapi/reduction.hpp
@@ -612,32 +612,15 @@ class reduction_impl_algo : public reduction_impl_common<T, BinaryOperation> {
     return {*MOutBufPtr, CGH};
   }
 
-  /// Returns user's USM pointer passed to reduction for editing.
-  template <bool IsOneWG, bool _IsUSM = is_usm>
-  std::enable_if_t<IsOneWG && _IsUSM, result_type *>
-  getWriteMemForPartialReds(size_t, handler &) {
-    return getUserRedVar();
-  }
-
-  /// Returns user's accessor passed to reduction for editing if that is
-  /// the read-write accessor. Otherwise, create a new buffer and return
-  /// read-write accessor to it.
-  template <bool IsOneWG, bool _IsUSM = is_usm>
-  std::enable_if_t<IsOneWG && !_IsUSM, rw_accessor_type>
-  getWriteMemForPartialReds(size_t, handler &CGH) {
-    if constexpr (is_rw_acc)
-      return MRedOut;
-    return getWriteMemForPartialReds<false>(1, CGH);
-  }
-
-  /// Constructs a new temporary buffer to hold partial sums and returns
-  /// the accessor for that buffer.
   template <bool IsOneWG>
-  std::enable_if_t<!IsOneWG, rw_accessor_type>
-  getWriteMemForPartialReds(size_t Size, handler &CGH) {
-    MOutBufPtr = std::make_shared<buffer<T, buffer_dim>>(range<1>(Size));
-    CGH.addReduction(MOutBufPtr);
-    return createHandlerWiredReadWriteAccessor(CGH, *MOutBufPtr);
+  auto getWriteMemForPartialReds(size_t Size, handler &CGH) {
+    if constexpr (IsOneWG && !is_dw_acc) {
+      return MRedOut;
+    } else {
+      MOutBufPtr = std::make_shared<buffer<T, buffer_dim>>(range<1>(Size));
+      CGH.addReduction(MOutBufPtr);
+      return createHandlerWiredReadWriteAccessor(CGH, *MOutBufPtr);
+    }
   }
 
   /// Returns an accessor accessing the memory that will hold the reduction

From 8967bf337c441f48f1aa26dee48bfd10b8b56fcc Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Fri, 15 Jul 2022 17:25:46 -0700
Subject: [PATCH 20/24] Update associateWithHandler uses

---
 sycl/include/sycl/ext/oneapi/reduction.hpp | 42 +++++++---------------
 1 file changed, 12 insertions(+), 30 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/reduction.hpp b/sycl/include/sycl/ext/oneapi/reduction.hpp
index 7997735fa87ab..78735fc571c79 100644
--- a/sycl/include/sycl/ext/oneapi/reduction.hpp
+++ b/sycl/include/sycl/ext/oneapi/reduction.hpp
@@ -586,15 +586,6 @@ class reduction_impl_algo : public reduction_impl_common<T, BinaryOperation> {
                       RedOutVar RedOut)
       : base(Identity, BinaryOp, Init), MRedOut(std::move(RedOut)){};
 
-  /// Associates the reduction accessor to user's memory with \p CGH handler
-  /// to keep the accessor alive until the command group finishes the work.
-  /// This function does not do anything for USM reductions.
-  void associateWithHandler(handler &CGH) {
-    if constexpr (is_acc) {
-      CGH.associateWithHandler(&MRedOut, access::target::device);
-    }
-  }
-
   /// Creates and returns a local accessor with the \p Size elements.
   /// By default the local accessor elements are of the same type as the
   /// elements processed by the reduction, but may it be altered by specifying
@@ -632,7 +623,7 @@ class reduction_impl_algo : public reduction_impl_common<T, BinaryOperation> {
   rw_accessor_type getWriteAccForPartialReds(size_t Size, handler &CGH) {
     if constexpr (is_rw_acc) {
       if (Size == 1) {
-        associateWithHandler(CGH);
+        CGH.associateWithHandler(&MRedOut, access::target::device);
         return MRedOut;
       }
     }
@@ -808,7 +799,7 @@ class reduction_impl
   reduction_impl(RedOutVar &Acc, handler &CGH, bool InitializeToIdentity)
       : algo(reducer_type::getIdentity(), BinaryOperation(),
              InitializeToIdentity, Acc) {
-    algo::associateWithHandler(CGH);
+    associateWithHandler(CGH, &Acc, access::target::device);
     if (Acc.size() != 1)
       throw sycl::runtime_error(errc::invalid,
                                 "Reduction variable must be a scalar.",
@@ -838,7 +829,7 @@ class reduction_impl
   reduction_impl(RedOutVar &Acc, handler &CGH, const T &Identity,
                  BinaryOperation BOp, bool InitializeToIdentity)
       : algo(chooseIdentity(Identity), BOp, InitializeToIdentity, Acc) {
-    algo::associateWithHandler(CGH);
+    associateWithHandler(CGH, &Acc, access::target::device);
     if (Acc.size() != 1)
       throw sycl::runtime_error(errc::invalid,
                                 "Reduction variable must be a scalar.",
@@ -1561,7 +1552,7 @@ template <typename KernelName, class Reduction>
 std::enable_if_t<!Reduction::is_usm>
 reduSaveFinalResultToUserMem(handler &CGH, Reduction &Redu) {
   auto InAcc = Redu.getReadAccToPreviousPartialReds(CGH);
-  Redu.associateWithHandler(CGH);
+  associateWithHandler(CGH, &Redu.getUserRedVar(), access::target::device);
   CGH.copy(InAcc, Redu.getUserRedVar());
 }
 
@@ -2089,26 +2080,16 @@ void reduCGFuncAtomic64(handler &CGH, KernelType KernelFunc,
       CGH, KernelFunc, Range, Redu, Out);
 }
 
-inline void associateReduAccsWithHandlerHelper(handler &) {}
-
-template <typename ReductionT>
-void associateReduAccsWithHandlerHelper(handler &CGH, ReductionT &Redu) {
-  Redu.associateWithHandler(CGH);
-}
-
-template <typename ReductionT, typename... RestT,
-          enable_if_t<(sizeof...(RestT) > 0), int> Z = 0>
-void associateReduAccsWithHandlerHelper(handler &CGH, ReductionT &Redu,
-                                        RestT &...Rest) {
-  Redu.associateWithHandler(CGH);
-  associateReduAccsWithHandlerHelper(CGH, Rest...);
-}
-
 template <typename... Reductions, size_t... Is>
 void associateReduAccsWithHandler(handler &CGH,
                                   std::tuple<Reductions...> &ReduTuple,
                                   std::index_sequence<Is...>) {
-  associateReduAccsWithHandlerHelper(CGH, std::get<Is>(ReduTuple)...);
+  auto ProcessOne = [&CGH](auto Redu) {
+    if constexpr (decltype(Redu)::is_acc) {
+      associateWithHandler(CGH, &Redu.getUserRedVar(), access::target::device);
+    }
+  };
+  (ProcessOne(std::get<Is>(ReduTuple)), ...);
 }
 
 /// All scalar reductions are processed together; there is one loop of log2(N)
@@ -2378,7 +2359,8 @@ void reduSaveFinalResultToUserMemHelper(
     event CopyEvent = withAuxHandler(Queue, IsHost, [&](handler &CopyHandler) {
       auto InAcc = Redu.getReadAccToPreviousPartialReds(CopyHandler);
       auto OutAcc = Redu.getUserRedVar();
-      Redu.associateWithHandler(CopyHandler);
+      associateWithHandler(CopyHandler, &Redu.getUserRedVar(),
+                           access::target::device);
       if (!Events.empty())
         CopyHandler.depends_on(Events.back());
       CopyHandler.copy(InAcc, OutAcc);

From 156f9ee0366b673f222368ea8464afd5c93822ab Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Mon, 18 Jul 2022 15:29:15 -0700
Subject: [PATCH 21/24] Remove useless single-callsite wrapper

---
 sycl/include/sycl/ext/oneapi/reduction.hpp | 21 ++++-----------------
 1 file changed, 4 insertions(+), 17 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/reduction.hpp b/sycl/include/sycl/ext/oneapi/reduction.hpp
index 78735fc571c79..211bf0cf2ab20 100644
--- a/sycl/include/sycl/ext/oneapi/reduction.hpp
+++ b/sycl/include/sycl/ext/oneapi/reduction.hpp
@@ -2032,6 +2032,7 @@ namespace main_krn {
 template <class KernelName> struct NDRangeAtomic64;
 } // namespace main_krn
 } // namespace reduction
+
 // Specialization for devices with the atomic64 aspect, which guarantees 64 (and
 // temporarily 32) bit floating point support for atomic add.
 // TODO 32 bit floating point atomics are eventually expected to be supported by
@@ -2039,9 +2040,9 @@ template <class KernelName> struct NDRangeAtomic64;
 // IsReduOptForAtomic64Add, as prescribed in its documentation, should then also
 // be made.
 template <typename KernelName, typename KernelType, int Dims, class Reduction>
-void reduCGFuncImplAtomic64(handler &CGH, KernelType KernelFunc,
-                            const nd_range<Dims> &Range, Reduction &,
-                            typename Reduction::rw_accessor_type Out) {
+void reduCGFuncAtomic64(handler &CGH, KernelType KernelFunc,
+                        const nd_range<Dims> &Range, Reduction &Redu) {
+  auto Out = Redu.getReadWriteAccessorToInitializedMem(CGH);
   static_assert(Reduction::has_atomic_add_float64,
                 "Only suitable for reductions that have FP64 atomic add.");
   constexpr size_t NElements = Reduction::num_elements;
@@ -2066,20 +2067,6 @@ void reduCGFuncImplAtomic64(handler &CGH, KernelType KernelFunc,
   });
 }
 
-// Specialization for devices with the atomic64 aspect, which guarantees 64 (and
-// temporarily 32) bit floating point support for atomic add.
-// TODO 32 bit floating point atomics are eventually expected to be supported by
-// the has_fast_atomics specialization. Corresponding changes to
-// IsReduOptForAtomic64Add, as prescribed in its documentation, should then also
-// be made.
-template <typename KernelName, typename KernelType, int Dims, class Reduction>
-void reduCGFuncAtomic64(handler &CGH, KernelType KernelFunc,
-                        const nd_range<Dims> &Range, Reduction &Redu) {
-  auto Out = Redu.getReadWriteAccessorToInitializedMem(CGH);
-  reduCGFuncImplAtomic64<KernelName, KernelType, Dims, Reduction>(
-      CGH, KernelFunc, Range, Redu, Out);
-}
-
 template <typename... Reductions, size_t... Is>
 void associateReduAccsWithHandler(handler &CGH,
                                   std::tuple<Reductions...> &ReduTuple,

From b6127e39ba897289c5bab186ed972ae3c69fdaff Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Fri, 22 Jul 2022 11:42:13 -0700
Subject: [PATCH 22/24] Fix bug

---
 sycl/include/sycl/ext/oneapi/reduction.hpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/sycl/include/sycl/ext/oneapi/reduction.hpp b/sycl/include/sycl/ext/oneapi/reduction.hpp
index 211bf0cf2ab20..63258312d7484 100644
--- a/sycl/include/sycl/ext/oneapi/reduction.hpp
+++ b/sycl/include/sycl/ext/oneapi/reduction.hpp
@@ -774,7 +774,7 @@ class reduction_impl
   template <typename _self = self,
             enable_if_t<_self::is_known_identity && _self::is_acc> * = nullptr>
   reduction_impl(RedOutVar &Acc)
-      : algo(reducer_type::getIdentity(), BinaryOperation(), false, Acc) {
+      : algo(reducer_type::getIdentity(), BinaryOperation(), is_dw_acc, Acc) {
     if (Acc.size() != 1)
       throw sycl::runtime_error(errc::invalid,
                                 "Reduction variable must be a scalar.",

From 914fcfa14980faf94b3dac7de337526db41d82d6 Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Wed, 27 Jul 2022 10:28:11 -0700
Subject: [PATCH 23/24] Rename the trait

---
 sycl/include/sycl/ext/oneapi/reduction.hpp | 7 ++++---
 1 file changed, 4 insertions(+), 3 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/reduction.hpp b/sycl/include/sycl/ext/oneapi/reduction.hpp
index 63258312d7484..2387e74b6c98b 100644
--- a/sycl/include/sycl/ext/oneapi/reduction.hpp
+++ b/sycl/include/sycl/ext/oneapi/reduction.hpp
@@ -519,13 +519,14 @@ struct is_placeholder_t<accessor<T, AccessorDims, Mode, access::target::device,
                                  access::placeholder::true_t, PropList>>
     : public std::true_type {};
 
-template <class T> struct accessor_dim_t {
+// Used for determining dimensions for temporary storage (mainly).
+template <class T> struct data_dim_t {
   static constexpr int value = 1;
 };
 
 template <class T, int AccessorDims, access::mode Mode,
           access::placeholder IsPH, typename PropList>
-struct accessor_dim_t<
+struct data_dim_t<
     accessor<T, AccessorDims, Mode, access::target::device, IsPH, PropList>> {
   static constexpr int value = AccessorDims;
 };
@@ -556,7 +557,7 @@ class reduction_impl_algo : public reduction_impl_common<T, BinaryOperation> {
   // Buffers and accessors always describe scalar reductions (i.e. Dims == 0)
   // The input buffer/accessor is allowed to have different dimensionality
   // AccessorDims also determines the dimensionality of some temp storage
-  static constexpr int accessor_dim = accessor_dim_t<RedOutVar>::value;
+  static constexpr int accessor_dim = data_dim_t<RedOutVar>::value;
   static constexpr int buffer_dim = (accessor_dim == 0) ? 1 : accessor_dim;
   static constexpr access::placeholder is_placeholder =
       is_placeholder_t<RedOutVar>::value ? access::placeholder::true_t

From a96dfac9257231d123289484cfc3191d8cf7dc37 Mon Sep 17 00:00:00 2001
From: "Elovikov, Andrei" <andrei.elovikov@intel.com>
Date: Wed, 27 Jul 2022 12:02:07 -0700
Subject: [PATCH 24/24] Add a comment

---
 sycl/include/sycl/ext/oneapi/reduction.hpp | 6 ++++++
 1 file changed, 6 insertions(+)

diff --git a/sycl/include/sycl/ext/oneapi/reduction.hpp b/sycl/include/sycl/ext/oneapi/reduction.hpp
index 2387e74b6c98b..08598b2540cd5 100644
--- a/sycl/include/sycl/ext/oneapi/reduction.hpp
+++ b/sycl/include/sycl/ext/oneapi/reduction.hpp
@@ -606,6 +606,12 @@ class reduction_impl_algo : public reduction_impl_common<T, BinaryOperation> {
 
   template <bool IsOneWG>
   auto getWriteMemForPartialReds(size_t Size, handler &CGH) {
+    // If there is only one WG we can avoid creation of temporary buffer with
+    // partial sums and write directly into user's reduction variable.
+    //
+    // Current implementation doesn't allow that in case of DW accessor used for
+    // reduction because C++ types for it and for temporary storage don't match,
+    // hence the second part of the check.
     if constexpr (IsOneWG && !is_dw_acc) {
       return MRedOut;
     } else {