117 lines
		
	
	
		
			3.9 KiB
		
	
	
	
		
			Plaintext
		
	
	
	
	
	
		
		
			
		
	
	
			117 lines
		
	
	
		
			3.9 KiB
		
	
	
	
		
			Plaintext
		
	
	
	
	
	
|   | //---------------------------------------------------------------------------// | ||
|  | // Copyright (c) 2014 Roshan <thisisroshansmail@gmail.com> | ||
|  | // | ||
|  | // Distributed under the Boost Software License, Version 1.0 | ||
|  | // See accompanying file LICENSE_1_0.txt or copy at | ||
|  | // http://www.boost.org/LICENSE_1_0.txt | ||
|  | // | ||
|  | // See http://boostorg.github.com/compute for more information. | ||
|  | //---------------------------------------------------------------------------// | ||
|  | 
 | ||
|  | #ifndef BOOST_COMPUTE_ALGORITHM_DETAIL_MERGE_PATH_HPP | ||
|  | #define BOOST_COMPUTE_ALGORITHM_DETAIL_MERGE_PATH_HPP | ||
|  | 
 | ||
|  | #include <iterator> | ||
|  | 
 | ||
|  | #include <boost/compute/algorithm/find_if.hpp> | ||
|  | #include <boost/compute/container/vector.hpp> | ||
|  | #include <boost/compute/detail/iterator_range_size.hpp> | ||
|  | #include <boost/compute/detail/meta_kernel.hpp> | ||
|  | #include <boost/compute/lambda.hpp> | ||
|  | #include <boost/compute/system.hpp> | ||
|  | 
 | ||
|  | namespace boost { | ||
|  | namespace compute { | ||
|  | namespace detail { | ||
|  | 
 | ||
|  | /// | ||
|  | /// \brief Merge Path kernel class | ||
|  | /// | ||
|  | /// Subclass of meta_kernel to break two sets into tiles according | ||
|  | /// to their merge path | ||
|  | /// | ||
|  | class merge_path_kernel : public meta_kernel | ||
|  | { | ||
|  | public: | ||
|  |     unsigned int tile_size; | ||
|  | 
 | ||
|  |     merge_path_kernel() : meta_kernel("merge_path") | ||
|  |     { | ||
|  |         tile_size = 4; | ||
|  |     } | ||
|  | 
 | ||
|  |     template<class InputIterator1, class InputIterator2, | ||
|  |              class OutputIterator1, class OutputIterator2, | ||
|  |              class Compare> | ||
|  |     void set_range(InputIterator1 first1, | ||
|  |                    InputIterator1 last1, | ||
|  |                    InputIterator2 first2, | ||
|  |                    InputIterator2 last2, | ||
|  |                    OutputIterator1 result_a, | ||
|  |                    OutputIterator2 result_b, | ||
|  |                    Compare comp) | ||
|  |     { | ||
|  |         m_a_count = iterator_range_size(first1, last1); | ||
|  |         m_a_count_arg = add_arg<uint_>("a_count"); | ||
|  | 
 | ||
|  |         m_b_count = iterator_range_size(first2, last2); | ||
|  |         m_b_count_arg = add_arg<uint_>("b_count"); | ||
|  | 
 | ||
|  |         *this << | ||
|  |             "uint i = get_global_id(0);\n" << | ||
|  |             "uint target = (i+1)*" << tile_size << ";\n" << | ||
|  |             "uint start = max(convert_int(0),convert_int(target)-convert_int(b_count));\n" << | ||
|  |             "uint end = min(target,a_count);\n" << | ||
|  |             "uint a_index, b_index;\n" << | ||
|  |             "while(start<end)\n" << | ||
|  |             "{\n" << | ||
|  |             "   a_index = (start + end)/2;\n" << | ||
|  |             "   b_index = target - a_index - 1;\n" << | ||
|  |             "   if(!(" << comp(first2[expr<uint_>("b_index")], | ||
|  |                               first1[expr<uint_>("a_index")]) << "))\n" << | ||
|  |             "       start = a_index + 1;\n" << | ||
|  |             "   else end = a_index;\n" << | ||
|  |             "}\n" << | ||
|  |             result_a[expr<uint_>("i")] << " = start;\n" << | ||
|  |             result_b[expr<uint_>("i")] << " = target - start;\n"; | ||
|  |     } | ||
|  | 
 | ||
|  |     template<class InputIterator1, class InputIterator2, | ||
|  |              class OutputIterator1, class OutputIterator2> | ||
|  |     void set_range(InputIterator1 first1, | ||
|  |                    InputIterator1 last1, | ||
|  |                    InputIterator2 first2, | ||
|  |                    InputIterator2 last2, | ||
|  |                    OutputIterator1 result_a, | ||
|  |                    OutputIterator2 result_b) | ||
|  |     { | ||
|  |         typedef typename std::iterator_traits<InputIterator1>::value_type value_type; | ||
|  |         ::boost::compute::less<value_type> less_than; | ||
|  |         set_range(first1, last1, first2, last2, result_a, result_b, less_than); | ||
|  |     } | ||
|  | 
 | ||
|  |     event exec(command_queue &queue) | ||
|  |     { | ||
|  |         if((m_a_count + m_b_count)/tile_size == 0) { | ||
|  |             return event(); | ||
|  |         } | ||
|  | 
 | ||
|  |         set_arg(m_a_count_arg, uint_(m_a_count)); | ||
|  |         set_arg(m_b_count_arg, uint_(m_b_count)); | ||
|  | 
 | ||
|  |         return exec_1d(queue, 0, (m_a_count + m_b_count)/tile_size); | ||
|  |     } | ||
|  | 
 | ||
|  | private: | ||
|  |     size_t m_a_count; | ||
|  |     size_t m_a_count_arg; | ||
|  |     size_t m_b_count; | ||
|  |     size_t m_b_count_arg; | ||
|  | }; | ||
|  | 
 | ||
|  | } //end detail namespace | ||
|  | } //end compute namespace | ||
|  | } //end boost namespace | ||
|  | 
 | ||
|  | #endif // BOOST_COMPUTE_ALGORITHM_DETAIL_MERGE_PATH_HPP |