1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
|
//---------------------------------------------------------------------------//
// 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
|
