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