Remove make_iter_mode and refactor access mode resolution

include/oneapi/dpl/experimental/kt/internal/esimd_radix_sort_utils.h

@@ -256,9 +256,9 @@ __rng_data(const _Rng& __rng)
 //      sycl::accessor::operator[] are supported only with -fsycl-esimd-force-stateless-mem.
 //      Otherwise, all memory accesses through an accessor are done via explicit APIs
 // TODO: rely on begin() once -fsycl-esimd-force-stateless-mem has been enabled by default
-template <typename _T, sycl::access::mode _M>
+template <typename _T, sycl::access::mode _M, bool _NoInit>
 auto
-__rng_data(const oneapi::dpl::__ranges::all_view<_T, _M>& __view)
+__rng_data(const oneapi::dpl::__ranges::all_view<_T, _M, _NoInit>& __view)
 {
     return __view.accessor();
 }

include/oneapi/dpl/internal/async_impl/async_impl_hetero.h

@@ -169,10 +169,8 @@ __pattern_fill_async(__hetero_tag<_BackendTag> __tag, _ExecutionPolicy&& __exec,
                      _ForwardIterator __last, const _T& __value)
 {
     return __pattern_walk1_async<__par_backend_hetero::access_mode::write, /*_IsNoInitRequested=*/true>(
-        __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});
+        __tag, std::forward<_ExecutionPolicy>(__exec), __first, __last,
+        oneapi::dpl::__internal::__brick_fill<__hetero_tag<_BackendTag>, _T>{__value});
 }
 
 //------------------------------------------------------------------------

include/oneapi/dpl/pstl/algorithm_fwd.h

@@ -74,15 +74,6 @@ template <class _IsVector, class _ExecutionPolicy, class _RandomAccessIterator,
 void
 __pattern_walk1(__parallel_tag<_IsVector>, _ExecutionPolicy&&, _RandomAccessIterator, _RandomAccessIterator, _Function);
 
-template <class _Tag, class _ExecutionPolicy, class _ForwardIterator, class _Brick>
-void
-__pattern_walk_brick(_Tag, _ExecutionPolicy&&, _ForwardIterator, _ForwardIterator, _Brick) noexcept;
-
-template <class _IsVector, class _ExecutionPolicy, class _RandomAccessIterator, class _Brick>
-void
-__pattern_walk_brick(__parallel_tag<_IsVector>, _ExecutionPolicy&&, _RandomAccessIterator, _RandomAccessIterator,
-                     _Brick);
-
 //------------------------------------------------------------------------
 // walk1_n
 //------------------------------------------------------------------------
@@ -103,14 +94,6 @@ template <class _IsVector, class _ExecutionPolicy, class _RandomAccessIterator,
 _RandomAccessIterator
 __pattern_walk1_n(__parallel_tag<_IsVector>, _ExecutionPolicy&&, _RandomAccessIterator, _Size, _Function);
 
-template <class _Tag, class _ExecutionPolicy, class _ForwardIterator, class _Size, class _Brick>
-_ForwardIterator
-__pattern_walk_brick_n(_Tag, _ExecutionPolicy&&, _ForwardIterator, _Size, _Brick) noexcept;
-
-template <class _IsVector, class _ExecutionPolicy, class _RandomAccessIterator, class _Size, class _Brick>
-_RandomAccessIterator
-__pattern_walk_brick_n(__parallel_tag<_IsVector>, _ExecutionPolicy&&, _RandomAccessIterator, _Size, _Brick);
-
 //------------------------------------------------------------------------
 // walk2 (pseudo)
 //

include/oneapi/dpl/pstl/algorithm_impl.h

@@ -177,30 +177,6 @@ __pattern_walk1(__parallel_tag<_IsVector>, _ExecutionPolicy&& __exec, _RandomAcc
     });
 }
 
-template <class _Tag, class _ExecutionPolicy, class _ForwardIterator, class _Brick>
-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 <class _IsVector, class _ExecutionPolicy, class _RandomAccessIterator, class _Brick>
-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 <class _Tag, class _ExecutionPolicy, class _ForwardIterator, class _Size, class _Brick>
-_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 <class _IsVector, class _ExecutionPolicy, class _RandomAccessIterator, class _Size, class _Brick>
-_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)
 //

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<decltype(__dispatch_tag), _ValueType>{_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<decltype(__dispatch_tag), _ValueType>{_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<std::decay_t<_ExecutionPolicy>>{});
     }
@@ -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<std::decay_t<_ExecutionPolicy>>{});
     }
@@ -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<decltype(__dispatch_tag), _ValueType>{_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<std::decay_t<_ExecutionPolicy>>{});
     }
@@ -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<decltype(__dispatch_tag), _ValueType>{_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<std::decay_t<_ExecutionPolicy>>{});
     }

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 <typename _Name>
-struct __walk_brick_wrapper;
-
-template <typename _BackendTag, typename _ExecutionPolicy, typename _ForwardIterator, typename _Function>
-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 <typename _Name>
-struct __walk_brick_n_wrapper;
-
-template <typename _BackendTag, typename _ExecutionPolicy, typename _ForwardIterator, typename _Size,
-          typename _Function>
-_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,28 +296,14 @@ __pattern_walk3_transform_if(__hetero_tag<_BackendTag> __tag, _ExecutionPolicy&&
 // fill
 //------------------------------------------------------------------------
 
-template <typename _SourceT>
-struct fill_functor
-{
-    _SourceT __value;
-    template <typename _TargetT>
-    void
-    operator()(_TargetT& __target) const
-    {
-        __target = __value;
-    }
-};
-
 template <typename _BackendTag, typename _ExecutionPolicy, typename _ForwardIterator, typename _T>
 _ForwardIterator
 __pattern_fill(__hetero_tag<_BackendTag> __tag, _ExecutionPolicy&& __exec, _ForwardIterator __first,
                _ForwardIterator __last, const _T& __value)
 {
     __pattern_hetero_walk1<__par_backend_hetero::access_mode::write, /*_IsNoInitRequested=*/true>(
-        __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});
+        __tag, std::forward<_ExecutionPolicy>(__exec), __first, __last,
+        __brick_fill<__hetero_tag<_BackendTag>, _T>{__value});
     return __last;
 }
 
@@ -389,10 +338,7 @@ __pattern_generate(__hetero_tag<_BackendTag> __tag, _ExecutionPolicy&& __exec, _
                    _ForwardIterator __last, _Generator __g)
 {
     __pattern_hetero_walk1<__par_backend_hetero::access_mode::write, /*_IsNoInitRequested=*/true>(
-        __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),
-        generate_functor<_Generator>{__g});
+        __tag, std::forward<_ExecutionPolicy>(__exec), __first, __last, generate_functor<_Generator>{__g});
     return __last;
 }
 
@@ -968,9 +914,7 @@ __pattern_partition_copy(__hetero_tag<_BackendTag>, _ExecutionPolicy&& __exec, _
     auto __keep1 = oneapi::dpl::__ranges::__get_sycl_range<__par_backend_hetero::access_mode::read>();
     auto __buf1 = __keep1(__first, __last);
 
-    auto __zipped_res = __par_backend_hetero::zip(
-        __par_backend_hetero::make_iter_mode<__par_backend_hetero::access_mode::write>(__result1),
-        __par_backend_hetero::make_iter_mode<__par_backend_hetero::access_mode::write>(__result2));
+    auto __zipped_res = __par_backend_hetero::zip(__result1, __result2);
 
     auto __keep2 = oneapi::dpl::__ranges::__get_sycl_range<__par_backend_hetero::access_mode::write,
                                                            /*_IsNoInitRequested=*/true>();
@@ -1271,12 +1215,7 @@ __pattern_inplace_merge(__hetero_tag<_BackendTag> __tag, _ExecutionPolicy&& __ex
     auto __copy_first = __buf.get();
     auto __copy_last = __copy_first + __n;
 
-    __pattern_merge(
-        __tag, __exec, __par_backend_hetero::make_iter_mode<__par_backend_hetero::access_mode::read>(__first),
-        __par_backend_hetero::make_iter_mode<__par_backend_hetero::access_mode::read>(__middle),
-        __par_backend_hetero::make_iter_mode<__par_backend_hetero::access_mode::read>(__middle),
-        __par_backend_hetero::make_iter_mode<__par_backend_hetero::access_mode::read>(__last),
-        __par_backend_hetero::make_iter_mode<__par_backend_hetero::access_mode::write>(__copy_first), __comp);
+    __pattern_merge(__tag, __exec, __first, __middle, __middle, __last, __copy_first, __comp);
 
     //TODO: optimize copy back depending on Iterator, i.e. set_final_data for host iterator/pointer
 
@@ -1517,11 +1456,8 @@ __pattern_partial_sort(__hetero_tag<_BackendTag>, _ExecutionPolicy&& __exec, _It
     if (__last - __first < 2)
         return;
 
-    __par_backend_hetero::__parallel_partial_sort(
-        _BackendTag{}, ::std::forward<_ExecutionPolicy>(__exec),
-        __par_backend_hetero::make_iter_mode<__par_backend_hetero::access_mode::read_write>(__first),
-        __par_backend_hetero::make_iter_mode<__par_backend_hetero::access_mode::read_write>(__mid),
-        __par_backend_hetero::make_iter_mode<__par_backend_hetero::access_mode::read_write>(__last), __comp)
+    __par_backend_hetero::__parallel_partial_sort(_BackendTag{}, ::std::forward<_ExecutionPolicy>(__exec), __first,
+                                                  __mid, __last, __comp)
         .__checked_deferrable_wait();
 }
 
@@ -1618,10 +1554,8 @@ __pattern_partial_sort_copy(__hetero_tag<_BackendTag> __tag, _ExecutionPolicy&&
         // the patterns: __pattern_hetero_walk2, __parallel_partial_sort and __pattern_hetero_walk2.
 
         __par_backend_hetero::__parallel_partial_sort(
-            _BackendTag{}, __par_backend_hetero::make_wrapped_policy<__partial_sort_2>(__exec),
-            __par_backend_hetero::make_iter_mode<__par_backend_hetero::access_mode::read_write>(__buf_first),
-            __par_backend_hetero::make_iter_mode<__par_backend_hetero::access_mode::read_write>(__buf_mid),
-            __par_backend_hetero::make_iter_mode<__par_backend_hetero::access_mode::read_write>(__buf_last), __comp);
+            _BackendTag{}, __par_backend_hetero::make_wrapped_policy<__partial_sort_2>(__exec), __buf_first, __buf_mid,
+            __buf_last, __comp);
 
         return __pattern_hetero_walk2<__par_backend_hetero::__deferrable_mode, __par_backend_hetero::access_mode::write,
                                       /*_IsOutNoInitRequested=*/true>(