[Perf] Break loop in __early_exit_find_or after first found element

include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h

@@ -1035,6 +1035,11 @@ struct __early_exit_find_or
                _LocalAtomic& __found_local, _BrickTag, _Ranges&&... __rngs) const
     {
         using __par_backend_hetero::__parallel_or_tag;
+
+        // There is 3 possible tag types here:
+        //  - __parallel_find_forward_tag : in case when we find the first value in the data;
+        //  - __parallel_find_backward_tag : in case when we find the last value in the data;
+        //  - __parallel_or_tag : in case when we find any value in the data.
         using _OrTagType = ::std::is_same<_BrickTag, __par_backend_hetero::__parallel_or_tag>;
         using _BackwardTagType = ::std::is_same<typename _BrickTag::_Compare, oneapi::dpl::__internal::__pstl_greater>;
 
@@ -1068,16 +1073,21 @@ struct __early_exit_find_or
             {
                 if constexpr (_OrTagType::value)
                 {
+                    // As far as we find any value of the data here, we can set the atomic value to 1
+                    // No additional actions required.
                     __found_local.store(1);
-                    break;
                 }
                 else
                 {
-                    for (auto __old = __found_local.load(); __comp(__shifted_idx, __old); __old = __found_local.load())
-                    {
-                        __found_local.compare_exchange_strong(__old, __shifted_idx);
-                    }
+                    // As far as we find the first/last value of the data here, we need to set the atomic value to the index of the found data.
+                    // But only in this case when this value is less (if we find the first value)/greater(if we find the last value) than the current value of the atomic.
+                    if constexpr (!_BackwardTagType::value)
+                        __found_local.fetch_min(__shifted_idx);
+                    else
+                        __found_local.fetch_max(__shifted_idx);
                 }
+
+                break;
             }
         }
     }
@@ -1101,7 +1111,9 @@ __parallel_find_or(oneapi::dpl::__internal::__device_backend_tag, _ExecutionPoli
         oneapi::dpl::__par_backend_hetero::__internal::__kernel_name_generator<__find_or_kernel, _CustomName, _Brick,
                                                                                _BrickTag, _Ranges...>;
 
-    constexpr bool __or_tag_check = ::std::is_same_v<_BrickTag, __parallel_or_tag>;
+    constexpr bool __or_tag_check = std::is_same_v<_BrickTag, __parallel_or_tag>;
+    using _BackwardTagType = std::is_same<typename _BrickTag::_Compare, oneapi::dpl::__internal::__pstl_greater>;
+
     auto __rng_n = oneapi::dpl::__ranges::__get_first_range_size(__rngs...);
     assert(__rng_n > 0);
 
@@ -1165,16 +1177,20 @@ __parallel_find_or(oneapi::dpl::__internal::__device_backend_tag, _ExecutionPoli
                     __dpl_sycl::__group_barrier(__item_id);
 
                     // Set local atomic value to global atomic
-                    if (__local_idx == 0 && __comp(__found_local.load(), __found.load()))
+                    if (__local_idx == 0)
                     {
-                        if constexpr (__or_tag_check)
-                            __found.store(1);
-                        else
+                        const auto __found_local_state = __found_local.load();
+
+                        if (__comp(__found_local_state, __found.load()))
                         {
-                            for (auto __old = __found.load(); __comp(__found_local.load(), __old);
-                                 __old = __found.load())
+                            if constexpr (__or_tag_check)
+                                __found.store(1);
+                            else
                             {
-                                __found.compare_exchange_strong(__old, __found_local.load());
+                                if constexpr (!_BackwardTagType::value)
+                                    __found.fetch_min(__found_local_state);
+                                else
+                                    __found.fetch_max(__found_local_state);
                             }
                         }
                     }