Skip to content

Commit

Permalink
Merge branch 'dev/skopienko/serial_merge_without_unroll' into dev/sko…
Browse files Browse the repository at this point in the history 
…pienko/serial_merge
  • Loading branch information
@SergeyKopienko
SergeyKopienko committed Dec 18, 2024
2 parents 726bd2e + ab54931 commit 764a494
Show file tree
Hide file tree
Showing 2 changed files with 9 additions and 15 deletions.
12 changes: 4 additions & 8 deletions include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_merge.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -80,8 +80,7 @@ __find_start_point(const _Rng1& __rng1, const _Rng2& __rng2, const _Index __i_el

// Do serial merge of the data from rng1 (starting from start1) and rng2 (starting from start2) and writing
// to rng3 (starting from start3) in 'chunk' steps, but do not exceed the total size of the sequences (n1 and n2)
template <unsigned int _UnrollFactor = 4, typename _Rng1, typename _Rng2, typename _Rng3, typename _Index,
typename _Compare>
template <typename _Rng1, typename _Rng2, typename _Rng3, typename _Index, typename _Compare>
void
__serial_merge(const _Rng1& __rng1, const _Rng2& __rng2, _Rng3& __rng3, _Index __start1, _Index __start2,
const _Index __start3, const std::uint8_t __chunk, const _Index __n1, const _Index __n2, _Compare __comp)
Expand All @@ -98,7 +97,6 @@ __serial_merge(const _Rng1& __rng1, const _Rng2& __rng2, _Rng3& __rng3, _Index _
_Index __rng2_idx = __start2;
_Index __rng3_idx = __start3;

#pragma unroll _UnrollFactor
for (_Index __rng3_idx = __start3; __rng3_idx < __rng3_idx_end; ++__rng3_idx)
{
const bool __rng1_idx_less__n1 = __rng1_idx < __rng1_idx_end;
Expand All @@ -118,8 +116,6 @@ struct __parallel_merge_submitter;
template <typename _IdType, typename... _Name>
struct __parallel_merge_submitter<_IdType, __internal::__optional_kernel_name<_Name...>>
{
static constexpr std::uint32_t __gpu_chunk = 4;

template <typename _ExecutionPolicy, typename _Range1, typename _Range2, typename _Range3, typename _Compare>
auto
operator()(_ExecutionPolicy&& __exec, _Range1&& __rng1, _Range2&& __rng2, _Range3&& __rng3, _Compare __comp) const
Expand All @@ -133,7 +129,7 @@ struct __parallel_merge_submitter<_IdType, __internal::__optional_kernel_name<_N
_PRINT_INFO_IN_DEBUG_MODE(__exec);

// Empirical number of values to process per work-item
const std::uint8_t __chunk = __exec.queue().get_device().is_cpu() ? 128 : __gpu_chunk;
const std::uint8_t __chunk = __exec.queue().get_device().is_cpu() ? 128 : 4;

const _IdType __steps = oneapi::dpl::__internal::__dpl_ceiling_div(__n, __chunk);

Expand All @@ -142,8 +138,8 @@ struct __parallel_merge_submitter<_IdType, __internal::__optional_kernel_name<_N
__cgh.parallel_for<_Name...>(sycl::range</*dim=*/1>(__steps), [=](sycl::item</*dim=*/1> __item_id) {
const _IdType __i_elem = __item_id.get_linear_id() * __chunk;
const auto __start = __find_start_point(__rng1, __rng2, __i_elem, __n1, __n2, __comp);
__serial_merge<__gpu_chunk>(__rng1, __rng2, __rng3, __start.first, __start.second, __i_elem, __chunk,
__n1, __n2, __comp);
__serial_merge(__rng1, __rng2, __rng3, __start.first, __start.second, __i_elem, __chunk, __n1, __n2,
__comp);
});
});
return __future(__event);
Expand Down
12 changes: 5 additions & 7 deletions include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_merge_sort.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -233,8 +233,6 @@ struct __merge_sort_global_submitter;
template <typename _IndexT, typename... _GlobalSortName>
struct __merge_sort_global_submitter<_IndexT, __internal::__optional_kernel_name<_GlobalSortName...>>
{
static constexpr std::uint32_t __gpu_chunk = 4;

template <typename _Range, typename _Compare, typename _TempBuf, typename _LeafSizeT>
std::pair<sycl::event, bool>
operator()(sycl::queue& __q, _Range& __rng, _Compare __comp, _LeafSizeT __leaf_size, _TempBuf& __temp_buf,
Expand All @@ -243,7 +241,7 @@ struct __merge_sort_global_submitter<_IndexT, __internal::__optional_kernel_name
const _IndexT __n = __rng.size();
_IndexT __n_sorted = __leaf_size;
const bool __is_cpu = __q.get_device().is_cpu();
const std::uint32_t __chunk = __is_cpu ? 32 : __gpu_chunk;
const std::uint32_t __chunk = __is_cpu ? 32 : 4;
const std::size_t __steps = oneapi::dpl::__internal::__dpl_ceiling_div(__n, __chunk);
bool __data_in_temp = false;

Expand Down Expand Up @@ -274,17 +272,17 @@ struct __merge_sort_global_submitter<_IndexT, __internal::__optional_kernel_name
const oneapi::dpl::__ranges::drop_view_simple __rng2(__dst, __offset + __n1);

const auto start = __find_start_point(__rng1, __rng2, __i_elem_local, __n1, __n2, __comp);
__serial_merge<__gpu_chunk>(__rng1, __rng2, __rng /*__rng3*/, start.first, start.second,
__i_elem, __chunk, __n1, __n2, __comp);
__serial_merge(__rng1, __rng2, __rng /*__rng3*/, start.first, start.second, __i_elem,
__chunk, __n1, __n2, __comp);
}
else
{
const oneapi::dpl::__ranges::drop_view_simple __rng1(__rng, __offset);
const oneapi::dpl::__ranges::drop_view_simple __rng2(__rng, __offset + __n1);

const auto start = __find_start_point(__rng1, __rng2, __i_elem_local, __n1, __n2, __comp);
__serial_merge<__gpu_chunk>(__rng1, __rng2, __dst /*__rng3*/, start.first, start.second,
__i_elem, __chunk, __n1, __n2, __comp);
__serial_merge(__rng1, __rng2, __dst /*__rng3*/, start.first, start.second, __i_elem,
__chunk, __n1, __n2, __comp);
}
});
});
Expand Down

0 comments on commit 764a494

Please sign in to comment.