[SYCL] Add reduction overloads accepting span (#6019)

A new static member function "num_elements" is added to all reduction
specializations to distinguish between scalar reductions and array reductions.
A scalar reduction (all existing implementations of reduction) always
has (num_elements == 1); an array reduction (this new specialization) has
(num_elements == Extent), where Extent is the extent of the span.

If (num_elements > 1), the implementation currently reduces each element
separately. This allows array reductions to use the same amount of work-group
local memory as a scalar reduction using the same T, but at the expense of
additional synchronization calls.

The notion of a reducer's "dimensionality" is now tied to the dimensionality of
the reduction being performed (i.e. 0 for scalars, 1 for spans) and not to the
dimensionality of the input accessor/buffer. This will simplify extending
reductions to true multi-dimensional array reductions (i.e. with md_span).

In future, the code handling the reduction packs could be generalized to
handle more reduction/reducer types while making fewer assumptions about
the reduction algorithm.

Signed-off-by: John Pennycook [email protected]

Author: Pennycook

sycl/include/CL/sycl/handler.hpp

@@ -240,9 +240,9 @@ class RoundedRangeKernelWithKH {
 namespace ext {
 namespace oneapi {
 namespace detail {
-template <typename T, class BinaryOperation, int Dims, bool IsUSM,
-          access::placeholder IsPlaceholder>
-class reduction_impl;
+template <typename T, class BinaryOperation, int Dims, size_t Extent,
+          class Algorithm>
+class reduction_impl_algo;
 
 using cl::sycl::detail::enable_if_t;
 using cl::sycl::detail::queue_impl;
@@ -2682,11 +2682,11 @@ class __SYCL_EXPORT handler {
   // Make stream class friend to be able to keep the list of associated streams
   friend class stream;
   friend class detail::stream_impl;
-  // Make reduction_impl friend to store buffers and arrays created for it
-  // in handler from reduction_impl methods.
-  template <typename T, class BinaryOperation, int Dims, bool IsUSM,
-            access::placeholder IsPlaceholder>
-  friend class ext::oneapi::detail::reduction_impl;
+  // 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>
+  friend class ext::oneapi::detail::reduction_impl_algo;
 
   // This method needs to call the method finalize().
   template <typename Reduction, typename... RestT>

sycl/include/CL/sycl/reduction.hpp

@@ -20,7 +20,9 @@ namespace sycl {
 template <typename T, typename AllocatorT, typename BinaryOperation>
 std::enable_if_t<has_known_identity<BinaryOperation, T>::value,
                  ext::oneapi::detail::reduction_impl<
-                     T, BinaryOperation, 1, false, access::placeholder::true_t>>
+                     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 = {}) {
   bool InitializeToIdentity =
@@ -35,7 +37,9 @@ reduction(buffer<T, 1, AllocatorT> Var, handler &CGH, BinaryOperation,
 template <typename T, typename AllocatorT, typename BinaryOperation>
 std::enable_if_t<!has_known_identity<BinaryOperation, T>::value,
                  ext::oneapi::detail::reduction_impl<
-                     T, BinaryOperation, 1, false, access::placeholder::true_t>>
+                     T, BinaryOperation, 0, 1,
+                     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.
@@ -49,9 +53,11 @@ 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, 1, true>>
+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 = {}) {
   bool InitializeToIdentity =
       PropList.has_property<property::reduction::initialize_to_identity>();
@@ -64,9 +70,11 @@ 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, 1, true>>
+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 *, BinaryOperation, const property_list &PropList = {}) {
   // TODO: implement reduction that works even when identity is not known.
   (void)PropList;
@@ -79,8 +87,10 @@ 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, 1, false,
-                                    access::placeholder::true_t>
+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 = {}) {
   bool InitializeToIdentity =
@@ -92,13 +102,72 @@ 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, 1, true>
+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 = {}) {
   bool InitializeToIdentity =
       PropList.has_property<property::reduction::initialize_to_identity>();
   return {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 = {}) {
+  bool InitializeToIdentity =
+      PropList.has_property<property::reduction::initialize_to_identity>();
+  return {Span, InitializeToIdentity};
+}
+
+/// Constructs a reduction object using the reduction variable referenced by
+/// the given sycl::span \p Span, 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, 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 = {}) {
+  // TODO: implement reduction that works even when identity is not known.
+  (void)PropList;
+  throw runtime_error("Identity-less reductions with unknown identity are not "
+                      "supported yet.",
+                      PI_INVALID_VALUE);
+}
+
+/// 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 = {}) {
+  bool InitializeToIdentity =
+      PropList.has_property<property::reduction::initialize_to_identity>();
+  return {Span, Identity, Combiner, InitializeToIdentity};
+}
+#endif
+
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)