Skip to content

Commit

Permalink
Split up experimental_row_operator_tests.cu to improve its compile ti…
Browse files Browse the repository at this point in the history 
…me (#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: #13382
  • Loading branch information
@davidwendt
davidwendt authored May 24, 2023
1 parent 707a480 commit 43c04e2
Show file tree
Hide file tree
Showing 4 changed files with 275 additions and 177 deletions.
2 changes: 1 addition & 1 deletion cpp/tests/CMakeLists.txt
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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
)

# ##################################################################################################
Expand Down
191 changes: 15 additions & 176 deletions cpp/tests/table/experimental_row_operator_tests.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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 {};
Expand All @@ -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)
{
Expand Down
Loading

0 comments on commit 43c04e2

Please sign in to comment.