Skip to content

Commit fe6cd52

Browse files
Apply GitHUB clang format
1 parent a719672 commit fe6cd52

8 files changed

+76
-73
lines changed

include/oneapi/dpl/internal/scan_by_segment_impl.h

+6-5
Original file line numberDiff line numberDiff line change
@@ -108,8 +108,9 @@ struct __sycl_scan_by_segment_impl
108108
template <typename _BackendTag, typename _ExecutionPolicy, typename _Range1, typename _Range2, typename _Range3,
109109
typename _BinaryPredicate, typename _BinaryOperator, typename _T>
110110
void
111-
operator()(_BackendTag, const _ExecutionPolicy& __exec, _Range1&& __keys, _Range2&& __values, _Range3&& __out_values,
112-
_BinaryPredicate __binary_pred, _BinaryOperator __binary_op, _T __init, _T __identity)
111+
operator()(_BackendTag, const _ExecutionPolicy& __exec, _Range1&& __keys, _Range2&& __values,
112+
_Range3&& __out_values, _BinaryPredicate __binary_pred, _BinaryOperator __binary_op, _T __init,
113+
_T __identity)
113114
{
114115
using _CustomName = oneapi::dpl::__internal::__policy_kernel_name<_ExecutionPolicy>;
115116

@@ -395,9 +396,9 @@ __scan_by_segment_impl_common(__internal::__hetero_tag<_BackendTag>, const Polic
395396

396397
constexpr iter_value_t identity = unseq_backend::__known_identity<Operator, iter_value_t>;
397398

398-
__sycl_scan_by_segment_impl<Inclusive::value>()(_BackendTag{}, policy, key_buf.all_view(),
399-
value_buf.all_view(), value_output_buf.all_view(), binary_pred,
400-
binary_op, init, identity);
399+
__sycl_scan_by_segment_impl<Inclusive::value>()(_BackendTag{}, policy, key_buf.all_view(), value_buf.all_view(),
400+
value_output_buf.all_view(), binary_pred, binary_op, init,
401+
identity);
401402
return result + n;
402403
}
403404

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

+46-47
Original file line numberDiff line numberDiff line change
@@ -232,8 +232,8 @@ struct __parallel_scan_submitter;
232232
template <typename _CustomName, typename... _PropagateScanName>
233233
struct __parallel_scan_submitter<_CustomName, __internal::__optional_kernel_name<_PropagateScanName...>>
234234
{
235-
template <typename _ExecutionPolicy, typename _Range1, typename _Range2, typename _InitType,
236-
typename _LocalScan, typename _GroupScan, typename _GlobalScan>
235+
template <typename _ExecutionPolicy, typename _Range1, typename _Range2, typename _InitType, typename _LocalScan,
236+
typename _GroupScan, typename _GlobalScan>
237237
auto
238238
operator()(const _ExecutionPolicy& __exec, _Range1&& __rng1, _Range2&& __rng2, _InitType __init,
239239
_LocalScan __local_scan, _GroupScan __group_scan, _GlobalScan __global_scan) const
@@ -594,8 +594,8 @@ __parallel_transform_scan_single_group(oneapi::dpl::__internal::__device_backend
594594
::std::integral_constant<::std::uint16_t, __wg_size>,
595595
::std::integral_constant<::std::uint16_t, __num_elems_per_item>, _BinaryOperation,
596596
/* _IsFullGroup= */ std::true_type, _Inclusive, _CustomName>>>()(
597-
__exec, std::forward<_InRng>(__in_rng),
598-
std::forward<_OutRng>(__out_rng), __n, __init, __binary_op, __unary_op);
597+
__exec, std::forward<_InRng>(__in_rng), std::forward<_OutRng>(__out_rng), __n, __init, __binary_op,
598+
__unary_op);
599599
else
600600
__event = __parallel_transform_scan_static_single_group_submitter<
601601
_Inclusive::value, __num_elems_per_item, __wg_size,
@@ -604,8 +604,8 @@ __parallel_transform_scan_single_group(oneapi::dpl::__internal::__device_backend
604604
::std::integral_constant<::std::uint16_t, __wg_size>,
605605
::std::integral_constant<::std::uint16_t, __num_elems_per_item>, _BinaryOperation,
606606
/* _IsFullGroup= */ ::std::false_type, _Inclusive, _CustomName>>>()(
607-
__exec, std::forward<_InRng>(__in_rng),
608-
std::forward<_OutRng>(__out_rng), __n, __init, __binary_op, __unary_op);
607+
__exec, std::forward<_InRng>(__in_rng), std::forward<_OutRng>(__out_rng), __n, __init, __binary_op,
608+
__unary_op);
609609
return __future(__event, __dummy_result_and_scratch);
610610
};
611611
if (__n <= 16)
@@ -638,8 +638,8 @@ __parallel_transform_scan_single_group(oneapi::dpl::__internal::__device_backend
638638

639639
auto __event =
640640
__parallel_transform_scan_dynamic_single_group_submitter<_Inclusive::value, _DynamicGroupScanKernel>()(
641-
__exec, std::forward<_InRng>(__in_rng),
642-
std::forward<_OutRng>(__out_rng), __n, __init, __binary_op, __unary_op, __max_wg_size);
641+
__exec, std::forward<_InRng>(__in_rng), std::forward<_OutRng>(__out_rng), __n, __init, __binary_op,
642+
__unary_op, __max_wg_size);
643643
return __future(__event, __dummy_result_and_scratch);
644644
}
645645
}
@@ -656,9 +656,9 @@ __parallel_transform_scan_base(oneapi::dpl::__internal::__device_backend_tag, co
656656
using _PropagateKernel =
657657
oneapi::dpl::__par_backend_hetero::__internal::__kernel_name_provider<__scan_propagate_kernel<_CustomName>>;
658658

659-
return __parallel_scan_submitter<_CustomName, _PropagateKernel>()(
660-
__exec, std::forward<_Range1>(__in_rng), std::forward<_Range2>(__out_rng),
661-
__init, __local_scan, __group_scan, __global_scan);
659+
return __parallel_scan_submitter<_CustomName, _PropagateKernel>()(__exec, std::forward<_Range1>(__in_rng),
660+
std::forward<_Range2>(__out_rng), __init,
661+
__local_scan, __group_scan, __global_scan);
662662
}
663663

664664
template <typename _Type>
@@ -1075,9 +1075,9 @@ __parallel_transform_scan(oneapi::dpl::__internal::__device_backend_tag __backen
10751075
std::size_t __single_group_upper_limit = __use_reduce_then_scan ? 2048 : 16384;
10761076
if (__group_scan_fits_in_slm<_Type>(__exec.queue(), __n, __n_uniform, __single_group_upper_limit))
10771077
{
1078-
return __parallel_transform_scan_single_group(
1079-
__backend_tag, __exec, std::forward<_Range1>(__in_rng),
1080-
std::forward<_Range2>(__out_rng), __n, __unary_op, __init, __binary_op, _Inclusive{});
1078+
return __parallel_transform_scan_single_group(__backend_tag, __exec, std::forward<_Range1>(__in_rng),
1079+
std::forward<_Range2>(__out_rng), __n, __unary_op, __init,
1080+
__binary_op, _Inclusive{});
10811081
}
10821082
}
10831083
if (__use_reduce_then_scan)
@@ -1088,10 +1088,10 @@ __parallel_transform_scan(oneapi::dpl::__internal::__device_backend_tag __backen
10881088

10891089
_GenInput __gen_transform{__unary_op};
10901090

1091-
return __parallel_transform_reduce_then_scan(
1092-
__backend_tag, __exec, std::forward<_Range1>(__in_rng),
1093-
std::forward<_Range2>(__out_rng), __gen_transform, __binary_op, __gen_transform, _ScanInputTransform{},
1094-
_WriteOp{}, __init, _Inclusive{}, /*_IsUniquePattern=*/std::false_type{});
1091+
return __parallel_transform_reduce_then_scan(__backend_tag, __exec, std::forward<_Range1>(__in_rng),
1092+
std::forward<_Range2>(__out_rng), __gen_transform, __binary_op,
1093+
__gen_transform, _ScanInputTransform{}, _WriteOp{}, __init,
1094+
_Inclusive{}, /*_IsUniquePattern=*/std::false_type{});
10951095
}
10961096
}
10971097

@@ -1148,8 +1148,8 @@ struct __invoke_single_group_copy_if
11481148
using _FullKernelName = oneapi::dpl::__par_backend_hetero::__internal::__kernel_name_provider<_FullKernel>;
11491149
return __par_backend_hetero::__parallel_copy_if_static_single_group_submitter<
11501150
_SizeType, __num_elems_per_item, __wg_size, true, _FullKernelName>()(
1151-
__exec, std::forward<_InRng>(__in_rng),
1152-
std::forward<_OutRng>(__out_rng), __n, _InitType{}, _ReduceOp{}, __pred, __assign);
1151+
__exec, std::forward<_InRng>(__in_rng), std::forward<_OutRng>(__out_rng), __n, _InitType{}, _ReduceOp{},
1152+
__pred, __assign);
11531153
}
11541154
else
11551155
{
@@ -1161,29 +1161,29 @@ struct __invoke_single_group_copy_if
11611161
oneapi::dpl::__par_backend_hetero::__internal::__kernel_name_provider<_NonFullKernel>;
11621162
return __par_backend_hetero::__parallel_copy_if_static_single_group_submitter<
11631163
_SizeType, __num_elems_per_item, __wg_size, false, _NonFullKernelName>()(
1164-
__exec, std::forward<_InRng>(__in_rng),
1165-
std::forward<_OutRng>(__out_rng), __n, _InitType{}, _ReduceOp{}, __pred, __assign);
1164+
__exec, std::forward<_InRng>(__in_rng), std::forward<_OutRng>(__out_rng), __n, _InitType{}, _ReduceOp{},
1165+
__pred, __assign);
11661166
}
11671167
}
11681168
};
11691169

11701170
template <typename _ExecutionPolicy, typename _InRng, typename _OutRng, typename _Size, typename _GenMask,
11711171
typename _WriteOp, typename _IsUniquePattern>
11721172
auto
1173-
__parallel_reduce_then_scan_copy(oneapi::dpl::__internal::__device_backend_tag __backend_tag, const _ExecutionPolicy& __exec,
1174-
_InRng&& __in_rng, _OutRng&& __out_rng, _Size, _GenMask __generate_mask,
1175-
_WriteOp __write_op, _IsUniquePattern __is_unique_pattern)
1173+
__parallel_reduce_then_scan_copy(oneapi::dpl::__internal::__device_backend_tag __backend_tag,
1174+
const _ExecutionPolicy& __exec, _InRng&& __in_rng, _OutRng&& __out_rng, _Size,
1175+
_GenMask __generate_mask, _WriteOp __write_op, _IsUniquePattern __is_unique_pattern)
11761176
{
11771177
using _GenReduceInput = oneapi::dpl::__par_backend_hetero::__gen_count_mask<_GenMask>;
11781178
using _ReduceOp = std::plus<_Size>;
11791179
using _GenScanInput = oneapi::dpl::__par_backend_hetero::__gen_expand_count_mask<_GenMask>;
11801180
using _ScanInputTransform = oneapi::dpl::__par_backend_hetero::__get_zeroth_element;
11811181

1182-
return __parallel_transform_reduce_then_scan(
1183-
__backend_tag, __exec, std::forward<_InRng>(__in_rng),
1184-
std::forward<_OutRng>(__out_rng), _GenReduceInput{__generate_mask}, _ReduceOp{}, _GenScanInput{__generate_mask},
1185-
_ScanInputTransform{}, __write_op, oneapi::dpl::unseq_backend::__no_init_value<_Size>{},
1186-
/*_Inclusive=*/std::true_type{}, __is_unique_pattern);
1182+
return __parallel_transform_reduce_then_scan(__backend_tag, __exec, std::forward<_InRng>(__in_rng),
1183+
std::forward<_OutRng>(__out_rng), _GenReduceInput{__generate_mask},
1184+
_ReduceOp{}, _GenScanInput{__generate_mask}, _ScanInputTransform{},
1185+
__write_op, oneapi::dpl::unseq_backend::__no_init_value<_Size>{},
1186+
/*_Inclusive=*/std::true_type{}, __is_unique_pattern);
11871187
}
11881188

11891189
template <typename _ExecutionPolicy, typename _InRng, typename _OutRng, typename _Size, typename _CreateMaskOp,
@@ -1379,9 +1379,9 @@ __parallel_copy_if(oneapi::dpl::__internal::__device_backend_tag __backend_tag,
13791379
template <typename _ExecutionPolicy, typename _Range1, typename _Range2, typename _Range3, typename _Compare,
13801380
typename _IsOpDifference>
13811381
auto
1382-
__parallel_set_reduce_then_scan(oneapi::dpl::__internal::__device_backend_tag __backend_tag, const _ExecutionPolicy& __exec,
1383-
_Range1&& __rng1, _Range2&& __rng2, _Range3&& __result, _Compare __comp,
1384-
_IsOpDifference)
1382+
__parallel_set_reduce_then_scan(oneapi::dpl::__internal::__device_backend_tag __backend_tag,
1383+
const _ExecutionPolicy& __exec, _Range1&& __rng1, _Range2&& __rng2, _Range3&& __result,
1384+
_Compare __comp, _IsOpDifference)
13851385
{
13861386
// fill in reduce then scan impl
13871387
using _GenMaskReduce = oneapi::dpl::__par_backend_hetero::__gen_set_mask<_IsOpDifference, _Compare>;
@@ -1469,9 +1469,9 @@ __parallel_set_op(oneapi::dpl::__internal::__device_backend_tag __backend_tag, _
14691469
{
14701470
if (oneapi::dpl::__par_backend_hetero::__is_gpu_with_reduce_then_scan_sg_sz(__exec))
14711471
{
1472-
return __parallel_set_reduce_then_scan(__backend_tag, __exec,
1473-
std::forward<_Range1>(__rng1), std::forward<_Range2>(__rng2),
1474-
std::forward<_Range3>(__result), __comp, __is_op_difference);
1472+
return __parallel_set_reduce_then_scan(__backend_tag, __exec, std::forward<_Range1>(__rng1),
1473+
std::forward<_Range2>(__rng2), std::forward<_Range3>(__result), __comp,
1474+
__is_op_difference);
14751475
}
14761476
else
14771477
{
@@ -1907,8 +1907,8 @@ __parallel_find_or(oneapi::dpl::__internal::__device_backend_tag, const _Executi
19071907

19081908
// Single WG implementation
19091909
__result = __parallel_find_or_impl_one_wg<__or_tag_check, __find_or_one_wg_kernel_name>()(
1910-
oneapi::dpl::__internal::__device_backend_tag{}, __exec, __brick_tag,
1911-
__rng_n, __wgroup_size, __init_value, __pred, std::forward<_Ranges>(__rngs)...);
1910+
oneapi::dpl::__internal::__device_backend_tag{}, __exec, __brick_tag, __rng_n, __wgroup_size, __init_value,
1911+
__pred, std::forward<_Ranges>(__rngs)...);
19121912
}
19131913
else
19141914
{
@@ -1920,8 +1920,8 @@ __parallel_find_or(oneapi::dpl::__internal::__device_backend_tag, const _Executi
19201920

19211921
// Multiple WG implementation
19221922
__result = __parallel_find_or_impl_multiple_wgs<__or_tag_check, __find_or_kernel_name>()(
1923-
oneapi::dpl::__internal::__device_backend_tag{}, __exec, __brick_tag,
1924-
__rng_n, __n_groups, __wgroup_size, __init_value, __pred, std::forward<_Ranges>(__rngs)...);
1923+
oneapi::dpl::__internal::__device_backend_tag{}, __exec, __brick_tag, __rng_n, __n_groups, __wgroup_size,
1924+
__init_value, __pred, std::forward<_Ranges>(__rngs)...);
19251925
}
19261926

19271927
if constexpr (__or_tag_check)
@@ -2174,8 +2174,8 @@ class __sort_global_kernel;
21742174

21752175
template <typename _ExecutionPolicy, typename _Range, typename _Merge, typename _Compare>
21762176
auto
2177-
__parallel_partial_sort_impl(oneapi::dpl::__internal::__device_backend_tag, const _ExecutionPolicy& __exec, _Range&& __rng,
2178-
_Merge __merge, _Compare __comp)
2177+
__parallel_partial_sort_impl(oneapi::dpl::__internal::__device_backend_tag, const _ExecutionPolicy& __exec,
2178+
_Range&& __rng, _Merge __merge, _Compare __comp)
21792179
{
21802180
using _CustomName = oneapi::dpl::__internal::__policy_kernel_name<_ExecutionPolicy>;
21812181
using _GlobalSortKernel =
@@ -2184,8 +2184,7 @@ __parallel_partial_sort_impl(oneapi::dpl::__internal::__device_backend_tag, cons
21842184
oneapi::dpl::__par_backend_hetero::__internal::__kernel_name_provider<__sort_copy_back_kernel<_CustomName>>;
21852185

21862186
return __parallel_partial_sort_submitter<_GlobalSortKernel, _CopyBackKernel>()(
2187-
oneapi::dpl::__internal::__device_backend_tag{}, __exec,
2188-
::std::forward<_Range>(__rng), __merge, __comp);
2187+
oneapi::dpl::__internal::__device_backend_tag{}, __exec, ::std::forward<_Range>(__rng), __merge, __comp);
21892188
}
21902189

21912190
//------------------------------------------------------------------------
@@ -2413,9 +2412,9 @@ __parallel_reduce_by_segment(oneapi::dpl::__internal::__device_backend_tag, _Exe
24132412
if (oneapi::dpl::__par_backend_hetero::__is_gpu_with_reduce_then_scan_sg_sz(__exec))
24142413
{
24152414
auto __res = oneapi::dpl::__par_backend_hetero::__parallel_reduce_by_segment_reduce_then_scan(
2416-
oneapi::dpl::__internal::__device_backend_tag{}, __exec,
2417-
std::forward<_Range1>(__keys), std::forward<_Range2>(__values), std::forward<_Range3>(__out_keys),
2418-
std::forward<_Range4>(__out_values), __binary_pred, __binary_op);
2415+
oneapi::dpl::__internal::__device_backend_tag{}, __exec, std::forward<_Range1>(__keys),
2416+
std::forward<_Range2>(__values), std::forward<_Range3>(__out_keys), std::forward<_Range4>(__out_values),
2417+
__binary_pred, __binary_op);
24192418
// Because our init type ends up being tuple<std::size_t, ValType>, return the first component which is the write index. Add 1 to return the
24202419
// past-the-end iterator pair of segmented reduction.
24212420
return std::get<0>(__res.get()) + 1;

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

+8-7
Original file line numberDiff line numberDiff line change
@@ -288,9 +288,10 @@ struct __histogram_general_registers_local_reduction_submitter<__iters_per_work_
288288
template <::std::uint16_t __iters_per_work_item, ::std::uint8_t __bins_per_work_item, typename _ExecutionPolicy,
289289
typename _Range1, typename _Range2, typename _BinHashMgr>
290290
auto
291-
__histogram_general_registers_local_reduction(oneapi::dpl::__internal::__device_backend_tag, const _ExecutionPolicy& __exec,
292-
const sycl::event& __init_event, ::std::uint16_t __work_group_size,
293-
_Range1&& __input, _Range2&& __bins, const _BinHashMgr& __binhash_manager)
291+
__histogram_general_registers_local_reduction(oneapi::dpl::__internal::__device_backend_tag,
292+
const _ExecutionPolicy& __exec, const sycl::event& __init_event,
293+
::std::uint16_t __work_group_size, _Range1&& __input, _Range2&& __bins,
294+
const _BinHashMgr& __binhash_manager)
294295
{
295296
using _CustomName = oneapi::dpl::__internal::__policy_kernel_name<_ExecutionPolicy>;
296297

@@ -481,10 +482,10 @@ struct __histogram_general_private_global_atomics_submitter<__internal::__option
481482
};
482483
template <typename _ExecutionPolicy, typename _Range1, typename _Range2, typename _BinHashMgr>
483484
auto
484-
__histogram_general_private_global_atomics(oneapi::dpl::__internal::__device_backend_tag, const _ExecutionPolicy& __exec,
485-
const sycl::event& __init_event, ::std::uint16_t __min_iters_per_work_item,
486-
::std::uint16_t __work_group_size, _Range1&& __input, _Range2&& __bins,
487-
const _BinHashMgr& __binhash_manager)
485+
__histogram_general_private_global_atomics(oneapi::dpl::__internal::__device_backend_tag,
486+
const _ExecutionPolicy& __exec, const sycl::event& __init_event,
487+
::std::uint16_t __min_iters_per_work_item, ::std::uint16_t __work_group_size,
488+
_Range1&& __input, _Range2&& __bins, const _BinHashMgr& __binhash_manager)
488489
{
489490
using _CustomName = oneapi::dpl::__internal::__policy_kernel_name<_ExecutionPolicy>;
490491

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

+2-1
Original file line numberDiff line numberDiff line change
@@ -206,7 +206,8 @@ struct __parallel_merge_submitter<_OutSizeLimit, _IdType, __internal::__optional
206206
{
207207
template <typename _ExecutionPolicy, typename _Range1, typename _Range2, typename _Range3, typename _Compare>
208208
auto
209-
operator()(const _ExecutionPolicy& __exec, _Range1&& __rng1, _Range2&& __rng2, _Range3&& __rng3, _Compare __comp) const
209+
operator()(const _ExecutionPolicy& __exec, _Range1&& __rng1, _Range2&& __rng2, _Range3&& __rng3,
210+
_Compare __comp) const
210211
{
211212
const _IdType __n1 = __rng1.size();
212213
const _IdType __n2 = __rng2.size();

0 commit comments

Comments
 (0)