Merge pull request #433 from haahh/dev_reverse_copy

reverse_copy() performance improvement

Author: kylelutz

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;

perf/CMakeLists.txt

@@ -41,6 +41,7 @@ set(BENCHMARKS
   partition_point
   prev_permutation
   reverse
+  reverse_copy
   rotate
   rotate_copy
   host_sort
@@ -80,6 +81,7 @@ set(STL_BENCHMARKS
   stl_partition
   stl_prev_permutation
   stl_reverse
+  stl_reverse_copy
   stl_rotate
   stl_rotate_copy
   stl_saxpy
@@ -123,6 +125,7 @@ if(${BOOST_COMPUTE_HAVE_CUDA})
     thrust_partial_sum
     thrust_partition
     thrust_reverse
+    thrust_reverse_copy
     thrust_rotate
     thrust_saxpy
     thrust_set_difference

perf/perf.py

@@ -123,6 +123,7 @@ def run_benchmark(name, sizes, vs=[]):
             "partial_sum",
             "partition",
             "reverse",
+            "reverse_copy",
             "rotate",
             "saxpy",
             "sort",
@@ -149,6 +150,7 @@ def run_benchmark(name, sizes, vs=[]):
             "partition_point",
             "prev_permutation",
             "reverse",
+            "reverse_copy",
             "rotate",
             "rotate_copy",
             "saxpy",

perf/perf_reverse_copy.cpp

@@ -0,0 +1,65 @@
+//---------------------------------------------------------------------------//
+// Copyright (c) 2015 Jakub Szuppe <j.szuppe@gmail.com>
+//
+// Distributed under the Boost Software License, Version 1.0
+// See accompanying file LICENSE_1_0.txt or copy at
+// http://www.boost.org/LICENSE_1_0.txt
+//
+// See http://kylelutz.github.com/compute for more information.
+//---------------------------------------------------------------------------//
+
+#include <algorithm>
+#include <iostream>
+#include <numeric>
+#include <vector>
+
+#include <boost/compute/system.hpp>
+#include <boost/compute/algorithm/reverse_copy.hpp>
+#include <boost/compute/container/vector.hpp>
+
+#include "perf.hpp"
+
+int rand_int()
+{
+    return static_cast<int>((rand() / double(RAND_MAX)) * 25.0);
+}
+
+int main(int argc, char *argv[])
+{
+    perf_parse_args(argc, argv);
+    std::cout << "size: " << PERF_N << std::endl;
+
+    // setup context and queue for the default device
+    boost::compute::device device = boost::compute::system::default_device();
+    boost::compute::context context(device);
+    boost::compute::command_queue queue(context, device);
+    std::cout << "device: " << device.name() << std::endl;
+
+    // create vector of random numbers on the host
+    std::vector<int> host_vector(PERF_N);
+    std::generate(host_vector.begin(), host_vector.end(), rand_int);
+
+    // create vector on the device and copy the data
+    boost::compute::vector<int> device_vector(PERF_N, context);
+    boost::compute::copy(
+        host_vector.begin(), host_vector.end(), device_vector.begin(), queue
+    );
+
+    // create vector on the device for reversed data
+    boost::compute::vector<int> device_reversed_vector(PERF_N, context);
+
+    perf_timer t;
+    for(size_t trial = 0; trial < PERF_TRIALS; trial++){
+        t.start();
+        boost::compute::reverse_copy(
+            device_vector.begin(), device_vector.end(),
+            device_reversed_vector.begin(),
+            queue
+        );
+        queue.finish();
+        t.stop();
+    }
+    std::cout << "time: " << t.min_time() / 1e6 << " ms" << std::endl;
+
+    return 0;
+}

perf/perf_stl_reverse_copy.cpp

@@ -0,0 +1,45 @@
+//---------------------------------------------------------------------------//
+// Copyright (c) 2015 Jakub Szuppe <j.szuppe@gmail.com>
+//
+// Distributed under the Boost Software License, Version 1.0
+// See accompanying file LICENSE_1_0.txt or copy at
+// http://www.boost.org/LICENSE_1_0.txt
+//
+// See http://kylelutz.github.com/compute for more information.
+//---------------------------------------------------------------------------//
+
+#include <algorithm>
+#include <iostream>
+#include <numeric>
+#include <vector>
+
+#include "perf.hpp"
+
+int rand_int()
+{
+    return static_cast<int>((rand() / double(RAND_MAX)) * 25.0);
+}
+
+int main(int argc, char *argv[])
+{
+    perf_parse_args(argc, argv);
+    std::cout << "size: " << PERF_N << std::endl;
+
+    // create vector of random numbers on the host
+    std::vector<int> host_vector(PERF_N);
+    std::generate(host_vector.begin(), host_vector.end(), rand_int);
+
+    // create vector for reversed data
+    std::vector<int> host_reversed_vector(PERF_N);
+
+    perf_timer t;
+    for(size_t trial = 0; trial < PERF_TRIALS; trial++){
+        t.start();
+        std::reverse_copy(host_vector.begin(), host_vector.end(),
+                          host_reversed_vector.begin());
+        t.stop();
+    }
+    std::cout << "time: " << t.min_time() / 1e6 << " ms" << std::endl;
+
+    return 0;
+}

perf/perf_thrust_reverse_copy.cu

@@ -0,0 +1,46 @@
+//---------------------------------------------------------------------------//
+// Copyright (c) 2015 Jakub Szuppe <j.szuppe@gmail.com>
+//
+// Distributed under the Boost Software License, Version 1.0
+// See accompanying file LICENSE_1_0.txt or copy at
+// http://www.boost.org/LICENSE_1_0.txt
+//
+// See http://kylelutz.github.com/compute for more information.
+//---------------------------------------------------------------------------//
+
+#include <algorithm>
+#include <cstdlib>
+
+#include <thrust/copy.h>
+#include <thrust/device_vector.h>
+#include <thrust/generate.h>
+#include <thrust/host_vector.h>
+#include <thrust/reverse.h>
+
+#include "perf.hpp"
+
+int main(int argc, char *argv[])
+{
+    perf_parse_args(argc, argv);
+
+    std::cout << "size: " << PERF_N << std::endl;
+    thrust::host_vector<int> h_vec = generate_random_vector<int>(PERF_N);
+
+	// transfer data to the device
+    thrust::device_vector<int> d_vec;   
+    d_vec = h_vec;
+    
+    // device vector for reversed data
+    thrust::device_vector<int> d_reversed_vec(PERF_N);
+    
+    perf_timer t;
+    for(size_t trial = 0; trial < PERF_TRIALS; trial++){
+        t.start();
+        thrust::reverse_copy(d_vec.begin(), d_vec.end(), d_reversed_vec.begin());
+        cudaDeviceSynchronize();
+        t.stop();
+    }
+    std::cout << "time: " << t.min_time() / 1e6 << " ms" << std::endl;
+
+    return 0;
+}

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)