diff --git a/include/oneapi/dpl/internal/async_impl/async_impl_hetero.h b/include/oneapi/dpl/internal/async_impl/async_impl_hetero.h index 37fc5ccb7b8..da8f4679ed1 100644 --- a/include/oneapi/dpl/internal/async_impl/async_impl_hetero.h +++ b/include/oneapi/dpl/internal/async_impl/async_impl_hetero.h @@ -172,7 +172,7 @@ __pattern_fill_async(__hetero_tag<_BackendTag> __tag, _ExecutionPolicy&& __exec, __tag, ::std::forward<_ExecutionPolicy>(__exec), __par_backend_hetero::make_iter_mode<__par_backend_hetero::access_mode::write>(__first), __par_backend_hetero::make_iter_mode<__par_backend_hetero::access_mode::write>(__last), - fill_functor<_T>{__value}); + oneapi::dpl::__internal::__brick_fill<__hetero_tag<_BackendTag>, _T>{__value}); } //------------------------------------------------------------------------ diff --git a/include/oneapi/dpl/pstl/algorithm_fwd.h b/include/oneapi/dpl/pstl/algorithm_fwd.h index 6af65da659c..0de06e55d9d 100644 --- a/include/oneapi/dpl/pstl/algorithm_fwd.h +++ b/include/oneapi/dpl/pstl/algorithm_fwd.h @@ -74,15 +74,6 @@ template , _ExecutionPolicy&&, _RandomAccessIterator, _RandomAccessIterator, _Function); -template -void -__pattern_walk_brick(_Tag, _ExecutionPolicy&&, _ForwardIterator, _ForwardIterator, _Brick) noexcept; - -template -void -__pattern_walk_brick(__parallel_tag<_IsVector>, _ExecutionPolicy&&, _RandomAccessIterator, _RandomAccessIterator, - _Brick); - //------------------------------------------------------------------------ // walk1_n //------------------------------------------------------------------------ @@ -103,14 +94,6 @@ template , _ExecutionPolicy&&, _RandomAccessIterator, _Size, _Function); -template -_ForwardIterator -__pattern_walk_brick_n(_Tag, _ExecutionPolicy&&, _ForwardIterator, _Size, _Brick) noexcept; - -template -_RandomAccessIterator -__pattern_walk_brick_n(__parallel_tag<_IsVector>, _ExecutionPolicy&&, _RandomAccessIterator, _Size, _Brick); - //------------------------------------------------------------------------ // walk2 (pseudo) // diff --git a/include/oneapi/dpl/pstl/algorithm_impl.h b/include/oneapi/dpl/pstl/algorithm_impl.h index f49b18f54a7..8891933dcc3 100644 --- a/include/oneapi/dpl/pstl/algorithm_impl.h +++ b/include/oneapi/dpl/pstl/algorithm_impl.h @@ -177,30 +177,6 @@ __pattern_walk1(__parallel_tag<_IsVector>, _ExecutionPolicy&& __exec, _RandomAcc }); } -template -void -__pattern_walk_brick(_Tag, _ExecutionPolicy&&, _ForwardIterator __first, _ForwardIterator __last, - _Brick __brick) noexcept -{ - static_assert(__is_serial_tag_v<_Tag> || __is_parallel_forward_tag_v<_Tag>); - - __brick(__first, __last, typename _Tag::__is_vector{}); -} - -template -void -__pattern_walk_brick(__parallel_tag<_IsVector>, _ExecutionPolicy&& __exec, _RandomAccessIterator __first, - _RandomAccessIterator __last, _Brick __brick) -{ - using __backend_tag = typename __parallel_tag<_IsVector>::__backend_tag; - - __internal::__except_handler([&]() { - __par_backend::__parallel_for( - __backend_tag{}, ::std::forward<_ExecutionPolicy>(__exec), __first, __last, - [__brick](_RandomAccessIterator __i, _RandomAccessIterator __j) { __brick(__i, __j, _IsVector{}); }); - }); -} - //------------------------------------------------------------------------ // walk1_n //------------------------------------------------------------------------ @@ -239,30 +215,6 @@ __pattern_walk1_n(__parallel_tag<_IsVector> __tag, _ExecutionPolicy&& __exec, _R return __first + __n; } -template -_ForwardIterator -__pattern_walk_brick_n(_Tag, _ExecutionPolicy&&, _ForwardIterator __first, _Size __n, _Brick __brick) noexcept -{ - static_assert(__is_serial_tag_v<_Tag> || __is_parallel_forward_tag_v<_Tag>); - - return __brick(__first, __n, typename _Tag::__is_vector{}); -} - -template -_RandomAccessIterator -__pattern_walk_brick_n(__parallel_tag<_IsVector>, _ExecutionPolicy&& __exec, _RandomAccessIterator __first, _Size __n, - _Brick __brick) -{ - using __backend_tag = typename __parallel_tag<_IsVector>::__backend_tag; - - return __internal::__except_handler([&]() { - __par_backend::__parallel_for( - __backend_tag{}, ::std::forward<_ExecutionPolicy>(__exec), __first, __first + __n, - [__brick](_RandomAccessIterator __i, _RandomAccessIterator __j) { __brick(__i, __j - __i, _IsVector{}); }); - return __first + __n; - }); -} - //------------------------------------------------------------------------ // walk2 (pseudo) // diff --git a/include/oneapi/dpl/pstl/glue_memory_impl.h b/include/oneapi/dpl/pstl/glue_memory_impl.h index cc2c4bc4f2c..637899759a7 100644 --- a/include/oneapi/dpl/pstl/glue_memory_impl.h +++ b/include/oneapi/dpl/pstl/glue_memory_impl.h @@ -21,6 +21,7 @@ #if _ONEDPL_HETERO_BACKEND # include "hetero/algorithm_impl_hetero.h" +# include "hetero/memory_impl_hetero.h" #endif #include "memory_fwd.h" @@ -147,13 +148,12 @@ uninitialized_fill(_ExecutionPolicy&& __exec, _ForwardIterator __first, _Forward if constexpr (oneapi::dpl::__internal::__trivial_uninitialized_fill<_ValueType, _Tp>) { - oneapi::dpl::__internal::__pattern_walk_brick( - __dispatch_tag, std::forward<_ExecutionPolicy>(__exec), __first, __last, - oneapi::dpl::__internal::__brick_fill{_ValueType(__value)}); + oneapi::dpl::__internal::__pattern_fill(__dispatch_tag, std::forward<_ExecutionPolicy>(__exec), __first, __last, + _ValueType(__value)); } else { - oneapi::dpl::__internal::__pattern_walk1( + oneapi::dpl::__internal::__pattern_uninitialized_walk1( __dispatch_tag, std::forward<_ExecutionPolicy>(__exec), __first, __last, oneapi::dpl::__internal::__op_uninitialized_fill<_Tp, std::decay_t<_ExecutionPolicy>>{__value}); } @@ -169,13 +169,12 @@ uninitialized_fill_n(_ExecutionPolicy&& __exec, _ForwardIterator __first, _Size if constexpr (oneapi::dpl::__internal::__trivial_uninitialized_fill<_ValueType, _Tp>) { - return oneapi::dpl::__internal::__pattern_walk_brick_n( - __dispatch_tag, std::forward<_ExecutionPolicy>(__exec), __first, __n, - oneapi::dpl::__internal::__brick_fill_n{_ValueType(__value)}); + return oneapi::dpl::__internal::__pattern_fill_n(__dispatch_tag, std::forward<_ExecutionPolicy>(__exec), + __first, __n, _ValueType(__value)); } else { - return oneapi::dpl::__internal::__pattern_walk1_n( + return oneapi::dpl::__internal::__pattern_uninitialized_walk1_n( __dispatch_tag, std::forward<_ExecutionPolicy>(__exec), __first, __n, oneapi::dpl::__internal::__op_uninitialized_fill<_Tp, std::decay_t<_ExecutionPolicy>>{__value}); } @@ -243,7 +242,7 @@ uninitialized_default_construct(_ExecutionPolicy&& __exec, _ForwardIterator __fi { const auto __dispatch_tag = oneapi::dpl::__internal::__select_backend(__exec, __first); - oneapi::dpl::__internal::__pattern_walk1( + oneapi::dpl::__internal::__pattern_uninitialized_walk1( __dispatch_tag, std::forward<_ExecutionPolicy>(__exec), __first, __last, oneapi::dpl::__internal::__op_uninitialized_default_construct>{}); } @@ -263,7 +262,7 @@ uninitialized_default_construct_n(_ExecutionPolicy&& __exec, _ForwardIterator __ { const auto __dispatch_tag = oneapi::dpl::__internal::__select_backend(__exec, __first); - return oneapi::dpl::__internal::__pattern_walk1_n( + return oneapi::dpl::__internal::__pattern_uninitialized_walk1_n( __dispatch_tag, std::forward<_ExecutionPolicy>(__exec), __first, __n, oneapi::dpl::__internal::__op_uninitialized_default_construct>{}); } @@ -281,13 +280,12 @@ uninitialized_value_construct(_ExecutionPolicy&& __exec, _ForwardIterator __firs if constexpr (oneapi::dpl::__internal::__trivial_uninitialized_value_construct<_ValueType>) { - oneapi::dpl::__internal::__pattern_walk_brick( - __dispatch_tag, std::forward<_ExecutionPolicy>(__exec), __first, __last, - oneapi::dpl::__internal::__brick_fill{_ValueType()}); + oneapi::dpl::__internal::__pattern_fill(__dispatch_tag, std::forward<_ExecutionPolicy>(__exec), __first, __last, + _ValueType()); } else { - oneapi::dpl::__internal::__pattern_walk1( + oneapi::dpl::__internal::__pattern_uninitialized_walk1( __dispatch_tag, std::forward<_ExecutionPolicy>(__exec), __first, __last, oneapi::dpl::__internal::__op_uninitialized_value_construct>{}); } @@ -303,13 +301,12 @@ uninitialized_value_construct_n(_ExecutionPolicy&& __exec, _ForwardIterator __fi if constexpr (oneapi::dpl::__internal::__trivial_uninitialized_value_construct<_ValueType>) { - return oneapi::dpl::__internal::__pattern_walk_brick_n( - __dispatch_tag, std::forward<_ExecutionPolicy>(__exec), __first, __n, - oneapi::dpl::__internal::__brick_fill_n{_ValueType()}); + return oneapi::dpl::__internal::__pattern_fill_n(__dispatch_tag, std::forward<_ExecutionPolicy>(__exec), + __first, __n, _ValueType()); } else { - return oneapi::dpl::__internal::__pattern_walk1_n( + return oneapi::dpl::__internal::__pattern_uninitialized_walk1_n( __dispatch_tag, std::forward<_ExecutionPolicy>(__exec), __first, __n, oneapi::dpl::__internal::__op_uninitialized_value_construct>{}); } diff --git a/include/oneapi/dpl/pstl/hetero/algorithm_impl_hetero.h b/include/oneapi/dpl/pstl/hetero/algorithm_impl_hetero.h index 6238a0f6121..7e92384040e 100644 --- a/include/oneapi/dpl/pstl/hetero/algorithm_impl_hetero.h +++ b/include/oneapi/dpl/pstl/hetero/algorithm_impl_hetero.h @@ -214,43 +214,6 @@ __pattern_walk3(__hetero_tag<_BackendTag>, _ExecutionPolicy&& __exec, _ForwardIt __f); } -//------------------------------------------------------------------------ -// walk_brick, walk_brick_n -//------------------------------------------------------------------------ - -template -struct __walk_brick_wrapper; - -template -void -__pattern_walk_brick(__hetero_tag<_BackendTag> __tag, _ExecutionPolicy&& __exec, _ForwardIterator __first, - _ForwardIterator __last, _Function __f) -{ - if (__last - __first <= 0) - return; - - __pattern_hetero_walk1<__par_backend_hetero::access_mode::read_write, false>( - __tag, - __par_backend_hetero::make_wrapped_policy<__walk_brick_wrapper>(::std::forward<_ExecutionPolicy>(__exec)), - __first, __last, __f); -} - -template -struct __walk_brick_n_wrapper; - -template -_ForwardIterator -__pattern_walk_brick_n(__hetero_tag<_BackendTag> __tag, _ExecutionPolicy&& __exec, _ForwardIterator __first, _Size __n, - _Function __f) -{ - __pattern_hetero_walk1<__par_backend_hetero::access_mode::read_write, false>( - __tag, - __par_backend_hetero::make_wrapped_policy<__walk_brick_n_wrapper>(::std::forward<_ExecutionPolicy>(__exec)), - __first, __first + __n, __f); - return __first + __n; -} - //------------------------------------------------------------------------ // walk2_brick, walk2_brick_n //------------------------------------------------------------------------ @@ -333,18 +296,6 @@ __pattern_walk3_transform_if(__hetero_tag<_BackendTag> __tag, _ExecutionPolicy&& // fill //------------------------------------------------------------------------ -template -struct fill_functor -{ - _SourceT __value; - template - void - operator()(_TargetT& __target) const - { - __target = __value; - } -}; - template _ForwardIterator __pattern_fill(__hetero_tag<_BackendTag> __tag, _ExecutionPolicy&& __exec, _ForwardIterator __first, @@ -354,7 +305,7 @@ __pattern_fill(__hetero_tag<_BackendTag> __tag, _ExecutionPolicy&& __exec, _Forw __tag, std::forward<_ExecutionPolicy>(__exec), __par_backend_hetero::make_iter_mode<__par_backend_hetero::access_mode::write>(__first), __par_backend_hetero::make_iter_mode<__par_backend_hetero::access_mode::write>(__last), - fill_functor<_T>{__value}); + __brick_fill<__hetero_tag<_BackendTag>, _T>{__value}); return __last; } diff --git a/include/oneapi/dpl/pstl/hetero/dpcpp/sycl_traits.h b/include/oneapi/dpl/pstl/hetero/dpcpp/sycl_traits.h index d176fa29440..d29b3e7d809 100644 --- a/include/oneapi/dpl/pstl/hetero/dpcpp/sycl_traits.h +++ b/include/oneapi/dpl/pstl/hetero/dpcpp/sycl_traits.h @@ -95,9 +95,6 @@ class __replace_functor; template class __replace_copy_functor; -template -struct fill_functor; - template struct generate_functor; @@ -248,12 +245,6 @@ struct sycl::is_device_copyable<_ONEDPL_SPECIALIZE_FOR(oneapi::dpl::__internal:: { }; -template -struct sycl::is_device_copyable<_ONEDPL_SPECIALIZE_FOR(oneapi::dpl::__internal::fill_functor, _SourceT)> - : oneapi::dpl::__internal::__are_all_device_copyable<_SourceT> -{ -}; - template struct sycl::is_device_copyable<_ONEDPL_SPECIALIZE_FOR(oneapi::dpl::__internal::generate_functor, _Generator)> : oneapi::dpl::__internal::__are_all_device_copyable<_Generator> diff --git a/include/oneapi/dpl/pstl/hetero/histogram_impl_hetero.h b/include/oneapi/dpl/pstl/hetero/histogram_impl_hetero.h index e93879cd637..ef784ea646e 100644 --- a/include/oneapi/dpl/pstl/hetero/histogram_impl_hetero.h +++ b/include/oneapi/dpl/pstl/hetero/histogram_impl_hetero.h @@ -134,7 +134,8 @@ __pattern_histogram(__hetero_tag<_BackendTag>, _ExecutionPolicy&& __exec, _Rando auto __bins_buf = __keep_bins(__histogram_first, __histogram_first + __num_bins); auto __bins = __bins_buf.all_view(); - auto __fill_func = oneapi::dpl::__internal::fill_functor<_global_histogram_type>{_global_histogram_type{0}}; + auto __fill_func = oneapi::dpl::__internal::__brick_fill<__hetero_tag<_BackendTag>, _global_histogram_type>{ + _global_histogram_type{0}}; //fill histogram bins with zeros auto __init_event = oneapi::dpl::__par_backend_hetero::__parallel_for( diff --git a/include/oneapi/dpl/pstl/hetero/memory_impl_hetero.h b/include/oneapi/dpl/pstl/hetero/memory_impl_hetero.h new file mode 100644 index 00000000000..a750170f441 --- /dev/null +++ b/include/oneapi/dpl/pstl/hetero/memory_impl_hetero.h @@ -0,0 +1,54 @@ +// -*- C++ -*- +//===----------------------------------------------------------------------===// +// +// Copyright (C) UXL Foundation Contributors +// +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#ifndef _ONEDPL_MEMORY_IMPL_HETERO_H +#define _ONEDPL_MEMORY_IMPL_HETERO_H + +#include "algorithm_impl_hetero.h" + +namespace oneapi +{ +namespace dpl +{ +namespace __internal +{ + +//------------------------------------------------------------------------ +// uninitialized_walk1 +//------------------------------------------------------------------------ + +template +void +__pattern_uninitialized_walk1(__hetero_tag<_BackendTag> tag, _ExecutionPolicy&& __exec, _ForwardIterator __first, + _ForwardIterator __last, _Function __f) +{ + oneapi::dpl::__internal::__pattern_hetero_walk1( + tag, std::forward<_ExecutionPolicy>(__exec), __first, __last, __f); +} + +//------------------------------------------------------------------------ +// uninitialized_walk1_n +//------------------------------------------------------------------------ + +template +_ForwardIterator +__pattern_uninitialized_walk1_n(__hetero_tag<_BackendTag> __tag, _ExecutionPolicy&& __exec, _ForwardIterator __first, + _Size __n, _Function __f) +{ + oneapi::dpl::__internal::__pattern_hetero_walk1( + __tag, std::forward<_ExecutionPolicy>(__exec), __first, __first + __n, __f); + return __first + __n; +} + +} // namespace __internal +} // namespace dpl +} // namespace oneapi + +#endif // _ONEDPL_MEMORY_IMPL_HETERO_H diff --git a/include/oneapi/dpl/pstl/memory_fwd.h b/include/oneapi/dpl/pstl/memory_fwd.h index 79493b9e68e..ef6b5b5a4a8 100644 --- a/include/oneapi/dpl/pstl/memory_fwd.h +++ b/include/oneapi/dpl/pstl/memory_fwd.h @@ -41,6 +41,14 @@ struct __op_uninitialized_default_construct; template struct __op_uninitialized_value_construct; +template +void +__pattern_uninitialized_walk1(_Tag, _ExecutionPolicy&&, _ForwardIterator, _ForwardIterator, _Function); + +template +_ForwardIterator +__pattern_uninitialized_walk1_n(_Tag, _ExecutionPolicy&&, _ForwardIterator, _Size, _Function); + } // namespace __internal } // namespace dpl } // namespace oneapi diff --git a/include/oneapi/dpl/pstl/memory_impl.h b/include/oneapi/dpl/pstl/memory_impl.h index 2ac5870a3ad..7111283586c 100644 --- a/include/oneapi/dpl/pstl/memory_impl.h +++ b/include/oneapi/dpl/pstl/memory_impl.h @@ -21,6 +21,7 @@ #include "memory_fwd.h" #include "unseq_backend_simd.h" +#include "algorithm_fwd.h" namespace oneapi { @@ -210,6 +211,21 @@ struct __op_uninitialized_value_construct<_ExecutionPolicy> } }; +template +void +__pattern_uninitialized_walk1(_Tag tag, _ExecutionPolicy&& __exec, _ForwardIterator __first, _ForwardIterator __last, + _Function __f) +{ + __pattern_walk1(tag, std::forward<_ExecutionPolicy>(__exec), __first, __last, __f); +} + +template +_ForwardIterator +__pattern_uninitialized_walk1_n(_Tag tag, _ExecutionPolicy&& __exec, _ForwardIterator __first, _Size __n, _Function __f) +{ + return __pattern_walk1_n(tag, std::forward<_ExecutionPolicy>(__exec), __first, __n, __f); +} + } // namespace __internal } // namespace dpl } // namespace oneapi diff --git a/include/oneapi/dpl/pstl/memory_ranges_impl.h b/include/oneapi/dpl/pstl/memory_ranges_impl.h index 39e8e6e41a6..9805d6dc329 100644 --- a/include/oneapi/dpl/pstl/memory_ranges_impl.h +++ b/include/oneapi/dpl/pstl/memory_ranges_impl.h @@ -54,7 +54,7 @@ __pattern_uninitialized_default_construct(_Tag __tag, _ExecutionPolicy&& __exec, if constexpr (!std::is_trivially_default_constructible_v<_ValueType>) { - oneapi::dpl::__internal::__pattern_walk1( + oneapi::dpl::__internal::__pattern_uninitialized_walk1( __tag, std::forward<_ExecutionPolicy>(__exec), __first, __last, oneapi::dpl::__internal::__op_uninitialized_default_construct>{}); } @@ -85,13 +85,12 @@ __pattern_uninitialized_value_construct(_Tag __tag, _ExecutionPolicy&& __exec, _ auto __last = __first + std::ranges::size(__r); if constexpr (oneapi::dpl::__internal::__trivial_uninitialized_value_construct<_ValueType>) { - oneapi::dpl::__internal::__pattern_walk_brick( - __tag, std::forward<_ExecutionPolicy>(__exec), __first, __last, - oneapi::dpl::__internal::__brick_fill<_Tag, _ValueType>{_ValueType()}); + oneapi::dpl::__internal::__pattern_fill(__tag, std::forward<_ExecutionPolicy>(__exec), __first, __last, + _ValueType()); } else { - oneapi::dpl::__internal::__pattern_walk1( + oneapi::dpl::__internal::__pattern_uninitialized_walk1( __tag, std::forward<_ExecutionPolicy>(__exec), __first, __last, oneapi::dpl::__internal::__op_uninitialized_value_construct>{}); } @@ -221,13 +220,12 @@ __pattern_uninitialized_fill(_Tag __tag, _ExecutionPolicy&& __exec, _R&& __r, co if constexpr (oneapi::dpl::__internal::__trivial_uninitialized_fill<_ValueType, _T>) { - oneapi::dpl::__internal::__pattern_walk_brick( - __tag, std::forward<_ExecutionPolicy>(__exec), __first, __last, - oneapi::dpl::__internal::__brick_fill<_Tag, _ValueType>{_ValueType(__value)}); + oneapi::dpl::__internal::__pattern_fill(__tag, std::forward<_ExecutionPolicy>(__exec), __first, __last, + _ValueType(__value)); } else { - oneapi::dpl::__internal::__pattern_walk1( + oneapi::dpl::__internal::__pattern_uninitialized_walk1( __tag, std::forward<_ExecutionPolicy>(__exec), __first, __last, oneapi::dpl::__internal::__op_uninitialized_fill<_T, std::decay_t<_ExecutionPolicy>>{__value}); } diff --git a/test/general/implementation_details/device_copyable.pass.cpp b/test/general/implementation_details/device_copyable.pass.cpp index 4af36305b68..fd963f0a676 100644 --- a/test/general/implementation_details/device_copyable.pass.cpp +++ b/test/general/implementation_details/device_copyable.pass.cpp @@ -338,9 +338,6 @@ test_device_copyable() static_assert(sycl::is_device_copyable_v< oneapi::dpl::__internal::__replace_copy_functor>, "__replace_copy_functor is not device copyable with device copyable types"); - //fill_functor - static_assert(sycl::is_device_copyable_v>, - "fill_functor is not device copyable with device copyable types"); //generate_functor static_assert(sycl::is_device_copyable_v>, "generate_functor is not device copyable with device copyable types"); @@ -676,10 +673,6 @@ test_non_device_copyable() oneapi::dpl::__internal::__replace_copy_functor>, "__replace_copy_functor is device copyable with non device copyable types"); - //fill_functor - static_assert(!sycl::is_device_copyable_v>, - "fill_functor is device copyable with non device copyable types"); - //generate_functor static_assert(!sycl::is_device_copyable_v>, "generate_functor is device copyable with non device copyable types");