[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

@@ -49,12 +49,12 @@ This extension also depends on the following other SYCL extensions:
 
 == Status
 
-This is a proposed extension specification, intended to gather community
-feedback.  Interfaces defined in this specification may not be implemented yet
-or may be in a preliminary state.  The specification itself may also change in
-incompatible ways before it is finalized.  *Shipping software products should
-not rely on APIs defined in this specification.*
-
+This is an experimental extension specification, intended to provide early
+access to features and gather community feedback.  Interfaces defined in this
+specification are implemented in {dpcpp}, but they are not finalized and may
+change incompatibly in future versions of {dpcpp} without prior notice.
+*Shipping software products should not rely on APIs defined in this
+specification.*
 
 == Overview
 
@@ -87,7 +87,8 @@ implementation supports.
 |Description
 
 |1
-|Initial version of this extension.
+|The APIs of this experimental extension are not versioned, so the
+ feature-test macro always has this value.
 |===
 
 === `reduction` overload
@@ -111,7 +112,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,28 +125,35 @@ 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);
 
 }
 ----
 
+_Constraints_: Available only when `PropertyList` is an instance of
+`sycl::ext::oneapi::experimental::properties` which contains no properties
+other than those listed below in the section "Reduction properties".
+
+The `reduction` functions that take no `identity` parameter have the following
+clause:
+
+_Mandates_: If `properties` contains the `initialize_to_identity` property,
+then the identity of the `BinaryOperation` operation must be identifiable via
+the `known_identity` trait class.
+
 === Reduction properties
 
 New `reduction` properties are introduced to allow developers to constrain
 reduction algorithm selection based on desired behavior(s). Compile-time
 properties corresponding to existing runtime properties are also introduced to
 ensure that all information can be passed via a single property list.
 
-If a reduction kernel is submitted to a device that cannot satisfy the
-request for specific reduction behavior(s), the implementation must throw an
-`exception` with the `errc::feature_not_supported` error code.
-
 [source,c++]
 ----
-namespace sycl::ext::oneapi {
+namespace sycl::ext::oneapi::experimental {
 
 struct deterministic_key {
   using value_t = property_value<deterministic_key>;
@@ -192,7 +200,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 +213,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,143 @@
+//==------- 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
+    : detail::compile_time_property_key<detail::PropKind::Deterministic> {
+  using value_t = property_value<deterministic_key>;
+};
+inline constexpr deterministic_key::value_t deterministic;
+
+struct initialize_to_identity_key
+    : detail::compile_time_property_key<
+          detail::PropKind::InitializeToIdentity> {
+  using value_t = property_value<initialize_to_identity_key>;
+};
+inline constexpr initialize_to_identity_key::value_t initialize_to_identity;
+
+} // 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 T, typename BinaryOperation, typename PropertyList>
+void CheckReductionIdentity(PropertyList properties) {
+  if constexpr (properties.template has_property<
+                    ext::oneapi::experimental::initialize_to_identity_key>()) {
+    static_assert(has_known_identity_v<BinaryOperation, T>,
+                  "initialize_to_identity requires an identity value.");
+  }
+}
+
+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) {
+  detail::CheckReductionIdentity<typename BufferT::value_type, BinaryOperation>(
+      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) {
+  detail::CheckReductionIdentity<T, BinaryOperation>(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) {
+  detail::CheckReductionIdentity<T, BinaryOperation>(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

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

sycl/include/sycl/known_identity.hpp

@@ -25,54 +25,48 @@ namespace sycl {
 inline namespace _V1 {
 namespace detail {
 
+// Forward declaration for deterministic reductions.
+template <typename BinaryOperation> struct DeterministicOperatorWrapper;
+
+template <typename T, class BinaryOperation,
+          template <typename> class... KnownOperation>
+using IsKnownOp = std::bool_constant<(
+    (std::is_same_v<BinaryOperation, KnownOperation<T>> ||
+     std::is_same_v<BinaryOperation, KnownOperation<void>> ||
+     std::is_same_v<BinaryOperation,
+                    DeterministicOperatorWrapper<KnownOperation<T>>> ||
+     std::is_same_v<BinaryOperation,
+                    DeterministicOperatorWrapper<KnownOperation<void>>>) ||
+    ...)>;
+
 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 = IsKnownOp<T, BinaryOperation, sycl::plus>;
 
 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>>>;
+using IsMultiplies = IsKnownOp<T, BinaryOperation, sycl::multiplies>;
 
 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>>>;
+using IsMinimum = IsKnownOp<T, BinaryOperation, sycl::minimum>;
 
 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>>>;
+using IsMaximum = IsKnownOp<T, BinaryOperation, sycl::maximum>;
 
 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>>>;
+using IsBitAND = IsKnownOp<T, BinaryOperation, sycl::bit_and>;
 
 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>>>;
+using IsBitOR = IsKnownOp<T, BinaryOperation, sycl::bit_or>;
 
 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>>>;
+using IsBitXOR = IsKnownOp<T, BinaryOperation, sycl::bit_xor>;
 
 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 =
+    IsKnownOp<T, BinaryOperation, std::logical_and, sycl::logical_and>;
 
 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>>>;
+    IsKnownOp<T, BinaryOperation, std::logical_or, sycl::logical_or>;
 
 // Use SFINAE so that the "true" branch could be implemented in
 // include/sycl/stl_wrappers/complex that would only be available if STL's