[compute] 29/49: Support for custom comparision function in find_extrema() and min/max_element()
Ghislain Vaillant
ghisvail-guest at moszumanska.debian.org
Fri Dec 18 17:58:19 UTC 2015
This is an automated email from the git hooks/post-receive script.
ghisvail-guest pushed a commit to branch master
in repository compute.
commit a50dae747d79a56f951f0838cd66737378c4c03e
Author: Jakub Szuppe <j.szuppe at gmail.com>
Date: Wed Aug 19 12:11:57 2015 +0200
Support for custom comparision function in find_extrema() and min/max_element()
---
.../compute/algorithm/detail/find_extrema.hpp | 15 +-
.../algorithm/detail/find_extrema_reduce.hpp | 182 +++++++++++++++------
.../algorithm/detail/find_extrema_with_atomics.hpp | 23 ++-
.../algorithm/detail/serial_find_extrema.hpp | 21 ++-
include/boost/compute/algorithm/max_element.hpp | 39 ++++-
include/boost/compute/algorithm/min_element.hpp | 41 ++++-
include/boost/compute/algorithm/minmax_element.hpp | 25 ++-
test/test_extrema.cpp | 34 ++++
8 files changed, 307 insertions(+), 73 deletions(-)
diff --git a/include/boost/compute/algorithm/detail/find_extrema.hpp b/include/boost/compute/algorithm/detail/find_extrema.hpp
index ca5d138..c9f3353 100644
--- a/include/boost/compute/algorithm/detail/find_extrema.hpp
+++ b/include/boost/compute/algorithm/detail/find_extrema.hpp
@@ -20,10 +20,11 @@ namespace boost {
namespace compute {
namespace detail {
-template<class InputIterator>
+template<class InputIterator, class Compare>
inline InputIterator find_extrema(InputIterator first,
InputIterator last,
- char sign,
+ Compare compare,
+ const bool find_minimum,
command_queue &queue)
{
size_t count = iterator_range_size(first, last);
@@ -37,23 +38,23 @@ inline InputIterator find_extrema(InputIterator first,
// use serial method for small inputs
// and when device is a CPU
- if(count < 64 || (device.type() & device::cpu)){
- return serial_find_extrema(first, last, sign, queue);
+ if(count < 512 || (device.type() & device::cpu)){
+ return serial_find_extrema(first, last, compare, find_minimum, queue);
}
// find_extrema_reduce() is used only if requirements are met
if(find_extrema_reduce_requirements_met(first, last, queue))
{
- return find_extrema_reduce(first, last, sign, queue);
+ return find_extrema_reduce(first, last, compare, find_minimum, queue);
}
// use serial method for OpenCL version 1.0 due to
// problems with atomic_cmpxchg()
#ifndef CL_VERSION_1_1
- return serial_find_extrema(first, last, sign, queue);
+ return serial_find_extrema(first, last, compare, find_minimum, queue);
#endif
- return find_extrema_with_atomics(first, last, sign, queue);
+ return find_extrema_with_atomics(first, last, compare, find_minimum, queue);
}
} // end detail namespace
diff --git a/include/boost/compute/algorithm/detail/find_extrema_reduce.hpp b/include/boost/compute/algorithm/detail/find_extrema_reduce.hpp
index 63a3fef..9a2350c 100644
--- a/include/boost/compute/algorithm/detail/find_extrema_reduce.hpp
+++ b/include/boost/compute/algorithm/detail/find_extrema_reduce.hpp
@@ -24,6 +24,7 @@
#include <boost/compute/memory/local_buffer.hpp>
#include <boost/compute/type_traits/type_name.hpp>
#include <boost/compute/utility/program_cache.hpp>
+#include <boost/compute/algorithm/detail/serial_find_extrema.hpp>
namespace boost {
namespace compute {
@@ -72,14 +73,15 @@ bool find_extrema_reduce_requirements_met(InputIterator first,
return ((required_local_mem_size * 4) <= local_mem_size);
}
-template<class InputIterator, class ResultIterator>
+template<class InputIterator, class ResultIterator, class Compare>
inline size_t find_extrema_reduce(InputIterator first,
size_t count,
ResultIterator result,
vector<uint_>::iterator result_idx,
size_t work_groups_no,
size_t work_group_size,
- char sign,
+ Compare compare,
+ const bool find_minimum,
command_queue &queue)
{
typedef typename std::iterator_traits<InputIterator>::value_type input_type;
@@ -115,7 +117,13 @@ inline size_t find_extrema_reduce(InputIterator first,
// Next element
k.decl<input_type>("next") << " = " << first[k.var<uint_>("idx")] << ";\n" <<
// Comparison between currently best element (acc) and next element
- "compare_result = acc " << sign << " next;\n" <<
+ "#ifndef BOOST_COMPUTE_FIND_MAXIMUM\n" <<
+ "compare_result = " << compare(k.var<input_type>("acc"),
+ k.var<input_type>("next")) << ";\n" <<
+ "#else\n" <<
+ "compare_result = " << compare(k.var<input_type>("next"),
+ k.var<input_type>("acc")) << ";\n" <<
+ "#endif\n" <<
"acc = compare_result ? acc : next;\n" <<
"acc_idx = compare_result ? acc_idx : idx;\n" <<
"idx += get_global_size(0);\n" <<
@@ -136,7 +144,13 @@ inline size_t find_extrema_reduce(InputIterator first,
"if((lid < offset) && ((lid + offset) < group_offset)) { \n" <<
k.decl<input_type>("mine") << " = block[lid];\n" <<
k.decl<input_type>("other") << " = block[lid+offset];\n" <<
- "compare_result = mine " << sign << " other;\n" <<
+ "#ifndef BOOST_COMPUTE_FIND_MAXIMUM\n" <<
+ "compare_result = " << compare(k.var<input_type>("mine"),
+ k.var<input_type>("other")) << ";\n" <<
+ "#else\n" <<
+ "compare_result = " << compare(k.var<input_type>("other"),
+ k.var<input_type>("mine")) << ";\n" <<
+ "#endif\n" <<
"block[lid] = compare_result ? mine : other;\n" <<
"block_idx[lid] = compare_result ? " <<
"block_idx[lid] : block_idx[lid+offset];\n" <<
@@ -150,7 +164,12 @@ inline size_t find_extrema_reduce(InputIterator first,
" output_idx[get_group_id(0)] = block_idx[0];\n" <<
"}";
- kernel kernel = k.compile(context);
+ std::string options;
+ if(!find_minimum){
+ options = "-DBOOST_COMPUTE_FIND_MAXIMUM";
+ }
+ kernel kernel = k.compile(context, options);
+
kernel.set_arg(count_arg, static_cast<uint_>(count));
kernel.set_arg(output_arg, result.get_buffer());
kernel.set_arg(output_idx_arg, result_idx.get_buffer());
@@ -165,10 +184,101 @@ inline size_t find_extrema_reduce(InputIterator first,
return 0;
}
+template<class InputIterator, class Compare>
+uint_ find_extrema_final(InputIterator candidates,
+ vector<uint_>::iterator candidates_idx,
+ const size_t count,
+ Compare compare,
+ const bool find_minimum,
+ const size_t work_group_size,
+ command_queue &queue)
+{
+ typedef typename std::iterator_traits<InputIterator>::value_type input_type;
+
+ const context &context = queue.get_context();
+
+ // device vectors for the result
+ vector<input_type> result(1, context);
+ vector<uint_> result_idx(1, context);
+
+ // get extremum from among the candidates
+ find_extrema_reduce(
+ candidates, count, result.begin(), result_idx.begin(),
+ 1, work_group_size, compare, find_minimum, queue
+ );
+
+ // get candidate index
+ const uint_ idx = (result_idx.begin()).read(queue);
+ // get extremum index
+ typename vector<uint_>::iterator extremum_idx = candidates_idx + idx;
+
+ // return extremum index
+ return extremum_idx.read(queue);
+}
+
template<class InputIterator>
+uint_ find_extrema_final(InputIterator candidates,
+ vector<uint_>::iterator candidates_idx,
+ const size_t count,
+ ::boost::compute::less<
+ typename std::iterator_traits<InputIterator>::value_type
+ > compare,
+ const bool find_minimum,
+ const size_t work_group_size,
+ command_queue &queue)
+{
+ (void) work_group_size;
+
+ typedef typename std::iterator_traits<InputIterator>::difference_type difference_type;
+ typedef typename std::iterator_traits<InputIterator>::value_type input_type;
+
+ // host vectors
+ std::vector<input_type> host_candidates(count);
+ std::vector<uint_> host_candidates_idx(count);
+
+ InputIterator candidates_last =
+ candidates + static_cast<difference_type>(count);
+ vector<uint_>::iterator candidates_idx_last =
+ candidates_idx + count;
+
+ // copying extremum candidates found by find_extrema_reduce(...) to host
+ ::boost::compute::copy(candidates_idx, candidates_idx_last,
+ host_candidates_idx.begin(), queue);
+ ::boost::compute::copy(candidates, candidates_last,
+ host_candidates.begin(), queue);
+
+ typename std::vector<input_type>::iterator i = host_candidates.begin();
+ std::vector<uint_>::iterator idx = host_candidates_idx.begin();
+ std::vector<uint_>::iterator extremum_idx = idx;
+ input_type extremum = *i;
+
+ // find extremum from among the candidates
+ if(!find_minimum) {
+ while(idx != host_candidates_idx.end()) {
+ bool compare_result = *i > extremum;
+ extremum = compare_result ? *i : extremum;
+ extremum_idx = compare_result ? idx : extremum_idx;
+ idx++, i++;
+ }
+ }
+ else {
+ while(idx != host_candidates_idx.end()) {
+ bool compare_result = *i < extremum;
+ extremum = compare_result ? *i : extremum;
+ extremum_idx = compare_result ? idx : extremum_idx;
+ idx++, i++;
+ }
+ }
+
+ // return extremum index
+ return (*extremum_idx);
+}
+
+template<class InputIterator, class Compare>
InputIterator find_extrema_reduce(InputIterator first,
InputIterator last,
- char sign,
+ Compare compare,
+ const bool find_minimum,
command_queue &queue)
{
typedef typename std::iterator_traits<InputIterator>::difference_type difference_type;
@@ -203,54 +313,22 @@ InputIterator find_extrema_reduce(InputIterator first,
static_cast<size_t>(std::ceil(float(count) / work_group_size)));
// device vectors for extremum candidates and their indices
- vector<input_type> results(work_groups_no, context);
- vector<uint_> results_idx(work_groups_no, context);
+ vector<input_type> candidates(work_groups_no, context);
+ vector<uint_> candidates_idx(work_groups_no, context);
// find extremum candidates and their indices
- find_extrema_reduce(first, count,
- results.begin(), results_idx.begin(),
- work_groups_no, work_group_size,
- sign,
- queue);
-
- // host vectors
- std::vector<input_type> host_results(work_groups_no);
- std::vector<uint_> host_results_idx(work_groups_no);
-
- // copying extremum candidates found by
- // find_extrema_reduce(...) to host
- copy(results_idx.begin(),
- results_idx.end(),
- host_results_idx.begin(), queue);
- copy(results.begin(),
- results.end(),
- host_results.begin(), queue);
-
- typename std::vector<input_type>::iterator i = host_results.begin();
- std::vector<uint_>::iterator idx = host_results_idx.begin();
- std::vector<uint_>::iterator extreme_idx = idx;
- input_type extreme = *i;
-
- // find extremum from candidates found by find_extrema_reduce(...)
- if(sign == '>') {
- while(idx != host_results_idx.end()) {
- bool compare_result = *i > extreme;
- extreme = compare_result ? *i : extreme;
- extreme_idx = compare_result ? idx : extreme_idx;
- idx++, i++;
- }
- }
- else {
- while(idx != host_results_idx.end()) {
- bool compare_result = *i < extreme;
- extreme = compare_result ? *i : extreme;
- extreme_idx = compare_result ? idx : extreme_idx;
- idx++, i++;
- }
- }
-
- // return iterator to extremum
- return first + static_cast<difference_type>(*extreme_idx);
+ find_extrema_reduce(
+ first, count, candidates.begin(), candidates_idx.begin(),
+ work_groups_no, work_group_size, compare, find_minimum, queue
+ );
+
+ // get extremum index
+ const uint_ extremum_idx = find_extrema_final(
+ candidates.begin(), candidates_idx.begin(), work_groups_no, compare,
+ find_minimum, work_group_size, queue
+ );
+
+ return first + static_cast<difference_type>(extremum_idx);
}
} // end detail namespace
diff --git a/include/boost/compute/algorithm/detail/find_extrema_with_atomics.hpp b/include/boost/compute/algorithm/detail/find_extrema_with_atomics.hpp
index 4e31b6c..5ebaf83 100644
--- a/include/boost/compute/algorithm/detail/find_extrema_with_atomics.hpp
+++ b/include/boost/compute/algorithm/detail/find_extrema_with_atomics.hpp
@@ -22,10 +22,11 @@ namespace boost {
namespace compute {
namespace detail {
-template<class InputIterator>
+template<class InputIterator, class Compare>
inline InputIterator find_extrema_with_atomics(InputIterator first,
InputIterator last,
- char sign,
+ Compare compare,
+ const bool find_minimum,
command_queue &queue)
{
typedef typename std::iterator_traits<InputIterator>::difference_type difference_type;
@@ -38,9 +39,15 @@ inline InputIterator find_extrema_with_atomics(InputIterator first,
k <<
"const uint gid = get_global_id(0);\n" <<
"uint old_index = *index;\n" <<
- "while(" << first[k.var<uint_>("gid")]
- << sign
- << first[k.var<uint_>("old_index")] << "){\n" <<
+
+ "#ifndef BOOST_COMPUTE_FIND_MAXIMUM\n" <<
+ "while(" << compare(first[k.var<uint_>("gid")],
+ first[k.var<uint_>("old_index")]) << "){\n" <<
+ "#else\n" <<
+ "while(" << compare(first[k.var<uint_>("old_index")],
+ first[k.var<uint_>("gid")]) << "){\n" <<
+ "#endif\n" <<
+
" if(" << atomic_cmpxchg_uint(k.var<uint_ *>("index"),
k.var<uint_>("old_index"),
k.var<uint_>("gid")) << " == old_index)\n" <<
@@ -51,7 +58,11 @@ inline InputIterator find_extrema_with_atomics(InputIterator first,
size_t index_arg_index = k.add_arg<uint_ *>(memory_object::global_memory, "index");
- kernel kernel = k.compile(context);
+ std::string options;
+ if(!find_minimum){
+ options = "-DBOOST_COMPUTE_FIND_MAXIMUM";
+ }
+ kernel kernel = k.compile(context, options);
// setup index buffer
scalar<uint_> index(context);
diff --git a/include/boost/compute/algorithm/detail/serial_find_extrema.hpp b/include/boost/compute/algorithm/detail/serial_find_extrema.hpp
index 11d9122..8407c88 100644
--- a/include/boost/compute/algorithm/detail/serial_find_extrema.hpp
+++ b/include/boost/compute/algorithm/detail/serial_find_extrema.hpp
@@ -21,10 +21,11 @@ namespace boost {
namespace compute {
namespace detail {
-template<class InputIterator>
+template<class InputIterator, class Compare>
inline InputIterator serial_find_extrema(InputIterator first,
InputIterator last,
- char sign,
+ Compare compare,
+ const bool find_minimum,
command_queue &queue)
{
typedef typename std::iterator_traits<InputIterator>::value_type value_type;
@@ -40,7 +41,15 @@ inline InputIterator serial_find_extrema(InputIterator first,
"for(uint i = 1; i < size; i++){\n" <<
" " << k.decl<value_type>("candidate") << "="
<< first[k.expr<uint_>("i")] << ";\n" <<
- " if(candidate" << sign << "value){\n" <<
+
+ "#ifndef BOOST_COMPUTE_FIND_MAXIMUM\n" <<
+ " if(" << compare(k.var<value_type>("candidate"),
+ k.var<value_type>("value")) << "){\n" <<
+ "#else\n" <<
+ " if(" << compare(k.var<value_type>("value"),
+ k.var<value_type>("candidate")) << "){\n" <<
+ "#endif\n" <<
+
" value = candidate;\n" <<
" value_index = i;\n" <<
" }\n" <<
@@ -50,7 +59,11 @@ inline InputIterator serial_find_extrema(InputIterator first,
size_t index_arg_index = k.add_arg<uint_ *>(memory_object::global_memory, "index");
size_t size_arg_index = k.add_arg<uint_>("size");
- kernel kernel = k.compile(context);
+ std::string options;
+ if(!find_minimum){
+ options = "-DBOOST_COMPUTE_FIND_MAXIMUM";
+ }
+ kernel kernel = k.compile(context, options);
// setup index buffer
scalar<uint_> index(context);
diff --git a/include/boost/compute/algorithm/max_element.hpp b/include/boost/compute/algorithm/max_element.hpp
index 8b805e3..4384404 100644
--- a/include/boost/compute/algorithm/max_element.hpp
+++ b/include/boost/compute/algorithm/max_element.hpp
@@ -13,6 +13,7 @@
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
+#include <boost/compute/functional.hpp>
#include <boost/compute/algorithm/detail/find_extrema.hpp>
namespace boost {
@@ -21,14 +22,50 @@ namespace compute {
/// Returns an iterator pointing to the element in the range
/// [\p first, \p last) with the maximum value.
///
+/// \param first first element in the input range
+/// \param last last element in the input range
+/// \param compare comparison function object which returns true if the first
+/// argument is less than (i.e. is ordered before) the second.
+/// \param queue command queue to perform the operation
+///
+/// For example, to find \c int2 value with maximum first component in given vector:
+/// \code
+/// // comparison function object
+/// BOOST_COMPUTE_FUNCTION(bool, compare_first, (const int2_ &a, const int2_ &b),
+/// {
+/// return a.x < b.x;
+/// });
+///
+/// // create vector
+/// boost::compute::vector<uint2_> data = ...
+///
+/// boost::compute::vector<uint2_>::iterator max =
+/// boost::compute::max_element(data.begin(), data.end(), compare_first, queue);
+/// \endcode
+///
/// \see min_element()
+template<class InputIterator, class Compare>
+inline InputIterator
+max_element(InputIterator first,
+ InputIterator last,
+ Compare compare,
+ command_queue &queue = system::default_queue())
+{
+ return detail::find_extrema(first, last, compare, false, queue);
+}
+
+///\overload
template<class InputIterator>
inline InputIterator
max_element(InputIterator first,
InputIterator last,
command_queue &queue = system::default_queue())
{
- return detail::find_extrema(first, last, '>', queue);
+ typedef typename std::iterator_traits<InputIterator>::value_type value_type;
+
+ return ::boost::compute::max_element(
+ first, last, ::boost::compute::less<value_type>(), queue
+ );
}
} // end compute namespace
diff --git a/include/boost/compute/algorithm/min_element.hpp b/include/boost/compute/algorithm/min_element.hpp
index 00bf361..97603af 100644
--- a/include/boost/compute/algorithm/min_element.hpp
+++ b/include/boost/compute/algorithm/min_element.hpp
@@ -13,22 +13,59 @@
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
+#include <boost/compute/functional.hpp>
#include <boost/compute/algorithm/detail/find_extrema.hpp>
namespace boost {
namespace compute {
/// Returns an iterator pointing to the element in range
-/// [\p first, \p last) with the minumum value.
+/// [\p first, \p last) with the minimum value.
+///
+/// \param first first element in the input range
+/// \param last last element in the input range
+/// \param compare comparison function object which returns true if the first
+/// argument is less than (i.e. is ordered before) the second.
+/// \param queue command queue to perform the operation
+///
+/// For example, to find \c int2 value with minimum first component in given vector:
+/// \code
+/// // comparison function object
+/// BOOST_COMPUTE_FUNCTION(bool, compare_first, (const int2_ &a, const int2_ &b),
+/// {
+/// return a.x < b.x;
+/// });
+///
+/// // create vector
+/// boost::compute::vector<uint2_> data = ...
+///
+/// boost::compute::vector<uint2_>::iterator min =
+/// boost::compute::min_element(data.begin(), data.end(), compare_first, queue);
+/// \endcode
///
/// \see max_element()
+template<class InputIterator, class Compare>
+inline InputIterator
+min_element(InputIterator first,
+ InputIterator last,
+ Compare compare,
+ command_queue &queue = system::default_queue())
+{
+ return detail::find_extrema(first, last, compare, true, queue);
+}
+
+///\overload
template<class InputIterator>
inline InputIterator
min_element(InputIterator first,
InputIterator last,
command_queue &queue = system::default_queue())
{
- return detail::find_extrema(first, last, '<', queue);
+ typedef typename std::iterator_traits<InputIterator>::value_type value_type;
+
+ return ::boost::compute::min_element(
+ first, last, ::boost::compute::less<value_type>(), queue
+ );
}
} // end compute namespace
diff --git a/include/boost/compute/algorithm/minmax_element.hpp b/include/boost/compute/algorithm/minmax_element.hpp
index eff5aef..600add7 100644
--- a/include/boost/compute/algorithm/minmax_element.hpp
+++ b/include/boost/compute/algorithm/minmax_element.hpp
@@ -25,8 +25,31 @@ namespace compute {
/// element and the second pointing to the maximum element in the range
/// [\p first, \p last).
///
+/// \param first first element in the input range
+/// \param last last element in the input range
+/// \param compare comparison function object which returns true if the first
+/// argument is less than (i.e. is ordered before) the second.
+/// \param queue command queue to perform the operation
+///
/// \see max_element(), min_element()
-template<class InputIterator>
+template<class InputIterator, class Compare>
+inline std::pair<InputIterator, InputIterator>
+minmax_element(InputIterator first,
+ InputIterator last,
+ Compare compare,
+ command_queue &queue = system::default_queue())
+{
+ if(first == last){
+ // empty range
+ return std::make_pair(first, first);
+ }
+
+ return std::make_pair(min_element(first, last, compare, queue),
+ max_element(first, last, compare, queue));
+}
+
+///\overload
+template<class InputIterator, class Compare>
inline std::pair<InputIterator, InputIterator>
minmax_element(InputIterator first,
InputIterator last,
diff --git a/test/test_extrema.cpp b/test/test_extrema.cpp
index a7cadbb..7d68667 100644
--- a/test/test_extrema.cpp
+++ b/test/test_extrema.cpp
@@ -39,6 +39,40 @@ BOOST_AUTO_TEST_CASE(int_min_max)
BOOST_CHECK_EQUAL(*max_iter, 15);
}
+BOOST_AUTO_TEST_CASE(int2_min_max_custom_comparision_function)
+{
+ using boost::compute::int2_;
+
+ boost::compute::vector<int2_> vector(context);
+ vector.push_back(int2_(1, 10), queue);
+ vector.push_back(int2_(2, -100), queue);
+ vector.push_back(int2_(3, 30), queue);
+ vector.push_back(int2_(4, 20), queue);
+ vector.push_back(int2_(5, 5), queue);
+ vector.push_back(int2_(6, -80), queue);
+ vector.push_back(int2_(7, 21), queue);
+ vector.push_back(int2_(8, -5), queue);
+
+ BOOST_COMPUTE_FUNCTION(bool, compare_second, (const int2_ a, const int2_ b),
+ {
+ return a.y < b.y;
+ });
+
+ boost::compute::vector<int2_>::iterator min_iter =
+ boost::compute::min_element(
+ vector.begin(), vector.end(), compare_second, queue
+ );
+ BOOST_CHECK(min_iter == vector.begin() + 1);
+ BOOST_CHECK_EQUAL(*min_iter, int2_(2, -100));
+
+ boost::compute::vector<int2_>::iterator max_iter =
+ boost::compute::max_element(
+ vector.begin(), vector.end(), compare_second, queue
+ );
+ BOOST_CHECK(max_iter == vector.begin() + 2);
+ BOOST_CHECK_EQUAL(*max_iter, int2_(3, 30));
+}
+
BOOST_AUTO_TEST_CASE(iota_min_max)
{
boost::compute::vector<int> vector(5000);
--
Alioth's /usr/local/bin/git-commit-notice on /srv/git.debian.org/git/debian-science/packages/compute.git
More information about the debian-science-commits
mailing list