Implement reverse in libcudf

cpp/CMakeLists.txt

@@ -168,6 +168,7 @@ add_library(cudf
     src/copying/gather.cu
     src/copying/get_element.cu
     src/copying/pack.cpp
+    src/copying/reverse.cu
     src/copying/sample.cu
     src/copying/scatter.cu
     src/copying/shift.cu

cpp/include/cudf/copying.hpp

@@ -81,6 +81,36 @@ std::unique_ptr<table> gather(
   out_of_bounds_policy bounds_policy  = out_of_bounds_policy::DONT_CHECK,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
+/**
+ * @brief Reverses the rows within a table.
+ * Creates a new table that is the reverse of @p source_table.
+ * Example:
+ * ```
+ * source = [[4,5,6], [7,8,9], [10,11,12]]
+ * return = [[6,5,4], [9,8,7], [12,11,10]]
+ * ```
+ *
+ * @param source_table Table that will be reversed
+ */
+std::unique_ptr<table> reverse(
+  table_view const& source_table,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+
+/**
+ * @brief Reverses the elements of a column
+ * Creates a new column that is the reverse of @p source_column.
+ * Example:
+ * ```
+ * source = [4,5,6]
+ * return = [6,5,4]
+ * ```
+ *
+ * @param source_column Column that will be reversed
+ */
+std::unique_ptr<column> reverse(
+  column_view const& source_column,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+
 /**
  * @brief Scatters the rows of the source table into a copy of the target table
  * according to a scatter map.

cpp/src/copying/reverse.cu

@@ -0,0 +1,68 @@
+/*
+ * Copyright (c) 2021, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include <cudf/column/column_view.hpp>
+#include <cudf/copying.hpp>
+#include <cudf/detail/gather.cuh>
+#include <cudf/detail/nvtx/ranges.hpp>
+#include <cudf/table/table.hpp>
+#include <cudf/table/table_view.hpp>
+
+#include <rmm/cuda_stream_view.hpp>
+#include <rmm/device_uvector.hpp>
+#include <rmm/exec_policy.hpp>
+#include <rmm/mr/device/per_device_resource.hpp>
+
+#include <thrust/iterator/constant_iterator.h>
+#include <thrust/iterator/counting_iterator.h>
+#include <thrust/iterator/transform_output_iterator.h>
+#include <thrust/scan.h>
+
+namespace cudf {
+namespace detail {
+std::unique_ptr<table> reverse(table_view const& source_table,
+                               rmm::cuda_stream_view stream,
+                               rmm::mr::device_memory_resource* mr)
+{
+  size_type num_rows = source_table.num_rows();
+  auto elements =
+    make_counting_transform_iterator(0, [num_rows] __device__(auto i) { return num_rows - i - 1; });
+  auto elements_end = elements + source_table.num_rows();
+
+  return gather(source_table, elements, elements_end, out_of_bounds_policy::DONT_CHECK, stream, mr);
+}
+
+std::unique_ptr<column> reverse(column_view const& source_column,
+                                rmm::cuda_stream_view stream,
+                                rmm::mr::device_memory_resource* mr)
+{
+  return std::move(cudf::reverse(table_view({source_column}))->release().front());
+}
+}  // namespace detail
+
+std::unique_ptr<table> reverse(table_view const& source_table, rmm::mr::device_memory_resource* mr)
+{
+  CUDF_FUNC_RANGE();
+  return detail::reverse(source_table, rmm::cuda_stream_default, mr);
+}
+
+std::unique_ptr<column> reverse(column_view const& source_column,
+                                rmm::mr::device_memory_resource* mr)
+{
+  CUDF_FUNC_RANGE();
+  return detail::reverse(source_column, rmm::cuda_stream_default, mr);
+}
+}  // namespace cudf

cpp/tests/CMakeLists.txt

@@ -218,7 +218,8 @@ ConfigureTest(COPYING_TEST
     copying/shift_tests.cpp
     copying/slice_tests.cpp
     copying/split_tests.cpp
-    copying/utility_tests.cpp)
+    copying/utility_tests.cpp
+    copying/reverse_tests.cpp)
 
 ###################################################################################################
 # - utilities tests -------------------------------------------------------------------------------

cpp/tests/copying/reverse_tests.cpp

@@ -0,0 +1,180 @@
+/*
+ * Copyright (c) 2021, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include <cudf_test/base_fixture.hpp>
+#include <cudf_test/column_utilities.hpp>
+#include <cudf_test/column_wrapper.hpp>
+#include <cudf_test/cudf_gtest.hpp>
+#include <cudf_test/type_lists.hpp>
+
+#include <cudf/copying.hpp>
+#include <cudf/detail/iterator.cuh>
+#include <cudf/scalar/scalar.hpp>
+#include <cudf/scalar/scalar_factories.hpp>
+#include <cudf/table/table.hpp>
+#include <cudf/table/table_view.hpp>
+
+#include <thrust/iterator/counting_iterator.h>
+#include <thrust/tabulate.h>
+
+template <typename T>
+class ReverseTypedTestFixture : public cudf::test::BaseFixture {
+};
+
+TYPED_TEST_CASE(ReverseTypedTestFixture, cudf::test::AllTypes);
+TYPED_TEST(ReverseTypedTestFixture, ReverseTable)
+{
+  using T = TypeParam;
+  constexpr cudf::size_type num_values{10};
+
+  auto input = cudf::test::fixed_width_column_wrapper<T, int32_t>(
+    thrust::make_counting_iterator(0), thrust::make_counting_iterator(0) + num_values);
+
+  auto expected_elements = cudf::detail::make_counting_transform_iterator(
+    0, [num_values] __device__(auto i) { return num_values - i - 1; });
+
+  auto expected =
+    cudf::test::fixed_width_column_wrapper<T, typename decltype(expected_elements)::value_type>(
+      expected_elements, expected_elements + num_values);
+
+  auto input_table = cudf::table_view{{input}};
+  auto const p_ret = cudf::reverse(input_table);
+
+  EXPECT_EQ(p_ret->num_columns(), 1);
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(p_ret->view().column(0), expected, true);
+}
+
+TYPED_TEST(ReverseTypedTestFixture, ReverseColumn)
+{
+  using T = TypeParam;
+  constexpr cudf::size_type num_values{10};
+
+  auto input = cudf::test::fixed_width_column_wrapper<T, int32_t>(
+    thrust::make_counting_iterator(0), thrust::make_counting_iterator(0) + num_values);
+
+  auto expected_elements = cudf::detail::make_counting_transform_iterator(
+    0, [num_values] __device__(auto i) { return num_values - i - 1; });
+
+  auto expected =
+    cudf::test::fixed_width_column_wrapper<T, typename decltype(expected_elements)::value_type>(
+      expected_elements, expected_elements + num_values);
+
+  auto const column_ret = cudf::reverse(input);
+
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(column_ret->view(), expected, true);
+}
+
+TYPED_TEST(ReverseTypedTestFixture, ReverseNullable)
+{
+  using T = TypeParam;
+  constexpr cudf::size_type num_values{20};
+
+  std::vector<int64_t> input_values(num_values);
+  std::iota(input_values.begin(), input_values.end(), 1);
+
+  thrust::host_vector<bool> input_valids(num_values);
+  thrust::tabulate(
+    thrust::seq, input_valids.begin(), input_valids.end(), [](auto i) { return not(i % 2); });
+
+  std::vector<T> expected_values(input_values.size());
+  thrust::host_vector<bool> expected_valids(input_valids.size());
+
+  std::transform(std::make_reverse_iterator(input_values.end()),
+                 std::make_reverse_iterator(input_values.begin()),
+                 expected_values.begin(),
+                 [](auto i) { return cudf::test::make_type_param_scalar<T>(i); });
+  std::reverse_copy(input_valids.begin(), input_valids.end(), expected_valids.begin());
+
+  cudf::test::fixed_width_column_wrapper<T, int64_t> input(
+    input_values.begin(), input_values.end(), input_valids.begin());
+
+  cudf::test::fixed_width_column_wrapper<T> expected(
+    expected_values.begin(), expected_values.end(), expected_valids.begin());
+
+  cudf::table_view input_table{{input}};
+  auto p_ret = cudf::reverse(input_table);
+
+  EXPECT_EQ(p_ret->num_columns(), 1);
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(p_ret->view().column(0), expected);
+}
+
+TYPED_TEST(ReverseTypedTestFixture, ZeroSizeInput)
+{
+  using T = TypeParam;
+  cudf::test::fixed_width_column_wrapper<T, int32_t> input(thrust::make_counting_iterator(0),
+                                                           thrust::make_counting_iterator(0));
+
+  cudf::test::fixed_width_column_wrapper<T, int32_t> expected(thrust::make_counting_iterator(0),
+                                                              thrust::make_counting_iterator(0));
+
+  cudf::table_view input_table{{input}};
+  auto p_ret = cudf::reverse(input_table);
+
+  EXPECT_EQ(p_ret->num_columns(), 1);
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(p_ret->view().column(0), expected);
+}
+
+class ReverseStringTestFixture : public cudf::test::BaseFixture {
+};
+
+TEST_F(ReverseStringTestFixture, ReverseNullable)
+{
+  constexpr cudf::size_type num_values{20};
+
+  std::vector<std::string> input_values(num_values);
+  thrust::host_vector<bool> input_valids(num_values);
+
+  thrust::tabulate(thrust::seq, input_values.begin(), input_values.end(), [](auto i) {
+    return "#" + std::to_string(i);
+  });
+  thrust::tabulate(
+    thrust::seq, input_valids.begin(), input_valids.end(), [](auto i) { return not(i % 2); });
+
+  std::vector<std::string> expected_values(input_values.size());
+  thrust::host_vector<bool> expected_valids(input_valids.size());
+
+  std::reverse_copy(input_values.begin(), input_values.end(), expected_values.begin());
+  std::reverse_copy(input_valids.begin(), input_valids.end(), expected_valids.begin());
+
+  auto input = cudf::test::strings_column_wrapper(
+    input_values.begin(), input_values.end(), input_valids.begin());
+
+  auto expected = cudf::test::strings_column_wrapper(
+    expected_values.begin(), expected_values.end(), expected_valids.begin());
+
+  cudf::table_view input_table{{input}};
+  auto p_ret = cudf::reverse(input_table);
+
+  EXPECT_EQ(p_ret->num_columns(), 1);
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(p_ret->view().column(0), expected);
+}
+
+TEST_F(ReverseStringTestFixture, ZeroSizeInput)
+{
+  std::vector<std::string> input_values{};
+  auto input = cudf::test::strings_column_wrapper(input_values.begin(), input_values.end());
+
+  auto count = cudf::test::fixed_width_column_wrapper<cudf::size_type>(
+    thrust::make_counting_iterator(0), thrust::make_counting_iterator(0));
+
+  auto expected = cudf::test::strings_column_wrapper(input_values.begin(), input_values.end());
+
+  cudf::table_view input_table{{input}};
+  auto p_ret = cudf::reverse(input_table);
+
+  EXPECT_EQ(p_ret->num_columns(), 1);
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(p_ret->view().column(0), expected);
+}

python/cudf/cudf/_lib/copying.pyx

@@ -276,6 +276,47 @@ def scatter(object input, object scatter_map, Table target,
         return _scatter_scalar(input, scatter_map, target, bounds_check)
 
 
+def _reverse_column(Column source_column):
+    cdef column_view reverse_column_view = source_column.view()
+
+    cdef unique_ptr[column] c_result
+
+    with nogil:
+        c_result = move(cpp_copying.reverse(
+            reverse_column_view
+        ))
+
+    return Column.from_unique_ptr(
+        move(c_result)
+    )
+
+
+def _reverse_table(Table source_table):
+    cdef table_view reverse_table_view = source_table.view()
+
+    cdef unique_ptr[table] c_result
+    with nogil:
+        c_result = move(cpp_copying.reverse(
+            reverse_table_view
+        ))
+
+    return Table.from_unique_ptr(
+        move(c_result),
+        column_names=source_table._column_names,
+        index_names=source_table._index_names
+    )
+
+
+def reverse(object source):
+    """
+    Reversing a column or a table
+    """
+    if isinstance(source, Column):
+        return _reverse_column(source)
+    else:
+        return _reverse_table(source)
+
+
 def column_empty_like(Column input_column):
 
     cdef column_view input_column_view = input_column.view()

python/cudf/cudf/_lib/cpp/copying.pxd

@@ -31,6 +31,14 @@ cdef extern from "cudf/copying.hpp" namespace "cudf" nogil:
         out_of_bounds_policy policy
     ) except +
 
+    cdef unique_ptr[table] reverse (
+        const table_view& source_table
+    ) except +
+
+    cdef unique_ptr[column] reverse (
+        const column_view& source_column
+    ) except +
+
     cdef unique_ptr[column] shift(
         const column_view& input,
         size_type offset,

python/cudf/cudf/core/column/column.py

@@ -350,6 +350,9 @@ def to_array(self, fillna=None) -> np.ndarray:
 
         return self.to_gpu_array(fillna=fillna).copy_to_host()
 
+    def _reverse(self):
+        return libcudf.copying.reverse(self)
+
     def _fill(
         self,
         fill_value: ScalarLike,

python/cudf/cudf/core/frame.py

@@ -1622,6 +1622,10 @@ def _repeat(self, count):
         result._copy_type_metadata(self)
         return result
 
+    def _reverse(self):
+        result = self.__class__._from_table(libcudf.copying.reverse(self))
+        return result
+
     def _fill(self, fill_values, begin, end, inplace):
         col_and_fill = zip(self._columns, fill_values)