Avoid specializations of the same submitters with the some policy type but with different type qualifiers (`l-value`, `r-value`) #2093

SergeyKopienko · 2025-03-03T11:14:46Z

In this PR we fix the issue #2083 and #2041

The problem statement.

Usually, oneDPL submitters which do sycl::queue::submit() calls are implemented in three ways:

operator() implementation inside some class / struct;
submit() function implementation inside some class / struct;
some free function.

In all cases, these functions are template functions parametrized by _ExecutioPolicy type. Also they have _ExecutioPolicy&& __exec in the list of functions parameters.
This approach means that we may have two specializations of these functions:

one specialization for the case when __exec is r-value;
another specialization for other cases.

The problems of this approach which already implemented in `oneDPL`:

we should generate unique Kernel-names using not only _ExecutionPolicy type but also their qualifiers (l-value, r-value);
we increase the amount of sycl-kernel code and their compile-time;
we increase the amount of host code in compile-time.

The solution

As solution in this PR proposed the idea of @akukanov - to declare _ExecutionPolicy in submitters and on one upper level as const _ExecutionPolicy& (const reference).
This approach fixes all point of the problem described below.

Some key changes in the code

the usage of _ExecutionPolicy is not required inside __kernel_name_generator template parameters anymore:
-- before:

    template <typename... _Name>
    using _SegScanPrefixPhase = __seg_scan_prefix_kernel<__is_inclusive, _Name...>;

    template <typename _BackendTag, typename _ExecutionPolicy, typename _Range1, typename _Range2, typename _Range3,
              typename _BinaryPredicate, typename _BinaryOperator, typename _T>
    void
    operator()(_BackendTag, _ExecutionPolicy&& __exec, _Range1&& __keys, _Range2&& __values, _Range3&& __out_values,
               _BinaryPredicate __binary_pred, _BinaryOperator __binary_op, _T __init, _T __identity)
    {
        using _CustomName = oneapi::dpl::__internal::__policy_kernel_name<_ExecutionPolicy>;

        using _SegScanWgKernel = oneapi::dpl::__par_backend_hetero::__internal::__kernel_name_generator<
            _SegScanWgPhase, _CustomName, _ExecutionPolicy, _Range1, _Range2, _Range3, _BinaryPredicate,
            _BinaryOperator>;

        // ...
    }

-- now:

    template <typename _BackendTag, typename _ExecutionPolicy, typename _Range1, typename _Range2, typename _Range3,
              typename _BinaryPredicate, typename _BinaryOperator, typename _T>
    void
    operator()(_BackendTag, const _ExecutionPolicy& __exec, _Range1&& __keys, _Range2&& __values, _Range3&& __out_values,
               _BinaryPredicate __binary_pred, _BinaryOperator __binary_op, _T __init, _T __identity)
    {
        using _CustomName = oneapi::dpl::__internal::__policy_kernel_name<_ExecutionPolicy>;

        // We should avoid using _ExecutionPolicy in __kernel_name_generator template params
        // because we always specialize this operator() calls only by _ExecutionPolicy as "const reference".
        // So, from this template param point of view, only one specialization is possible.
        using _SegScanWgKernel = oneapi::dpl::__par_backend_hetero::__internal::__kernel_name_generator<
            _SegScanWgPhase, _CustomName, _Range1, _Range2, _Range3, _BinaryPredicate, _BinaryOperator>;

        // ...
    }

Now we able to detect a lot of errors as described in #2041 - this test coverage implemented in the PR #2102

Another approach has been prepared in #2082

Please review this PR under switched off option "File changed - Gear button - Hide whitespace"

…allel_merge_submitter_large::operator()

…an_static_single_group_submitter::operator(), __parallel_transform_scan_dynamic_single_group_submitter::oeprator()

…ator()

…up, __parallel_transform_reduce_then_scan

…tatic_single_group_submitter::operator(), __parallel_copy_if_static_single_group_submitter::operator()

…hen_scan

…form_reduce_then_scan

…en_scan

…ral_registers_local_reduction_submitter::operator()

…omics_submitter::operator()

…_private_global_atomics_submitter::operator()

…rs_local_reduction, __histogram_general_local_atomics, __histogram_general_private_global_atomics

…impl, __parallel_transform_reduce_mid_impl, __parallel_transform_reduce_impl

…) call

…form_reduce_then_scan

….h - remove :: before std::decay_t in __kernel_name_generator

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

Insert example from #2041

…t case

SergeyKopienko mentioned this pull request Mar 3, 2025

Avoid specializations of the same submitters with the some policy type but with different qualifiers #2082

Closed

SergeyKopienko force-pushed the dev/skopienko/fix_exec_policy_type_in_submitters_V1 branch 4 times, most recently from 0976051 to 0102ab8 Compare March 3, 2025 16:27

SergeyKopienko changed the title ~~Avoid specializations of the same submitters with the some policy type but with different qualifiers - V1~~ Avoid specializations of the same submitters with the some policy type but with different qualifiers Mar 3, 2025

SergeyKopienko requested review from mmichel11, dmitriy-sobolev, danhoeflinger and akukanov and removed request for akukanov, dmitriy-sobolev, mmichel11 and danhoeflinger March 3, 2025 17:41

SergeyKopienko added this to the 2022.9.0 milestone Mar 3, 2025

SergeyKopienko added the enhancement label Mar 3, 2025

SergeyKopienko force-pushed the dev/skopienko/fix_exec_policy_type_in_submitters_V1 branch 4 times, most recently from 629a956 to e005551 Compare March 4, 2025 09:51

SergeyKopienko requested review from dmitriy-sobolev, mmichel11, danhoeflinger and rarutyun March 4, 2025 10:40

SergeyKopienko mentioned this pull request Mar 4, 2025

Some __kernel_name_generator usages doesn't depends from _ExecutionPolicy type #2083

Open

SergeyKopienko marked this pull request as ready for review March 4, 2025 10:46

SergeyKopienko changed the title ~~Avoid specializations of the same submitters with the some policy type but with different qualifiers~~ Avoid specializations of the same submitters with the some policy type but with different type qualifiers (l-value, r-value and etc.) Mar 4, 2025

SergeyKopienko force-pushed the dev/skopienko/fix_exec_policy_type_in_submitters_V1 branch 2 times, most recently from fe6cd52 to 1880b1e Compare March 4, 2025 12:32

SergeyKopienko and others added 29 commits March 6, 2025 18:57

Fix __parallel_merge -> __parallel_merge_submitter::operator(), __par…

47f5ee5

…allel_merge_submitter_large::operator()

Fix __merge_sort and stuff

ac1b9b8

Fix __radix_sort_count_submit

be8fe91

Fix __radix_sort_scan_submit

50288d6

Fix __radix_sort_reorder_submit

874a10e

Fix __parallel_transform_scan_single_group -> __parallel_transform_sc…

9185dbe

…an_static_single_group_submitter::operator(), __parallel_transform_scan_dynamic_single_group_submitter::oeprator()

Fix __parallel_transform_scan_base -> __parallel_scan_submitter::oper…

02903da

…ator()

Fix __parallel_transform_scan -> __parallel_transform_scan_single_gro…

d81d171

…up, __parallel_transform_reduce_then_scan

Fix __invoke_single_group_copy_if::operator() -> __parallel_copy_if_s…

6559b0c

…tatic_single_group_submitter::operator(), __parallel_copy_if_static_single_group_submitter::operator()

Fix __parallel_reduce_then_scan_copy -> __parallel_transform_reduce_t…

47d6c93

…hen_scan

Fix __parallel_reduce_by_segment_reduce_then_scan -> __parallel_trans…

6f30401

…form_reduce_then_scan

Fix __parallel_set_reduce_then_scan -> __parallel_transform_reduce_th…

c2daf0a

…en_scan

Fix __parallel_for -> __parallel_for_fpga_submitter::operator()

99ce9b5

Fix __histogram_general_registers_local_reduction -> __histogram_gene…

638c2ff

…ral_registers_local_reduction_submitter::operator()

Fix __histogram_general_private_global_atomics_submitter::operator()

ebbb160

Fix __histogram_general_local_atomics -> __histogram_general_local_at…

d3e6fb2

…omics_submitter::operator()

Fix __histogram_general_private_global_atomics -> __histogram_general…

a266177

…_private_global_atomics_submitter::operator()

Fix __parallel_histogram_select_kernel -> __histogram_general_registe…

59f1831

…rs_local_reduction, __histogram_general_local_atomics, __histogram_general_private_global_atomics

Fix __parallel_histogram -> __parallel_histogram_select_kernel

fab428e

Fix __parallel_transform_reduce -> __parallel_transform_reduce_small_…

d99a2cd

…impl, __parallel_transform_reduce_mid_impl, __parallel_transform_reduce_impl

Fix __parallel_radix_sort -> __parallel_radix_sort_iteration::submit(…

b150823

…) call

Fix __parallel_reduce_then_scan_copy

6327040

fix __parallel_reduce_by_segment_reduce_then_scan -> __parallel_trans…

89370f8

…form_reduce_then_scan

include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_radix_sort…

d2cd3a2

….h - remove :: before std::decay_t in __kernel_name_generator

Apply GitHUB clang format

85bc86f

Update include/oneapi/dpl/internal/scan_by_segment_impl.h

bdb73a2

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

Fix review comments

83339c8

test/general/lambda_naming.pass.cpp - expand test coverage

1b13d6f

Insert example from #2041

include/oneapi/dpl/pstl/hetero/numeric_impl_hetero.h - fix broken tes…

Loading
Loading status checks…

ca5f2da

…t case

SergeyKopienko force-pushed the dev/skopienko/fix_exec_policy_type_in_submitters_V1 branch from 3d9988c to ca5f2da Compare March 6, 2025 17:57

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Avoid specializations of the same submitters with the some policy type but with different type qualifiers (`l-value`, `r-value`) #2093

Avoid specializations of the same submitters with the some policy type but with different type qualifiers (`l-value`, `r-value`) #2093

SergeyKopienko commented Mar 3, 2025 •

edited

Loading

Avoid specializations of the same submitters with the some policy type but with different type qualifiers (l-value, r-value) #2093

Are you sure you want to change the base?

Avoid specializations of the same submitters with the some policy type but with different type qualifiers (l-value, r-value) #2093

Conversation

SergeyKopienko commented Mar 3, 2025 • edited Loading

The problem statement.

The problems of this approach which already implemented in oneDPL:

The solution

Some key changes in the code

Please review this PR under switched off option "File changed - Gear button - Hide whitespace"

Avoid specializations of the same submitters with the some policy type but with different type qualifiers (`l-value`, `r-value`) #2093

Avoid specializations of the same submitters with the some policy type but with different type qualifiers (`l-value`, `r-value`) #2093

SergeyKopienko commented Mar 3, 2025 •

edited

Loading

The problems of this approach which already implemented in `oneDPL`: