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

@@ -171,11 +171,8 @@ auto
 __pattern_fill_async(__hetero_tag<_BackendTag> __tag, _ExecutionPolicy&& __exec, _ForwardIterator __first,
                      _ForwardIterator __last, const _T& __value)
 {
-    return __pattern_walk1_async(
-        __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),
-        oneapi::dpl::__internal::__brick_fill<__hetero_tag<_BackendTag>, _T>{__value});
+    return __pattern_walk1_async(__tag, ::std::forward<_ExecutionPolicy>(__exec), __first, __last,
+                                 oneapi::dpl::__internal::__brick_fill<__hetero_tag<_BackendTag>, _T>{__value});
 }
 
 //------------------------------------------------------------------------

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

@@ -291,9 +291,7 @@ _ForwardIterator
 __pattern_fill(__hetero_tag<_BackendTag> __tag, _ExecutionPolicy&& __exec, _ForwardIterator __first,
                _ForwardIterator __last, const _T& __value)
 {
-    __pattern_walk1(__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),
+    __pattern_walk1(__tag, ::std::forward<_ExecutionPolicy>(__exec), __first, __last,
                     __brick_fill<__hetero_tag<_BackendTag>, _T>{__value});
     return __last;
 }
@@ -328,9 +326,7 @@ _ForwardIterator
 __pattern_generate(__hetero_tag<_BackendTag> __tag, _ExecutionPolicy&& __exec, _ForwardIterator __first,
                    _ForwardIterator __last, _Generator __g)
 {
-    __pattern_walk1(__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),
+    __pattern_walk1(__tag, ::std::forward<_ExecutionPolicy>(__exec), __first,__last,
                     generate_functor<_Generator>{__g});
     return __last;
 }
@@ -908,9 +904,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, /*_NoInit=*/true>();
@@ -1211,12 +1205,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
 
@@ -1454,11 +1443,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();
 }
 
@@ -1551,10 +1537,8 @@ __pattern_partial_sort_copy(__hetero_tag<_BackendTag> __tag, _ExecutionPolicy&&
         // the patterns: __pattern_walk2, __parallel_partial_sort and __pattern_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_walk2(
             __tag, __par_backend_hetero::make_wrapped_policy<__copy_back>(::std::forward<_ExecutionPolicy>(__exec)),

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

@@ -66,136 +66,6 @@ namespace dpl
 namespace __par_backend_hetero
 {
 
-//-----------------------------------------------------------------------------
-//- iter_mode_resolver
-//-----------------------------------------------------------------------------
-
-// iter_mode_resolver resolves the situations when
-// the access mode provided by a user differs (inMode) from
-// the access mode required by an algorithm (outMode).
-// In general case iter_mode_resolver accepts the only situations
-// when inMode == outMode,
-// whereas the template specializations describe cases with specific
-// inMode and outMode and the preferred access mode between the two.
-template <access_mode inMode, access_mode outMode>
-struct iter_mode_resolver
-{
-    static_assert(inMode == outMode, "Access mode provided by user conflicts with the one required by the algorithm");
-    static constexpr access_mode value = inMode;
-};
-
-template <>
-struct iter_mode_resolver<access_mode::read, access_mode::read_write>
-{
-    static constexpr access_mode value = access_mode::read;
-};
-
-template <>
-struct iter_mode_resolver<access_mode::write, access_mode::read_write>
-{
-    static constexpr access_mode value = access_mode::write;
-};
-
-template <>
-struct iter_mode_resolver<access_mode::read_write, access_mode::read>
-{
-    //TODO: warn user that the access mode is changed
-    static constexpr access_mode value = access_mode::read;
-};
-
-template <>
-struct iter_mode_resolver<access_mode::read_write, access_mode::write>
-{
-    //TODO: warn user that the access mode is changed
-    static constexpr access_mode value = access_mode::write;
-};
-
-template <>
-struct iter_mode_resolver<access_mode::discard_write, access_mode::write>
-{
-    static constexpr access_mode value = access_mode::discard_write;
-};
-
-template <>
-struct iter_mode_resolver<access_mode::discard_read_write, access_mode::write>
-{
-    //TODO: warn user that the access mode is changed
-    static constexpr access_mode value = access_mode::write;
-};
-
-template <>
-struct iter_mode_resolver<access_mode::discard_read_write, access_mode::read_write>
-{
-    static constexpr access_mode value = access_mode::discard_read_write;
-};
-
-//-----------------------------------------------------------------------------
-//- iter_mode
-//-----------------------------------------------------------------------------
-
-// create iterator with different access mode
-template <access_mode outMode>
-struct iter_mode
-{
-    // for common heterogeneous iterator
-    template <template <access_mode, typename...> class Iter, access_mode inMode, typename... Types>
-    Iter<iter_mode_resolver<inMode, outMode>::value, Types...>
-    operator()(const Iter<inMode, Types...>& it)
-    {
-        constexpr access_mode preferredMode = iter_mode_resolver<inMode, outMode>::value;
-        if (inMode == preferredMode)
-            return it;
-        return Iter<preferredMode, Types...>(it);
-    }
-    // for counting_iterator
-    template <typename T>
-    oneapi::dpl::counting_iterator<T>
-    operator()(const oneapi::dpl::counting_iterator<T>& it)
-    {
-        return it;
-    }
-    // for zip_iterator
-    template <typename... Iters>
-    auto
-    operator()(const oneapi::dpl::zip_iterator<Iters...>& it)
-        -> decltype(oneapi::dpl::__internal::map_zip(*this, it.base()))
-    {
-        return oneapi::dpl::__internal::map_zip(*this, it.base());
-    }
-    // for common iterator
-    template <typename Iter>
-    Iter
-    operator()(const Iter& it1)
-    {
-        return it1;
-    }
-    // for raw pointers
-    template <typename T>
-    T*
-    operator()(T* ptr)
-    {
-        // it does not have any iter mode because of two factors:
-        //   - since it is a raw pointer, kernel can read/write despite of access_mode
-        //   - access_mode also serves for implicit synchronization for buffers to build graph dependency
-        //     and since usm have only explicit synchronization and does not provide dependency resolution mechanism
-        //     it does not require access_mode
-        return ptr;
-    }
-
-    template <typename T>
-    const T*
-    operator()(const T* ptr)
-    {
-        return ptr;
-    }
-};
-
-template <access_mode outMode, typename _Iterator>
-auto
-make_iter_mode(const _Iterator& __it) -> decltype(iter_mode<outMode>()(__it))
-{
-    return iter_mode<outMode>()(__it);
-}
 
 // set of class templates to name kernels