[SYCL] Implement SYCL-2020 reductions with read_write to reduction vars (#3315)

* [SYCL] Implement SYCL-2020 reductions with read_write access to reduction variable

LIT tests: intel/llvm-test-suite#170

This patch:
- adds SYCL-2020 property::reduction::initialize_to_identity
- implements support for SYCL-2020 reductions for which initialize_to_identity
  is NOT used. Supported reductions correspond to read_write and
  USM ONEAPI::reductions.
- sycl::reduction re-uses ONEAPI::reduction implementation/classes and
  automatically creates placeholder accessors for sycl::reduction called
  with sycl::buffer argument.
- adds operator++ for reducer class defined in SYCL-2020
- fixes 2 errors in an ONEAPI::reduction used with placeholder accessors.

The attribute "initialize_to_identity" is added but NOT supported yet.
The multi-dimensional reductions are NOT supported yet.
Reductions accepting sycl::span are not supported yet.
Reductions cannot be yet used in parallel_for() accepting sycl::range

Signed-off-by: Vyacheslav N Klochkov <vyacheslav.n.klochkov@intel.com>

Author: v-klochkov

sycl/include/CL/sycl.hpp

@@ -45,6 +45,7 @@
 #include <CL/sycl/properties/all_properties.hpp>
 #include <CL/sycl/queue.hpp>
 #include <CL/sycl/range.hpp>
+#include <CL/sycl/reduction.hpp>
 #include <CL/sycl/sampler.hpp>
 #include <CL/sycl/stream.hpp>
 #include <CL/sycl/types.hpp>

sycl/include/CL/sycl/ONEAPI/reduction.hpp

@@ -221,6 +221,20 @@ class reducer {
 
   T getIdentity() const { return MIdentity; }
 
+  template <typename _T = T>
+  enable_if_t<IsReduPlus<_T, BinaryOperation>::value &&
+              sycl::detail::is_geninteger<_T>::value>
+  operator++() {
+    combine(static_cast<T>(1));
+  }
+
+  template <typename _T = T>
+  enable_if_t<IsReduPlus<_T, BinaryOperation>::value &&
+              sycl::detail::is_geninteger<_T>::value>
+  operator++(int) {
+    combine(static_cast<T>(1));
+  }
+
   template <typename _T = T>
   enable_if_t<IsReduPlus<_T, BinaryOperation>::value>
   operator+=(const _T &Partial) {
@@ -293,6 +307,20 @@ class reducer<T, BinaryOperation,
     return known_identity_impl<_BinaryOperation, _T>::value;
   }
 
+  template <typename _T = T>
+  enable_if_t<IsReduPlus<_T, BinaryOperation>::value &&
+              sycl::detail::is_geninteger<_T>::value>
+  operator++() {
+    combine(static_cast<T>(1));
+  }
+
+  template <typename _T = T>
+  enable_if_t<IsReduPlus<_T, BinaryOperation>::value &&
+              sycl::detail::is_geninteger<_T>::value>
+  operator++(int) {
+    combine(static_cast<T>(1));
+  }
+
   template <typename _T = T>
   enable_if_t<IsReduPlus<_T, BinaryOperation>::value>
   operator+=(const _T &Partial) {
@@ -419,7 +447,7 @@ class reduction_impl : private reduction_impl_base {
                ONEAPI::accessor_property_list<>>;
   using rw_accessor_type =
       accessor<T, Dims, access::mode::read_write, access::target::global_buffer,
-               IsPlaceholder, ONEAPI::accessor_property_list<>>;
+               access::placeholder::false_t, ONEAPI::accessor_property_list<>>;
   static constexpr access::mode accessor_mode = AccMode;
   static constexpr int accessor_dim = Dims;
   static constexpr int buffer_dim = (Dims == 0) ? 1 : Dims;
@@ -455,6 +483,20 @@ class reduction_impl : private reduction_impl_base {
     return MIdentity;
   }
 
+  /// SYCL-2020.
+  /// Constructs reduction_impl when the identity value is statically known.
+  template <typename _T, typename AllocatorT,
+            std::enable_if_t<IsKnownIdentityOp<_T, BinaryOperation>::value> * =
+                nullptr>
+  reduction_impl(buffer<_T, 1, AllocatorT> Buffer, handler &CGH)
+      : MAcc(std::make_shared<accessor_type>(Buffer)),
+        MIdentity(getIdentity()) {
+    associateWithHandler(CGH);
+    if (Buffer.get_count() != 1)
+      throw runtime_error("Reduction variable must be a scalar.",
+                          PI_INVALID_VALUE);
+  }
+
   /// Constructs reduction_impl when the identity value is statically known.
   // Note that aliasing constructor was used to initialize MAcc to avoid
   // destruction of the object referenced by the parameter Acc.
@@ -465,8 +507,36 @@ class reduction_impl : private reduction_impl_base {
       : MAcc(shared_ptr_class<accessor_type>(shared_ptr_class<accessor_type>{},
                                              &Acc)),
         MIdentity(getIdentity()) {
-    assert(Acc.get_count() == 1 &&
-           "Only scalar/1-element reductions are supported now.");
+    if (Acc.get_count() != 1)
+      throw runtime_error("Reduction variable must be a scalar.",
+                          PI_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<IsKnownIdentityOp<_T, BinaryOperation>::value> * = nullptr>
+  reduction_impl(buffer<_T, 1, AllocatorT> Buffer, handler &CGH,
+                 const T & /*Identity*/, BinaryOperation)
+      : MAcc(std::make_shared<accessor_type>(Buffer)),
+        MIdentity(getIdentity()) {
+    associateWithHandler(CGH);
+    if (Buffer.get_count() != 1)
+      throw runtime_error("Reduction variable must be a scalar.",
+                          PI_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,
@@ -476,13 +546,13 @@ class reduction_impl : private reduction_impl_base {
   template <
       typename _T = T, class _BinaryOperation = BinaryOperation,
       enable_if_t<IsKnownIdentityOp<_T, _BinaryOperation>::value> * = nullptr>
-  reduction_impl(accessor_type &Acc, const T &Identity, BinaryOperation)
+  reduction_impl(accessor_type &Acc, const T & /*Identity*/, BinaryOperation)
       : MAcc(shared_ptr_class<accessor_type>(shared_ptr_class<accessor_type>{},
                                              &Acc)),
         MIdentity(getIdentity()) {
-    (void)Identity;
-    assert(Acc.get_count() == 1 &&
-           "Only scalar/1-element reductions are supported now.");
+    if (Acc.get_count() != 1)
+      throw runtime_error("Reduction variable must be a scalar.",
+                          PI_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
@@ -496,6 +566,21 @@ class reduction_impl : private reduction_impl_base {
     // 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<!IsKnownIdentityOp<_T, BinaryOperation>::value> * = nullptr>
+  reduction_impl(buffer<_T, 1, AllocatorT> Buffer, handler &CGH,
+                 const T &Identity, BinaryOperation BOp)
+      : MAcc(std::make_shared<accessor_type>(Buffer)), MIdentity(Identity),
+        MBinaryOp(BOp) {
+    associateWithHandler(CGH);
+    if (Buffer.get_count() != 1)
+      throw runtime_error("Reduction variable must be a scalar.",
+                          PI_INVALID_VALUE);
+  }
+
   /// Constructs reduction_impl when the identity value is unknown.
   // Note that aliasing constructor was used to initialize MAcc to avoid
   // destruction of the object referenced by the parameter Acc.
@@ -506,8 +591,9 @@ class reduction_impl : private reduction_impl_base {
       : MAcc(shared_ptr_class<accessor_type>(shared_ptr_class<accessor_type>{},
                                              &Acc)),
         MIdentity(Identity), MBinaryOp(BOp) {
-    assert(Acc.get_count() == 1 &&
-           "Only scalar/1-element reductions are supported now.");
+    if (Acc.get_count() != 1)
+      throw runtime_error("Reduction variable must be a scalar.",
+                          PI_INVALID_VALUE);
   }
 
   /// Constructs reduction_impl when the identity value is statically known.
@@ -587,15 +673,29 @@ class reduction_impl : private reduction_impl_base {
   }
 
   /// Constructs a new temporary buffer to hold partial sums and returns
-  /// the accessor that that buffer.
-  template <bool IsOneWG>
-  std::enable_if_t<!IsOneWG, accessor_type>
+  /// the accessor for that buffer. Non-placeholder case.
+  template <bool IsOneWG, access::placeholder _IsPlaceholder = IsPlaceholder>
+  std::enable_if_t<!IsOneWG && _IsPlaceholder == access::placeholder::false_t,
+                   accessor_type>
   getWriteMemForPartialReds(size_t Size, handler &CGH) {
     MOutBufPtr = std::make_shared<buffer<T, buffer_dim>>(range<1>(Size));
     CGH.addReduction(MOutBufPtr);
     return accessor_type(*MOutBufPtr, CGH);
   }
 
+  /// Constructs a new temporary buffer to hold partial sums and returns
+  /// the accessor for that buffer. Placeholder case.
+  template <bool IsOneWG, access::placeholder _IsPlaceholder = IsPlaceholder>
+  std::enable_if_t<!IsOneWG && _IsPlaceholder == access::placeholder::true_t,
+                   accessor_type>
+  getWriteMemForPartialReds(size_t Size, handler &CGH) {
+    MOutBufPtr = std::make_shared<buffer<T, buffer_dim>>(range<1>(Size));
+    CGH.addReduction(MOutBufPtr);
+    accessor_type Acc(*MOutBufPtr);
+    CGH.require(Acc);
+    return Acc;
+  }
+
   template <access::placeholder _IsPlaceholder = IsPlaceholder>
   enable_if_t<_IsPlaceholder == access::placeholder::false_t, accessor_type>
   getWriteAccForPartialReds(size_t Size, handler &CGH) {
@@ -624,8 +724,7 @@ class reduction_impl : private reduction_impl_base {
 
   /// Creates 1-element global buffer initialized with identity value and
   /// returns an accessor to that buffer.
-  accessor<T, Dims, access::mode::read_write, access::target::global_buffer>
-  getReadWriteScalarAcc(handler &CGH) const {
+  rw_accessor_type getReadWriteScalarAcc(handler &CGH) const {
     auto RWReduVal = std::make_shared<T>(MIdentity);
     CGH.addReduction(RWReduVal);
     auto RWReduBuf =
@@ -1576,7 +1675,6 @@ template <typename T, class BinaryOperation, int Dims, access::mode AccMode,
 detail::reduction_impl<T, BinaryOperation, Dims, false, AccMode, IsPH>
 reduction(accessor<T, Dims, AccMode, access::target::global_buffer, IsPH> &Acc,
           const T &Identity, BinaryOperation BOp) {
-  // The Combiner argument was needed only to define the BinaryOperation param.
   return detail::reduction_impl<T, BinaryOperation, Dims, false, AccMode, IsPH>(
       Acc, Identity, BOp);
 }
@@ -1592,7 +1690,6 @@ std::enable_if_t<
     detail::reduction_impl<T, BinaryOperation, Dims, false, AccMode, IsPH>>
 reduction(accessor<T, Dims, AccMode, access::target::global_buffer, IsPH> &Acc,
           BinaryOperation) {
-  // The Combiner argument was needed only to define the BinaryOperation param.
   return detail::reduction_impl<T, BinaryOperation, Dims, false, AccMode, IsPH>(
       Acc);
 }
@@ -1643,29 +1740,5 @@ inline constexpr AccumulatorT known_identity_v =
     known_identity<BinaryOperation, AccumulatorT>::value;
 #endif
 } // namespace ONEAPI
-
-// Currently, the type traits defined below correspond to SYCL 1.2.1 ONEAPI
-// reduction extension. That may be changed later when SYCL 2020 reductions
-// are implemented.
-template <typename BinaryOperation, typename AccumulatorT>
-struct has_known_identity
-    : ONEAPI::has_known_identity<BinaryOperation, AccumulatorT> {};
-
-#if __cplusplus >= 201703L
-template <typename BinaryOperation, typename AccumulatorT>
-inline constexpr bool has_known_identity_v =
-    has_known_identity<BinaryOperation, AccumulatorT>::value;
-#endif
-
-template <typename BinaryOperation, typename AccumulatorT>
-struct known_identity : ONEAPI::known_identity<BinaryOperation, AccumulatorT> {
-};
-
-#if __cplusplus >= 201703L
-template <typename BinaryOperation, typename AccumulatorT>
-inline constexpr AccumulatorT known_identity_v =
-    known_identity<BinaryOperation, AccumulatorT>::value;
-#endif
-
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

sycl/include/CL/sycl/detail/properties_traits.def

@@ -8,3 +8,4 @@ __SYCL_PARAM_TRAITS_SPEC(sycl::ext::oneapi::property::buffer::use_pinned_host_me
 __SYCL_PARAM_TRAITS_SPEC(sycl::property::noinit)
 __SYCL_PARAM_TRAITS_SPEC(sycl::property::context::cuda::use_primary_context)
 __SYCL_PARAM_TRAITS_SPEC(sycl::property::queue::in_order)
+__SYCL_PARAM_TRAITS_SPEC(sycl::property::reduction::initialize_to_identity)

sycl/include/CL/sycl/detail/property_helper.hpp

@@ -31,7 +31,8 @@ enum DataLessPropKind {
   NoInit = 4,
   BufferUsePinnedHostMemory = 5,
   UsePrimaryContext = 6,
-  DataLessPropKindSize = 7
+  DataLessPropKindSize = 7,
+  InitializeToIdentity = 8
 };
 
 // List of all properties with data IDs

sycl/include/CL/sycl/properties/all_properties.hpp

@@ -3,3 +3,4 @@
 #include <CL/sycl/properties/context_properties.hpp>
 #include <CL/sycl/properties/image_properties.hpp>
 #include <CL/sycl/properties/queue_properties.hpp>
+#include <CL/sycl/properties/reduction_properties.hpp>

sycl/include/CL/sycl/properties/reduction_properties.hpp

@@ -0,0 +1,23 @@
+//==------- reduction_properties.hpp --- SYCL reduction properties ---------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <CL/sycl/context.hpp>
+#include <CL/sycl/detail/property_helper.hpp>
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace property {
+namespace reduction {
+class initialize_to_identity
+    : public detail::DataLessProperty<detail::InitializeToIdentity> {};
+} // namespace reduction
+} // namespace property
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)