Skip to content

Commit 11b3094

Browse files
aelovikov-intelaelovikov-intel
authored
[SYCL][Reduction] Use core group_algorithm.hpp instead of an extension (#6419)
1 parent 4baeb81 commit 11b3094

File tree

1 file changed

+76
-79
lines changed

1 file changed

+76
-79
lines changed

‎sycl/include/sycl/ext/oneapi/reduction.hpp

Lines changed: 76 additions & 79 deletions
Original file line numberDiff line numberDiff line change
@@ -13,12 +13,13 @@
1313

1414
#include <CL/sycl/accessor.hpp>
1515
#include <CL/sycl/atomic.hpp>
16+
#include <CL/sycl/atomic_ref.hpp>
1617
#include <CL/sycl/detail/tuple.hpp>
18+
#include <CL/sycl/group_algorithm.hpp>
1719
#include <CL/sycl/handler.hpp>
1820
#include <CL/sycl/kernel.hpp>
1921
#include <CL/sycl/known_identity.hpp>
2022
#include <sycl/ext/oneapi/accessor_property_list.hpp>
21-
#include <sycl/ext/oneapi/group_algorithm.hpp>
2223

2324
#include <tuple>
2425

@@ -222,8 +223,8 @@ template <class Reducer> class combiner {
222223
auto reducer = static_cast<const Reducer *>(this);
223224
for (size_t E = 0; E < Extent; ++E) {
224225
auto AtomicRef =
225-
atomic_ref<T, memory_order::relaxed, getMemoryScope<Space>(), Space>(
226-
multi_ptr<T, Space>(ReduVarPtr)[E]);
226+
sycl::atomic_ref<T, memory_order::relaxed, getMemoryScope<Space>(),
227+
Space>(multi_ptr<T, Space>(ReduVarPtr)[E]);
227228
Functor(AtomicRef, reducer->getElement(E));
228229
}
229230
}
@@ -312,13 +313,15 @@ template <class Reducer> class combiner {
312313
/// reduction.
313314
template <typename T, class BinaryOperation, int Dims, size_t Extent,
314315
class Algorithm, bool View>
315-
class reducer<T, BinaryOperation, Dims, Extent, Algorithm, View,
316-
enable_if_t<Dims == 0 && Extent == 1 && View == false &&
317-
!IsKnownIdentityOp<T, BinaryOperation>::value>>
316+
class reducer<
317+
T, BinaryOperation, Dims, Extent, Algorithm, View,
318+
enable_if_t<Dims == 0 && Extent == 1 && View == false &&
319+
!sycl::detail::IsKnownIdentityOp<T, BinaryOperation>::value>>
318320
: public combiner<
319321
reducer<T, BinaryOperation, Dims, Extent, Algorithm, View,
320322
enable_if_t<Dims == 0 && Extent == 1 && View == false &&
321-
!IsKnownIdentityOp<T, BinaryOperation>::value>>> {
323+
!sycl::detail::IsKnownIdentityOp<
324+
T, BinaryOperation>::value>>> {
322325
public:
323326
reducer(const T &Identity, BinaryOperation BOp)
324327
: MValue(Identity), MIdentity(Identity), MBinaryOp(BOp) {}
@@ -343,13 +346,15 @@ class reducer<T, BinaryOperation, Dims, Extent, Algorithm, View,
343346
/// the identity field inside it and allows to add a default constructor.
344347
template <typename T, class BinaryOperation, int Dims, size_t Extent,
345348
class Algorithm, bool View>
346-
class reducer<T, BinaryOperation, Dims, Extent, Algorithm, View,
347-
enable_if_t<Dims == 0 && Extent == 1 && View == false &&
348-
IsKnownIdentityOp<T, BinaryOperation>::value>>
349+
class reducer<
350+
T, BinaryOperation, Dims, Extent, Algorithm, View,
351+
enable_if_t<Dims == 0 && Extent == 1 && View == false &&
352+
sycl::detail::IsKnownIdentityOp<T, BinaryOperation>::value>>
349353
: public combiner<
350354
reducer<T, BinaryOperation, Dims, Extent, Algorithm, View,
351355
enable_if_t<Dims == 0 && Extent == 1 && View == false &&
352-
IsKnownIdentityOp<T, BinaryOperation>::value>>> {
356+
sycl::detail::IsKnownIdentityOp<
357+
T, BinaryOperation>::value>>> {
353358
public:
354359
reducer() : MValue(getIdentity()) {}
355360
reducer(const T & /* Identity */, BinaryOperation) : MValue(getIdentity()) {}
@@ -360,7 +365,7 @@ class reducer<T, BinaryOperation, Dims, Extent, Algorithm, View,
360365
}
361366

362367
static T getIdentity() {
363-
return known_identity_impl<BinaryOperation, T>::value;
368+
return sycl::detail::known_identity_impl<BinaryOperation, T>::value;
364369
}
365370

366371
T &getElement(size_t) { return MValue; }
@@ -390,13 +395,14 @@ class reducer<T, BinaryOperation, Dims, Extent, Algorithm, View,
390395
/// subscript operator.
391396
template <typename T, class BinaryOperation, int Dims, size_t Extent,
392397
class Algorithm, bool View>
393-
class reducer<T, BinaryOperation, Dims, Extent, Algorithm, View,
394-
enable_if_t<Dims == 1 && View == false &&
395-
!IsKnownIdentityOp<T, BinaryOperation>::value>>
396-
: public combiner<
397-
reducer<T, BinaryOperation, Dims, Extent, Algorithm, View,
398-
enable_if_t<Dims == 1 && View == false &&
399-
!IsKnownIdentityOp<T, BinaryOperation>::value>>> {
398+
class reducer<
399+
T, BinaryOperation, Dims, Extent, Algorithm, View,
400+
enable_if_t<Dims == 1 && View == false &&
401+
!sycl::detail::IsKnownIdentityOp<T, BinaryOperation>::value>>
402+
: public combiner<reducer<T, BinaryOperation, Dims, Extent, Algorithm, View,
403+
enable_if_t<Dims == 1 && View == false &&
404+
!sycl::detail::IsKnownIdentityOp<
405+
T, BinaryOperation>::value>>> {
400406
public:
401407
reducer(const T &Identity, BinaryOperation BOp)
402408
: MValue(Identity), MIdentity(Identity), MBinaryOp(BOp) {}
@@ -420,13 +426,14 @@ class reducer<T, BinaryOperation, Dims, Extent, Algorithm, View,
420426
/// in cases where the identity value is known.
421427
template <typename T, class BinaryOperation, int Dims, size_t Extent,
422428
class Algorithm, bool View>
423-
class reducer<T, BinaryOperation, Dims, Extent, Algorithm, View,
424-
enable_if_t<Dims == 1 && View == false &&
425-
IsKnownIdentityOp<T, BinaryOperation>::value>>
426-
: public combiner<
427-
reducer<T, BinaryOperation, Dims, Extent, Algorithm, View,
428-
enable_if_t<Dims == 1 && View == false &&
429-
IsKnownIdentityOp<T, BinaryOperation>::value>>> {
429+
class reducer<
430+
T, BinaryOperation, Dims, Extent, Algorithm, View,
431+
enable_if_t<Dims == 1 && View == false &&
432+
sycl::detail::IsKnownIdentityOp<T, BinaryOperation>::value>>
433+
: public combiner<reducer<T, BinaryOperation, Dims, Extent, Algorithm, View,
434+
enable_if_t<Dims == 1 && View == false &&
435+
sycl::detail::IsKnownIdentityOp<
436+
T, BinaryOperation>::value>>> {
430437
public:
431438
reducer() : MValue(getIdentity()) {}
432439
reducer(const T & /* Identity */, BinaryOperation) : MValue(getIdentity()) {}
@@ -439,7 +446,7 @@ class reducer<T, BinaryOperation, Dims, Extent, Algorithm, View,
439446
}
440447

441448
static T getIdentity() {
442-
return known_identity_impl<BinaryOperation, T>::value;
449+
return sycl::detail::known_identity_impl<BinaryOperation, T>::value;
443450
}
444451

445452
T &getElement(size_t E) { return MValue[E]; }
@@ -464,14 +471,14 @@ template <typename T, class BinaryOperation> class reduction_impl_common {
464471
public:
465472
/// Returns the statically known identity value.
466473
template <typename _T = T, class _BinaryOperation = BinaryOperation>
467-
enable_if_t<IsKnownIdentityOp<_T, _BinaryOperation>::value,
474+
enable_if_t<sycl::detail::IsKnownIdentityOp<_T, _BinaryOperation>::value,
468475
_T> constexpr getIdentity() {
469-
return known_identity_impl<_BinaryOperation, _T>::value;
476+
return sycl::detail::known_identity_impl<_BinaryOperation, _T>::value;
470477
}
471478

472479
/// Returns the identity value given by user.
473480
template <typename _T = T, class _BinaryOperation = BinaryOperation>
474-
enable_if_t<!IsKnownIdentityOp<_T, _BinaryOperation>::value, _T>
481+
enable_if_t<!sycl::detail::IsKnownIdentityOp<_T, _BinaryOperation>::value, _T>
475482
getIdentity() {
476483
return MIdentity;
477484
}
@@ -752,8 +759,8 @@ class reduction_impl
752759
/// SYCL-2020.
753760
/// Constructs reduction_impl when the identity value is statically known.
754761
template <typename _T, typename AllocatorT,
755-
std::enable_if_t<IsKnownIdentityOp<_T, BinaryOperation>::value> * =
756-
nullptr>
762+
std::enable_if_t<sycl::detail::IsKnownIdentityOp<
763+
_T, BinaryOperation>::value> * = nullptr>
757764
reduction_impl(buffer<_T, 1, AllocatorT> Buffer, handler &CGH,
758765
bool InitializeToIdentity)
759766
: algo(reducer_type::getIdentity(), BinaryOperation(),
@@ -766,9 +773,8 @@ class reduction_impl
766773
}
767774

768775
/// Constructs reduction_impl when the identity value is statically known.
769-
template <
770-
typename _T = T,
771-
enable_if_t<IsKnownIdentityOp<_T, BinaryOperation>::value> * = nullptr>
776+
template <typename _T = T, enable_if_t<sycl::detail::IsKnownIdentityOp<
777+
_T, BinaryOperation>::value> * = nullptr>
772778
reduction_impl(rw_accessor_type &Acc)
773779
: algo(reducer_type::getIdentity(), BinaryOperation(), false,
774780
std::make_shared<rw_accessor_type>(Acc)) {
@@ -779,9 +785,8 @@ class reduction_impl
779785
}
780786

781787
/// Constructs reduction_impl when the identity value is statically known.
782-
template <
783-
typename _T = T,
784-
enable_if_t<IsKnownIdentityOp<_T, BinaryOperation>::value> * = nullptr>
788+
template <typename _T = T, enable_if_t<sycl::detail::IsKnownIdentityOp<
789+
_T, BinaryOperation>::value> * = nullptr>
785790
reduction_impl(dw_accessor_type &Acc)
786791
: algo(reducer_type::getIdentity(), BinaryOperation(), true,
787792
std::make_shared<dw_accessor_type>(Acc)) {
@@ -796,7 +801,8 @@ class reduction_impl
796801
/// and user still passed the identity value.
797802
template <
798803
typename _T, typename AllocatorT,
799-
enable_if_t<IsKnownIdentityOp<_T, BinaryOperation>::value> * = nullptr>
804+
enable_if_t<sycl::detail::IsKnownIdentityOp<_T, BinaryOperation>::value>
805+
* = nullptr>
800806
reduction_impl(buffer<_T, 1, AllocatorT> Buffer, handler &CGH,
801807
const T & /*Identity*/, BinaryOperation,
802808
bool InitializeToIdentity)
@@ -822,9 +828,8 @@ class reduction_impl
822828

823829
/// Constructs reduction_impl when the identity value is statically known,
824830
/// and user still passed the identity value.
825-
template <
826-
typename _T = T,
827-
enable_if_t<IsKnownIdentityOp<_T, BinaryOperation>::value> * = nullptr>
831+
template <typename _T = T, enable_if_t<sycl::detail::IsKnownIdentityOp<
832+
_T, BinaryOperation>::value> * = nullptr>
828833
reduction_impl(rw_accessor_type &Acc, const T & /*Identity*/, BinaryOperation)
829834
: algo(reducer_type::getIdentity(), BinaryOperation(), false,
830835
std::make_shared<rw_accessor_type>(Acc)) {
@@ -847,9 +852,8 @@ class reduction_impl
847852

848853
/// Constructs reduction_impl when the identity value is statically known,
849854
/// and user still passed the identity value.
850-
template <
851-
typename _T = T,
852-
enable_if_t<IsKnownIdentityOp<_T, BinaryOperation>::value> * = nullptr>
855+
template <typename _T = T, enable_if_t<sycl::detail::IsKnownIdentityOp<
856+
_T, BinaryOperation>::value> * = nullptr>
853857
reduction_impl(dw_accessor_type &Acc, const T & /*Identity*/, BinaryOperation)
854858
: algo(reducer_type::getIdentity(), BinaryOperation(), true,
855859
std::make_shared<dw_accessor_type>(Acc)) {
@@ -874,7 +878,8 @@ class reduction_impl
874878
/// Constructs reduction_impl when the identity value is NOT known statically.
875879
template <
876880
typename _T, typename AllocatorT,
877-
enable_if_t<!IsKnownIdentityOp<_T, BinaryOperation>::value> * = nullptr>
881+
enable_if_t<!sycl::detail::IsKnownIdentityOp<_T, BinaryOperation>::value>
882+
* = nullptr>
878883
reduction_impl(buffer<_T, 1, AllocatorT> Buffer, handler &CGH,
879884
const T &Identity, BinaryOperation BOp,
880885
bool InitializeToIdentity)
@@ -888,9 +893,8 @@ class reduction_impl
888893
}
889894

890895
/// Constructs reduction_impl when the identity value is unknown.
891-
template <
892-
typename _T = T,
893-
enable_if_t<!IsKnownIdentityOp<_T, BinaryOperation>::value> * = nullptr>
896+
template <typename _T = T, enable_if_t<!sycl::detail::IsKnownIdentityOp<
897+
_T, BinaryOperation>::value> * = nullptr>
894898
reduction_impl(rw_accessor_type &Acc, const T &Identity, BinaryOperation BOp)
895899
: algo(Identity, BOp, false, std::make_shared<rw_accessor_type>(Acc)) {
896900
if (Acc.size() != 1)
@@ -900,9 +904,8 @@ class reduction_impl
900904
}
901905

902906
/// Constructs reduction_impl when the identity value is unknown.
903-
template <
904-
typename _T = T,
905-
enable_if_t<!IsKnownIdentityOp<_T, BinaryOperation>::value> * = nullptr>
907+
template <typename _T = T, enable_if_t<!sycl::detail::IsKnownIdentityOp<
908+
_T, BinaryOperation>::value> * = nullptr>
906909
reduction_impl(dw_accessor_type &Acc, const T &Identity, BinaryOperation BOp)
907910
: algo(Identity, BOp, true, std::make_shared<dw_accessor_type>(Acc)) {
908911
if (Acc.size() != 1)
@@ -915,9 +918,8 @@ class reduction_impl
915918
/// The \param VarPtr is a USM pointer to memory, to where the computed
916919
/// reduction value is added using BinaryOperation, i.e. it is expected that
917920
/// the memory is pre-initialized with some meaningful value.
918-
template <
919-
typename _T = T,
920-
enable_if_t<IsKnownIdentityOp<_T, BinaryOperation>::value> * = nullptr>
921+
template <typename _T = T, enable_if_t<sycl::detail::IsKnownIdentityOp<
922+
_T, BinaryOperation>::value> * = nullptr>
921923
reduction_impl(T *VarPtr, bool InitializeToIdentity = false)
922924
: algo(reducer_type::getIdentity(), BinaryOperation(),
923925
InitializeToIdentity, VarPtr) {}
@@ -927,9 +929,8 @@ class reduction_impl
927929
/// The \param VarPtr is a USM pointer to memory, to where the computed
928930
/// reduction value is added using BinaryOperation, i.e. it is expected that
929931
/// the memory is pre-initialized with some meaningful value.
930-
template <
931-
typename _T = T,
932-
enable_if_t<IsKnownIdentityOp<_T, BinaryOperation>::value> * = nullptr>
932+
template <typename _T = T, enable_if_t<sycl::detail::IsKnownIdentityOp<
933+
_T, BinaryOperation>::value> * = nullptr>
933934
reduction_impl(T *VarPtr, const T &Identity, BinaryOperation,
934935
bool InitializeToIdentity = false)
935936
: algo(Identity, BinaryOperation(), InitializeToIdentity, VarPtr) {
@@ -950,35 +951,31 @@ class reduction_impl
950951
/// The \param VarPtr is a USM pointer to memory, to where the computed
951952
/// reduction value is added using BinaryOperation, i.e. it is expected that
952953
/// the memory is pre-initialized with some meaningful value.
953-
template <
954-
typename _T = T,
955-
enable_if_t<!IsKnownIdentityOp<_T, BinaryOperation>::value> * = nullptr>
954+
template <typename _T = T, enable_if_t<!sycl::detail::IsKnownIdentityOp<
955+
_T, BinaryOperation>::value> * = nullptr>
956956
reduction_impl(T *VarPtr, const T &Identity, BinaryOperation BOp,
957957
bool InitializeToIdentity = false)
958958
: algo(Identity, BOp, InitializeToIdentity, VarPtr) {}
959959

960960
/// Constructs reduction_impl when the identity value is statically known
961-
template <
962-
typename _T = T,
963-
enable_if_t<IsKnownIdentityOp<_T, BinaryOperation>::value> * = nullptr>
961+
template <typename _T = T, enable_if_t<sycl::detail::IsKnownIdentityOp<
962+
_T, BinaryOperation>::value> * = nullptr>
964963
reduction_impl(span<_T, Extent> Span, bool InitializeToIdentity = false)
965964
: algo(reducer_type::getIdentity(), BinaryOperation(),
966965
InitializeToIdentity, Span.data()) {}
967966

968967
/// Constructs reduction_impl when the identity value is statically known
969968
/// and user passed an identity value anyway
970-
template <
971-
typename _T = T,
972-
enable_if_t<IsKnownIdentityOp<_T, BinaryOperation>::value> * = nullptr>
969+
template <typename _T = T, enable_if_t<sycl::detail::IsKnownIdentityOp<
970+
_T, BinaryOperation>::value> * = nullptr>
973971
reduction_impl(span<_T, Extent> Span, const T & /* Identity */,
974972
BinaryOperation BOp, bool InitializeToIdentity = false)
975973
: algo(reducer_type::getIdentity(), BOp, InitializeToIdentity,
976974
Span.data()) {}
977975

978976
/// Constructs reduction_impl when the identity value is not statically known
979-
template <
980-
typename _T = T,
981-
enable_if_t<!IsKnownIdentityOp<_T, BinaryOperation>::value> * = nullptr>
977+
template <typename _T = T, enable_if_t<!sycl::detail::IsKnownIdentityOp<
978+
_T, BinaryOperation>::value> * = nullptr>
982979
reduction_impl(span<T, Extent> Span, const T &Identity, BinaryOperation BOp,
983980
bool InitializeToIdentity = false)
984981
: algo(Identity, BOp, InitializeToIdentity, Span.data()) {}
@@ -1105,8 +1102,8 @@ void reduCGFuncForRangeFastReduce(handler &CGH, KernelType KernelFunc,
11051102
// Signal this work-group has finished after all values are reduced
11061103
if (LID == 0) {
11071104
auto NFinished =
1108-
atomic_ref<int, memory_order::relaxed, memory_scope::device,
1109-
access::address_space::global_space>(
1105+
sycl::atomic_ref<int, memory_order::relaxed, memory_scope::device,
1106+
access::address_space::global_space>(
11101107
NWorkGroupsFinished[0]);
11111108
DoReducePartialSumsInLastWG[0] =
11121109
++NFinished == NWorkGroups && NWorkGroups > 1;
@@ -1204,8 +1201,8 @@ void reduCGFuncForRangeBasic(handler &CGH, KernelType KernelFunc,
12041201
// Signal this work-group has finished after all values are reduced
12051202
if (LID == 0) {
12061203
auto NFinished =
1207-
atomic_ref<int, memory_order::relaxed, memory_scope::device,
1208-
access::address_space::global_space>(
1204+
sycl::atomic_ref<int, memory_order::relaxed, memory_scope::device,
1205+
access::address_space::global_space>(
12091206
NWorkGroupsFinished[0]);
12101207
DoReducePartialSumsInLastWG[0] =
12111208
++NFinished == NWorkGroups && NWorkGroups > 1;
@@ -1302,7 +1299,7 @@ void reduCGFuncForNDRangeBothFastReduceAndAtomics(
13021299
typename Reduction::binary_operation BOp;
13031300
for (int E = 0; E < NElements; ++E) {
13041301
Reducer.getElement(E) =
1305-
ext::oneapi::reduce(NDIt.get_group(), Reducer.getElement(E), BOp);
1302+
reduce_over_group(NDIt.get_group(), Reducer.getElement(E), BOp);
13061303
}
13071304
if (NDIt.get_local_linear_id() == 0)
13081305
Reducer.atomic_combine(Reduction::getOutPointer(Out));
@@ -1426,7 +1423,7 @@ void reduCGFuncForNDRangeFastReduceOnly(
14261423
for (int E = 0; E < NElements; ++E) {
14271424
typename Reduction::result_type PSum;
14281425
PSum = Reducer.getElement(E);
1429-
PSum = ext::oneapi::reduce(NDIt.get_group(), PSum, BOp);
1426+
PSum = reduce_over_group(NDIt.get_group(), PSum, BOp);
14301427
if (NDIt.get_local_linear_id() == 0) {
14311428
if (IsUpdateOfUserVar)
14321429
PSum = BOp(Reduction::getOutPointer(Out)[E], PSum);
@@ -1556,7 +1553,7 @@ void reduAuxCGFuncFastReduceImpl(handler &CGH, bool UniformWG,
15561553
(UniformWG || (GID < NWorkItems))
15571554
? In[GID * NElements + E]
15581555
: Reduction::reducer_type::getIdentity();
1559-
PSum = ext::oneapi::reduce(NDIt.get_group(), PSum, BOp);
1556+
PSum = reduce_over_group(NDIt.get_group(), PSum, BOp);
15601557
if (NDIt.get_local_linear_id() == 0) {
15611558
if (IsUpdateOfUserVar)
15621559
PSum = BOp(Reduction::getOutPointer(Out)[E], PSum);
@@ -2587,7 +2584,7 @@ reduction(accessor<T, Dims, AccMode, access::target::device, IsPH> &Acc,
25872584
/// The identity value is not passed to this version as it is statically known.
25882585
template <typename T, class BinaryOperation, int Dims, access::mode AccMode,
25892586
access::placeholder IsPH>
2590-
std::enable_if_t<detail::IsKnownIdentityOp<T, BinaryOperation>::value,
2587+
std::enable_if_t<sycl::detail::IsKnownIdentityOp<T, BinaryOperation>::value,
25912588
detail::reduction_impl<
25922589
T, BinaryOperation, 0, 1,
25932590
detail::default_reduction_algorithm<false, IsPH, Dims>>>
@@ -2615,7 +2612,7 @@ reduction(T *VarPtr, const T &Identity, BinaryOperation BOp) {
26152612
/// The identity value is not passed to this version as it is statically known.
26162613
template <typename T, class BinaryOperation>
26172614
std::enable_if_t<
2618-
detail::IsKnownIdentityOp<T, BinaryOperation>::value,
2615+
sycl::detail::IsKnownIdentityOp<T, BinaryOperation>::value,
26192616
detail::reduction_impl<T, BinaryOperation, 0, 1,
26202617
detail::default_reduction_algorithm<
26212618
true, access::placeholder::false_t, 1>>>

0 commit comments

Comments
 (0)