algorithm/detail/find_extrema_with_atomics.hpp

0001 //---------------------------------------------------------------------------//
0002 // Copyright (c) 2013 Kyle Lutz <kyle.r.lutz@gmail.com>
0003 //
0004 // Distributed under the Boost Software License, Version 1.0
0005 // See accompanying file LICENSE_1_0.txt or copy at
0006 // http://www.boost.org/LICENSE_1_0.txt
0007 //
0008 // See http://boostorg.github.com/compute for more information.
0009 //---------------------------------------------------------------------------//
0010
0011 #ifndef BOOST_COMPUTE_ALGORITHM_DETAIL_FIND_EXTREMA_WITH_ATOMICS_HPP
0012 #define BOOST_COMPUTE_ALGORITHM_DETAIL_FIND_EXTREMA_WITH_ATOMICS_HPP
0013
0014 #include <boost/compute/types.hpp>
0015 #include <boost/compute/command_queue.hpp>
0016 #include <boost/compute/container/detail/scalar.hpp>
0017 #include <boost/compute/functional/atomic.hpp>
0018 #include <boost/compute/detail/meta_kernel.hpp>
0019 #include <boost/compute/detail/iterator_range_size.hpp>
0020
0021 namespace boost {
0022 namespace compute {
0023 namespace detail {
0024
0025 template<class InputIterator, class Compare>
0026 inline InputIterator find_extrema_with_atomics(InputIterator first,
0027                                                InputIterator last,
0028                                                Compare compare,
0029                                                const bool find_minimum,
0030                                                command_queue &queue)
0031 {
0032     typedef typename std::iterator_traits<InputIterator>::value_type value_type;
0033     typedef typename std::iterator_traits<InputIterator>::difference_type difference_type;
0034
0035     const context &context = queue.get_context();
0036
0037     meta_kernel k("find_extrema");
0038     atomic_cmpxchg<uint_> atomic_cmpxchg_uint;
0039
0040     k <<
0041         "const uint gid = get_global_id(0);\n" <<
0042         "uint old_index = *index;\n" <<
0043
0044         k.decl<value_type>("old") <<
0045             " = " << first[k.var<uint_>("old_index")] << ";\n" <<
0046         k.decl<value_type>("new") <<
0047             " = " << first[k.var<uint_>("gid")] << ";\n" <<
0048
0049         k.decl<bool>("compare_result") << ";\n" <<
0050         "#ifdef BOOST_COMPUTE_FIND_MAXIMUM\n" <<
0051         "while(" <<
0052             "(compare_result = " << compare(k.var<value_type>("old"),
0053                                             k.var<value_type>("new")) << ")" <<
0054             " || (!(compare_result" <<
0055                       " || " << compare(k.var<value_type>("new"),
0056                                         k.var<value_type>("old")) << ") "
0057                   "&& gid < old_index)){\n" <<
0058         "#else\n" <<
0059         // while condition explained for minimum case with less (<)
0060         // as comparison function:
0061         // while(new_value < old_value
0062         //       OR (new_value == old_value AND new_index < old_index))
0063         "while(" <<
0064             "(compare_result = " << compare(k.var<value_type>("new"),
0065                                             k.var<value_type>("old"))  << ")" <<
0066             " || (!(compare_result" <<
0067                       " || " << compare(k.var<value_type>("old"),
0068                                         k.var<value_type>("new")) << ") "
0069                   "&& gid < old_index)){\n" <<
0070         "#endif\n" <<
0071
0072         "  if(" << atomic_cmpxchg_uint(k.var<uint_ *>("index"),
0073                                        k.var<uint_>("old_index"),
0074                                        k.var<uint_>("gid")) << " == old_index)\n" <<
0075         "      break;\n" <<
0076         "  else\n" <<
0077         "    old_index = *index;\n" <<
0078         "old = " << first[k.var<uint_>("old_index")] << ";\n" <<
0079         "}\n";
0080
0081     size_t index_arg_index = k.add_arg<uint_ *>(memory_object::global_memory, "index");
0082
0083     std::string options;
0084     if(!find_minimum){
0085         options = "-DBOOST_COMPUTE_FIND_MAXIMUM";
0086     }
0087     kernel kernel = k.compile(context, options);
0088
0089     // setup index buffer
0090     scalar<uint_> index(context);
0091     kernel.set_arg(index_arg_index, index.get_buffer());
0092
0093     // initialize index
0094     index.write(0, queue);
0095
0096     // run kernel
0097     size_t count = iterator_range_size(first, last);
0098     queue.enqueue_1d_range_kernel(kernel, 0, count, 0);
0099
0100     // read index and return iterator
0101     return first + static_cast<difference_type>(index.read(queue));
0102 }
0103
0104 } // end detail namespace
0105 } // end compute namespace
0106 } // end boost namespace
0107
0108 #endif // BOOST_COMPUTE_ALGORITHM_DETAIL_FIND_EXTREMA_WITH_ATOMICS_HPP