Make __parallel_copy_if blocking instead of returning a future

Author: akukanov

include/oneapi/dpl/pstl/hetero/algorithm_impl_hetero.h

@@ -918,22 +918,19 @@ _Iterator2
 __pattern_copy_if(__hetero_tag<_BackendTag>, _ExecutionPolicy&& __exec, _Iterator1 __first, _Iterator1 __last,
                   _Iterator2 __result_first, _Predicate __pred)
 {
-    using _It1DifferenceType = typename ::std::iterator_traits<_Iterator1>::difference_type;
-
     if (__first == __last)
         return __result_first;
 
-    _It1DifferenceType __n = __last - __first;
+    typename std::iterator_traits<_Iterator1>::difference_type __n = __last - __first;
 
     auto __keep1 = oneapi::dpl::__ranges::__get_sycl_range<__par_backend_hetero::access_mode::read, _Iterator1>();
     auto __buf1 = __keep1(__first, __last);
     auto __keep2 = oneapi::dpl::__ranges::__get_sycl_range<__par_backend_hetero::access_mode::write, _Iterator2>();
     auto __buf2 = __keep2(__result_first, __result_first + __n);
 
-    auto __res = __par_backend_hetero::__parallel_copy_if(_BackendTag{}, ::std::forward<_ExecutionPolicy>(__exec),
-                                                          __buf1.all_view(), __buf2.all_view(), __n, __n, __pred);
+    std::size_t __num_copied = __par_backend_hetero::__parallel_copy_if(
+        _BackendTag{}, std::forward<_ExecutionPolicy>(__exec), __buf1.all_view(), __buf2.all_view(), __n, __n, __pred);
 
-    ::std::size_t __num_copied = __res.get(); //is a blocking call
     return __result_first + __num_copied;
 }
 

include/oneapi/dpl/pstl/hetero/algorithm_ranges_impl_hetero.h

@@ -659,12 +659,10 @@ __pattern_copy_if(__hetero_tag<_BackendTag>, _ExecutionPolicy&& __exec, _Range1&
     if (__n == 0 || __m == 0)
         return 0;
 
-    auto __res = oneapi::dpl::__par_backend_hetero::__parallel_copy_if(
+    return oneapi::dpl::__par_backend_hetero::__parallel_copy_if(
         _BackendTag{}, std::forward<_ExecutionPolicy>(__exec),
         oneapi::dpl::__ranges::__get_subscription_view(std::forward<_Range1>(__rng1)),
         oneapi::dpl::__ranges::__get_subscription_view(std::forward<_Range2>(__rng2)), __n, __m, __pred, __assign);
-
-    return __res.get(); //is a blocking call
 }
 
 #if _ONEDPL_CPP20_RANGES_PRESENT

include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h

@@ -889,7 +889,7 @@ __parallel_partition_copy(oneapi::dpl::__internal::__device_backend_tag, _Execut
 
 template <typename _ExecutionPolicy, typename _InRng, typename _OutRng, typename _Size, typename _Pred,
           typename _Assign = oneapi::dpl::__internal::__pstl_assign>
-__future<sycl::event, __result_and_scratch_storage<_Size>>
+_Size
 __parallel_copy_if(oneapi::dpl::__internal::__device_backend_tag, _ExecutionPolicy&& __exec, _InRng&& __in_rng,
                    _OutRng&& __out_rng, _Size __n, _Size __m, _Pred __pred, _Assign __assign = _Assign{})
 // TODO: __parallel_copy_if should return two stop positions
@@ -919,7 +919,8 @@ __parallel_copy_if(oneapi::dpl::__internal::__device_backend_tag, _ExecutionPoli
             __scan_copy_single_wg_kernel<_CustomName>>;
         return __par_backend_hetero::__parallel_copy_if_static_single_group_submitter<_Size, _KernelName>()(
             __q_local, std::forward<_InRng>(__in_rng), std::forward<_OutRng>(__out_rng), __n, __m, __pred, __assign,
-            static_cast<std::uint16_t>(__n_uniform), static_cast<std::uint16_t>(std::min(__n_uniform, __max_wg_size)));
+            static_cast<std::uint16_t>(__n_uniform), static_cast<std::uint16_t>(std::min(__n_uniform, __max_wg_size)))
+            .get();
     }
     else if (oneapi::dpl::__par_backend_hetero::__is_gpu_with_reduce_then_scan_sg_sz(__q_local))
     {
@@ -929,7 +930,7 @@ __parallel_copy_if(oneapi::dpl::__internal::__device_backend_tag, _ExecutionPoli
         return __parallel_reduce_then_scan_copy<_CustomName>(__q_local, std::forward<_InRng>(__in_rng),
                                                              std::forward<_OutRng>(__out_rng), __n,
                                                              _GenMask{__pred, {}}, _WriteOp{__assign},
-                                                             /*_IsUniquePattern=*/std::false_type{});
+                                                             /*_IsUniquePattern=*/std::false_type{}).get();
     }
     else
     {
@@ -940,7 +941,7 @@ __parallel_copy_if(oneapi::dpl::__internal::__device_backend_tag, _ExecutionPoli
 
         return __parallel_scan_copy<_CustomName>(__q_local, std::forward<_InRng>(__in_rng),
                                                  std::forward<_OutRng>(__out_rng), __n, _CreateOp{__pred},
-                                                 _CopyOp{_ReduceOp{}, __assign});
+                                                 _CopyOp{_ReduceOp{}, __assign}).get();
     }
 }
 
@@ -2230,8 +2231,7 @@ __parallel_reduce_by_segment_fallback(oneapi::dpl::__internal::__device_backend_
         oneapi::dpl::__internal::__device_backend_tag{},
         oneapi::dpl::__par_backend_hetero::make_wrapped_policy<__assign_key1_wrapper>(__exec), __view1, __view2, __n,
         __n, __internal::__parallel_reduce_by_segment_fallback_fn1<_BinaryPredicate>{__binary_pred, __wgroup_size},
-        unseq_backend::__brick_assign_key_position{})
-        .get();
+        unseq_backend::__brick_assign_key_position{});
 
     //reduce by segment
     oneapi::dpl::__par_backend_hetero::__parallel_for(
@@ -2269,8 +2269,7 @@ __parallel_reduce_by_segment_fallback(oneapi::dpl::__internal::__device_backend_
         oneapi::dpl::__par_backend_hetero::make_wrapped_policy<__assign_key2_wrapper>(__exec), __view3, __view4,
         __view3.size(), __view3.size(),
         __internal::__parallel_reduce_by_segment_fallback_fn2<_BinaryPredicate>{__binary_pred},
-        unseq_backend::__brick_assign_key_position{})
-        .get();
+        unseq_backend::__brick_assign_key_position{});
 
     //reduce by segment
     oneapi::dpl::__par_backend_hetero::__parallel_for(