Skip to content

Commit

Permalink
Remove free function __stride_recommender
Browse files Browse the repository at this point in the history
Signed-off-by: Matthew Michel <[email protected]>
  • Loading branch information
mmichel11 committed Nov 6, 2024
1 parent 8f9c5bb commit 55623b2
Showing 1 changed file with 0 additions and 33 deletions.
33 changes: 0 additions & 33 deletions include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_utils.h
Original file line number Diff line number Diff line change
Expand Up @@ -834,39 +834,6 @@ class __static_monotonic_dispatcher<::std::integer_sequence<::std::uint16_t, _X,
}
};

// Utility to recommend a stride for the best-performing memory access pattern from empirical testing on different
// devices. This utility can only be called from the device.
//
// SPIR-V compilation targets show best performance with a stride of the sub-group size.
// Other compilation targets perform best with a work-group size stride.
inline std::tuple<std::size_t, std::size_t, bool>
__stride_recommender(const sycl::nd_item<1>& __item, std::size_t __count, std::size_t __iters_per_work_item,
std::size_t __work_group_size)
{
if constexpr (oneapi::dpl::__internal::__is_spirv_target_v)
{
const __dpl_sycl::__sub_group __sub_group = __item.get_sub_group();
const std::uint32_t __sub_group_size = __sub_group.get_local_linear_range();
const std::uint32_t __sub_group_id = __sub_group.get_group_linear_id();
const std::uint32_t __sub_group_local_id = __sub_group.get_local_linear_id();
const std::size_t __work_group_id = __item.get_group().get_group_linear_id();

const std::size_t __sub_group_start_idx =
__iters_per_work_item * (__work_group_id * __work_group_size + __sub_group_size * __sub_group_id);
const bool __is_full_sub_group = __sub_group_start_idx + __iters_per_work_item * __sub_group_size <= __count;
const std::size_t __work_item_idx = __sub_group_start_idx + __sub_group_local_id;
return std::make_tuple(__work_item_idx, __sub_group_size, __is_full_sub_group);
}
else
{
const std::size_t __work_group_start_idx =
__item.get_group().get_group_linear_id() * __work_group_size * __iters_per_work_item;
const std::size_t __work_item_idx = __work_group_start_idx + __item.get_local_linear_id();
const bool __is_full_work_group = __work_group_start_idx + __iters_per_work_item * __work_group_size <= __count;
return std::make_tuple(__work_item_idx, __work_group_size, __is_full_work_group);
}
}

} // namespace __par_backend_hetero
} // namespace dpl
} // namespace oneapi
Expand Down

0 comments on commit 55623b2

Please sign in to comment.