From d5db68e018d720094489325d5547a7ed82f22b0d Mon Sep 17 00:00:00 2001 From: David Wendt <45795991+davidwendt@users.noreply.github.com> Date: Mon, 29 Jan 2024 08:47:31 -0500 Subject: [PATCH] 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: https://github.com/rapidsai/cudf/pull/14826 --- cpp/tests/CMakeLists.txt | 9 ++- .../table/row_operator_tests_utilities.cu | 51 +------------- .../table/row_operator_tests_utilities.hpp | 8 ++- .../table/row_operator_tests_utilities2.cu | 69 +++++++++++++++++++ 4 files changed, 84 insertions(+), 53 deletions(-) create mode 100644 cpp/tests/table/row_operator_tests_utilities2.cu diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index caa92d60151..8b0e625fecf 100644 --- 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 index d1f918cc7af..cfffa1cdd54 100644 --- 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 #include -#include #include #include -#include #include #include #include -template -std::unique_ptr self_comparison(cudf::table_view input, - std::vector 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(), - table_comparator.less(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(), - table_comparator.less(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 self_comparison( - cudf::table_view input, - std::vector const& column_order, - physical_comparator_t comparator); -template std::unique_ptr self_comparison( - cudf::table_view input, - std::vector const& column_order, - sorting_comparator_t comparator); - template std::unique_ptr two_table_comparison(cudf::table_view lhs, cudf::table_view rhs, @@ -199,9 +153,6 @@ std::unique_ptr 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 two_table_equality( 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 index b34bf65d176..023a54669b4 100644 --- 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 +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 std::unique_ptr self_comparison(cudf::table_view input, std::vector const& column_order, diff --git a/cpp/tests/table/row_operator_tests_utilities2.cu b/cpp/tests/table/row_operator_tests_utilities2.cu new file mode 100644 index 00000000000..057d9ee1004 --- /dev/null +++ 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 +#include + +#include +#include + +#include +#include + +// Including this declaration/defintion in row_operator_tests_utilities.cu causes +// the nvcc compiler to segfault when built with the debug (-g) flag. + +template +std::unique_ptr self_comparison(cudf::table_view input, + std::vector 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(), + table_comparator.less(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(), + table_comparator.less(cudf::nullate::NO{}, comparator)); + } + return output; +} + +template std::unique_ptr self_comparison( + cudf::table_view input, + std::vector const& column_order, + physical_comparator_t comparator); +template std::unique_ptr self_comparison( + cudf::table_view input, + std::vector const& column_order, + sorting_comparator_t comparator);