Skip to content

Commit

Permalink
include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_reduce.h -…
Browse files Browse the repository at this point in the history
… special fix for __parallel_transform_reduce_work_group_kernel_submitter::operator()

Signed-off-by: Sergey Kopienko <[email protected]>
  • Loading branch information
SergeyKopienko committed Sep 25, 2024
1 parent 384d0d4 commit 1b91850
Showing 1 changed file with 5 additions and 4 deletions.
Original file line number Diff line number Diff line change
Expand Up @@ -237,9 +237,9 @@ struct __parallel_transform_reduce_work_group_kernel_submitter<_Tp, _Commutative
template <typename _ExecutionPolicy, typename _Size, typename _ReduceOp, typename _InitType,
typename _ExecutionPolicy2>
__future<sycl::event, __result_and_scratch_storage<_ExecutionPolicy2, _Tp>>
operator()(oneapi::dpl::__internal::__device_backend_tag, _ExecutionPolicy&& __exec, sycl::event& __reduce_event,
operator()(oneapi::dpl::__internal::__device_backend_tag, _ExecutionPolicy&& __exec, sycl::event&& __reduce_event,
const _Size __n, const _Size __work_group_size, const _Size __iters_per_work_item, _ReduceOp __reduce_op,
_InitType __init, const __result_and_scratch_storage<_ExecutionPolicy2, _Tp>& __scratch_container) const
_InitType __init, __result_and_scratch_storage<_ExecutionPolicy2, _Tp>& __scratch_container) const
{
using _NoOpFunctor = unseq_backend::walk_n<_ExecutionPolicy, oneapi::dpl::__internal::__no_op>;
auto __transform_pattern =
Expand Down Expand Up @@ -269,7 +269,8 @@ struct __parallel_transform_reduce_work_group_kernel_submitter<_Tp, _Commutative
});
});

return __future(__reduce_event, __scratch_container);
return __future<sycl::event, __result_and_scratch_storage<_ExecutionPolicy2, _Tp>>{std::move(__reduce_event),
std::move(__scratch_container)};
}
}; // struct __parallel_transform_reduce_work_group_kernel_submitter

Expand Down Expand Up @@ -301,7 +302,7 @@ __parallel_transform_reduce_mid_impl(oneapi::dpl::__internal::__device_backend_t
// __n_groups preliminary results from the device kernel.
return __parallel_transform_reduce_work_group_kernel_submitter<_Tp, _Commutative, _VecSize,
_ReduceWorkGroupKernel>()(
__backend_tag, std::forward<_ExecutionPolicy>(__exec), __reduce_event, __n_groups, __work_group_size,
__backend_tag, std::forward<_ExecutionPolicy>(__exec), std::move(__reduce_event), __n_groups, __work_group_size,
__iters_per_work_item_work_group_kernel, __reduce_op, __init, __scratch_container);
}

Expand Down

0 comments on commit 1b91850

Please sign in to comment.