Skip to content

Commit

Permalink
Apply GitHUB CLang format
Browse files Browse the repository at this point in the history
Signed-off-by: Sergey Kopienko <sergey.kopienko@intel.com>
  • Loading branch information
SergeyKopienko committed Sep 6, 2024
1 parent 1251e87 commit 73ac0a5
Show file tree
Hide file tree
Showing 5 changed files with 18 additions and 15 deletions.
16 changes: 8 additions & 8 deletions include/oneapi/dpl/internal/reduce_by_segment_impl.h
Original file line number Diff line number Diff line change
Expand Up @@ -281,9 +281,9 @@ __sycl_reduce_by_segment(__internal::__hetero_tag<_BackendTag>, _ExecutionPolicy
auto __seg_end_identification = __exec.queue().submit([&](sycl::handler& __cgh) {
oneapi::dpl::__ranges::__require_access(__cgh, __keys);
auto __seg_ends_acc = __seg_ends.template get_access<sycl::access_mode::write>(__cgh);
#if _ONEDPL_COMPILE_KERNEL
# if _ONEDPL_COMPILE_KERNEL
__cgh.use_kernel_bundle(__seg_reduce_count_kernel.get_kernel_bundle());
#endif
# endif
__cgh.parallel_for<_SegReduceCountKernel>(
sycl::nd_range<1>{__n_groups * __wgroup_size, __wgroup_size}, [=](sycl::nd_item<1> __item) {
auto __group = __item.get_group();
Expand Down Expand Up @@ -315,9 +315,9 @@ __sycl_reduce_by_segment(__internal::__hetero_tag<_BackendTag>, _ExecutionPolicy
__cgh.depends_on(__seg_end_identification);
auto __seg_ends_acc = __seg_ends.template get_access<sycl::access_mode::read>(__cgh);
auto __seg_ends_scan_acc = __seg_ends_scanned.template get_access<sycl::access_mode::read_write>(__cgh);
#if _ONEDPL_COMPILE_KERNEL
# if _ONEDPL_COMPILE_KERNEL
__cgh.use_kernel_bundle(__seg_reduce_offset_kernel.get_kernel_bundle());
#endif
# endif
__cgh.parallel_for<_SegReduceOffsetKernel>(
sycl::nd_range<1>{__wgroup_size, __wgroup_size}, [=](sycl::nd_item<1> __item) {
auto __beg = __dpl_sycl::__get_accessor_ptr(__seg_ends_acc);
Expand All @@ -335,9 +335,9 @@ __sycl_reduce_by_segment(__internal::__hetero_tag<_BackendTag>, _ExecutionPolicy
auto __partials_acc = __partials.template get_access<sycl::access_mode::read_write>(__cgh);
auto __seg_ends_scan_acc = __seg_ends_scanned.template get_access<sycl::access_mode::read>(__cgh);
__dpl_sycl::__local_accessor<__val_type> __loc_acc(2 * __wgroup_size, __cgh);
#if _ONEDPL_COMPILE_KERNEL
# if _ONEDPL_COMPILE_KERNEL
__cgh.use_kernel_bundle(__seg_reduce_wg_kernel.get_kernel_bundle());
#endif
# endif
__cgh.parallel_for<_SegReduceWgKernel>(
sycl::nd_range<1>{__n_groups * __wgroup_size, __wgroup_size}, [=](sycl::nd_item<1> __item) {
::std::array<__val_type, __vals_per_item> __loc_partials;
Expand Down Expand Up @@ -455,9 +455,9 @@ __sycl_reduce_by_segment(__internal::__hetero_tag<_BackendTag>, _ExecutionPolicy
__dpl_sycl::__local_accessor<__diff_type> __loc_seg_ends_acc(__wgroup_size, __cgh);

__cgh.depends_on(__wg_reduce);
#if _ONEDPL_COMPILE_KERNEL
# if _ONEDPL_COMPILE_KERNEL
__cgh.use_kernel_bundle(__seg_reduce_prefix_kernel.get_kernel_bundle());
#endif
# endif
__cgh.parallel_for<_SegReducePrefixKernel>(
sycl::nd_range<1>{__n_groups * __wgroup_size, __wgroup_size}, [=](sycl::nd_item<1> __item) {
auto __group = __item.get_group();
Expand Down
8 changes: 4 additions & 4 deletions include/oneapi/dpl/internal/scan_by_segment_impl.h
Original file line number Diff line number Diff line change
Expand Up @@ -164,9 +164,9 @@ struct __sycl_scan_by_segment_impl

__dpl_sycl::__local_accessor<__val_type> __loc_acc(2 * __wgroup_size, __cgh);

#if _ONEDPL_COMPILE_KERNEL
# if _ONEDPL_COMPILE_KERNEL
__cgh.use_kernel_bundle(__seg_scan_wg_kernel.get_kernel_bundle());
#endif
# endif
__cgh.parallel_for<_SegScanWgKernel>(
sycl::nd_range<1>{__n_groups * __wgroup_size, __wgroup_size}, [=](sycl::nd_item<1> __item) {
__val_type __accumulator = __identity;
Expand Down Expand Up @@ -265,9 +265,9 @@ struct __sycl_scan_by_segment_impl
__dpl_sycl::__local_accessor<__val_type> __loc_partials_acc(__wgroup_size, __cgh);

__dpl_sycl::__local_accessor<bool> __loc_seg_ends_acc(__wgroup_size, __cgh);
#if _ONEDPL_COMPILE_KERNEL
# if _ONEDPL_COMPILE_KERNEL
__cgh.use_kernel_bundle(__seg_scan_prefix_kernel.get_kernel_bundle());
#endif
# endif
__cgh.parallel_for<_SegScanPrefixKernel>(
sycl::nd_range<1>{__n_groups * __wgroup_size, __wgroup_size}, [=](sycl::nd_item<1> __item) {
auto __group = __item.get_group();
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -596,7 +596,11 @@ __radix_sort_reorder_submit(_ExecutionPolicy&& __exec, ::std::size_t __segments,
if (__residual > 0)
{
//_ValueT may not have a default constructor, so we create just a storage via union type
union __storage { _ValueT __v; __storage(){} } __in_val;
union __storage
{
_ValueT __v;
__storage() {}
} __in_val;

::std::uint32_t __bucket = __radix_states; // greater than any actual radix state
if (__self_lidx < __residual)
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -267,7 +267,6 @@ class __kernel_compiler
static_assert(__kernel_count > 0, "At least one kernel name should be provided");

public:

template <typename _Exec>
static auto
__compile(_Exec&& __exec)
Expand Down
2 changes: 1 addition & 1 deletion include/oneapi/dpl/pstl/hetero/dpcpp/unseq_backend_sycl.h
Original file line number Diff line number Diff line change
Expand Up @@ -60,7 +60,7 @@ using __has_known_identity = ::std::conditional_t<
::std::is_same<::std::decay_t<_BinaryOp>, __dpl_sycl::__minimum<void>>,
::std::is_same<::std::decay_t<_BinaryOp>, __dpl_sycl::__maximum<_Tp>>,
::std::is_same<::std::decay_t<_BinaryOp>, __dpl_sycl::__maximum<void>>>>>,
::std::false_type>; // This is for the case of __can_use_known_identity<_Tp>==false
::std::false_type>; // This is for the case of __can_use_known_identity<_Tp>==false

#else //_USE_GROUP_ALGOS && defined(SYCL_IMPLEMENTATION_INTEL)

Expand Down

0 comments on commit 73ac0a5

Please sign in to comment.