Replace __reduce_over_group to __any_of_group for __parallel_or_tag

Signed-off-by: Sergey Kopienko <sergey.kopienko@intel.com>
uxlfoundation · Jul 10, 2024 · 678ad57 · 678ad57
1 parent 49ce3bb
commit 678ad57
Show file tree

Hide file tree

Showing 2 changed files with 6 additions and 11 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
@@ -1035,8 +1035,6 @@ struct __parallel_or_tag
 {
     using _AtomicType = int32_t;
 
-    using _LocalResultsReduceOp = __dpl_sycl::__bit_or<_AtomicType>;
-
     // The template parameter is intended to unify __init_value in tags.
     template <typename _DiffType>
     constexpr static _AtomicType __init_value(_DiffType)
@@ -1233,10 +1231,13 @@ __parallel_find_or(oneapi::dpl::__internal::__device_backend_tag, _ExecutionPoli
 
                     // 3. Reduce over group: find __dpl_sycl::__minimum (for the __parallel_find_forward_tag),
                     // find __dpl_sycl::__maximum (for the __parallel_find_backward_tag)
-                    // or update state with __dpl_sycl::__bit_or (for the __parallel_or_tag)
+                    // or update state with __dpl_sycl::__any_of_group (for the __parallel_or_tag)
                     // inside all our group items
-                    __found_local = __dpl_sycl::__reduce_over_group(__item_id.get_group(), __found_local,
-                                                                    typename _BrickTag::_LocalResultsReduceOp{});
+                    if constexpr (__or_tag_check)
+                        __found_local = __dpl_sycl::__any_of_group(__item_id.get_group(), __found_local);
+                    else
+                        __found_local = __dpl_sycl::__reduce_over_group(__item_id.get_group(), __found_local,
+                                                                        typename _BrickTag::_LocalResultsReduceOp{});
 
                     // Set local found state value value to global atomic
                     if (__local_idx == 0 && __found_local != __init_value)

diff --git a/include/oneapi/dpl/pstl/hetero/dpcpp/sycl_defs.h b/include/oneapi/dpl/pstl/hetero/dpcpp/sycl_defs.h
@@ -129,9 +129,6 @@ using __maximum = sycl::maximum<_T>;
 
 template <typename _T = void>
 using __minimum = sycl::minimum<_T>;
-
-template <typename _T = void>
-using __bit_or = sycl::bit_or<_T>;
 #else  // _ONEDPL_SYCL2020_FUNCTIONAL_OBJECTS_PRESENT
 template <typename _T>
 using __plus = sycl::ONEAPI::plus<_T>;
@@ -141,9 +138,6 @@ using __maximum = sycl::ONEAPI::maximum<_T>;
 
 template <typename _T>
 using __minimum = sycl::ONEAPI::minimum<_T>;
-
-template <typename _T = void>
-using __bit_or = sycl::ONEAPI::bit_or<_T>;
 #endif // _ONEDPL_SYCL2020_FUNCTIONAL_OBJECTS_PRESENT
 
 template <typename _Buffer>