From 55623b23aab00449ef2b248e5a6659862dea7c7d Mon Sep 17 00:00:00 2001 From: Matthew Michel Date: Wed, 25 Sep 2024 09:19:31 -0500 Subject: [PATCH] Remove free function __stride_recommender Signed-off-by: Matthew Michel --- .../dpcpp/parallel_backend_sycl_utils.h | 33 ------------------- 1 file changed, 33 deletions(-) diff --git a/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_utils.h b/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_utils.h index a3ce6390d3..e7fbfb7ae7 100644 --- a/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_utils.h +++ b/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_utils.h @@ -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 -__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