diff --git a/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h b/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h index 5f6b77d2b2a..d226d647d97 100644 --- a/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h +++ b/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h @@ -396,7 +396,7 @@ __group_scan_fits_in_slm(const sycl::queue& __q, std::size_t __n, std::size_t __ template -__future> +auto /*__future*/ __parallel_transform_scan(oneapi::dpl::__internal::__device_backend_tag, _ExecutionPolicy&& __exec, _Range1&& __in_rng, _Range2&& __out_rng, std::size_t __n, _UnaryOperation __unary_op, _InitType __init, _BinaryOperation __binary_op, _Inclusive) @@ -405,6 +405,7 @@ __parallel_transform_scan(oneapi::dpl::__internal::__device_backend_tag, _Execut using _Type = typename _InitType::__value_type; sycl::queue __q_local = __exec.queue(); + __transform_scan_storage_holder_simple<_Type> __holder(__q_local); // The single work-group implementation requires a fundamental type which must be trivially copyable. if constexpr (std::is_trivially_copyable_v<_Type>) @@ -421,15 +422,11 @@ __parallel_transform_scan(oneapi::dpl::__internal::__device_backend_tag, _Execut const std::size_t __single_group_upper_limit = __q_local.get_device().is_gpu() ? 2048 : 16384; if (__group_scan_fits_in_slm<_Type>(__q_local, __n, __n_uniform, __single_group_upper_limit)) { - auto __event = __parallel_transform_scan_single_group<_CustomName>( + sycl::event __event = __parallel_transform_scan_single_group<_CustomName>( __q_local, std::forward<_Range1>(__in_rng), std::forward<_Range2>(__out_rng), __n, __unary_op, __init, __binary_op, _Inclusive{}); - // Although we do not actually need result storage in this case, we need to construct - // a placeholder here to match the return type of the non-single-work-group implementation - __result_and_scratch_storage<_Type> __dummy_result_and_scratch{__q_local, 0}; - - return __future{std::move(__event), std::move(__dummy_result_and_scratch)}; + return __future{std::move(__event), std::move(__holder).__extract()}; } } } @@ -443,17 +440,17 @@ __parallel_transform_scan(oneapi::dpl::__internal::__device_backend_tag, _Execut // the size of the scanned type produced by __unary_op, is what determines the input footprint of a block. constexpr std::uint32_t __bytes_per_work_item_iter = sizeof(oneapi::dpl::__internal::__value_t<_Range1>); - auto&& [__event, __payload] = __parallel_transform_reduce_then_scan< + sycl::event __event = __parallel_transform_reduce_then_scan< /*_Bounded*/ false, __bytes_per_work_item_iter, _CustomName>( __q_local, __n, std::forward<_Range1>(__in_rng), std::forward<_Range2>(__out_rng), __gen_transform, __binary_op, - __gen_transform, _ScanInputTransform{}, _WriteOp{}, __init, _Inclusive{}, + __gen_transform, _ScanInputTransform{}, _WriteOp{}, __init, __holder, _Inclusive{}, /*_IsUniquePattern=*/std::false_type{}); - return __create_future(std::move(__event), std::move(__payload)); + return __future(std::move(__event), std::move(__holder).__extract()); } template -__transform_reduce_then_scan_result_t<_Bounded, _Size, _Size> +std::array<_Size, 2> __parallel_reduce_then_scan_copy(sycl::queue& __q, _InRng&& __in_rng, _OutRng&& __out_rng, _Size __n, _GenMask __generate_mask, _WriteOp __write_op, _IsUniquePattern __is_unique_pattern) { @@ -463,14 +460,25 @@ __parallel_reduce_then_scan_copy(sycl::queue& __q, _InRng&& __in_rng, _OutRng&& using _GenScanInput = oneapi::dpl::__par_backend_hetero::__gen_expand_count_mask<_GenMask, _Size>; using _ScanInputTransform = oneapi::dpl::__par_backend_hetero::__get_zeroth_element; + std::array<_Size, 2> __ret{}; + __transform_scan_storage_holder<_Bounded, _Size, _Size> __holder(__q); + // Each work-item iteration reads a single input element to evaluate the mask and to copy it to the output. constexpr std::uint32_t __bytes_per_work_item_iter = sizeof(oneapi::dpl::__internal::__value_t<_InRng>); - return __parallel_transform_reduce_then_scan<_Bounded, __bytes_per_work_item_iter, _CustomName>( + sycl::event __event = __parallel_transform_reduce_then_scan<_Bounded, __bytes_per_work_item_iter, _CustomName>( __q, __n, std::forward<_InRng>(__in_rng), std::forward<_OutRng>(__out_rng), _GenReduceInput{__generate_mask}, _ReduceOp{}, _GenScanInput{__generate_mask}, _ScanInputTransform{}, __write_op, - oneapi::dpl::unseq_backend::__no_init_value<_Size>{}, /*_Inclusive=*/std::true_type{}, __is_unique_pattern, - /*__stop_pos_initial_state=*/__n); + oneapi::dpl::unseq_backend::__no_init_value<_Size>{}, __holder, /*_Inclusive=*/std::true_type{}, + __is_unique_pattern, /*__stop_pos_initial_state=*/__n); + __event.wait_and_throw(); + + __holder.template __copy_result<0>(__ret.data(), 1); + if constexpr (_Bounded) + __holder.template __copy_result<1>(__ret.data() + 1, 1); + else + __ret[1] = __n; + return __ret; } template ; using _WriteOp = oneapi::dpl::__par_backend_hetero::__write_to_id_if<1, _Assign>; - std::tuple __res = __parallel_reduce_then_scan_copy<_Bounded, _CustomName>( + __ret = __parallel_reduce_then_scan_copy<_Bounded, _CustomName>( __q_local, std::forward<_Range1>(__rng), std::forward<_Range2>(__result), __n, _GenMask{__pred}, _WriteOp{std::size_t(__n_out)}, /*_IsUniquePattern=*/std::true_type{}); - - std::get<0>(__res).wait_and_throw(); - std::get<1>(__res).__copy_result(__ret.data(), 1); - if constexpr (_Bounded) - __ret[1] = __load_result(std::get<2>(__res)); - else - __ret[1] = __n; } assert(__ret[0] >= 1 && __n_out >= __ret[0]); @@ -523,8 +524,7 @@ __parallel_unique_copy(oneapi::dpl::__internal::__device_backend_tag, _Execution template -__future>>> +auto /*__future*/ __parallel_reduce_by_segment_reduce_then_scan(sycl::queue& __q, _Range1&& __keys, _Range2&& __values, _Range3&& __out_keys, _Range4&& __out_values, _BinaryPredicate __binary_pred, _BinaryOperator __binary_op) @@ -541,23 +541,24 @@ __parallel_reduce_by_segment_reduce_then_scan(sycl::queue& __q, _Range1&& __keys using _WriteOp = __write_red_by_seg<_BinaryPredicate>; using _KeyType = oneapi::dpl::__internal::__value_t<_Range1>; using _ValueType = oneapi::dpl::__internal::__value_t<_Range2>; + using _ResultType = oneapi::dpl::__internal::tuple; + std::size_t __n = oneapi::dpl::__ranges::__size(__keys); // __gen_red_by_seg_scan_input requires that __n > 1 assert(__n > 1); - + __transform_scan_storage_holder_simple<_ResultType> __holder(__q); // Each work-item iteration reads one key and one value from the zipped input. The comparison against the previous // key is not counted separately, as that key is read by the adjacent index's iteration. constexpr std::uint32_t __bytes_per_work_item_iter = sizeof(_KeyType) + sizeof(_ValueType); - auto&& [__event, __payload] = __parallel_transform_reduce_then_scan< + sycl::event __event = __parallel_transform_reduce_then_scan< /*_Bounded*/ false, __bytes_per_work_item_iter, _CustomName>( __q, __n, oneapi::dpl::__ranges::make_zip_view(std::forward<_Range1>(__keys), std::forward<_Range2>(__values)), oneapi::dpl::__ranges::make_zip_view(std::forward<_Range3>(__out_keys), std::forward<_Range4>(__out_values)), _GenReduceInput{__binary_pred}, _ReduceOp{__binary_op}, _GenScanInput{__binary_pred, __n}, - _ScanInputTransform{}, _WriteOp{__binary_pred, __n}, - oneapi::dpl::unseq_backend::__no_init_value>{}, - /*Inclusive*/ std::true_type{}, /*_IsUniquePattern=*/std::false_type{}); - return __create_future(std::move(__event), std::move(__payload)); + _ScanInputTransform{}, _WriteOp{__binary_pred, __n}, oneapi::dpl::unseq_backend::__no_init_value<_ResultType>{}, + __holder, /*Inclusive*/ std::true_type{}, /*_IsUniquePattern=*/std::false_type{}); + return __future(std::move(__event), std::move(__holder).__extract()); } template (__out_true), std::forward<_Range3>(__out_false)); sycl::queue __q_local = __exec.queue(); + __transform_scan_storage_holder<_Bounded, diff_t, _WriteOp::__position_type> __holder(__q_local); // Each work-item iteration reads a single input element to evaluate the mask and to copy it to the output. constexpr std::uint32_t __bytes_per_work_item_iter = sizeof(oneapi::dpl::__internal::__value_t<_Range1>); - std::tuple __res = __parallel_transform_reduce_then_scan<_Bounded, __bytes_per_work_item_iter, _CustomName>( + sycl::event __event = __parallel_transform_reduce_then_scan<_Bounded, __bytes_per_work_item_iter, _CustomName>( __q_local, __n, std::forward<_Range1>(__rng), std::move(__zipped_output), _GenReduceInput{_GenMask{__pred}}, std::plus{}, _GenScanInput{_GenMask{__pred}}, _ScanInputTransform{}, _WriteOp{__n_out1, __n_out2}, - oneapi::dpl::unseq_backend::__no_init_value{}, /*_Inclusive=*/std::true_type{}, - /*_IsUniquePattern=*/std::false_type{}, __write_partitioned::__position_type{__n, __n_out1}); + oneapi::dpl::unseq_backend::__no_init_value{}, __holder, /*_Inclusive=*/std::true_type{}, + /*_IsUniquePattern=*/std::false_type{}, _WriteOp::__position_type{__n, __n_out1}); + __event.wait_and_throw(); std::array __ret{}; - std::get<0>(__res).wait_and_throw(); - std::get<1>(__res).__copy_result(__ret.data(), 1); + std::tuple __results = __get_results(__holder); + __ret[0] = std::get<0>(__results); if constexpr (_Bounded) { - auto [__stop_in, __stop_out1] = __load_result(std::get<2>(__res)); + auto [__stop_in, __stop_out1] = std::get<1>(__results); __ret[1] = __stop_in; if (__stop_out1 < __ret[0]) __ret[0] = __stop_out1; @@ -640,16 +643,9 @@ __parallel_copy_if(oneapi::dpl::__internal::__device_backend_tag, _ExecutionPoli using _GenMask = oneapi::dpl::__par_backend_hetero::__gen_mask<_Pred>; using _WriteOp = oneapi::dpl::__par_backend_hetero::__write_to_id_if<0, _Assign>; - std::tuple __res = __parallel_reduce_then_scan_copy<_Bounded, _CustomName>( + __ret = __parallel_reduce_then_scan_copy<_Bounded, _CustomName>( __q_local, std::forward<_InRng>(__in_rng), std::forward<_OutRng>(__out_rng), __n, _GenMask{__pred}, _WriteOp{std::size_t(__n_out), __assign}, /*_IsUniquePattern=*/std::false_type{}); - - std::get<0>(__res).wait_and_throw(); - std::get<1>(__res).__copy_result(__ret.data(), 1); - if constexpr (_Bounded) - __ret[1] = __load_result(std::get<2>(__res)); - else - __ret[1] = __n; } assert(__ret[0] >= 0 && __n_out >= __ret[0]); @@ -662,8 +658,7 @@ __parallel_copy_if(oneapi::dpl::__internal::__device_backend_tag, _ExecutionPoli // balanced path template -__transform_reduce_then_scan_result_t<_Bounded, oneapi::dpl::__internal::__difference_t<_Range3>, - _SetOpFinalAndOOBPosType<_Range1, _Range2>> +auto /*std::tuple<...>*/ __parallel_set_write_a_b_op(_SetTag __set_tag, sycl::queue& __q, _Range1&& __rng1, _Range2&& __rng2, _Range3&& __result, _Compare __comp, _Proj1 __proj1, _Proj2 __proj2) { @@ -731,14 +726,18 @@ __parallel_set_write_a_b_op(_SetTag __set_tag, sycl::queue& __q, _Range1&& __rng } // Initial stop pos state - const auto __stop_pos_initial_state = __create_initial_final_and_oob_pos_state<_Bounded>(__set_tag, __rng1, __rng2); + auto __stop_pos_initial_state = __create_initial_final_and_oob_pos_state<_Bounded>(__set_tag, __rng1, __rng2); + __transform_scan_storage_holder<_Bounded, _Size, decltype(__stop_pos_initial_state)> __holder(__q); - return __parallel_transform_reduce_then_scan<_Bounded, __bytes_per_work_item_iter, _CustomName>( + __partition_event = __parallel_transform_reduce_then_scan<_Bounded, __bytes_per_work_item_iter, _CustomName>( __q, __num_diagonals, std::move(__in_in_tmp_rng), std::forward<_Range3>(__result), __gen_reduce_input, _ReduceOp{}, _GenScanInput{_SetOperation{}, __diagonal_spacing, __comp, __proj1, __proj2}, - _ScanInputTransform{}, _WriteOp{__n_out}, oneapi::dpl::unseq_backend::__no_init_value<_Size>{}, + _ScanInputTransform{}, _WriteOp{__n_out}, oneapi::dpl::unseq_backend::__no_init_value<_Size>{}, __holder, /*_Inclusive=*/std::true_type{}, /*__is_unique_pattern=*/std::false_type{}, __stop_pos_initial_state, __partition_event); + __partition_event.wait_and_throw(); + + return __get_results(__holder); } template (__rng1), std::forward<_Range2>(__rng2), std::forward<_Range3>(__result), __comp, __proj1, __proj2); - std::get<0>(__res).wait_and_throw(); - // Load stop position in the output range - const oneapi::dpl::__internal::__difference_t<_Range3> __stop_pos3 = __load_result(std::get<1>(__res)); + const oneapi::dpl::__internal::__difference_t<_Range3> __stop_pos3 = std::get<0>(__res); // Load stop positions in the input ranges oneapi::dpl::__internal::__difference_t<_Range1> __stop_pos1 = {}; oneapi::dpl::__internal::__difference_t<_Range2> __stop_pos2 = {}; if constexpr (_Bounded) - std::tie(__stop_pos1, __stop_pos2) = __load_result(std::get<2>(__res)).__compute_stop_pos(); + std::tie(__stop_pos1, __stop_pos2) = std::get<1>(__res).__compute_stop_pos(); return __create_set_op_impl_result<_Bounded, _Range1, _Range2, _Range3>(__stop_pos1, __stop_pos2, __stop_pos3); } @@ -1681,8 +1678,8 @@ __parallel_reduce_by_segment(oneapi::dpl::__internal::__device_backend_tag, _Exe auto __res = oneapi::dpl::__par_backend_hetero::__parallel_reduce_by_segment_reduce_then_scan<_CustomName>( __q_local, std::forward<_Range1>(__keys), std::forward<_Range2>(__values), std::forward<_Range3>(__out_keys), std::forward<_Range4>(__out_values), __binary_pred, __binary_op); - // Because our init type ends up being tuple, return the first component which is the write index. Add 1 to return the - // past-the-end iterator pair of segmented reduction. + // Because our init type ends up being tuple, return the first component which is + // the write index. Add 1 to return the past-the-end iterator pair of segmented reduction. return std::get<0>(__res.get()) + 1; #else using __val_type = oneapi::dpl::__internal::__value_t<_Range2>; @@ -1700,8 +1697,7 @@ __parallel_reduce_by_segment(oneapi::dpl::__internal::__device_backend_tag, _Exe //------------------------------------------------------------------------ template -__future>>> +auto /*__future*/ __parallel_scan_by_segment_reduce_then_scan(sycl::queue& __q, _Range1&& __keys, _Range2&& __values, _Range3&& __out_values, _BinaryPredicate __binary_pred, _BinaryOperator __binary_op, [[maybe_unused]] _InitType __init) @@ -1726,14 +1722,15 @@ __parallel_scan_by_segment_reduce_then_scan(sycl::queue& __q, _Range1&& __keys, // Each work-item iteration reads one key and one value from the zipped input. constexpr std::uint32_t __bytes_per_work_item_iter = sizeof(_KeyType) + sizeof(_ValueType); + __transform_scan_storage_holder_simple<_PackedFlagValueType> __holder(__q); - auto&& [__event, __payload] = __parallel_transform_reduce_then_scan< + sycl::event __event = __parallel_transform_reduce_then_scan< /*_Bounded*/ false, __bytes_per_work_item_iter, _CustomName>( __q, __n, oneapi::dpl::__ranges::make_zip_view(std::forward<_Range1>(__keys), std::forward<_Range2>(__values)), std::forward<_Range3>(__out_values), _GenReduceInput{__binary_pred}, _ReduceOp{__binary_op}, _GenScanInput{}, - _ScanInputTransform{}, _WriteOp{__init, __binary_op}, __placeholder_no_init, + _ScanInputTransform{}, _WriteOp{__init, __binary_op}, __placeholder_no_init, __holder, /*Inclusive*/ std::bool_constant<__is_inclusive>{}, /*_IsUniquePattern=*/std::false_type{}); - return __create_future(std::move(__event), std::move(__payload)); + return __future(std::move(__event), std::move(__holder).__extract()); } template -using __transform_reduce_then_scan_result_t = - std::conditional_t<_Bounded, - std::tuple, __result_storage<_StopPosType>>, - std::tuple>>; - template @@ -2279,6 +2273,12 @@ struct __parallel_reduce_then_scan_scan_submitter<_Bounded, __is_inclusive, __is const bool __use_subgroup_ops; }; +template +using __transform_scan_storage_holder = std::conditional_t<_Bounded, __storage_holder<1, _ValueType, _StopPosType>, + __storage_holder<1, _ValueType>>; +template +using __transform_scan_storage_holder_simple = __storage_holder<1, _ValueType>; + // Helper for __parallel_transform_reduce_then_scan templated on the choice of sub-group communication // strategy via _ScanOpsTag, which selects which communication path(s) are compiled into the kernel. The // runtime __use_subgroup_ops flag then chooses between them when both are available. @@ -2286,13 +2286,14 @@ template -__transform_reduce_then_scan_result_t<_Bounded, typename _InitType::__value_type, _StopPosInitState> -__parallel_transform_reduce_then_scan_impl(sycl::queue& __q, const std::size_t __n, _InRng&& __in_rng, - _OutRng&& __out_rng, _GenReduceInput __gen_reduce_input, - _ReduceOp __reduce_op, _GenScanInput __gen_scan_input, - _ScanInputTransform __scan_input_transform, _WriteOp __write_op, - _InitType __init, _Inclusive, _IsUniquePattern, bool __use_subgroup_ops, - _StopPosInitState __stop_pos_initial_state, sycl::event __prior_event) +sycl::event +__parallel_transform_reduce_then_scan_impl( + sycl::queue& __q, const std::size_t __n, _InRng&& __in_rng, _OutRng&& __out_rng, _GenReduceInput __gen_reduce_input, + _ReduceOp __reduce_op, _GenScanInput __gen_scan_input, _ScanInputTransform __scan_input_transform, + _WriteOp __write_op, _InitType __init, + __transform_scan_storage_holder<_Bounded, typename _InitType::__value_type, _StopPosInitState>& __holder, + _Inclusive, _IsUniquePattern, bool __use_subgroup_ops, _StopPosInitState __stop_pos_initial_state, + sycl::event __prior_event) { using _ReduceKernel = oneapi::dpl::__par_backend_hetero::__internal::__kernel_name_provider< __reduce_then_scan_reduce_kernel<_ScanOpsTag, _CustomName>>; @@ -2475,10 +2476,10 @@ __parallel_transform_reduce_then_scan_impl(sycl::queue& __q, const std::size_t _ } } + __holder.template __store<0>(std::move(__result_and_scratch)); if constexpr (_Bounded) - return {std::move(__prior_event), std::move(__result_and_scratch), std::move(__stop_pos_storage)}; - else - return {std::move(__prior_event), std::move(__result_and_scratch)}; + __holder.template __store<1>(std::move(__stop_pos_storage)); + return __prior_event; } // General scan-like algorithm helpers @@ -2499,13 +2500,14 @@ template > -__transform_reduce_then_scan_result_t<_Bounded, typename _InitType::__value_type, _StopPosInitState> -__parallel_transform_reduce_then_scan(sycl::queue& __q, const std::size_t __n, _InRng&& __in_rng, _OutRng&& __out_rng, - _GenReduceInput __gen_reduce_input, _ReduceOp __reduce_op, - _GenScanInput __gen_scan_input, _ScanInputTransform __scan_input_transform, - _WriteOp __write_op, _InitType __init, _Inclusive __inclusive, - _IsUniquePattern __is_unique_pattern, - _StopPosInitState __stop_pos_initial_state = {}, sycl::event __prior_event = {}) +sycl::event +__parallel_transform_reduce_then_scan( + sycl::queue& __q, const std::size_t __n, _InRng&& __in_rng, _OutRng&& __out_rng, _GenReduceInput __gen_reduce_input, + _ReduceOp __reduce_op, _GenScanInput __gen_scan_input, _ScanInputTransform __scan_input_transform, + _WriteOp __write_op, _InitType __init, + __transform_scan_storage_holder<_Bounded, typename _InitType::__value_type, _StopPosInitState>& __holder, + _Inclusive __inclusive, _IsUniquePattern __is_unique_pattern, _StopPosInitState __stop_pos_initial_state = {}, + sycl::event __prior_event = {}) { using _ValueType = typename _InitType::__value_type; @@ -2523,7 +2525,7 @@ __parallel_transform_reduce_then_scan(sycl::queue& __q, const std::size_t __n, _ return __parallel_transform_reduce_then_scan_impl<_Bounded, __slm_or_subgroup_tag<_ValueType>, __bytes_per_work_item_iter, _CustomName>( __q, __n, std::forward<_InRng>(__in_rng), std::forward<_OutRng>(__out_rng), __gen_reduce_input, __reduce_op, - __gen_scan_input, __scan_input_transform, __write_op, __init, __inclusive, __is_unique_pattern, + __gen_scan_input, __scan_input_transform, __write_op, __init, __holder, __inclusive, __is_unique_pattern, __use_subgroup_ops, __stop_pos_initial_state, std::move(__prior_event)); } else @@ -2531,7 +2533,7 @@ __parallel_transform_reduce_then_scan(sycl::queue& __q, const std::size_t __n, _ return __parallel_transform_reduce_then_scan_impl<_Bounded, __slm_only_tag<_ValueType>, __bytes_per_work_item_iter, _CustomName>( __q, __n, std::forward<_InRng>(__in_rng), std::forward<_OutRng>(__out_rng), __gen_reduce_input, __reduce_op, - __gen_scan_input, __scan_input_transform, __write_op, __init, __inclusive, __is_unique_pattern, + __gen_scan_input, __scan_input_transform, __write_op, __init, __holder, __inclusive, __is_unique_pattern, /*__use_subgroup_ops=*/false, __stop_pos_initial_state, std::move(__prior_event)); } } diff --git a/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_utils.h b/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_utils.h index 5cf20ca7154..4f238d8623e 100644 --- a/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_utils.h +++ b/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_utils.h @@ -627,6 +627,14 @@ class __future : private std::tuple<_Args...> return {__buf[0], __buf[1]}; } + template + _T + __wait_and_get_value(const __internal::__copyable_storage_state<_T>& __st) + { + __internal::__copyable_storage_state<_T> __st_copy = __st; + return __wait_and_get_value(__result_and_scratch_storage(std::move(__st_copy))); + } + template _T __wait_and_get_value(const _T& __val) @@ -698,13 +706,6 @@ class __future : private std::tuple<_Args...> } }; -template -auto -__create_future(sycl::event&& __event, __combined_storage<_ValueType>&& __payload) -{ - return __future(std::move(__event), __result_and_scratch_storage<_ValueType>(std::move(__payload).__move_state())); -} - struct __scalar_load_op { oneapi::dpl::__internal::__pstl_assign __assigner; diff --git a/include/oneapi/dpl/pstl/hetero/dpcpp/utils_storage_sycl.h b/include/oneapi/dpl/pstl/hetero/dpcpp/utils_storage_sycl.h index 54ee375552a..9d463492831 100644 --- a/include/oneapi/dpl/pstl/hetero/dpcpp/utils_storage_sycl.h +++ b/include/oneapi/dpl/pstl/hetero/dpcpp/utils_storage_sycl.h @@ -190,7 +190,6 @@ struct __no_result_needed_tag // The type to exchange information between storage types. // Useful for the interoperability during the transition period -// TODO: afterwards, remove together with __combined_storage::__move_state template struct __copyable_storage_state { @@ -225,7 +224,10 @@ __move_state(__result_raw_state<_T>&& __raw_st, const sycl::queue& __q) __raw_st.__result_sz, __raw_st.__offset, __raw_st.__kind}; if (__raw_st.__usm_ptr) { - __state.__result_buf = std::shared_ptr<_T>(__raw_st.__usm_ptr, __sycl_usm_free{__q}); + if (__raw_st.__kind == sycl::usm::alloc::host) + __state.__result_buf = std::shared_ptr<_T>(__raw_st.__usm_ptr, __sycl_usm_free{__q}); + else + __state.__scratch_buf = std::shared_ptr<_T>(__raw_st.__usm_ptr, __sycl_usm_free{__q}); __raw_st.__usm_ptr = nullptr; } __raw_st.__sycl_buf.reset(); @@ -523,24 +525,8 @@ struct __combined_storage : public __device_storage<_T> __cgh, __st.__sycl_buf, __st.__usm_buf.get(), /*offset*/ __st.__sz, __st.__result_sz, __prop_list); } } - - __internal::__copyable_storage_state<_T> - __move_state() && - { - return {std::move(__result_buf), std::move(this->__usm_buf), std::move(this->__sycl_buf), - __result_sz, __sz, __kind}; - } }; -template typename _Storage> -std::enable_if_t, _T> -__load_result(_Storage<_T>& __storage) -{ - _T __result{}; - __storage.__copy_result(&__result, 1); - return __result; -} - template class __storage_holder { @@ -550,6 +536,9 @@ class __storage_holder std::array<__internal::__scratch_raw_state, _NScratch> __scratch_slots = {}; std::size_t __scratch_count = 0; + template + using __result_value_t = std::tuple_element_t<_I, std::tuple<_ResultTypes...>>; + template auto __extract_impl(std::index_sequence<_ResultIs...>, std::index_sequence<_ScratchIs...>) && @@ -593,7 +582,7 @@ class __storage_holder { for (auto& __slot : __scratch_slots) __internal::__free_usm(__q, __slot.__usm_ptr); - std::apply([this](auto&... __slot) { + std::apply([&__q = this->__q](auto&... __slot) { ((__internal::__free_usm(__q, __slot.__usm_ptr)), ...); }, __result_slots); } @@ -611,7 +600,7 @@ class __storage_holder __store(__result_storage<_T>&& __st) { static_assert(_I < sizeof...(_ResultTypes), "Result slot index out of range"); - static_assert(std::is_same_v<_T, std::tuple_element_t<_I, std::tuple<_ResultTypes...>>>); + static_assert(std::is_same_v<_T, __result_value_t<_I>>); auto& __slot = std::get<_I>(__result_slots); assert(__slot.__usm_ptr == nullptr && !__slot.__sycl_buf.has_value()); std::move(__st).__move_state_to(__slot); @@ -622,7 +611,7 @@ class __storage_holder __store(__combined_storage<_T>&& __st) { static_assert(_I < sizeof...(_ResultTypes), "Result index out of range"); - static_assert(std::is_same_v<_T, std::tuple_element_t<_I, std::tuple<_ResultTypes...>>>); + static_assert(std::is_same_v<_T, __result_value_t<_I>>); auto& __slot = std::get<_I>(__result_slots); assert(__slot.__usm_ptr == nullptr && !__slot.__sycl_buf.has_value()); void* __scratch_ptr = std::move(__st).__move_state_to(__slot); @@ -635,7 +624,7 @@ class __storage_holder template void - __copy_result(std::tuple_element_t<_I, std::tuple<_ResultTypes...>>* __dst, std::size_t __n) + __copy_result(__result_value_t<_I>* __dst, std::size_t __n) { __internal::__copy_n(__dst, __n, std::get<_I>(__result_slots), __q); } @@ -646,8 +635,26 @@ class __storage_holder return std::move(*this).__extract_impl(std::index_sequence_for<_ResultTypes...>{}, std::make_index_sequence<_NScratch>{}); } + + template + friend std::enable_if_t<(std::is_default_constructible_v<_Types> && ...), std::tuple<_Types...>> + __get_results(__storage_holder<_N, _Types...>&); }; +template +std::enable_if_t<(std::is_default_constructible_v<_Types> && ...), std::tuple<_Types...>> +__get_results(__storage_holder<_NScratch, _Types...>& __h) +{ + return std::apply([&__q = __h.__q](auto&... __slots) { + auto __load_one = [&](auto& __rs, auto* __p) { + std::remove_pointer_t __dst{}; + __internal::__copy_n(&__dst, 1, __rs, __q); + return __dst; + }; + return std::tuple<_Types...>{__load_one(__slots, static_cast<_Types*>(nullptr))...}; + }, __h.__result_slots); +} + } // namespace __par_backend_hetero } // namespace oneapi::dpl