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..3b6d1a3bf04 100644
--- a/include/oneapi/dpl/internal/async_impl/async_impl_hetero.h
+++ b/include/oneapi/dpl/internal/async_impl/async_impl_hetero.h
@@ -169,9 +169,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/pstl/hetero/algorithm_impl_hetero.h b/include/oneapi/dpl/pstl/hetero/algorithm_impl_hetero.h
index 1422749848a..fbb72a56f61 100644
--- a/include/oneapi/dpl/pstl/hetero/algorithm_impl_hetero.h
+++ b/include/oneapi/dpl/pstl/hetero/algorithm_impl_hetero.h
@@ -302,9 +302,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 +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;
 }
 
@@ -919,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>();
@@ -1222,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
 
@@ -1468,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();
 }
 
@@ -1569,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>(
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..494080da56a 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, 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{sycl::property::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,92 @@ struct is_sycl_iterator<oneapi::dpl::__internal::sycl_iterator<Mode, Types...>>
 {
 };
 
+// Resolves conflicts between user-specified access mode (inMode, from sycl_iterator)
+// and algorithm-required access mode (outMode).
+// Primary template - undefined (invalid combinations fail to compile)
+template <sycl::access_mode inMode, sycl::access_mode outMode, bool _LocalNoInit = false>
+struct __iter_mode_resolver;
+
+// Partial specialization when modes match exactly
+template <sycl::access_mode _Mode, bool _LocalNoInit>
+struct __iter_mode_resolver<_Mode, _Mode, _LocalNoInit>
+{
+    static constexpr sycl::access_mode value = _Mode;
+    static constexpr bool no_init = _LocalNoInit;
+};
+
+// read_write can satisfy read requirement (downgrade to read)
+template <bool _LocalNoInit>
+struct __iter_mode_resolver<sycl::access_mode::read_write, sycl::access_mode::read, _LocalNoInit>
+{
+    static constexpr sycl::access_mode value = sycl::access_mode::read;
+    static constexpr bool no_init = _LocalNoInit;
+};
+
+// read_write can satisfy write requirement (downgrade to write)
+template <bool _LocalNoInit>
+struct __iter_mode_resolver<sycl::access_mode::read_write, sycl::access_mode::write, _LocalNoInit>
+{
+    static constexpr sycl::access_mode value = sycl::access_mode::write;
+    static constexpr bool no_init = _LocalNoInit;
+};
+
+// read can satisfy read_write requirement (demoteable with hints read_write to serve for_each use case)
+template <bool _LocalNoInit>
+struct __iter_mode_resolver<sycl::access_mode::read, sycl::access_mode::read_write, _LocalNoInit>
+{
+    static constexpr sycl::access_mode value = sycl::access_mode::read;
+    static constexpr bool no_init = _LocalNoInit;
+};
+
+// write can satisfy read_write requirement (demoteable with hints read_write to serve for_each use case)
+template <bool _LocalNoInit>
+struct __iter_mode_resolver<sycl::access_mode::write, sycl::access_mode::read_write, _LocalNoInit>
+{
+    static constexpr sycl::access_mode value = sycl::access_mode::write;
+    static constexpr bool no_init = _LocalNoInit;
+};
+
+// discard_write can satisfy write requirement (discard modes inherently imply no_init)
+// if algorithm requires initialization, we ignore user hints
+template <bool _LocalNoInit>
+struct __iter_mode_resolver<sycl::access_mode::discard_write, sycl::access_mode::write, _LocalNoInit>
+{
+    static constexpr sycl::access_mode value = sycl::access_mode::write;
+    static constexpr bool no_init = _LocalNoInit;
+};
+
+// discard_read_write can satisfy write or read_write requirement (discard modes inherently imply no_init)
+// if algorithm requires initialization, we ignore user hints
+template <sycl::access_mode outMode, bool _LocalNoInit>
+struct __iter_mode_resolver<sycl::access_mode::discard_read_write, outMode, _LocalNoInit>
+{
+    static constexpr sycl::access_mode value = outMode;
+    static constexpr bool no_init = _LocalNoInit;
+};
+
+// Helper to check if a mode combination is resolvable
+template <sycl::access_mode inMode, sycl::access_mode outMode, bool noInit, typename = void>
+struct __is_iter_mode_resolvable : std::false_type
+{
+};
+
+template <sycl::access_mode inMode, sycl::access_mode outMode, bool noInit>
+struct __is_iter_mode_resolvable<inMode, outMode, noInit,
+                                 std::void_t<decltype(__iter_mode_resolver<inMode, outMode, noInit>::value)>>
+    : std::true_type
+{
+};
+
+template <sycl::access_mode inMode, sycl::access_mode outMode, bool noInit>
+inline constexpr bool __is_iter_mode_resolvable_v = __is_iter_mode_resolvable<inMode, outMode, noInit>::value;
+
+template <sycl::access_mode inMode, sycl::access_mode outMode, bool noInit>
+inline constexpr sycl::access_mode __iter_mode_resolver_v = __iter_mode_resolver<inMode, outMode, noInit>::value;
+
+template <sycl::access_mode inMode, sycl::access_mode outMode, bool noInit>
+inline constexpr bool __iter_mode_resolver_no_init_v = __iter_mode_resolver<inMode, outMode, noInit>::no_init;
+
 template <typename Iter, typename Void = void>
 struct is_hetero_legacy_trait : ::std::false_type
 {
@@ -286,9 +375,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);
 }
@@ -608,12 +697,44 @@ struct __get_sycl_range
             oneapi::dpl::__ranges::guard_view<_Iter>{__first, __last - __first}};
     }
 
-    //specialization for hetero iterator
+    //specialization for sycl_iterator (applies access mode resolution)
+    template <sycl::access::mode _LocalAccMode, bool _LocalNoInit, typename _Iter,
+              std::enable_if_t<is_sycl_iterator<_Iter>::value, int> = 0>
+    auto
+    __process_input_iter(_Iter __first, _Iter __last)
+    {
+        static_assert(!(_LocalAccMode == sycl::access::mode::read && _LocalNoInit),
+                      "Read mode cannot be used with no_init property.");
+        // Check mode compatibility with a clear error message
+        static_assert(__is_iter_mode_resolvable_v<_Iter::mode, _LocalAccMode, _LocalNoInit>,
+                      "Access mode provided by user conflicts with the one required by the algorithm");
+
+        assert(__first < __last);
+        using value_type = val_t<_Iter>;
+
+        // Resolve the access mode and no_init: iterator's embedded mode vs. algorithm's required mode
+        static constexpr sycl::access_mode _ResolvedMode =
+            __iter_mode_resolver_v<_Iter::mode, _LocalAccMode, _LocalNoInit>;
+        static constexpr bool _ResolvedNoInit =
+            __iter_mode_resolver_no_init_v<_Iter::mode, _LocalAccMode, _LocalNoInit>;
+
+        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);
+
+        using _View = oneapi::dpl::__ranges::all_view<value_type, _ResolvedMode, _ResolvedNoInit,
+                                                      __dpl_sycl::__target_device, sycl::access::placeholder::true_t>;
+
+        return __range_holder<_View>{_View(__first.get_buffer(), __offset, __n)};
+    }
+
+    //specialization for other hetero iterators (non-sycl_iterator that sets is_hetero trait)
     template <sycl::access::mode _LocalAccMode, bool _LocalNoInit, typename _Iter>
     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>>>
+        -> std::enable_if_t<oneapi::dpl::__ranges::is_hetero_iterator_v<_Iter> && !is_sycl_iterator<_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.");
@@ -627,23 +748,23 @@ 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, _LocalAccMode, _LocalNoInit>;
+
+        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>>>
+    __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, _LocalNoInit>>>
     {
         static_assert(!(_LocalAccMode == sycl::access::mode::read && _LocalNoInit),
                       "Read mode cannot be used with no_init property.");
@@ -671,10 +792,10 @@ struct __get_sycl_range
     //SFINAE-overload for non-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>>>
+    __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, _LocalNoInit>>>
     {
         using _T = val_t<_Iter>;
 
@@ -718,8 +839,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 +885,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
diff --git a/test/general/implementation_details/get_sycl_range.pass.cpp b/test/general/implementation_details/get_sycl_range.pass.cpp
index c33c8680e7a..f8331f52930 100644
--- a/test/general/implementation_details/get_sycl_range.pass.cpp
+++ b/test/general/implementation_details/get_sycl_range.pass.cpp
@@ -76,12 +76,10 @@ test_is_copy_back_v()
     using read_write_mode = oneapi::dpl::__ranges::__get_sycl_range<sycl::access::mode::read_write>;
 
     // Test: read mode -> no copy back (false)
-    static_assert(read_mode::__is_copy_back_v<sycl::access::mode::read> == false,
-                  "read mode should not copy back");
+    static_assert(read_mode::__is_copy_back_v<sycl::access::mode::read> == false, "read mode should not copy back");
 
     // Test: write mode -> copy back (true)
-    static_assert(write_mode::__is_copy_back_v<sycl::access::mode::write> == true,
-                  "write mode should copy back");
+    static_assert(write_mode::__is_copy_back_v<sycl::access::mode::write> == true, "write mode should copy back");
 
     // Test: read_write mode -> copy back (true)
     static_assert(read_write_mode::__is_copy_back_v<sycl::access::mode::read_write> == true,
@@ -149,6 +147,105 @@ test_default_template_parameter()
                   "default _IsNoInitRequested should be false");
 }
 
+void
+test_is_iter_mode_resolvable_v()
+{
+    using oneapi::dpl::__ranges::__is_iter_mode_resolvable_v;
+    using oneapi::dpl::__ranges::__iter_mode_resolver_v;
+    using oneapi::dpl::__ranges::__iter_mode_resolver_no_init_v;
+    namespace sa = sycl::access;
+
+    // Standard modes are always resolvable regardless of noInit
+    static_assert(__is_iter_mode_resolvable_v<sa::mode::read, sa::mode::read, false> &&
+                  __iter_mode_resolver_v<sa::mode::read, sa::mode::read, false> == sa::mode::read &&
+                  __iter_mode_resolver_no_init_v<sa::mode::read, sa::mode::read, false> == false,
+                  "read_init/read_init should resolve to read_init");
+
+    static_assert(__is_iter_mode_resolvable_v<sa::mode::read, sa::mode::read, true> &&
+                  __iter_mode_resolver_v<sa::mode::read, sa::mode::read, true> == sa::mode::read &&
+                  __iter_mode_resolver_no_init_v<sa::mode::read, sa::mode::read, true> == true,
+                  "read_init/read_no_init should resolve to read_no_init");
+    
+    static_assert(__is_iter_mode_resolvable_v<sa::mode::write, sa::mode::write, false> &&
+                  __iter_mode_resolver_v<sa::mode::write, sa::mode::write, false> == sa::mode::write &&
+                  __iter_mode_resolver_no_init_v<sa::mode::write, sa::mode::write, false> == false,
+                  "write_init/write_init should resolve to write_init");
+
+    static_assert(__is_iter_mode_resolvable_v<sa::mode::write, sa::mode::write, true> &&
+                  __iter_mode_resolver_v<sa::mode::write, sa::mode::write, true> == sa::mode::write &&
+                  __iter_mode_resolver_no_init_v<sa::mode::write, sa::mode::write, true> == true,
+                  "write_init/write_no_init should resolve to write_no_init");
+    
+
+    static_assert(__is_iter_mode_resolvable_v<sa::mode::read_write, sa::mode::read_write, false> &&
+                  __iter_mode_resolver_v<sa::mode::read_write, sa::mode::read_write, false> == sa::mode::read_write &&
+                  __iter_mode_resolver_no_init_v<sa::mode::read_write, sa::mode::read_write, false> == false,
+                  "read_write_init/read_write_init should resolve to read_write_init");
+
+    static_assert(__is_iter_mode_resolvable_v<sa::mode::read_write, sa::mode::read_write, true> &&
+                  __iter_mode_resolver_v<sa::mode::read_write, sa::mode::read_write, true> == sa::mode::read_write &&
+                  __iter_mode_resolver_no_init_v<sa::mode::read_write, sa::mode::read_write, true> == true,
+                  "read_write_init/read_write_no_init should resolve to read_write_no_init");
+
+
+    // read_write iterator can be downgraded to satisfy read or write requirements
+    static_assert(__is_iter_mode_resolvable_v<sa::mode::read_write, sa::mode::read, false> &&
+                  __iter_mode_resolver_v<sa::mode::read_write, sa::mode::read, false> == sa::mode::read &&
+                  __iter_mode_resolver_no_init_v<sa::mode::read_write, sa::mode::read, false> == false,
+                  "read_write_init/read_init should resolve to read_init");
+
+    static_assert(__is_iter_mode_resolvable_v<sa::mode::read_write, sa::mode::write, false> &&
+                  __iter_mode_resolver_v<sa::mode::read_write, sa::mode::write, false> == sa::mode::write &&
+                  __iter_mode_resolver_no_init_v<sa::mode::read_write, sa::mode::write, false> == false,
+                  "read_write_init/write_init should resolve to write_init");
+
+    // Incompatible: read iterator cannot satisfy write or read_write
+    static_assert(!__is_iter_mode_resolvable_v<sa::mode::read, sa::mode::write, false>, "read cannot satisfy write");
+    
+    // Incompatible: write iterator cannot satisfy read or read_write
+    static_assert(!__is_iter_mode_resolvable_v<sa::mode::write, sa::mode::read, false>, "write cannot satisfy read");
+    
+    // demoteable read_write can be satisfied to serve for_each use case where read_write can be demoted 
+    // to read or write by user hints
+    // TODO: look at creating a specifically demotable access mode tag instead of using read_write for this purpose
+    static_assert(__is_iter_mode_resolvable_v<sa::mode::read, sa::mode::read_write, false> &&
+                  __iter_mode_resolver_v<sa::mode::read, sa::mode::read_write, false> == sa::mode::read &&
+                  __iter_mode_resolver_no_init_v<sa::mode::read, sa::mode::read_write, false> == false,
+                  "read_init / read_write_init should resolve to read_init to serve demoteable for_each");
+    static_assert(__is_iter_mode_resolvable_v<sa::mode::write, sa::mode::read_write, false> &&
+                  __iter_mode_resolver_v<sa::mode::write, sa::mode::read_write, false> == sa::mode::write &&
+                  __iter_mode_resolver_no_init_v<sa::mode::write, sa::mode::read_write, false> == false,
+                  "write_init / read_write_init should resolve to write_init to serve demoteable for_each");
+
+    // Discard modes are resolvable with write when algorithm allows no_init
+    static_assert(__is_iter_mode_resolvable_v<sa::mode::discard_write, sa::mode::write, true> &&
+                  __iter_mode_resolver_v<sa::mode::discard_write, sa::mode::write, true> == sa::mode::write &&
+                  __iter_mode_resolver_no_init_v<sa::mode::discard_write, sa::mode::write, true> == true,
+                "write_no_init / write_no_init should resolve to write_no_init");
+    static_assert(__is_iter_mode_resolvable_v<sa::mode::discard_read_write, sa::mode::write, true> && 
+                  __iter_mode_resolver_v<sa::mode::discard_read_write, sa::mode::write, true> == sa::mode::write &&
+                  __iter_mode_resolver_no_init_v<sa::mode::discard_read_write, sa::mode::write, true> == true,
+                  "read_write_no_init/ write_no_init should resolve to write_no_init");
+    static_assert(__is_iter_mode_resolvable_v<sa::mode::discard_read_write, sa::mode::read_write, true> && 
+                  __iter_mode_resolver_v<sa::mode::discard_read_write, sa::mode::read_write, true> == sa::mode::read_write &&
+                  __iter_mode_resolver_no_init_v<sa::mode::discard_read_write, sa::mode::read_write, true> == true,
+                  "read_write_no_init/ read_write_no_init should resolve to read_write_no_init");
+
+    // Discard modes are resolvable (and hints ignored) when algorithm requires initialization (noInit=false)
+    static_assert(__is_iter_mode_resolvable_v<sa::mode::discard_write, sa::mode::write, false> &&
+                  __iter_mode_resolver_v<sa::mode::discard_write, sa::mode::write, false> == sa::mode::write &&
+                  __iter_mode_resolver_no_init_v<sa::mode::discard_write, sa::mode::write, false> == false,
+                  "write_no_init/ write_init should resolve to write_init (ignoring hint)");
+    static_assert(__is_iter_mode_resolvable_v<sa::mode::discard_read_write, sa::mode::write, false> &&
+                  __iter_mode_resolver_v<sa::mode::discard_read_write, sa::mode::write, false> == sa::mode::write &&
+                  __iter_mode_resolver_no_init_v<sa::mode::discard_read_write, sa::mode::write, false> == false,
+                  "read_write_no_init/ write_init should resolve to write_init (ignoring hint)");
+    static_assert(__is_iter_mode_resolvable_v<sa::mode::discard_read_write, sa::mode::read_write, false> &&
+                  __iter_mode_resolver_v<sa::mode::discard_read_write, sa::mode::read_write, false> == sa::mode::read_write &&
+                  __iter_mode_resolver_no_init_v<sa::mode::discard_read_write, sa::mode::read_write, false> == false,
+                  "read_write_no_init/ read_write_init should resolve to read_write_init (ignoring hint)");
+}
+
 #endif // TEST_DPCPP_BACKEND_PRESENT
 
 int
@@ -159,6 +256,7 @@ main()
     test_is_copy_back_v();
     test_traits_use_local_parameters();
     test_default_template_parameter();
+    test_is_iter_mode_resolvable_v();
 #endif
 
     return TestUtils::done(TEST_DPCPP_BACKEND_PRESENT);