]>
Commit | Line | Data |
---|---|---|
7c673cae FG |
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> | |
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 |