w4kek/WSJT-X
1
0
Fork 0
mirror of https://github.com/saitohirga/WSJT-X.git synced 2026-06-07 08:24:53 -04:00
Code Issues Packages Projects Releases Wiki Activity

Squashed 'boost/' changes from d9443bc48..c27aa31f0

Browse Source 
c27aa31f0 Updated Boost to v1.70.0 including iterator range math numeric crc circular_buffer multi_index intrusive

git-subtree-dir: boost
git-subtree-split: c27aa31f06ebf1a91b3fa3ae9df9b5efdf14ec9f
This commit is contained in:
Bill Somerville
2019-07-02 23:38:24 +01:00
parent edd0930758
commit d361e123c6
6021 changed files with 118782 additions and 872011 deletions
Download Patch File Download Diff File Expand all files Collapse all files
boost/compute/algorithm/accumulate.hpp
+9
View File
@@ -11,6 +11,7 @@
#ifndef BOOST_COMPUTE_ALGORITHM_ACCUMULATE_HPP
#define BOOST_COMPUTE_ALGORITHM_ACCUMULATE_HPP
#include <boost/static_assert.hpp>
#include <boost/preprocessor/seq/for_each.hpp>
#include <boost/compute/system.hpp>
@@ -20,12 +21,14 @@
#include <boost/compute/algorithm/detail/serial_accumulate.hpp>
#include <boost/compute/container/array.hpp>
#include <boost/compute/container/vector.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
#include <boost/compute/detail/iterator_range_size.hpp>
namespace boost {
namespace compute {
namespace detail {
// Space complexity O(1)
template<class InputIterator, class T, class BinaryFunction>
inline T generic_accumulate(InputIterator first,
InputIterator last,
@@ -155,6 +158,9 @@ inline T dispatch_accumulate(InputIterator first,
/// reduce(vec.begin(), vec.end(), &result, plus<float>()); // fast
/// \endcode
///
/// Space complexity: \Omega(1)<br>
/// Space complexity when optimized to \c reduce(): \Omega(n)
///
/// \see reduce()
template<class InputIterator, class T, class BinaryFunction>
inline T accumulate(InputIterator first,
@@ -163,6 +169,8 @@ inline T accumulate(InputIterator first,
BinaryFunction function,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
return detail::dispatch_accumulate(first, last, init, function, queue);
}
@@ -173,6 +181,7 @@ inline T accumulate(InputIterator first,
T init,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
typedef typename std::iterator_traits<InputIterator>::value_type IT;
return detail::dispatch_accumulate(first, last, init, plus<IT>(), queue);
boost/compute/algorithm/adjacent_difference.hpp
+10
View File
@@ -13,12 +13,15 @@
#include <iterator>
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/detail/meta_kernel.hpp>
#include <boost/compute/detail/iterator_range_size.hpp>
#include <boost/compute/functional/operator.hpp>
#include <boost/compute/container/vector.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -64,6 +67,9 @@ dispatch_adjacent_difference(InputIterator first,
///
/// \return \c OutputIterator to the end of the result range
///
/// Space complexity: \Omega(1)<br>
/// Space complexity when \p result == \p first: \Omega(n)
///
/// \see adjacent_find()
template<class InputIterator, class OutputIterator, class BinaryFunction>
inline OutputIterator
@@ -73,6 +79,8 @@ adjacent_difference(InputIterator first,
BinaryFunction op,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
typedef typename std::iterator_traits<InputIterator>::value_type value_type;
if(first == last) {
@@ -103,6 +111,8 @@ adjacent_difference(InputIterator first,
OutputIterator result,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
typedef typename std::iterator_traits<InputIterator>::value_type value_type;
return ::boost::compute::adjacent_difference(
boost/compute/algorithm/adjacent_find.hpp
+7
View File
@@ -13,6 +13,8 @@
#include <iterator>
#include <boost/static_assert.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/lambda.hpp>
#include <boost/compute/system.hpp>
@@ -21,6 +23,7 @@
#include <boost/compute/detail/meta_kernel.hpp>
#include <boost/compute/functional/operator.hpp>
#include <boost/compute/type_traits/vector_size.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -114,6 +117,8 @@ adjacent_find_with_atomics(InputIterator first,
/// \return \c InputIteratorm to the first element which compares equal
/// to the following element. If none are equal, returns \c last.
///
/// Space complexity: \Omega(1)
///
/// \see find(), adjacent_difference()
template<class InputIterator, class Compare>
inline InputIterator
@@ -122,6 +127,7 @@ adjacent_find(InputIterator first,
Compare compare,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
size_t count = detail::iterator_range_size(first, last);
if(count < 32){
return detail::serial_adjacent_find(first, last, compare, queue);
@@ -138,6 +144,7 @@ adjacent_find(InputIterator first,
InputIterator last,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
typedef typename std::iterator_traits<InputIterator>::value_type value_type;
using ::boost::compute::lambda::_1;
boost/compute/algorithm/all_of.hpp
+6
View File
@@ -11,8 +11,11 @@
#ifndef BOOST_COMPUTE_ALGORITHM_ALL_OF_HPP
#define BOOST_COMPUTE_ALGORITHM_ALL_OF_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/algorithm/find_if_not.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -20,6 +23,8 @@ namespace compute {
/// Returns \c true if \p predicate returns \c true for all of the elements in
/// the range [\p first, \p last).
///
/// Space complexity: \Omega(1)
///
/// \see any_of(), none_of()
template<class InputIterator, class UnaryPredicate>
inline bool all_of(InputIterator first,
@@ -27,6 +32,7 @@ inline bool all_of(InputIterator first,
UnaryPredicate predicate,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
return ::boost::compute::find_if_not(first, last, predicate, queue) == last;
}
boost/compute/algorithm/any_of.hpp
+6
View File
@@ -11,8 +11,11 @@
#ifndef BOOST_COMPUTE_ALGORITHM_ANY_OF_HPP
#define BOOST_COMPUTE_ALGORITHM_ANY_OF_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/algorithm/find_if.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -24,6 +27,8 @@ namespace compute {
///
/// \snippet test/test_any_all_none_of.cpp any_of
///
/// Space complexity: \Omega(1)
///
/// \see all_of(), none_of()
template<class InputIterator, class UnaryPredicate>
inline bool any_of(InputIterator first,
@@ -31,6 +36,7 @@ inline bool any_of(InputIterator first,
UnaryPredicate predicate,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
return ::boost::compute::find_if(first, last, predicate, queue) != last;
}
boost/compute/algorithm/binary_search.hpp
+6
View File
@@ -11,21 +11,27 @@
#ifndef BOOST_COMPUTE_ALGORITHM_BINARY_SEARCH_HPP
#define BOOST_COMPUTE_ALGORITHM_BINARY_SEARCH_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/algorithm/lower_bound.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
/// Returns \c true if \p value is in the sorted range [\p first,
/// \p last).
///
/// Space complexity: \Omega(1)
template<class InputIterator, class T>
inline bool binary_search(InputIterator first,
InputIterator last,
const T &value,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
InputIterator position = lower_bound(first, last, value, queue);
return position != last && position.read(queue) == value;
boost/compute/algorithm/copy.hpp
+57 -32
View File
@@ -103,6 +103,7 @@ dispatch_copy_async(InputIterator first,
InputIterator last,
OutputIterator result,
command_queue &queue,
const wait_list &events,
typename boost::enable_if<
mpl::and_<
mpl::not_<
@@ -118,7 +119,7 @@ dispatch_copy_async(InputIterator first,
"copy_async() is only supported for contiguous host iterators"
);
return copy_to_device_async(first, last, result, queue);
return copy_to_device_async(first, last, result, queue, events);
}
// host -> device (async)
@@ -129,6 +130,7 @@ dispatch_copy_async(InputIterator first,
InputIterator last,
OutputIterator result,
command_queue &queue,
const wait_list &events,
typename boost::enable_if<
mpl::and_<
mpl::not_<
@@ -167,7 +169,7 @@ dispatch_copy_async(InputIterator first,
context
);
return copy_on_device_async(
mapped_host.begin(), mapped_host.end(), result, queue
mapped_host.begin(), mapped_host.end(), result, queue, events
);
}
@@ -179,6 +181,7 @@ dispatch_copy(InputIterator first,
InputIterator last,
OutputIterator result,
command_queue &queue,
const wait_list &events,
typename boost::enable_if<
mpl::and_<
mpl::not_<
@@ -190,7 +193,7 @@ dispatch_copy(InputIterator first,
>
>::type* = 0)
{
return copy_to_device(first, last, result, queue);
return copy_to_device(first, last, result, queue, events);
}
// host -> device
@@ -202,6 +205,7 @@ dispatch_copy(InputIterator first,
InputIterator last,
OutputIterator result,
command_queue &queue,
const wait_list &events,
typename boost::enable_if<
mpl::and_<
mpl::not_<
@@ -227,8 +231,8 @@ dispatch_copy(InputIterator first,
boost::shared_ptr<parameter_cache> parameters =
detail::parameter_cache::get_global_cache(device);
size_t map_copy_threshold;
size_t direct_copy_threshold;
uint_ map_copy_threshold;
uint_ direct_copy_threshold;
// calculate default values of thresholds
if (device.type() & device::gpu) {
@@ -258,13 +262,15 @@ dispatch_copy(InputIterator first,
// [0; map_copy_threshold) -> copy_to_device_map()
if(input_size_bytes < map_copy_threshold) {
return copy_to_device_map(first, last, result, queue);
return copy_to_device_map(first, last, result, queue, events);
}
// [map_copy_threshold; direct_copy_threshold) -> convert [first; last)
// on host and then perform copy_to_device()
else if(input_size_bytes < direct_copy_threshold) {
std::vector<output_type> vector(first, last);
return copy_to_device(vector.begin(), vector.end(), result, queue);
return copy_to_device(
vector.begin(), vector.end(), result, queue, events
);
}
// [direct_copy_threshold; inf) -> map [first; last) to device and
@@ -275,7 +281,7 @@ dispatch_copy(InputIterator first,
// return the result.
// At this point we are sure that count > 1 (first != last), so event
// returned by dispatch_copy_async() must be valid.
return dispatch_copy_async(first, last, result, queue).get();
return dispatch_copy_async(first, last, result, queue, events).get();
}
// host -> device
@@ -286,6 +292,7 @@ dispatch_copy(InputIterator first,
InputIterator last,
OutputIterator result,
command_queue &queue,
const wait_list &events,
typename boost::enable_if<
mpl::and_<
mpl::not_<
@@ -310,8 +317,8 @@ dispatch_copy(InputIterator first,
boost::shared_ptr<parameter_cache> parameters =
detail::parameter_cache::get_global_cache(device);
size_t map_copy_threshold;
size_t direct_copy_threshold;
uint_ map_copy_threshold;
uint_ direct_copy_threshold;
// calculate default values of thresholds
if (device.type() & device::gpu) {
@@ -345,12 +352,12 @@ dispatch_copy(InputIterator first,
// copy_to_device_map() is used for every input
if(input_size_bytes < map_copy_threshold
|| direct_copy_threshold <= map_copy_threshold) {
return copy_to_device_map(first, last, result, queue);
return copy_to_device_map(first, last, result, queue, events);
}
// [map_copy_threshold; inf) -> convert [first; last)
// on host and then perform copy_to_device()
std::vector<output_type> vector(first, last);
return copy_to_device(vector.begin(), vector.end(), result, queue);
return copy_to_device(vector.begin(), vector.end(), result, queue, events);
}
// device -> host (async)
@@ -360,6 +367,7 @@ dispatch_copy_async(InputIterator first,
InputIterator last,
OutputIterator result,
command_queue &queue,
const wait_list &events,
typename boost::enable_if<
mpl::and_<
is_device_iterator<InputIterator>,
@@ -375,7 +383,7 @@ dispatch_copy_async(InputIterator first,
"copy_async() is only supported for contiguous host iterators"
);
return copy_to_host_async(first, last, result, queue);
return copy_to_host_async(first, last, result, queue, events);
}
// device -> host (async)
@@ -386,6 +394,7 @@ dispatch_copy_async(InputIterator first,
InputIterator last,
OutputIterator result,
command_queue &queue,
const wait_list &events,
typename boost::enable_if<
mpl::and_<
is_device_iterator<InputIterator>,
@@ -426,7 +435,8 @@ dispatch_copy_async(InputIterator first,
first,
last,
make_buffer_iterator<output_type>(mapped_host),
queue
queue,
events
);
// update host memory asynchronously by maping and unmaping memory
event map_event;
@@ -451,6 +461,7 @@ dispatch_copy(InputIterator first,
InputIterator last,
OutputIterator result,
command_queue &queue,
const wait_list &events,
typename boost::enable_if<
mpl::and_<
is_device_iterator<InputIterator>,
@@ -465,7 +476,7 @@ dispatch_copy(InputIterator first,
>
>::type* = 0)
{
return copy_to_host(first, last, result, queue);
return copy_to_host(first, last, result, queue, events);
}
// device -> host
@@ -478,6 +489,7 @@ dispatch_copy(InputIterator first,
InputIterator last,
OutputIterator result,
command_queue &queue,
const wait_list &events,
typename boost::enable_if<
mpl::and_<
is_device_iterator<InputIterator>,
@@ -505,8 +517,8 @@ dispatch_copy(InputIterator first,
boost::shared_ptr<parameter_cache> parameters =
detail::parameter_cache::get_global_cache(device);
size_t map_copy_threshold;
size_t direct_copy_threshold;
uint_ map_copy_threshold;
uint_ direct_copy_threshold;
// calculate default values of thresholds
if (device.type() & device::gpu) {
@@ -540,12 +552,12 @@ dispatch_copy(InputIterator first,
// copy_to_host_map() is used for every input
if(input_size_bytes < map_copy_threshold
|| direct_copy_threshold <= map_copy_threshold) {
return copy_to_host_map(first, last, result, queue);
return copy_to_host_map(first, last, result, queue, events);
}
// [map_copy_threshold; inf) -> copy [first;last) to temporary vector
// then copy (and convert) to result using std::copy()
std::vector<input_type> vector(count);
copy_to_host(first, last, vector.begin(), queue);
copy_to_host(first, last, vector.begin(), queue, events);
return std::copy(vector.begin(), vector.end(), result);
}
@@ -559,6 +571,7 @@ dispatch_copy(InputIterator first,
InputIterator last,
OutputIterator result,
command_queue &queue,
const wait_list &events,
typename boost::enable_if<
mpl::and_<
is_device_iterator<InputIterator>,
@@ -587,8 +600,8 @@ dispatch_copy(InputIterator first,
boost::shared_ptr<parameter_cache> parameters =
detail::parameter_cache::get_global_cache(device);
size_t map_copy_threshold;
size_t direct_copy_threshold;
uint_ map_copy_threshold;
uint_ direct_copy_threshold;
// calculate default values of thresholds
if (device.type() & device::gpu) {
@@ -618,13 +631,13 @@ dispatch_copy(InputIterator first,
// [0; map_copy_threshold) -> copy_to_host_map()
if(input_size_bytes < map_copy_threshold) {
return copy_to_host_map(first, last, result, queue);
return copy_to_host_map(first, last, result, queue, events);
}
// [map_copy_threshold; direct_copy_threshold) -> copy [first;last) to
// temporary vector then copy (and convert) to result using std::copy()
else if(input_size_bytes < direct_copy_threshold) {
std::vector<input_type> vector(count);
copy_to_host(first, last, vector.begin(), queue);
copy_to_host(first, last, vector.begin(), queue, events);
return std::copy(vector.begin(), vector.end(), result);
}
@@ -636,7 +649,7 @@ dispatch_copy(InputIterator first,
// return the result.
// At this point we are sure that count > 1 (first != last), so event
// returned by dispatch_copy_async() must be valid.
return dispatch_copy_async(first, last, result, queue).get();
return dispatch_copy_async(first, last, result, queue, events).get();
}
// device -> device
@@ -646,6 +659,7 @@ dispatch_copy(InputIterator first,
InputIterator last,
OutputIterator result,
command_queue &queue,
const wait_list &events,
typename boost::enable_if<
mpl::and_<
is_device_iterator<InputIterator>,
@@ -658,7 +672,7 @@ dispatch_copy(InputIterator first,
>
>::type* = 0)
{
return copy_on_device(first, last, result, queue);
return copy_on_device(first, last, result, queue, events);
}
// device -> device (specialization for buffer iterators)
@@ -668,6 +682,7 @@ dispatch_copy(InputIterator first,
InputIterator last,
OutputIterator result,
command_queue &queue,
const wait_list &events,
typename boost::enable_if<
mpl::and_<
is_device_iterator<InputIterator>,
@@ -691,7 +706,8 @@ dispatch_copy(InputIterator first,
result.get_buffer(),
first.get_index() * sizeof(value_type),
result.get_index() * sizeof(value_type),
static_cast<size_t>(n) * sizeof(value_type));
static_cast<size_t>(n) * sizeof(value_type),
events);
return result + n;
}
@@ -702,6 +718,7 @@ dispatch_copy_async(InputIterator first,
InputIterator last,
OutputIterator result,
command_queue &queue,
const wait_list &events,
typename boost::enable_if<
mpl::and_<
is_device_iterator<InputIterator>,
@@ -714,7 +731,7 @@ dispatch_copy_async(InputIterator first,
>
>::type* = 0)
{
return copy_on_device_async(first, last, result, queue);
return copy_on_device_async(first, last, result, queue, events);
}
// device -> device (async, specialization for buffer iterators)
@@ -724,6 +741,7 @@ dispatch_copy_async(InputIterator first,
InputIterator last,
OutputIterator result,
command_queue &queue,
const wait_list &events,
typename boost::enable_if<
mpl::and_<
is_device_iterator<InputIterator>,
@@ -749,7 +767,8 @@ dispatch_copy_async(InputIterator first,
result.get_buffer(),
first.get_index() * sizeof(value_type),
result.get_index() * sizeof(value_type),
static_cast<size_t>(n) * sizeof(value_type)
static_cast<size_t>(n) * sizeof(value_type),
events
);
return make_future(result + n, event_);
@@ -762,12 +781,14 @@ dispatch_copy(InputIterator first,
InputIterator last,
OutputIterator result,
command_queue &queue,
const wait_list &events,
typename boost::enable_if_c<
!is_device_iterator<InputIterator>::value &&
!is_device_iterator<OutputIterator>::value
>::type* = 0)
{
(void) queue;
(void) events;
return std::copy(first, last, result);
}
@@ -826,14 +847,17 @@ dispatch_copy(InputIterator first,
/// );
/// \endcode
///
/// Space complexity: \Omega(1)
///
/// \see copy_n(), copy_if(), copy_async()
template<class InputIterator, class OutputIterator>
inline OutputIterator copy(InputIterator first,
InputIterator last,
OutputIterator result,
command_queue &queue = system::default_queue())
command_queue &queue = system::default_queue(),
const wait_list &events = wait_list())
{
return detail::dispatch_copy(first, last, result, queue);
return detail::dispatch_copy(first, last, result, queue, events);
}
/// Copies the values in the range [\p first, \p last) to the range
@@ -845,9 +869,10 @@ inline future<OutputIterator>
copy_async(InputIterator first,
InputIterator last,
OutputIterator result,
command_queue &queue = system::default_queue())
command_queue &queue = system::default_queue(),
const wait_list &events = wait_list())
{
return detail::dispatch_copy_async(first, last, result, queue);
return detail::dispatch_copy_async(first, last, result, queue, events);
}
} // end compute namespace
boost/compute/algorithm/copy_if.hpp
+9
View File
@@ -11,8 +11,11 @@
#ifndef BOOST_COMPUTE_ALGORITHM_COPY_IF_HPP
#define BOOST_COMPUTE_ALGORITHM_COPY_IF_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/algorithm/transform_if.hpp>
#include <boost/compute/functional/identity.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -27,6 +30,8 @@ inline OutputIterator copy_index_if(InputIterator first,
Predicate predicate,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
typedef typename std::iterator_traits<InputIterator>::value_type T;
return detail::transform_if_impl(
@@ -38,6 +43,8 @@ inline OutputIterator copy_index_if(InputIterator first,
/// Copies each element in the range [\p first, \p last) for which
/// \p predicate returns \c true to the range beginning at \p result.
///
/// Space complexity: \Omega(2n)
template<class InputIterator, class OutputIterator, class Predicate>
inline OutputIterator copy_if(InputIterator first,
InputIterator last,
@@ -45,6 +52,8 @@ inline OutputIterator copy_if(InputIterator first,
Predicate predicate,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
typedef typename std::iterator_traits<InputIterator>::value_type T;
return ::boost::compute::transform_if(
boost/compute/algorithm/copy_n.hpp
+6 -2
View File
@@ -30,19 +30,23 @@ namespace compute {
/// boost::compute::copy_n(values, 4, vec.begin(), queue);
/// \endcode
///
/// Space complexity: \Omega(1)
///
/// \see copy()
template<class InputIterator, class Size, class OutputIterator>
inline OutputIterator copy_n(InputIterator first,
Size count,
OutputIterator result,
command_queue &queue = system::default_queue())
command_queue &queue = system::default_queue(),
const wait_list &events = wait_list())
{
typedef typename std::iterator_traits<InputIterator>::difference_type difference_type;
return ::boost::compute::copy(first,
first + static_cast<difference_type>(count),
result,
queue);
queue,
events);
}
} // end compute namespace
boost/compute/algorithm/count.hpp
+7
View File
@@ -11,11 +11,14 @@
#ifndef BOOST_COMPUTE_ALGORITHM_COUNT_HPP
#define BOOST_COMPUTE_ALGORITHM_COUNT_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/lambda.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/algorithm/count_if.hpp>
#include <boost/compute/type_traits/vector_size.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -23,6 +26,9 @@ namespace compute {
/// Returns the number of occurrences of \p value in the range
/// [\p first, \p last).
///
/// Space complexity on CPUs: \Omega(1)<br>
/// Space complexity on GPUs: \Omega(n)
///
/// \see count_if()
template<class InputIterator, class T>
inline size_t count(InputIterator first,
@@ -30,6 +36,7 @@ inline size_t count(InputIterator first,
const T &value,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
typedef typename std::iterator_traits<InputIterator>::value_type value_type;
using ::boost::compute::_1;
boost/compute/algorithm/count_if.hpp
+7
View File
@@ -11,6 +11,8 @@
#ifndef BOOST_COMPUTE_ALGORITHM_COUNT_IF_HPP
#define BOOST_COMPUTE_ALGORITHM_COUNT_IF_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/device.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
@@ -19,18 +21,23 @@
#include <boost/compute/algorithm/detail/count_if_with_threads.hpp>
#include <boost/compute/algorithm/detail/serial_count_if.hpp>
#include <boost/compute/detail/iterator_range_size.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
/// Returns the number of elements in the range [\p first, \p last)
/// for which \p predicate returns \c true.
///
/// Space complexity on CPUs: \Omega(1)<br>
/// Space complexity on GPUs: \Omega(n)
template<class InputIterator, class Predicate>
inline size_t count_if(InputIterator first,
InputIterator last,
Predicate predicate,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
const device &device = queue.get_device();
size_t input_size = detail::iterator_range_size(first, last);
boost/compute/algorithm/detail/copy_on_device.hpp
+27 -18
View File
@@ -32,7 +32,8 @@ template<class InputIterator, class OutputIterator>
inline event copy_on_device_cpu(InputIterator first,
OutputIterator result,
size_t count,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
meta_kernel k("copy");
const device& device = queue.get_device();
@@ -52,14 +53,15 @@ inline event copy_on_device_cpu(InputIterator first,
size_t global_work_size = device.compute_units();
if(count <= 1024) global_work_size = 1;
return k.exec_1d(queue, 0, global_work_size);
return k.exec_1d(queue, 0, global_work_size, events);
}
template<class InputIterator, class OutputIterator>
inline event copy_on_device_gpu(InputIterator first,
OutputIterator result,
size_t count,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
typedef typename std::iterator_traits<InputIterator>::value_type input_type;
@@ -86,14 +88,15 @@ inline event copy_on_device_gpu(InputIterator first,
k.add_set_arg<const uint_>("count", static_cast<uint_>(count));
size_t global_work_size = calculate_work_size(count, vpt, tpb);
return k.exec_1d(queue, 0, global_work_size, tpb);
return k.exec_1d(queue, 0, global_work_size, tpb, events);
}
template<class InputIterator, class OutputIterator>
inline event dispatch_copy_on_device(InputIterator first,
InputIterator last,
OutputIterator result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
const size_t count = detail::iterator_range_size(first, last);
@@ -108,18 +111,19 @@ inline event dispatch_copy_on_device(InputIterator first,
// See https://github.com/boostorg/compute/pull/626
if((device.type() & device::cpu) && !is_apple_platform_device(device))
{
return copy_on_device_cpu(first, result, count, queue);
return copy_on_device_cpu(first, result, count, queue, events);
}
return copy_on_device_gpu(first, result, count, queue);
return copy_on_device_gpu(first, result, count, queue, events);
}
template<class InputIterator, class OutputIterator>
inline OutputIterator copy_on_device(InputIterator first,
InputIterator last,
OutputIterator result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
dispatch_copy_on_device(first, last, result, queue);
dispatch_copy_on_device(first, last, result, queue, events);
return result + std::distance(first, last);
}
@@ -127,9 +131,11 @@ template<class InputIterator>
inline discard_iterator copy_on_device(InputIterator first,
InputIterator last,
discard_iterator result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
(void) queue;
(void) events;
return result + std::distance(first, last);
}
@@ -138,19 +144,21 @@ template<class InputIterator, class OutputIterator>
inline future<OutputIterator> copy_on_device_async(InputIterator first,
InputIterator last,
OutputIterator result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
event event_ = dispatch_copy_on_device(first, last, result, queue);
event event_ = dispatch_copy_on_device(first, last, result, queue, events);
return make_future(result + std::distance(first, last), event_);
}
#ifdef CL_VERSION_2_0
#ifdef BOOST_COMPUTE_CL_VERSION_2_0
// copy_on_device() specialization for svm_ptr
template<class T>
inline svm_ptr<T> copy_on_device(svm_ptr<T> first,
svm_ptr<T> last,
svm_ptr<T> result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
size_t count = iterator_range_size(first, last);
if(count == 0){
@@ -158,7 +166,7 @@ inline svm_ptr<T> copy_on_device(svm_ptr<T> first,
}
queue.enqueue_svm_memcpy(
result.get(), first.get(), count * sizeof(T)
result.get(), first.get(), count * sizeof(T), events
);
return result + count;
@@ -168,7 +176,8 @@ template<class T>
inline future<svm_ptr<T> > copy_on_device_async(svm_ptr<T> first,
svm_ptr<T> last,
svm_ptr<T> result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
size_t count = iterator_range_size(first, last);
if(count == 0){
@@ -176,12 +185,12 @@ inline future<svm_ptr<T> > copy_on_device_async(svm_ptr<T> first,
}
event event_ = queue.enqueue_svm_memcpy_async(
result.get(), first.get(), count * sizeof(T)
result.get(), first.get(), count * sizeof(T), events
);
return make_future(result + count, event_);
}
#endif // CL_VERSION_2_0
#endif // BOOST_COMPUTE_CL_VERSION_2_0
} // end detail namespace
} // end compute namespace
boost/compute/algorithm/detail/copy_to_device.hpp
+25 -14
View File
@@ -28,7 +28,8 @@ template<class HostIterator, class DeviceIterator>
inline DeviceIterator copy_to_device(HostIterator first,
HostIterator last,
DeviceIterator result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
typedef typename
std::iterator_traits<DeviceIterator>::value_type
@@ -47,7 +48,8 @@ inline DeviceIterator copy_to_device(HostIterator first,
queue.enqueue_write_buffer(result.get_buffer(),
offset * sizeof(value_type),
count * sizeof(value_type),
::boost::addressof(*first));
::boost::addressof(*first),
events);
return result + static_cast<difference_type>(count);
}
@@ -56,7 +58,8 @@ template<class HostIterator, class DeviceIterator>
inline DeviceIterator copy_to_device_map(HostIterator first,
HostIterator last,
DeviceIterator result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
typedef typename
std::iterator_traits<DeviceIterator>::value_type
@@ -78,7 +81,8 @@ inline DeviceIterator copy_to_device_map(HostIterator first,
result.get_buffer(),
CL_MAP_WRITE,
offset * sizeof(value_type),
count * sizeof(value_type)
count * sizeof(value_type),
events
)
);
@@ -99,7 +103,8 @@ template<class HostIterator, class DeviceIterator>
inline future<DeviceIterator> copy_to_device_async(HostIterator first,
HostIterator last,
DeviceIterator result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
typedef typename
std::iterator_traits<DeviceIterator>::value_type
@@ -119,18 +124,20 @@ inline future<DeviceIterator> copy_to_device_async(HostIterator first,
queue.enqueue_write_buffer_async(result.get_buffer(),
offset * sizeof(value_type),
count * sizeof(value_type),
::boost::addressof(*first));
::boost::addressof(*first),
events);
return make_future(result + static_cast<difference_type>(count), event_);
}
#ifdef CL_VERSION_2_0
#ifdef BOOST_COMPUTE_CL_VERSION_2_0
// copy_to_device() specialization for svm_ptr
template<class HostIterator, class T>
inline svm_ptr<T> copy_to_device(HostIterator first,
HostIterator last,
svm_ptr<T> result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
size_t count = iterator_range_size(first, last);
if(count == 0){
@@ -138,7 +145,7 @@ inline svm_ptr<T> copy_to_device(HostIterator first,
}
queue.enqueue_svm_memcpy(
result.get(), ::boost::addressof(*first), count * sizeof(T)
result.get(), ::boost::addressof(*first), count * sizeof(T), events
);
return result + count;
@@ -148,7 +155,8 @@ template<class HostIterator, class T>
inline future<svm_ptr<T> > copy_to_device_async(HostIterator first,
HostIterator last,
svm_ptr<T> result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
size_t count = iterator_range_size(first, last);
if(count == 0){
@@ -156,7 +164,7 @@ inline future<svm_ptr<T> > copy_to_device_async(HostIterator first,
}
event event_ = queue.enqueue_svm_memcpy_async(
result.get(), ::boost::addressof(*first), count * sizeof(T)
result.get(), ::boost::addressof(*first), count * sizeof(T), events
);
return make_future(result + count, event_);
@@ -166,7 +174,8 @@ template<class HostIterator, class T>
inline svm_ptr<T> copy_to_device_map(HostIterator first,
HostIterator last,
svm_ptr<T> result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
size_t count = iterator_range_size(first, last);
if(count == 0){
@@ -174,7 +183,9 @@ inline svm_ptr<T> copy_to_device_map(HostIterator first,
}
// map
queue.enqueue_svm_map(result.get(), count * sizeof(T), CL_MAP_WRITE);
queue.enqueue_svm_map(
result.get(), count * sizeof(T), CL_MAP_WRITE, events
);
// copy [first; last) to result buffer
std::copy(first, last, static_cast<T*>(result.get()));
@@ -184,7 +195,7 @@ inline svm_ptr<T> copy_to_device_map(HostIterator first,
return result + count;
}
#endif // CL_VERSION_2_0
#endif // BOOST_COMPUTE_CL_VERSION_2_0
} // end detail namespace
} // end compute namespace
boost/compute/algorithm/detail/copy_to_host.hpp
+23 -14
View File
@@ -29,7 +29,8 @@ template<class DeviceIterator, class HostIterator>
inline HostIterator copy_to_host(DeviceIterator first,
DeviceIterator last,
HostIterator result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
typedef typename
std::iterator_traits<DeviceIterator>::value_type
@@ -46,7 +47,8 @@ inline HostIterator copy_to_host(DeviceIterator first,
queue.enqueue_read_buffer(buffer,
offset * sizeof(value_type),
count * sizeof(value_type),
::boost::addressof(*result));
::boost::addressof(*result),
events);
return iterator_plus_distance(result, count);
}
@@ -55,7 +57,8 @@ template<class DeviceIterator, class HostIterator>
inline HostIterator copy_to_host_map(DeviceIterator first,
DeviceIterator last,
HostIterator result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
typedef typename
std::iterator_traits<DeviceIterator>::value_type
@@ -77,7 +80,8 @@ inline HostIterator copy_to_host_map(DeviceIterator first,
first.get_buffer(),
CL_MAP_READ,
offset * sizeof(value_type),
count * sizeof(value_type)
count * sizeof(value_type),
events
)
);
@@ -102,7 +106,8 @@ template<class DeviceIterator, class HostIterator>
inline future<HostIterator> copy_to_host_async(DeviceIterator first,
DeviceIterator last,
HostIterator result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
typedef typename
std::iterator_traits<DeviceIterator>::value_type
@@ -120,18 +125,20 @@ inline future<HostIterator> copy_to_host_async(DeviceIterator first,
queue.enqueue_read_buffer_async(buffer,
offset * sizeof(value_type),
count * sizeof(value_type),
::boost::addressof(*result));
::boost::addressof(*result),
events);
return make_future(iterator_plus_distance(result, count), event_);
}
#ifdef CL_VERSION_2_0
#ifdef BOOST_COMPUTE_CL_VERSION_2_0
// copy_to_host() specialization for svm_ptr
template<class T, class HostIterator>
inline HostIterator copy_to_host(svm_ptr<T> first,
svm_ptr<T> last,
HostIterator result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
size_t count = iterator_range_size(first, last);
if(count == 0){
@@ -139,7 +146,7 @@ inline HostIterator copy_to_host(svm_ptr<T> first,
}
queue.enqueue_svm_memcpy(
::boost::addressof(*result), first.get(), count * sizeof(T)
::boost::addressof(*result), first.get(), count * sizeof(T), events
);
return result + count;
@@ -149,7 +156,8 @@ template<class T, class HostIterator>
inline future<HostIterator> copy_to_host_async(svm_ptr<T> first,
svm_ptr<T> last,
HostIterator result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
size_t count = iterator_range_size(first, last);
if(count == 0){
@@ -157,7 +165,7 @@ inline future<HostIterator> copy_to_host_async(svm_ptr<T> first,
}
event event_ = queue.enqueue_svm_memcpy_async(
::boost::addressof(*result), first.get(), count * sizeof(T)
::boost::addressof(*result), first.get(), count * sizeof(T), events
);
return make_future(iterator_plus_distance(result, count), event_);
@@ -167,7 +175,8 @@ template<class T, class HostIterator>
inline HostIterator copy_to_host_map(svm_ptr<T> first,
svm_ptr<T> last,
HostIterator result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
size_t count = iterator_range_size(first, last);
if(count == 0){
@@ -175,7 +184,7 @@ inline HostIterator copy_to_host_map(svm_ptr<T> first,
}
// map
queue.enqueue_svm_map(first.get(), count * sizeof(T), CL_MAP_READ);
queue.enqueue_svm_map(first.get(), count * sizeof(T), CL_MAP_READ, events);
// copy [first; last) to result
std::copy(
@@ -189,7 +198,7 @@ inline HostIterator copy_to_host_map(svm_ptr<T> first,
return iterator_plus_distance(result, count);
}
#endif // CL_VERSION_2_0
#endif // BOOST_COMPUTE_CL_VERSION_2_0
} // end detail namespace
} // end compute namespace
boost/compute/algorithm/detail/find_extrema.hpp
+1 -1
View File
@@ -56,7 +56,7 @@ inline InputIterator find_extrema(InputIterator first,
// use serial method for OpenCL version 1.0 due to
// problems with atomic_cmpxchg()
#ifndef CL_VERSION_1_1
#ifndef BOOST_COMPUTE_CL_VERSION_1_1
return serial_find_extrema(first, last, compare, find_minimum, queue);
#endif
boost/compute/algorithm/detail/find_extrema_with_reduce.hpp
+1
View File
@@ -246,6 +246,7 @@ inline void find_extrema_with_reduce(InputIterator input,
);
}
// Space complexity: \Omega(2 * work-group-size * work-groups-per-compute-unit)
template<class InputIterator, class Compare>
InputIterator find_extrema_with_reduce(InputIterator first,
InputIterator last,
boost/compute/algorithm/detail/find_if_with_atomics.hpp
+1
View File
@@ -153,6 +153,7 @@ inline InputIterator find_if_with_atomics_multiple_vpt(InputIterator first,
return first + static_cast<difference_type>(index.read(queue));
}
// Space complexity: O(1)
template<class InputIterator, class UnaryPredicate>
inline InputIterator find_if_with_atomics(InputIterator first,
InputIterator last,
boost/compute/algorithm/detail/merge_sort_on_gpu.hpp
+15 -3
View File
@@ -91,6 +91,7 @@ inline size_t bitonic_block_sort(KeyIterator keys_first,
command_queue &queue)
{
typedef typename std::iterator_traits<KeyIterator>::value_type key_type;
typedef typename std::iterator_traits<ValueIterator>::value_type value_type;
meta_kernel k("bitonic_block_sort");
size_t count_arg = k.add_arg<const uint_>("count");
@@ -169,8 +170,12 @@ inline size_t bitonic_block_sort(KeyIterator keys_first,
k.decl<bool>("compare") << " = " <<
compare(k.var<key_type>("sibling_key"),
k.var<key_type>("my_key")) << ";\n" <<
k.decl<bool>("equal") << " = !(compare || " <<
compare(k.var<key_type>("my_key"),
k.var<key_type>("sibling_key")) << ");\n" <<
k.decl<bool>("swap") <<
" = compare ^ (sibling_idx < lid) ^ direction;\n" <<
"swap = equal ? false : swap;\n" <<
"my_key = swap ? sibling_key : my_key;\n";
if(sort_by_key)
{
@@ -219,8 +224,12 @@ inline size_t bitonic_block_sort(KeyIterator keys_first,
k.decl<bool>("compare") << " = " <<
compare(k.var<key_type>("sibling_key"),
k.var<key_type>("my_key")) << ";\n" <<
k.decl<bool>("equal") << " = !(compare || " <<
compare(k.var<key_type>("my_key"),
k.var<key_type>("sibling_key")) << ");\n" <<
k.decl<bool>("swap") <<
" = compare ^ (sibling_idx < lid);\n" <<
"swap = equal ? false : swap;\n" <<
"my_key = swap ? sibling_key : my_key;\n";
if(sort_by_key)
{
@@ -249,8 +258,11 @@ inline size_t bitonic_block_sort(KeyIterator keys_first,
k.var<key_type>("my_key") << ";\n";
if(sort_by_key)
{
k << values_first[k.var<const uint_>("gid")] << " = " <<
values_first[k.var<const uint_>("offset + my_index")] << ";\n";
k <<
k.decl<value_type>("my_value") << " = " <<
values_first[k.var<const uint_>("offset + my_index")] << ";\n" <<
"barrier(CLK_GLOBAL_MEM_FENCE);\n" <<
values_first[k.var<const uint_>("gid")] << " = my_value;\n";
}
k <<
// end if
@@ -418,7 +430,7 @@ inline void merge_blocks_on_gpu(KeyIterator keys_first,
");\n" <<
"left_idx = equal ? mid_idx + 1 : left_idx + 1;\n" <<
"right_idx = equal ? right_idx : mid_idx;\n" <<
"upper_key = equal ? upper_key : " <<
"upper_key = " <<
keys_first[k.var<const uint_>("left_idx")] << ";\n" <<
"}\n" <<
"}\n" <<
boost/compute/algorithm/detail/radix_sort.hpp
+9 -1
View File
@@ -17,6 +17,9 @@
#include <boost/type_traits/is_signed.hpp>
#include <boost/type_traits/is_floating_point.hpp>
#include <boost/mpl/and.hpp>
#include <boost/mpl/not.hpp>
#include <boost/compute/kernel.hpp>
#include <boost/compute/program.hpp>
#include <boost/compute/command_queue.hpp>
@@ -173,8 +176,10 @@ const char radix_sort_source[] =
" uint sum = 0;\n"
" for(uint i = 0; i < K2_BITS; i++){\n"
" uint x = global_offsets[i] + last_block_offsets[i];\n"
" mem_fence(CLK_GLOBAL_MEM_FENCE);\n" // work around the RX 500/Vega bug, see #811
" global_offsets[i] = sum;\n"
" sum += x;\n"
" mem_fence(CLK_GLOBAL_MEM_FENCE);\n" // work around the RX Vega bug, see #811
" }\n"
"}\n"
@@ -305,9 +310,12 @@ inline void radix_sort_impl(const buffer_iterator<T> first,
options << " -DASC";
}
// get type definition if it is a custom struct
std::string custom_type_def = boost::compute::type_definition<T2>() + "\n";
// load radix sort program
program radix_sort_program = cache->get_or_build(
cache_key, options.str(), radix_sort_source, context
cache_key, options.str(), custom_type_def + radix_sort_source, context
);
kernel count_kernel(radix_sort_program, "count");
boost/compute/algorithm/detail/scan_on_cpu.hpp
+4 -6
View File
@@ -76,10 +76,10 @@ inline OutputIterator scan_on_cpu(InputIterator first,
k.add_arg<output_type *>(memory_object::global_memory, "block_partial_sums");
k <<
"uint block = " <<
"(uint)ceil(((float)count)/(get_global_size(0) + 1));\n" <<
"uint block = (count + get_global_size(0))/(get_global_size(0) + 1);\n" <<
"uint index = get_global_id(0) * block;\n" <<
"uint end = min(count, index + block);\n";
"uint end = min(count, index + block);\n" <<
"if(index >= end) return;\n";
if(!exclusive){
k <<
@@ -154,11 +154,9 @@ inline OutputIterator scan_on_cpu(InputIterator first,
l.add_arg<output_type *>(memory_object::global_memory, "block_partial_sums");
l <<
"uint block = " <<
"(uint)ceil(((float)count)/(get_global_size(0) + 1));\n" <<
"uint block = (count + get_global_size(0))/(get_global_size(0) + 1);\n" <<
"uint index = block + get_global_id(0) * block;\n" <<
"uint end = min(count, index + block);\n" <<
k.decl<output_type>("sum") << " = block_partial_sums[0];\n" <<
"for(uint i = 0; i < get_global_id(0); i++) {\n" <<
"sum = " << op(k.var<output_type>("sum"),
boost/compute/algorithm/detail/scan_on_gpu.hpp
+1 -1
View File
@@ -108,7 +108,7 @@ public:
// store sum for the block
if(exclusive){
*this <<
"if(lid == block_size - 1){\n" <<
"if(lid == block_size - 1 && gid < count) {\n" <<
" block_sums[get_group_id(0)] = " <<
op(first[expr<cl_uint>("gid")], var<T>("scratch[lid]")) <<
";\n" <<
boost/compute/algorithm/detail/serial_reduce.hpp
+1
View File
@@ -20,6 +20,7 @@ namespace boost {
namespace compute {
namespace detail {
// Space complexity: O(1)
template<class InputIterator, class OutputIterator, class BinaryFunction>
inline void serial_reduce(InputIterator first,
InputIterator last,
boost/compute/algorithm/detail/serial_reduce_by_key.hpp
+2 -4
View File
@@ -55,11 +55,9 @@ inline size_t serial_reduce_by_key(InputKeyIterator keys_first,
size_t result_size_arg = k.add_arg<uint_ *>(memory_object::global_memory,
"result_size");
convert<result_type> to_result_type;
k <<
k.decl<result_type>("result") <<
" = " << to_result_type(values_first[0]) << ";\n" <<
" = " << values_first[0] << ";\n" <<
k.decl<key_type>("previous_key") << " = " << keys_first[0] << ";\n" <<
k.decl<result_type>("value") << ";\n" <<
k.decl<key_type>("key") << ";\n" <<
@@ -70,7 +68,7 @@ inline size_t serial_reduce_by_key(InputKeyIterator keys_first,
values_result[0] << " = result;\n" <<
"for(ulong i = 1; i < count; i++) {\n" <<
" value = " << to_result_type(values_first[k.var<uint_>("i")]) << ";\n" <<
" value = " << values_first[k.var<uint_>("i")] << ";\n" <<
" key = " << keys_first[k.var<uint_>("i")] << ";\n" <<
" if (" << predicate(k.var<key_type>("previous_key"),
k.var<key_type>("key")) << ") {\n" <<
boost/compute/algorithm/equal.hpp
+9
View File
@@ -11,21 +11,28 @@
#ifndef BOOST_COMPUTE_ALGORITHM_EQUAL_HPP
#define BOOST_COMPUTE_ALGORITHM_EQUAL_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/algorithm/mismatch.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
/// Returns \c true if the range [\p first1, \p last1) and the range
/// beginning at \p first2 are equal.
///
/// Space complexity: \Omega(1)
template<class InputIterator1, class InputIterator2>
inline bool equal(InputIterator1 first1,
InputIterator1 last1,
InputIterator2 first2,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator1>::value);
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator2>::value);
return ::boost::compute::mismatch(first1,
last1,
first2,
@@ -40,6 +47,8 @@ inline bool equal(InputIterator1 first1,
InputIterator2 last2,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator1>::value);
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator2>::value);
if(std::distance(first1, last1) != std::distance(first2, last2)){
return false;
}
boost/compute/algorithm/equal_range.hpp
+6
View File
@@ -13,16 +13,21 @@
#include <utility>
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/algorithm/lower_bound.hpp>
#include <boost/compute/algorithm/upper_bound.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
/// Returns a pair of iterators containing the range of values equal
/// to \p value in the sorted range [\p first, \p last).
///
/// Space complexity: \Omega(1)
template<class InputIterator, class T>
inline std::pair<InputIterator, InputIterator>
equal_range(InputIterator first,
@@ -30,6 +35,7 @@ equal_range(InputIterator first,
const T &value,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
return std::make_pair(
::boost::compute::lower_bound(first, last, value, queue),
::boost::compute::upper_bound(first, last, value, queue)
boost/compute/algorithm/exclusive_scan.hpp
+13
View File
@@ -11,10 +11,13 @@
#ifndef BOOST_COMPUTE_ALGORITHM_EXCLUSIVE_SCAN_HPP
#define BOOST_COMPUTE_ALGORITHM_EXCLUSIVE_SCAN_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/functional.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/algorithm/detail/scan.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -44,6 +47,10 @@ namespace compute {
///
/// \snippet test/test_scan.cpp exclusive_scan_int_multiplies
///
/// Space complexity on GPUs: \Omega(n)<br>
/// Space complexity on GPUs when \p first == \p result: \Omega(2n)<br>
/// Space complexity on CPUs: \Omega(1)
///
/// \see inclusive_scan()
template<class InputIterator, class OutputIterator, class T, class BinaryOperator>
inline OutputIterator
@@ -54,6 +61,8 @@ exclusive_scan(InputIterator first,
BinaryOperator binary_op,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
return detail::scan(first, last, result, true, init, binary_op, queue);
}
@@ -66,6 +75,8 @@ exclusive_scan(InputIterator first,
T init,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
typedef typename
std::iterator_traits<OutputIterator>::value_type output_type;
@@ -82,6 +93,8 @@ exclusive_scan(InputIterator first,
OutputIterator result,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
typedef typename
std::iterator_traits<OutputIterator>::value_type output_type;
boost/compute/algorithm/fill.hpp
+11 -4
View File
@@ -13,6 +13,7 @@
#include <iterator>
#include <boost/static_assert.hpp>
#include <boost/mpl/int.hpp>
#include <boost/mpl/vector.hpp>
#include <boost/mpl/contains.hpp>
@@ -27,6 +28,8 @@
#include <boost/compute/iterator/discard_iterator.hpp>
#include <boost/compute/detail/is_buffer_iterator.hpp>
#include <boost/compute/detail/iterator_range_size.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -64,7 +67,7 @@ inline future<void> fill_async_with_copy(BufferIterator first,
);
}
#if defined(CL_VERSION_1_2)
#if defined(BOOST_COMPUTE_CL_VERSION_1_2)
// meta-function returing true if Iterator points to a range of values
// that can be filled using clEnqueueFillBuffer(). to meet this criteria
@@ -172,7 +175,7 @@ dispatch_fill_async(BufferIterator first,
return future<void>(event_);
}
#ifdef CL_VERSION_2_0
#ifdef BOOST_COMPUTE_CL_VERSION_2_0
// specializations for svm_ptr<T>
template<class T>
inline void dispatch_fill(svm_ptr<T> first,
@@ -205,7 +208,7 @@ inline future<void> dispatch_fill_async(svm_ptr<T> first,
return future<void>(event_);
}
#endif // CL_VERSION_2_0
#endif // BOOST_COMPUTE_CL_VERSION_2_0
// default implementations
template<class BufferIterator, class T>
@@ -251,7 +254,7 @@ inline future<void> dispatch_fill_async(BufferIterator first,
{
return fill_async_with_copy(first, count, value, queue);
}
#endif // !defined(CL_VERSION_1_2)
#endif // !defined(BOOST_COMPUTE_CL_VERSION_1_2)
} // end detail namespace
@@ -271,6 +274,8 @@ inline future<void> dispatch_fill_async(BufferIterator first,
/// boost::compute::fill(vec.begin(), vec.end(), 7, queue);
/// \endcode
///
/// Space complexity: \Omega(1)
///
/// \see boost::compute::fill_n()
template<class BufferIterator, class T>
inline void fill(BufferIterator first,
@@ -278,6 +283,7 @@ inline void fill(BufferIterator first,
const T &value,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<BufferIterator>::value);
size_t count = detail::iterator_range_size(first, last);
if(count == 0){
return;
@@ -292,6 +298,7 @@ inline future<void> fill_async(BufferIterator first,
const T &value,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(detail::is_buffer_iterator<BufferIterator>::value);
size_t count = detail::iterator_range_size(first, last);
if(count == 0){
return future<void>();
boost/compute/algorithm/fill_n.hpp
+6
View File
@@ -11,15 +11,20 @@
#ifndef BOOST_COMPUTE_ALGORITHM_FILL_N_HPP
#define BOOST_COMPUTE_ALGORITHM_FILL_N_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/algorithm/fill.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
/// Fills the range [\p first, \p first + count) with \p value.
///
/// Space complexity: \Omega(1)
///
/// \see fill()
template<class BufferIterator, class Size, class T>
inline void fill_n(BufferIterator first,
@@ -27,6 +32,7 @@ inline void fill_n(BufferIterator first,
const T &value,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<BufferIterator>::value);
::boost::compute::fill(first, first + count, value, queue);
}
boost/compute/algorithm/find.hpp
+6
View File
@@ -11,23 +11,29 @@
#ifndef BOOST_COMPUTE_ALGORITHM_FIND_HPP
#define BOOST_COMPUTE_ALGORITHM_FIND_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/lambda.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/algorithm/find_if.hpp>
#include <boost/compute/type_traits/vector_size.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
/// Returns an iterator pointing to the first element in the range
/// [\p first, \p last) that equals \p value.
///
/// Space complexity: \Omega(1)
template<class InputIterator, class T>
inline InputIterator find(InputIterator first,
InputIterator last,
const T &value,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
typedef typename std::iterator_traits<InputIterator>::value_type value_type;
using ::boost::compute::_1;
boost/compute/algorithm/find_end.hpp
+10 -2
View File
@@ -11,6 +11,8 @@
#ifndef BOOST_COMPUTE_ALGORITHM_FIND_END_HPP
#define BOOST_COMPUTE_ALGORITHM_FIND_END_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/algorithm/copy.hpp>
#include <boost/compute/algorithm/detail/search_all.hpp>
#include <boost/compute/container/detail/scalar.hpp>
@@ -18,6 +20,7 @@
#include <boost/compute/detail/iterator_range_size.hpp>
#include <boost/compute/detail/meta_kernel.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -26,8 +29,8 @@ namespace detail {
///
/// \brief Helper function for find_end
///
/// Basically a copy of find_if which returns last occurence
/// instead of first occurence
/// Basically a copy of find_if which returns last occurrence
/// instead of first occurrence
///
template<class InputIterator, class UnaryPredicate>
inline InputIterator find_end_helper(InputIterator first,
@@ -90,6 +93,8 @@ inline InputIterator find_end_helper(InputIterator first,
/// \param p_last Iterator pointing to end of pattern
/// \param queue Queue on which to execute
///
/// Space complexity: \Omega(n)
///
template<class TextIterator, class PatternIterator>
inline TextIterator find_end(TextIterator t_first,
TextIterator t_last,
@@ -97,6 +102,9 @@ inline TextIterator find_end(TextIterator t_first,
PatternIterator p_last,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<TextIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<PatternIterator>::value);
const context &context = queue.get_context();
// there is no need to check if pattern starts at last n - 1 indices
boost/compute/algorithm/find_if.hpp
+6
View File
@@ -11,21 +11,27 @@
#ifndef BOOST_COMPUTE_ALGORITHM_FIND_IF_HPP
#define BOOST_COMPUTE_ALGORITHM_FIND_IF_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/algorithm/detail/find_if_with_atomics.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
/// Returns an iterator pointing to the first element in the range
/// [\p first, \p last) for which \p predicate returns \c true.
///
/// Space complexity: \Omega(1)
template<class InputIterator, class UnaryPredicate>
inline InputIterator find_if(InputIterator first,
InputIterator last,
UnaryPredicate predicate,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
return detail::find_if_with_atomics(first, last, predicate, queue);
}
boost/compute/algorithm/find_if_not.hpp
+6
View File
@@ -11,10 +11,13 @@
#ifndef BOOST_COMPUTE_ALGORITHM_FIND_IF_NOT_HPP
#define BOOST_COMPUTE_ALGORITHM_FIND_IF_NOT_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/functional.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/algorithm/find_if.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -22,6 +25,8 @@ namespace compute {
/// Returns an iterator pointing to the first element in the range
/// [\p first, \p last) for which \p predicate returns \c false.
///
/// Space complexity: \Omega(1)
///
/// \see find_if()
template<class InputIterator, class UnaryPredicate>
inline InputIterator find_if_not(InputIterator first,
@@ -29,6 +34,7 @@ inline InputIterator find_if_not(InputIterator first,
UnaryPredicate predicate,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
return ::boost::compute::find_if(
first,
last,
boost/compute/algorithm/for_each.hpp
+7
View File
@@ -11,10 +11,13 @@
#ifndef BOOST_COMPUTE_ALGORITHM_FOR_EACH_HPP
#define BOOST_COMPUTE_ALGORITHM_FOR_EACH_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/detail/meta_kernel.hpp>
#include <boost/compute/detail/iterator_range_size.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -45,6 +48,8 @@ struct for_each_kernel : public meta_kernel
/// Calls \p function on each element in the range [\p first, \p last).
///
/// Space complexity: \Omega(1)
///
/// \see transform()
template<class InputIterator, class UnaryFunction>
inline UnaryFunction for_each(InputIterator first,
@@ -52,6 +57,8 @@ inline UnaryFunction for_each(InputIterator first,
UnaryFunction function,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
detail::for_each_kernel<InputIterator, UnaryFunction> kernel(first, last, function);
kernel.exec(queue);
boost/compute/algorithm/for_each_n.hpp
+6
View File
@@ -11,7 +11,10 @@
#ifndef BOOST_COMPUTE_ALGORITHM_FOR_EACH_N_HPP
#define BOOST_COMPUTE_ALGORITHM_FOR_EACH_N_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/algorithm/for_each.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -19,6 +22,8 @@ namespace compute {
/// Calls \p function on each element in the range [\p first, \p first
/// \c + \p count).
///
/// Space complexity: \Omega(1)
///
/// \see for_each()
template<class InputIterator, class Size, class UnaryFunction>
inline UnaryFunction for_each_n(InputIterator first,
@@ -26,6 +31,7 @@ inline UnaryFunction for_each_n(InputIterator first,
UnaryFunction function,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
return ::boost::compute::for_each(first, first + count, function, queue);
}
boost/compute/algorithm/gather.hpp
+9
View File
@@ -11,6 +11,8 @@
#ifndef BOOST_COMPUTE_ALGORITHM_GATHER_HPP
#define BOOST_COMPUTE_ALGORITHM_GATHER_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/detail/iterator_range_size.hpp>
#include <boost/compute/detail/meta_kernel.hpp>
@@ -18,6 +20,7 @@
#include <boost/compute/iterator/buffer_iterator.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/type_traits/type_name.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -62,6 +65,8 @@ private:
/// to the range beginning at \p result using the input values from the range
/// beginning at \p input.
///
/// Space complexity: \Omega(1)
///
/// \see scatter()
template<class InputIterator, class MapIterator, class OutputIterator>
inline void gather(MapIterator first,
@@ -70,6 +75,10 @@ inline void gather(MapIterator first,
OutputIterator result,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<MapIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
detail::gather_kernel<InputIterator, MapIterator, OutputIterator> kernel;
kernel.set_range(first, last, input, result);
boost/compute/algorithm/generate.hpp
+6
View File
@@ -11,23 +11,29 @@
#ifndef BOOST_COMPUTE_ALGORITHM_GENERATE_HPP
#define BOOST_COMPUTE_ALGORITHM_GENERATE_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/algorithm/copy.hpp>
#include <boost/compute/iterator/function_input_iterator.hpp>
#include <boost/compute/detail/iterator_range_size.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
/// Stores the result of \p generator for each element in the range
/// [\p first, \p last).
///
/// Space complexity: \Omega(1)
template<class OutputIterator, class Generator>
inline void generate(OutputIterator first,
OutputIterator last,
Generator generator,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
size_t count = detail::iterator_range_size(first, last);
if(count == 0){
return;
boost/compute/algorithm/generate_n.hpp
+2
View File
@@ -20,6 +20,8 @@ namespace compute {
/// Stores the result of \p generator for each element in the range
/// [\p first, \p first + \p count).
///
/// Space complexity: \Omega(1)
template<class OutputIterator, class Size, class Generator>
inline void generate_n(OutputIterator first,
Size count,
boost/compute/algorithm/includes.hpp
+7
View File
@@ -13,6 +13,8 @@
#include <iterator>
#include <boost/static_assert.hpp>
#include <boost/compute/algorithm/detail/balanced_path.hpp>
#include <boost/compute/algorithm/fill_n.hpp>
#include <boost/compute/algorithm/find.hpp>
@@ -21,6 +23,7 @@
#include <boost/compute/detail/meta_kernel.hpp>
#include <boost/compute/detail/read_write_single_value.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -110,6 +113,7 @@ private:
/// \param last2 Iterator pointing to end of second set
/// \param queue Queue on which to execute
///
/// Space complexity: \Omega(distance(\p first1, \p last1) + distance(\p first2, \p last2))
template<class InputIterator1, class InputIterator2>
inline bool includes(InputIterator1 first1,
InputIterator1 last1,
@@ -117,6 +121,9 @@ inline bool includes(InputIterator1 first1,
InputIterator2 last2,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator1>::value);
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator2>::value);
size_t tile_size = 1024;
size_t count1 = detail::iterator_range_size(first1, last1);
boost/compute/algorithm/inclusive_scan.hpp
+11
View File
@@ -11,10 +11,13 @@
#ifndef BOOST_COMPUTE_ALGORITHM_INCLUSIVE_SCAN_HPP
#define BOOST_COMPUTE_ALGORITHM_INCLUSIVE_SCAN_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/functional.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/algorithm/detail/scan.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -42,6 +45,10 @@ namespace compute {
///
/// \snippet test/test_scan.cpp inclusive_scan_int_multiplies
///
/// Space complexity on GPUs: \Omega(n)<br>
/// Space complexity on GPUs when \p first == \p result: \Omega(2n)<br>
/// Space complexity on CPUs: \Omega(1)
///
/// \see exclusive_scan()
template<class InputIterator, class OutputIterator, class BinaryOperator>
inline OutputIterator
@@ -51,6 +58,8 @@ inclusive_scan(InputIterator first,
BinaryOperator binary_op,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
typedef typename
std::iterator_traits<OutputIterator>::value_type output_type;
@@ -67,6 +76,8 @@ inclusive_scan(InputIterator first,
OutputIterator result,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
typedef typename
std::iterator_traits<OutputIterator>::value_type output_type;
boost/compute/algorithm/inner_product.hpp
+10
View File
@@ -11,6 +11,8 @@
#ifndef BOOST_COMPUTE_ALGORITHM_INNER_PRODUCT_HPP
#define BOOST_COMPUTE_ALGORITHM_INNER_PRODUCT_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/functional.hpp>
#include <boost/compute/command_queue.hpp>
@@ -19,6 +21,7 @@
#include <boost/compute/iterator/transform_iterator.hpp>
#include <boost/compute/iterator/zip_iterator.hpp>
#include <boost/compute/functional/detail/unpack.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -26,6 +29,9 @@ namespace compute {
/// Returns the inner product of the elements in the range
/// [\p first1, \p last1) with the elements in the range beginning
/// at \p first2.
///
/// Space complexity: \Omega(1)<br>
/// Space complexity when binary operator is recognized as associative: \Omega(n)
template<class InputIterator1, class InputIterator2, class T>
inline T inner_product(InputIterator1 first1,
InputIterator1 last1,
@@ -33,6 +39,8 @@ inline T inner_product(InputIterator1 first1,
T init,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator1>::value);
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator2>::value);
typedef typename std::iterator_traits<InputIterator1>::value_type input_type;
ptrdiff_t n = std::distance(first1, last1);
@@ -69,6 +77,8 @@ inline T inner_product(InputIterator1 first1,
BinaryTransformFunction transform_function,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator1>::value);
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator2>::value);
typedef typename std::iterator_traits<InputIterator1>::value_type value_type;
size_t count = detail::iterator_range_size(first1, last1);
boost/compute/algorithm/iota.hpp
+6
View File
@@ -11,10 +11,13 @@
#ifndef BOOST_COMPUTE_ALGORITHM_IOTA_HPP
#define BOOST_COMPUTE_ALGORITHM_IOTA_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/algorithm/copy.hpp>
#include <boost/compute/iterator/counting_iterator.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -26,12 +29,15 @@ namespace compute {
/// \snippet test/test_iota.cpp iota
///
/// Will fill \c vec with the values (\c 0, \c 1, \c 2, \c ...).
///
/// Space complexity: \Omega(1)
template<class BufferIterator, class T>
inline void iota(BufferIterator first,
BufferIterator last,
const T &value,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<BufferIterator>::value);
T count = static_cast<T>(detail::iterator_range_size(first, last));
copy(
boost/compute/algorithm/is_partitioned.hpp
+6
View File
@@ -11,22 +11,28 @@
#ifndef BOOST_COMPUTE_ALGORITHM_IS_PARTITIONED_HPP
#define BOOST_COMPUTE_ALGORITHM_IS_PARTITIONED_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/algorithm/find_if.hpp>
#include <boost/compute/algorithm/find_if_not.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
/// Returns \c true if the values in the range [\p first, \p last)
/// are partitioned according to \p predicate.
///
/// Space complexity: \Omega(1)
template<class InputIterator, class UnaryPredicate>
inline bool is_partitioned(InputIterator first,
InputIterator last,
UnaryPredicate predicate,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
return ::boost::compute::find_if(
::boost::compute::find_if_not(first,
last,
boost/compute/algorithm/is_permutation.hpp
+6
View File
@@ -13,12 +13,15 @@
#include <iterator>
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/container/vector.hpp>
#include <boost/compute/detail/iterator_range_size.hpp>
#include <boost/compute/algorithm/equal.hpp>
#include <boost/compute/algorithm/sort.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -36,6 +39,7 @@ namespace compute {
/// \param last2 Iterator pointing to end of second range
/// \param queue Queue on which to execute
///
/// Space complexity: \Omega(distance(\p first1, \p last1) + distance(\p first2, \p last2))
template<class InputIterator1, class InputIterator2>
inline bool is_permutation(InputIterator1 first1,
InputIterator1 last1,
@@ -43,6 +47,8 @@ inline bool is_permutation(InputIterator1 first1,
InputIterator2 last2,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator1>::value);
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator2>::value);
typedef typename std::iterator_traits<InputIterator1>::value_type value_type1;
typedef typename std::iterator_traits<InputIterator2>::value_type value_type2;
boost/compute/algorithm/is_sorted.hpp
+7
View File
@@ -11,11 +11,14 @@
#ifndef BOOST_COMPUTE_ALGORITHM_IS_SORTED_HPP
#define BOOST_COMPUTE_ALGORITHM_IS_SORTED_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/functional/bind.hpp>
#include <boost/compute/functional/operator.hpp>
#include <boost/compute/algorithm/adjacent_find.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -30,6 +33,8 @@ namespace compute {
///
/// \return \c true if the range [\p first, \p last) is sorted
///
/// Space complexity: \Omega(1)
///
/// \see sort()
template<class InputIterator, class Compare>
inline bool is_sorted(InputIterator first,
@@ -37,6 +42,7 @@ inline bool is_sorted(InputIterator first,
Compare compare,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
using ::boost::compute::placeholders::_1;
using ::boost::compute::placeholders::_2;
@@ -51,6 +57,7 @@ inline bool is_sorted(InputIterator first,
InputIterator last,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
typedef typename std::iterator_traits<InputIterator>::value_type value_type;
return ::boost::compute::is_sorted(
boost/compute/algorithm/lexicographical_compare.hpp
+13 -4
View File
@@ -8,12 +8,15 @@
// See http://boostorg.github.com/compute for more information.
//---------------------------------------------------------------------------//
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/context.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/algorithm/any_of.hpp>
#include <boost/compute/container/vector.hpp>
#include <boost/compute/utility/program_cache.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -42,10 +45,10 @@ const char lexicographical_compare_source[] =
template<class InputIterator1, class InputIterator2>
inline bool dispatch_lexicographical_compare(InputIterator1 first1,
InputIterator1 last1,
InputIterator2 first2,
InputIterator2 last2,
command_queue &queue)
InputIterator1 last1,
InputIterator2 first2,
InputIterator2 last2,
command_queue &queue)
{
const boost::compute::context &context = queue.get_context();
@@ -103,6 +106,9 @@ inline bool dispatch_lexicographical_compare(InputIterator1 first1,
/// Checks if the first range [first1, last1) is lexicographically
/// less than the second range [first2, last2).
///
/// Space complexity:
/// \Omega(max(distance(\p first1, \p last1), distance(\p first2, \p last2)))
template<class InputIterator1, class InputIterator2>
inline bool lexicographical_compare(InputIterator1 first1,
InputIterator1 last1,
@@ -110,6 +116,9 @@ inline bool lexicographical_compare(InputIterator1 first1,
InputIterator2 last2,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator1>::value);
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator2>::value);
return detail::dispatch_lexicographical_compare(first1, last1, first2, last2, queue);
}
boost/compute/algorithm/lower_bound.hpp
+6
View File
@@ -11,10 +11,13 @@
#ifndef BOOST_COMPUTE_ALGORITHM_LOWER_BOUND_HPP
#define BOOST_COMPUTE_ALGORITHM_LOWER_BOUND_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/lambda.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/algorithm/detail/binary_find.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -22,6 +25,8 @@ namespace compute {
/// Returns an iterator pointing to the first element in the sorted
/// range [\p first, \p last) that is not less than \p value.
///
/// Space complexity: \Omega(1)
///
/// \see upper_bound()
template<class InputIterator, class T>
inline InputIterator
@@ -30,6 +35,7 @@ lower_bound(InputIterator first,
const T &value,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
using ::boost::compute::_1;
InputIterator position =
boost/compute/algorithm/max_element.hpp
+8
View File
@@ -11,10 +11,13 @@
#ifndef BOOST_COMPUTE_ALGORITHM_MAX_ELEMENT_HPP
#define BOOST_COMPUTE_ALGORITHM_MAX_ELEMENT_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/functional.hpp>
#include <boost/compute/algorithm/detail/find_extrema.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -43,6 +46,9 @@ namespace compute {
/// boost::compute::max_element(data.begin(), data.end(), compare_first, queue);
/// \endcode
///
/// Space complexity on CPUs: \Omega(1)<br>
/// Space complexity on GPUs: \Omega(N)
///
/// \see min_element()
template<class InputIterator, class Compare>
inline InputIterator
@@ -51,6 +57,7 @@ max_element(InputIterator first,
Compare compare,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
return detail::find_extrema(first, last, compare, false, queue);
}
@@ -61,6 +68,7 @@ max_element(InputIterator first,
InputIterator last,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
typedef typename std::iterator_traits<InputIterator>::value_type value_type;
return ::boost::compute::max_element(
boost/compute/algorithm/merge.hpp
+11
View File
@@ -11,6 +11,8 @@
#ifndef BOOST_COMPUTE_ALGORITHM_MERGE_HPP
#define BOOST_COMPUTE_ALGORITHM_MERGE_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/algorithm/copy.hpp>
@@ -18,6 +20,7 @@
#include <boost/compute/algorithm/detail/serial_merge.hpp>
#include <boost/compute/detail/iterator_range_size.hpp>
#include <boost/compute/detail/parameter_cache.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -37,6 +40,8 @@ namespace compute {
///
/// \return \c OutputIterator to the end of the result range
///
/// Space complexity: \Omega(distance(\p first1, \p last1) + distance(\p first2, \p last2))
///
/// \see inplace_merge()
template<class InputIterator1,
class InputIterator2,
@@ -50,6 +55,9 @@ inline OutputIterator merge(InputIterator1 first1,
Compare comp,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator1>::value);
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator2>::value);
BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
typedef typename std::iterator_traits<InputIterator1>::value_type input1_type;
typedef typename std::iterator_traits<InputIterator2>::value_type input2_type;
typedef typename std::iterator_traits<OutputIterator>::value_type output_type;
@@ -94,6 +102,9 @@ inline OutputIterator merge(InputIterator1 first1,
OutputIterator result,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator1>::value);
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator2>::value);
BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
typedef typename std::iterator_traits<InputIterator1>::value_type value_type;
less<value_type> less_than;
return merge(first1, last1, first2, last2, result, less_than, queue);
boost/compute/algorithm/min_element.hpp
+8
View File
@@ -11,10 +11,13 @@
#ifndef BOOST_COMPUTE_ALGORITHM_MIN_ELEMENT_HPP
#define BOOST_COMPUTE_ALGORITHM_MIN_ELEMENT_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/functional.hpp>
#include <boost/compute/algorithm/detail/find_extrema.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -43,6 +46,9 @@ namespace compute {
/// boost::compute::min_element(data.begin(), data.end(), compare_first, queue);
/// \endcode
///
/// Space complexity on CPUs: \Omega(1)<br>
/// Space complexity on GPUs: \Omega(N)
///
/// \see max_element()
template<class InputIterator, class Compare>
inline InputIterator
@@ -51,6 +57,7 @@ min_element(InputIterator first,
Compare compare,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
return detail::find_extrema(first, last, compare, true, queue);
}
@@ -61,6 +68,7 @@ min_element(InputIterator first,
InputIterator last,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
typedef typename std::iterator_traits<InputIterator>::value_type value_type;
return ::boost::compute::min_element(
boost/compute/algorithm/minmax_element.hpp
+8
View File
@@ -13,10 +13,13 @@
#include <utility>
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/algorithm/max_element.hpp>
#include <boost/compute/algorithm/min_element.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -31,6 +34,9 @@ namespace compute {
/// argument is less than (i.e. is ordered before) the second.
/// \param queue command queue to perform the operation
///
/// Space complexity on CPUs: \Omega(1)<br>
/// Space complexity on GPUs: \Omega(N)
///
/// \see max_element(), min_element()
template<class InputIterator, class Compare>
inline std::pair<InputIterator, InputIterator>
@@ -39,6 +45,7 @@ minmax_element(InputIterator first,
Compare compare,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
if(first == last){
// empty range
return std::make_pair(first, first);
@@ -55,6 +62,7 @@ minmax_element(InputIterator first,
InputIterator last,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
if(first == last){
// empty range
return std::make_pair(first, first);
boost/compute/algorithm/mismatch.hpp
+9
View File
@@ -14,6 +14,8 @@
#include <iterator>
#include <utility>
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/functional.hpp>
#include <boost/compute/command_queue.hpp>
@@ -21,6 +23,7 @@
#include <boost/compute/iterator/transform_iterator.hpp>
#include <boost/compute/iterator/zip_iterator.hpp>
#include <boost/compute/functional/detail/unpack.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -28,6 +31,8 @@ namespace compute {
/// Returns a pair of iterators pointing to the first position where the
/// range [\p first1, \p last1) and the range starting at \p first2
/// differ.
///
/// Space complexity: \Omega(1)
template<class InputIterator1, class InputIterator2>
inline std::pair<InputIterator1, InputIterator2>
mismatch(InputIterator1 first1,
@@ -35,6 +40,8 @@ mismatch(InputIterator1 first1,
InputIterator2 first2,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator1>::value);
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator2>::value);
typedef typename std::iterator_traits<InputIterator1>::value_type value_type;
::boost::compute::equal_to<value_type> op;
@@ -73,6 +80,8 @@ mismatch(InputIterator1 first1,
InputIterator2 last2,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator1>::value);
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator2>::value);
if(std::distance(first1, last1) < std::distance(first2, last2)){
return ::boost::compute::mismatch(first1, last1, first2, queue);
}
boost/compute/algorithm/next_permutation.hpp
+5
View File
@@ -13,10 +13,13 @@
#include <iterator>
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/container/detail/scalar.hpp>
#include <boost/compute/algorithm/reverse.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -131,11 +134,13 @@ inline InputIterator np_ceiling(InputIterator first,
/// \param last Iterator pointing to end of range
/// \param queue Queue on which to execute
///
/// Space complexity: \Omega(1)
template<class InputIterator>
inline bool next_permutation(InputIterator first,
InputIterator last,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
typedef typename std::iterator_traits<InputIterator>::value_type value_type;
if(first == last) return false;
boost/compute/algorithm/none_of.hpp
+6
View File
@@ -11,8 +11,11 @@
#ifndef BOOST_COMPUTE_ALGORITHM_NONE_OF_HPP
#define BOOST_COMPUTE_ALGORITHM_NONE_OF_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/algorithm/find_if.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -20,6 +23,8 @@ namespace compute {
/// Returns \c true if \p predicate returns \c true for none of the elements in
/// the range [\p first, \p last).
///
/// Space complexity: \Omega(1)
///
/// \see all_of(), any_of()
template<class InputIterator, class UnaryPredicate>
inline bool none_of(InputIterator first,
@@ -27,6 +32,7 @@ inline bool none_of(InputIterator first,
UnaryPredicate predicate,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
return ::boost::compute::find_if(first, last, predicate, queue) == last;
}
boost/compute/algorithm/partial_sum.hpp
+9
View File
@@ -11,9 +11,12 @@
#ifndef BOOST_COMPUTE_ALGORITHM_PARTIAL_SUM_HPP
#define BOOST_COMPUTE_ALGORITHM_PARTIAL_SUM_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/algorithm/inclusive_scan.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -21,6 +24,10 @@ namespace compute {
/// Calculates the cumulative sum of the elements in the range [\p first,
/// \p last) and writes the resulting values to the range beginning at
/// \p result.
///
/// Space complexity on GPUs: \Omega(n)<br>
/// Space complexity on GPUs when \p first == \p result: \Omega(2n)<br>
/// Space complexity on CPUs: \Omega(1)
template<class InputIterator, class OutputIterator>
inline OutputIterator
partial_sum(InputIterator first,
@@ -28,6 +35,8 @@ partial_sum(InputIterator first,
OutputIterator result,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
return ::boost::compute::inclusive_scan(first, last, result, queue);
}
boost/compute/algorithm/partition.hpp
+6
View File
@@ -11,9 +11,12 @@
#ifndef BOOST_COMPUTE_ALGORITHM_PARTITION_HPP
#define BOOST_COMPUTE_ALGORITHM_PARTITION_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/algorithm/stable_partition.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -22,6 +25,8 @@ namespace compute {
/// Partitions the elements in the range [\p first, \p last) according to
/// \p predicate. Order of the elements need not be preserved.
///
/// Space complexity: \Omega(3n)
///
/// \see is_partitioned() and stable_partition()
///
template<class Iterator, class UnaryPredicate>
@@ -30,6 +35,7 @@ inline Iterator partition(Iterator first,
UnaryPredicate predicate,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<Iterator>::value);
return stable_partition(first, last, predicate, queue);
}
boost/compute/algorithm/partition_copy.hpp
+9
View File
@@ -11,10 +11,13 @@
#ifndef BOOST_COMPUTE_ALGORITHM_PARTITION_COPY_HPP
#define BOOST_COMPUTE_ALGORITHM_PARTITION_COPY_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/functional.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/algorithm/copy_if.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -24,6 +27,8 @@ namespace compute {
/// and all of the elements for which \p predicate returns \c false to
/// the range beginning at \p first_false.
///
/// Space complexity: \Omega(2n)
///
/// \see partition()
template<class InputIterator,
class OutputIterator1,
@@ -37,6 +42,10 @@ partition_copy(InputIterator first,
UnaryPredicate predicate,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator1>::value);
BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator2>::value);
// copy true values
OutputIterator1 last_true =
::boost::compute::copy_if(first,
boost/compute/algorithm/partition_point.hpp
+9 -3
View File
@@ -11,9 +11,12 @@
#ifndef BOOST_COMPUTE_ALGORITHM_PARTITION_POINT_HPP
#define BOOST_COMPUTE_ALGORITHM_PARTITION_POINT_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/algorithm/detail/binary_find.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -29,14 +32,17 @@ namespace compute {
/// \param predicate Unary predicate to be applied on each element
/// \param queue Queue on which to execute
///
/// Space complexity: \Omega(1)
///
/// \see partition() and stable_partition()
///
template<class InputIterator, class UnaryPredicate>
inline InputIterator partition_point(InputIterator first,
InputIterator last,
UnaryPredicate predicate,
command_queue &queue = system::default_queue())
InputIterator last,
UnaryPredicate predicate,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
return detail::binary_find(first, last, not1(predicate), queue);
}
boost/compute/algorithm/prev_permutation.hpp
+5
View File
@@ -13,10 +13,13 @@
#include <iterator>
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/container/detail/scalar.hpp>
#include <boost/compute/algorithm/reverse.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -131,11 +134,13 @@ inline InputIterator pp_floor(InputIterator first,
/// \param last Iterator pointing to end of range
/// \param queue Queue on which to execute
///
/// Space complexity: \Omega(1)
template<class InputIterator>
inline bool prev_permutation(InputIterator first,
InputIterator last,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
typedef typename std::iterator_traits<InputIterator>::value_type value_type;
if(first == last) return false;
boost/compute/algorithm/random_shuffle.hpp
+15
View File
@@ -14,6 +14,11 @@
#include <vector>
#include <algorithm>
#ifdef BOOST_COMPUTE_USE_CPP11
#include <random>
#endif
#include <boost/static_assert.hpp>
#include <boost/range/algorithm_ext/iota.hpp>
#include <boost/compute/system.hpp>
@@ -22,18 +27,22 @@
#include <boost/compute/container/vector.hpp>
#include <boost/compute/algorithm/scatter.hpp>
#include <boost/compute/detail/iterator_range_size.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
/// Randomly shuffles the elements in the range [\p first, \p last).
///
/// Space complexity: \Omega(2n)
///
/// \see scatter()
template<class Iterator>
inline void random_shuffle(Iterator first,
Iterator last,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<Iterator>::value);
typedef typename std::iterator_traits<Iterator>::value_type value_type;
size_t count = detail::iterator_range_size(first, last);
@@ -44,7 +53,13 @@ inline void random_shuffle(Iterator first,
// generate shuffled indices on the host
std::vector<cl_uint> random_indices(count);
boost::iota(random_indices, 0);
#ifdef BOOST_COMPUTE_USE_CPP11
std::random_device nondeterministic_randomness;
std::default_random_engine random_engine(nondeterministic_randomness());
std::shuffle(random_indices.begin(), random_indices.end(), random_engine);
#else
std::random_shuffle(random_indices.begin(), random_indices.end());
#endif
// copy random indices to the device
const context &context = queue.get_context();
boost/compute/algorithm/reduce.hpp
+9
View File
@@ -13,6 +13,8 @@
#include <iterator>
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/functional.hpp>
#include <boost/compute/detail/meta_kernel.hpp>
@@ -26,6 +28,7 @@
#include <boost/compute/detail/iterator_range_size.hpp>
#include <boost/compute/memory/local_buffer.hpp>
#include <boost/compute/type_traits/result_of.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -153,6 +156,7 @@ block_reduce(InputIterator first,
return result_vector;
}
// Space complexity: O( ceil(n / 2 / 256) )
template<class InputIterator, class OutputIterator, class BinaryFunction>
inline void generic_reduce(InputIterator first,
InputIterator last,
@@ -264,6 +268,9 @@ inline void dispatch_reduce(InputIterator first,
/// efficient on parallel hardware. For more information, see the documentation
/// on the \c accumulate() algorithm.
///
/// Space complexity on GPUs: \Omega(n)<br>
/// Space complexity on CPUs: \Omega(1)
///
/// \see accumulate()
template<class InputIterator, class OutputIterator, class BinaryFunction>
inline void reduce(InputIterator first,
@@ -272,6 +279,7 @@ inline void reduce(InputIterator first,
BinaryFunction function,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
if(first == last){
return;
}
@@ -286,6 +294,7 @@ inline void reduce(InputIterator first,
OutputIterator result,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
typedef typename std::iterator_traits<InputIterator>::value_type T;
if(first == last){
boost/compute/algorithm/reduce_by_key.hpp
+19
View File
@@ -14,11 +14,14 @@
#include <iterator>
#include <utility>
#include <boost/static_assert.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/device.hpp>
#include <boost/compute/functional.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/algorithm/detail/reduce_by_key.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -51,6 +54,9 @@ namespace compute {
///
/// \snippet test/test_reduce_by_key.cpp reduce_by_key_int
///
/// Space complexity on GPUs: \Omega(2n)<br>
/// Space complexity on CPUs: \Omega(1)
///
/// \see reduce()
template<class InputKeyIterator, class InputValueIterator,
class OutputKeyIterator, class OutputValueIterator,
@@ -65,6 +71,11 @@ reduce_by_key(InputKeyIterator keys_first,
BinaryPredicate predicate,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputKeyIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<InputValueIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<OutputKeyIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<OutputValueIterator>::value);
return detail::dispatch_reduce_by_key(keys_first, keys_last, values_first,
keys_result, values_result,
function, predicate,
@@ -84,6 +95,10 @@ reduce_by_key(InputKeyIterator keys_first,
BinaryFunction function,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputKeyIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<InputValueIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<OutputKeyIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<OutputValueIterator>::value);
typedef typename std::iterator_traits<InputKeyIterator>::value_type key_type;
return reduce_by_key(keys_first, keys_last, values_first,
@@ -103,6 +118,10 @@ reduce_by_key(InputKeyIterator keys_first,
OutputValueIterator values_result,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputKeyIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<InputValueIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<OutputKeyIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<OutputValueIterator>::value);
typedef typename std::iterator_traits<InputKeyIterator>::value_type key_type;
typedef typename std::iterator_traits<InputValueIterator>::value_type value_type;
boost/compute/algorithm/remove.hpp
+6
View File
@@ -11,10 +11,13 @@
#ifndef BOOST_COMPUTE_ALGORITHM_REMOVE_HPP
#define BOOST_COMPUTE_ALGORITHM_REMOVE_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/lambda.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/algorithm/remove_if.hpp>
#include <boost/compute/type_traits/vector_size.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -22,6 +25,8 @@ namespace compute {
/// Removes each element equal to \p value in the range [\p first,
/// \p last).
///
/// Space complexity: \Omega(3n)
///
/// \see remove_if()
template<class Iterator, class T>
inline Iterator remove(Iterator first,
@@ -29,6 +34,7 @@ inline Iterator remove(Iterator first,
const T &value,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<Iterator>::value);
typedef typename std::iterator_traits<Iterator>::value_type value_type;
using ::boost::compute::_1;
boost/compute/algorithm/remove_if.hpp
+6
View File
@@ -11,10 +11,13 @@
#ifndef BOOST_COMPUTE_ALGORITHM_REMOVE_IF_HPP
#define BOOST_COMPUTE_ALGORITHM_REMOVE_IF_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/algorithm/copy_if.hpp>
#include <boost/compute/container/vector.hpp>
#include <boost/compute/functional/logical.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -22,6 +25,8 @@ namespace compute {
/// Removes each element for which \p predicate returns \c true in the
/// range [\p first, \p last).
///
/// Space complexity: \Omega(3n)
///
/// \see remove()
template<class Iterator, class Predicate>
inline Iterator remove_if(Iterator first,
@@ -29,6 +34,7 @@ inline Iterator remove_if(Iterator first,
Predicate predicate,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<Iterator>::value);
typedef typename std::iterator_traits<Iterator>::value_type value_type;
// temporary storage for the input data
boost/compute/algorithm/replace.hpp
+6
View File
@@ -11,10 +11,13 @@
#ifndef BOOST_COMPUTE_ALGORITHM_REPLACE_HPP
#define BOOST_COMPUTE_ALGORITHM_REPLACE_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/detail/meta_kernel.hpp>
#include <boost/compute/detail/iterator_range_size.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -68,6 +71,8 @@ private:
/// Replaces each instance of \p old_value in the range [\p first,
/// \p last) with \p new_value.
///
/// Space complexity: \Omega(1)
template<class Iterator, class T>
inline void replace(Iterator first,
Iterator last,
@@ -75,6 +80,7 @@ inline void replace(Iterator first,
const T &new_value,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<Iterator>::value);
detail::replace_kernel<Iterator, T> kernel;
kernel.set_range(first, last);
boost/compute/algorithm/replace_copy.hpp
+8
View File
@@ -13,10 +13,13 @@
#include <iterator>
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/algorithm/copy.hpp>
#include <boost/compute/algorithm/replace.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -25,6 +28,8 @@ namespace compute {
/// beginning at \p result while replacing each instance of \p old_value
/// with \p new_value.
///
/// Space complexity: \Omega(1)
///
/// \see replace()
template<class InputIterator, class OutputIterator, class T>
inline OutputIterator
@@ -35,6 +40,9 @@ replace_copy(InputIterator first,
const T &new_value,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
typedef typename std::iterator_traits<OutputIterator>::difference_type difference_type;
difference_type count = std::distance(first, last);
boost/compute/algorithm/reverse.hpp
+6
View File
@@ -11,10 +11,13 @@
#ifndef BOOST_COMPUTE_ALGORITHM_REVERSE_HPP
#define BOOST_COMPUTE_ALGORITHM_REVERSE_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/detail/meta_kernel.hpp>
#include <boost/compute/detail/iterator_range_size.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -52,12 +55,15 @@ struct reverse_kernel : public meta_kernel
/// Reverses the elements in the range [\p first, \p last).
///
/// Space complexity: \Omega(1)
///
/// \see reverse_copy()
template<class Iterator>
inline void reverse(Iterator first,
Iterator last,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<Iterator>::value);
size_t count = detail::iterator_range_size(first, last);
if(count < 2){
return;
boost/compute/algorithm/reverse_copy.hpp
+8
View File
@@ -13,10 +13,13 @@
#include <iterator>
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/algorithm/copy.hpp>
#include <boost/compute/algorithm/reverse.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -51,6 +54,8 @@ struct reverse_copy_kernel : public meta_kernel
/// Copies the elements in the range [\p first, \p last) in reversed
/// order to the range beginning at \p result.
///
/// Space complexity: \Omega(1)
///
/// \see reverse()
template<class InputIterator, class OutputIterator>
inline OutputIterator
@@ -59,6 +64,9 @@ reverse_copy(InputIterator first,
OutputIterator result,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
typedef typename std::iterator_traits<OutputIterator>::difference_type difference_type;
difference_type count = std::distance(first, last);
boost/compute/algorithm/rotate.hpp
+2
View File
@@ -21,6 +21,8 @@ namespace compute {
/// Performs left rotation such that element at \p n_first comes to the
/// beginning.
///
/// Space complexity: \Omega(distance(\p first, \p last))
///
/// \see rotate_copy()
template<class InputIterator>
inline void rotate(InputIterator first,
boost/compute/algorithm/rotate_copy.hpp
+2
View File
@@ -20,6 +20,8 @@ namespace compute {
/// Performs left rotation such that element at n_first comes to the
/// beginning and the output is stored in range starting at result.
///
/// Space complexity: \Omega(1)
///
/// \see rotate()
template<class InputIterator, class OutputIterator>
inline void rotate_copy(InputIterator first,
boost/compute/algorithm/scatter.hpp
+8
View File
@@ -11,6 +11,7 @@
#ifndef BOOST_COMPUTE_ALGORITHM_SCATTER_HPP
#define BOOST_COMPUTE_ALGORITHM_SCATTER_HPP
#include <boost/static_assert.hpp>
#include <boost/algorithm/string/replace.hpp>
#include <boost/compute/system.hpp>
@@ -20,6 +21,7 @@
#include <boost/compute/type_traits/type_name.hpp>
#include <boost/compute/detail/iterator_range_size.hpp>
#include <boost/compute/detail/meta_kernel.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -79,6 +81,8 @@ private:
/// beginning at \p result using the output indices from the range beginning
/// at \p map.
///
/// Space complexity: \Omega(1)
///
/// \see gather()
template<class InputIterator, class MapIterator, class OutputIterator>
inline void scatter(InputIterator first,
@@ -87,6 +91,10 @@ inline void scatter(InputIterator first,
OutputIterator result,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<MapIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
detail::scatter_kernel<InputIterator, MapIterator, OutputIterator> kernel;
kernel.set_range(first, last, map, result);
boost/compute/algorithm/search.hpp
+7
View File
@@ -11,12 +11,15 @@
#ifndef BOOST_COMPUTE_ALGORITHM_SEARCH_HPP
#define BOOST_COMPUTE_ALGORITHM_SEARCH_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/algorithm/detail/search_all.hpp>
#include <boost/compute/algorithm/find.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/system.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -34,6 +37,7 @@ namespace compute {
/// \param p_last Iterator pointing to end of pattern
/// \param queue Queue on which to execute
///
/// Space complexity: \Omega(distance(\p t_first, \p t_last))
template<class TextIterator, class PatternIterator>
inline TextIterator search(TextIterator t_first,
TextIterator t_last,
@@ -41,6 +45,9 @@ inline TextIterator search(TextIterator t_first,
PatternIterator p_last,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<TextIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<PatternIterator>::value);
// there is no need to check if pattern starts at last n - 1 indices
vector<uint_> matching_indices(
detail::iterator_range_size(t_first, t_last)
boost/compute/algorithm/search_n.hpp
+6
View File
@@ -13,11 +13,14 @@
#include <iterator>
#include <boost/static_assert.hpp>
#include <boost/compute/algorithm/find.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/system.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -102,6 +105,7 @@ private:
/// \param value Value which repeats
/// \param queue Queue on which to execute
///
/// Space complexity: \Omega(distance(\p t_first, \p t_last))
template<class TextIterator, class ValueType>
inline TextIterator search_n(TextIterator t_first,
TextIterator t_last,
@@ -109,6 +113,8 @@ inline TextIterator search_n(TextIterator t_first,
ValueType value,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<TextIterator>::value);
// there is no need to check if pattern starts at last n - 1 indices
vector<uint_> matching_indices(
detail::iterator_range_size(t_first, t_last) + 1 - n,
boost/compute/algorithm/set_difference.hpp
+9
View File
@@ -13,6 +13,8 @@
#include <iterator>
#include <boost/static_assert.hpp>
#include <boost/compute/algorithm/detail/compact.hpp>
#include <boost/compute/algorithm/detail/balanced_path.hpp>
#include <boost/compute/algorithm/exclusive_scan.hpp>
@@ -21,6 +23,7 @@
#include <boost/compute/detail/iterator_range_size.hpp>
#include <boost/compute/detail/meta_kernel.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -122,6 +125,8 @@ private:
/// will be stored
/// \param queue Queue on which to execute
///
/// Space complexity:
/// \Omega(2(distance(\p first1, \p last1) + distance(\p first2, \p last2)))
template<class InputIterator1, class InputIterator2, class OutputIterator>
inline OutputIterator set_difference(InputIterator1 first1,
InputIterator1 last1,
@@ -130,6 +135,10 @@ inline OutputIterator set_difference(InputIterator1 first1,
OutputIterator result,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator1>::value);
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator2>::value);
BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
typedef typename std::iterator_traits<InputIterator1>::value_type value_type;
int tile_size = 1024;
boost/compute/algorithm/set_intersection.hpp
+9
View File
@@ -13,6 +13,8 @@
#include <iterator>
#include <boost/static_assert.hpp>
#include <boost/compute/algorithm/detail/compact.hpp>
#include <boost/compute/algorithm/detail/balanced_path.hpp>
#include <boost/compute/algorithm/exclusive_scan.hpp>
@@ -21,6 +23,7 @@
#include <boost/compute/detail/iterator_range_size.hpp>
#include <boost/compute/detail/meta_kernel.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -110,6 +113,8 @@ private:
/// will be stored
/// \param queue Queue on which to execute
///
/// Space complexity:
/// \Omega(2(distance(\p first1, \p last1) + distance(\p first2, \p last2)))
template<class InputIterator1, class InputIterator2, class OutputIterator>
inline OutputIterator set_intersection(InputIterator1 first1,
InputIterator1 last1,
@@ -118,6 +123,10 @@ inline OutputIterator set_intersection(InputIterator1 first1,
OutputIterator result,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator1>::value);
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator2>::value);
BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
typedef typename std::iterator_traits<InputIterator1>::value_type value_type;
int tile_size = 1024;
boost/compute/algorithm/set_symmetric_difference.hpp
+16 -6
View File
@@ -13,6 +13,8 @@
#include <iterator>
#include <boost/static_assert.hpp>
#include <boost/compute/algorithm/detail/compact.hpp>
#include <boost/compute/algorithm/detail/balanced_path.hpp>
#include <boost/compute/algorithm/exclusive_scan.hpp>
@@ -21,6 +23,7 @@
#include <boost/compute/detail/iterator_range_size.hpp>
#include <boost/compute/detail/meta_kernel.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -133,14 +136,21 @@ private:
/// difference will be stored
/// \param queue Queue on which to execute
///
/// Space complexity:
/// \Omega(2(distance(\p first1, \p last1) + distance(\p first2, \p last2)))
template<class InputIterator1, class InputIterator2, class OutputIterator>
inline OutputIterator set_symmetric_difference(InputIterator1 first1,
InputIterator1 last1,
InputIterator2 first2,
InputIterator2 last2,
OutputIterator result,
command_queue &queue = system::default_queue())
inline OutputIterator
set_symmetric_difference(InputIterator1 first1,
InputIterator1 last1,
InputIterator2 first2,
InputIterator2 last2,
OutputIterator result,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator1>::value);
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator2>::value);
BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
typedef typename std::iterator_traits<InputIterator1>::value_type value_type;
int tile_size = 1024;
boost/compute/algorithm/set_union.hpp
+9
View File
@@ -13,6 +13,8 @@
#include <iterator>
#include <boost/static_assert.hpp>
#include <boost/compute/algorithm/detail/balanced_path.hpp>
#include <boost/compute/algorithm/detail/compact.hpp>
#include <boost/compute/algorithm/exclusive_scan.hpp>
@@ -21,6 +23,7 @@
#include <boost/compute/detail/iterator_range_size.hpp>
#include <boost/compute/detail/meta_kernel.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -135,6 +138,8 @@ private:
/// will be stored
/// \param queue Queue on which to execute
///
/// Space complexity:
/// \Omega(2(distance(\p first1, \p last1) + distance(\p first2, \p last2)))
template<class InputIterator1, class InputIterator2, class OutputIterator>
inline OutputIterator set_union(InputIterator1 first1,
InputIterator1 last1,
@@ -143,6 +148,10 @@ inline OutputIterator set_union(InputIterator1 first1,
OutputIterator result,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator1>::value);
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator2>::value);
BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
typedef typename std::iterator_traits<InputIterator1>::value_type value_type;
int tile_size = 1024;
boost/compute/algorithm/sort.hpp
+2
View File
@@ -176,6 +176,8 @@ inline void dispatch_sort(Iterator first,
/// boost::compute::sort(data.begin(), data.end(), queue);
/// \endcode
///
/// Space complexity: \Omega(n)
///
/// \see is_sorted()
template<class Iterator, class Compare>
inline void sort(Iterator first,
boost/compute/algorithm/sort_by_key.hpp
+8 -1
View File
@@ -13,6 +13,7 @@
#include <iterator>
#include <boost/static_assert.hpp>
#include <boost/utility/enable_if.hpp>
#include <boost/compute/system.hpp>
@@ -23,10 +24,10 @@
#include <boost/compute/algorithm/detail/radix_sort.hpp>
#include <boost/compute/algorithm/reverse.hpp>
#include <boost/compute/detail/iterator_range_size.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
namespace detail {
template<class KeyIterator, class ValueIterator>
@@ -128,6 +129,8 @@ inline void dispatch_sort_by_key(KeyIterator keys_first,
///
/// If no compare function is specified, \c less is used.
///
/// Space complexity: \Omega(2n)
///
/// \see sort()
template<class KeyIterator, class ValueIterator, class Compare>
inline void sort_by_key(KeyIterator keys_first,
@@ -136,6 +139,8 @@ inline void sort_by_key(KeyIterator keys_first,
Compare compare,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<KeyIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<ValueIterator>::value);
::boost::compute::detail::dispatch_sort_by_key(
keys_first, keys_last, values_first, compare, queue
);
@@ -148,6 +153,8 @@ inline void sort_by_key(KeyIterator keys_first,
ValueIterator values_first,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<KeyIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<ValueIterator>::value);
typedef typename std::iterator_traits<KeyIterator>::value_type key_type;
::boost::compute::sort_by_key(
boost/compute/algorithm/stable_partition.hpp
+6
View File
@@ -11,12 +11,15 @@
#ifndef BOOST_COMPUTE_ALGORITHM_STABLE_PARTITION_HPP
#define BOOST_COMPUTE_ALGORITHM_STABLE_PARTITION_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/context.hpp>
#include <boost/compute/functional.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/algorithm/copy_if.hpp>
#include <boost/compute/container/vector.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -33,6 +36,8 @@ namespace compute {
/// \param predicate Unary predicate to be applied on each element
/// \param queue Queue on which to execute
///
/// Space complexity: \Omega(3n)
///
/// \see is_partitioned() and partition()
///
template<class Iterator, class UnaryPredicate>
@@ -41,6 +46,7 @@ inline Iterator stable_partition(Iterator first,
UnaryPredicate predicate,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<Iterator>::value);
typedef typename std::iterator_traits<Iterator>::value_type value_type;
// make temporary copy of the input
boost/compute/algorithm/stable_sort.hpp
+8
View File
@@ -13,6 +13,8 @@
#include <iterator>
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/algorithm/detail/merge_sort_on_cpu.hpp>
@@ -22,6 +24,7 @@
#include <boost/compute/algorithm/reverse.hpp>
#include <boost/compute/functional/operator.hpp>
#include <boost/compute/detail/iterator_range_size.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -72,6 +75,8 @@ dispatch_gpu_stable_sort(buffer_iterator<T> first,
/// Sorts the values in the range [\p first, \p last) according to
/// \p compare. The relative order of identical values is preserved.
///
/// Space complexity: \Omega(n)
///
/// \see sort(), is_sorted()
template<class Iterator, class Compare>
inline void stable_sort(Iterator first,
@@ -79,6 +84,8 @@ inline void stable_sort(Iterator first,
Compare compare,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<Iterator>::value);
if(queue.get_device().type() & device::gpu) {
::boost::compute::detail::dispatch_gpu_stable_sort(
first, last, compare, queue
@@ -94,6 +101,7 @@ inline void stable_sort(Iterator first,
Iterator last,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<Iterator>::value);
typedef typename std::iterator_traits<Iterator>::value_type value_type;
::boost::compute::less<value_type> less;
boost/compute/algorithm/stable_sort_by_key.hpp
+9
View File
@@ -13,10 +13,13 @@
#include <iterator>
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/algorithm/sort_by_key.hpp>
#include <boost/compute/detail/iterator_range_size.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -126,6 +129,8 @@ inline void dispatch_ssort_by_key(KeyIterator keys_first,
///
/// If no compare function is specified, \c less is used.
///
/// Space complexity: \Omega(2n)
///
/// \see sort()
template<class KeyIterator, class ValueIterator, class Compare>
inline void stable_sort_by_key(KeyIterator keys_first,
@@ -134,6 +139,8 @@ inline void stable_sort_by_key(KeyIterator keys_first,
Compare compare,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<KeyIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<ValueIterator>::value);
::boost::compute::detail::dispatch_ssort_by_key(
keys_first, keys_last, values_first, compare, queue
);
@@ -146,6 +153,8 @@ inline void stable_sort_by_key(KeyIterator keys_first,
ValueIterator values_first,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<KeyIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<ValueIterator>::value);
typedef typename std::iterator_traits<KeyIterator>::value_type key_type;
::boost::compute::stable_sort_by_key(
boost/compute/algorithm/swap_ranges.hpp
+8
View File
@@ -11,22 +11,30 @@
#ifndef BOOST_COMPUTE_ALGORITHM_SWAP_RANGES_HPP
#define BOOST_COMPUTE_ALGORITHM_SWAP_RANGES_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/algorithm/copy.hpp>
#include <boost/compute/container/vector.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
/// Swaps the elements in the range [\p first1, \p last1) with the
/// elements in the range beginning at \p first2.
///
/// Space complexity: \Omega(distance(\p first1, \p last1))
template<class Iterator1, class Iterator2>
inline Iterator2 swap_ranges(Iterator1 first1,
Iterator1 last1,
Iterator2 first2,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<Iterator1>::value);
BOOST_STATIC_ASSERT(is_device_iterator<Iterator2>::value);
typedef typename std::iterator_traits<Iterator1>::value_type value_type;
Iterator2 last2 = first2 + std::distance(first1, last1);
boost/compute/algorithm/transform.hpp
+13 -2
View File
@@ -11,12 +11,15 @@
#ifndef BOOST_COMPUTE_ALGORITHM_TRANSFORM_HPP
#define BOOST_COMPUTE_ALGORITHM_TRANSFORM_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/algorithm/copy.hpp>
#include <boost/compute/iterator/transform_iterator.hpp>
#include <boost/compute/iterator/zip_iterator.hpp>
#include <boost/compute/functional/detail/unpack.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -29,6 +32,8 @@ namespace compute {
///
/// \snippet test/test_transform.cpp transform_abs
///
/// Space complexity: \Omega(1)
///
/// \see copy()
template<class InputIterator, class OutputIterator, class UnaryOperator>
inline OutputIterator transform(InputIterator first,
@@ -37,6 +42,8 @@ inline OutputIterator transform(InputIterator first,
UnaryOperator op,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
return copy(
::boost::compute::make_transform_iterator(first, op),
::boost::compute::make_transform_iterator(last, op),
@@ -57,13 +64,17 @@ inline OutputIterator transform(InputIterator1 first1,
BinaryOperator op,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator1>::value);
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator2>::value);
BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
typedef typename std::iterator_traits<InputIterator1>::difference_type difference_type;
difference_type n = std::distance(first1, last1);
return transform(
make_zip_iterator(boost::make_tuple(first1, first2)),
make_zip_iterator(boost::make_tuple(last1, first2 + n)),
::boost::compute::make_zip_iterator(boost::make_tuple(first1, first2)),
::boost::compute::make_zip_iterator(boost::make_tuple(last1, first2 + n)),
result,
detail::unpack(op),
queue
boost/compute/algorithm/transform_if.hpp
+10 -4
View File
@@ -11,6 +11,8 @@
#ifndef BOOST_COMPUTE_ALGORITHM_TRANSFORM_IF_HPP
#define BOOST_COMPUTE_ALGORITHM_TRANSFORM_IF_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/cl.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
@@ -21,11 +23,13 @@
#include <boost/compute/detail/meta_kernel.hpp>
#include <boost/compute/detail/iterator_range_size.hpp>
#include <boost/compute/iterator/discard_iterator.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
namespace detail {
// Space complexity: O(2n)
template<class InputIterator, class OutputIterator, class UnaryFunction, class Predicate>
inline OutputIterator transform_if_impl(InputIterator first,
InputIterator last,
@@ -53,14 +57,12 @@ inline OutputIterator transform_if_impl(InputIterator first,
<< predicate(first[k1.get_global_id(0)]) << " ? 1 : 0;\n";
k1.exec_1d(queue, 0, count);
// count number of elements to be copied
size_t copied_element_count =
::boost::compute::count(indices.begin(), indices.end(), 1, queue);
// scan indices
size_t copied_element_count = (indices.cend() - 1).read(queue);
::boost::compute::exclusive_scan(
indices.begin(), indices.end(), indices.begin(), queue
);
copied_element_count += (indices.cend() - 1).read(queue); // last scan element plus last mask element
// copy values
::boost::compute::detail::meta_kernel k2("transform_if_do_copy");
@@ -98,6 +100,8 @@ inline discard_iterator transform_if_impl(InputIterator first,
/// Copies each element in the range [\p first, \p last) for which
/// \p predicate returns \c true to the range beginning at \p result.
///
/// Space complexity: O(2n)
template<class InputIterator, class OutputIterator, class UnaryFunction, class Predicate>
inline OutputIterator transform_if(InputIterator first,
InputIterator last,
@@ -106,6 +110,8 @@ inline OutputIterator transform_if(InputIterator first,
Predicate predicate,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
return detail::transform_if_impl(
first, last, result, function, predicate, false, queue
);
boost/compute/algorithm/transform_reduce.hpp
+11
View File
@@ -11,12 +11,15 @@
#ifndef BOOST_COMPUTE_ALGORITHM_TRANSFORM_REDUCE_HPP
#define BOOST_COMPUTE_ALGORITHM_TRANSFORM_REDUCE_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/algorithm/reduce.hpp>
#include <boost/compute/iterator/transform_iterator.hpp>
#include <boost/compute/iterator/zip_iterator.hpp>
#include <boost/compute/functional/detail/unpack.hpp>
#include <boost/compute/detail/iterator_range_size.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -30,6 +33,9 @@ namespace compute {
///
/// \snippet test/test_transform_reduce.cpp sum_abs_int
///
/// Space complexity on GPUs: \Omega(n)<br>
/// Space complexity on CPUs: \Omega(1)
///
/// \see reduce(), inner_product()
template<class InputIterator,
class OutputIterator,
@@ -42,6 +48,7 @@ inline void transform_reduce(InputIterator first,
BinaryReduceFunction reduce_function,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
::boost::compute::reduce(
::boost::compute::make_transform_iterator(first, transform_function),
::boost::compute::make_transform_iterator(last, transform_function),
@@ -65,6 +72,10 @@ inline void transform_reduce(InputIterator1 first1,
BinaryReduceFunction reduce_function,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator1>::value);
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator2>::value);
BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
typedef typename std::iterator_traits<InputIterator1>::difference_type difference_type;
difference_type n = std::distance(first1, last1);
boost/compute/algorithm/unique.hpp
+7
View File
@@ -11,11 +11,14 @@
#ifndef BOOST_COMPUTE_ALGORITHM_UNIQUE_HPP
#define BOOST_COMPUTE_ALGORITHM_UNIQUE_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/algorithm/unique_copy.hpp>
#include <boost/compute/container/vector.hpp>
#include <boost/compute/functional/operator.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -31,6 +34,8 @@ namespace compute {
///
/// \return \c InputIterator to the new logical end of the range
///
/// Space complexity: \Omega(4n)
///
/// \see unique_copy()
template<class InputIterator, class BinaryPredicate>
inline InputIterator unique(InputIterator first,
@@ -38,6 +43,7 @@ inline InputIterator unique(InputIterator first,
BinaryPredicate op,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
typedef typename std::iterator_traits<InputIterator>::value_type value_type;
vector<value_type> temp(first, last, queue);
@@ -53,6 +59,7 @@ inline InputIterator unique(InputIterator first,
InputIterator last,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
typedef typename std::iterator_traits<InputIterator>::value_type value_type;
return ::boost::compute::unique(
boost/compute/algorithm/unique_copy.hpp
+11
View File
@@ -11,6 +11,8 @@
#ifndef BOOST_COMPUTE_ALGORITHM_UNIQUE_COPY_HPP
#define BOOST_COMPUTE_ALGORITHM_UNIQUE_COPY_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/lambda.hpp>
#include <boost/compute/system.hpp>
@@ -21,6 +23,7 @@
#include <boost/compute/detail/iterator_range_size.hpp>
#include <boost/compute/detail/meta_kernel.hpp>
#include <boost/compute/functional/operator.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -127,6 +130,8 @@ inline OutputIterator unique_copy(InputIterator first,
///
/// \return \c OutputIterator to the end of the result range
///
/// Space complexity: \Omega(4n)
///
/// \see unique()
template<class InputIterator, class OutputIterator, class BinaryPredicate>
inline OutputIterator unique_copy(InputIterator first,
@@ -135,6 +140,9 @@ inline OutputIterator unique_copy(InputIterator first,
BinaryPredicate op,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
size_t count = detail::iterator_range_size(first, last);
if(count < 32){
return detail::serial_unique_copy(first, last, result, op, queue);
@@ -151,6 +159,9 @@ inline OutputIterator unique_copy(InputIterator first,
OutputIterator result,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
typedef typename std::iterator_traits<InputIterator>::value_type value_type;
return ::boost::compute::unique_copy(
boost/compute/algorithm/upper_bound.hpp
+6
View File
@@ -11,10 +11,13 @@
#ifndef BOOST_COMPUTE_ALGORITHM_UPPER_BOUND_HPP
#define BOOST_COMPUTE_ALGORITHM_UPPER_BOUND_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/lambda.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/algorithm/detail/binary_find.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
@@ -22,6 +25,8 @@ namespace compute {
/// Returns an iterator pointing to the first element in the sorted
/// range [\p first, \p last) that is not less than or equal to
/// \p value.
///
/// Space complexity: \Omega(1)
template<class InputIterator, class T>
inline InputIterator
upper_bound(InputIterator first,
@@ -29,6 +34,7 @@ upper_bound(InputIterator first,
const T &value,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
using ::boost::compute::_1;
InputIterator position =