Imported Upstream version 1.65.0upstream/1.65.0

Change-Id: Icf8400b375482cb11bcf77440a6934ba360d6ba4
Signed-off-by: DongHun Kwak <dh0128.kwak@samsung.com>

89 files changed, 217 insertions, 36 deletions

diff --git a/boost/compute/algorithm/accumulate.hpp b/boost/compute/algorithm/accumulate.hpp
index 328420a07c..be20bee60e 100644
--- a/boost/compute/algorithm/accumulate.hpp
+++ b/boost/compute/algorithm/accumulate.hpp
@@ -26,6 +26,7 @@ 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 +156,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,
diff --git a/boost/compute/algorithm/adjacent_difference.hpp b/boost/compute/algorithm/adjacent_difference.hpp
index ef13970754..c3b0e7d191 100644
--- a/boost/compute/algorithm/adjacent_difference.hpp
+++ b/boost/compute/algorithm/adjacent_difference.hpp
@@ -64,6 +64,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
diff --git a/boost/compute/algorithm/adjacent_find.hpp b/boost/compute/algorithm/adjacent_find.hpp
index 992a01eddc..a71a817f57 100644
--- a/boost/compute/algorithm/adjacent_find.hpp
+++ b/boost/compute/algorithm/adjacent_find.hpp
@@ -114,6 +114,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
diff --git a/boost/compute/algorithm/all_of.hpp b/boost/compute/algorithm/all_of.hpp
index 34d7518f32..56c5809992 100644
--- a/boost/compute/algorithm/all_of.hpp
+++ b/boost/compute/algorithm/all_of.hpp
@@ -20,6 +20,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,
diff --git a/boost/compute/algorithm/any_of.hpp b/boost/compute/algorithm/any_of.hpp
index b07779597c..54031fbac5 100644
--- a/boost/compute/algorithm/any_of.hpp
+++ b/boost/compute/algorithm/any_of.hpp
@@ -24,6 +24,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,
diff --git a/boost/compute/algorithm/binary_search.hpp b/boost/compute/algorithm/binary_search.hpp
index 6e19498790..880f3628fb 100644
--- a/boost/compute/algorithm/binary_search.hpp
+++ b/boost/compute/algorithm/binary_search.hpp
@@ -20,6 +20,8 @@ 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,
diff --git a/boost/compute/algorithm/copy.hpp b/boost/compute/algorithm/copy.hpp
index 7779277b82..4866726f6e 100644
--- a/boost/compute/algorithm/copy.hpp
+++ b/boost/compute/algorithm/copy.hpp
@@ -826,6 +826,8 @@ 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,
diff --git a/boost/compute/algorithm/copy_if.hpp b/boost/compute/algorithm/copy_if.hpp
index 3cd08ef293..bdedcb8536 100644
--- a/boost/compute/algorithm/copy_if.hpp
+++ b/boost/compute/algorithm/copy_if.hpp
@@ -38,6 +38,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,
diff --git a/boost/compute/algorithm/copy_n.hpp b/boost/compute/algorithm/copy_n.hpp
index f0989edc67..5280193497 100644
--- a/boost/compute/algorithm/copy_n.hpp
+++ b/boost/compute/algorithm/copy_n.hpp
@@ -30,6 +30,8 @@ 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,
diff --git a/boost/compute/algorithm/count.hpp b/boost/compute/algorithm/count.hpp
index 140d67379f..7a502c6791 100644
--- a/boost/compute/algorithm/count.hpp
+++ b/boost/compute/algorithm/count.hpp
@@ -23,6 +23,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,
diff --git a/boost/compute/algorithm/count_if.hpp b/boost/compute/algorithm/count_if.hpp
index c9381ce5d4..81996dc828 100644
--- a/boost/compute/algorithm/count_if.hpp
+++ b/boost/compute/algorithm/count_if.hpp
@@ -25,6 +25,9 @@ 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,
diff --git a/boost/compute/algorithm/detail/copy_on_device.hpp b/boost/compute/algorithm/detail/copy_on_device.hpp
index 8738c8c0b4..034b3bc212 100644
--- a/boost/compute/algorithm/detail/copy_on_device.hpp
+++ b/boost/compute/algorithm/detail/copy_on_device.hpp
@@ -144,7 +144,7 @@ inline future<OutputIterator> copy_on_device_async(InputIterator first,
     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,
@@ -181,7 +181,7 @@ inline future<svm_ptr<T> > copy_on_device_async(svm_ptr<T> first,
 
     return make_future(result + count, event_);
 }
-#endif // CL_VERSION_2_0
+#endif // BOOST_COMPUTE_CL_VERSION_2_0
 
 } // end detail namespace
 } // end compute namespace
diff --git a/boost/compute/algorithm/detail/copy_to_device.hpp b/boost/compute/algorithm/detail/copy_to_device.hpp
index bce5975f53..8601bb20ec 100644
--- a/boost/compute/algorithm/detail/copy_to_device.hpp
+++ b/boost/compute/algorithm/detail/copy_to_device.hpp
@@ -124,7 +124,7 @@ inline future<DeviceIterator> copy_to_device_async(HostIterator first,
     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,
@@ -184,7 +184,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
diff --git a/boost/compute/algorithm/detail/copy_to_host.hpp b/boost/compute/algorithm/detail/copy_to_host.hpp
index d770a996ef..89b57174fa 100644
--- a/boost/compute/algorithm/detail/copy_to_host.hpp
+++ b/boost/compute/algorithm/detail/copy_to_host.hpp
@@ -125,7 +125,7 @@ inline future<HostIterator> copy_to_host_async(DeviceIterator first,
     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,
@@ -189,7 +189,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
diff --git a/boost/compute/algorithm/detail/find_extrema.hpp b/boost/compute/algorithm/detail/find_extrema.hpp
index eef2e36c3c..734b75aa90 100644
--- a/boost/compute/algorithm/detail/find_extrema.hpp
+++ b/boost/compute/algorithm/detail/find_extrema.hpp
@@ -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
 
diff --git a/boost/compute/algorithm/detail/find_extrema_with_reduce.hpp b/boost/compute/algorithm/detail/find_extrema_with_reduce.hpp
index 8f2a83c38b..515d7cc6da 100644
--- a/boost/compute/algorithm/detail/find_extrema_with_reduce.hpp
+++ b/boost/compute/algorithm/detail/find_extrema_with_reduce.hpp
@@ -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,
diff --git a/boost/compute/algorithm/detail/find_if_with_atomics.hpp b/boost/compute/algorithm/detail/find_if_with_atomics.hpp
index 112c34cf00..e14fd12ae0 100644
--- a/boost/compute/algorithm/detail/find_if_with_atomics.hpp
+++ b/boost/compute/algorithm/detail/find_if_with_atomics.hpp
@@ -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,
diff --git a/boost/compute/algorithm/detail/merge_sort_on_gpu.hpp b/boost/compute/algorithm/detail/merge_sort_on_gpu.hpp
index e62c6beb8d..d5e1a2d8c9 100644
--- a/boost/compute/algorithm/detail/merge_sort_on_gpu.hpp
+++ b/boost/compute/algorithm/detail/merge_sort_on_gpu.hpp
@@ -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");
@@ -249,8 +250,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 +422,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" <<
diff --git a/boost/compute/algorithm/detail/radix_sort.hpp b/boost/compute/algorithm/detail/radix_sort.hpp
index 8e6d5f9c0a..53b1205c70 100644
--- a/boost/compute/algorithm/detail/radix_sort.hpp
+++ b/boost/compute/algorithm/detail/radix_sort.hpp
@@ -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>
@@ -305,9 +308,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");
diff --git a/boost/compute/algorithm/detail/serial_reduce.hpp b/boost/compute/algorithm/detail/serial_reduce.hpp
index 53aaf140fe..8b121274b9 100644
--- a/boost/compute/algorithm/detail/serial_reduce.hpp
+++ b/boost/compute/algorithm/detail/serial_reduce.hpp
@@ -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,
diff --git a/boost/compute/algorithm/detail/serial_reduce_by_key.hpp b/boost/compute/algorithm/detail/serial_reduce_by_key.hpp
index f9bda8e476..6fb04baa6d 100644
--- a/boost/compute/algorithm/detail/serial_reduce_by_key.hpp
+++ b/boost/compute/algorithm/detail/serial_reduce_by_key.hpp
@@ -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" <<
diff --git a/boost/compute/algorithm/equal.hpp b/boost/compute/algorithm/equal.hpp
index 35d0c5f0ea..c3c8053b71 100644
--- a/boost/compute/algorithm/equal.hpp
+++ b/boost/compute/algorithm/equal.hpp
@@ -20,6 +20,8 @@ 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,
diff --git a/boost/compute/algorithm/equal_range.hpp b/boost/compute/algorithm/equal_range.hpp
index fd82177324..d7008e3cf4 100644
--- a/boost/compute/algorithm/equal_range.hpp
+++ b/boost/compute/algorithm/equal_range.hpp
@@ -23,6 +23,8 @@ 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,
diff --git a/boost/compute/algorithm/exclusive_scan.hpp b/boost/compute/algorithm/exclusive_scan.hpp
index 205d3de658..806a172cf4 100644
--- a/boost/compute/algorithm/exclusive_scan.hpp
+++ b/boost/compute/algorithm/exclusive_scan.hpp
@@ -44,6 +44,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
diff --git a/boost/compute/algorithm/fill.hpp b/boost/compute/algorithm/fill.hpp
index c711f46b94..646d8acda4 100644
--- a/boost/compute/algorithm/fill.hpp
+++ b/boost/compute/algorithm/fill.hpp
@@ -64,7 +64,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 +172,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 +205,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 +251,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 +271,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,
diff --git a/boost/compute/algorithm/fill_n.hpp b/boost/compute/algorithm/fill_n.hpp
index 18a8f706a5..6be2d280a6 100644
--- a/boost/compute/algorithm/fill_n.hpp
+++ b/boost/compute/algorithm/fill_n.hpp
@@ -20,6 +20,8 @@ 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,
diff --git a/boost/compute/algorithm/find.hpp b/boost/compute/algorithm/find.hpp
index ef3ebf0c47..a6225b8c99 100644
--- a/boost/compute/algorithm/find.hpp
+++ b/boost/compute/algorithm/find.hpp
@@ -22,6 +22,8 @@ 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,
diff --git a/boost/compute/algorithm/find_end.hpp b/boost/compute/algorithm/find_end.hpp
index 265a1da542..a0a1b2e8c9 100644
--- a/boost/compute/algorithm/find_end.hpp
+++ b/boost/compute/algorithm/find_end.hpp
@@ -26,8 +26,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 +90,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,
diff --git a/boost/compute/algorithm/find_if.hpp b/boost/compute/algorithm/find_if.hpp
index db99cc0396..074b47e280 100644
--- a/boost/compute/algorithm/find_if.hpp
+++ b/boost/compute/algorithm/find_if.hpp
@@ -20,6 +20,8 @@ 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,
diff --git a/boost/compute/algorithm/find_if_not.hpp b/boost/compute/algorithm/find_if_not.hpp
index 61de050d31..a008a99469 100644
--- a/boost/compute/algorithm/find_if_not.hpp
+++ b/boost/compute/algorithm/find_if_not.hpp
@@ -22,6 +22,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,
diff --git a/boost/compute/algorithm/for_each.hpp b/boost/compute/algorithm/for_each.hpp
index 3ed399e6e9..7afba2b5f5 100644
--- a/boost/compute/algorithm/for_each.hpp
+++ b/boost/compute/algorithm/for_each.hpp
@@ -45,6 +45,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,
diff --git a/boost/compute/algorithm/for_each_n.hpp b/boost/compute/algorithm/for_each_n.hpp
index d0be784bf7..77932ab209 100644
--- a/boost/compute/algorithm/for_each_n.hpp
+++ b/boost/compute/algorithm/for_each_n.hpp
@@ -19,6 +19,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,
diff --git a/boost/compute/algorithm/gather.hpp b/boost/compute/algorithm/gather.hpp
index 24c5c727ae..62442587f7 100644
--- a/boost/compute/algorithm/gather.hpp
+++ b/boost/compute/algorithm/gather.hpp
@@ -62,6 +62,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,
diff --git a/boost/compute/algorithm/generate.hpp b/boost/compute/algorithm/generate.hpp
index c70a542683..9ac76a3dca 100644
--- a/boost/compute/algorithm/generate.hpp
+++ b/boost/compute/algorithm/generate.hpp
@@ -22,6 +22,8 @@ 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,
diff --git a/boost/compute/algorithm/generate_n.hpp b/boost/compute/algorithm/generate_n.hpp
index 6d8e607b64..066a831ddf 100644
--- a/boost/compute/algorithm/generate_n.hpp
+++ b/boost/compute/algorithm/generate_n.hpp
@@ -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,
diff --git a/boost/compute/algorithm/includes.hpp b/boost/compute/algorithm/includes.hpp
index c4e7c793e7..cfef9540a7 100644
--- a/boost/compute/algorithm/includes.hpp
+++ b/boost/compute/algorithm/includes.hpp
@@ -110,6 +110,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,
diff --git a/boost/compute/algorithm/inclusive_scan.hpp b/boost/compute/algorithm/inclusive_scan.hpp
index 9f98beaf7c..84f1b8cbf7 100644
--- a/boost/compute/algorithm/inclusive_scan.hpp
+++ b/boost/compute/algorithm/inclusive_scan.hpp
@@ -42,6 +42,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
diff --git a/boost/compute/algorithm/inner_product.hpp b/boost/compute/algorithm/inner_product.hpp
index 614611f91e..0aeaf9110e 100644
--- a/boost/compute/algorithm/inner_product.hpp
+++ b/boost/compute/algorithm/inner_product.hpp
@@ -26,6 +26,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,
diff --git a/boost/compute/algorithm/inplace_merge.hpp b/boost/compute/algorithm/inplace_merge.hpp
index 3080950df5..91f5be5335 100644
--- a/boost/compute/algorithm/inplace_merge.hpp
+++ b/boost/compute/algorithm/inplace_merge.hpp
@@ -23,6 +23,8 @@ namespace compute {
 
 /// Merges the sorted values in the range [\p first, \p middle) with
 /// the sorted values in the range [\p middle, \p last) in-place.
+///
+/// Space complexity: \Omega(n)
 template<class Iterator>
 inline void inplace_merge(Iterator first,
                           Iterator middle,
diff --git a/boost/compute/algorithm/iota.hpp b/boost/compute/algorithm/iota.hpp
index 084c3d8d97..4cd7aa9c7b 100644
--- a/boost/compute/algorithm/iota.hpp
+++ b/boost/compute/algorithm/iota.hpp
@@ -26,6 +26,8 @@ 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,
diff --git a/boost/compute/algorithm/is_partitioned.hpp b/boost/compute/algorithm/is_partitioned.hpp
index 3916825057..6ad24f240f 100644
--- a/boost/compute/algorithm/is_partitioned.hpp
+++ b/boost/compute/algorithm/is_partitioned.hpp
@@ -21,6 +21,8 @@ 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,
diff --git a/boost/compute/algorithm/is_permutation.hpp b/boost/compute/algorithm/is_permutation.hpp
index 1e502efb37..88b89b7973 100644
--- a/boost/compute/algorithm/is_permutation.hpp
+++ b/boost/compute/algorithm/is_permutation.hpp
@@ -36,6 +36,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,
diff --git a/boost/compute/algorithm/is_sorted.hpp b/boost/compute/algorithm/is_sorted.hpp
index a605159ac3..7441620978 100644
--- a/boost/compute/algorithm/is_sorted.hpp
+++ b/boost/compute/algorithm/is_sorted.hpp
@@ -30,6 +30,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,
diff --git a/boost/compute/algorithm/lexicographical_compare.hpp b/boost/compute/algorithm/lexicographical_compare.hpp
index c4f7120807..952e678a68 100644
--- a/boost/compute/algorithm/lexicographical_compare.hpp
+++ b/boost/compute/algorithm/lexicographical_compare.hpp
@@ -42,10 +42,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 +103,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,
diff --git a/boost/compute/algorithm/lower_bound.hpp b/boost/compute/algorithm/lower_bound.hpp
index b2011c66ef..f78bbd6364 100644
--- a/boost/compute/algorithm/lower_bound.hpp
+++ b/boost/compute/algorithm/lower_bound.hpp
@@ -22,6 +22,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
diff --git a/boost/compute/algorithm/max_element.hpp b/boost/compute/algorithm/max_element.hpp
index 55f2f7ffbf..f9df37420c 100644
--- a/boost/compute/algorithm/max_element.hpp
+++ b/boost/compute/algorithm/max_element.hpp
@@ -43,6 +43,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
diff --git a/boost/compute/algorithm/merge.hpp b/boost/compute/algorithm/merge.hpp
index 875a283044..ff3e6e879a 100644
--- a/boost/compute/algorithm/merge.hpp
+++ b/boost/compute/algorithm/merge.hpp
@@ -37,6 +37,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,
diff --git a/boost/compute/algorithm/min_element.hpp b/boost/compute/algorithm/min_element.hpp
index 62744efb98..b52e2670cb 100644
--- a/boost/compute/algorithm/min_element.hpp
+++ b/boost/compute/algorithm/min_element.hpp
@@ -43,6 +43,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
diff --git a/boost/compute/algorithm/minmax_element.hpp b/boost/compute/algorithm/minmax_element.hpp
index 3f44c09eaf..4b2aae6dee 100644
--- a/boost/compute/algorithm/minmax_element.hpp
+++ b/boost/compute/algorithm/minmax_element.hpp
@@ -31,6 +31,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>
diff --git a/boost/compute/algorithm/mismatch.hpp b/boost/compute/algorithm/mismatch.hpp
index e7db883004..ff31f49f97 100644
--- a/boost/compute/algorithm/mismatch.hpp
+++ b/boost/compute/algorithm/mismatch.hpp
@@ -28,6 +28,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,
diff --git a/boost/compute/algorithm/next_permutation.hpp b/boost/compute/algorithm/next_permutation.hpp
index e81fbd2ee8..061ea1efe9 100644
--- a/boost/compute/algorithm/next_permutation.hpp
+++ b/boost/compute/algorithm/next_permutation.hpp
@@ -131,6 +131,7 @@ 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,
diff --git a/boost/compute/algorithm/none_of.hpp b/boost/compute/algorithm/none_of.hpp
index c25dd12a87..fc3ed94bc1 100644
--- a/boost/compute/algorithm/none_of.hpp
+++ b/boost/compute/algorithm/none_of.hpp
@@ -20,6 +20,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,
diff --git a/boost/compute/algorithm/nth_element.hpp b/boost/compute/algorithm/nth_element.hpp
index 68f7a3dbc0..93344271dd 100644
--- a/boost/compute/algorithm/nth_element.hpp
+++ b/boost/compute/algorithm/nth_element.hpp
@@ -23,6 +23,8 @@ namespace compute {
 
 /// Rearranges the elements in the range [\p first, \p last) such that
 /// the \p nth element would be in that position in a sorted sequence.
+///
+/// Space complexity: \Omega(3n)
 template<class Iterator, class Compare>
 inline void nth_element(Iterator first,
                         Iterator nth,
diff --git a/boost/compute/algorithm/partial_sum.hpp b/boost/compute/algorithm/partial_sum.hpp
index d440369a5a..53d36a9db0 100644
--- a/boost/compute/algorithm/partial_sum.hpp
+++ b/boost/compute/algorithm/partial_sum.hpp
@@ -21,6 +21,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,
diff --git a/boost/compute/algorithm/partition.hpp b/boost/compute/algorithm/partition.hpp
index 7860350e0d..59d0c78f7e 100644
--- a/boost/compute/algorithm/partition.hpp
+++ b/boost/compute/algorithm/partition.hpp
@@ -22,6 +22,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>
diff --git a/boost/compute/algorithm/partition_copy.hpp b/boost/compute/algorithm/partition_copy.hpp
index 80a2c6475f..3215ec0736 100644
--- a/boost/compute/algorithm/partition_copy.hpp
+++ b/boost/compute/algorithm/partition_copy.hpp
@@ -24,6 +24,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,
diff --git a/boost/compute/algorithm/partition_point.hpp b/boost/compute/algorithm/partition_point.hpp
index 3cc2bc0ca6..748824512d 100644
--- a/boost/compute/algorithm/partition_point.hpp
+++ b/boost/compute/algorithm/partition_point.hpp
@@ -29,6 +29,8 @@ 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>
diff --git a/boost/compute/algorithm/prev_permutation.hpp b/boost/compute/algorithm/prev_permutation.hpp
index 03c01bf8f4..ea20835caa 100644
--- a/boost/compute/algorithm/prev_permutation.hpp
+++ b/boost/compute/algorithm/prev_permutation.hpp
@@ -131,6 +131,7 @@ 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,
diff --git a/boost/compute/algorithm/random_shuffle.hpp b/boost/compute/algorithm/random_shuffle.hpp
index 7d2d46a133..8e020830a5 100644
--- a/boost/compute/algorithm/random_shuffle.hpp
+++ b/boost/compute/algorithm/random_shuffle.hpp
@@ -28,6 +28,8 @@ 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,
diff --git a/boost/compute/algorithm/reduce.hpp b/boost/compute/algorithm/reduce.hpp
index 19d070019f..e71d90fe24 100644
--- a/boost/compute/algorithm/reduce.hpp
+++ b/boost/compute/algorithm/reduce.hpp
@@ -153,6 +153,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 +265,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,
diff --git a/boost/compute/algorithm/reduce_by_key.hpp b/boost/compute/algorithm/reduce_by_key.hpp
index 87c73e887f..1a233c7dd4 100644
--- a/boost/compute/algorithm/reduce_by_key.hpp
+++ b/boost/compute/algorithm/reduce_by_key.hpp
@@ -51,6 +51,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,
diff --git a/boost/compute/algorithm/remove.hpp b/boost/compute/algorithm/remove.hpp
index 98feb1f9d8..086ba8cc7f 100644
--- a/boost/compute/algorithm/remove.hpp
+++ b/boost/compute/algorithm/remove.hpp
@@ -22,6 +22,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,
diff --git a/boost/compute/algorithm/remove_if.hpp b/boost/compute/algorithm/remove_if.hpp
index 5e416bef88..9aece18bbd 100644
--- a/boost/compute/algorithm/remove_if.hpp
+++ b/boost/compute/algorithm/remove_if.hpp
@@ -22,6 +22,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,
diff --git a/boost/compute/algorithm/replace.hpp b/boost/compute/algorithm/replace.hpp
index fd649a2fad..336c1d3e0f 100644
--- a/boost/compute/algorithm/replace.hpp
+++ b/boost/compute/algorithm/replace.hpp
@@ -68,6 +68,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,
diff --git a/boost/compute/algorithm/replace_copy.hpp b/boost/compute/algorithm/replace_copy.hpp
index 7224bd3ae6..34f61b514f 100644
--- a/boost/compute/algorithm/replace_copy.hpp
+++ b/boost/compute/algorithm/replace_copy.hpp
@@ -25,6 +25,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
diff --git a/boost/compute/algorithm/reverse.hpp b/boost/compute/algorithm/reverse.hpp
index b6a9e8098c..15fe5533ac 100644
--- a/boost/compute/algorithm/reverse.hpp
+++ b/boost/compute/algorithm/reverse.hpp
@@ -52,6 +52,8 @@ 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,
diff --git a/boost/compute/algorithm/reverse_copy.hpp b/boost/compute/algorithm/reverse_copy.hpp
index c839f44651..9fda9d4e27 100644
--- a/boost/compute/algorithm/reverse_copy.hpp
+++ b/boost/compute/algorithm/reverse_copy.hpp
@@ -51,6 +51,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
diff --git a/boost/compute/algorithm/rotate.hpp b/boost/compute/algorithm/rotate.hpp
index 54cb073cc2..715699340d 100644
--- a/boost/compute/algorithm/rotate.hpp
+++ b/boost/compute/algorithm/rotate.hpp
@@ -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,
diff --git a/boost/compute/algorithm/rotate_copy.hpp b/boost/compute/algorithm/rotate_copy.hpp
index fa1b44c5e5..679b3c998b 100644
--- a/boost/compute/algorithm/rotate_copy.hpp
+++ b/boost/compute/algorithm/rotate_copy.hpp
@@ -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,
diff --git a/boost/compute/algorithm/scatter.hpp b/boost/compute/algorithm/scatter.hpp
index bea4201628..8ae5a99443 100644
--- a/boost/compute/algorithm/scatter.hpp
+++ b/boost/compute/algorithm/scatter.hpp
@@ -79,6 +79,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,
diff --git a/boost/compute/algorithm/scatter_if.hpp b/boost/compute/algorithm/scatter_if.hpp
index 159edd8c86..c7db51d3be 100644
--- a/boost/compute/algorithm/scatter_if.hpp
+++ b/boost/compute/algorithm/scatter_if.hpp
@@ -83,7 +83,7 @@ private:
 /// at \p map if stencil is resolved to true. By default the predicate is
 /// an identity
 ///
-///
+/// Space complexity: \Omega(1)
 template<class InputIterator, class MapIterator, class StencilIterator, class OutputIterator,
          class Predicate>
 inline void scatter_if(InputIterator first,
diff --git a/boost/compute/algorithm/search.hpp b/boost/compute/algorithm/search.hpp
index 3d3d035b3c..a1f3dece62 100644
--- a/boost/compute/algorithm/search.hpp
+++ b/boost/compute/algorithm/search.hpp
@@ -34,6 +34,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,
diff --git a/boost/compute/algorithm/search_n.hpp b/boost/compute/algorithm/search_n.hpp
index 9e03111bb0..86ff64dfd9 100644
--- a/boost/compute/algorithm/search_n.hpp
+++ b/boost/compute/algorithm/search_n.hpp
@@ -102,6 +102,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,
diff --git a/boost/compute/algorithm/set_difference.hpp b/boost/compute/algorithm/set_difference.hpp
index 17ce7bd3f6..85a846ba13 100644
--- a/boost/compute/algorithm/set_difference.hpp
+++ b/boost/compute/algorithm/set_difference.hpp
@@ -122,6 +122,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,
diff --git a/boost/compute/algorithm/set_intersection.hpp b/boost/compute/algorithm/set_intersection.hpp
index 50f291e84a..74d46f57c6 100644
--- a/boost/compute/algorithm/set_intersection.hpp
+++ b/boost/compute/algorithm/set_intersection.hpp
@@ -110,6 +110,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,
diff --git a/boost/compute/algorithm/set_symmetric_difference.hpp b/boost/compute/algorithm/set_symmetric_difference.hpp
index 6e60b38511..34d280daa3 100644
--- a/boost/compute/algorithm/set_symmetric_difference.hpp
+++ b/boost/compute/algorithm/set_symmetric_difference.hpp
@@ -133,13 +133,16 @@ 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())
 {
     typedef typename std::iterator_traits<InputIterator1>::value_type value_type;
 
diff --git a/boost/compute/algorithm/set_union.hpp b/boost/compute/algorithm/set_union.hpp
index c61f7b29b3..6b405a0905 100644
--- a/boost/compute/algorithm/set_union.hpp
+++ b/boost/compute/algorithm/set_union.hpp
@@ -135,6 +135,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,
diff --git a/boost/compute/algorithm/sort.hpp b/boost/compute/algorithm/sort.hpp
index 7e0a583e3e..b8fa90f335 100644
--- a/boost/compute/algorithm/sort.hpp
+++ b/boost/compute/algorithm/sort.hpp
@@ -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,
diff --git a/boost/compute/algorithm/sort_by_key.hpp b/boost/compute/algorithm/sort_by_key.hpp
index c39bcf9890..fdd2d1c481 100644
--- a/boost/compute/algorithm/sort_by_key.hpp
+++ b/boost/compute/algorithm/sort_by_key.hpp
@@ -128,6 +128,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,
diff --git a/boost/compute/algorithm/stable_partition.hpp b/boost/compute/algorithm/stable_partition.hpp
index 283b068283..2b07f034b9 100644
--- a/boost/compute/algorithm/stable_partition.hpp
+++ b/boost/compute/algorithm/stable_partition.hpp
@@ -33,6 +33,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>
diff --git a/boost/compute/algorithm/stable_sort.hpp b/boost/compute/algorithm/stable_sort.hpp
index 381fc81bc0..0857d75dc9 100644
--- a/boost/compute/algorithm/stable_sort.hpp
+++ b/boost/compute/algorithm/stable_sort.hpp
@@ -72,6 +72,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,
diff --git a/boost/compute/algorithm/stable_sort_by_key.hpp b/boost/compute/algorithm/stable_sort_by_key.hpp
index 878f999f44..ce8811ef19 100644
--- a/boost/compute/algorithm/stable_sort_by_key.hpp
+++ b/boost/compute/algorithm/stable_sort_by_key.hpp
@@ -126,6 +126,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,
diff --git a/boost/compute/algorithm/swap_ranges.hpp b/boost/compute/algorithm/swap_ranges.hpp
index 6ff3e14f6a..a706df7a61 100644
--- a/boost/compute/algorithm/swap_ranges.hpp
+++ b/boost/compute/algorithm/swap_ranges.hpp
@@ -21,6 +21,8 @@ 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,
diff --git a/boost/compute/algorithm/transform.hpp b/boost/compute/algorithm/transform.hpp
index 68750a6523..9137604d55 100644
--- a/boost/compute/algorithm/transform.hpp
+++ b/boost/compute/algorithm/transform.hpp
@@ -29,6 +29,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,
diff --git a/boost/compute/algorithm/transform_if.hpp b/boost/compute/algorithm/transform_if.hpp
index 0eb0fd434e..9a98102d27 100644
--- a/boost/compute/algorithm/transform_if.hpp
+++ b/boost/compute/algorithm/transform_if.hpp
@@ -26,6 +26,7 @@ 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 +54,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 +97,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,
diff --git a/boost/compute/algorithm/transform_reduce.hpp b/boost/compute/algorithm/transform_reduce.hpp
index fbeee5a691..a59a76aefd 100644
--- a/boost/compute/algorithm/transform_reduce.hpp
+++ b/boost/compute/algorithm/transform_reduce.hpp
@@ -30,6 +30,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,
diff --git a/boost/compute/algorithm/unique.hpp b/boost/compute/algorithm/unique.hpp
index faa36bad9d..8b7e2a0d0d 100644
--- a/boost/compute/algorithm/unique.hpp
+++ b/boost/compute/algorithm/unique.hpp
@@ -31,6 +31,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,
diff --git a/boost/compute/algorithm/unique_copy.hpp b/boost/compute/algorithm/unique_copy.hpp
index 2ce60a9359..d5fffd4ff9 100644
--- a/boost/compute/algorithm/unique_copy.hpp
+++ b/boost/compute/algorithm/unique_copy.hpp
@@ -127,6 +127,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,
diff --git a/boost/compute/algorithm/upper_bound.hpp b/boost/compute/algorithm/upper_bound.hpp
index a5a82d301c..f592c79b9a 100644
--- a/boost/compute/algorithm/upper_bound.hpp
+++ b/boost/compute/algorithm/upper_bound.hpp
@@ -22,6 +22,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,