1 //---------------------------------------------------------------------------//
2 // Copyright (c) 2013 Kyle Lutz <kyle.r.lutz@gmail.com>
3 //
4 // Distributed under the Boost Software License, Version 1.0
5 // See accompanying file LICENSE_1_0.txt or copy at
6 // http://www.boost.org/LICENSE_1_0.txt
7 //
8 // See http://boostorg.github.com/compute for more information.
9 //---------------------------------------------------------------------------//
10
11 #ifndef BOOST_COMPUTE_ALGORITHM_DETAIL_FIND_EXTREMA_WITH_ATOMICS_HPP
12 #define BOOST_COMPUTE_ALGORITHM_DETAIL_FIND_EXTREMA_WITH_ATOMICS_HPP
13
14 #include <boost/compute/types.hpp>
15 #include <boost/compute/command_queue.hpp>
16 #include <boost/compute/container/detail/scalar.hpp>
17 #include <boost/compute/functional/atomic.hpp>
18 #include <boost/compute/detail/meta_kernel.hpp>
19 #include <boost/compute/detail/iterator_range_size.hpp>
20
21 namespace boost {
22 namespace compute {
23 namespace detail {
24
25 template<class InputIterator, class Compare>
find_extrema_with_atomics(InputIterator first,InputIterator last,Compare compare,const bool find_minimum,command_queue & queue)26 inline InputIterator find_extrema_with_atomics(InputIterator first,
27 InputIterator last,
28 Compare compare,
29 const bool find_minimum,
30 command_queue &queue)
31 {
32 typedef typename std::iterator_traits<InputIterator>::value_type value_type;
33 typedef typename std::iterator_traits<InputIterator>::difference_type difference_type;
34
35 const context &context = queue.get_context();
36
37 meta_kernel k("find_extrema");
38 atomic_cmpxchg<uint_> atomic_cmpxchg_uint;
39
40 k <<
41 "const uint gid = get_global_id(0);\n" <<
42 "uint old_index = *index;\n" <<
43
44 k.decl<value_type>("old") <<
45 " = " << first[k.var<uint_>("old_index")] << ";\n" <<
46 k.decl<value_type>("new") <<
47 " = " << first[k.var<uint_>("gid")] << ";\n" <<
48
49 k.decl<bool>("compare_result") << ";\n" <<
50 "#ifdef BOOST_COMPUTE_FIND_MAXIMUM\n" <<
51 "while(" <<
52 "(compare_result = " << compare(k.var<value_type>("old"),
53 k.var<value_type>("new")) << ")" <<
54 " || (!(compare_result" <<
55 " || " << compare(k.var<value_type>("new"),
56 k.var<value_type>("old")) << ") "
57 "&& gid < old_index)){\n" <<
58 "#else\n" <<
59 // while condition explained for minimum case with less (<)
60 // as comparison function:
61 // while(new_value < old_value
62 // OR (new_value == old_value AND new_index < old_index))
63 "while(" <<
64 "(compare_result = " << compare(k.var<value_type>("new"),
65 k.var<value_type>("old")) << ")" <<
66 " || (!(compare_result" <<
67 " || " << compare(k.var<value_type>("old"),
68 k.var<value_type>("new")) << ") "
69 "&& gid < old_index)){\n" <<
70 "#endif\n" <<
71
72 " if(" << atomic_cmpxchg_uint(k.var<uint_ *>("index"),
73 k.var<uint_>("old_index"),
74 k.var<uint_>("gid")) << " == old_index)\n" <<
75 " break;\n" <<
76 " else\n" <<
77 " old_index = *index;\n" <<
78 "old = " << first[k.var<uint_>("old_index")] << ";\n" <<
79 "}\n";
80
81 size_t index_arg_index = k.add_arg<uint_ *>(memory_object::global_memory, "index");
82
83 std::string options;
84 if(!find_minimum){
85 options = "-DBOOST_COMPUTE_FIND_MAXIMUM";
86 }
87 kernel kernel = k.compile(context, options);
88
89 // setup index buffer
90 scalar<uint_> index(context);
91 kernel.set_arg(index_arg_index, index.get_buffer());
92
93 // initialize index
94 index.write(0, queue);
95
96 // run kernel
97 size_t count = iterator_range_size(first, last);
98 queue.enqueue_1d_range_kernel(kernel, 0, count, 0);
99
100 // read index and return iterator
101 return first + static_cast<difference_type>(index.read(queue));
102 }
103
104 } // end detail namespace
105 } // end compute namespace
106 } // end boost namespace
107
108 #endif // BOOST_COMPUTE_ALGORITHM_DETAIL_FIND_EXTREMA_WITH_ATOMICS_HPP
109