includes.hpp 5.5 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162
  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_INCLUDES_HPP
  11. #define BOOST_COMPUTE_ALGORITHM_INCLUDES_HPP
  12. #include <iterator>
  13. #include <boost/static_assert.hpp>
  14. #include <boost/compute/algorithm/detail/balanced_path.hpp>
  15. #include <boost/compute/algorithm/fill_n.hpp>
  16. #include <boost/compute/algorithm/find.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/detail/read_write_single_value.hpp>
  21. #include <boost/compute/system.hpp>
  22. #include <boost/compute/type_traits/is_device_iterator.hpp>
  23. namespace boost {
  24. namespace compute {
  25. namespace detail {
  26. ///
  27. /// \brief Serial includes kernel class
  28. ///
  29. /// Subclass of meta_kernel to perform includes operation after tiling
  30. ///
  31. class serial_includes_kernel : meta_kernel
  32. {
  33. public:
  34. serial_includes_kernel() : meta_kernel("includes")
  35. {
  36. }
  37. template<class InputIterator1, class InputIterator2,
  38. class InputIterator3, class InputIterator4,
  39. class OutputIterator>
  40. void set_range(InputIterator1 first1,
  41. InputIterator2 first2,
  42. InputIterator3 tile_first1,
  43. InputIterator3 tile_last1,
  44. InputIterator4 tile_first2,
  45. OutputIterator result)
  46. {
  47. m_count = iterator_range_size(tile_first1, tile_last1) - 1;
  48. *this <<
  49. "uint i = get_global_id(0);\n" <<
  50. "uint start1 = " << tile_first1[expr<uint_>("i")] << ";\n" <<
  51. "uint end1 = " << tile_first1[expr<uint_>("i+1")] << ";\n" <<
  52. "uint start2 = " << tile_first2[expr<uint_>("i")] << ";\n" <<
  53. "uint end2 = " << tile_first2[expr<uint_>("i+1")] << ";\n" <<
  54. "uint includes = 1;\n" <<
  55. "while(start1<end1 && start2<end2)\n" <<
  56. "{\n" <<
  57. " if(" << first1[expr<uint_>("start1")] << " == " <<
  58. first2[expr<uint_>("start2")] << ")\n" <<
  59. " {\n" <<
  60. " start1++; start2++;\n" <<
  61. " }\n" <<
  62. " else if(" << first1[expr<uint_>("start1")] << " < " <<
  63. first2[expr<uint_>("start2")] << ")\n" <<
  64. " start1++;\n" <<
  65. " else\n" <<
  66. " {\n" <<
  67. " includes = 0;\n" <<
  68. " break;\n" <<
  69. " }\n" <<
  70. "}\n" <<
  71. "if(start2<end2)\n" <<
  72. " includes = 0;\n" <<
  73. result[expr<uint_>("i")] << " = includes;\n";
  74. }
  75. event exec(command_queue &queue)
  76. {
  77. if(m_count == 0) {
  78. return event();
  79. }
  80. return exec_1d(queue, 0, m_count);
  81. }
  82. private:
  83. size_t m_count;
  84. };
  85. } //end detail namespace
  86. ///
  87. /// \brief Includes algorithm
  88. ///
  89. /// Finds if the sorted range [first1, last1) includes the sorted
  90. /// range [first2, last2). In other words, it checks if [first1, last1) is
  91. /// a superset of [first2, last2).
  92. ///
  93. /// \return True, if [first1, last1) includes [first2, last2). False otherwise.
  94. ///
  95. /// \param first1 Iterator pointing to start of first set
  96. /// \param last1 Iterator pointing to end of first set
  97. /// \param first2 Iterator pointing to start of second set
  98. /// \param last2 Iterator pointing to end of second set
  99. /// \param queue Queue on which to execute
  100. ///
  101. /// Space complexity: \Omega(distance(\p first1, \p last1) + distance(\p first2, \p last2))
  102. template<class InputIterator1, class InputIterator2>
  103. inline bool includes(InputIterator1 first1,
  104. InputIterator1 last1,
  105. InputIterator2 first2,
  106. InputIterator2 last2,
  107. command_queue &queue = system::default_queue())
  108. {
  109. BOOST_STATIC_ASSERT(is_device_iterator<InputIterator1>::value);
  110. BOOST_STATIC_ASSERT(is_device_iterator<InputIterator2>::value);
  111. size_t tile_size = 1024;
  112. size_t count1 = detail::iterator_range_size(first1, last1);
  113. size_t count2 = detail::iterator_range_size(first2, last2);
  114. vector<uint_> tile_a((count1+count2+tile_size-1)/tile_size+1, queue.get_context());
  115. vector<uint_> tile_b((count1+count2+tile_size-1)/tile_size+1, queue.get_context());
  116. // Tile the sets
  117. detail::balanced_path_kernel tiling_kernel;
  118. tiling_kernel.tile_size = static_cast<unsigned int>(tile_size);
  119. tiling_kernel.set_range(first1, last1, first2, last2,
  120. tile_a.begin()+1, tile_b.begin()+1);
  121. fill_n(tile_a.begin(), 1, uint_(0), queue);
  122. fill_n(tile_b.begin(), 1, uint_(0), queue);
  123. tiling_kernel.exec(queue);
  124. fill_n(tile_a.end()-1, 1, static_cast<uint_>(count1), queue);
  125. fill_n(tile_b.end()-1, 1, static_cast<uint_>(count2), queue);
  126. vector<uint_> result((count1+count2+tile_size-1)/tile_size, queue.get_context());
  127. // Find individually
  128. detail::serial_includes_kernel includes_kernel;
  129. includes_kernel.set_range(first1, first2, tile_a.begin(), tile_a.end(),
  130. tile_b.begin(), result.begin());
  131. includes_kernel.exec(queue);
  132. return find(result.begin(), result.end(), 0, queue) == result.end();
  133. }
  134. } //end compute namespace
  135. } //end boost namespace
  136. #endif // BOOST_COMPUTE_ALGORITHM_SET_UNION_HPP