Fix debug build by splitting row_operator_tests_utilities.cu

cpp/tests/CMakeLists.txt

@@ -433,8 +433,13 @@ 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/row_operator_tests_utilities.cu
+  TABLE_TEST
+  table/table_tests.cpp
+  table/table_view_tests.cu
+  table/row_operators_tests.cpp
+  table/experimental_row_operator_tests.cu
+  table/row_operator_tests_utilities.cu
+  table/row_operator_tests_utilities2.cu
 )
 
 # ##################################################################################################

cpp/tests/table/row_operator_tests_utilities.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2022-2023, NVIDIA CORPORATION.
+ * Copyright (c) 2022-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -18,60 +18,14 @@
 
 #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,
@@ -199,9 +153,6 @@ std::unique_ptr<cudf::column> two_table_equality(cudf::table_view lhs,
   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,

cpp/tests/table/row_operator_tests_utilities.hpp

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2023, NVIDIA CORPORATION.
+ * Copyright (c) 2023-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -22,6 +22,12 @@
 
 #include <vector>
 
+using physical_comparator_t = cudf::experimental::row::lexicographic::physical_element_comparator;
+using sorting_comparator_t =
+  cudf::experimental::row::lexicographic::sorting_physical_element_comparator;
+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 <typename PhysicalElementComparator>
 std::unique_ptr<cudf::column> self_comparison(cudf::table_view input,
                                               std::vector<cudf::order> const& column_order,

cpp/tests/table/row_operator_tests_utilities2.cu

@@ -0,0 +1,69 @@
+/*
+ * Copyright (c) 2024, 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 <rmm/cuda_stream_view.hpp>
+#include <rmm/exec_policy.hpp>
+
+#include <thrust/iterator/counting_iterator.h>
+#include <thrust/transform.h>
+
+// Including this declaration/defintion in row_operator_tests_utilities.cu causes
+// the nvcc compiler to segfault when built with the debug (-g) flag.
+
+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;
+}
+
+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);