[Reduce_then_scan refactor pt 1] Relaxing requirement of trivial copyable types

[danhoeflinger]

This pull request relaxes the requirement of trivially copyable value types for reduce_then_scan, taking the first step toward enabling it for all cases.

• Implements SLM and barrier-based fallback implementations for __group_broadcast and __shift_group_right. This requires extra SLM data to be allocated and passed.
• To avoid many repeated branches in inner loops, this extracts the kernel bodies into member functions to branch once for SLM vs sycl-builtin implementations. This is the biggest change from a LOC standpoint, but is mostly just moving code.
• Removes trivially copyable restriction in code dispatching to reduce_then_scan
• Adds explicit return type to __gen_count_mask, __gen_expand_count_mask, __gen_set_balanced_path,
  __gen_set_op_from_known_balanced_path to correct intermediate type to be better controlled and resolve ambiguity with SLM pointer type.

Full picture:

• Relax trivially copyable requirement <------
• Relax subgroup size requirements
• Implement output limited scan_copy
• Remove dead code implementations (scan_then_propagate, reduce_by_segment_fallback, scan_by_segment_fallback)

[Copilot]

Pull request overview

This PR refactors the SYCL reduce-then-scan implementation to support non-trivially-copyable value types by introducing SLM/barrier-based sub-group communication fallbacks, and updates related generator types/traits to use explicit return types for better type control.

Changes:

• Adds SLM/barrier-based fallback implementations for sub-group communication primitives used by reduce-then-scan.
• Removes the “trivially copyable” restriction in dispatch logic so reduce-then-scan can be selected for more value types.
• Adds explicit return-type template parameters to several generator types and updates associated device-copyable trait specializations and tests.

Reviewed changes

Copilot reviewed 8 out of 8 changed files in this pull request and generated 7 comments.

Show a summary per file

File	Description
test/general/implementation_details/device_copyable.pass.cpp	Updates static-asserts for new generator template parameters (explicit return types).
test/general/implementation_details/balanced_path_unit_tests.pass.cpp	Updates balanced-path generator instantiation to include explicit return type.
include/oneapi/dpl/pstl/hetero/dpcpp/sycl_traits.h	Updates forward declarations and sycl::is_device_copyable specializations for new generator signatures.
include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h	Removes trivially-copyable gating for reduce-then-scan selection in scan/segmented paths.
include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_reduce_then_scan.h	Implements SLM-based sub-group communication fallbacks and refactors kernels to branch once per kernel body.
include/oneapi/dpl/experimental/kt/internal/work_group/work_group_scan.h	Updates calls into the shared sub-group scan primitives to match new signatures.
include/oneapi/dpl/experimental/kt/internal/sub_group/sub_group_scan.h	Updates calls into the shared sub-group scan primitives to match new signatures.
include/oneapi/dpl/experimental/kt/internal/cooperative_lookback.h	Updates calls into the shared sub-group scan primitives to match new signatures.

---

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