diff --git a/sycl/include/sycl/ext/oneapi/reduction.hpp b/sycl/include/sycl/ext/oneapi/reduction.hpp
index 3005dc3417944..08598b2540cd5 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()};
   }
 
@@ -504,78 +496,96 @@ template <typename T, class BinaryOperation> class reduction_impl_common {
   bool InitializeToIdentity;
 };
 
-/// Types representing specific reduction algorithms
-/// Enables reduction_impl_algo to take additional algorithm-specific templates
-template <bool IsUSM, access::placeholder IsPlaceholder, int AccessorDims>
-class default_reduction_algorithm {};
+template <class T> struct is_rw_acc_t : public std::false_type {};
 
-/// Templated class for implementations of specific reduction algorithms
-template <typename T, class BinaryOperation, int Dims, size_t Extent,
-          class Algorithm>
-class reduction_impl_algo;
+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,
+          typename PropList>
+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, 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 {};
+
+// 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 data_dim_t<
+    accessor<T, AccessorDims, Mode, access::target::device, IsPH, PropList>> {
+  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;
+};
 
-/// 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>
-class reduction_impl_algo<
-    T, BinaryOperation, Dims, Extent,
-    default_reduction_algorithm<IsUSM, IsPlaceholder, AccessorDims>>
-    : 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, RedOutVar>;
 
 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;
 
   // 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;
-  using rw_accessor_type = accessor<T, AccessorDims, access::mode::read_write,
-                                    access::target::device, IsPlaceholder,
+  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
+                                         : access::placeholder::false_t;
+  using rw_accessor_type = accessor<T, accessor_dim, access::mode::read_write,
+                                    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 =
       IsReduOptForFastAtomicFetch<T, BinaryOperation>::value;
   static constexpr bool has_fast_reduce =
       IsReduOptForFastReduce<T, BinaryOperation>::value;
-  static constexpr bool is_usm = IsUSM;
-  static constexpr bool is_placeholder =
-      (IsPlaceholder == access::placeholder::true_t);
+
+  static constexpr bool is_usm = std::is_same_v<RedOutVar, T *>;
+
+  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, "Can be only one at once!");
 
   static constexpr size_t dims = Dims;
   static constexpr size_t num_elements = Extent;
 
   reduction_impl_algo(const T &Identity, BinaryOperation BinaryOp, bool Init,
-                      std::shared_ptr<rw_accessor_type> AccPointer)
-      : base(Identity, BinaryOp, Init), MRWAcc(AccPointer){};
-  reduction_impl_algo(const T &Identity, BinaryOperation BinaryOp, bool Init,
-                      std::shared_ptr<dw_accessor_type> AccPointer)
-      : base(Identity, BinaryOp, Init), MDWAcc(AccPointer){};
-  reduction_impl_algo(const T &Identity, BinaryOperation BinaryOp, bool Init,
-                      T *USMPointer)
-      : base(Identity, BinaryOp, Init), MUSMPointer(USMPointer){};
-
-  /// 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 (MRWAcc)
-      CGH.associateWithHandler(MRWAcc.get(), access::target::device);
-    else if (MDWAcc)
-      CGH.associateWithHandler(MDWAcc.get(), access::target::device);
-  }
+                      RedOutVar RedOut)
+      : base(Identity, BinaryOp, Init), MRedOut(std::move(RedOut)){};
 
   /// Creates and returns a local accessor with the \p Size elements.
   /// By default the local accessor elements are of the same type as the
@@ -594,32 +604,21 @@ class reduction_impl_algo<
     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 getUSMPointer();
-  }
-
-  /// 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 (MRWAcc)
-      return *MRWAcc;
-    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 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 {
+      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
@@ -629,9 +628,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 (is_rw_acc) {
+      if (Size == 1) {
+        CGH.associateWithHandler(&MRedOut, access::target::device);
+        return MRedOut;
+      }
     }
 
     // Create a new output buffer and return an accessor to it.
@@ -647,8 +648,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 (is_rw_acc) {
+      if (!base::initializeToIdentity())
+        return MRedOut;
+    }
+    assert(!(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]>();
@@ -675,22 +680,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();
@@ -701,29 +691,27 @@ class reduction_impl_algo<
   }
 
 private:
-  template <typename BufferT, access::placeholder IsPH = IsPlaceholder>
-  std::enable_if_t<IsPH == 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>
-  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};
+    }
   }
 
-  /// 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;
+  /// 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
@@ -741,82 +729,22 @@ template <typename T> struct AreAllButLastReductions<T> {
   static constexpr bool value =
       !std::is_base_of<reduction_impl_base, std::remove_reference_t<T>>::value;
 };
-
 /// 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 reduction_impl
     : private reduction_impl_base,
-      public reduction_impl_algo<T, BinaryOperation, Dims, Extent, Algorithm> {
+      public reduction_impl_algo<T, BinaryOperation, Dims, Extent, RedOutVar> {
 private:
-  using algo = reduction_impl_algo<T, BinaryOperation, Dims, Extent, Algorithm>;
+  using algo = reduction_impl_algo<T, BinaryOperation, Dims, Extent, RedOutVar>;
+  using self = reduction_impl<T, BinaryOperation, Dims, Extent, RedOutVar>;
 
-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;
+  static constexpr bool is_known_identity =
+      sycl::detail::IsKnownIdentityOp<T, BinaryOperation>::value;
 
-  // 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 _T, typename AllocatorT,
-            std::enable_if_t<sycl::detail::IsKnownIdentityOp<
-                _T, BinaryOperation>::value> * = nullptr>
-  reduction_impl(buffer<_T, 1, AllocatorT> Buffer, handler &CGH,
-                 bool InitializeToIdentity)
-      : algo(reducer_type::getIdentity(), BinaryOperation(),
-             InitializeToIdentity, std::make_shared<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);
-  }
-
-  /// 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(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> * = nullptr>
-  reduction_impl(dw_accessor_type &Acc)
-      : algo(reducer_type::getIdentity(), BinaryOperation(), true,
-             std::make_shared<dw_accessor_type>(Acc)) {
-    if (Acc.size() != 1)
-      throw sycl::runtime_error(errc::invalid,
-                                "Reduction variable must be a scalar.",
-                                PI_ERROR_INVALID_VALUE);
-  }
-
-  /// 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>
-          * = nullptr>
-  reduction_impl(buffer<_T, 1, AllocatorT> Buffer, handler &CGH,
-                 const T & /*Identity*/, BinaryOperation,
-                 bool InitializeToIdentity)
-      : algo(reducer_type::getIdentity(), BinaryOperation(),
-             InitializeToIdentity, std::make_shared<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);
+  // 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
@@ -828,79 +756,57 @@ class reduction_impl
     // 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();
 
-  /// 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>
-  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.
+    } else {
+      return Identity;
+    }
   }
 
-  /// 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>
-  reduction_impl(dw_accessor_type &Acc, const T & /*Identity*/, BinaryOperation)
-      : algo(reducer_type::getIdentity(), BinaryOperation(), true,
-             std::make_shared<dw_accessor_type>(Acc)) {
+public:
+  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;
+
+  // Only scalar and 1D array reductions are supported by SYCL 2020.
+  static_assert(Dims <= 1, "Multi-dimensional reductions are not supported.");
+
+  /// Constructs reduction_impl when the identity value is statically known.
+  template <typename _self = self,
+            enable_if_t<_self::is_known_identity && _self::is_acc> * = nullptr>
+  reduction_impl(RedOutVar &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.",
                                 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>
-          * = nullptr>
-  reduction_impl(buffer<_T, 1, AllocatorT> Buffer, handler &CGH,
-                 const T &Identity, BinaryOperation BOp,
-                 bool InitializeToIdentity)
-      : algo(Identity, BOp, InitializeToIdentity,
-             std::make_shared<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);
-  }
+  /// 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::is_usm> * = nullptr>
+  reduction_impl(RedOutVar VarPtr, bool InitializeToIdentity = false)
+      : algo(reducer_type::getIdentity(), BinaryOperation(),
+             InitializeToIdentity, VarPtr) {}
 
-  /// Constructs reduction_impl when the identity value is unknown.
-  template <typename _T = T, enable_if_t<!sycl::detail::IsKnownIdentityOp<
-                                 _T, BinaryOperation>::value> * = nullptr>
-  reduction_impl(rw_accessor_type &Acc, const T &Identity, BinaryOperation BOp)
-      : algo(Identity, BOp, false, std::make_shared<rw_accessor_type>(Acc)) {
+  /// 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::is_rw_acc> * =
+                nullptr>
+  reduction_impl(RedOutVar &Acc, handler &CGH, bool InitializeToIdentity)
+      : algo(reducer_type::getIdentity(), BinaryOperation(),
+             InitializeToIdentity, Acc) {
+    associateWithHandler(CGH, &Acc, access::target::device);
     if (Acc.size() != 1)
       throw sycl::runtime_error(errc::invalid,
                                 "Reduction variable must be a scalar.",
@@ -908,83 +814,44 @@ 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>
-  reduction_impl(dw_accessor_type &Acc, const T &Identity, BinaryOperation BOp)
-      : algo(Identity, BOp, true, std::make_shared<dw_accessor_type>(Acc)) {
+  template <typename _self = self, enable_if_t<_self::is_acc> * = nullptr>
+  reduction_impl(RedOutVar &Acc, const T &Identity, BinaryOperation BOp)
+      : algo(chooseIdentity(Identity), BOp, is_dw_acc, 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.
-  /// 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>
-  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> * = nullptr>
-  reduction_impl(T *VarPtr, const T &Identity, BinaryOperation,
+  template <typename _self = self, enable_if_t<_self::is_usm> * = nullptr>
+  reduction_impl(RedOutVar VarPtr, const T &Identity, BinaryOperation BOp,
                  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.
+      : algo(chooseIdentity(Identity), BOp, InitializeToIdentity, VarPtr) {}
+
+  /// For placeholder accessor.
+  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) {
+    associateWithHandler(CGH, &Acc, access::target::device);
+    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.
-  /// 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>
-  reduction_impl(T *VarPtr, const T &Identity, BinaryOperation BOp,
-                 bool InitializeToIdentity = false)
-      : algo(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)
-      : 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> * = 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> * = nullptr>
-  reduction_impl(span<T, Extent> Span, const T &Identity, BinaryOperation BOp,
-                 bool InitializeToIdentity = false)
-      : algo(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>
@@ -1387,7 +1254,6 @@ void reduCGFuncForNDRangeFastAtomicsOnly(
     if (LID == 0) {
       Reducer.atomic_combine(Reduction::getOutPointer(Out));
     }
-
   });
 }
 
@@ -1693,11 +1559,8 @@ 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);
-  if (Redu.hasUserDiscardWriteAccessor())
-    CGH.copy(InAcc, Redu.getUserDiscardWriteAccessor());
-  else
-    CGH.copy(InAcc, Redu.getUserReadWriteAccessor());
+  associateWithHandler(CGH, &Redu.getUserRedVar(), access::target::device);
+  CGH.copy(InAcc, Redu.getUserRedVar());
 }
 
 // This method is used for implementation of parallel_for accepting 1 reduction.
@@ -1710,7 +1573,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>([=] {
@@ -1918,7 +1781,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 {
@@ -2174,6 +2039,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
@@ -2181,9 +2047,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;
@@ -2208,40 +2074,16 @@ 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);
-}
-
-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)
@@ -2507,21 +2349,17 @@ 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::is_dw_acc) {
+    event CopyEvent = withAuxHandler(Queue, IsHost, [&](handler &CopyHandler) {
+      auto InAcc = Redu.getReadAccToPreviousPartialReds(CopyHandler);
+      auto OutAcc = Redu.getUserRedVar();
+      associateWithHandler(CopyHandler, &Redu.getUserRedVar(),
+                           access::target::device);
+      if (!Events.empty())
+        CopyHandler.depends_on(Events.back());
+      CopyHandler.copy(InAcc, OutAcc);
+    });
+    Events.push_back(CopyEvent);
   }
   reduSaveFinalResultToUserMemHelper(Events, Queue, IsHost, Rest...);
 }
@@ -2575,11 +2413,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,
-                       detail::default_reduction_algorithm<false, IsPH, Dims>>
-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.
@@ -2587,14 +2423,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,
-                     detail::default_reduction_algorithm<false, IsPH, Dims>>>
-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.
@@ -2602,11 +2436,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<true, access::placeholder::false_t, 1>>
-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.
@@ -2614,14 +2445,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,
-                           detail::default_reduction_algorithm<
-                               true, access::placeholder::false_t, 1>>>
-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
@@ -2646,15 +2474,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/handler.hpp b/sycl/include/sycl/handler.hpp
index c6ee4bc17c675..6425d59a9b5fe 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 reduction_impl_algo;
 
 using cl::sycl::detail::enable_if_t;
@@ -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::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::is_dw_acc) {
       MLastEvent = withAuxHandler(QueueCopy, [&](handler &CopyHandler) {
         ext::oneapi::detail::reduSaveFinalResultToUserMem<KernelName>(
             CopyHandler, Redu);
@@ -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>
   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..1ff6682db335e 100644
--- a/sycl/include/sycl/reduction.hpp
+++ b/sycl/include/sycl/reduction.hpp
@@ -20,29 +20,29 @@ 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>>>
-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{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,
-                     ext::oneapi::detail::default_reduction_algorithm<
-                         false, access::placeholder::true_t, 1>>>
+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.
@@ -55,16 +55,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,
-                     ext::oneapi::detail::default_reduction_algorithm<
-                         true, access::placeholder::false_t, 1>>>
-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
@@ -72,12 +70,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,
-                     ext::oneapi::detail::default_reduction_algorithm<
-                         true, access::placeholder::false_t, 1>>>
+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,48 +86,40 @@ 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,
-    ext::oneapi::detail::default_reduction_algorithm<
-        false, access::placeholder::true_t, 1>>
-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 {Var, CGH, Identity, Combiner, InitializeToIdentity};
+  return ext::oneapi::detail::make_reduction<BinaryOperation, 0, 1>(
+      accessor{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,
-    ext::oneapi::detail::default_reduction_algorithm<
-        true, access::placeholder::false_t, 1>>
-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,
-                     ext::oneapi::detail::default_reduction_algorithm<
-                         true, access::placeholder::false_t, 1>>>
-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
@@ -139,13 +127,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,
-                     ext::oneapi::detail::default_reduction_algorithm<
-                         true, access::placeholder::false_t, 1>>>
+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.
@@ -158,17 +144,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,
-                     ext::oneapi::detail::default_reduction_algorithm<
-                         true, access::placeholder::false_t, 1>>>
-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