Restrict __future to the async API boundary

[SergeyKopienko]

Summary

__future used to be created deep inside the SYCL backend and then carried up through every internal algorithm layer, even for fully synchronous calls. This PR removes __future from all internal layers: backends and pattern implementations now return a backend-neutral event (__hetero_event<_BackendTag>) or a std::tuple of such an event plus storage, and __future is constructed only at the async API boundary.

What changed

__hetero_event<_BackendTag> and backend-tag specializations (*)

• Introduced __hetero_event<_BackendTag> — a thin, backend-tagged wrapper over the native completion object. It exposes wait(), wait_and_throw(), a nested __type alias for the native event, and an implicit conversion to it.
• __hetero_event<__device_backend_tag> is specialized over sycl::event; __hetero_event<__fpga_backend_tag> (under _ONEDPL_FPGA_DEVICE) derives from the device specialization and reuses its constructors.
• Added the trait __is_hetero_event / __is_hetero_event_v, which is true for sycl::event as well as for both __hetero_event specializations. It drives the overload resolution of __finalize_call.
• Pattern layers now spell the returned type explicitly as __hetero_event<_BackendTag> (or __hetero_event<__device_backend_tag> where the tag is fixed), which keeps the internal contract visible and independent of __future.
• __future is now parameterized by the backend tag (__future<_BackendTag, _Args...>) and stores an __hetero_event<_BackendTag>; operator __native_event_t() preserves the previous conversion to sycl::event.

Finalization and result reading

• Added __finalize_call<_WaitModeTag>(...) as the single place implementing the previous wait semantics (__sync_mode / __deferrable_mode / __async_mode, ONEDPL_ALLOW_DEFERRED_WAITING). It has two overloads: one for anything satisfying __is_hetero_event_v, and one for a std::tuple whose first element is such an event.
• The tuple overload resolves the effective wait mode through __resolve_wait_mode and the __wait_required_of_finalize_sycl_call trait: if the tuple carries __device_storage, __result_storage or __combined_storage, deferred/async modes are upgraded to __sync_mode, because the storage must outlive the kernel.
• The tuple overload takes a non-const lvalue reference on purpose: the payload must outlive the waiting, so passing a temporary is prohibited by construction.
• Added __load_result(...) for reading a host result out of the storage returned by a backend call. It does not require the value type to be default-constructible: the no-holder overload uses __lazy_ctor_storage<_T>, exactly as the former __result_and_scratch_storage::__get_value() did, so device-copyable types with a deleted default constructor keep working (min_element, minmax_element, reduce_by_segment, the corresponding range algorithms). A second overload accepts an already constructed result holder (used by transform_reduce, where the storage keeps _RepackedTp).

Future creation and payloads

• __create_future(std::tuple, _ExtraArgs&&...) is now the only future factory; __make_future and the extra overloads were removed. At most one extra payload item is accepted, and it is placed first, so it is the one returned by __future::get().
• Payload conversion goes through __to_future_payload:
  • __combined_storage → __result_and_scratch_storage (the payload the result is read from);
  • __device_storage → __lifetime_payload<__device_storage> (a lifetime-only payload);
  • __result_storage is explicitly rejected by a static_assert to prevent a silent slicing to the __device_storage base class.
• Added __lifetime_payload: a copyable, shared-ownership wrapper for payloads that must outlive the kernel but carry no algorithm result. __future::get() for such a payload only waits and returns void, and the future stays copyable as before.

Storage cleanup

• __combined_storage now accepts a zero-sized scratch request and allocates no scratch memory in that case; it is default-constructible, so "no storage" states are expressible without pointers.
• __result_and_scratch_storage is reduced to a value-reading payload: the __result_and_scratch_storage_base virtual indirection, __get_data / __fill_data, __get_result_acc / __get_scratch_acc and __get_usm_or_buffer_accessor_ptr were removed. Kernels use __get_accessor / __get_result_accessor and acc.__data() instead.
• Removed dead code: __future::wait(_WaitModeTag) and __future::__checked_deferrable_wait().

Merge and merge sort return data

• Split points are kept in __device_storage<_split_point_t<_IdType>> (_split_points_device_storage_t, with the 32/64 aliases) instead of __result_and_scratch_storage behind a shared_ptr<__result_and_scratch_storage_base>.
• __parallel_merge_return_data_t<_OutSizeLimit, _Range1, _Range2> describes the merge return as a tuple: event, 32-bit split points storage, 64-bit split points storage, and — only when _OutSizeLimit is true — a __result_storage for the stop position. __create_parallel_merge_return_data, __get_parallel_merge_sp_storage and __get_parallel_merge_stop_pos_accessor_opt (with the __no_parallel_merge_stop_pos_acc_tag sentinel) build and access it.
• __parallel_sort_return_t gives merge sort the same shape (event plus the two split-point storages), so the index-type branch is expressed in the type instead of a type-erased base pointer.
• Added __difference_tuple_t<_Range...> in utils_ranges.h and reused it in the merge and set-operation return types.

Before: __future created inside the backend

flowchart TD
    U1["User call: sync"] --> A1["__pattern_*"]
    U2["User call: async"] --> A2["__pattern_*_async"]
    A1 --> B["parallel_backend_sycl"]
    A2 --> B
    B --> F["__future created here"]
    F --> E["sycl::event + storage"]
    F -.->|"wait() / get() on every layer"| A1
    F -.->|"returned as is"| A2

Loading

After: __future created only at the async boundary

flowchart TD
    U1["User call: sync"] --> A1["__pattern_*"]
    U2["User call: async"] --> A2["__pattern_*_async"]
    A1 --> B["parallel_backend_sycl"]
    A2 --> B
    B --> E["__hetero_event or<br/>tuple of event and storage"]
    E --> S["__finalize_call<br/>__load_result"]
    E --> C["__create_future<br/>the only place"]
    S --> R1["plain result"]
    C --> R2["__future for the user"]

Loading

Backend-tagged event

flowchart TD
    H["__hetero_event&lt;_BackendTag&gt;"] --> D["__device_backend_tag:<br/>wraps sycl::event"]
    H --> F["__fpga_backend_tag:<br/>derives from device one"]
    D --> T["__is_hetero_event_v"]
    F --> T
    T --> FC["__finalize_call&lt;_WaitModeTag&gt;"]
    T --> FU["__future&lt;_BackendTag, ...&gt;"]

Loading

Payload conversion in __create_future

flowchart TD
    T["tuple of event and storage"] --> P["__to_future_payload"]
    P --> P1["__combined_storage"]
    P --> P2["__device_storage"]
    P --> P3["plain value"]
    P --> P4["__result_storage"]
    P1 --> C1["__result_and_scratch_storage"]
    P2 --> C2["__lifetime_payload"]
    P3 --> C3["kept as is"]
    P4 --> C4["static_assert:<br/>not supported"]
    C1 --> G1["get(): wait and<br/>return the result"]
    C3 --> G1
    C2 --> G2["get(): wait only,<br/>return void"]

Loading

Replacement mapping

Old (via __future)	New
backend returns __future	backend returns sycl::event / __hetero_event<_BackendTag> / std::tuple
sycl::event spelled directly in internal layers	__hetero_event<_BackendTag> with tag specializations
__future::wait() in internal layers	__finalize_call<__sync_mode>(...)
__future::wait(__deferrable_mode) / __checked_deferrable_wait()	__finalize_call<__deferrable_mode>(...)
__future::get() in internal layers	__load_result(...) / __result_and_scratch_storage::__wait_and_get_value
__make_future(...) / multiple __create_future overloads	single __create_future(std::tuple, _ExtraArgs&&...)
std::shared_ptr<__result_and_scratch_storage_base> as a lifetime payload	__lifetime_payload<__device_storage>
__result_and_scratch_storage for split points	__device_storage<_split_point_t<_IdType>> + __parallel_merge_return_data_t / __parallel_sort_return_t
__get_scratch_acc / __get_result_acc / __get_usm_or_buffer_accessor_ptr	__get_accessor / __get_result_accessor + acc.__data()
storage carried as __future payload	__to_future_payload performs the conversion

Tests

• asynch.pass.cpp now covers the lifetime-only payload contract: the future returned by sort_async is copied (shared ownership) and its get() is checked to return void. A dedicated comparator functor is passed to sort_async so that the merge sort path (the one returning a lifetime-only payload) is taken in all configurations instead of the radix sort path.

Notes

• __future remains only as the public async return type; its wait(), get(), event() and the conversion to the native event are still used by the async API and tests.
• __to_future_payload(__combined_storage&&) asserts __result_sz <= 1, tolerating the placeholder storage returned by the single-work-group transform-scan path.
• Wait semantics were re-checked against main: places that historically used .wait() were kept as .wait(), while the paths that previously went through __future::wait() keep wait_and_throw() (e.g. the empty-input branch in __pattern_histogram).
• No public API or standard-version requirements changed; the code stays within the existing C++17 baseline of the headers.

[Copilot]

Pull request overview

Moves __future construction to asynchronous API boundaries while simplifying internal SYCL backend return types and synchronization.

Changes:

• Replaces internal futures with events and event/payload tuples.
• Centralizes waiting and result extraction.
• Updates algorithms and storage handling for the new return model.

Reviewed changes

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

Show a summary per file

File	Description
include/oneapi/dpl/pstl/utils_ranges.h	Adds range difference tuple alias.
include/oneapi/dpl/pstl/hetero/numeric_ranges_impl_hetero.h	Finalizes numeric range operations explicitly.
include/oneapi/dpl/pstl/hetero/numeric_impl_hetero.h	Adapts numeric algorithms to event returns.
include/oneapi/dpl/pstl/hetero/histogram_impl_hetero.h	Adapts histogram synchronization.
include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h	Converts core backend returns to events/tuples.
include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_utils.h	Adds finalization, payload conversion, and result loading.
include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_reduce.h	Migrates reduction storage and returns.
include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_reduce_then_scan_pos_tools.h	Uses the difference tuple alias.
include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_radix_sort.h	Returns radix-sort events directly.
include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_merge.h	Reworks merge result and temporary storage.
include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_merge_sort.h	Reworks merge-sort temporary storage.
include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_histogram.h	Returns histogram events directly.
include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_fpga.h	Returns FPGA backend events directly.
include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_for.h	Returns parallel-for events directly.
include/oneapi/dpl/pstl/hetero/algorithm_ranges_impl_hetero.h	Updates range algorithms for tuple results.
include/oneapi/dpl/pstl/hetero/algorithm_impl_hetero.h	Updates iterator algorithms for explicit finalization.
include/oneapi/dpl/internal/binary_search_impl.h	Finalizes binary-search kernels explicitly.
include/oneapi/dpl/internal/async_impl/glue_async_impl.h	Creates sort futures at the API boundary.
include/oneapi/dpl/internal/async_impl/async_impl_hetero.h	Creates asynchronous futures from backend tuples.

---

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

[Copilot]

Pull request overview

Copilot reviewed 19 out of 19 changed files in this pull request and generated 1 comment.

[Copilot]

Pull request overview

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

Suppressed comments (1)

include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_utils.h:999

• The new lifetime-only payload contract is not covered by the existing async tests: asynch.pass.cpp only waits on sort_async futures and never copies one or calls get(). Add a non-radix comparator case that copies the returned future and verifies that get() returns void after sorting, so both shared ownership and the dedicated overload remain protected.

// __device_storage is a move-only payload which is required to keep the data alive until the kernel completes.
// It carries no algorithm result, but __future must stay copyable, so such a payload is kept by a shared ownership.
template <typename _T>
__lifetime_payload<__device_storage<_T>>
__to_future_payload(__device_storage<_T>&& __ds)
{
    return __lifetime_payload<__device_storage<_T>>{std::make_shared<__device_storage<_T>>(std::move(__ds))};

[Copilot]

Pull request overview

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

Suppressed comments (2)

include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_reduce.h:145

• The single-work-group reduction has no scratch data, but this now allocates one scratch element solely to satisfy __combined_storage. On Level Zero GPUs, where the result uses host USM, __combined_storage consequently performs an additional device-USM allocation for every small reduction; the old {__q, 0} storage allocated only the result. Allow zero-sized scratch storage (while retaining any fake accessor buffer required by the compatibility macro) so this common path does not add an allocation.

        __combined_storage<_Tp> __result(__q, /*__scratch_n, ==1 just for compatibility*/ 1, 1);

test/parallel_api/experimental/asynch.pass.cpp:150

• This test does not reliably exercise the lifetime-only payload it describes. With radix sort enabled, int plus std::greater<int> selects the std::tuple<sycl::event> radix path, so the copy/get assertions never instantiate __lifetime_payload. Use an unrecognized comparator type to force the merge-sort return tuple and cover the regression in all configurations.

    // The future returned by sort_async carries a lifetime-only payload: it must stay copyable
    // and its get() must not expose any internal storage.
    auto delta_copy = delta;
    static_assert(std::is_void_v<decltype(delta_copy.get())>, "sort_async future must return void from get()");

[danhoeflinger]

Are we OK that sycl specifics are spilling out into algorithm_impl_hetero (sycl::event)?

I think our general disposition to this point has been that this layer should be agnostic to backend specifics and written in a way which would not be tailored to a sycl backend necessarily.

We may want some general "event" contract or something like that. Future has been this for us so far.

[SergeyKopienko]

Are we OK that sycl specifics are spilling out into algorithm_impl_hetero (sycl::event)?

I think our general disposition to this point has been that this layer should be agnostic to backend specifics and written in a way which would not be tailored to a sycl backend necessarily.

We may want some general "event" contract or something like that. Future has been this for us so far.

You mean lower_bound_impl, upper_bound_impl and binary_search_impl ?

I think they are simply implemented in the wrong place.
The evidence - they get __hetero_tag as the first parameter.
I think we can move them later into some proper place and no reasons for now mask under auto that they are operating by sycl::event inside.

[SergeyKopienko]

Are we OK that sycl specifics are spilling out into algorithm_impl_hetero (sycl::event)?

I think our general disposition to this point has been that this layer should be agnostic to backend specifics and written in a way which would not be tailored to a sycl backend necessarily.

We may want some general "event" contract or something like that. Future has been this for us so far.

Fixed - let's still use auto

[Copilot]

Pull request overview

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

Suppressed comments (1)

include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_reduce.h:145

• This single-work-group path never uses scratch storage, but passing 1 makes __combined_storage allocate it anyway. In the host-USM result path, __combined_storage performs a separate device-USM allocation for this element (parallel_backend_sycl_utils.h:767-773), regressing the former zero-scratch path and adding an allocation to every small reduction. Please allow __scratch_n == 0 in __combined_storage and pass 0 here.

        __combined_storage<_Tp> __result(__q, /*__scratch_n, ==1 just for compatibility*/ 1, 1);

[MikeDvorskiy]

It looks this PR breaks the asynchrony of the oneDPL async API.

[MikeDvorskiy]

I am afraid this PR breaks the asynchrony of the oneDPL async API because:
As far as I can see, this PR makes the

1. As far as I can see, this PR makes the __par_backend_hetero SYCL backend (SYCL patterns) synchronous instead of asynchronous.
   __par_backend_hetero SYCL backend (SYCL patterns) synchronous instead of asynchronous.
2. No new asynchronous SYCL patterns are added.
3. The oneDPL async code layer simply reuses the SYCL backend, which has become synchronous and uses blocking calls.

[MikeDvorskiy]

Please correct me if I am wrong.

[SergeyKopienko]

A little bit later I going to prepare some presentation and discuss about this architecture design for this PR.
But in short - yes, you are incorrect.

Missed some details during review the changes.