Skip to content

Commit 22e5ced

Browse files
authored
[SYCL] Implement reduction properties extension (#15804)
Adds support for initialize_to_identity and deterministic properties. Since this extension is only experimental, the implementation here avoids making significant changes to reduction-related classes (e.g, reducer). A more straightforward implementation that attaches a compile-time property list to these classes is possible, but may be considered an ABI break. --- A note to reviewers: it occurred to me that `IsDeterministicOperator<BinaryOperation>` might actually be better as `IsDeterministicOperator<T, BinaryOperation>`, so that the implementation can infer that it's safe to use the faster reductions for certain types (e.g., `int` and `sycl::plus<>()` can safely use atomics) but I wanted to get your feedback on the proposed approach before going any further. **EDIT**: To clarify, I think such changes to `IsDeterministicOperator` could be left until a future PR, since the behavior of the `deterministic` property as implemented here should match the legacy `SYCL_DETERMINISTIC_REDUCTION` macro. We could implement optimizations later. --------- Signed-off-by: John Pennycook <john.pennycook@intel.com>
1 parent 07942fc commit 22e5ced

File tree

8 files changed

+303
-72
lines changed

8 files changed

+303
-72
lines changed

sycl/doc/extensions/proposed/sycl_ext_oneapi_reduction_properties.asciidoc renamed to sycl/doc/extensions/experimental/sycl_ext_oneapi_reduction_properties.asciidoc

Lines changed: 25 additions & 17 deletions
Original file line numberDiff line numberDiff line change
@@ -49,12 +49,12 @@ This extension also depends on the following other SYCL extensions:
4949

5050
== Status
5151

52-
This is a proposed extension specification, intended to gather community
53-
feedback. Interfaces defined in this specification may not be implemented yet
54-
or may be in a preliminary state. The specification itself may also change in
55-
incompatible ways before it is finalized. *Shipping software products should
56-
not rely on APIs defined in this specification.*
57-
52+
This is an experimental extension specification, intended to provide early
53+
access to features and gather community feedback. Interfaces defined in this
54+
specification are implemented in {dpcpp}, but they are not finalized and may
55+
change incompatibly in future versions of {dpcpp} without prior notice.
56+
*Shipping software products should not rely on APIs defined in this
57+
specification.*
5858

5959
== Overview
6060

@@ -87,7 +87,8 @@ implementation supports.
8787
|Description
8888

8989
|1
90-
|Initial version of this extension.
90+
|The APIs of this experimental extension are not versioned, so the
91+
feature-test macro always has this value.
9192
|===
9293

9394
=== `reduction` overload
@@ -111,7 +112,7 @@ template <typename T, typename BinaryOperation, typename PropertyList>
111112
__unspecified__ reduction(T* var, BinaryOperation combiner,
112113
PropertyList properties);
113114
114-
template <typename T, typename Extent, typename BinaryOperation, typename PropertyList>
115+
template <typename T, size_t Extent, typename BinaryOperation, typename PropertyList>
115116
__unspecified__ reduction(span<T, Extent> vars, BinaryOperation combiner,
116117
PropertyList properties);
117118
@@ -124,28 +125,35 @@ template <typename T, typename BinaryOperation, typename PropertyList>
124125
__unspecified__ reduction(T* var, const T& identity, BinaryOperation combiner,
125126
PropertyList properties);
126127
127-
template <typename T, typename Extent, typename BinaryOperation, typename PropertyList>
128+
template <typename T, size_t Extent, typename BinaryOperation, typename PropertyList>
128129
__unspecified__ reduction(span<T, Extent> vars, const T& identity,
129130
BinaryOperation combiner,
130131
PropertyList properties);
131132
132133
}
133134
----
134135

136+
_Constraints_: Available only when `PropertyList` is an instance of
137+
`sycl::ext::oneapi::experimental::properties` which contains no properties
138+
other than those listed below in the section "Reduction properties".
139+
140+
The `reduction` functions that take no `identity` parameter have the following
141+
clause:
142+
143+
_Mandates_: If `properties` contains the `initialize_to_identity` property,
144+
then the identity of the `BinaryOperation` operation must be identifiable via
145+
the `known_identity` trait class.
146+
135147
=== Reduction properties
136148

137149
New `reduction` properties are introduced to allow developers to constrain
138150
reduction algorithm selection based on desired behavior(s). Compile-time
139151
properties corresponding to existing runtime properties are also introduced to
140152
ensure that all information can be passed via a single property list.
141153

142-
If a reduction kernel is submitted to a device that cannot satisfy the
143-
request for specific reduction behavior(s), the implementation must throw an
144-
`exception` with the `errc::feature_not_supported` error code.
145-
146154
[source,c++]
147155
----
148-
namespace sycl::ext::oneapi {
156+
namespace sycl::ext::oneapi::experimental {
149157
150158
struct deterministic_key {
151159
using value_t = property_value<deterministic_key>;
@@ -192,7 +200,7 @@ use of atomic operations, etc. _{endnote}_]
192200

193201
[source,c++]
194202
----
195-
using syclex = sycl::ext::oneapi::experimental;
203+
namespace syclex = sycl::ext::oneapi::experimental;
196204
197205
float sum(sycl::queue q, float* input, size_t N) {
198206
@@ -205,10 +213,10 @@ float sum(sycl::queue q, float* input, size_t N) {
205213
h.parallel_for(N, reduction, [=](size_t i, auto& reducer) {
206214
reducer += input[i];
207215
});
208-
}
216+
});
209217
}
210218
return result;
211-
219+
212220
}
213221
214222
...
Lines changed: 143 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,143 @@
1+
//==------- properties.hpp - SYCL properties associated with reductions ----==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#pragma once
10+
#define SYCL_EXT_ONEAPI_REDUCTION_PROPERTIES
11+
12+
#include <sycl/ext/oneapi/properties/property.hpp>
13+
#include <sycl/ext/oneapi/properties/property_value.hpp>
14+
#include <sycl/reduction.hpp>
15+
16+
namespace sycl {
17+
inline namespace _V1 {
18+
namespace ext {
19+
namespace oneapi {
20+
namespace experimental {
21+
22+
struct deterministic_key
23+
: detail::compile_time_property_key<detail::PropKind::Deterministic> {
24+
using value_t = property_value<deterministic_key>;
25+
};
26+
inline constexpr deterministic_key::value_t deterministic;
27+
28+
struct initialize_to_identity_key
29+
: detail::compile_time_property_key<
30+
detail::PropKind::InitializeToIdentity> {
31+
using value_t = property_value<initialize_to_identity_key>;
32+
};
33+
inline constexpr initialize_to_identity_key::value_t initialize_to_identity;
34+
35+
} // namespace experimental
36+
} // namespace oneapi
37+
} // namespace ext
38+
39+
namespace detail {
40+
41+
template <typename BinaryOperation, typename PropertyList>
42+
auto WrapOp(BinaryOperation combiner, PropertyList properties) {
43+
if constexpr (properties.template has_property<
44+
ext::oneapi::experimental::deterministic_key>()) {
45+
return DeterministicOperatorWrapper(combiner);
46+
} else {
47+
return combiner;
48+
}
49+
}
50+
51+
template <typename T, typename BinaryOperation, typename PropertyList>
52+
void CheckReductionIdentity(PropertyList properties) {
53+
if constexpr (properties.template has_property<
54+
ext::oneapi::experimental::initialize_to_identity_key>()) {
55+
static_assert(has_known_identity_v<BinaryOperation, T>,
56+
"initialize_to_identity requires an identity value.");
57+
}
58+
}
59+
60+
template <typename PropertyList>
61+
property_list GetReductionPropertyList(PropertyList properties) {
62+
if constexpr (properties.template has_property<
63+
ext::oneapi::experimental::initialize_to_identity_key>()) {
64+
return sycl::property::reduction::initialize_to_identity{};
65+
}
66+
return {};
67+
}
68+
69+
template <typename BinaryOperation> struct DeterministicOperatorWrapper {
70+
71+
DeterministicOperatorWrapper(BinaryOperation BinOp = BinaryOperation())
72+
: BinOp(BinOp) {}
73+
74+
template <typename... Args>
75+
std::invoke_result_t<BinaryOperation, Args...> operator()(Args... args) {
76+
return BinOp(std::forward<Args>(args)...);
77+
}
78+
79+
BinaryOperation BinOp;
80+
};
81+
82+
template <typename BinaryOperation>
83+
struct IsDeterministicOperator<DeterministicOperatorWrapper<BinaryOperation>>
84+
: std::true_type {};
85+
86+
} // namespace detail
87+
88+
template <typename BufferT, typename BinaryOperation, typename PropertyList>
89+
auto reduction(BufferT vars, handler &cgh, BinaryOperation combiner,
90+
PropertyList properties) {
91+
detail::CheckReductionIdentity<typename BufferT::value_type, BinaryOperation>(
92+
properties);
93+
auto WrappedOp = detail::WrapOp(combiner, properties);
94+
auto RuntimeProps = detail::GetReductionPropertyList(properties);
95+
return reduction(vars, cgh, WrappedOp, RuntimeProps);
96+
}
97+
98+
template <typename T, typename BinaryOperation, typename PropertyList>
99+
auto reduction(T *var, BinaryOperation combiner, PropertyList properties) {
100+
detail::CheckReductionIdentity<T, BinaryOperation>(properties);
101+
auto WrappedOp = detail::WrapOp(combiner, properties);
102+
auto RuntimeProps = detail::GetReductionPropertyList(properties);
103+
return reduction(var, WrappedOp, RuntimeProps);
104+
}
105+
106+
template <typename T, size_t Extent, typename BinaryOperation,
107+
typename PropertyList>
108+
auto reduction(span<T, Extent> vars, BinaryOperation combiner,
109+
PropertyList properties) {
110+
detail::CheckReductionIdentity<T, BinaryOperation>(properties);
111+
auto WrappedOp = detail::WrapOp(combiner, properties);
112+
auto RuntimeProps = detail::GetReductionPropertyList(properties);
113+
return reduction(vars, WrappedOp, RuntimeProps);
114+
}
115+
116+
template <typename BufferT, typename BinaryOperation, typename PropertyList>
117+
auto reduction(BufferT vars, handler &cgh,
118+
const typename BufferT::value_type &identity,
119+
BinaryOperation combiner, PropertyList properties) {
120+
auto WrappedOp = detail::WrapOp(combiner, properties);
121+
auto RuntimeProps = detail::GetReductionPropertyList(properties);
122+
return reduction(vars, cgh, identity, WrappedOp, RuntimeProps);
123+
}
124+
125+
template <typename T, typename BinaryOperation, typename PropertyList>
126+
auto reduction(T *var, const T &identity, BinaryOperation combiner,
127+
PropertyList properties) {
128+
auto WrappedOp = detail::WrapOp(combiner, properties);
129+
auto RuntimeProps = detail::GetReductionPropertyList(properties);
130+
return reduction(var, identity, WrappedOp, RuntimeProps);
131+
}
132+
133+
template <typename T, size_t Extent, typename BinaryOperation,
134+
typename PropertyList>
135+
auto reduction(span<T, Extent> vars, const T &identity,
136+
BinaryOperation combiner, PropertyList properties) {
137+
auto WrappedOp = detail::WrapOp(combiner, properties);
138+
auto RuntimeProps = detail::GetReductionPropertyList(properties);
139+
return reduction(vars, identity, WrappedOp, RuntimeProps);
140+
}
141+
142+
} // namespace _V1
143+
} // namespace sycl

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

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -215,8 +215,10 @@ enum PropKind : uint32_t {
215215
MaxWorkGroupSize = 74,
216216
MaxLinearWorkGroupSize = 75,
217217
Prefetch = 76,
218+
Deterministic = 77,
219+
InitializeToIdentity = 78,
218220
// PropKindSize must always be the last value.
219-
PropKindSize = 77,
221+
PropKindSize = 79,
220222
};
221223

222224
struct property_key_base_tag {};

sycl/include/sycl/known_identity.hpp

Lines changed: 24 additions & 30 deletions
Original file line numberDiff line numberDiff line change
@@ -25,54 +25,48 @@ namespace sycl {
2525
inline namespace _V1 {
2626
namespace detail {
2727

28+
// Forward declaration for deterministic reductions.
29+
template <typename BinaryOperation> struct DeterministicOperatorWrapper;
30+
31+
template <typename T, class BinaryOperation,
32+
template <typename> class... KnownOperation>
33+
using IsKnownOp = std::bool_constant<(
34+
(std::is_same_v<BinaryOperation, KnownOperation<T>> ||
35+
std::is_same_v<BinaryOperation, KnownOperation<void>> ||
36+
std::is_same_v<BinaryOperation,
37+
DeterministicOperatorWrapper<KnownOperation<T>>> ||
38+
std::is_same_v<BinaryOperation,
39+
DeterministicOperatorWrapper<KnownOperation<void>>>) ||
40+
...)>;
41+
2842
template <typename T, class BinaryOperation>
29-
using IsPlus =
30-
std::bool_constant<std::is_same_v<BinaryOperation, sycl::plus<T>> ||
31-
std::is_same_v<BinaryOperation, sycl::plus<void>>>;
43+
using IsPlus = IsKnownOp<T, BinaryOperation, sycl::plus>;
3244

3345
template <typename T, class BinaryOperation>
34-
using IsMultiplies =
35-
std::bool_constant<std::is_same_v<BinaryOperation, sycl::multiplies<T>> ||
36-
std::is_same_v<BinaryOperation, sycl::multiplies<void>>>;
46+
using IsMultiplies = IsKnownOp<T, BinaryOperation, sycl::multiplies>;
3747

3848
template <typename T, class BinaryOperation>
39-
using IsMinimum =
40-
std::bool_constant<std::is_same_v<BinaryOperation, sycl::minimum<T>> ||
41-
std::is_same_v<BinaryOperation, sycl::minimum<void>>>;
49+
using IsMinimum = IsKnownOp<T, BinaryOperation, sycl::minimum>;
4250

4351
template <typename T, class BinaryOperation>
44-
using IsMaximum =
45-
std::bool_constant<std::is_same_v<BinaryOperation, sycl::maximum<T>> ||
46-
std::is_same_v<BinaryOperation, sycl::maximum<void>>>;
52+
using IsMaximum = IsKnownOp<T, BinaryOperation, sycl::maximum>;
4753

4854
template <typename T, class BinaryOperation>
49-
using IsBitAND =
50-
std::bool_constant<std::is_same_v<BinaryOperation, sycl::bit_and<T>> ||
51-
std::is_same_v<BinaryOperation, sycl::bit_and<void>>>;
55+
using IsBitAND = IsKnownOp<T, BinaryOperation, sycl::bit_and>;
5256

5357
template <typename T, class BinaryOperation>
54-
using IsBitOR =
55-
std::bool_constant<std::is_same_v<BinaryOperation, sycl::bit_or<T>> ||
56-
std::is_same_v<BinaryOperation, sycl::bit_or<void>>>;
58+
using IsBitOR = IsKnownOp<T, BinaryOperation, sycl::bit_or>;
5759

5860
template <typename T, class BinaryOperation>
59-
using IsBitXOR =
60-
std::bool_constant<std::is_same_v<BinaryOperation, sycl::bit_xor<T>> ||
61-
std::is_same_v<BinaryOperation, sycl::bit_xor<void>>>;
61+
using IsBitXOR = IsKnownOp<T, BinaryOperation, sycl::bit_xor>;
6262

6363
template <typename T, class BinaryOperation>
64-
using IsLogicalAND = std::bool_constant<
65-
std::is_same_v<BinaryOperation, std::logical_and<T>> ||
66-
std::is_same_v<BinaryOperation, std::logical_and<void>> ||
67-
std::is_same_v<BinaryOperation, sycl::logical_and<T>> ||
68-
std::is_same_v<BinaryOperation, sycl::logical_and<void>>>;
64+
using IsLogicalAND =
65+
IsKnownOp<T, BinaryOperation, std::logical_and, sycl::logical_and>;
6966

7067
template <typename T, class BinaryOperation>
7168
using IsLogicalOR =
72-
std::bool_constant<std::is_same_v<BinaryOperation, std::logical_or<T>> ||
73-
std::is_same_v<BinaryOperation, std::logical_or<void>> ||
74-
std::is_same_v<BinaryOperation, sycl::logical_or<T>> ||
75-
std::is_same_v<BinaryOperation, sycl::logical_or<void>>>;
69+
IsKnownOp<T, BinaryOperation, std::logical_or, sycl::logical_or>;
7670

7771
// Use SFINAE so that the "true" branch could be implemented in
7872
// include/sycl/stl_wrappers/complex that would only be available if STL's

0 commit comments

Comments
 (0)