Skip to content

Commit 111ea2f

Browse files
PennycookPennycook
committed
[SYCL][NFC] Refactor to avoid code duplication
This commit refactors the reduction implementation to avoid code duplication, and additionally makes a number of changes to the class hierarchy for future extensibility. All functionality and member variables expected to be common to all reductions is moved to a common base class (reduction_impl_common). The existing reduction_impl_base is unsuitable for this purpose because it was deliberately designed not to be a template class. The reduction_impl is now templated on a reduction algorithm, with any functionality related to the current (default) algorithm encapsulated in the default_reduction_algorithm class. This template is carried from the reduction to any reducers it creates, enabling future specialization of both reduction and reducer for interesting combinations of type, extent and properties. The reducer class is simplified using CRTP to avoid duplicate definitions of combine() and atomic_combine() for scalar and array reductions. The notion of a reducer's "dimensionality" is now tied to the dimensionality of the reduction being performed (i.e. 0 for scalars, 1 for spans) and not to the dimensionality of the input accessor/buffer. This will simplify extending reductions to true multi-dimensional array reductions (i.e. with md_span). Signed-off-by: John Pennycook <[email protected]>
1 parent 807440c commit 111ea2f

File tree

3 files changed

+572
-827
lines changed

3 files changed

+572
-827
lines changed

‎sycl/include/CL/sycl/handler.hpp

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -240,9 +240,9 @@ class RoundedRangeKernelWithKH {
240240
namespace ext {
241241
namespace oneapi {
242242
namespace detail {
243-
template <typename T, class BinaryOperation, int Dims, bool IsUSM,
244-
access::placeholder IsPlaceholder>
245-
class reduction_impl;
243+
template <typename T, class BinaryOperation, int Dims, size_t Extent,
244+
class Algorithm>
245+
class reduction_impl_algo;
246246

247247
using cl::sycl::detail::enable_if_t;
248248
using cl::sycl::detail::queue_impl;
@@ -2685,11 +2685,11 @@ class __SYCL_EXPORT handler {
26852685
// Make stream class friend to be able to keep the list of associated streams
26862686
friend class stream;
26872687
friend class detail::stream_impl;
2688-
// Make reduction_impl friend to store buffers and arrays created for it
2689-
// in handler from reduction_impl methods.
2690-
template <typename T, class BinaryOperation, int Dims, bool IsUSM,
2691-
access::placeholder IsPlaceholder>
2692-
friend class ext::oneapi::detail::reduction_impl;
2688+
// Make reduction friends to store buffers and arrays created for it
2689+
// in handler from reduction methods.
2690+
template <typename T, class BinaryOperation, int Dims, size_t Extent,
2691+
class Algorithm>
2692+
friend class ext::oneapi::detail::reduction_impl_algo;
26932693

26942694
// This method needs to call the method finalize().
26952695
template <typename Reduction, typename... RestT>

‎sycl/include/CL/sycl/reduction.hpp

Lines changed: 36 additions & 17 deletions
Original file line numberDiff line numberDiff line change
@@ -20,7 +20,9 @@ namespace sycl {
2020
template <typename T, typename AllocatorT, typename BinaryOperation>
2121
std::enable_if_t<has_known_identity<BinaryOperation, T>::value,
2222
ext::oneapi::detail::reduction_impl<
23-
T, BinaryOperation, 1, false, access::placeholder::true_t>>
23+
T, BinaryOperation, 0, 1,
24+
ext::oneapi::detail::default_reduction_algorithm<
25+
false, access::placeholder::true_t, 1>>>
2426
reduction(buffer<T, 1, AllocatorT> Var, handler &CGH, BinaryOperation,
2527
const property_list &PropList = {}) {
2628
bool InitializeToIdentity =
@@ -35,7 +37,9 @@ reduction(buffer<T, 1, AllocatorT> Var, handler &CGH, BinaryOperation,
3537
template <typename T, typename AllocatorT, typename BinaryOperation>
3638
std::enable_if_t<!has_known_identity<BinaryOperation, T>::value,
3739
ext::oneapi::detail::reduction_impl<
38-
T, BinaryOperation, 1, false, access::placeholder::true_t>>
40+
T, BinaryOperation, 0, 1,
41+
ext::oneapi::detail::default_reduction_algorithm<
42+
false, access::placeholder::true_t, 1>>>
3943
reduction(buffer<T, 1, AllocatorT>, handler &, BinaryOperation,
4044
const property_list &PropList = {}) {
4145
// TODO: implement reduction that works even when identity is not known.
@@ -49,9 +53,11 @@ reduction(buffer<T, 1, AllocatorT>, handler &, BinaryOperation,
4953
/// the given USM pointer \p Var, handler \p CGH, reduction operation
5054
/// \p Combiner, and optional reduction properties.
5155
template <typename T, typename BinaryOperation>
52-
std::enable_if_t<
53-
has_known_identity<BinaryOperation, T>::value,
54-
ext::oneapi::detail::reduction_impl<T, BinaryOperation, 1, true>>
56+
std::enable_if_t<has_known_identity<BinaryOperation, T>::value,
57+
ext::oneapi::detail::reduction_impl<
58+
T, BinaryOperation, 0, 1,
59+
ext::oneapi::detail::default_reduction_algorithm<
60+
true, access::placeholder::false_t, 1>>>
5561
reduction(T *Var, BinaryOperation, const property_list &PropList = {}) {
5662
bool InitializeToIdentity =
5763
PropList.has_property<property::reduction::initialize_to_identity>();
@@ -64,9 +70,11 @@ reduction(T *Var, BinaryOperation, const property_list &PropList = {}) {
6470
/// The reduction algorithm may be less efficient for this variant as the
6571
/// reduction identity is not known statically and it is not provided by user.
6672
template <typename T, typename BinaryOperation>
67-
std::enable_if_t<
68-
!has_known_identity<BinaryOperation, T>::value,
69-
ext::oneapi::detail::reduction_impl<T, BinaryOperation, 1, true>>
73+
std::enable_if_t<!has_known_identity<BinaryOperation, T>::value,
74+
ext::oneapi::detail::reduction_impl<
75+
T, BinaryOperation, 0, 1,
76+
ext::oneapi::detail::default_reduction_algorithm<
77+
true, access::placeholder::false_t, 1>>>
7078
reduction(T *, BinaryOperation, const property_list &PropList = {}) {
7179
// TODO: implement reduction that works even when identity is not known.
7280
(void)PropList;
@@ -79,8 +87,10 @@ reduction(T *, BinaryOperation, const property_list &PropList = {}) {
7987
/// reduction identity value \p Identity, reduction operation \p Combiner,
8088
/// and optional reduction properties.
8189
template <typename T, typename AllocatorT, typename BinaryOperation>
82-
ext::oneapi::detail::reduction_impl<T, BinaryOperation, 1, false,
83-
access::placeholder::true_t>
90+
ext::oneapi::detail::reduction_impl<
91+
T, BinaryOperation, 0, 1,
92+
ext::oneapi::detail::default_reduction_algorithm<
93+
false, access::placeholder::true_t, 1>>
8494
reduction(buffer<T, 1, AllocatorT> Var, handler &CGH, const T &Identity,
8595
BinaryOperation Combiner, const property_list &PropList = {}) {
8696
bool InitializeToIdentity =
@@ -92,7 +102,10 @@ reduction(buffer<T, 1, AllocatorT> Var, handler &CGH, const T &Identity,
92102
/// the given USM pointer \p Var, reduction identity value \p Identity,
93103
/// binary operation \p Combiner, and optional reduction properties.
94104
template <typename T, typename BinaryOperation>
95-
ext::oneapi::detail::reduction_impl<T, BinaryOperation, 1, true>
105+
ext::oneapi::detail::reduction_impl<
106+
T, BinaryOperation, 0, 1,
107+
ext::oneapi::detail::default_reduction_algorithm<
108+
true, access::placeholder::false_t, 1>>
96109
reduction(T *Var, const T &Identity, BinaryOperation Combiner,
97110
const property_list &PropList = {}) {
98111
bool InitializeToIdentity =
@@ -109,8 +122,10 @@ reduction(T *Var, const T &Identity, BinaryOperation Combiner,
109122
template <typename T, size_t Extent, typename BinaryOperation>
110123
std::enable_if_t<Extent != dynamic_extent &&
111124
has_known_identity<BinaryOperation, T>::value,
112-
ext::oneapi::detail::reduction_impl<span<T, Extent>,
113-
BinaryOperation, 1, true>>
125+
ext::oneapi::detail::reduction_impl<
126+
T, BinaryOperation, 1, Extent,
127+
ext::oneapi::detail::default_reduction_algorithm<
128+
true, access::placeholder::false_t, 1>>>
114129
reduction(span<T, Extent> Span, BinaryOperation,
115130
const property_list &PropList = {}) {
116131
bool InitializeToIdentity =
@@ -126,8 +141,10 @@ reduction(span<T, Extent> Span, BinaryOperation,
126141
template <typename T, size_t Extent, typename BinaryOperation>
127142
std::enable_if_t<Extent != dynamic_extent &&
128143
!has_known_identity<BinaryOperation, T>::value,
129-
ext::oneapi::detail::reduction_impl<span<T, Extent>,
130-
BinaryOperation, 1, true>>
144+
ext::oneapi::detail::reduction_impl<
145+
T, BinaryOperation, 1, Extent,
146+
ext::oneapi::detail::default_reduction_algorithm<
147+
true, access::placeholder::false_t, 1>>>
131148
reduction(span<T, Extent> Span, BinaryOperation,
132149
const property_list &PropList = {}) {
133150
// TODO: implement reduction that works even when identity is not known.
@@ -142,8 +159,10 @@ reduction(span<T, Extent> Span, BinaryOperation,
142159
/// reduction operation \p Combiner, and optional reduction properties.
143160
template <typename T, size_t Extent, typename BinaryOperation>
144161
std::enable_if_t<Extent != dynamic_extent,
145-
ext::oneapi::detail::reduction_impl<span<T, Extent>,
146-
BinaryOperation, 1, true>>
162+
ext::oneapi::detail::reduction_impl<
163+
T, BinaryOperation, 1, Extent,
164+
ext::oneapi::detail::default_reduction_algorithm<
165+
true, access::placeholder::false_t, 1>>>
147166
reduction(span<T, Extent> Span, const T &Identity, BinaryOperation Combiner,
148167
const property_list &PropList = {}) {
149168
bool InitializeToIdentity =

0 commit comments

Comments
 (0)