From 371b123513d72c38e7490d66e0ee477215d1d2e8 Mon Sep 17 00:00:00 2001
From: Dan Hoeflinger <dan.hoeflinger@intel.com>
Date: Fri, 19 Dec 2025 17:00:35 -0500
Subject: [PATCH 01/11] refactor away redundant fill functor; fix access modes
 for uninitialized apis remove unnecessary pattern

Signed-off-by: Dan Hoeflinger <dan.hoeflinger@intel.com>
---
 include/oneapi/dpl/pstl/algorithm_impl.h | 10 ++++++----
 1 file changed, 6 insertions(+), 4 deletions(-)

diff --git a/include/oneapi/dpl/pstl/algorithm_impl.h b/include/oneapi/dpl/pstl/algorithm_impl.h
index 20d7043f30f..06c5645502e 100644
--- a/include/oneapi/dpl/pstl/algorithm_impl.h
+++ b/include/oneapi/dpl/pstl/algorithm_impl.h
@@ -428,8 +428,8 @@ __pattern_walk2_brick_n(__parallel_tag<_IsVector>, _ExecutionPolicy&& __exec, _R
 template <class _Tag, class _ExecutionPolicy, class _ForwardIterator1, class _Size, class _ForwardIterator2,
           class _Brick>
 _ForwardIterator2
-__pattern_walk2_brick_n(_Tag, _ExecutionPolicy&&, _ForwardIterator1 __first1, _Size __n,
-                        _ForwardIterator2 __first2, _Brick __brick) noexcept
+__pattern_walk2_brick_n(_Tag, _ExecutionPolicy&&, _ForwardIterator1 __first1, _Size __n, _ForwardIterator2 __first2,
+                        _Brick __brick) noexcept
 {
     static_assert(__is_serial_tag_v<_Tag> || __is_parallel_forward_tag_v<_Tag>);
 
@@ -3417,11 +3417,13 @@ __parallel_set_op(__parallel_tag<_IsVector>, _ExecutionPolicy&& __exec, _RandomA
 
                 //try searching for the first element which not equal to *__b
                 if (__b != __first1)
-                    __b += __internal::__pstl_upper_bound(__b, _DifferenceType1{0}, __last1 - __b, __b, __comp, __proj1, __proj1);
+                    __b += __internal::__pstl_upper_bound(__b, _DifferenceType1{0}, __last1 - __b, __b, __comp, __proj1,
+                                                          __proj1);
 
                 //try searching for the first element which not equal to *__e
                 if (__e != __last1)
-                    __e += __internal::__pstl_upper_bound(__e, _DifferenceType1{0}, __last1 - __e, __e, __comp, __proj1, __proj1);
+                    __e += __internal::__pstl_upper_bound(__e, _DifferenceType1{0}, __last1 - __e, __e, __comp, __proj1,
+                                                          __proj1);
 
                 //check is [__b; __e) empty
                 if (__e - __b < 1)

From a7ecfc28b5aadf343a3596c4e91f330ba7ef076b Mon Sep 17 00:00:00 2001
From: Dan Hoeflinger <dan.hoeflinger@intel.com>
Date: Fri, 19 Dec 2025 17:00:35 -0500
Subject: [PATCH 02/11] refactor away redundant fill functor; fix access modes
 for uninitialized apis remove unnecessary pattern

Signed-off-by: Dan Hoeflinger <dan.hoeflinger@intel.com>

From aeda78008211ebd1f3c0c90c20689510d3b80915 Mon Sep 17 00:00:00 2001
From: Dan Hoeflinger <dan.hoeflinger@intel.com>
Date: Mon, 29 Dec 2025 11:54:19 -0500
Subject: [PATCH 03/11] refactor sycl_iter resolver (remove make_iter_mode)

Signed-off-by: Dan Hoeflinger <dan.hoeflinger@intel.com>
---
 .../internal/async_impl/async_impl_hetero.h   |   7 +-
 .../dpl/pstl/hetero/algorithm_impl_hetero.h   |  28 +---
 .../pstl/hetero/dpcpp/parallel_backend_sycl.h | 130 ------------------
 .../dpl/pstl/hetero/dpcpp/utils_ranges_sycl.h | 105 +++++++++++++-
 .../get_sycl_range.pass.cpp                   |  44 ++++++
 5 files changed, 155 insertions(+), 159 deletions(-)

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 faac0b2ff1c..93018aab76a 100644
--- a/include/oneapi/dpl/internal/async_impl/async_impl_hetero.h
+++ b/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});
 }
 
 //------------------------------------------------------------------------
diff --git a/include/oneapi/dpl/pstl/hetero/algorithm_impl_hetero.h b/include/oneapi/dpl/pstl/hetero/algorithm_impl_hetero.h
index 4ee8acc5a30..ea57ff8609e 100644
--- a/include/oneapi/dpl/pstl/hetero/algorithm_impl_hetero.h
+++ b/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
 
@@ -1455,10 +1444,7 @@ __pattern_partial_sort(__hetero_tag<_BackendTag>, _ExecutionPolicy&& __exec, _It
         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)
+        _BackendTag{}, ::std::forward<_ExecutionPolicy>(__exec), __first, __mid, __last, __comp)
         .__checked_deferrable_wait();
 }
 
@@ -1552,9 +1538,7 @@ __pattern_partial_sort_copy(__hetero_tag<_BackendTag> __tag, _ExecutionPolicy&&
 
         __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);
+            __buf_first, __buf_mid, __buf_last, __comp);
 
         return __pattern_walk2(
             __tag, __par_backend_hetero::make_wrapped_policy<__copy_back>(::std::forward<_ExecutionPolicy>(__exec)),
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 dbb77588847..36dc1a86743 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 b3a7f887c25..1e2febc127c 100644
--- a/include/oneapi/dpl/pstl/hetero/dpcpp/utils_ranges_sycl.h
+++ b/include/oneapi/dpl/pstl/hetero/dpcpp/utils_ranges_sycl.h
@@ -210,6 +210,78 @@ struct is_sycl_iterator<oneapi::dpl::__internal::sycl_iterator<Mode, Types...>>
 {
 };
 
+// Trait to check if mode is a "discard" mode (deprecated SYCL 2020 modes)
+template <sycl::access_mode _Mode>
+inline constexpr bool __is_discard_mode_v =
+    _Mode == sycl::access_mode::discard_write || _Mode == sycl::access_mode::discard_read_write;
+
+//-----------------------------------------------------------------------------
+// iter_mode_resolver
+//-----------------------------------------------------------------------------
+// Resolves conflicts between user-specified access mode (inMode, from sycl_iterator)
+// and algorithm-required access mode (outMode).
+// In general case, inMode must equal outMode.
+// Specializations handle specific compatible combinations.
+// Also tracks whether no_init semantics should be applied (from discard_* modes).
+
+template <sycl::access_mode inMode, sycl::access_mode outMode, bool _LocalNoInit = false>
+struct iter_mode_resolver
+{
+    static_assert(inMode == outMode,
+                  "Access mode provided by user conflicts with the one required by the algorithm");
+    static constexpr sycl::access_mode value = inMode;
+    static constexpr bool no_init = _LocalNoInit;
+};
+
+// NOTE: read + read_write and write + read_write are NOT compatible.
+// If algorithm needs read_write, a read-only or write-only iterator cannot satisfy it.
+
+// 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;
+};
+
+// discard_write can satisfy write requirement with no_init
+template <>
+struct iter_mode_resolver<sycl::access_mode::discard_write, sycl::access_mode::write, true>
+{
+    static constexpr sycl::access_mode value = sycl::access_mode::write;
+    static constexpr bool no_init = true;
+};
+
+// discard_read_write can satisfy any requirement with no_init
+template <sycl::access_mode outMode>
+struct iter_mode_resolver<sycl::access_mode::discard_read_write, outMode, true>
+{
+    static constexpr sycl::access_mode value = outMode;
+    static constexpr bool no_init = true;
+};
+
+// NOTE: discard_read_write + read_write is NOT compatible.
+// User says "don't copy original data" but algorithm needs to read it!
+
+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;
+
+// Checks if iterator mode is compatible with algorithm mode.
+// Returns false if user specified a discard mode (no_init) but algorithm requires initialization.
+template <sycl::access_mode _IterMode, sycl::access_mode _AlgoMode, bool _AlgoNoInit>
+inline constexpr bool __is_iter_mode_compatible_v = _AlgoNoInit || !iter_mode_resolver_no_init_v<_IterMode, _AlgoMode, _AlgoNoInit>;
+
 template <typename Iter, typename Void = void>
 struct is_hetero_legacy_trait : ::std::false_type
 {
@@ -608,11 +680,40 @@ 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>
+    auto
+    __process_input_iter(_Iter __first, _Iter __last)
+        -> std::enable_if_t<
+            is_sycl_iterator<_Iter>::value,
+            __range_holder<oneapi::dpl::__ranges::all_view<val_t<_Iter>,
+                                                           iter_mode_resolver_v<_Iter::mode, _LocalAccMode, _LocalNoInit>>>>
+    {
+        assert(__first < __last);
+        using value_type = val_t<_Iter>;
+
+        // Resolve the access mode: iterator's embedded mode vs. algorithm's required mode
+        static constexpr sycl::access_mode _ResolvedMode = iter_mode_resolver_v<_Iter::mode, _LocalAccMode, _LocalNoInit>;
+
+        // If user specified a discard mode (no_init), algorithm must also allow no_init
+        static_assert(__is_iter_mode_compatible_v<_Iter::mode, _LocalAccMode, _LocalNoInit>,
+                      "User specified discard mode (no_init) but algorithm requires initialization");
+
+        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);
+        assert(__offset + __n <= __size);
+
+        return __range_holder<oneapi::dpl::__ranges::all_view<value_type, _ResolvedMode>>{
+            oneapi::dpl::__ranges::all_view<value_type, _ResolvedMode>(__first.get_buffer() /* buffer */,
+                                                                       __offset /* offset*/, __n /* size*/)};
+    }
+
+    //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>,
+        -> 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>>>
     {
         static_assert(!(_LocalAccMode == sycl::access::mode::read && _LocalNoInit),
diff --git a/test/general/implementation_details/get_sycl_range.pass.cpp b/test/general/implementation_details/get_sycl_range.pass.cpp
index 35d5d85e473..76114b395c4 100644
--- a/test/general/implementation_details/get_sycl_range.pass.cpp
+++ b/test/general/implementation_details/get_sycl_range.pass.cpp
@@ -146,6 +146,49 @@ test_default_template_parameter()
                   "default _NoInit should be false");
 }
 
+void
+test_is_iter_mode_compatible_v()
+{
+    using oneapi::dpl::__ranges::__is_iter_mode_compatible_v;
+    namespace sa = sycl::access;
+
+    // Standard modes are always compatible regardless of _AlgoNoInit
+    static_assert(__is_iter_mode_compatible_v<sa::mode::read, sa::mode::read, false>,
+                  "read/read should be compatible");
+    static_assert(__is_iter_mode_compatible_v<sa::mode::read, sa::mode::read, true>,
+                  "read/read should be compatible with no_init");
+    static_assert(__is_iter_mode_compatible_v<sa::mode::write, sa::mode::write, false>,
+                  "write/write should be compatible");
+    static_assert(__is_iter_mode_compatible_v<sa::mode::write, sa::mode::write, true>,
+                  "write/write should be compatible with no_init");
+    static_assert(__is_iter_mode_compatible_v<sa::mode::read_write, sa::mode::read_write, false>,
+                  "read_write/read_write should be compatible");
+    static_assert(__is_iter_mode_compatible_v<sa::mode::read_write, sa::mode::read_write, true>,
+                  "read_write/read_write should be compatible with no_init");
+
+    // read_write iterator can be downgraded to satisfy read or write requirements
+    static_assert(__is_iter_mode_compatible_v<sa::mode::read_write, sa::mode::read, false>,
+                  "read_write can satisfy read");
+    static_assert(__is_iter_mode_compatible_v<sa::mode::read_write, sa::mode::write, false>,
+                  "read_write can satisfy write");
+
+    // Discard modes are compatible with write when algorithm allows no_init
+    static_assert(__is_iter_mode_compatible_v<sa::mode::discard_write, sa::mode::write, true>,
+                  "discard_write compatible with write when algo allows no_init");
+    static_assert(__is_iter_mode_compatible_v<sa::mode::discard_read_write, sa::mode::write, true>,
+                  "discard_read_write compatible with write when algo allows no_init");
+
+    // Discard modes are NOT compatible when algorithm requires initialization
+    //TODO: Fix the trait to enable testing of this but keep good error message for user
+    //static_assert(!__is_iter_mode_compatible_v<sa::mode::discard_write, sa::mode::write, false>,
+    //              "discard_write incompatible with write when algo requires init");
+    //static_assert(!__is_iter_mode_compatible_v<sa::mode::discard_read_write, sa::mode::write, false>,
+    //              "discard_read_write incompatible with write when algo requires init");
+
+    // NOTE: discard_* + read_write would fail at iter_mode_resolver level (static_assert)
+    // because user says "don't copy data" but algorithm needs to read it
+}
+
 #endif // TEST_DPCPP_BACKEND_PRESENT
 
 int
@@ -156,6 +199,7 @@ main()
     test_is_copy_back_v();
     test_traits_use_local_parameters();
     test_default_template_parameter();
+    test_is_iter_mode_compatible_v();
 #endif
 
     return TestUtils::done(TEST_DPCPP_BACKEND_PRESENT);

From 6f4da44bad4a91cc9d5bdd9df6545860156e012b Mon Sep 17 00:00:00 2001
From: Dan Hoeflinger <dan.hoeflinger@intel.com>
Date: Mon, 29 Dec 2025 12:07:41 -0500
Subject: [PATCH 04/11] fix trait (maintain good error msg)

Signed-off-by: Dan Hoeflinger <dan.hoeflinger@intel.com>
---
 .../dpl/pstl/hetero/dpcpp/utils_ranges_sycl.h | 43 +++++-----
 .../get_sycl_range.pass.cpp                   | 78 +++++++++++--------
 2 files changed, 67 insertions(+), 54 deletions(-)

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 1e2febc127c..ecfba9bebcf 100644
--- a/include/oneapi/dpl/pstl/hetero/dpcpp/utils_ranges_sycl.h
+++ b/include/oneapi/dpl/pstl/hetero/dpcpp/utils_ranges_sycl.h
@@ -220,22 +220,19 @@ inline constexpr bool __is_discard_mode_v =
 //-----------------------------------------------------------------------------
 // Resolves conflicts between user-specified access mode (inMode, from sycl_iterator)
 // and algorithm-required access mode (outMode).
-// In general case, inMode must equal outMode.
-// Specializations handle specific compatible combinations.
-// Also tracks whether no_init semantics should be applied (from discard_* modes).
 
+// Primary template - undefined (invalid combinations fail to compile)
 template <sycl::access_mode inMode, sycl::access_mode outMode, bool _LocalNoInit = false>
-struct iter_mode_resolver
+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_assert(inMode == outMode,
-                  "Access mode provided by user conflicts with the one required by the algorithm");
-    static constexpr sycl::access_mode value = inMode;
+    static constexpr sycl::access_mode value = _Mode;
     static constexpr bool no_init = _LocalNoInit;
 };
 
-// NOTE: read + read_write and write + read_write are NOT compatible.
-// If algorithm needs read_write, a read-only or write-only iterator cannot satisfy it.
-
 // 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>
@@ -268,20 +265,24 @@ struct iter_mode_resolver<sycl::access_mode::discard_read_write, outMode, true>
     static constexpr bool no_init = true;
 };
 
-// NOTE: discard_read_write + read_write is NOT compatible.
-// User says "don't copy original data" but algorithm needs to read it!
+// 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;
+
+// Access helpers (only valid when resolvable)
 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;
 
-// Checks if iterator mode is compatible with algorithm mode.
-// Returns false if user specified a discard mode (no_init) but algorithm requires initialization.
-template <sycl::access_mode _IterMode, sycl::access_mode _AlgoMode, bool _AlgoNoInit>
-inline constexpr bool __is_iter_mode_compatible_v = _AlgoNoInit || !iter_mode_resolver_no_init_v<_IterMode, _AlgoMode, _AlgoNoInit>;
-
 template <typename Iter, typename Void = void>
 struct is_hetero_legacy_trait : ::std::false_type
 {
@@ -689,16 +690,16 @@ struct __get_sycl_range
             __range_holder<oneapi::dpl::__ranges::all_view<val_t<_Iter>,
                                                            iter_mode_resolver_v<_Iter::mode, _LocalAccMode, _LocalNoInit>>>>
     {
+        // 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: iterator's embedded mode vs. algorithm's required mode
         static constexpr sycl::access_mode _ResolvedMode = iter_mode_resolver_v<_Iter::mode, _LocalAccMode, _LocalNoInit>;
 
-        // If user specified a discard mode (no_init), algorithm must also allow no_init
-        static_assert(__is_iter_mode_compatible_v<_Iter::mode, _LocalAccMode, _LocalNoInit>,
-                      "User specified discard mode (no_init) but algorithm requires initialization");
-
         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);
diff --git a/test/general/implementation_details/get_sycl_range.pass.cpp b/test/general/implementation_details/get_sycl_range.pass.cpp
index 76114b395c4..b46af6c751e 100644
--- a/test/general/implementation_details/get_sycl_range.pass.cpp
+++ b/test/general/implementation_details/get_sycl_range.pass.cpp
@@ -147,46 +147,58 @@ test_default_template_parameter()
 }
 
 void
-test_is_iter_mode_compatible_v()
+test_is_iter_mode_resolvable_v()
 {
-    using oneapi::dpl::__ranges::__is_iter_mode_compatible_v;
+    using oneapi::dpl::__ranges::__is_iter_mode_resolvable_v;
     namespace sa = sycl::access;
 
-    // Standard modes are always compatible regardless of _AlgoNoInit
-    static_assert(__is_iter_mode_compatible_v<sa::mode::read, sa::mode::read, false>,
-                  "read/read should be compatible");
-    static_assert(__is_iter_mode_compatible_v<sa::mode::read, sa::mode::read, true>,
-                  "read/read should be compatible with no_init");
-    static_assert(__is_iter_mode_compatible_v<sa::mode::write, sa::mode::write, false>,
-                  "write/write should be compatible");
-    static_assert(__is_iter_mode_compatible_v<sa::mode::write, sa::mode::write, true>,
-                  "write/write should be compatible with no_init");
-    static_assert(__is_iter_mode_compatible_v<sa::mode::read_write, sa::mode::read_write, false>,
-                  "read_write/read_write should be compatible");
-    static_assert(__is_iter_mode_compatible_v<sa::mode::read_write, sa::mode::read_write, true>,
-                  "read_write/read_write should be compatible with no_init");
+    // Standard modes are always resolvable regardless of noInit
+    static_assert(__is_iter_mode_resolvable_v<sa::mode::read, sa::mode::read, false>,
+                  "read/read should be resolvable");
+    static_assert(__is_iter_mode_resolvable_v<sa::mode::read, sa::mode::read, true>,
+                  "read/read should be resolvable with no_init");
+    static_assert(__is_iter_mode_resolvable_v<sa::mode::write, sa::mode::write, false>,
+                  "write/write should be resolvable");
+    static_assert(__is_iter_mode_resolvable_v<sa::mode::write, sa::mode::write, true>,
+                  "write/write should be resolvable with no_init");
+    static_assert(__is_iter_mode_resolvable_v<sa::mode::read_write, sa::mode::read_write, false>,
+                  "read_write/read_write should be resolvable");
+    static_assert(__is_iter_mode_resolvable_v<sa::mode::read_write, sa::mode::read_write, true>,
+                  "read_write/read_write should be resolvable with no_init");
 
     // read_write iterator can be downgraded to satisfy read or write requirements
-    static_assert(__is_iter_mode_compatible_v<sa::mode::read_write, sa::mode::read, false>,
+    static_assert(__is_iter_mode_resolvable_v<sa::mode::read_write, sa::mode::read, false>,
                   "read_write can satisfy read");
-    static_assert(__is_iter_mode_compatible_v<sa::mode::read_write, sa::mode::write, false>,
+    static_assert(__is_iter_mode_resolvable_v<sa::mode::read_write, sa::mode::write, false>,
                   "read_write can satisfy write");
 
-    // Discard modes are compatible with write when algorithm allows no_init
-    static_assert(__is_iter_mode_compatible_v<sa::mode::discard_write, sa::mode::write, true>,
-                  "discard_write compatible with write when algo allows no_init");
-    static_assert(__is_iter_mode_compatible_v<sa::mode::discard_read_write, sa::mode::write, true>,
-                  "discard_read_write compatible with write when algo allows no_init");
-
-    // Discard modes are NOT compatible when algorithm requires initialization
-    //TODO: Fix the trait to enable testing of this but keep good error message for user
-    //static_assert(!__is_iter_mode_compatible_v<sa::mode::discard_write, sa::mode::write, false>,
-    //              "discard_write incompatible with write when algo requires init");
-    //static_assert(!__is_iter_mode_compatible_v<sa::mode::discard_read_write, sa::mode::write, false>,
-    //              "discard_read_write incompatible with write when algo requires init");
-
-    // NOTE: discard_* + read_write would fail at iter_mode_resolver level (static_assert)
-    // because user says "don't copy data" but algorithm needs to read it
+    // 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");
+    static_assert(!__is_iter_mode_resolvable_v<sa::mode::read, sa::mode::read_write, false>,
+                  "read cannot satisfy read_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");
+    static_assert(!__is_iter_mode_resolvable_v<sa::mode::write, sa::mode::read_write, false>,
+                  "write cannot satisfy read_write");
+
+    // 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>,
+                  "discard_write resolvable with write when algo allows no_init");
+    static_assert(__is_iter_mode_resolvable_v<sa::mode::discard_read_write, sa::mode::write, true>,
+                  "discard_read_write resolvable with write when algo allows no_init");
+    static_assert(__is_iter_mode_resolvable_v<sa::mode::discard_read_write, sa::mode::read_write, true>,
+                  "discard_read_write resolvable with read_write when algo allows no_init");
+
+    // Discard modes are NOT resolvable when algorithm requires initialization (noInit=false)
+    static_assert(!__is_iter_mode_resolvable_v<sa::mode::discard_write, sa::mode::write, false>,
+                  "discard_write not resolvable with write when algo requires init");
+    static_assert(!__is_iter_mode_resolvable_v<sa::mode::discard_read_write, sa::mode::write, false>,
+                  "discard_read_write not resolvable with write when algo requires init");
+    static_assert(!__is_iter_mode_resolvable_v<sa::mode::discard_read_write, sa::mode::read_write, false>,
+                  "discard_read_write not resolvable with read_write when algo requires init");
 }
 
 #endif // TEST_DPCPP_BACKEND_PRESENT
@@ -199,7 +211,7 @@ main()
     test_is_copy_back_v();
     test_traits_use_local_parameters();
     test_default_template_parameter();
-    test_is_iter_mode_compatible_v();
+    test_is_iter_mode_resolvable_v();
 #endif
 
     return TestUtils::done(TEST_DPCPP_BACKEND_PRESENT);

From 38c18e1c5971e900cf709e846ecf3aee01355699 Mon Sep 17 00:00:00 2001
From: Dan Hoeflinger <dan.hoeflinger@intel.com>
Date: Wed, 31 Dec 2025 16:48:11 -0500
Subject: [PATCH 05/11] remove unused trait, rework comments

Signed-off-by: Dan Hoeflinger <dan.hoeflinger@intel.com>
---
 include/oneapi/dpl/pstl/hetero/dpcpp/utils_ranges_sycl.h | 9 ---------
 1 file changed, 9 deletions(-)

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 ecfba9bebcf..04263e7326b 100644
--- a/include/oneapi/dpl/pstl/hetero/dpcpp/utils_ranges_sycl.h
+++ b/include/oneapi/dpl/pstl/hetero/dpcpp/utils_ranges_sycl.h
@@ -210,17 +210,8 @@ struct is_sycl_iterator<oneapi::dpl::__internal::sycl_iterator<Mode, Types...>>
 {
 };
 
-// Trait to check if mode is a "discard" mode (deprecated SYCL 2020 modes)
-template <sycl::access_mode _Mode>
-inline constexpr bool __is_discard_mode_v =
-    _Mode == sycl::access_mode::discard_write || _Mode == sycl::access_mode::discard_read_write;
-
-//-----------------------------------------------------------------------------
-// iter_mode_resolver
-//-----------------------------------------------------------------------------
 // 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;

From 35b5791beb9994246fcd118c62f80ddbdc4d78bc Mon Sep 17 00:00:00 2001
From: Dan Hoeflinger <dan.hoeflinger@intel.com>
Date: Wed, 31 Dec 2025 16:50:07 -0500
Subject: [PATCH 06/11] connecting no_init to accessors

Signed-off-by: Dan Hoeflinger <dan.hoeflinger@intel.com>
---
 .../kt/internal/esimd_radix_sort_utils.h      |  4 +-
 .../dpl/pstl/hetero/dpcpp/utils_ranges_sycl.h | 82 ++++++++++---------
 2 files changed, 46 insertions(+), 40 deletions(-)

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 027541ff761..b45acb75398 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
@@ -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();
 }
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 04263e7326b..ce5e88a70df 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,7 @@ 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 +108,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 +155,20 @@ 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"
@@ -214,11 +214,11 @@ struct is_sycl_iterator<oneapi::dpl::__internal::sycl_iterator<Mode, Types...>>
 // 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;
+struct __iter_mode_resolver;
 
 // Partial specialization when modes match exactly
 template <sycl::access_mode _Mode, bool _LocalNoInit>
-struct iter_mode_resolver<_Mode, _Mode, _LocalNoInit>
+struct __iter_mode_resolver<_Mode, _Mode, _LocalNoInit>
 {
     static constexpr sycl::access_mode value = _Mode;
     static constexpr bool no_init = _LocalNoInit;
@@ -226,7 +226,7 @@ struct iter_mode_resolver<_Mode, _Mode, _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>
+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;
@@ -234,7 +234,7 @@ struct iter_mode_resolver<sycl::access_mode::read_write, sycl::access_mode::read
 
 // 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>
+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;
@@ -242,7 +242,7 @@ struct iter_mode_resolver<sycl::access_mode::read_write, sycl::access_mode::writ
 
 // discard_write can satisfy write requirement with no_init
 template <>
-struct iter_mode_resolver<sycl::access_mode::discard_write, sycl::access_mode::write, true>
+struct __iter_mode_resolver<sycl::access_mode::discard_write, sycl::access_mode::write, true>
 {
     static constexpr sycl::access_mode value = sycl::access_mode::write;
     static constexpr bool no_init = true;
@@ -250,7 +250,7 @@ struct iter_mode_resolver<sycl::access_mode::discard_write, sycl::access_mode::w
 
 // discard_read_write can satisfy any requirement with no_init
 template <sycl::access_mode outMode>
-struct iter_mode_resolver<sycl::access_mode::discard_read_write, outMode, true>
+struct __iter_mode_resolver<sycl::access_mode::discard_read_write, outMode, true>
 {
     static constexpr sycl::access_mode value = outMode;
     static constexpr bool no_init = true;
@@ -262,17 +262,17 @@ 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 {};
+    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;
 
 // Access helpers (only valid when resolvable)
 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;
+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;
+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
@@ -350,9 +350,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);
 }
@@ -678,8 +678,10 @@ struct __get_sycl_range
     __process_input_iter(_Iter __first, _Iter __last)
         -> std::enable_if_t<
             is_sycl_iterator<_Iter>::value,
-            __range_holder<oneapi::dpl::__ranges::all_view<val_t<_Iter>,
-                                                           iter_mode_resolver_v<_Iter::mode, _LocalAccMode, _LocalNoInit>>>>
+            __range_holder<oneapi::dpl::__ranges::all_view<
+                val_t<_Iter>, __iter_mode_resolver_v<_Iter::mode, _LocalAccMode, _LocalNoInit>,
+                __iter_mode_resolver_no_init_v<_Iter::mode, _LocalAccMode, _LocalNoInit>,
+                __dpl_sycl::__target_device, sycl::access::placeholder::true_t>>>
     {
         // Check mode compatibility with a clear error message
         static_assert(__is_iter_mode_resolvable_v<_Iter::mode, _LocalAccMode, _LocalNoInit>,
@@ -688,17 +690,20 @@ struct __get_sycl_range
         assert(__first < __last);
         using value_type = val_t<_Iter>;
 
-        // Resolve the access mode: iterator's embedded mode vs. algorithm's required mode
-        static constexpr sycl::access_mode _ResolvedMode = iter_mode_resolver_v<_Iter::mode, _LocalAccMode, _LocalNoInit>;
+        // 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 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);
         assert(__offset + __n <= __size);
 
-        return __range_holder<oneapi::dpl::__ranges::all_view<value_type, _ResolvedMode>>{
-            oneapi::dpl::__ranges::all_view<value_type, _ResolvedMode>(__first.get_buffer() /* buffer */,
-                                                                       __offset /* offset*/, __n /* size*/)};
+        return __range_holder<oneapi::dpl::__ranges::all_view<value_type, _ResolvedMode, _ResolvedNoInit, __dpl_sycl::__target_device,
+                                                               sycl::access::placeholder::true_t>>{
+            oneapi::dpl::__ranges::all_view<value_type, _ResolvedMode, _ResolvedNoInit, __dpl_sycl::__target_device,
+                                            sycl::access::placeholder::true_t>(
+                __first.get_buffer() /* buffer */, __offset /* offset*/, __n /* size*/)};
     }
 
     //specialization for other hetero iterators (non-sycl_iterator that sets is_hetero trait)
@@ -706,7 +711,7 @@ struct __get_sycl_range
     auto
     __process_input_iter(_Iter __first, _Iter __last)
         -> 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>>>
+                            __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.");
@@ -725,8 +730,8 @@ struct __get_sycl_range
         const auto __n = ::std::min(decltype(__size)(__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 */,
+        return __range_holder<oneapi::dpl::__ranges::all_view<value_type, _LocalAccMode, _LocalNoInit>>{
+            oneapi::dpl::__ranges::all_view<value_type, _LocalAccMode, _LocalNoInit>(__first.get_buffer() /* buffer */,
                                                                        __offset /* offset*/, __n /* size*/)};
     }
 
@@ -736,7 +741,7 @@ struct __get_sycl_range
     __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>>>
+                              __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.");
@@ -767,7 +772,7 @@ struct __get_sycl_range
     __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>>>
+                              __range_holder<oneapi::dpl::__ranges::all_view<val_t<_Iter>, _LocalAccMode, _LocalNoInit>>>
     {
         using _T = val_t<_Iter>;
 
@@ -811,8 +816,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:
@@ -857,8 +862,9 @@ __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

From 7cfdcc63cf46f835671ce060c37429858bd126e0 Mon Sep 17 00:00:00 2001
From: Dan Hoeflinger <dan.hoeflinger@intel.com>
Date: Wed, 31 Dec 2025 16:59:41 -0500
Subject: [PATCH 07/11] formatting changes

Signed-off-by: Dan Hoeflinger <dan.hoeflinger@intel.com>
---
 .../dpl/pstl/hetero/algorithm_impl_hetero.h   | 10 +--
 .../dpl/pstl/hetero/dpcpp/utils_ranges_sycl.h | 75 +++++++++++--------
 .../get_sycl_range.pass.cpp                   | 15 ++--
 3 files changed, 53 insertions(+), 47 deletions(-)

diff --git a/include/oneapi/dpl/pstl/hetero/algorithm_impl_hetero.h b/include/oneapi/dpl/pstl/hetero/algorithm_impl_hetero.h
index ea57ff8609e..e8e676bebd2 100644
--- a/include/oneapi/dpl/pstl/hetero/algorithm_impl_hetero.h
+++ b/include/oneapi/dpl/pstl/hetero/algorithm_impl_hetero.h
@@ -326,7 +326,7 @@ _ForwardIterator
 __pattern_generate(__hetero_tag<_BackendTag> __tag, _ExecutionPolicy&& __exec, _ForwardIterator __first,
                    _ForwardIterator __last, _Generator __g)
 {
-    __pattern_walk1(__tag, ::std::forward<_ExecutionPolicy>(__exec), __first, __last,
+    __pattern_walk1(__tag, ::std::forward<_ExecutionPolicy>(__exec), __first,__last,
                     generate_functor<_Generator>{__g});
     return __last;
 }
@@ -1443,8 +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), __first, __mid, __last, __comp)
+    __par_backend_hetero::__parallel_partial_sort(_BackendTag{}, ::std::forward<_ExecutionPolicy>(__exec), __first,
+                                                  __mid, __last, __comp)
         .__checked_deferrable_wait();
 }
 
@@ -1537,8 +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),
-            __buf_first, __buf_mid, __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)),
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 ce5e88a70df..2b20bb48be3 100644
--- a/include/oneapi/dpl/pstl/hetero/dpcpp/utils_ranges_sycl.h
+++ b/include/oneapi/dpl/pstl/hetero/dpcpp/utils_ranges_sycl.h
@@ -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, _NoInit ? sycl::property_list{sycl::property::no_init{}} : sycl::property_list{}))
+        : __m_acc(__internal::__create_accessor<__accessor_t>(
+              __buf, __offset, __n, _NoInit ? sycl::property_list{sycl::property::no_init{}} : sycl::property_list{}))
     {
     }
 
@@ -155,20 +156,22 @@ struct all_host_view_fn
 
 namespace views
 {
-inline constexpr all_view_fn<sycl::access::mode::read_write, /*_NoInit=*/ false, __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, /*_NoInit=*/ false, __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, /*_NoInit=*/ false, __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, /*_NoInit=*/ false, __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"
@@ -258,11 +261,16 @@ struct __iter_mode_resolver<sycl::access_mode::discard_read_write, outMode, true
 
 // 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 {};
+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 {};
+                                 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;
@@ -676,12 +684,11 @@ 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_sycl_iterator<_Iter>::value,
-            __range_holder<oneapi::dpl::__ranges::all_view<
-                val_t<_Iter>, __iter_mode_resolver_v<_Iter::mode, _LocalAccMode, _LocalNoInit>,
-                __iter_mode_resolver_no_init_v<_Iter::mode, _LocalAccMode, _LocalNoInit>,
-                __dpl_sycl::__target_device, sycl::access::placeholder::true_t>>>
+        -> std::enable_if_t<is_sycl_iterator<_Iter>::value,
+                            __range_holder<oneapi::dpl::__ranges::all_view<
+                                val_t<_Iter>, __iter_mode_resolver_v<_Iter::mode, _LocalAccMode, _LocalNoInit>,
+                                __iter_mode_resolver_no_init_v<_Iter::mode, _LocalAccMode, _LocalNoInit>,
+                                __dpl_sycl::__target_device, sycl::access::placeholder::true_t>>>
     {
         // Check mode compatibility with a clear error message
         static_assert(__is_iter_mode_resolvable_v<_Iter::mode, _LocalAccMode, _LocalNoInit>,
@@ -691,19 +698,22 @@ struct __get_sycl_range
         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>;
+        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 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);
         assert(__offset + __n <= __size);
 
-        return __range_holder<oneapi::dpl::__ranges::all_view<value_type, _ResolvedMode, _ResolvedNoInit, __dpl_sycl::__target_device,
-                                                               sycl::access::placeholder::true_t>>{
+        return __range_holder<
             oneapi::dpl::__ranges::all_view<value_type, _ResolvedMode, _ResolvedNoInit, __dpl_sycl::__target_device,
-                                            sycl::access::placeholder::true_t>(
-                __first.get_buffer() /* buffer */, __offset /* offset*/, __n /* size*/)};
+                                            sycl::access::placeholder::true_t>>{
+            oneapi::dpl::__ranges::all_view<value_type, _ResolvedMode, _ResolvedNoInit, __dpl_sycl::__target_device,
+                                            sycl::access::placeholder::true_t>(__first.get_buffer() /* buffer */,
+                                                                               __offset /* offset*/, __n /* size*/)};
     }
 
     //specialization for other hetero iterators (non-sycl_iterator that sets is_hetero trait)
@@ -731,17 +741,17 @@ struct __get_sycl_range
         assert(__offset + __n <= __size);
 
         return __range_holder<oneapi::dpl::__ranges::all_view<value_type, _LocalAccMode, _LocalNoInit>>{
-            oneapi::dpl::__ranges::all_view<value_type, _LocalAccMode, _LocalNoInit>(__first.get_buffer() /* buffer */,
-                                                                       __offset /* offset*/, __n /* size*/)};
+            oneapi::dpl::__ranges::all_view<value_type, _LocalAccMode, _LocalNoInit>(
+                __first.get_buffer() /* buffer */, __offset /* offset*/, __n /* size*/)};
     }
 
     //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, _LocalNoInit>>>
+    __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.");
@@ -769,10 +779,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, _LocalNoInit>>>
+    __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>;
 
@@ -862,9 +872,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, 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;
+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 b46af6c751e..7336b2d153c 100644
--- a/test/general/implementation_details/get_sycl_range.pass.cpp
+++ b/test/general/implementation_details/get_sycl_range.pass.cpp
@@ -75,12 +75,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,
@@ -153,8 +151,7 @@ test_is_iter_mode_resolvable_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>,
-                  "read/read should be resolvable");
+    static_assert(__is_iter_mode_resolvable_v<sa::mode::read, sa::mode::read, false>, "read/read should be resolvable");
     static_assert(__is_iter_mode_resolvable_v<sa::mode::read, sa::mode::read, true>,
                   "read/read should be resolvable with no_init");
     static_assert(__is_iter_mode_resolvable_v<sa::mode::write, sa::mode::write, false>,
@@ -173,14 +170,12 @@ test_is_iter_mode_resolvable_v()
                   "read_write can satisfy write");
 
     // 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");
+    static_assert(!__is_iter_mode_resolvable_v<sa::mode::read, sa::mode::write, false>, "read cannot satisfy write");
     static_assert(!__is_iter_mode_resolvable_v<sa::mode::read, sa::mode::read_write, false>,
                   "read cannot satisfy read_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");
+    static_assert(!__is_iter_mode_resolvable_v<sa::mode::write, sa::mode::read, false>, "write cannot satisfy read");
     static_assert(!__is_iter_mode_resolvable_v<sa::mode::write, sa::mode::read_write, false>,
                   "write cannot satisfy read_write");
 

From 888045980ade4f2f73bd034bef2e335bde8e702d Mon Sep 17 00:00:00 2001
From: Dan Hoeflinger <dan.hoeflinger@intel.com>
Date: Fri, 2 Jan 2026 08:38:49 -0500
Subject: [PATCH 08/11] removing unnecessary comment

Signed-off-by: Dan Hoeflinger <dan.hoeflinger@intel.com>
---
 include/oneapi/dpl/pstl/hetero/dpcpp/utils_ranges_sycl.h | 1 -
 1 file changed, 1 deletion(-)

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 2b20bb48be3..99eeea04886 100644
--- a/include/oneapi/dpl/pstl/hetero/dpcpp/utils_ranges_sycl.h
+++ b/include/oneapi/dpl/pstl/hetero/dpcpp/utils_ranges_sycl.h
@@ -275,7 +275,6 @@ struct __is_iter_mode_resolvable<inMode, outMode, noInit,
 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;
 
-// Access helpers (only valid when resolvable)
 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;
 

From c0563c6fa0c7cb7bf24c483b44e77db2de9ee672 Mon Sep 17 00:00:00 2001
From: Dan Hoeflinger <dan.hoeflinger@intel.com>
Date: Fri, 2 Jan 2026 17:35:35 -0500
Subject: [PATCH 09/11] adding static assertion for read + no_init

Signed-off-by: Dan Hoeflinger <dan.hoeflinger@intel.com>
---
 include/oneapi/dpl/pstl/hetero/dpcpp/utils_ranges_sycl.h | 2 ++
 1 file changed, 2 insertions(+)

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 99eeea04886..883c04f613b 100644
--- a/include/oneapi/dpl/pstl/hetero/dpcpp/utils_ranges_sycl.h
+++ b/include/oneapi/dpl/pstl/hetero/dpcpp/utils_ranges_sycl.h
@@ -689,6 +689,8 @@ struct __get_sycl_range
                                 __iter_mode_resolver_no_init_v<_Iter::mode, _LocalAccMode, _LocalNoInit>,
                                 __dpl_sycl::__target_device, sycl::access::placeholder::true_t>>>
     {
+        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");

From 3684ca2b89a325b7d07b7442413f87c6c4f4a2a5 Mon Sep 17 00:00:00 2001
From: Dan Hoeflinger <dan.hoeflinger@intel.com>
Date: Tue, 13 Jan 2026 09:02:15 -0500
Subject: [PATCH 10/11] revert formatting only changes

Signed-off-by: Dan Hoeflinger <dan.hoeflinger@intel.com>
---
 include/oneapi/dpl/pstl/algorithm_impl.h | 6 ++----
 1 file changed, 2 insertions(+), 4 deletions(-)

diff --git a/include/oneapi/dpl/pstl/algorithm_impl.h b/include/oneapi/dpl/pstl/algorithm_impl.h
index 06c5645502e..81b55789f30 100644
--- a/include/oneapi/dpl/pstl/algorithm_impl.h
+++ b/include/oneapi/dpl/pstl/algorithm_impl.h
@@ -3417,13 +3417,11 @@ __parallel_set_op(__parallel_tag<_IsVector>, _ExecutionPolicy&& __exec, _RandomA
 
                 //try searching for the first element which not equal to *__b
                 if (__b != __first1)
-                    __b += __internal::__pstl_upper_bound(__b, _DifferenceType1{0}, __last1 - __b, __b, __comp, __proj1,
-                                                          __proj1);
+                    __b += __internal::__pstl_upper_bound(__b, _DifferenceType1{0}, __last1 - __b, __b, __comp, __proj1, __proj1);
 
                 //try searching for the first element which not equal to *__e
                 if (__e != __last1)
-                    __e += __internal::__pstl_upper_bound(__e, _DifferenceType1{0}, __last1 - __e, __e, __comp, __proj1,
-                                                          __proj1);
+                    __e += __internal::__pstl_upper_bound(__e, _DifferenceType1{0}, __last1 - __e, __e, __comp, __proj1, __proj1);
 
                 //check is [__b; __e) empty
                 if (__e - __b < 1)

From 267d03c01f4babab627d2f2a591d0b5475d0ed7f Mon Sep 17 00:00:00 2001
From: Dan Hoeflinger <dan.hoeflinger@intel.com>
Date: Tue, 13 Jan 2026 09:03:14 -0500
Subject: [PATCH 11/11] revert formatting only changes

Signed-off-by: Dan Hoeflinger <dan.hoeflinger@intel.com>
---
 include/oneapi/dpl/pstl/algorithm_impl.h | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/include/oneapi/dpl/pstl/algorithm_impl.h b/include/oneapi/dpl/pstl/algorithm_impl.h
index 81b55789f30..20d7043f30f 100644
--- a/include/oneapi/dpl/pstl/algorithm_impl.h
+++ b/include/oneapi/dpl/pstl/algorithm_impl.h
@@ -428,8 +428,8 @@ __pattern_walk2_brick_n(__parallel_tag<_IsVector>, _ExecutionPolicy&& __exec, _R
 template <class _Tag, class _ExecutionPolicy, class _ForwardIterator1, class _Size, class _ForwardIterator2,
           class _Brick>
 _ForwardIterator2
-__pattern_walk2_brick_n(_Tag, _ExecutionPolicy&&, _ForwardIterator1 __first1, _Size __n, _ForwardIterator2 __first2,
-                        _Brick __brick) noexcept
+__pattern_walk2_brick_n(_Tag, _ExecutionPolicy&&, _ForwardIterator1 __first1, _Size __n,
+                        _ForwardIterator2 __first2, _Brick __brick) noexcept
 {
     static_assert(__is_serial_tag_v<_Tag> || __is_parallel_forward_tag_v<_Tag>);