Skip to content

Conversation

@danhoeflinger
Copy link
Contributor

@danhoeflinger danhoeflinger commented Nov 18, 2025

Align __get_sycl_range with SYCL runtime behavior for write access mode

Fixes #1272

Summary

This PR aligns __get_sycl_range with SYCL semantics by adding no_init property support and making write mode perform copy-in by default (consistent with SYCL standard). It also optimizes write-only algorithms and fixes access mode workarounds.

Key Changes

Core Implementation

  • Added bool _NoInit = false template parameter to __get_sycl_range
  • Updated __is_copy_direct_v to make write mode copy-in by default unless no_init is specified
  • Removed unused _Iterator template parameter
  • Updated existing write mode callsites to use no_init=true, preserving current behavior

Pattern API Enhancements

Added _NoInit template parameters to __pattern_walk1/2/3 (and access mode for __pattern_walk1), enabling fine-grained copy-in control over access modes for output sequences. Removed (unsupported with vector) access modes for input sequences, they must be read without NoInit.

Optimizations

  • fill and generate: Now use write, no_init=true to avoid unnecessary copy-in
  • binary_search, lower_bound, upper_bound: changed to use write, no_init=true for output to avoid unnecessary copy-in

Fixes

  • transform_if: Changed from read_write workaround to proper write mode (without no_init) to preserve non-transformed elements
  • histogram: Changed from write workaround to read_write + no_init, correctly expressing atomic update semantics
  • unique: copy back changed from read_write for both input and output to defaulted read, then write with no_init, it is unclear why this was the case, but it is unnecessary.

@danhoeflinger danhoeflinger marked this pull request as draft November 18, 2025 17:25
Copy link
Contributor

Copilot AI left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Pull Request Overview

This PR refactors __get_sycl_range to align with SYCL runtime semantics for the write access mode. The primary change introduces a _NoInit template parameter to control copy-in behavior, making write mode perform copy-in by default (SYCL-compliant) unless explicitly suppressed.

Key Changes:

  • Added _NoInit template parameter to __get_sycl_range to control copy-in behavior for write access mode
  • Updated transform_if patterns to use proper write access mode instead of read_write workaround
  • Fixed histogram pattern to use read_write + no_init instead of write workaround

Reviewed Changes

Copilot reviewed 11 out of 11 changed files in this pull request and generated 3 comments.

Show a summary per file
File Description
utils_ranges_sycl.h Core implementation: added _NoInit parameter, removed unused _Iterator parameter, updated __is_copy_direct_v logic
algorithm_impl_hetero.h Updated all callsites to remove _Iterator parameter; added /*_NoInit=*/true to preserve existing behavior for write mode; fixed transform_if patterns
numeric_impl_hetero.h Updated callsites to remove _Iterator parameter and add /*_NoInit=*/true for write mode
histogram_impl_hetero.h Fixed histogram to use read_write + no_init instead of write workaround; removed _Iterator parameter from callsites
parallel_backend_sycl.h Updated set operation temporary buffers with /*_NoInit=*/true; removed _Iterator parameter
binary_search_impl.h Removed unused _Iterator template parameter from all __get_sycl_range calls
async_impl_hetero.h Updated async operations with /*_NoInit=*/true for write mode
glue_async_impl.h Removed _Iterator parameter from sort_async
single_pass_scan.h Updated scan kernel template with /*_NoInit=*/true
esimd_radix_sort_dispatchers.h Removed _Iterator parameter from radix sort dispatcher
esimd_radix_sort.h Removed _Iterator parameter from all radix sort variants

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

danhoeflinger and others added 5 commits December 17, 2025 09:28
 * separated no_init from write
 * remove unnecessary type specification

Signed-off-by: Dan Hoeflinger <[email protected]>
Signed-off-by: Dan Hoeflinger <[email protected]>
Signed-off-by: Dan Hoeflinger <[email protected]>
Signed-off-by: Dan Hoeflinger <[email protected]>
@danhoeflinger danhoeflinger force-pushed the dev/dhoeflin/align_write_no_init branch from 3625a7e to 8f0adfc Compare December 17, 2025 14:29
@danhoeflinger danhoeflinger marked this pull request as ready for review December 17, 2025 20:51
Signed-off-by: Dan Hoeflinger <[email protected]>
Copy link
Contributor

Copilot AI left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Pull request overview

Copilot reviewed 11 out of 11 changed files in this pull request and generated no new comments.


💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

Signed-off-by: Dan Hoeflinger <[email protected]>
Signed-off-by: Dan Hoeflinger <[email protected]>
Signed-off-by: Dan Hoeflinger <[email protected]>
//------------------------------------------------------------------------

template <typename _BackendTag, typename _ExecutionPolicy, typename _ForwardIterator, typename _Function>
template <__par_backend_hetero::access_mode __acc_mode = __par_backend_hetero::access_mode::read_write,
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thinking about the template defaults, they make sense for the common case across __pattern_walk 1,2,3. But is having different default template arguments for __pattern_walk1 from the others confusing?

Copy link
Contributor Author

@danhoeflinger danhoeflinger Dec 18, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think it does makes sense to have these defaults (read_write) for __pattern_walk1, partially because this has been the only mode prior to this PR for __pattern_walk1. However, I wouldn't object to simply requiring the access mode always (no default).

We could have the default be write & no_init, like the others, but I would worry a little about that because it changes the semantics of walk1 if no arguments are given.

While looking at this, I realize that uninitialized_fill and uninitialized_value_construct, uninitialized_default_construct also dont need read_write, but currently use it because they go through a wrapper of pattern_walk1. I should probably propagate these template options up to all the wrappers.

std::vector<std::unique_ptr<oneapi::dpl::__internal::__lifetime_keeper_base>> m_buffers;

template <sycl::access::mode _LocalAccMode>
template <sycl::access::mode _LocalAccMode, bool _LocalNoInit>
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think no reasons to have here bool _LocalNoInit and not use template parameter of bool _NoInit from this struct struct __get_sycl_range. This approach gives more changes than we may have.
Or somewhere we have some specialization of __is_copy_direct_v which initialized by something not equal to bool _NoInit.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Similar to _LocalAccMode, I believe the main reason for this is with permutation iterator and its map iterator.
Regardless of the exterior access mode, a permutation iterator's map is always read, and always without no-init.

It would cause problems to have permutation_iterator<usm_ptr, host_iter_map> with write + no_init and then use read + no_init for the host_iter_map after recursing.

This also leaves room for other future types which might require similar mechanisms, since the recursion of get_sycl_range for fancy iterators is all within a single instance of the struct.

@SergeyKopienko
Copy link
Contributor

Probably we forgot to change somehow the __pattern_walk1_async() call in

template <typename _BackendTag, typename _ExecutionPolicy, typename _ForwardIterator, typename _T>
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),
        fill_functor<_T>{__value});
}

@SergeyKopienko
Copy link
Contributor

One more consideration.
For example we have the code like

template <typename _BackendTag, typename _ExecutionPolicy, typename _ForwardIterator, typename _T>
_ForwardIterator
__pattern_fill(__hetero_tag<_BackendTag> __tag, _ExecutionPolicy&& __exec, _ForwardIterator __first,
               _ForwardIterator __last, const _T& __value)
{
    __pattern_walk1<__par_backend_hetero::access_mode::write, /*_NoInit=*/true>(
        __tag, ::std::forward<_ExecutionPolicy>(__exec),
        __par_backend_hetero::make_iter_mode<__par_backend_hetero::access_mode::write>(__first),
        __par_backend_hetero::make_iter_mode<__par_backend_hetero::access_mode::write>(__last),
        fill_functor<_T>{__value});
    return __last;
}

So we have two iterators initialized in __par_backend_hetero::access_mode::write mode.
If all iterators has this __par_backend_hetero::access_mode::write mode, it it not enough to make the same decision as you directly specify by /*_NoInit=*/true ?
I think if it really possible to extract fris mode from iterators, that better to avoid one more new template argument.

@danhoeflinger
Copy link
Contributor Author

One more consideration. For example we have the code like

template <typename _BackendTag, typename _ExecutionPolicy, typename _ForwardIterator, typename _T>
_ForwardIterator
__pattern_fill(__hetero_tag<_BackendTag> __tag, _ExecutionPolicy&& __exec, _ForwardIterator __first,
               _ForwardIterator __last, const _T& __value)
{
    __pattern_walk1<__par_backend_hetero::access_mode::write, /*_NoInit=*/true>(
        __tag, ::std::forward<_ExecutionPolicy>(__exec),
        __par_backend_hetero::make_iter_mode<__par_backend_hetero::access_mode::write>(__first),
        __par_backend_hetero::make_iter_mode<__par_backend_hetero::access_mode::write>(__last),
        fill_functor<_T>{__value});
    return __last;
}

So we have two iterators initialized in __par_backend_hetero::access_mode::write mode. If all iterators has this __par_backend_hetero::access_mode::write mode, it it not enough to make the same decision as you directly specify by /*_NoInit=*/true ? I think if it really possible to extract fris mode from iterators, that better to avoid one more new template argument.

Its a good consideration...
It may be possible to switch to something like this and removing the template arguments, but it needs more investigation. I don't like the way it is currently decoupled either. If I remember correctly, make_iter_mode is to have the access mode for the accessor in the case of a buffer. However, we could possibly utilize that if it is required to wrap iterators like this on their way in.

If we did switch to something like this we would want to ensure at compile time that the "input" iterators must be read (only). This is to support vector instructions which do not store back "input" iterators for walk2/3.

I'll investigate.

@danhoeflinger
Copy link
Contributor Author

Probably we forgot to change somehow the __pattern_walk1_async() call in

template <typename _BackendTag, typename _ExecutionPolicy, typename _ForwardIterator, typename _T>
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),
        fill_functor<_T>{__value});
}

I originally chose not to extend this to async patterns as we have not wanted to focus our efforts there, to limit the changes but I plan to extend the changes to include some of the wrapper patterns around __pattern_walk1/2/3 so I can do the async ones as well.

@SergeyKopienko
Copy link
Contributor

Probably we forgot to change somehow the __pattern_walk1_async() call in

template <typename _BackendTag, typename _ExecutionPolicy, typename _ForwardIterator, typename _T>
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),
        fill_functor<_T>{__value});
}

I originally chose not to extend this to async patterns as we have not wanted to focus our efforts there, to limit the changes but I plan to extend the changes to include some of the wrapper patterns around __pattern_walk1/2/3 so I can do the async ones as well.

I think it make sense to fix all places in async patterns too.

@danhoeflinger
Copy link
Contributor Author

danhoeflinger commented Dec 19, 2025

I'll investigate.

OK, I think I have an understanding of the make_iter_mode calls now. It is specifically for handling and resolving of embedded access modes within sycl_iterators vs algorithmic needs.

There are some issues with it that led me to create this issue:
#2550
Resolving this issue I think would remove all these make_iter_mode calls in favor of using __get_sycl_range to do any required resolution of embedded access modes in sycl_iterators, and remove the redundancy.

This means that I don't think that extending the wrapping of iterators in this way is a better way to handle access mode communication as compared to the direct template arguments for the walk functions.

Signed-off-by: Dan Hoeflinger <[email protected]>
@danhoeflinger
Copy link
Contributor Author

@SergeyKopienko @mmichel11
I did not want to grow the scope of this PR too large. I did add the async routines, but the uninitialized APIs required more refactoring, so I want to handle that in a separate PR: #2549 which will follow this one.

Signed-off-by: Dan Hoeflinger <[email protected]>
Signed-off-by: Dan Hoeflinger <[email protected]>
Copy link
Contributor

Copilot AI left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Pull request overview

Copilot reviewed 12 out of 12 changed files in this pull request and generated 2 comments.


💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

Signed-off-by: Dan Hoeflinger <[email protected]>
Signed-off-by: Dan Hoeflinger <[email protected]>
(check mangled output past range)

Signed-off-by: Dan Hoeflinger <[email protected]>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

Align __get_sycl_range with SYCL runtime in it's treatment of write access mode and no_init{}

3 participants