[SYCL] Implement reduction properties extension

sycl/doc/extensions/proposed/sycl_ext_oneapi_reduction_properties.asciidoc → sycl/doc/extensions/experimental/sycl_ext_oneapi_reduction_properties.asciidoc

@@ -111,7 +111,7 @@ template <typename T, typename BinaryOperation, typename PropertyList>
 __unspecified__ reduction(T* var, BinaryOperation combiner,
                           PropertyList properties);
 
-template <typename T, typename Extent, typename BinaryOperation, typename PropertyList>
+template <typename T, size_t Extent, typename BinaryOperation, typename PropertyList>
 __unspecified__ reduction(span<T, Extent> vars, BinaryOperation combiner,
                           PropertyList properties);
 
@@ -124,7 +124,7 @@ template <typename T, typename BinaryOperation, typename PropertyList>
 __unspecified__ reduction(T* var, const T& identity, BinaryOperation combiner,
                           PropertyList properties);
 
-template <typename T, typename Extent, typename BinaryOperation, typename PropertyList>
+template <typename T, size_t Extent, typename BinaryOperation, typename PropertyList>
 __unspecified__ reduction(span<T, Extent> vars, const T& identity,
                           BinaryOperation combiner,
                           PropertyList properties);
@@ -192,7 +192,7 @@ use of atomic operations, etc. _{endnote}_]
 
 [source,c++]
 ----
-using syclex = sycl::ext::oneapi::experimental;
+namespace syclex = sycl::ext::oneapi::experimental;
 
 float sum(sycl::queue q, float* input, size_t N) {
 
@@ -205,10 +205,10 @@ float sum(sycl::queue q, float* input, size_t N) {
       h.parallel_for(N, reduction, [=](size_t i, auto& reducer) {
         reducer += input[i];
       });
-    }
+    });
   }
   return result;
-  
+
 }
 
 ...

sycl/include/sycl/ext/oneapi/experimental/reduction_properties.hpp

@@ -0,0 +1,145 @@
+//==------- properties.hpp - SYCL properties associated with reductions ----==//
+//
+// 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
+#define SYCL_EXT_ONEAPI_REDUCTION_PROPERTIES
+
+#include <sycl/ext/oneapi/properties/property.hpp>
+#include <sycl/ext/oneapi/properties/property_value.hpp>
+#include <sycl/reduction.hpp>
+
+namespace sycl {
+inline namespace _V1 {
+namespace ext {
+namespace oneapi {
+namespace experimental {
+
+struct deterministic_key {
+  using value_t = property_value<deterministic_key>;
+};
+inline constexpr deterministic_key::value_t deterministic;
+
+struct initialize_to_identity_key {
+  using value_t = property_value<initialize_to_identity_key>;
+};
+inline constexpr initialize_to_identity_key::value_t initialize_to_identity;
+
+template <> struct is_property_key<deterministic_key> : std::true_type {};
+template <>
+struct is_property_key<initialize_to_identity_key> : std::true_type {};
+
+namespace detail {
+
+template <> struct IsCompileTimeProperty<deterministic_key> : std::true_type {};
+template <>
+struct IsCompileTimeProperty<initialize_to_identity_key> : std::true_type {};
+
+template <> struct PropertyToKind<deterministic_key> {
+  static constexpr PropKind Kind = PropKind::Deterministic;
+};
+template <> struct PropertyToKind<initialize_to_identity_key> {
+  static constexpr PropKind Kind = PropKind::InitializeToIdentity;
+};
+
+} // namespace detail
+} // namespace experimental
+} // namespace oneapi
+} // namespace ext
+
+namespace detail {
+
+template <typename BinaryOperation, typename PropertyList>
+auto WrapOp(BinaryOperation combiner, PropertyList properties) {
+  if constexpr (properties.template has_property<
+                    ext::oneapi::experimental::deterministic_key>()) {
+    return DeterministicOperatorWrapper(combiner);
+  } else {
+    return combiner;
+  }
+}
+
+template <typename PropertyList>
+property_list GetReductionPropertyList(PropertyList properties) {
+  if constexpr (properties.template has_property<
+                    ext::oneapi::experimental::initialize_to_identity_key>()) {
+    return sycl::property::reduction::initialize_to_identity{};
+  }
+  return {};
+}
+
+template <typename BinaryOperation> struct DeterministicOperatorWrapper {
+
+  DeterministicOperatorWrapper(BinaryOperation BinOp = BinaryOperation())
+      : BinOp(BinOp) {}
+
+  template <typename... Args>
+  std::invoke_result_t<BinaryOperation, Args...> operator()(Args... args) {
+    return BinOp(std::forward<Args>(args)...);
+  }
+
+  BinaryOperation BinOp;
+};
+
+template <typename BinaryOperation>
+struct IsDeterministicOperator<DeterministicOperatorWrapper<BinaryOperation>>
+    : std::true_type {};
+
+} // namespace detail
+
+template <typename BufferT, typename BinaryOperation, typename PropertyList>
+auto reduction(BufferT vars, handler &cgh, BinaryOperation combiner,
+               PropertyList properties) {
+  auto WrappedOp = detail::WrapOp(combiner, properties);
+  auto RuntimeProps = detail::GetReductionPropertyList(properties);
+  return reduction(vars, cgh, WrappedOp, RuntimeProps);
+}
+
+template <typename T, typename BinaryOperation, typename PropertyList>
+auto reduction(T *var, BinaryOperation combiner, PropertyList properties) {
+  auto WrappedOp = detail::WrapOp(combiner, properties);
+  auto RuntimeProps = detail::GetReductionPropertyList(properties);
+  return reduction(var, WrappedOp, RuntimeProps);
+}
+
+template <typename T, size_t Extent, typename BinaryOperation,
+          typename PropertyList>
+auto reduction(span<T, Extent> vars, BinaryOperation combiner,
+               PropertyList properties) {
+  auto WrappedOp = detail::WrapOp(combiner, properties);
+  auto RuntimeProps = detail::GetReductionPropertyList(properties);
+  return reduction(vars, WrappedOp, RuntimeProps);
+}
+
+template <typename BufferT, typename BinaryOperation, typename PropertyList>
+auto reduction(BufferT vars, handler &cgh,
+               const typename BufferT::value_type &identity,
+               BinaryOperation combiner, PropertyList properties) {
+  auto WrappedOp = detail::WrapOp(combiner, properties);
+  auto RuntimeProps = detail::GetReductionPropertyList(properties);
+  return reduction(vars, cgh, identity, WrappedOp, RuntimeProps);
+}
+
+template <typename T, typename BinaryOperation, typename PropertyList>
+auto reduction(T *var, const T &identity, BinaryOperation combiner,
+               PropertyList properties) {
+  auto WrappedOp = detail::WrapOp(combiner, properties);
+  auto RuntimeProps = detail::GetReductionPropertyList(properties);
+  return reduction(var, identity, WrappedOp, RuntimeProps);
+}
+
+template <typename T, size_t Extent, typename BinaryOperation,
+          typename PropertyList>
+auto reduction(span<T, Extent> vars, const T &identity,
+               BinaryOperation combiner, PropertyList properties) {
+  auto WrappedOp = detail::WrapOp(combiner, properties);
+  auto RuntimeProps = detail::GetReductionPropertyList(properties);
+  return reduction(vars, identity, WrappedOp, RuntimeProps);
+}
+
+} // namespace _V1
+} // namespace sycl

sycl/include/sycl/ext/oneapi/properties/property.hpp

@@ -214,8 +214,10 @@ enum PropKind : uint32_t {
   ResponseCapacity = 73,
   MaxWorkGroupSize = 74,
   MaxLinearWorkGroupSize = 75,
+  Deterministic = 76,
+  InitializeToIdentity = 77,
   // PropKindSize must always be the last value.
-  PropKindSize = 76,
+  PropKindSize = 78,
 };
 
 struct property_key_base_tag {};

sycl/include/sycl/known_identity.hpp

@@ -25,54 +25,67 @@ namespace sycl {
 inline namespace _V1 {
 namespace detail {
 
+// Forward declaration for deterministic reductions.
+template <typename BinaryOperation> struct DeterministicOperatorWrapper;
+
+template <typename T, typename U> struct is_like : public std::is_same<T, U> {};
+
+template <typename T, typename U>
+constexpr bool is_like_v = is_like<T, U>::value;
+
+template <typename T, typename U>
+struct is_like<DeterministicOperatorWrapper<T>, U> : std::is_same<T, U> {};
+
+template <typename T, typename U>
+struct is_like<T, DeterministicOperatorWrapper<U>> : std::is_same<T, U> {};
+
 template <typename T, class BinaryOperation>
-using IsPlus =
-    std::bool_constant<std::is_same_v<BinaryOperation, sycl::plus<T>> ||
-                       std::is_same_v<BinaryOperation, sycl::plus<void>>>;
+using IsPlus = std::bool_constant<is_like_v<BinaryOperation, sycl::plus<T>> ||
+                                  is_like_v<BinaryOperation, sycl::plus<void>>>;
 
 template <typename T, class BinaryOperation>
 using IsMultiplies =
-    std::bool_constant<std::is_same_v<BinaryOperation, sycl::multiplies<T>> ||
-                       std::is_same_v<BinaryOperation, sycl::multiplies<void>>>;
+    std::bool_constant<is_like_v<BinaryOperation, sycl::multiplies<T>> ||
+                       is_like_v<BinaryOperation, sycl::multiplies<void>>>;
 
 template <typename T, class BinaryOperation>
 using IsMinimum =
-    std::bool_constant<std::is_same_v<BinaryOperation, sycl::minimum<T>> ||
-                       std::is_same_v<BinaryOperation, sycl::minimum<void>>>;
+    std::bool_constant<is_like_v<BinaryOperation, sycl::minimum<T>> ||
+                       is_like_v<BinaryOperation, sycl::minimum<void>>>;
 
 template <typename T, class BinaryOperation>
 using IsMaximum =
-    std::bool_constant<std::is_same_v<BinaryOperation, sycl::maximum<T>> ||
-                       std::is_same_v<BinaryOperation, sycl::maximum<void>>>;
+    std::bool_constant<is_like_v<BinaryOperation, sycl::maximum<T>> ||
+                       is_like_v<BinaryOperation, sycl::maximum<void>>>;
 
 template <typename T, class BinaryOperation>
 using IsBitAND =
-    std::bool_constant<std::is_same_v<BinaryOperation, sycl::bit_and<T>> ||
-                       std::is_same_v<BinaryOperation, sycl::bit_and<void>>>;
+    std::bool_constant<is_like_v<BinaryOperation, sycl::bit_and<T>> ||
+                       is_like_v<BinaryOperation, sycl::bit_and<void>>>;
 
 template <typename T, class BinaryOperation>
 using IsBitOR =
-    std::bool_constant<std::is_same_v<BinaryOperation, sycl::bit_or<T>> ||
-                       std::is_same_v<BinaryOperation, sycl::bit_or<void>>>;
+    std::bool_constant<is_like_v<BinaryOperation, sycl::bit_or<T>> ||
+                       is_like_v<BinaryOperation, sycl::bit_or<void>>>;
 
 template <typename T, class BinaryOperation>
 using IsBitXOR =
-    std::bool_constant<std::is_same_v<BinaryOperation, sycl::bit_xor<T>> ||
-                       std::is_same_v<BinaryOperation, sycl::bit_xor<void>>>;
+    std::bool_constant<is_like_v<BinaryOperation, sycl::bit_xor<T>> ||
+                       is_like_v<BinaryOperation, sycl::bit_xor<void>>>;
 
 template <typename T, class BinaryOperation>
-using IsLogicalAND = std::bool_constant<
-    std::is_same_v<BinaryOperation, std::logical_and<T>> ||
-    std::is_same_v<BinaryOperation, std::logical_and<void>> ||
-    std::is_same_v<BinaryOperation, sycl::logical_and<T>> ||
-    std::is_same_v<BinaryOperation, sycl::logical_and<void>>>;
+using IsLogicalAND =
+    std::bool_constant<is_like_v<BinaryOperation, std::logical_and<T>> ||
+                       is_like_v<BinaryOperation, std::logical_and<void>> ||
+                       is_like_v<BinaryOperation, sycl::logical_and<T>> ||
+                       is_like_v<BinaryOperation, sycl::logical_and<void>>>;
 
 template <typename T, class BinaryOperation>
 using IsLogicalOR =
-    std::bool_constant<std::is_same_v<BinaryOperation, std::logical_or<T>> ||
-                       std::is_same_v<BinaryOperation, std::logical_or<void>> ||
-                       std::is_same_v<BinaryOperation, sycl::logical_or<T>> ||
-                       std::is_same_v<BinaryOperation, sycl::logical_or<void>>>;
+    std::bool_constant<is_like_v<BinaryOperation, std::logical_or<T>> ||
+                       is_like_v<BinaryOperation, std::logical_or<void>> ||
+                       is_like_v<BinaryOperation, sycl::logical_or<T>> ||
+                       is_like_v<BinaryOperation, sycl::logical_or<void>>>;
 
 // Use SFINAE so that the "true" branch could be implemented in
 // include/sycl/stl_wrappers/complex that would only be available if STL's

sycl/include/sycl/reduction.hpp

@@ -77,20 +77,24 @@ template <typename T, class BinaryOperation, int Dims, size_t Extent,
 class reducer;
 
 namespace detail {
+
+#ifdef SYCL_DETERMINISTIC_REDUCTION
+// Act as if all operators require determinism.
+template <typename T> struct IsDeterministicOperator : std::true_type {};
+#else
+// Each operator declares whether determinism is required.
+template <typename T> struct IsDeterministicOperator : std::false_type {};
+#endif
+
 // This type trait is used to detect if the atomic operation BinaryOperation
 // used with operands of the type T is available for using in reduction.
 // The order in which the atomic operations are performed may be arbitrary and
 // thus may cause different results from run to run even on the same elements
-// and on same device. The macro SYCL_REDUCTION_DETERMINISTIC prohibits using
-// atomic operations for reduction and helps to produce stable results.
-// SYCL_REDUCTION_DETERMINISTIC is a short term solution, which perhaps become
-// deprecated eventually and is replaced by a sycl property passed to reduction.
+// and on same device.
 template <typename T, class BinaryOperation>
 using IsReduOptForFastAtomicFetch =
-#ifdef SYCL_REDUCTION_DETERMINISTIC
-    std::bool_constant<false>;
-#else
-    std::bool_constant<((is_sgenfloat_v<T> && sizeof(T) == 4) ||
+    std::bool_constant<!IsDeterministicOperator<BinaryOperation>::value &&
+                       ((is_sgenfloat_v<T> && sizeof(T) == 4) ||
                         is_sgeninteger_v<T>) &&
                        IsValidAtomicType<T>::value &&
                        (IsPlus<T, BinaryOperation>::value ||
@@ -99,44 +103,33 @@ using IsReduOptForFastAtomicFetch =
                         IsBitOR<T, BinaryOperation>::value ||
                         IsBitXOR<T, BinaryOperation>::value ||
                         IsBitAND<T, BinaryOperation>::value)>;
-#endif
 
 // This type trait is used to detect if the atomic operation BinaryOperation
 // used with operands of the type T is available for using in reduction, in
 // addition to the cases covered by "IsReduOptForFastAtomicFetch", if the device
 // has the atomic64 aspect. This type trait should only be used if the device
 // has the atomic64 aspect.  Note that this type trait is currently a subset of
-// IsReduOptForFastReduce. The macro SYCL_REDUCTION_DETERMINISTIC prohibits
-// using the reduce_over_group() algorithm to produce stable results across same
-// type devices.
+// IsReduOptForFastReduce.
 template <typename T, class BinaryOperation>
 using IsReduOptForAtomic64Op =
-#ifdef SYCL_REDUCTION_DETERMINISTIC
-    std::bool_constant<false>;
-#else
-    std::bool_constant<(IsPlus<T, BinaryOperation>::value ||
+    std::bool_constant<!IsDeterministicOperator<BinaryOperation>::value &&
+                       (IsPlus<T, BinaryOperation>::value ||
                         IsMinimum<T, BinaryOperation>::value ||
                         IsMaximum<T, BinaryOperation>::value) &&
                        is_sgenfloat_v<T> && sizeof(T) == 8>;
-#endif
 
 // This type trait is used to detect if the group algorithm reduce() used with
 // operands of the type T and the operation BinaryOperation is available
 // for using in reduction.
-// The macro SYCL_REDUCTION_DETERMINISTIC prohibits using the reduce() algorithm
-// to produce stable results across same type devices.
 template <typename T, class BinaryOperation>
 using IsReduOptForFastReduce =
-#ifdef SYCL_REDUCTION_DETERMINISTIC
-    std::bool_constant<false>;
-#else
-    std::bool_constant<((is_sgeninteger_v<T> &&
+    std::bool_constant<!IsDeterministicOperator<BinaryOperation>::value &&
+                       ((is_sgeninteger_v<T> &&
                          (sizeof(T) == 4 || sizeof(T) == 8)) ||
                         is_sgenfloat_v<T>) &&
                        (IsPlus<T, BinaryOperation>::value ||
                         IsMinimum<T, BinaryOperation>::value ||
                         IsMaximum<T, BinaryOperation>::value)>;
-#endif
 
 // std::tuple seems to be a) too heavy and b) not copyable to device now
 // Thus sycl::detail::tuple is used instead.

sycl/include/sycl/sycl.hpp

@@ -98,6 +98,7 @@
 #include <sycl/ext/oneapi/experimental/prefetch.hpp>
 #include <sycl/ext/oneapi/experimental/profiling_tag.hpp>
 #include <sycl/ext/oneapi/experimental/raw_kernel_arg.hpp>
+#include <sycl/ext/oneapi/experimental/reduction_properties.hpp>
 #include <sycl/ext/oneapi/experimental/root_group.hpp>
 #include <sycl/ext/oneapi/experimental/tangle_group.hpp>
 #include <sycl/ext/oneapi/filter_selector.hpp>