[PROTOTYPE] Add fallback to legacy scan implementation for CPU devices and devices that lack size 32 sub-groups #1749

mmichel11 · 2024-07-31T16:13:37Z

This PR adds a fallback from reduce-then-scan to the legacy scan implementation in two separate cases:

As a fallback if the device does not support size 32 sub-groups. We explicitly require this size and must fallback if it's not supported.
For CPU devices. The sub-group level operations do not translate well to CPU devices and the legacy implementation is significantly faster on CPUs.

Test is currently compiling. None of the real device code has been integrated, but the preliminary host code and general structure has been defined. Signed-off-by: Matthew Michel <matthew.michel@intel.com>

…implementation Signed-off-by: Matthew Michel <matthew.michel@intel.com>

Signed-off-by: Matthew Michel <matthew.michel@intel.com>

… the init computation Signed-off-by: Matthew Michel <matthew.michel@intel.com>

Signed-off-by: Matthew Michel <matthew.michel@intel.com>

This PR changes the two pass algorithm to be more generalized for use with other scan-like algorithms like copy_if. This PR adds copy_if as an example --------- Signed-off-by: Dan Hoeflinger <dan.hoeflinger@intel.com> Signed-off-by: Matthew Michel <matthew.michel@intel.com> Co-authored-by: Adam Fidel <110841220+adamfidel@users.noreply.github.com> Co-authored-by: Matthew Michel <106704043+mmichel11@users.noreply.github.com>

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

… fix use-after free error (#1712)

Signed-off-by: Matthew Michel <matthew.michel@intel.com>

…:tuple in zip_iterator.pass (#1714) * Fix for zip_iterator.pass in copy_if assignment Signed-off-by: Matthew Michel <matthew.michel@intel.com> * Add similar fix to __simple_write_to_idx Signed-off-by: Matthew Michel <matthew.michel@intel.com> --------- Signed-off-by: Matthew Michel <matthew.michel@intel.com>

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

…last active subgroup otherwise Signed-off-by: Matthew Michel <matthew.michel@intel.com>

Signed-off-by: Matthew Michel <matthew.michel@intel.com>

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

We fallback to the old implementation if the device does not support sub-group sizes of 32, or if the device is a CPU target. Signed-off-by: Matthew Michel <matthew.michel@intel.com>

Signed-off-by: Matthew Michel <matthew.michel@intel.com>

…gacy_scan_fallback Signed-off-by: Matthew Michel <matthew.michel@intel.com>

Signed-off-by: Matthew Michel <matthew.michel@intel.com>

mmichel11 · 2024-07-31T16:25:13Z

Moving back to draft. Some issues occurred when merging with the target branch.

Update: The issues have been fixed.

…cessary check Signed-off-by: Matthew Michel <matthew.michel@intel.com>

Signed-off-by: Matthew Michel <matthew.michel@intel.com>

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

Co-authored-by: Dan Hoeflinger <109972525+danhoeflinger@users.noreply.github.com>

adamfidel · 2024-07-31T19:50:11Z

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

@@ -769,6 +770,15 @@ class __static_monotonic_dispatcher<::std::integer_sequence<::std::uint16_t, _X,
    }
 };

+template <typename _ExecutionPolicy, typename _SizeType>
+bool
+__supports_sub_group_size(const _ExecutionPolicy& __exec, _SizeType __target_size)


Since the only usage of __target_size in this function casts it to std::size_t, would it make sense to take the parameter by std::size_t directly instead of using a template?

Suggested change

__supports_sub_group_size(const _ExecutionPolicy& __exec, _SizeType __target_size)

__supports_sub_group_size(const _ExecutionPolicy& __exec, std::size_t __target_size)

Good point. I've added this and removed the static_cast below.

danhoeflinger

one minor nitpick, otherwise, I think LGTM

danhoeflinger · 2024-07-31T19:26:33Z

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

@@ -1052,6 +1064,23 @@ __parallel_copy_if(oneapi::dpl::__internal::__device_backend_tag __backend_tag,
            oneapi::dpl::unseq_backend::__no_init_value<_Size>{},
            /*_Inclusive=*/std::true_type{});
    }
+    else
+    {
+        using _ReduceOp = ::std::plus<_Size>;


::std->std

Signed-off-by: Matthew Michel <matthew.michel@intel.com>

danhoeflinger

LGTM

…s and devices that lack size 32 sub-groups (#1749) Signed-off-by: Matthew Michel <matthew.michel@intel.com> Signed-off-by: Dan Hoeflinger <dan.hoeflinger@intel.com> Co-authored-by: Adam Fidel <adam.fidel@intel.com> Co-authored-by: Dan Hoeflinger <109972525+danhoeflinger@users.noreply.github.com> Co-authored-by: Adam Fidel <110841220+adamfidel@users.noreply.github.com> Co-authored-by: Dan Hoeflinger <dan.hoeflinger@intel.com>

mmichel11 and others added 22 commits July 24, 2024 11:20

Checkpoint for reduce then scan integration

1eaaf7d

Test is currently compiling. None of the real device code has been integrated, but the preliminary host code and general structure has been defined. Signed-off-by: Matthew Michel <matthew.michel@intel.com>

Introduce a parallel_backend_sycl_reduce_then_scan.h file to contain …

e7a909c

…implementation Signed-off-by: Matthew Michel <matthew.michel@intel.com>

Port of kernels from two-pass scan KT branch

b8c08c1

Signed-off-by: Matthew Michel <matthew.michel@intel.com>

Move the single-element last element storage for exclusive_scan after…

40b5f02

… the init computation Signed-off-by: Matthew Michel <matthew.michel@intel.com>

Use init value type for init processing helper

4ae5144

Lower single work-group upper limit to 2048 elements (empirically found)

a3034de

Signed-off-by: Matthew Michel <matthew.michel@intel.com>

bug fix for global race on block carry-out

e16a5af

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

bugfix for elements to process in partial subgroup scan

d90a348

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

[PROTOTYPE] Add unused temporary storage to single work-group scan to…

7f2dca1

… fix use-after free error (#1712)

Add temporary work-group size cap for FPGA_EMU testing

61295cf

Signed-off-by: Matthew Michel <matthew.michel@intel.com>

Use __dpl_sycl::__local_accessor

6585460

bugfix for overruning input for small non multiples of subgroup size

33d816d

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

Check if a subgroup is active before fetching its carry and grab the …

c155b58

…last active subgroup otherwise Signed-off-by: Matthew Michel <matthew.michel@intel.com>

Comment out std::complex tests in scan_by_segment tests

23f7886

Signed-off-by: Matthew Michel <matthew.michel@intel.com>

renaming __out as it seems to be a keyword

bd971db

fixing device copyable for helpers

5337317

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

Add legacy scan fallback case

e60aa75

We fallback to the old implementation if the device does not support sub-group sizes of 32, or if the device is a CPU target. Signed-off-by: Matthew Michel <matthew.michel@intel.com>

Clang-format

f2037cc

Signed-off-by: Matthew Michel <matthew.michel@intel.com>

Merge branch 'dev/shared/reduce_then_scan_impl' into dev/mmichel11/le…

ce98569

…gacy_scan_fallback Signed-off-by: Matthew Michel <matthew.michel@intel.com>

Remove unneeded #endif

8a0a7a9

Signed-off-by: Matthew Michel <matthew.michel@intel.com>

mmichel11 requested review from danhoeflinger and adamfidel July 31, 2024 16:13

mmichel11 marked this pull request as draft July 31, 2024 16:24

mmichel11 added 2 commits July 31, 2024 11:30

Remove bad code after merge and pull single-wg invocation out of unne…

e768892

…cessary check Signed-off-by: Matthew Michel <matthew.michel@intel.com>

Cleanup

7b69c66

Signed-off-by: Matthew Michel <matthew.michel@intel.com>

mmichel11 marked this pull request as ready for review July 31, 2024 16:40

danhoeflinger reviewed Jul 31, 2024

View reviewed changes

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

Check __target_size instead of 32 in __supports_sub_group_size

e11df04

Co-authored-by: Dan Hoeflinger <109972525+danhoeflinger@users.noreply.github.com>

adamfidel reviewed Jul 31, 2024

View reviewed changes

danhoeflinger reviewed Jul 31, 2024

View reviewed changes

mmichel11 added 2 commits July 31, 2024 15:08

::std:: -> std::

ef96d5a

Signed-off-by: Matthew Michel <matthew.michel@intel.com>

Remove unnecessary templated _SizeType

5942d73

Signed-off-by: Matthew Michel <matthew.michel@intel.com>

danhoeflinger approved these changes Jul 31, 2024

View reviewed changes

mmichel11 merged commit cdd7ceb into dev/shared/reduce_then_scan_impl Jul 31, 2024

mmichel11 deleted the dev/mmichel11/legacy_scan_fallback branch July 31, 2024 21:58

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[PROTOTYPE] Add fallback to legacy scan implementation for CPU devices and devices that lack size 32 sub-groups #1749

[PROTOTYPE] Add fallback to legacy scan implementation for CPU devices and devices that lack size 32 sub-groups #1749

mmichel11 commented Jul 31, 2024

mmichel11 commented Jul 31, 2024 •

edited

Loading

adamfidel Jul 31, 2024

mmichel11 Jul 31, 2024

danhoeflinger left a comment

danhoeflinger Jul 31, 2024

mmichel11 Jul 31, 2024

danhoeflinger left a comment

	__supports_sub_group_size(const _ExecutionPolicy& __exec, _SizeType __target_size)
	__supports_sub_group_size(const _ExecutionPolicy& __exec, std::size_t __target_size)

[PROTOTYPE] Add fallback to legacy scan implementation for CPU devices and devices that lack size 32 sub-groups #1749

[PROTOTYPE] Add fallback to legacy scan implementation for CPU devices and devices that lack size 32 sub-groups #1749

Conversation

mmichel11 commented Jul 31, 2024

mmichel11 commented Jul 31, 2024 • edited Loading

adamfidel Jul 31, 2024

Choose a reason for hiding this comment

mmichel11 Jul 31, 2024

Choose a reason for hiding this comment

danhoeflinger left a comment

Choose a reason for hiding this comment

danhoeflinger Jul 31, 2024

Choose a reason for hiding this comment

mmichel11 Jul 31, 2024

Choose a reason for hiding this comment

danhoeflinger left a comment

Choose a reason for hiding this comment

mmichel11 commented Jul 31, 2024 •

edited

Loading