Fix debug build by splitting row_operator_tests_utilities.cu (#14826)

Splitting up `row_operator_tests_utilities.cu` works around a nvcc segfault when building with Debug. The segfault is found to occur on at least 11.8 and 12.0. ``` Building CUDA object tests/CMakeFiles/TABLE_TEST.dir/table/row_operator_tests_utilities.cu.o FAILED: tests/CMakeFiles/TABLE_TEST.dir/table/row_operator_tests_utilities.cu.o /usr/local/bin/nvcc -forward-unknown-to-host-compiler -DFMT_HEADER_ONLY=1 -DLIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE -DSPDLOG_ACTIVE_LEVEL=SPDLOG_LEVEL_INFO -DSPDLOG_FMT_EXTERNAL -DTHRUST_DEVICE_SYSTEM=THRUST_DEVICE_SYSTEM_CUDA -DTHRUST_HOST_SYSTEM=THRUST_HOST_SYSTEM_CPP -I/cudf/cpp -I/cudf/cpp/src -I/cudf/cpp/build/_deps/dlpack-src/include -I/cudf/cpp/build/_deps/jitify-src -I/cudf/cpp/include -I/cudf/cpp/build/include -I/cudf/cpp/build/_deps/cccl-src/thrust/thrust/cmake/../.. -I/cudf/cpp/build/_deps/cccl-src/libcudacxx/lib/cmake/libcudacxx/../../../include -I/cudf/cpp/build/_deps/cccl-src/cub/cub/cmake/../.. -isystem /conda/envs/rapids/include -isystem /usr/local/cuda/targets/x86_64-linux/include -Xcompiler=-fdiagnostics-color=always -g -std=c++17 "--generate-code=arch=compute_70,code=[sm_70]" -Xcompiler=-fPIE --expt-extended-lambda --expt-relaxed-constexpr -Werror=all-warnings -Xcompiler=-Wall,-Werror,-Wno-error=deprecated-declarations -Xfatbin=-compress-all -Xcompiler=-rdynamic -MD -MT tests/CMakeFiles/TABLE_TEST.dir/table/row_operator_tests_utilities.cu.o -MF tests/CMakeFiles/TABLE_TEST.dir/table/row_operator_tests_utilities.cu.o.d -x cu -c /cudf/cpp/tests/table/row_operator_tests_utilities.cu -o tests/CMakeFiles/TABLE_TEST.dir/table/row_operator_tests_utilities.cu.o Segmentation fault (core dumped) ninja: build stopped: subcommand failed. ``` This PR has been verified to workaround the error in both versions. Authors: - David Wendt (https://github.com/davidwendt) Approvers: - Karthikeyan (https://github.com/karthikeyann) - Bradley Dice (https://github.com/bdice) URL: #14826
rapidsai · Jan 29, 2024 · d5db68e · d5db68e
1 parent b47f5ee
commit d5db68e
Show file tree

Hide file tree

Showing 4 changed files with 84 additions and 53 deletions.
diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt
@@ -439,8 +439,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
 )
 
 # ##################################################################################################

diff --git a/cpp/tests/table/row_operator_tests_utilities.cu b/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,

diff --git a/cpp/tests/table/row_operator_tests_utilities.hpp b/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,

diff --git a/cpp/tests/table/row_operator_tests_utilities2.cu b/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);