]>
Commit | Line | Data |
---|---|---|
7c673cae FG |
1 | //---------------------------------------------------------------------------// |
2 | // Copyright (c) 2014 Roshan <thisisroshansmail@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_FIND_END_HPP | |
12 | #define BOOST_COMPUTE_ALGORITHM_FIND_END_HPP | |
13 | ||
14 | #include <boost/compute/algorithm/copy.hpp> | |
15 | #include <boost/compute/algorithm/detail/search_all.hpp> | |
16 | #include <boost/compute/container/detail/scalar.hpp> | |
17 | #include <boost/compute/container/vector.hpp> | |
18 | #include <boost/compute/detail/iterator_range_size.hpp> | |
19 | #include <boost/compute/detail/meta_kernel.hpp> | |
20 | #include <boost/compute/system.hpp> | |
21 | ||
22 | namespace boost { | |
23 | namespace compute { | |
24 | namespace detail { | |
25 | ||
26 | /// | |
27 | /// \brief Helper function for find_end | |
28 | /// | |
29 | /// Basically a copy of find_if which returns last occurence | |
30 | /// instead of first occurence | |
31 | /// | |
32 | template<class InputIterator, class UnaryPredicate> | |
33 | inline InputIterator find_end_helper(InputIterator first, | |
34 | InputIterator last, | |
35 | UnaryPredicate predicate, | |
36 | command_queue &queue) | |
37 | { | |
38 | typedef typename std::iterator_traits<InputIterator>::value_type value_type; | |
39 | typedef typename std::iterator_traits<InputIterator>::difference_type difference_type; | |
40 | ||
41 | size_t count = detail::iterator_range_size(first, last); | |
42 | if(count == 0){ | |
43 | return last; | |
44 | } | |
45 | ||
46 | const context &context = queue.get_context(); | |
47 | ||
48 | detail::meta_kernel k("find_end"); | |
49 | size_t index_arg = k.add_arg<int *>(memory_object::global_memory, "index"); | |
50 | atomic_max<int_> atomic_max_int; | |
51 | ||
52 | k << k.decl<const int_>("i") << " = get_global_id(0);\n" | |
53 | << k.decl<const value_type>("value") << "=" | |
54 | << first[k.var<const int_>("i")] << ";\n" | |
55 | << "if(" << predicate(k.var<const value_type>("value")) << "){\n" | |
56 | << " " << atomic_max_int(k.var<int_ *>("index"), k.var<int_>("i")) << ";\n" | |
57 | << "}\n"; | |
58 | ||
59 | kernel kernel = k.compile(context); | |
60 | ||
61 | scalar<int_> index(context); | |
62 | kernel.set_arg(index_arg, index.get_buffer()); | |
63 | ||
64 | index.write(static_cast<int_>(-1), queue); | |
65 | ||
66 | queue.enqueue_1d_range_kernel(kernel, 0, count, 0); | |
67 | ||
68 | int result = static_cast<int>(index.read(queue)); | |
69 | ||
70 | if(result == -1){ | |
71 | return last; | |
72 | } | |
73 | else { | |
74 | return first + static_cast<difference_type>(result); | |
75 | } | |
76 | } | |
77 | ||
78 | } // end detail namespace | |
79 | ||
80 | /// | |
81 | /// \brief Substring matching algorithm | |
82 | /// | |
83 | /// Searches for the last match of the pattern [p_first, p_last) | |
84 | /// in text [t_first, t_last). | |
85 | /// \return Iterator pointing to beginning of last occurence | |
86 | /// | |
87 | /// \param t_first Iterator pointing to start of text | |
88 | /// \param t_last Iterator pointing to end of text | |
89 | /// \param p_first Iterator pointing to start of pattern | |
90 | /// \param p_last Iterator pointing to end of pattern | |
91 | /// \param queue Queue on which to execute | |
92 | /// | |
93 | template<class TextIterator, class PatternIterator> | |
94 | inline TextIterator find_end(TextIterator t_first, | |
95 | TextIterator t_last, | |
96 | PatternIterator p_first, | |
97 | PatternIterator p_last, | |
98 | command_queue &queue = system::default_queue()) | |
99 | { | |
100 | const context &context = queue.get_context(); | |
101 | ||
102 | // there is no need to check if pattern starts at last n - 1 indices | |
103 | vector<uint_> matching_indices( | |
104 | detail::iterator_range_size(t_first, t_last) | |
105 | + 1 - detail::iterator_range_size(p_first, p_last), | |
106 | context | |
107 | ); | |
108 | ||
109 | detail::search_kernel<PatternIterator, | |
110 | TextIterator, | |
111 | vector<uint_>::iterator> kernel; | |
112 | ||
113 | kernel.set_range(p_first, p_last, t_first, t_last, matching_indices.begin()); | |
114 | kernel.exec(queue); | |
115 | ||
116 | using boost::compute::_1; | |
117 | ||
118 | vector<uint_>::iterator index = | |
119 | detail::find_end_helper( | |
120 | matching_indices.begin(), | |
121 | matching_indices.end(), | |
122 | _1 == 1, | |
123 | queue | |
124 | ); | |
125 | ||
126 | // pattern was not found | |
127 | if(index == matching_indices.end()) | |
128 | return t_last; | |
129 | ||
130 | return t_first + detail::iterator_range_size(matching_indices.begin(), index); | |
131 | } | |
132 | ||
133 | } //end compute namespace | |
134 | } //end boost namespace | |
135 | ||
136 | #endif // BOOST_COMPUTE_ALGORITHM_FIND_END_HPP |