From acc4f2048bfd1eb4a036077f4fbb7ed2d195cd41 Mon Sep 17 00:00:00 2001 From: Sergey Kopienko Date: Thu, 14 Nov 2024 14:23:26 +0100 Subject: [PATCH] include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h - remove usage of __kernel_name_generator as not required Signed-off-by: Sergey Kopienko --- .../pstl/hetero/dpcpp/parallel_backend_sycl.h | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) diff --git a/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h b/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h index 77b2e6c35f..08497a7ce4 100644 --- a/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h +++ b/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h @@ -1756,12 +1756,6 @@ __parallel_find_or(oneapi::dpl::__internal::__device_backend_tag, _ExecutionPoli _BrickTag __brick_tag, _Ranges&&... __rngs) { using _CustomName = oneapi::dpl::__internal::__policy_kernel_name<_ExecutionPolicy>; - using _FindOrKernelOneWG = - oneapi::dpl::__par_backend_hetero::__internal::__kernel_name_generator<__find_or_kernel_one_wg, _CustomName, - _Brick, _BrickTag, _Ranges...>; - using _FindOrKernel = - oneapi::dpl::__par_backend_hetero::__internal::__kernel_name_generator<__find_or_kernel, _CustomName, _Brick, - _BrickTag, _Ranges...>; auto __rng_n = oneapi::dpl::__ranges::__get_first_range_size(__rngs...); assert(__rng_n > 0); @@ -1784,8 +1778,11 @@ __parallel_find_or(oneapi::dpl::__internal::__device_backend_tag, _ExecutionPoli // We shouldn't have any restrictions for _AtomicType type here // because we have a single work-group and we don't need to use atomics for inter-work-group communication. + using _KernelName = oneapi::dpl::__par_backend_hetero::__internal::__kernel_name_provider< + __find_or_kernel_one_wg<_CustomName, _Brick, _BrickTag, _Ranges...>>; + // Single WG implementation - __result = __parallel_find_or_impl_one_wg<_FindOrKernelOneWG, __or_tag_check>( + __result = __parallel_find_or_impl_one_wg<_KernelName, __or_tag_check>( oneapi::dpl::__internal::__device_backend_tag{}, std::forward<_ExecutionPolicy>(__exec), __brick_tag, __rng_n, __wgroup_size, __init_value, __pred, std::forward<_Ranges>(__rngs)...); } @@ -1794,8 +1791,11 @@ __parallel_find_or(oneapi::dpl::__internal::__device_backend_tag, _ExecutionPoli assert("This device does not support 64-bit atomics" && (sizeof(_AtomicType) < 8 || __exec.queue().get_device().has(sycl::aspect::atomic64))); + using _KernelName = oneapi::dpl::__par_backend_hetero::__internal::__kernel_name_provider< + __find_or_kernel<_CustomName, _Brick, _BrickTag, _Ranges...>>; + // Multiple WG implementation - __result = __parallel_find_or_impl_multiple_wgs<_FindOrKernel, __or_tag_check>( + __result = __parallel_find_or_impl_multiple_wgs<_KernelName, __or_tag_check>( oneapi::dpl::__internal::__device_backend_tag{}, std::forward<_ExecutionPolicy>(__exec), __brick_tag, __rng_n, __n_groups, __wgroup_size, __init_value, __pred, std::forward<_Ranges>(__rngs)...); }