@@ -626,7 +626,7 @@ class reduction_impl_algo : public reduction_impl_common<T, BinaryOperation> {
626
626
Func (Mem);
627
627
628
628
reduction::withAuxHandler (CGH, [&](handler &CopyHandler) {
629
- // MSVC (19.32.31329) has problems compiling the line below when used
629
+ // MSVC (19.32.31329) has problems compiling the line below when used
630
630
// as a host compiler in c++17 mode (but not in c++latest)
631
631
// accessor Mem{*Buf, CopyHandler};
632
632
// so use the old-style API.
@@ -751,67 +751,26 @@ class reduction_impl
751
751
752
752
// / Constructs reduction_impl when the identity value is statically known.
753
753
template <typename _self = self,
754
- enable_if_t <_self::is_known_identity && !_self::is_usm> * = nullptr >
755
- reduction_impl (RedOutVar &Acc)
756
- : algo(reducer_type::getIdentity(), BinaryOperation(), false , Acc) {
757
- if (Acc.size () != 1 )
758
- throw sycl::runtime_error (errc::invalid,
759
- " Reduction variable must be a scalar." ,
760
- PI_ERROR_INVALID_VALUE);
761
- }
762
-
763
- // / Constructs reduction_impl when the identity value is statically known.
764
- // / The \param VarPtr is a USM pointer to memory, to where the computed
765
- // / reduction value is added using BinaryOperation, i.e. it is expected that
766
- // / the memory is pre-initialized with some meaningful value.
767
- template <typename _self = self,
768
- enable_if_t <_self::is_known_identity && _self::is_usm> * = nullptr >
769
- reduction_impl (RedOutVar VarPtr, bool InitializeToIdentity = false )
754
+ enable_if_t <_self::is_known_identity> * = nullptr >
755
+ reduction_impl (RedOutVar Var, bool InitializeToIdentity = false )
770
756
: algo(reducer_type::getIdentity(), BinaryOperation(),
771
- InitializeToIdentity, VarPtr) {}
772
-
773
- // / SYCL-2020.
774
- // / Constructs reduction_impl when the identity value is statically known.
775
- template <typename _self = self, std::enable_if_t <_self::is_known_identity &&
776
- !_self::is_usm> * = nullptr >
777
- reduction_impl (RedOutVar &Acc, handler &CGH, bool InitializeToIdentity)
778
- : algo(reducer_type::getIdentity(), BinaryOperation(),
779
- InitializeToIdentity, Acc) {
780
- associateWithHandler (CGH, &Acc, access::target::device);
781
- if (Acc.size () != 1 )
782
- throw sycl::runtime_error (errc::invalid,
783
- " Reduction variable must be a scalar." ,
784
- PI_ERROR_INVALID_VALUE);
757
+ InitializeToIdentity, Var) {
758
+ if constexpr (!is_usm)
759
+ if (Var.size () != 1 )
760
+ throw sycl::runtime_error (errc::invalid,
761
+ " Reduction variable must be a scalar." ,
762
+ PI_ERROR_INVALID_VALUE);
785
763
}
786
764
787
765
// / Constructs reduction_impl when the identity value is unknown.
788
- template <typename _self = self, enable_if_t <!_self::is_usm> * = nullptr >
789
- reduction_impl (RedOutVar &Acc, const T &Identity, BinaryOperation BOp)
790
- : algo(chooseIdentity(Identity), BOp, false , Acc) {
791
- if (Acc.size () != 1 )
792
- throw sycl::runtime_error (errc::invalid,
793
- " Reduction variable must be a scalar." ,
794
- PI_ERROR_INVALID_VALUE);
795
- }
796
-
797
- // / The \param VarPtr is a USM pointer to memory, to where the computed
798
- // / reduction value is added using BinaryOperation, i.e. it is expected that
799
- // / the memory is pre-initialized with some meaningful value.
800
- template <typename _self = self, enable_if_t <_self::is_usm> * = nullptr >
801
- reduction_impl (RedOutVar VarPtr, const T &Identity, BinaryOperation BOp,
802
- bool InitializeToIdentity = false )
803
- : algo(chooseIdentity(Identity), BOp, InitializeToIdentity, VarPtr) {}
804
-
805
- // / For placeholder accessor.
806
- template <typename _self = self, enable_if_t <!_self::is_usm> * = nullptr >
807
- reduction_impl (RedOutVar &Acc, handler &CGH, const T &Identity,
808
- BinaryOperation BOp, bool InitializeToIdentity)
809
- : algo(chooseIdentity(Identity), BOp, InitializeToIdentity, Acc) {
810
- associateWithHandler (CGH, &Acc, access::target::device);
811
- if (Acc.size () != 1 )
812
- throw sycl::runtime_error (errc::invalid,
813
- " Reduction variable must be a scalar." ,
814
- PI_ERROR_INVALID_VALUE);
766
+ reduction_impl (RedOutVar &Var, const T &Identity, BinaryOperation BOp,
767
+ bool InitializeToIdentity)
768
+ : algo(chooseIdentity(Identity), BOp, InitializeToIdentity, Var) {
769
+ if constexpr (!is_usm)
770
+ if (Var.size () != 1 )
771
+ throw sycl::runtime_error (errc::invalid,
772
+ " Reduction variable must be a scalar." ,
773
+ PI_ERROR_INVALID_VALUE);
815
774
}
816
775
};
817
776
@@ -2356,7 +2315,7 @@ auto reduction(buffer<T, 1, AllocatorT> Var, handler &CGH, BinaryOperation,
2356
2315
const property_list &PropList = {}) {
2357
2316
bool InitializeToIdentity =
2358
2317
PropList.has_property <property::reduction::initialize_to_identity>();
2359
- return detail::make_reduction<BinaryOperation, 0 , 1 >(accessor{Var, CGH}, CGH,
2318
+ return detail::make_reduction<BinaryOperation, 0 , 1 >(accessor{Var, CGH},
2360
2319
InitializeToIdentity);
2361
2320
}
2362
2321
@@ -2420,7 +2379,7 @@ auto reduction(buffer<T, 1, AllocatorT> Var, handler &CGH, const T &Identity,
2420
2379
bool InitializeToIdentity =
2421
2380
PropList.has_property <property::reduction::initialize_to_identity>();
2422
2381
return detail::make_reduction<BinaryOperation, 0 , 1 >(
2423
- accessor{Var, CGH}, CGH, Identity, Combiner, InitializeToIdentity);
2382
+ accessor{Var, CGH}, Identity, Combiner, InitializeToIdentity);
2424
2383
}
2425
2384
2426
2385
// / Constructs a reduction object using the reduction variable referenced by
0 commit comments