find_end.hpp 4.7 KB

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