Skip to content

Commit b94cb3f

Browse files
PennycookPennycook
committed
[SYCL] Implement reduction properties extension
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. Signed-off-by: John Pennycook <[email protected]>
1 parent 761d45d commit b94cb3f

File tree

5 files changed

+202
-49
lines changed

5 files changed

+202
-49
lines changed

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

Lines changed: 144 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,144 @@
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+
15+
namespace sycl {
16+
inline namespace _V1 {
17+
namespace ext {
18+
namespace oneapi {
19+
namespace experimental {
20+
21+
struct deterministic_key {
22+
using value_t = property_value<deterministic_key>;
23+
};
24+
inline constexpr deterministic_key::value_t deterministic;
25+
26+
struct initialize_to_identity_key {
27+
using value_t = property_value<initialize_to_identity_key>;
28+
};
29+
inline constexpr initialize_to_identity_key::value_t initialize_to_identity;
30+
31+
template <> struct is_property_key<deterministic_key> : std::true_type {};
32+
template <>
33+
struct is_property_key<initialize_to_identity_key> : std::true_type {};
34+
35+
namespace detail {
36+
37+
template <> struct IsCompileTimeProperty<deterministic_key> : std::true_type {};
38+
template <>
39+
struct IsCompileTimeProperty<initialize_to_identity_key> : std::true_type {};
40+
41+
template <> struct PropertyToKind<deterministic_key> {
42+
static constexpr PropKind Kind = PropKind::Deterministic;
43+
};
44+
template <> struct PropertyToKind<initialize_to_identity_key> {
45+
static constexpr PropKind Kind = PropKind::InitializeToIdentity;
46+
};
47+
48+
} // namespace detail
49+
} // namespace experimental
50+
} // namespace oneapi
51+
} // namespace ext
52+
53+
namespace detail {
54+
55+
template <typename BinaryOperation, typename PropertyList>
56+
auto WrapOp(BinaryOperation combiner, PropertyList properties) {
57+
if constexpr (properties.template has_property<
58+
ext::oneapi::experimental::deterministic_key>()) {
59+
return DeterministicOperatorWrapper(combiner);
60+
} else {
61+
return combiner;
62+
}
63+
}
64+
65+
template <typename PropertyList>
66+
property_list GetReductionPropertyList(PropertyList properties) {
67+
if constexpr (properties.template has_property<
68+
ext::oneapi::experimental::initialize_to_identity_key>()) {
69+
return sycl::property::reduction::initialize_to_identity{};
70+
}
71+
return {};
72+
}
73+
74+
template <typename BinaryOperation> struct DeterministicOperatorWrapper {
75+
76+
DeterministicOperatorWrapper(BinaryOperation BinOp = BinaryOperation())
77+
: BinOp(BinOp) {}
78+
79+
template <typename... Args>
80+
std::invoke_result_t<BinaryOperation, Args...> operator()(Args... args) {
81+
return BinOp(std::forward<Args>(args)...);
82+
}
83+
84+
BinaryOperation BinOp;
85+
};
86+
87+
template <typename BinaryOperation>
88+
struct IsDeterministicOperator<DeterministicOperatorWrapper<BinaryOperation>>
89+
: std::true_type {};
90+
91+
} // namespace detail
92+
93+
template <typename BufferT, typename BinaryOperation, typename PropertyList>
94+
auto reduction(BufferT vars, handler &cgh, BinaryOperation combiner,
95+
PropertyList properties) {
96+
auto WrappedOp = detail::WrapOp(combiner, properties);
97+
auto RuntimeProps = detail::GetReductionPropertyList(properties);
98+
return reduction(vars, cgh, WrappedOp, RuntimeProps);
99+
}
100+
101+
template <typename T, typename BinaryOperation, typename PropertyList>
102+
auto reduction(T *var, BinaryOperation combiner, PropertyList properties) {
103+
auto WrappedOp = detail::WrapOp(combiner, properties);
104+
auto RuntimeProps = detail::GetReductionPropertyList(properties);
105+
return reduction(var, WrappedOp, RuntimeProps);
106+
}
107+
108+
template <typename T, size_t Extent, typename BinaryOperation,
109+
typename PropertyList>
110+
auto reduction(span<T, Extent> vars, BinaryOperation combiner,
111+
PropertyList properties) {
112+
auto WrappedOp = detail::WrapOp(combiner, properties);
113+
auto RuntimeProps = detail::GetReductionPropertyList(properties);
114+
return reduction(vars, WrappedOp, RuntimeProps);
115+
}
116+
117+
template <typename BufferT, typename BinaryOperation, typename PropertyList>
118+
auto reduction(BufferT vars, handler &cgh,
119+
const typename BufferT::value_type &identity,
120+
BinaryOperation combiner, PropertyList properties) {
121+
auto WrappedOp = detail::WrapOp(combiner, properties);
122+
auto RuntimeProps = detail::GetReductionPropertyList(properties);
123+
return reduction(vars, cgh, identity, WrappedOp, RuntimeProps);
124+
}
125+
126+
template <typename T, typename BinaryOperation, typename PropertyList>
127+
auto reduction(T *var, const T &identity, BinaryOperation combiner,
128+
PropertyList properties) {
129+
auto WrappedOp = detail::WrapOp(combiner, properties);
130+
auto RuntimeProps = detail::GetReductionPropertyList(properties);
131+
return reduction(var, identity, WrappedOp, RuntimeProps);
132+
}
133+
134+
template <typename T, size_t Extent, typename BinaryOperation,
135+
typename PropertyList>
136+
auto reduction(span<T, Extent> vars, const T &identity,
137+
BinaryOperation combiner, PropertyList properties) {
138+
auto WrappedOp = detail::WrapOp(combiner, properties);
139+
auto RuntimeProps = detail::GetReductionPropertyList(properties);
140+
return reduction(vars, identity, WrappedOp, RuntimeProps);
141+
}
142+
143+
} // namespace _V1
144+
} // namespace sycl

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

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -214,8 +214,10 @@ enum PropKind : uint32_t {
214214
ResponseCapacity = 73,
215215
MaxWorkGroupSize = 74,
216216
MaxLinearWorkGroupSize = 75,
217+
Deterministic = 76,
218+
InitializeToIdentity = 77,
217219
// PropKindSize must always be the last value.
218-
PropKindSize = 76,
220+
PropKindSize = 78,
219221
};
220222

221223
struct property_key_base_tag {};

‎sycl/include/sycl/known_identity.hpp‎

Lines changed: 37 additions & 24 deletions
Original file line numberDiff line numberDiff line change
@@ -25,54 +25,67 @@ 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, typename U> struct is_like : public std::is_same<T, U> {};
32+
33+
template <typename T, typename U>
34+
constexpr bool is_like_v = is_like<T, U>::value;
35+
36+
template <typename T, typename U>
37+
struct is_like<DeterministicOperatorWrapper<T>, U> : std::is_same<T, U> {};
38+
39+
template <typename T, typename U>
40+
struct is_like<T, DeterministicOperatorWrapper<U>> : std::is_same<T, U> {};
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 = std::bool_constant<is_like_v<BinaryOperation, sycl::plus<T>> ||
44+
is_like_v<BinaryOperation, sycl::plus<void>>>;
3245

3346
template <typename T, class BinaryOperation>
3447
using IsMultiplies =
35-
std::bool_constant<std::is_same_v<BinaryOperation, sycl::multiplies<T>> ||
36-
std::is_same_v<BinaryOperation, sycl::multiplies<void>>>;
48+
std::bool_constant<is_like_v<BinaryOperation, sycl::multiplies<T>> ||
49+
is_like_v<BinaryOperation, sycl::multiplies<void>>>;
3750

3851
template <typename T, class BinaryOperation>
3952
using IsMinimum =
40-
std::bool_constant<std::is_same_v<BinaryOperation, sycl::minimum<T>> ||
41-
std::is_same_v<BinaryOperation, sycl::minimum<void>>>;
53+
std::bool_constant<is_like_v<BinaryOperation, sycl::minimum<T>> ||
54+
is_like_v<BinaryOperation, sycl::minimum<void>>>;
4255

4356
template <typename T, class BinaryOperation>
4457
using IsMaximum =
45-
std::bool_constant<std::is_same_v<BinaryOperation, sycl::maximum<T>> ||
46-
std::is_same_v<BinaryOperation, sycl::maximum<void>>>;
58+
std::bool_constant<is_like_v<BinaryOperation, sycl::maximum<T>> ||
59+
is_like_v<BinaryOperation, sycl::maximum<void>>>;
4760

4861
template <typename T, class BinaryOperation>
4962
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>>>;
63+
std::bool_constant<is_like_v<BinaryOperation, sycl::bit_and<T>> ||
64+
is_like_v<BinaryOperation, sycl::bit_and<void>>>;
5265

5366
template <typename T, class BinaryOperation>
5467
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>>>;
68+
std::bool_constant<is_like_v<BinaryOperation, sycl::bit_or<T>> ||
69+
is_like_v<BinaryOperation, sycl::bit_or<void>>>;
5770

5871
template <typename T, class BinaryOperation>
5972
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>>>;
73+
std::bool_constant<is_like_v<BinaryOperation, sycl::bit_xor<T>> ||
74+
is_like_v<BinaryOperation, sycl::bit_xor<void>>>;
6275

6376
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>>>;
77+
using IsLogicalAND =
78+
std::bool_constant<is_like_v<BinaryOperation, std::logical_and<T>> ||
79+
is_like_v<BinaryOperation, std::logical_and<void>> ||
80+
is_like_v<BinaryOperation, sycl::logical_and<T>> ||
81+
is_like_v<BinaryOperation, sycl::logical_and<void>>>;
6982

7083
template <typename T, class BinaryOperation>
7184
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>>>;
85+
std::bool_constant<is_like_v<BinaryOperation, std::logical_or<T>> ||
86+
is_like_v<BinaryOperation, std::logical_or<void>> ||
87+
is_like_v<BinaryOperation, sycl::logical_or<T>> ||
88+
is_like_v<BinaryOperation, sycl::logical_or<void>>>;
7689

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

‎sycl/include/sycl/reduction.hpp‎

Lines changed: 17 additions & 24 deletions
Original file line numberDiff line numberDiff line change
@@ -77,20 +77,24 @@ template <typename T, class BinaryOperation, int Dims, size_t Extent,
7777
class reducer;
7878

7979
namespace detail {
80+
81+
#ifdef SYCL_DETERMINISTIC_REDUCTION
82+
// Act as if all operators require determinism.
83+
template <typename T> struct IsDeterministicOperator : std::true_type {};
84+
#else
85+
// Each operator declares whether determinism is required.
86+
template <typename T> struct IsDeterministicOperator : std::false_type {};
87+
#endif
88+
8089
// This type trait is used to detect if the atomic operation BinaryOperation
8190
// used with operands of the type T is available for using in reduction.
8291
// The order in which the atomic operations are performed may be arbitrary and
8392
// thus may cause different results from run to run even on the same elements
84-
// and on same device. The macro SYCL_REDUCTION_DETERMINISTIC prohibits using
85-
// atomic operations for reduction and helps to produce stable results.
86-
// SYCL_REDUCTION_DETERMINISTIC is a short term solution, which perhaps become
87-
// deprecated eventually and is replaced by a sycl property passed to reduction.
93+
// and on same device.
8894
template <typename T, class BinaryOperation>
8995
using IsReduOptForFastAtomicFetch =
90-
#ifdef SYCL_REDUCTION_DETERMINISTIC
91-
std::bool_constant<false>;
92-
#else
93-
std::bool_constant<((is_sgenfloat_v<T> && sizeof(T) == 4) ||
96+
std::bool_constant<not IsDeterministicOperator<BinaryOperation>::value &&
97+
((is_sgenfloat_v<T> && sizeof(T) == 4) ||
9498
is_sgeninteger_v<T>) &&
9599
IsValidAtomicType<T>::value &&
96100
(IsPlus<T, BinaryOperation>::value ||
@@ -99,44 +103,33 @@ using IsReduOptForFastAtomicFetch =
99103
IsBitOR<T, BinaryOperation>::value ||
100104
IsBitXOR<T, BinaryOperation>::value ||
101105
IsBitAND<T, BinaryOperation>::value)>;
102-
#endif
103106

104107
// This type trait is used to detect if the atomic operation BinaryOperation
105108
// used with operands of the type T is available for using in reduction, in
106109
// addition to the cases covered by "IsReduOptForFastAtomicFetch", if the device
107110
// has the atomic64 aspect. This type trait should only be used if the device
108111
// has the atomic64 aspect. Note that this type trait is currently a subset of
109-
// IsReduOptForFastReduce. The macro SYCL_REDUCTION_DETERMINISTIC prohibits
110-
// using the reduce_over_group() algorithm to produce stable results across same
111-
// type devices.
112+
// IsReduOptForFastReduce.
112113
template <typename T, class BinaryOperation>
113114
using IsReduOptForAtomic64Op =
114-
#ifdef SYCL_REDUCTION_DETERMINISTIC
115-
std::bool_constant<false>;
116-
#else
117-
std::bool_constant<(IsPlus<T, BinaryOperation>::value ||
115+
std::bool_constant<not IsDeterministicOperator<BinaryOperation>::value &&
116+
(IsPlus<T, BinaryOperation>::value ||
118117
IsMinimum<T, BinaryOperation>::value ||
119118
IsMaximum<T, BinaryOperation>::value) &&
120119
is_sgenfloat_v<T> && sizeof(T) == 8>;
121-
#endif
122120

123121
// This type trait is used to detect if the group algorithm reduce() used with
124122
// operands of the type T and the operation BinaryOperation is available
125123
// for using in reduction.
126-
// The macro SYCL_REDUCTION_DETERMINISTIC prohibits using the reduce() algorithm
127-
// to produce stable results across same type devices.
128124
template <typename T, class BinaryOperation>
129125
using IsReduOptForFastReduce =
130-
#ifdef SYCL_REDUCTION_DETERMINISTIC
131-
std::bool_constant<false>;
132-
#else
133-
std::bool_constant<((is_sgeninteger_v<T> &&
126+
std::bool_constant<not IsDeterministicOperator<BinaryOperation>::value &&
127+
((is_sgeninteger_v<T> &&
134128
(sizeof(T) == 4 || sizeof(T) == 8)) ||
135129
is_sgenfloat_v<T>) &&
136130
(IsPlus<T, BinaryOperation>::value ||
137131
IsMinimum<T, BinaryOperation>::value ||
138132
IsMaximum<T, BinaryOperation>::value)>;
139-
#endif
140133

141134
// std::tuple seems to be a) too heavy and b) not copyable to device now
142135
// Thus sycl::detail::tuple is used instead.

‎sycl/include/sycl/sycl.hpp‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -98,6 +98,7 @@
9898
#include <sycl/ext/oneapi/experimental/prefetch.hpp>
9999
#include <sycl/ext/oneapi/experimental/profiling_tag.hpp>
100100
#include <sycl/ext/oneapi/experimental/raw_kernel_arg.hpp>
101+
#include <sycl/ext/oneapi/experimental/reduction_properties.hpp>
101102
#include <sycl/ext/oneapi/experimental/root_group.hpp>
102103
#include <sycl/ext/oneapi/experimental/tangle_group.hpp>
103104
#include <sycl/ext/oneapi/filter_selector.hpp>

0 commit comments

Comments
 (0)