diff --git a/include/oneapi/dpl/pstl/execution_impl.h b/include/oneapi/dpl/pstl/execution_impl.h index 79ed46570bf..c6f32f17495 100644 --- a/include/oneapi/dpl/pstl/execution_impl.h +++ b/include/oneapi/dpl/pstl/execution_impl.h @@ -106,25 +106,25 @@ __select_backend(oneapi::dpl::execution::parallel_unsequenced_policy, _IteratorT namespace __ranges { -::oneapi::dpl::__internal::__serial_tag +inline ::oneapi::dpl::__internal::__serial_tag __select_backend(oneapi::dpl::execution::sequenced_policy) { return {}; } -::oneapi::dpl::__internal::__serial_tag //vectorization allowed +inline ::oneapi::dpl::__internal::__serial_tag //vectorization allowed __select_backend(oneapi::dpl::execution::unsequenced_policy) { return {}; } -::oneapi::dpl::__internal::__parallel_tag +inline ::oneapi::dpl::__internal::__parallel_tag __select_backend(oneapi::dpl::execution::parallel_policy) { return {}; } -::oneapi::dpl::__internal::__parallel_tag //vectorization allowed +inline ::oneapi::dpl::__internal::__parallel_tag //vectorization allowed __select_backend(oneapi::dpl::execution::parallel_unsequenced_policy) { return {}; diff --git a/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_radix_sort_one_wg.h b/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_radix_sort_one_wg.h index 8bc09527384..6dd3b193a08 100644 --- a/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_radix_sort_one_wg.h +++ b/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_radix_sort_one_wg.h @@ -146,7 +146,6 @@ struct __subgroup_radix_sort auto operator()(sycl::queue __q, _RangeIn&& __src, _Proj __proj, _SLM_tag_val, _SLM_counter) { - constexpr std::uint16_t __req_sub_group_size = 16; uint16_t __n = __src.size(); assert(__n <= __block_size * __wg_size); @@ -164,9 +163,12 @@ struct __subgroup_radix_sort auto __counter_lacc = __buf_count.get_acc(__cgh); __cgh.parallel_for<_Name...>( - __range, - ([=](sycl::nd_item<1> __it)[[_ONEDPL_SYCL_REQD_SUB_GROUP_SIZE_IF_SUPPORTED(__req_sub_group_size)]] { - union __storage { _ValT __v[__block_size]; __storage(){} } __values; + __range, ([=](sycl::nd_item<1> __it) [[_ONEDPL_SYCL_REQD_SUB_GROUP_SIZE_IF_SUPPORTED(16)]] { + union __storage + { + _ValT __v[__block_size]; + __storage() {} + } __values; uint16_t __wi = __it.get_local_linear_id(); uint16_t __begin_bit = 0; constexpr uint16_t __end_bit = sizeof(_KeyT) * ::std::numeric_limits::digits;