Add reduce then scan algorithm for transform scan family #1762

danhoeflinger · 2024-08-05T17:30:53Z

This PR Adds a generic reduce_then_scan algorithm for scan-like algorithms, and uses this algorithm for transform scan family of scan-like algorithms where it is beneficial.

This PR is targeted after a pair of other PRs #1769 + #1770, which lay some plumbing for these changes. Those two PRs were separated from this one to provide more focus to this change.

It is important to understand that subsequent PRs will add other families of scan-like algorithms (copy_if, partition, unique, ...), which may justify some aspects of the implementation of the reduce_then_scan kernels that could seem overcomplicated for the transform scan family alone.

This PR is targeted to #1770, to allow for a clean diff, and is a part of the following sequence of PRs meant to be merged in order:

#1769 [MERGED] ~~Relocate __lazy_ctor_storage to utils header~~
#1770 [MERGED] ~~Use __result_and_scratch_storage within scan kernels~~
#1762 Add reduce_then_scan algorithm for transform scan family (This PR)
#1763 Make Copy_if family of APIs use reduce_then_scan algorithm
#1764 Make Partition family of APIs use reduce_then_scan algorithm
#1765 Make Unique family of APIs use reduce_then_scan algorithm

This work is a collaboration between @mmichel11 @adamfidel and @danhoeflinger, and based upon an original prototype by Ted Painter.

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

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

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

julianmi

I've had another look and I think it is in a good state. Clang-format is red, but I prefer the current version since it is easier to read.

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

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

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

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

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

MikeDvorskiy · 2024-08-27T15:58:18Z

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

+{
+    template <typename _OutRng, typename ValueType>
+    void
+    operator()(const _OutRng& __out_rng, std::size_t __id, const ValueType& __v) const


The constant reference const _OutRng& __out_rng looks suspicious, because the range is output.

Good point, I removed the const here this and fixed a couple other range types in the kernel. I will go through the remaining PRs to check for similar issues. (I think output ranges for the other helpers probably have the same issue)

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

julianmi

LGTM. Please wait for another approval since there are partial reviews of several people.

adamfidel · 2024-08-29T15:45:11Z

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

+                //    and then write back the final values to memory
+                if (__sub_group_id == 0)
+                {
+                    // step 1) load to Xe SLM the WG-local S prefix sums


Should we remove the reference to Xe here since this algorithm is designed to work on non-Xe-based devices as well?

Yes, I think so.

danhoeflinger

a few minor things I found.

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

danhoeflinger · 2024-08-29T15:51:57Z

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

+                //    and then write back the final values to memory
+                if (__sub_group_id == 0)
+                {
+                    // step 1) load to Xe SLM the WG-local S prefix sums


Yes, I think so.

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

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

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

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

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

Co-authored-by: Matthew Michel <106704043+mmichel11@users.noreply.github.com>

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

adamfidel

This PR is in a good state and is okay to be merged.

mmichel11

LGTM. We have checked tests and performance with the latest version.

akukanov · 2024-08-24T17:13:55Z

include/oneapi/dpl/pstl/hetero/numeric_ranges_impl_hetero.h

@@ -91,35 +91,15 @@ oneapi::dpl::__internal::__difference_t<_Range2>
 __pattern_transform_scan_base(__hetero_tag<_BackendTag>, _ExecutionPolicy&& __exec, _Range1&& __rng1, _Range2&& __rng2,
                              _UnaryOperation __unary_op, _InitType __init, _BinaryOperation __binary_op, _Inclusive)
 {
-    if (__rng1.empty())
+    oneapi::dpl::__internal::__difference_t<_Range1> __n = __rng1.size();


A comment I forgot to send earlier:
__n is also the return value, and the function is declared to return __difference_t<_Range2>, not that of _Range1. In the previous code, we took the size of __rng1 and converted it to the return type. But I wonder if we should just take the size of __rng2 instead (if we assume ranges being of equal size), or maybe change the return type?

Thanks for pointing this out. We created PR #1811 to address this issue.

danhoeflinger requested review from mmichel11 and adamfidel August 5, 2024 17:30

mmichel11 reviewed Aug 6, 2024

View reviewed changes

include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_reduce_then_scan.h Outdated Show resolved Hide resolved

danhoeflinger commented Aug 6, 2024

View reviewed changes

include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h Outdated Show resolved Hide resolved

danhoeflinger commented Aug 6, 2024

View reviewed changes

include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h Outdated Show resolved Hide resolved

adamfidel mentioned this pull request Aug 6, 2024

Relocate __lazy_ctor_storage to utils header #1769

Merged

danhoeflinger force-pushed the dev/shared/transform_reduce_then_scan branch from 9eb3099 to fc43983 Compare August 6, 2024 17:16

danhoeflinger changed the base branch from main to dev/shared/scan_result_and_scratch_storage August 6, 2024 17:21

danhoeflinger added this to the 2022.7.0 milestone Aug 6, 2024

danhoeflinger force-pushed the dev/shared/scan_result_and_scratch_storage branch from c8a274a to 4e4c0bc Compare August 6, 2024 18:59

danhoeflinger force-pushed the dev/shared/transform_reduce_then_scan branch 2 times, most recently from 3d277c6 to 4df758a Compare August 6, 2024 20:58

danhoeflinger marked this pull request as ready for review August 6, 2024 22:02

danhoeflinger requested review from julianmi, timmiesmith, akukanov, MikeDvorskiy, dmitriy-sobolev, SergeyKopienko, reble and rarutyun August 6, 2024 22:03

danhoeflinger changed the title ~~[Draft] Add reduce then scan algorithm for transform scan family~~ Add reduce then scan algorithm for transform scan family Aug 6, 2024

danhoeflinger force-pushed the dev/shared/transform_reduce_then_scan branch 3 times, most recently from ea28133 to d01aeb1 Compare August 8, 2024 12:27

danhoeflinger force-pushed the dev/shared/scan_result_and_scratch_storage branch from 47a227c to c471ed5 Compare August 8, 2024 14:06

adamfidel and others added 5 commits August 21, 2024 07:58

Add alias for __dpl_sycl::__sub_group and replace templates

6096e7a

auto -> real types and formatting

60c8516

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

fixing type of subgroup id returns

8121d67

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

shrinking subgroup size id types

48724db

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

adjust type to depend on input range

3cc61db

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

julianmi reviewed Aug 22, 2024

View reviewed changes

include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h Outdated Show resolved Hide resolved

include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_reduce_then_scan.h Show resolved Hide resolved

idx -> id

c2c7e35

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

julianmi reviewed Aug 22, 2024

View reviewed changes

shrinking types, switch branch to min, remove double deref

ff7b256

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

danhoeflinger mentioned this pull request Aug 22, 2024

Potential performance improvements for reduce-then-scan algorithm #1792

Open

adamfidel and others added 3 commits August 22, 2024 17:05

Adjust block size for reduce-then-scan based on input type (#1782)

8a36d5a

shrinking missed types

9520f3c

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

bugfix for windows

5a928fd

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

MikeDvorskiy reviewed Aug 27, 2024

View reviewed changes

fixing range types

e57573f

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

julianmi previously approved these changes Aug 29, 2024

View reviewed changes

adamfidel reviewed Aug 29, 2024

View reviewed changes

danhoeflinger commented Aug 29, 2024

View reviewed changes

minor comments from review + formatting

93189b0

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

danhoeflinger dismissed julianmi’s stale review via 93189b0 August 29, 2024 16:16

mmichel11 reviewed Aug 29, 2024

View reviewed changes

danhoeflinger and others added 2 commits August 29, 2024 16:13

Apply std:: suggestions

4e4568e

Co-authored-by: Matthew Michel <106704043+mmichel11@users.noreply.github.com>

rounding workgroup size down to mult of subgroup size

af82182

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

adamfidel approved these changes Aug 30, 2024

View reviewed changes

mmichel11 approved these changes Aug 30, 2024

View reviewed changes

julianmi approved these changes Aug 30, 2024

View reviewed changes

danhoeflinger merged commit f320dbd into main Aug 30, 2024
20 of 21 checks passed

danhoeflinger deleted the dev/shared/transform_reduce_then_scan branch August 30, 2024 15:42

akukanov reviewed Aug 30, 2024

View reviewed changes

adamfidel mentioned this pull request Aug 30, 2024

Fix return type of transform scan and copy if patterns #1811

Merged

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Add reduce then scan algorithm for transform scan family #1762

Add reduce then scan algorithm for transform scan family #1762

danhoeflinger commented Aug 5, 2024 •

edited

Loading

julianmi left a comment

MikeDvorskiy Aug 27, 2024 •

edited

Loading

danhoeflinger Aug 27, 2024

julianmi left a comment

adamfidel Aug 29, 2024

danhoeflinger Aug 29, 2024

danhoeflinger left a comment

danhoeflinger Aug 29, 2024

adamfidel left a comment

mmichel11 left a comment

akukanov Aug 24, 2024

adamfidel Aug 30, 2024

Add reduce then scan algorithm for transform scan family #1762

Add reduce then scan algorithm for transform scan family #1762

Conversation

danhoeflinger commented Aug 5, 2024 • edited Loading

julianmi left a comment

Choose a reason for hiding this comment

MikeDvorskiy Aug 27, 2024 • edited Loading

Choose a reason for hiding this comment

danhoeflinger Aug 27, 2024

Choose a reason for hiding this comment

julianmi left a comment

Choose a reason for hiding this comment

adamfidel Aug 29, 2024

Choose a reason for hiding this comment

danhoeflinger Aug 29, 2024

Choose a reason for hiding this comment

danhoeflinger left a comment

Choose a reason for hiding this comment

danhoeflinger Aug 29, 2024

Choose a reason for hiding this comment

adamfidel left a comment

Choose a reason for hiding this comment

mmichel11 left a comment

Choose a reason for hiding this comment

akukanov Aug 24, 2024

Choose a reason for hiding this comment

adamfidel Aug 30, 2024

Choose a reason for hiding this comment

danhoeflinger commented Aug 5, 2024 •

edited

Loading

MikeDvorskiy Aug 27, 2024 •

edited

Loading