Separate kernel for reverse_copy()

Separate kernel for reverse_copy() algorithm improves its performance,
new tests for reverse_copy_int

Author: jszuppe

include/boost/compute/algorithm/reverse_copy.hpp

@@ -20,6 +20,35 @@
 
 namespace boost {
 namespace compute {
+namespace detail {
+
+template<class Iterator, class OutputIterator>
+struct reverse_copy_kernel : public meta_kernel
+{
+    reverse_copy_kernel(Iterator first, Iterator last, OutputIterator result)
+        : meta_kernel("reverse_copy")
+    {
+        typedef typename std::iterator_traits<Iterator>::value_type value_type;
+
+        // store size of the range
+        m_size = detail::iterator_range_size(first, last);
+        add_set_arg<const cl_uint>("size", static_cast<const cl_uint>(m_size));
+
+        *this <<
+            decl<cl_uint>("i") << " = get_global_id(0);\n" <<
+            decl<cl_uint>("j") << " = size - get_global_id(0) - 1;\n" <<
+            result[var<cl_uint>("j")] << "=" << first[var<cl_uint>("i")] << ";\n";
+    }
+
+    void exec(command_queue &queue)
+    {
+        exec_1d(queue, 0, m_size);
+    }
+
+    size_t m_size;
+};
+
+} // end detail namespace
 
 /// Copies the elements in the range [\p first, \p last) in reversed
 /// order to the range beginning at \p result.
@@ -36,11 +65,11 @@ reverse_copy(InputIterator first,
 
     difference_type count = std::distance(first, last);
 
-    // copy data to result
-    ::boost::compute::copy(first, last, result, queue);
+    detail::reverse_copy_kernel<InputIterator, OutputIterator>
+        kernel(first, last, result);
 
-    // reverse result
-    ::boost::compute::reverse(result, result + count, queue);
+    // run kernel
+    kernel.exec(queue);
 
     // return iterator to the end of result
     return result + count;

test/test_reverse.cpp

@@ -68,6 +68,14 @@ BOOST_AUTO_TEST_CASE(reverse_copy_int)
         bc::reverse_copy(a.begin(), a.end(), b.begin());
     BOOST_CHECK(iter == b.end());
     CHECK_RANGE_EQUAL(int, 5, b, (4, 3, 2, 1, 0));
+
+    iter = bc::reverse_copy(b.begin() + 1, b.end(), a.begin() + 1);
+    BOOST_CHECK(iter == a.end());
+    CHECK_RANGE_EQUAL(int, 5, a, (0, 0, 1, 2, 3));
+
+    iter = bc::reverse_copy(a.begin(), a.end() - 1, b.begin());
+    BOOST_CHECK(iter == (b.end() - 1));
+    CHECK_RANGE_EQUAL(int, 5, b, (2, 1, 0, 0, 0));
 }
 
 BOOST_AUTO_TEST_CASE(reverse_copy_counting_iterator)