Skip to content

Commit

Permalink
Revert "@@@ include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_syc…
Browse files Browse the repository at this point in the history 
…l_merge.h - debug code under DUMP_DATA_LOADING"

This reverts commit 952871e.
  • Loading branch information
@SergeyKopienko
SergeyKopienko committed Nov 20, 2024
1 parent 3dced51 commit 39b68e4
Showing 1 changed file with 0 additions and 63 deletions.
63 changes: 0 additions & 63 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 @@ -26,8 +26,6 @@
#include "sycl_defs.h"
#include "parallel_backend_sycl_utils.h"

//#define DUMP_DATA_LOADING 1

namespace oneapi
{
namespace dpl
Expand Down Expand Up @@ -282,20 +280,6 @@ struct __parallel_merge_submitter_large<_IdType, _CustomName,
__internal::__optional_kernel_name<_DiagonalsKernelName...>,
__internal::__optional_kernel_name<_MergeKernelName...>>
{
#if DUMP_DATA_LOADING
template <typename _Range, typename _Index, typename _Data>
static void
__load_item_into_slm(_Range&& __rng, _Index __idx_from, _Data* __slm, _Index __idx_to, std::size_t __range_index,
bool __b_check, std::size_t __group_linear_id, std::size_t __local_id)
{
// BP
// condition: __b_check
// action: __range_index = {__range_index}, __rng[{__idx_from}] -> __slm[{__idx_to}], __group_linear_id = {__group_linear_id}, __local_id = {__local_id}
// action: {__range_index}, {__idx_from}, {__idx_to}, {__group_linear_id}, {__local_id}
__slm[__idx_to] = __rng[__idx_from];
}
#endif

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 @@ -310,18 +294,6 @@ struct __parallel_merge_submitter_large<_IdType, _CustomName,
const _IdType __n2 = __rng2.size();
const _IdType __n = __n1 + __n2;

#if DUMP_DATA_LOADING
//const bool __b_check = __n1 == 16144 && __n2 == 8072;
//const bool __b_check = __n1 == 50716 && __n2 == 25358; // __wi_in_one_wg = 51 __wg_count = 12
const bool __b_check = false;

if (__b_check)
{
int i = 0;
i = i;
}
#endif

assert(__n1 > 0 || __n2 > 0);

_PRINT_INFO_IN_DEBUG_MODE(__exec);
Expand Down Expand Up @@ -447,11 +419,7 @@ struct __parallel_merge_submitter_large<_IdType, _CustomName,

_ONEDPL_PRAGMA_UNROLL
for (_IdType __idx = __idx_begin; __idx < __idx_end; ++__idx)
#if !DUMP_DATA_LOADING
__rng1_cache_slm[__idx] = __rng1[__sp_base_left_global.first + __idx];
#else
__load_item_into_slm(__rng1, __sp_base_left_global.first + __idx, __rng1_cache_slm, __idx, 1, __b_check, __group_linear_id, __local_id);
#endif
}
}

Expand All @@ -467,44 +435,13 @@ struct __parallel_merge_submitter_large<_IdType, _CustomName,

_ONEDPL_PRAGMA_UNROLL
for (_IdType __idx = __idx_begin; __idx < __idx_end; ++__idx)
#if !DUMP_DATA_LOADING
__rng2_cache_slm[__idx] = __rng2[__sp_base_left_global.second + __idx];
#else
__load_item_into_slm(__rng2, __sp_base_left_global.second + __idx, __rng2_cache_slm, __idx, 2, __b_check, __group_linear_id, __local_id);
#endif
}
}

// Wait until all the data is loaded
__dpl_sycl::__group_barrier(__nd_item);

#if DUMP_DATA_LOADING
if (__local_id == 0)
{
for (auto i = __sp_base_left_global.first; i < __sp_base_right_global.first; ++i)
{
auto _idx_slm = i - __sp_base_left_global.first;
if (__rng1_cache_slm[_idx_slm] != __rng1[i])
{
auto __group_linear_id_tmp = __group_linear_id;
__group_linear_id_tmp = __group_linear_id_tmp;
assert(false);
}
}

for (auto i = __sp_base_left_global.second; i < __sp_base_right_global.second; ++i)
{
auto _idx_slm = i - __sp_base_left_global.second;
if (__rng2_cache_slm[_idx_slm] != __rng2[i])
{
auto __group_linear_id_tmp = __group_linear_id;
__group_linear_id_tmp = __group_linear_id_tmp;
assert(false);
}
}
}
#endif

// Current diagonal inside of the merge matrix?
if (__global_linear_id * __chunk < __n)
{
Expand Down

0 comments on commit 39b68e4

Please sign in to comment.