Replace SYCL backend `reduce_by_segment` implementation with reduce-then-scan call #1915

mmichel11 · 2024-10-22T19:18:24Z

Summary

This PR implements a SYCL backend reduce_by_segment by using higher level calls to reduce-then-scan along with new specialty functors to achieve a segmented reduction. This PR is an initial step of porting the implementation to reduce-then-scan with optimization likely to follow. Future efforts may include additional modification to reduce-then-scan kernels.

Performance improves for all input sizes. For small inputs, we see 3-5x improvements and for very large sizes ~1.25x on GPU Series Max 1550. Please contact me if you would like to see performance data.

Description of changes

The SYCL reduce_by_segment implementation that was previously handwritten is replaced by a higher level call to our reduce-then-scan kernels. Several new callback functors for the reduce-then-scan kernel have been made to achieve this operation.
reduce_by_segment.pass was encountering linker crashes due to the large number of test cases being compiled growing past the maximum size of the binary's data region. SYCL testing has been trimmed down with regards to USM device and shared testing which resolves this issue. Instead of running each test with a device and shared USM allocation, every other test switches the USM type.
ONEDPL_WORKAROUND_FOR_IGPU_64BIT_REDUCTION has been removed as the SYCL implementation has been replaced, and we are no longer impacted by this issue.
The legacy reduce_by_segment implementation is used as a fallback for when the sub-group size, device, and trivial copyability constraints cannot be satisfied.

Future work

Future efforts on reduce_by_segment may built on top of this implementation and the reduce-then-scan kernels to better handle first and last element cases.

mmichel11 · 2024-11-05T20:30:03Z

I have made some design changes based on offline discussion. There is quite a bit of code movement that has happened, so here is a summary of the recently made changes:

An iterator-based __pattern_reduce_by_segment has been added to algorithm_impl_hetero.h. Previously, we just had a range-based version. This resolves the issue of calling range-based patterns from iterator-based algorithms.
The fallback reduce_by_segment implementation based on high-level copy_if and parallel_for calls has been moved down a level from algorithm_ranges_impl_hetero.h to dpcpp/parallel_backend_sycl.h so that we can fallback on this implementation when reduce-then-scan cannot be used in __parallel_reduce_by_segment. This pattern has been implemented synchronously as parallel pattern calls cannot currently depend on each other.
Due to observed performance issues for compilers prior to icpx 2025.0, the reduce-then-scan path must be disabled in this case. The known-identity based implementation has been added back to avoid introducing a performance regression for icpx 2024.2.1 and prior. It has been moved into dpcpp/parallel_backend_sycl_reduce_by_segment.h so the implementations are not split across several directories.

include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h

danhoeflinger

I like this implementation and thing it is basically good to go.
I think more comments are necessary, and there is some potential for future gains.

I probably want to look a bit further in to minor details before approving with another pass but at a high level I think this is in good shape.

include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h

test/general/implementation_details/device_copyable.pass.cpp

test/parallel_api/numeric/numeric.ops/reduce_by_segment.pass.cpp

danhoeflinger · 2024-11-21T19:28:04Z

include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_reduce_by_segment.h

+namespace __par_backend_hetero
+{
+
+template <typename... Name>


I confirmed in an outside editor that the only changes to this from the previous location are only cosmetic. I'm not looking deeply into this files changes otherwise since it has already been reviewed and it was just moved.

include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h

danhoeflinger · 2024-11-21T19:50:03Z

include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h

+        {
+            const _KeyType& __next_key = __in_keys[__id + 1];
+            return oneapi::dpl::__internal::make_tuple(
+                oneapi::dpl::__internal::make_tuple(std::size_t{0}, _ValueType{__in_vals[__id]}),
+                !__binary_pred(__current_key, __next_key), __next_key, __current_key);
+        }


It may be possible to avoid the __id == 0 case, in a similar way to unique. It is a little more complicated because we would need to set up the carry-in appropriately, but I think its possible and could provide some branch avoiding (and tuple shrinking) gains in the helpers.
If you think its possible to do this, lets leave it as an issue to be explored in a follow up.

I agree and have opened #1958.

adamfidel

First pass of comments. I have looked at primarily the fallback algorithms and intend to focus on the reduce-then-scan implementation next.

include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h

adamfidel · 2024-11-21T17:35:23Z

include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_reduce_by_segment.h

+                                                                              sycl::nd_item<1> __item) {
+                auto __group = __item.get_group();
+                std::size_t __group_id = __item.get_group(0);
+                std::size_t __local_id = __item.get_local_id(0);


This could probably safely be a uint32_t. Or, you can remove this variable and replace its only usage with if (__group.leader().

I changed this type to be std::uint32_t

include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_reduce_by_segment.h

test/general/implementation_details/device_copyable.pass.cpp

danhoeflinger

At a high level I agree with the changes in this PR, but there are still a few remaining nit picks outstanding.

I have run out of time before my time off to get into the small details like sizes of types and forwarding of references, things like that. The clang format suggestions can be ignored as of now.

So, I wont hit approve officially but I think this is very close and trust @adamfidel / others to be able to get it across the finish line and have no objections to merging with another approval.

adamfidel

A few more comments.

I think overall the approach is good and I don't mind approving once these few small comments have been addressed.

include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h

adamfidel

LGTM overall. The clang-format changes can be ignored in my opinion.

mmichel11 · 2024-12-10T14:55:56Z

I have fixed a compilation issue identified in CI. With some configurations, specifying the full namespace oneapi::dpl::__ranges when calling __get_sycl_range is needed to avoid compilation errors.

I am rerunning CI to make sure there are no more issues.