From 8099e07f40f2f8a6f678f3e2b2b209461b976c15 Mon Sep 17 00:00:00 2001
From: Jakub Szuppe <j.szuppe@gmail.com>
Date: Sat, 3 Oct 2015 16:07:18 +0200
Subject: [PATCH] Fix find_extrema_with_atomics

Now find_extrema_with_atomics always return the first extremum
in the given range.
---
 .../detail/find_extrema_with_atomics.hpp      | 32 ++++++++++++++++---
 1 file changed, 27 insertions(+), 5 deletions(-)

diff --git a/include/boost/compute/algorithm/detail/find_extrema_with_atomics.hpp b/include/boost/compute/algorithm/detail/find_extrema_with_atomics.hpp
index 5ebaf835..406d1bec 100644
--- a/include/boost/compute/algorithm/detail/find_extrema_with_atomics.hpp
+++ b/include/boost/compute/algorithm/detail/find_extrema_with_atomics.hpp
@@ -29,6 +29,7 @@ inline InputIterator find_extrema_with_atomics(InputIterator first,
                                                const bool find_minimum,
                                                command_queue &queue)
 {
+    typedef typename std::iterator_traits<InputIterator>::value_type value_type;
     typedef typename std::iterator_traits<InputIterator>::difference_type difference_type;
 
     const context &context = queue.get_context();
@@ -40,12 +41,32 @@ inline InputIterator find_extrema_with_atomics(InputIterator first,
         "const uint gid = get_global_id(0);\n" <<
         "uint old_index = *index;\n" <<
 
-        "#ifndef BOOST_COMPUTE_FIND_MAXIMUM\n" <<
-        "while(" << compare(first[k.var<uint_>("gid")],
-                            first[k.var<uint_>("old_index")]) << "){\n" <<
+        k.decl<value_type>("old") <<
+            " = " << first[k.var<uint_>("old_index")] << ";\n" <<
+        k.decl<value_type>("new") <<
+            " = " << first[k.var<uint_>("gid")] << ";\n" <<
+
+        k.decl<bool>("compare_result") << ";\n" <<
+        "#ifdef BOOST_COMPUTE_FIND_MAXIMUM\n" <<
+        "while(" <<
+            "(compare_result = " << compare(k.var<value_type>("old"),
+                                            k.var<value_type>("new")) << ")" <<
+            " || (!(compare_result" <<
+                      " || " << compare(k.var<value_type>("new"),
+                                        k.var<value_type>("old")) << ") "
+                  "&& gid < old_index)){\n" <<
         "#else\n" <<
-        "while(" << compare(first[k.var<uint_>("old_index")],
-                            first[k.var<uint_>("gid")]) << "){\n" <<
+        // while condition explained for minimum case with less (<)
+        // as comparison function:
+        // while(new_value < old_value
+        //       OR (new_value == old_value AND new_index < old_index))
+        "while(" <<
+            "(compare_result = " << compare(k.var<value_type>("new"),
+                                            k.var<value_type>("old"))  << ")" <<
+            " || (!(compare_result" <<
+                      " || " << compare(k.var<value_type>("old"),
+                                        k.var<value_type>("new")) << ") "
+                  "&& gid < old_index)){\n" <<
         "#endif\n" <<
 
         "  if(" << atomic_cmpxchg_uint(k.var<uint_ *>("index"),
@@ -54,6 +75,7 @@ inline InputIterator find_extrema_with_atomics(InputIterator first,
         "      break;\n" <<
         "  else\n" <<
         "    old_index = *index;\n" <<
+        "old = " << first[k.var<uint_>("old_index")] << ";\n" <<
         "}\n";
 
     size_t index_arg_index = k.add_arg<uint_ *>(memory_object::global_memory, "index");