diff --git a/include/oneapi/dpl/experimental/kt/internal/esimd_radix_sort_utils.h b/include/oneapi/dpl/experimental/kt/internal/esimd_radix_sort_utils.h
index b2a291822c4..7333fd6ac9d 100644
--- a/include/oneapi/dpl/experimental/kt/internal/esimd_radix_sort_utils.h
+++ b/include/oneapi/dpl/experimental/kt/internal/esimd_radix_sort_utils.h
@@ -219,9 +219,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();
 }
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 da8f4679ed1..dc687f893d5 100644
--- a/include/oneapi/dpl/internal/async_impl/async_impl_hetero.h
+++ b/include/oneapi/dpl/internal/async_impl/async_impl_hetero.h
@@ -30,8 +30,8 @@ namespace dpl
 namespace __internal
 {
 
-template <__par_backend_hetero::access_mode __acc_mode, bool _IsNoInitRequested, typename _BackendTag,
-          typename _ExecutionPolicy, typename _ForwardIterator, typename _Function>
+template <__par_backend_hetero::access_mode __acc_mode, bool _IsNoInitRequested, bool _DeferToUserHint = false,
+          typename _BackendTag, typename _ExecutionPolicy, typename _ForwardIterator, typename _Function>
 auto
 __pattern_walk1_async(__hetero_tag<_BackendTag>, _ExecutionPolicy&& __exec, _ForwardIterator __first,
                       _ForwardIterator __last, _Function __f)
@@ -39,7 +39,7 @@ __pattern_walk1_async(__hetero_tag<_BackendTag>, _ExecutionPolicy&& __exec, _For
     auto __n = __last - __first;
     assert(__n > 0);
 
-    auto __keep = oneapi::dpl::__ranges::__get_sycl_range<__acc_mode, _IsNoInitRequested>();
+    auto __keep = oneapi::dpl::__ranges::__get_sycl_range<__acc_mode, _IsNoInitRequested, _DeferToUserHint>();
     auto __buf = __keep(__first, __last);
 
     auto __future_obj = oneapi::dpl::__par_backend_hetero::__parallel_for(
@@ -108,6 +108,16 @@ __pattern_walk2_brick_async(__hetero_tag<_BackendTag> __tag, _ExecutionPolicy&&
         __first1, __last1, __first2, __brick);
 }
 
+template <typename _BackendTag, typename _ExecutionPolicy, typename _ForwardIterator, typename _Function>
+auto
+__pattern_for_each_async(__hetero_tag<_BackendTag>, _ExecutionPolicy&& __exec, _ForwardIterator __first,
+                         _ForwardIterator __last, _Function __f)
+{
+    return __pattern_walk1_async<__par_backend_hetero::access_mode::read_write, /*_IsNoInitRequested=*/false,
+                                 /*_DeferToUserHint=*/true>(
+        __hetero_tag<_BackendTag>{}, std::forward<_ExecutionPolicy>(__exec), __first, __last, __f);
+}
+
 //------------------------------------------------------------------------
 // transform_reduce (version with two binary functions)
 //------------------------------------------------------------------------
@@ -169,9 +179,7 @@ __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),
+        __tag, std::forward<_ExecutionPolicy>(__exec), __first, __last,
         oneapi::dpl::__internal::__brick_fill<__hetero_tag<_BackendTag>, _T>{__value});
 }
 
diff --git a/include/oneapi/dpl/internal/async_impl/glue_async_impl.h b/include/oneapi/dpl/internal/async_impl/glue_async_impl.h
index 5acb764abf3..b120f5e35b5 100644
--- a/include/oneapi/dpl/internal/async_impl/glue_async_impl.h
+++ b/include/oneapi/dpl/internal/async_impl/glue_async_impl.h
@@ -130,9 +130,8 @@ for_each_async(_ExecutionPolicy&& __exec, _ForwardIterator __first, _ForwardIter
     const auto __dispatch_tag = oneapi::dpl::__internal::__select_backend(__exec, __first);
 
     wait_for_all(::std::forward<_Events>(__dependencies)...);
-    auto ret_val = oneapi::dpl::__internal::__pattern_walk1_async<__par_backend_hetero::access_mode::read_write,
-                                                                  /*_IsNoInitRequested=*/false>(
-        __dispatch_tag, ::std::forward<_ExecutionPolicy>(__exec), __first, __last, __f);
+    auto ret_val = oneapi::dpl::__internal::__pattern_for_each_async(
+        __dispatch_tag, std::forward<_ExecutionPolicy>(__exec), __first, __last, __f);
     return ret_val;
 }
 
diff --git a/include/oneapi/dpl/pstl/algorithm_fwd.h b/include/oneapi/dpl/pstl/algorithm_fwd.h
index 0de06e55d9d..f23482ad57d 100644
--- a/include/oneapi/dpl/pstl/algorithm_fwd.h
+++ b/include/oneapi/dpl/pstl/algorithm_fwd.h
@@ -74,6 +74,10 @@ 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 _Function>
+void
+__pattern_for_each(_Tag, _ExecutionPolicy&&, _ForwardIterator, _ForwardIterator, _Function);
+
 //------------------------------------------------------------------------
 // walk1_n
 //------------------------------------------------------------------------
@@ -94,6 +98,10 @@ 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 _Function>
+_ForwardIterator
+__pattern_for_each_n(_Tag, _ExecutionPolicy&&, _ForwardIterator, _Size, _Function);
+
 //------------------------------------------------------------------------
 // walk2 (pseudo)
 //
diff --git a/include/oneapi/dpl/pstl/algorithm_impl.h b/include/oneapi/dpl/pstl/algorithm_impl.h
index 7cc71a9d8b1..290b0a7ee9f 100644
--- a/include/oneapi/dpl/pstl/algorithm_impl.h
+++ b/include/oneapi/dpl/pstl/algorithm_impl.h
@@ -178,6 +178,13 @@ __pattern_walk1(__parallel_tag<_IsVector>, _ExecutionPolicy&& __exec, _RandomAcc
     });
 }
 
+template <class _Tag, class _ExecutionPolicy, class _ForwardIterator, class _Function>
+void
+__pattern_for_each(_Tag, _ExecutionPolicy&& __exec, _ForwardIterator __first, _ForwardIterator __last, _Function __f)
+{
+    __pattern_walk1(_Tag{}, std::forward<_ExecutionPolicy>(__exec), __first, __last, __f);
+}
+
 //------------------------------------------------------------------------
 // walk1_n
 //------------------------------------------------------------------------
@@ -216,6 +223,13 @@ __pattern_walk1_n(__parallel_tag<_IsVector> __tag, _ExecutionPolicy&& __exec, _R
     return __first + __n;
 }
 
+template <class _Tag, class _ExecutionPolicy, class _ForwardIterator, class _Size, class _Function>
+_ForwardIterator
+__pattern_for_each_n(_Tag, _ExecutionPolicy&& __exec, _ForwardIterator __first, _Size __n, _Function __f)
+{
+    return __pattern_walk1_n(_Tag{}, std::forward<_ExecutionPolicy>(__exec), __first, __n, __f);
+}
+
 //------------------------------------------------------------------------
 // walk2 (pseudo)
 //
diff --git a/include/oneapi/dpl/pstl/glue_algorithm_impl.h b/include/oneapi/dpl/pstl/glue_algorithm_impl.h
index 4f480acdd0d..b1d5fc7ef52 100644
--- a/include/oneapi/dpl/pstl/glue_algorithm_impl.h
+++ b/include/oneapi/dpl/pstl/glue_algorithm_impl.h
@@ -78,8 +78,8 @@ for_each(_ExecutionPolicy&& __exec, _ForwardIterator __first, _ForwardIterator _
 {
     const auto __dispatch_tag = oneapi::dpl::__internal::__select_backend(__exec, __first);
 
-    oneapi::dpl::__internal::__pattern_walk1(__dispatch_tag, ::std::forward<_ExecutionPolicy>(__exec), __first, __last,
-                                             __f);
+    oneapi::dpl::__internal::__pattern_for_each(__dispatch_tag, ::std::forward<_ExecutionPolicy>(__exec), __first,
+                                                __last, __f);
 }
 
 template <class _ExecutionPolicy, class _ForwardIterator, class _Size, class _Function>
@@ -88,8 +88,8 @@ for_each_n(_ExecutionPolicy&& __exec, _ForwardIterator __first, _Size __n, _Func
 {
     const auto __dispatch_tag = oneapi::dpl::__internal::__select_backend(__exec, __first);
 
-    return oneapi::dpl::__internal::__pattern_walk1_n(__dispatch_tag, ::std::forward<_ExecutionPolicy>(__exec), __first,
-                                                      __n, __f);
+    return oneapi::dpl::__internal::__pattern_for_each_n(__dispatch_tag, ::std::forward<_ExecutionPolicy>(__exec),
+                                                         __first, __n, __f);
 }
 
 // [alg.find]
diff --git a/include/oneapi/dpl/pstl/hetero/algorithm_impl_hetero.h b/include/oneapi/dpl/pstl/hetero/algorithm_impl_hetero.h
index 1422749848a..8667343ea0f 100644
--- a/include/oneapi/dpl/pstl/hetero/algorithm_impl_hetero.h
+++ b/include/oneapi/dpl/pstl/hetero/algorithm_impl_hetero.h
@@ -43,8 +43,8 @@ namespace __internal
 // walk1
 //------------------------------------------------------------------------
 
-template <__par_backend_hetero::access_mode __acc_mode, bool _IsNoInitRequested, typename _BackendTag,
-          typename _ExecutionPolicy, typename _ForwardIterator, typename _Function>
+template <__par_backend_hetero::access_mode __acc_mode, bool _IsNoInitRequested, bool _DeferToUserHint = false,
+          typename _BackendTag, typename _ExecutionPolicy, typename _ForwardIterator, typename _Function>
 void
 __pattern_hetero_walk1(__hetero_tag<_BackendTag>, _ExecutionPolicy&& __exec, _ForwardIterator __first,
                        _ForwardIterator __last, _Function __f)
@@ -53,7 +53,7 @@ __pattern_hetero_walk1(__hetero_tag<_BackendTag>, _ExecutionPolicy&& __exec, _Fo
     if (__n <= 0)
         return;
 
-    auto __keep = oneapi::dpl::__ranges::__get_sycl_range<__acc_mode, _IsNoInitRequested>();
+    auto __keep = oneapi::dpl::__ranges::__get_sycl_range<__acc_mode, _IsNoInitRequested, _DeferToUserHint>();
     auto __buf = __keep(__first, __last);
 
     oneapi::dpl::__par_backend_hetero::__parallel_for(
@@ -71,6 +71,16 @@ __pattern_walk1(__hetero_tag<_BackendTag>, _ExecutionPolicy&& __exec, _ForwardIt
         __hetero_tag<_BackendTag>{}, std::forward<_ExecutionPolicy>(__exec), __first, __last, __f);
 }
 
+template <class _BackendTag, class _ExecutionPolicy, class _ForwardIterator, class _Function>
+void
+__pattern_for_each(__hetero_tag<_BackendTag>, _ExecutionPolicy&& __exec, _ForwardIterator __first,
+                   _ForwardIterator __last, _Function __f)
+{
+    __pattern_hetero_walk1<__par_backend_hetero::access_mode::read_write, /*_IsNoInitRequested=*/false,
+                           /*_DeferToUserHint=*/true>(__hetero_tag<_BackendTag>{},
+                                                      std::forward<_ExecutionPolicy>(__exec), __first, __last, __f);
+}
+
 //------------------------------------------------------------------------
 // walk1_n
 //------------------------------------------------------------------------
@@ -86,6 +96,18 @@ __pattern_walk1_n(__hetero_tag<_BackendTag> __tag, _ExecutionPolicy&& __exec, _F
     return __first + __n;
 }
 
+template <typename _BackendTag, typename _ExecutionPolicy, typename _ForwardIterator, typename _Size,
+          typename _Function>
+_ForwardIterator
+__pattern_for_each_n(__hetero_tag<_BackendTag> __tag, _ExecutionPolicy&& __exec, _ForwardIterator __first, _Size __n,
+                     _Function __f)
+{
+    __pattern_hetero_walk1<__par_backend_hetero::access_mode::read_write, /*_IsNoInitRequested=*/false,
+                           /*_DeferToUserHint=*/true>(__tag, std::forward<_ExecutionPolicy>(__exec), __first,
+                                                      __first + __n, __f);
+    return __first + __n;
+}
+
 //------------------------------------------------------------------------
 // walk2
 //------------------------------------------------------------------------
@@ -302,9 +324,7 @@ __pattern_fill(__hetero_tag<_BackendTag> __tag, _ExecutionPolicy&& __exec, _Forw
                _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),
+        __tag, std::forward<_ExecutionPolicy>(__exec), __first, __last,
         __brick_fill<__hetero_tag<_BackendTag>, _T>{__value});
     return __last;
 }
@@ -340,10 +360,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;
 }
 
@@ -919,9 +936,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>();
@@ -1222,12 +1237,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
 
@@ -1468,11 +1478,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();
 }
 
@@ -1569,10 +1576,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>(
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 f1da2184988..aca43ba4211 100644
--- a/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h
+++ b/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
 
diff --git a/include/oneapi/dpl/pstl/hetero/dpcpp/utils_ranges_sycl.h b/include/oneapi/dpl/pstl/hetero/dpcpp/utils_ranges_sycl.h
index 1fa148e79ad..0459d3350f7 100644
--- a/include/oneapi/dpl/pstl/hetero/dpcpp/utils_ranges_sycl.h
+++ b/include/oneapi/dpl/pstl/hetero/dpcpp/utils_ranges_sycl.h
@@ -36,7 +36,7 @@ namespace __internal
 {
 template <typename _AccessorType, typename _BufferType, typename _DiffType>
 static _AccessorType
-__create_accessor(_BufferType& __buf, _DiffType __offset, _DiffType __n)
+__create_accessor(_BufferType& __buf, _DiffType __offset, _DiffType __n, const sycl::property_list& __properties = {})
 {
     auto __n_buf = __dpl_sycl::__get_buffer_size(__buf);
     auto __n_acc = (__n > 0 ? __n : __n_buf);
@@ -44,12 +44,12 @@ __create_accessor(_BufferType& __buf, _DiffType __offset, _DiffType __n)
     assert(__offset + __n_acc <= __n_buf &&
            "The sum of accessRange and accessOffset should not exceed the range of buffer");
 
-    return {__buf, sycl::range<1>(__n_acc), __offset};
+    return {__buf, sycl::range<1>(__n_acc), __offset, __properties};
 }
 } // namespace __internal
 
 //A SYCL range over SYCL buffer
-template <typename _T, sycl::access::mode _AccMode = sycl::access::mode::read,
+template <typename _T, sycl::access::mode _AccMode = sycl::access::mode::read, bool _NoInit = false,
           __dpl_sycl::__target _Target = __dpl_sycl::__target_device,
           sycl::access::placeholder _Placeholder = sycl::access::placeholder::true_t>
 class all_view
@@ -62,7 +62,8 @@ class all_view
     using value_type = _T;
 
     all_view(sycl::buffer<_T, 1> __buf = sycl::buffer<_T, 1>(0), __diff_type __offset = 0, __diff_type __n = 0)
-        : __m_acc(__internal::__create_accessor<__accessor_t>(__buf, __offset, __n))
+        : __m_acc(__internal::__create_accessor<__accessor_t>(
+              __buf, __offset, __n, _NoInit ? sycl::property_list{__dpl_sycl::__no_init{}} : sycl::property_list{}))
     {
     }
 
@@ -108,17 +109,17 @@ class all_view
     __accessor_t __m_acc;
 };
 
-template <sycl::access::mode AccMode = sycl::access::mode::read_write,
+template <sycl::access::mode AccMode = sycl::access::mode::read_write, bool _NoInit = false,
           __dpl_sycl::__target _Target = __dpl_sycl::__target_device,
           sycl::access::placeholder _Placeholder = sycl::access::placeholder::true_t>
 struct all_view_fn
 {
     template <typename _T>
-    constexpr oneapi::dpl::__ranges::all_view<_T, AccMode, _Target, _Placeholder>
+    constexpr oneapi::dpl::__ranges::all_view<_T, AccMode, _NoInit, _Target, _Placeholder>
     operator()(sycl::buffer<_T, 1> __buf, typename ::std::iterator_traits<_T*>::difference_type __offset = 0,
                typename ::std::iterator_traits<_T*>::difference_type __n = 0) const
     {
-        return oneapi::dpl::__ranges::all_view<_T, AccMode, _Target, _Placeholder>(__buf, __offset, __n);
+        return oneapi::dpl::__ranges::all_view<_T, AccMode, _NoInit, _Target, _Placeholder>(__buf, __offset, __n);
     }
 
     template <typename _R>
@@ -155,20 +156,22 @@ struct all_host_view_fn
 
 namespace views
 {
-inline constexpr all_view_fn<sycl::access::mode::read_write, __dpl_sycl::__target_device,
+inline constexpr all_view_fn<sycl::access::mode::read_write, /*_NoInit=*/false, __dpl_sycl::__target_device,
                              sycl::access::placeholder::true_t>
     all;
 
-inline constexpr all_view_fn<sycl::access::mode::read, __dpl_sycl::__target_device, sycl::access::placeholder::true_t>
+inline constexpr all_view_fn<sycl::access::mode::read, /*_NoInit=*/false, __dpl_sycl::__target_device,
+                             sycl::access::placeholder::true_t>
     all_read;
 
-inline constexpr all_view_fn<sycl::access::mode::write, __dpl_sycl::__target_device, sycl::access::placeholder::true_t>
+inline constexpr all_view_fn<sycl::access::mode::write, /*_NoInit=*/false, __dpl_sycl::__target_device,
+                             sycl::access::placeholder::true_t>
     all_write;
 
 #if _ONEDPL_SYCL2020_HOST_ACCESSOR_PRESENT
 inline constexpr all_host_view_fn
 #elif _ONEDPL_LIBSYCL_VERSION_LESS_THAN(60200)
-inline constexpr all_view_fn<sycl::access::mode::read_write, __dpl_sycl::__host_target,
+inline constexpr all_view_fn<sycl::access::mode::read_write, /*_NoInit=*/false, __dpl_sycl::__host_target,
                              sycl::access::placeholder::false_t>
 #else
 #    error "sycl::host_accessor is not supported, and no alternative is available"
@@ -210,6 +213,38 @@ struct is_sycl_iterator<oneapi::dpl::__internal::sycl_iterator<Mode, Types...>>
 {
 };
 
+template <sycl::access_mode _UserHint>
+struct __extract_user_hint_data_impl
+{
+    static constexpr sycl::access_mode __value = _UserHint;
+    static constexpr bool __no_init = false;
+};
+
+template <>
+struct __extract_user_hint_data_impl<sycl::access_mode::discard_write>
+{
+    static constexpr sycl::access_mode __value = sycl::access_mode::write;
+    static constexpr bool __no_init = true;
+};
+
+template <>
+struct __extract_user_hint_data_impl<sycl::access_mode::discard_read_write>
+{
+    static constexpr sycl::access_mode __value = sycl::access_mode::read_write;
+    static constexpr bool __no_init = true;
+};
+
+template <typename ForwardIterator>
+struct __extract_user_hint_data : public __extract_user_hint_data_impl<sycl::access_mode::read_write>
+{
+};
+
+template <sycl::access_mode _AccessMode, typename T, typename Allocator>
+struct __extract_user_hint_data<oneapi::dpl::__internal::sycl_iterator<_AccessMode, T, Allocator>>
+    : public __extract_user_hint_data_impl<_AccessMode>
+{
+};
+
 template <typename Iter, typename Void = void>
 struct is_hetero_legacy_trait : ::std::false_type
 {
@@ -286,9 +321,9 @@ __require_access(sycl::handler&)
 {
 }
 
-template <typename T, sycl::access::mode M>
+template <typename T, sycl::access::mode M, bool _NoInit>
 void
-__require_access_range(sycl::handler& __cgh, oneapi::dpl::__ranges::all_view<T, M>& sycl_view)
+__require_access_range(sycl::handler& __cgh, oneapi::dpl::__ranges::all_view<T, M, _NoInit>& sycl_view)
 {
     sycl_view.require_access(__cgh);
 }
@@ -406,7 +441,7 @@ struct __range_holder
     }
 };
 
-template <sycl::access::mode AccMode, bool _IsNoInitRequested = false>
+template <sycl::access::mode AccMode, bool _IsNoInitRequested = false, bool _DeferToUserHint = false>
 struct __get_sycl_range
 {
     __get_sycl_range()
@@ -424,6 +459,22 @@ struct __get_sycl_range
         _LocalAccMode == sycl::access::mode::read_write || _LocalAccMode == sycl::access::mode::write;
 
   private:
+    // When _DeferToUserHint is true and _Iter is a sycl_iterator, resolve the access mode and no_init
+    // from the iterator's embedded user hint. Otherwise, pass through the provided mode.
+    template <sycl::access::mode _Mode, bool _NoInit>
+    struct __passthrough_access_mode
+    {
+        static constexpr sycl::access::mode __value = _Mode;
+        static constexpr bool __no_init = _NoInit;
+    };
+
+    template <sycl::access::mode _LocalAccMode, bool _LocalNoInit, typename _Iter>
+    struct __resolve_access_mode
+        : std::conditional_t<_DeferToUserHint && is_sycl_iterator<_Iter>::value, __extract_user_hint_data<_Iter>,
+                             __passthrough_access_mode<_LocalAccMode, _LocalNoInit>>
+    {
+    };
+
     // We have to keep sycl buffer(s) instance here by sync reasons;
     std::vector<std::unique_ptr<oneapi::dpl::__internal::__lifetime_keeper_base>> m_buffers;
 
@@ -613,9 +664,15 @@ struct __get_sycl_range
     auto
     __process_input_iter(_Iter __first, _Iter __last)
         -> std::enable_if_t<oneapi::dpl::__ranges::is_hetero_iterator_v<_Iter>,
-                            __range_holder<oneapi::dpl::__ranges::all_view<val_t<_Iter>, _LocalAccMode>>>
+                            __range_holder<oneapi::dpl::__ranges::all_view<
+                                val_t<_Iter>, __resolve_access_mode<_LocalAccMode, _LocalNoInit, _Iter>::__value,
+                                __resolve_access_mode<_LocalAccMode, _LocalNoInit, _Iter>::__no_init>>>
     {
-        static_assert(!(_LocalAccMode == sycl::access::mode::read && _LocalNoInit),
+        constexpr sycl::access::mode __resolved_mode =
+            __resolve_access_mode<_LocalAccMode, _LocalNoInit, _Iter>::__value;
+        constexpr bool __resolved_no_init = __resolve_access_mode<_LocalAccMode, _LocalNoInit, _Iter>::__no_init;
+
+        static_assert(!(__resolved_mode == sycl::access::mode::read && __resolved_no_init),
                       "Read mode cannot be used with no_init property.");
         assert(__first < __last);
         using value_type = val_t<_Iter>;
@@ -627,23 +684,24 @@ struct __get_sycl_range
         //  __first is not guaranteed to be a sycl_iterator, it may be another type which sets the trait
         //   is_hetero = ::std::true_type. We use get_idx() to get the buffer offset, use get_buffer() to get the
         //   buffer and use those to create the range.
-        const auto __offset = __first.get_idx();
-        const auto __size = __dpl_sycl::__get_buffer_size(__first.get_buffer());
-        const auto __n = ::std::min(decltype(__size)(__last - __first), __size);
+        const std::size_t __offset = __first.get_idx();
+        const std::size_t __size = __dpl_sycl::__get_buffer_size(__first.get_buffer());
+        const std::size_t __n = std::min<std::size_t>(__last - __first, __size);
         assert(__offset + __n <= __size);
 
-        return __range_holder<oneapi::dpl::__ranges::all_view<value_type, _LocalAccMode>>{
-            oneapi::dpl::__ranges::all_view<value_type, _LocalAccMode>(__first.get_buffer() /* buffer */,
-                                                                       __offset /* offset*/, __n /* size*/)};
+        using _View = oneapi::dpl::__ranges::all_view<value_type, __resolved_mode, __resolved_no_init>;
+
+        return __range_holder<_View>{_View(__first.get_buffer(), __offset, __n)};
     }
 
     //SFINAE-overload for a contiguous host iterator
     template <sycl::access::mode _LocalAccMode, bool _LocalNoInit, typename _Iter>
     auto
     __process_input_iter(_Iter __first, _Iter __last)
-        -> ::std::enable_if_t<is_temp_buff<_Iter>::value && __is_addressable_v<_Iter> && !is_zip<_Iter>::value &&
-                                  !is_permutation<_Iter>::value,
-                              __range_holder<oneapi::dpl::__ranges::all_view<val_t<_Iter>, _LocalAccMode>>>
+        -> ::std::enable_if_t<
+            is_temp_buff<_Iter>::value && __is_addressable_v<_Iter> && !is_zip<_Iter>::value &&
+                !is_permutation<_Iter>::value,
+            __range_holder<oneapi::dpl::__ranges::all_view<val_t<_Iter>, _LocalAccMode, _LocalNoInit>>>
     {
         static_assert(!(_LocalAccMode == sycl::access::mode::read && _LocalNoInit),
                       "Read mode cannot be used with no_init property.");
@@ -672,9 +730,10 @@ struct __get_sycl_range
     template <sycl::access::mode _LocalAccMode, bool _LocalNoInit, typename _Iter>
     auto
     __process_input_iter(_Iter __first, _Iter __last)
-        -> ::std::enable_if_t<is_temp_buff<_Iter>::value && !__is_addressable_v<_Iter> && !is_zip<_Iter>::value &&
-                                  !is_permutation<_Iter>::value,
-                              __range_holder<oneapi::dpl::__ranges::all_view<val_t<_Iter>, _LocalAccMode>>>
+        -> ::std::enable_if_t<
+            is_temp_buff<_Iter>::value && !__is_addressable_v<_Iter> && !is_zip<_Iter>::value &&
+                !is_permutation<_Iter>::value,
+            __range_holder<oneapi::dpl::__ranges::all_view<val_t<_Iter>, _LocalAccMode, _LocalNoInit>>>
     {
         using _T = val_t<_Iter>;
 
@@ -718,8 +777,8 @@ struct __get_sycl_range
         m_buffers.push_back(std::make_unique<BufferType>(__buf));
 
         using _T = val_t<_Iter>;
-        return __range_holder<oneapi::dpl::__ranges::all_view<_T, _LocalAccMode>>{
-            oneapi::dpl::__ranges::all_view<_T, _LocalAccMode>(__buf)};
+        return __range_holder<oneapi::dpl::__ranges::all_view<_T, _LocalAccMode, _LocalNoInit>>{
+            oneapi::dpl::__ranges::all_view<_T, _LocalAccMode, _LocalNoInit>(__buf)};
     }
 
   public:
@@ -764,8 +823,10 @@ __select_backend(const execution::fpga_policy<_Factor, _KernelName>&, _Ranges&&.
 
 #if _ONEDPL_CPP20_RANGES_PRESENT
 //A specialization for enable_view to true because oneapi::dpl::__ranges::all_view models a view (see C++ standard)
-template <typename _T, sycl::access::mode _AccMode, sycl::target _Target, sycl::access::placeholder _Placeholder>
-inline constexpr bool std::ranges::enable_view<oneapi::dpl::__ranges::all_view<_T, _AccMode, _Target, _Placeholder>> = true;
+template <typename _T, sycl::access::mode _AccMode, bool _NoInit, sycl::target _Target,
+          sycl::access::placeholder _Placeholder>
+inline constexpr bool
+    std::ranges::enable_view<oneapi::dpl::__ranges::all_view<_T, _AccMode, _NoInit, _Target, _Placeholder>> = true;
 #endif
 
 #endif // _ONEDPL_UTILS_RANGES_SYCL_H