From fa8f639651089b9ca2823a1722dd86e18d879406 Mon Sep 17 00:00:00 2001 From: divyegala Date: Sun, 19 Feb 2023 14:56:20 -0800 Subject: [PATCH] trying legacy again --- cpp/benchmarks/CMakeLists.txt | 3 +- cpp/src/join/hash_join.cu | 119 ++++++++++++++++++----------- cpp/src/join/join_common_utils.cuh | 36 ++++++++- 3 files changed, 109 insertions(+), 49 deletions(-) diff --git a/cpp/benchmarks/CMakeLists.txt b/cpp/benchmarks/CMakeLists.txt index c5ae3345da5..2fe1b36281c 100644 --- a/cpp/benchmarks/CMakeLists.txt +++ b/cpp/benchmarks/CMakeLists.txt @@ -155,7 +155,8 @@ ConfigureNVBench( # ################################################################################################## # * join benchmark -------------------------------------------------------------------------------- ConfigureBench(JOIN_BENCH join/left_join.cu join/conditional_join.cu) -ConfigureNVBench(JOIN_NVBENCH join/join.cu join/mixed_join.cu) +ConfigureNVBench(JOIN_NVBENCH join/join.cu) +# join/mixed_join.cu) # ################################################################################################## # * iterator benchmark ---------------------------------------------------------------------------- diff --git a/cpp/src/join/hash_join.cu b/cpp/src/join/hash_join.cu index 6e7ba2a5773..e440a5c4871 100644 --- a/cpp/src/join/hash_join.cu +++ b/cpp/src/join/hash_join.cu @@ -43,6 +43,40 @@ namespace cudf { namespace detail { namespace { + +auto get_legacy_comparator(table_view build_table, + table_view probe_table, + bool const has_nulls, + cudf::null_equality nulls_equal, + rmm::cuda_stream_view stream) { + auto build_table_d = cudf::table_device_view::create(build_table, stream); + auto probe_table_d = cudf::table_device_view::create(probe_table, stream); + + auto probe_nulls = cudf::nullate::DYNAMIC{has_nulls}; + + return row_equality{probe_nulls, *probe_table_d, *build_table_d, nulls_equal}; +} + +auto get_experimental_comparator(table_view build_table, + table_view probe_table, + bool const has_nulls, + cudf::null_equality const nulls_equal, + rmm::cuda_stream_view stream) { + auto const probe_nulls = cudf::nullate::DYNAMIC{has_nulls}; + + auto const preprocessed_probe = + cudf::experimental::row::hash::preprocessed_table::create(probe_table, stream); + + auto const preprocessed_build = + cudf::experimental::row::equality::preprocessed_table::create(build_table, stream); + auto const row_comparator = + cudf::experimental::row::equality::two_table_comparator{preprocessed_probe, preprocessed_build}; + + auto const device_comparator = row_comparator.equal_to(probe_nulls, nulls_equal); + + return device_comparator; +} + /** * @brief Calculates the exact size of the join output produced when * joining two tables together. @@ -88,21 +122,17 @@ std::size_t compute_join_output_size(table_view build_table, auto const probe_nulls = cudf::nullate::DYNAMIC{has_nulls}; - auto const preprocessed_probe = - cudf::experimental::row::hash::preprocessed_table::create(probe_table, stream); - auto row_hash = cudf::experimental::row::hash::row_hasher{preprocessed_probe}; - auto hash_probe = row_hash.device_hasher(probe_nulls); + // auto const preprocessed_probe = + // cudf::experimental::row::hash::preprocessed_table::create(probe_table, stream); + // auto row_hash = cudf::experimental::row::hash::row_hasher{preprocessed_probe}; + // auto hash_probe = row_hash.device_hasher(probe_nulls); + auto probe_table_d = cudf::table_device_view::create(probe_table, stream); + row_hash hash_probe{probe_nulls, *probe_table_d}; + auto const empty_key_sentinel = hash_table.get_empty_key_sentinel(); make_pair_function pair_func{hash_probe, empty_key_sentinel}; - auto const preprocessed_build = - cudf::experimental::row::equality::preprocessed_table::create(build_table, stream); - auto const row_comparator = - cudf::experimental::row::equality::two_table_comparator{preprocessed_probe, preprocessed_build}; - - auto const comparator_helper = [&](auto const device_comparator) { - pair_equality equality{device_comparator}; - + auto const comparator_helper = [&](auto const equality) { std::size_t size; auto iter = cudf::detail::make_counting_transform_iterator(0, pair_func); if constexpr (JoinKind == join_kind::LEFT_JOIN) { @@ -116,11 +146,13 @@ std::size_t compute_join_output_size(table_view build_table, }; if (cudf::detail::has_nested_columns(probe_table)) { - auto const device_comparator = row_comparator.equal_to(has_nulls, nulls_equal); - return comparator_helper(device_comparator); + auto const device_comparator = get_experimental_comparator(build_table, probe_table, has_nulls, nulls_equal, stream); + experimental::pair_equality equality{device_comparator}; + return comparator_helper(equality); } else { - auto const device_comparator = row_comparator.equal_to(has_nulls, nulls_equal); - return comparator_helper(device_comparator); + auto const device_comparator = get_legacy_comparator(build_table, probe_table, has_nulls, nulls_equal, stream); + pair_equality equality(device_comparator); + return comparator_helper(equality); } } @@ -174,19 +206,16 @@ probe_join_hash_table(cudf::table_view build_table, auto const probe_nulls = cudf::nullate::DYNAMIC{has_nulls}; - auto const preprocessed_probe = - cudf::experimental::row::hash::preprocessed_table::create(probe_table, stream); - auto row_hash = cudf::experimental::row::hash::row_hasher{preprocessed_probe}; - auto hash_probe = row_hash.device_hasher(probe_nulls); + // auto const preprocessed_probe = + // cudf::experimental::row::hash::preprocessed_table::create(probe_table, stream); + // auto row_hash = cudf::experimental::row::hash::row_hasher{preprocessed_probe}; + // auto hash_probe = row_hash.device_hasher(probe_nulls); + auto probe_table_d = cudf::table_device_view::create(probe_table, stream); + row_hash hash_probe{probe_nulls, *probe_table_d}; auto const empty_key_sentinel = hash_table.get_empty_key_sentinel(); make_pair_function pair_func{hash_probe, empty_key_sentinel}; - auto const preprocessed_build = - cudf::experimental::row::equality::preprocessed_table::create(build_table, stream); - auto const row_comparator = - cudf::experimental::row::equality::two_table_comparator{preprocessed_probe, preprocessed_build}; - auto const comparator_helper = [&](auto const device_comparator) { - pair_equality equality{device_comparator}; + auto const comparator_helper = [&](auto const equality) { const cudf::size_type probe_table_num_rows = probe_table.num_rows(); @@ -222,11 +251,13 @@ probe_join_hash_table(cudf::table_view build_table, }; if (cudf::detail::has_nested_columns(probe_table)) { - auto device_comparator = row_comparator.equal_to(probe_nulls, compare_nulls); - comparator_helper(device_comparator); + auto const device_comparator = get_experimental_comparator(build_table, probe_table, has_nulls, compare_nulls, stream); + experimental::pair_equality equality{device_comparator}; + comparator_helper(equality); } else { - auto device_comparator = row_comparator.equal_to(probe_nulls, compare_nulls); - comparator_helper(device_comparator); + auto const device_comparator = get_legacy_comparator(build_table, probe_table, has_nulls, compare_nulls, stream); + pair_equality equality{device_comparator}; + comparator_helper(equality); } return std::pair(std::move(left_indices), std::move(right_indices)); @@ -266,19 +297,16 @@ std::size_t get_full_join_size(cudf::table_view build_table, auto const probe_nulls = cudf::nullate::DYNAMIC{has_nulls}; - auto const preprocessed_probe = - cudf::experimental::row::hash::preprocessed_table::create(probe_table, stream); - auto row_hash = cudf::experimental::row::hash::row_hasher{preprocessed_probe}; - auto hash_probe = row_hash.device_hasher(probe_nulls); + // auto const preprocessed_probe = + // cudf::experimental::row::hash::preprocessed_table::create(probe_table, stream); + // auto row_hash = cudf::experimental::row::hash::row_hasher{preprocessed_probe}; + // auto hash_probe = row_hash.device_hasher(probe_nulls); + auto probe_table_d = cudf::table_device_view::create(probe_table, stream); + row_hash hash_probe{probe_nulls, *probe_table_d}; auto const empty_key_sentinel = hash_table.get_empty_key_sentinel(); make_pair_function pair_func{hash_probe, empty_key_sentinel}; - auto const preprocessed_build = - cudf::experimental::row::equality::preprocessed_table::create(build_table, stream); - auto const row_comparator = - cudf::experimental::row::equality::two_table_comparator{preprocessed_probe, preprocessed_build}; - auto const comparator_helper = [&](auto const device_comparator) { - pair_equality equality{device_comparator}; + auto const comparator_helper = [&](auto const equality) { auto iter = cudf::detail::make_counting_transform_iterator(0, pair_func); @@ -293,11 +321,14 @@ std::size_t get_full_join_size(cudf::table_view build_table, iter, iter + probe_table_num_rows, out1_zip_begin, out2_zip_begin, equality, stream.value()); }; if (cudf::detail::has_nested_columns(probe_table)) { - auto const device_comparator = row_comparator.equal_to(probe_nulls, compare_nulls); - comparator_helper(device_comparator); + auto const device_comparator = get_experimental_comparator(build_table, probe_table, has_nulls, compare_nulls, stream); + experimental::pair_equality equality{device_comparator}; + comparator_helper(equality); } else { - auto const device_comparator = row_comparator.equal_to(probe_nulls, compare_nulls); - comparator_helper(device_comparator); + + auto const device_comparator = get_legacy_comparator(build_table, probe_table, has_nulls, compare_nulls, stream); + pair_equality equality{device_comparator}; + comparator_helper(equality); } // Release intermediate memory allocation diff --git a/cpp/src/join/join_common_utils.cuh b/cpp/src/join/join_common_utils.cuh index 5fdc33420f5..42a052ebbf8 100644 --- a/cpp/src/join/join_common_utils.cuh +++ b/cpp/src/join/join_common_utils.cuh @@ -101,6 +101,31 @@ class row_is_valid { * * @tparam Comparator The row comparator type to perform row equality comparison from row indices. */ + + template + class pair_equality { + public: + pair_equality(table_device_view lhs, + table_device_view rhs, + nullate::DYNAMIC has_nulls, + null_equality nulls_are_equal = null_equality::EQUAL) + : _check_row_equality{has_nulls, lhs, rhs, nulls_are_equal} + { + } + + pair_equality(Comparator const d_eqcomp) : _check_row_equality{std::move(d_eqcomp)} {} + + template + __device__ __forceinline__ bool operator()(LhsPair const& lhs, RhsPair const& rhs) const noexcept + { + return lhs.first == rhs.first and _check_row_equality(rhs.second, lhs.second); + } + + private: + Comparator _check_row_equality; + }; + +namespace experimental { template class pair_equality { public: @@ -109,8 +134,8 @@ class pair_equality { template __device__ __forceinline__ bool operator()(LhsPair const& lhs, RhsPair const& rhs) const noexcept { - using experimental::row::lhs_index_type; - using experimental::row::rhs_index_type; + using cudf::experimental::row::lhs_index_type; + using cudf::experimental::row::rhs_index_type; // printf("lhs_index: %d, rhs_index: %d, hash: %d, equality: %d\n", lhs.second, rhs.second, // lhs.first == rhs.first, _check_row_equality(rhs_index_type{rhs.second}, // lhs_index_type{lhs.second})); @@ -121,6 +146,7 @@ class pair_equality { private: DeviceComparator _check_row_equality; }; +} /** * @brief Computes the trivial left join operation for the case when the @@ -165,8 +191,10 @@ void build_join_hash_table(cudf::table_view const& build, CUDF_EXPECTS(0 != build.num_columns(), "Selected build dataset is empty"); CUDF_EXPECTS(0 != build.num_rows(), "Build side table has no rows"); - auto row_hash = experimental::row::hash::row_hasher{build, stream}; - auto hash_build = row_hash.device_hasher(nullate::DYNAMIC{cudf::has_nested_nulls(build)}); + auto build_table_ptr = cudf::table_device_view::create(build, stream); + row_hash hash_build{nullate::DYNAMIC{cudf::has_nulls(build)}, *build_table_ptr}; + // auto row_hash = cudf::experimental::row::hash::row_hasher{build, stream}; + // auto hash_build = row_hash.device_hasher(nullate::DYNAMIC{cudf::has_nested_nulls(build)}); auto const empty_key_sentinel = hash_table.get_empty_key_sentinel(); make_pair_function pair_func{hash_build, empty_key_sentinel};