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/static_assert.hpp>
15
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>
23 #include <boost/compute/type_traits/is_device_iterator.hpp>
24
25 namespace boost {
26 namespace compute {
27 namespace detail {
28
29 ///
30 /// \brief Helper function for find_end
31 ///
32 /// Basically a copy of find_if which returns last occurrence
33 /// instead of first occurrence
34 ///
35 template<class InputIterator, class UnaryPredicate>
find_end_helper(InputIterator first,InputIterator last,UnaryPredicate predicate,command_queue & queue)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 ///
96 /// Space complexity: \Omega(n)
97 ///
98 template<class TextIterator, class PatternIterator>
find_end(TextIterator t_first,TextIterator t_last,PatternIterator p_first,PatternIterator p_last,command_queue & queue=system::default_queue ())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 {
105 BOOST_STATIC_ASSERT(is_device_iterator<TextIterator>::value);
106 BOOST_STATIC_ASSERT(is_device_iterator<PatternIterator>::value);
107
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
145