@@ -103,72 +103,82 @@ inline constexpr bool is_valid_reduction_prop_list =
103103 ext::oneapi::experimental::detail::reduction_property_check_anchor,
104104 PropertyList>;
105105
106+ template <typename BinaryOperation, typename PropertyList, typename ... Args>
107+ auto convert_reduction_properties (BinaryOperation combiner,
108+ PropertyList properties, Args &&...args) {
109+ if constexpr (is_valid_reduction_prop_list<PropertyList>) {
110+ auto WrappedOp = WrapOp (combiner, properties);
111+ auto RuntimeProps = GetReductionPropertyList (properties);
112+ return sycl::reduction (std::forward<Args>(args)..., WrappedOp,
113+ RuntimeProps);
114+ } else {
115+ // Invalid, will be disabled by SFINAE at the caller side. Make sure no hard
116+ // error is emitted from here.
117+ }
118+ }
106119} // namespace detail
107120
108- template <typename BufferT, typename BinaryOperation, typename PropertyList,
109- typename = std::enable_if_t <
110- detail::is_valid_reduction_prop_list<PropertyList>>>
121+ template <typename BufferT, typename BinaryOperation, typename PropertyList>
111122auto reduction (BufferT vars, handler &cgh, BinaryOperation combiner,
112- PropertyList properties) {
123+ PropertyList properties)
124+ -> std::enable_if_t<detail::is_valid_reduction_prop_list<PropertyList>,
125+ decltype(detail::convert_reduction_properties(
126+ combiner, properties, vars, cgh))> {
113127 detail::CheckReductionIdentity<typename BufferT::value_type, BinaryOperation>(
114128 properties);
115- auto WrappedOp = detail::WrapOp (combiner, properties);
116- auto RuntimeProps = detail::GetReductionPropertyList (properties);
117- return reduction (vars, cgh, WrappedOp, RuntimeProps);
129+ return detail::convert_reduction_properties (combiner, properties, vars, cgh);
118130}
119131
120- template <typename T, typename BinaryOperation, typename PropertyList,
121- typename = std::enable_if_t <
122- detail::is_valid_reduction_prop_list<PropertyList>>>
123- auto reduction (T *var, BinaryOperation combiner, PropertyList properties) {
132+ template <typename T, typename BinaryOperation, typename PropertyList>
133+ auto reduction (T *var, BinaryOperation combiner, PropertyList properties)
134+ -> std::enable_if_t<detail::is_valid_reduction_prop_list<PropertyList>,
135+ decltype(detail::convert_reduction_properties(
136+ combiner, properties, var))> {
124137 detail::CheckReductionIdentity<T, BinaryOperation>(properties);
125- auto WrappedOp = detail::WrapOp (combiner, properties);
126- auto RuntimeProps = detail::GetReductionPropertyList (properties);
127- return reduction (var, WrappedOp, RuntimeProps);
138+ return detail::convert_reduction_properties (combiner, properties, var);
128139}
129140
130- template <
131- typename T, size_t Extent, typename BinaryOperation, typename PropertyList,
132- typename =
133- std::enable_if_t <detail::is_valid_reduction_prop_list<PropertyList>>>
141+ template <typename T, size_t Extent, typename BinaryOperation,
142+ typename PropertyList>
134143auto reduction (span<T, Extent> vars, BinaryOperation combiner,
135- PropertyList properties) {
144+ PropertyList properties)
145+ -> std::enable_if_t<detail::is_valid_reduction_prop_list<PropertyList>,
146+ decltype(detail::convert_reduction_properties(
147+ combiner, properties, vars))> {
136148 detail::CheckReductionIdentity<T, BinaryOperation>(properties);
137- auto WrappedOp = detail::WrapOp (combiner, properties);
138- auto RuntimeProps = detail::GetReductionPropertyList (properties);
139- return reduction (vars, WrappedOp, RuntimeProps);
149+ return detail::convert_reduction_properties (combiner, properties, vars);
140150}
141151
142- template <typename BufferT, typename BinaryOperation, typename PropertyList,
143- typename = std::enable_if_t <
144- detail::is_valid_reduction_prop_list<PropertyList>>>
152+ template <typename BufferT, typename BinaryOperation, typename PropertyList>
145153auto reduction (BufferT vars, handler &cgh,
146154 const typename BufferT::value_type &identity,
147- BinaryOperation combiner, PropertyList properties) {
148- auto WrappedOp = detail::WrapOp (combiner, properties);
149- auto RuntimeProps = detail::GetReductionPropertyList (properties);
150- return reduction (vars, cgh, identity, WrappedOp, RuntimeProps);
155+ BinaryOperation combiner, PropertyList properties)
156+ -> std::enable_if_t<detail::is_valid_reduction_prop_list<PropertyList>,
157+ decltype(detail::convert_reduction_properties(
158+ combiner, properties, vars, cgh, identity))> {
159+ return detail::convert_reduction_properties (combiner, properties, vars, cgh,
160+ identity);
151161}
152162
153- template <typename T, typename BinaryOperation, typename PropertyList,
154- typename = std::enable_if_t <
155- detail::is_valid_reduction_prop_list<PropertyList>>>
163+ template <typename T, typename BinaryOperation, typename PropertyList>
156164auto reduction (T *var, const T &identity, BinaryOperation combiner,
157- PropertyList properties) {
158- auto WrappedOp = detail::WrapOp (combiner, properties);
159- auto RuntimeProps = detail::GetReductionPropertyList (properties);
160- return reduction (var, identity, WrappedOp, RuntimeProps);
165+ PropertyList properties)
166+ -> std::enable_if_t<detail::is_valid_reduction_prop_list<PropertyList>,
167+ decltype(detail::convert_reduction_properties(
168+ combiner, properties, var, identity))> {
169+ return detail::convert_reduction_properties (combiner, properties, var,
170+ identity);
161171}
162172
163- template <
164- typename T, size_t Extent, typename BinaryOperation, typename PropertyList,
165- typename =
166- std::enable_if_t <detail::is_valid_reduction_prop_list<PropertyList>>>
173+ template <typename T, size_t Extent, typename BinaryOperation,
174+ typename PropertyList>
167175auto reduction (span<T, Extent> vars, const T &identity,
168- BinaryOperation combiner, PropertyList properties) {
169- auto WrappedOp = detail::WrapOp (combiner, properties);
170- auto RuntimeProps = detail::GetReductionPropertyList (properties);
171- return reduction (vars, identity, WrappedOp, RuntimeProps);
176+ BinaryOperation combiner, PropertyList properties)
177+ -> std::enable_if_t<detail::is_valid_reduction_prop_list<PropertyList>,
178+ decltype(detail::convert_reduction_properties(
179+ combiner, properties, vars, identity))> {
180+ return detail::convert_reduction_properties (combiner, properties, vars,
181+ identity);
172182}
173183
174184} // namespace _V1
0 commit comments