From 43c04e2021f3b1df65e5c585becbef1172c03c84 Mon Sep 17 00:00:00 2001
From: David Wendt <45795991+davidwendt@users.noreply.github.com>
Date: Wed, 24 May 2023 10:36:06 -0400
Subject: [PATCH] Split up experimental_row_operator_tests.cu to improve its
 compile time (#13382)

Noticed the `experimental_row_operator_tests.cu` build time increased significantly in the last few weeks.
This PR splits the thrust calls into a separate .cu file. The remaining code is C++ but could not be converted to .cpp because the `row_operators.cuh` files contain device code that g++ cannot compile.

Authors:
  - David Wendt (https://github.com/davidwendt)

Approvers:
  - Nghia Truong (https://github.com/ttnghia)
  - Vyas Ramasubramani (https://github.com/vyasr)

URL: https://github.com/rapidsai/cudf/pull/13382
---
 cpp/tests/CMakeLists.txt                      |   2 +-
 .../table/experimental_row_operator_tests.cu  | 191 ++--------------
 .../table/row_operator_tests_utilities.cu     | 214 ++++++++++++++++++
 .../table/row_operator_tests_utilities.hpp    |  45 ++++
 4 files changed, 275 insertions(+), 177 deletions(-)
 create mode 100644 cpp/tests/table/row_operator_tests_utilities.cu
 create mode 100644 cpp/tests/table/row_operator_tests_utilities.hpp

diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt
index 1262e065041..c68dd3b4ffb 100644
--- a/cpp/tests/CMakeLists.txt
+++ b/cpp/tests/CMakeLists.txt
@@ -385,7 +385,7 @@ ConfigureTest(
 # * table tests -----------------------------------------------------------------------------------
 ConfigureTest(
   TABLE_TEST table/table_tests.cpp table/table_view_tests.cu table/row_operators_tests.cpp
-  table/experimental_row_operator_tests.cu
+  table/experimental_row_operator_tests.cu table/row_operator_tests_utilities.cu
 )
 
 # ##################################################################################################
diff --git a/cpp/tests/table/experimental_row_operator_tests.cu b/cpp/tests/table/experimental_row_operator_tests.cu
index 5ae1c7d9729..896cc7a82d4 100644
--- a/cpp/tests/table/experimental_row_operator_tests.cu
+++ b/cpp/tests/table/experimental_row_operator_tests.cu
@@ -14,28 +14,16 @@
  * limitations under the License.
  */
 
+#include "row_operator_tests_utilities.hpp"
+
 #include <cudf_test/base_fixture.hpp>
 #include <cudf_test/column_utilities.hpp>
 #include <cudf_test/column_wrapper.hpp>
 #include <cudf_test/type_lists.hpp>
 
-#include <cudf/column/column_factories.hpp>
-#include <cudf/column/column_view.hpp>
 #include <cudf/table/experimental/row_operators.cuh>
-#include <cudf/table/row_operators.cuh>
-#include <cudf/table/table_view.hpp>
-#include <cudf/utilities/default_stream.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
-#include <rmm/exec_policy.hpp>
-
-#include <thrust/iterator/counting_iterator.h>
-#include <thrust/sequence.h>
-#include <thrust/sort.h>
-#include <thrust/transform.h>
-
-#include <cmath>
-#include <vector>
 
 template <typename T>
 struct TypedTableViewTest : public cudf::test::BaseFixture {};
@@ -45,175 +33,26 @@ using NumericTypesNotBool =
 TYPED_TEST_SUITE(TypedTableViewTest, NumericTypesNotBool);
 
 template <typename PhysicalElementComparator>
-auto self_comparison(cudf::table_view input,
-                     std::vector<cudf::order> const& column_order,
-                     PhysicalElementComparator comparator)
-{
-  rmm::cuda_stream_view stream{cudf::get_default_stream()};
-
-  auto const table_comparator =
-    cudf::experimental::row::lexicographic::self_comparator{input, column_order, {}, stream};
-
-  auto output = cudf::make_numeric_column(
-    cudf::data_type(cudf::type_id::BOOL8), input.num_rows(), cudf::mask_state::UNALLOCATED);
-
-  if (cudf::detail::has_nested_columns(input)) {
-    thrust::transform(rmm::exec_policy(stream),
-                      thrust::make_counting_iterator(0),
-                      thrust::make_counting_iterator(input.num_rows()),
-                      thrust::make_counting_iterator(0),
-                      output->mutable_view().data<bool>(),
-                      table_comparator.less<true>(cudf::nullate::NO{}, comparator));
-  } else {
-    thrust::transform(rmm::exec_policy(stream),
-                      thrust::make_counting_iterator(0),
-                      thrust::make_counting_iterator(input.num_rows()),
-                      thrust::make_counting_iterator(0),
-                      output->mutable_view().data<bool>(),
-                      table_comparator.less<false>(cudf::nullate::NO{}, comparator));
-  }
-  return output;
-}
-
+std::unique_ptr<cudf::column> self_comparison(cudf::table_view input,
+                                              std::vector<cudf::order> const& column_order,
+                                              PhysicalElementComparator comparator);
 template <typename PhysicalElementComparator>
-auto two_table_comparison(cudf::table_view lhs,
-                          cudf::table_view rhs,
-                          std::vector<cudf::order> const& column_order,
-                          PhysicalElementComparator comparator)
-{
-  rmm::cuda_stream_view stream{cudf::get_default_stream()};
-
-  auto const table_comparator = cudf::experimental::row::lexicographic::two_table_comparator{
-    lhs, rhs, column_order, {}, stream};
-  auto const lhs_it = cudf::experimental::row::lhs_iterator(0);
-  auto const rhs_it = cudf::experimental::row::rhs_iterator(0);
-
-  auto output = cudf::make_numeric_column(
-    cudf::data_type(cudf::type_id::BOOL8), lhs.num_rows(), cudf::mask_state::UNALLOCATED);
-
-  if (cudf::detail::has_nested_columns(lhs) || cudf::detail::has_nested_columns(rhs)) {
-    thrust::transform(rmm::exec_policy(stream),
-                      lhs_it,
-                      lhs_it + lhs.num_rows(),
-                      rhs_it,
-                      output->mutable_view().data<bool>(),
-                      table_comparator.less<true>(cudf::nullate::NO{}, comparator));
-  } else {
-    thrust::transform(rmm::exec_policy(stream),
-                      lhs_it,
-                      lhs_it + lhs.num_rows(),
-                      rhs_it,
-                      output->mutable_view().data<bool>(),
-                      table_comparator.less<false>(cudf::nullate::NO{}, comparator));
-  }
-  return output;
-}
-
+std::unique_ptr<cudf::column> two_table_comparison(cudf::table_view lhs,
+                                                   cudf::table_view rhs,
+                                                   std::vector<cudf::order> const& column_order,
+                                                   PhysicalElementComparator comparator);
 template <typename PhysicalElementComparator>
-auto self_equality(cudf::table_view input,
-                   std::vector<cudf::order> const& column_order,
-                   PhysicalElementComparator comparator)
-{
-  rmm::cuda_stream_view stream{cudf::get_default_stream()};
-
-  auto const table_comparator = cudf::experimental::row::equality::self_comparator{input, stream};
-
-  auto output = cudf::make_numeric_column(
-    cudf::data_type(cudf::type_id::BOOL8), input.num_rows(), cudf::mask_state::UNALLOCATED);
-
-  if (cudf::detail::has_nested_columns(input)) {
-    auto const equal_comparator =
-      table_comparator.equal_to<true>(cudf::nullate::NO{}, cudf::null_equality::EQUAL, comparator);
-
-    thrust::transform(rmm::exec_policy(stream),
-                      thrust::make_counting_iterator(0),
-                      thrust::make_counting_iterator(input.num_rows()),
-                      thrust::make_counting_iterator(0),
-                      output->mutable_view().data<bool>(),
-                      equal_comparator);
-  } else {
-    auto const equal_comparator =
-      table_comparator.equal_to<false>(cudf::nullate::NO{}, cudf::null_equality::EQUAL, comparator);
-
-    thrust::transform(rmm::exec_policy(stream),
-                      thrust::make_counting_iterator(0),
-                      thrust::make_counting_iterator(input.num_rows()),
-                      thrust::make_counting_iterator(0),
-                      output->mutable_view().data<bool>(),
-                      equal_comparator);
-  }
-
-  return output;
-}
-
-template <typename PhysicalElementComparator>
-auto two_table_equality(cudf::table_view lhs,
-                        cudf::table_view rhs,
-                        std::vector<cudf::order> const& column_order,
-                        PhysicalElementComparator comparator)
-{
-  rmm::cuda_stream_view stream{cudf::get_default_stream()};
-
-  auto const table_comparator =
-    cudf::experimental::row::equality::two_table_comparator{lhs, rhs, stream};
-
-  auto const lhs_it = cudf::experimental::row::lhs_iterator(0);
-  auto const rhs_it = cudf::experimental::row::rhs_iterator(0);
-
-  auto output = cudf::make_numeric_column(
-    cudf::data_type(cudf::type_id::BOOL8), lhs.num_rows(), cudf::mask_state::UNALLOCATED);
-
-  if (cudf::detail::has_nested_columns(lhs) or cudf::detail::has_nested_columns(rhs)) {
-    auto const equal_comparator =
-      table_comparator.equal_to<true>(cudf::nullate::NO{}, cudf::null_equality::EQUAL, comparator);
-
-    thrust::transform(rmm::exec_policy(stream),
-                      lhs_it,
-                      lhs_it + lhs.num_rows(),
-                      rhs_it,
-                      output->mutable_view().data<bool>(),
-                      equal_comparator);
-  } else {
-    auto const equal_comparator =
-      table_comparator.equal_to<false>(cudf::nullate::NO{}, cudf::null_equality::EQUAL, comparator);
-
-    thrust::transform(rmm::exec_policy(stream),
-                      lhs_it,
-                      lhs_it + lhs.num_rows(),
-                      rhs_it,
-                      output->mutable_view().data<bool>(),
-                      equal_comparator);
-  }
-  return output;
-}
-
+std::unique_ptr<cudf::column> two_table_equality(cudf::table_view lhs,
+                                                 cudf::table_view rhs,
+                                                 std::vector<cudf::order> const& column_order,
+                                                 PhysicalElementComparator comparator);
 template <typename PhysicalElementComparator>
-auto sorted_order(
+std::unique_ptr<cudf::column> sorted_order(
   std::shared_ptr<cudf::experimental::row::lexicographic::preprocessed_table> preprocessed_input,
   cudf::size_type num_rows,
   bool has_nested,
   PhysicalElementComparator comparator,
-  rmm::cuda_stream_view stream)
-{
-  auto output = cudf::make_numeric_column(cudf::data_type(cudf::type_to_id<cudf::size_type>()),
-                                          num_rows,
-                                          cudf::mask_state::UNALLOCATED,
-                                          stream);
-  auto const out_begin = output->mutable_view().begin<cudf::size_type>();
-  thrust::sequence(rmm::exec_policy(stream), out_begin, out_begin + num_rows, 0);
-
-  auto const table_comparator =
-    cudf::experimental::row::lexicographic::self_comparator{preprocessed_input};
-  if (has_nested) {
-    auto const comp = table_comparator.less<true>(cudf::nullate::NO{}, comparator);
-    thrust::stable_sort(rmm::exec_policy(stream), out_begin, out_begin + num_rows, comp);
-  } else {
-    auto const comp = table_comparator.less<false>(cudf::nullate::NO{}, comparator);
-    thrust::stable_sort(rmm::exec_policy(stream), out_begin, out_begin + num_rows, comp);
-  }
-
-  return output;
-}
+  rmm::cuda_stream_view stream);
 
 TYPED_TEST(TypedTableViewTest, TestLexicographicalComparatorTwoTables)
 {
diff --git a/cpp/tests/table/row_operator_tests_utilities.cu b/cpp/tests/table/row_operator_tests_utilities.cu
new file mode 100644
index 00000000000..d1f918cc7af
--- /dev/null
+++ b/cpp/tests/table/row_operator_tests_utilities.cu
@@ -0,0 +1,214 @@
+/*
+ * Copyright (c) 2022-2023, 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 "row_operator_tests_utilities.hpp"
+
+#include <cudf/column/column_factories.hpp>
+#include <cudf/column/column_view.hpp>
+#include <cudf/table/experimental/row_operators.cuh>
+
+#include <rmm/cuda_stream_view.hpp>
+#include <rmm/exec_policy.hpp>
+
+#include <thrust/iterator/counting_iterator.h>
+#include <thrust/sequence.h>
+#include <thrust/sort.h>
+#include <thrust/transform.h>
+
+template <typename PhysicalElementComparator>
+std::unique_ptr<cudf::column> self_comparison(cudf::table_view input,
+                                              std::vector<cudf::order> const& column_order,
+                                              PhysicalElementComparator comparator)
+{
+  rmm::cuda_stream_view stream{cudf::get_default_stream()};
+
+  auto const table_comparator =
+    cudf::experimental::row::lexicographic::self_comparator{input, column_order, {}, stream};
+
+  auto output = cudf::make_numeric_column(
+    cudf::data_type(cudf::type_id::BOOL8), input.num_rows(), cudf::mask_state::UNALLOCATED);
+
+  if (cudf::detail::has_nested_columns(input)) {
+    thrust::transform(rmm::exec_policy(stream),
+                      thrust::make_counting_iterator(0),
+                      thrust::make_counting_iterator(input.num_rows()),
+                      thrust::make_counting_iterator(0),
+                      output->mutable_view().data<bool>(),
+                      table_comparator.less<true>(cudf::nullate::NO{}, comparator));
+  } else {
+    thrust::transform(rmm::exec_policy(stream),
+                      thrust::make_counting_iterator(0),
+                      thrust::make_counting_iterator(input.num_rows()),
+                      thrust::make_counting_iterator(0),
+                      output->mutable_view().data<bool>(),
+                      table_comparator.less<false>(cudf::nullate::NO{}, comparator));
+  }
+  return output;
+}
+
+using physical_comparator_t = cudf::experimental::row::lexicographic::physical_element_comparator;
+using sorting_comparator_t =
+  cudf::experimental::row::lexicographic::sorting_physical_element_comparator;
+
+template std::unique_ptr<cudf::column> self_comparison<physical_comparator_t>(
+  cudf::table_view input,
+  std::vector<cudf::order> const& column_order,
+  physical_comparator_t comparator);
+template std::unique_ptr<cudf::column> self_comparison<sorting_comparator_t>(
+  cudf::table_view input,
+  std::vector<cudf::order> const& column_order,
+  sorting_comparator_t comparator);
+
+template <typename PhysicalElementComparator>
+std::unique_ptr<cudf::column> two_table_comparison(cudf::table_view lhs,
+                                                   cudf::table_view rhs,
+                                                   std::vector<cudf::order> const& column_order,
+                                                   PhysicalElementComparator comparator)
+{
+  rmm::cuda_stream_view stream{cudf::get_default_stream()};
+
+  auto const table_comparator = cudf::experimental::row::lexicographic::two_table_comparator{
+    lhs, rhs, column_order, {}, stream};
+  auto const lhs_it = cudf::experimental::row::lhs_iterator(0);
+  auto const rhs_it = cudf::experimental::row::rhs_iterator(0);
+
+  auto output = cudf::make_numeric_column(
+    cudf::data_type(cudf::type_id::BOOL8), lhs.num_rows(), cudf::mask_state::UNALLOCATED);
+
+  if (cudf::detail::has_nested_columns(lhs) || cudf::detail::has_nested_columns(rhs)) {
+    thrust::transform(rmm::exec_policy(stream),
+                      lhs_it,
+                      lhs_it + lhs.num_rows(),
+                      rhs_it,
+                      output->mutable_view().data<bool>(),
+                      table_comparator.less<true>(cudf::nullate::NO{}, comparator));
+  } else {
+    thrust::transform(rmm::exec_policy(stream),
+                      lhs_it,
+                      lhs_it + lhs.num_rows(),
+                      rhs_it,
+                      output->mutable_view().data<bool>(),
+                      table_comparator.less<false>(cudf::nullate::NO{}, comparator));
+  }
+  return output;
+}
+
+template std::unique_ptr<cudf::column> two_table_comparison<physical_comparator_t>(
+  cudf::table_view lhs,
+  cudf::table_view rhs,
+  std::vector<cudf::order> const& column_order,
+  physical_comparator_t comparator);
+template std::unique_ptr<cudf::column> two_table_comparison<sorting_comparator_t>(
+  cudf::table_view lhs,
+  cudf::table_view rhs,
+  std::vector<cudf::order> const& column_order,
+  sorting_comparator_t comparator);
+
+template <typename PhysicalElementComparator>
+std::unique_ptr<cudf::column> sorted_order(
+  std::shared_ptr<cudf::experimental::row::lexicographic::preprocessed_table> preprocessed_input,
+  cudf::size_type num_rows,
+  bool has_nested,
+  PhysicalElementComparator comparator,
+  rmm::cuda_stream_view stream)
+{
+  auto output = cudf::make_numeric_column(cudf::data_type(cudf::type_to_id<cudf::size_type>()),
+                                          num_rows,
+                                          cudf::mask_state::UNALLOCATED,
+                                          stream);
+  auto const out_begin = output->mutable_view().begin<cudf::size_type>();
+  thrust::sequence(rmm::exec_policy(stream), out_begin, out_begin + num_rows, 0);
+
+  auto const table_comparator =
+    cudf::experimental::row::lexicographic::self_comparator{preprocessed_input};
+  if (has_nested) {
+    auto const comp = table_comparator.less<true>(cudf::nullate::NO{}, comparator);
+    thrust::stable_sort(rmm::exec_policy(stream), out_begin, out_begin + num_rows, comp);
+  } else {
+    auto const comp = table_comparator.less<false>(cudf::nullate::NO{}, comparator);
+    thrust::stable_sort(rmm::exec_policy(stream), out_begin, out_begin + num_rows, comp);
+  }
+
+  return output;
+}
+
+template std::unique_ptr<cudf::column> sorted_order<physical_comparator_t>(
+  std::shared_ptr<cudf::experimental::row::lexicographic::preprocessed_table> preprocessed_input,
+  cudf::size_type num_rows,
+  bool has_nested,
+  physical_comparator_t comparator,
+  rmm::cuda_stream_view stream);
+template std::unique_ptr<cudf::column> sorted_order<sorting_comparator_t>(
+  std::shared_ptr<cudf::experimental::row::lexicographic::preprocessed_table> preprocessed_input,
+  cudf::size_type num_rows,
+  bool has_nested,
+  sorting_comparator_t comparator,
+  rmm::cuda_stream_view stream);
+
+template <typename PhysicalElementComparator>
+std::unique_ptr<cudf::column> two_table_equality(cudf::table_view lhs,
+                                                 cudf::table_view rhs,
+                                                 std::vector<cudf::order> const& column_order,
+                                                 PhysicalElementComparator comparator)
+{
+  rmm::cuda_stream_view stream{cudf::get_default_stream()};
+
+  auto const table_comparator =
+    cudf::experimental::row::equality::two_table_comparator{lhs, rhs, stream};
+
+  auto const lhs_it = cudf::experimental::row::lhs_iterator(0);
+  auto const rhs_it = cudf::experimental::row::rhs_iterator(0);
+
+  auto output = cudf::make_numeric_column(
+    cudf::data_type(cudf::type_id::BOOL8), lhs.num_rows(), cudf::mask_state::UNALLOCATED);
+
+  if (cudf::detail::has_nested_columns(lhs) or cudf::detail::has_nested_columns(rhs)) {
+    auto const equal_comparator =
+      table_comparator.equal_to<true>(cudf::nullate::NO{}, cudf::null_equality::EQUAL, comparator);
+
+    thrust::transform(rmm::exec_policy(stream),
+                      lhs_it,
+                      lhs_it + lhs.num_rows(),
+                      rhs_it,
+                      output->mutable_view().data<bool>(),
+                      equal_comparator);
+  } else {
+    auto const equal_comparator =
+      table_comparator.equal_to<false>(cudf::nullate::NO{}, cudf::null_equality::EQUAL, comparator);
+
+    thrust::transform(rmm::exec_policy(stream),
+                      lhs_it,
+                      lhs_it + lhs.num_rows(),
+                      rhs_it,
+                      output->mutable_view().data<bool>(),
+                      equal_comparator);
+  }
+  return output;
+}
+
+using physical_equality_t = cudf::experimental::row::equality::physical_equality_comparator;
+using nan_equality_t = cudf::experimental::row::equality::nan_equal_physical_equality_comparator;
+
+template std::unique_ptr<cudf::column> two_table_equality<physical_equality_t>(
+  cudf::table_view lhs,
+  cudf::table_view rhs,
+  std::vector<cudf::order> const& column_order,
+  physical_equality_t comparator);
+template std::unique_ptr<cudf::column> two_table_equality<nan_equality_t>(
+  cudf::table_view lhs,
+  cudf::table_view rhs,
+  std::vector<cudf::order> const& column_order,
+  nan_equality_t comparator);
diff --git a/cpp/tests/table/row_operator_tests_utilities.hpp b/cpp/tests/table/row_operator_tests_utilities.hpp
new file mode 100644
index 00000000000..b34bf65d176
--- /dev/null
+++ b/cpp/tests/table/row_operator_tests_utilities.hpp
@@ -0,0 +1,45 @@
+/*
+ * Copyright (c) 2023, 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.
+ */
+#pragma once
+
+#include <cudf/table/experimental/row_operators.cuh>
+#include <cudf/table/table_view.hpp>
+
+#include <rmm/cuda_stream_view.hpp>
+
+#include <vector>
+
+template <typename PhysicalElementComparator>
+std::unique_ptr<cudf::column> self_comparison(cudf::table_view input,
+                                              std::vector<cudf::order> const& column_order,
+                                              PhysicalElementComparator comparator);
+template <typename PhysicalElementComparator>
+std::unique_ptr<cudf::column> two_table_comparison(cudf::table_view lhs,
+                                                   cudf::table_view rhs,
+                                                   std::vector<cudf::order> const& column_order,
+                                                   PhysicalElementComparator comparator);
+template <typename PhysicalElementComparator>
+std::unique_ptr<cudf::column> two_table_equality(cudf::table_view lhs,
+                                                 cudf::table_view rhs,
+                                                 std::vector<cudf::order> const& column_order,
+                                                 PhysicalElementComparator comparator);
+template <typename PhysicalElementComparator>
+std::unique_ptr<cudf::column> sorted_order(
+  std::shared_ptr<cudf::experimental::row::lexicographic::preprocessed_table> preprocessed_input,
+  cudf::size_type num_rows,
+  bool has_nested,
+  PhysicalElementComparator comparator,
+  rmm::cuda_stream_view stream);