From 70bba0742f2ef36944c06e293e8f8d9f098171b3 Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Thu, 21 Jul 2022 09:00:18 -0700 Subject: [PATCH 1/5] Complete refactor --- cpp/src/search/contains_table.cu | 168 ++++++++++++++++++++++--------- 1 file changed, 118 insertions(+), 50 deletions(-) diff --git a/cpp/src/search/contains_table.cu b/cpp/src/search/contains_table.cu index 46280d4ff5f..fa823e3bf29 100644 --- a/cpp/src/search/contains_table.cu +++ b/cpp/src/search/contains_table.cu @@ -22,12 +22,12 @@ #include #include +#include + #include #include -#include - -#include +#include namespace cudf::detail { @@ -36,6 +36,66 @@ namespace { using cudf::experimental::row::lhs_index_type; using cudf::experimental::row::rhs_index_type; +/** + * @brief Functor to remap hash values to new values if they are equal to the specified sentinel + * value. + */ +template +struct remap_sentinel_hasher { + remap_sentinel_hasher(Hasher&& row_hasher, hash_value_type const sentinel) + : _hasher{std::move(row_hasher)}, _sentinel(sentinel) + { + } + + template + __device__ inline auto operator()(T const idx) const noexcept + { + return remap_sentinel_hash(_hasher(static_cast(idx)), _sentinel); + } + + private: + Hasher _hasher; + hash_value_type const _sentinel; +}; + +/** + * @brief A comparator adapter so that the underlying self comparator can work with strong index + * types. + */ +template +struct strong_index_self_comparator_adapter { + strong_index_self_comparator_adapter(Comparator const& comparator) : _comparator{comparator} {} + + template + __device__ inline auto operator()(T const lhs_index, T const rhs_index) const noexcept + { + return _comparator(static_cast(lhs_index), static_cast(rhs_index)); + } + + private: + Comparator const _comparator; +}; + +/** + * @brief Invoke an `operator()` template with a row equality comparator based on the specified + * `compare_nans` parameter. + * + * @param compare_nans The flag to specify whether NaNs should be compared equal or not + * @param func The input functor to invoke + */ +template +void dispatch_nan_comparator(nan_equality compare_nans, Func&& func) +{ + if (compare_nans == nan_equality::ALL_EQUAL) { + using nan_equal_comparator = + cudf::experimental::row::equality::nan_equal_physical_equality_comparator; + func(nan_equal_comparator{}); + } else { + using nan_unequal_comparator = cudf::experimental::row::equality::physical_equality_comparator; + func(nan_unequal_comparator{}); + } +} + } // namespace rmm::device_uvector contains(table_view const& haystack, @@ -45,20 +105,16 @@ rmm::device_uvector contains(table_view const& haystack, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - // Use a hash map with key type is row hash values and map value type is `lhs_index_type` to store - // all row indices of the haystack table. - using static_multimap = - cuco::static_multimap>, - cuco::double_hashing>; - - auto map = static_multimap(compute_hash_table_size(haystack.num_rows()), - cuco::sentinel::empty_key{std::numeric_limits::max()}, - cuco::sentinel::empty_value{lhs_index_type{detail::JoinNoneValue}}, - stream.value(), - detail::hash_table_allocator_type{default_allocator{}, stream}); + using static_map = cuco::static_map>>; + + auto map = static_map(compute_hash_table_size(haystack.num_rows()), + cuco::sentinel::empty_key{std::numeric_limits::max()}, + cuco::sentinel::empty_value{detail::JoinNoneValue}, + detail::hash_table_allocator_type{default_allocator{}, stream}, + stream.value()); auto const haystack_has_nulls = has_nested_nulls(haystack); auto const needles_has_nulls = has_nested_nulls(needles); @@ -66,13 +122,16 @@ rmm::device_uvector contains(table_view const& haystack, // Insert all row hash values and indices of the haystack table. { - auto const hasher = cudf::experimental::row::hash::row_hasher(haystack, stream); - auto const d_hasher = hasher.device_hasher(nullate::DYNAMIC{has_any_nulls}); + auto const haystack_it = cudf::detail::make_counting_transform_iterator( + size_type{0}, + [] __device__(auto const idx) { return cuco::make_pair(lhs_index_type{idx}, 0); }); - using make_pair_fn = make_pair_function; + auto const hasher = cudf::experimental::row::hash::row_hasher(haystack, stream); + auto const d_hasher = + remap_sentinel_hasher(hasher.device_hasher(nullate::DYNAMIC{has_any_nulls}), + static_cast(map.get_empty_key_sentinel())); - auto const haystack_it = cudf::detail::make_counting_transform_iterator( - size_type{0}, make_pair_fn{d_hasher, map.get_empty_key_sentinel()}); + auto const comparator = cudf::experimental::row::equality::self_comparator(haystack, stream); // If the haystack table has nulls but they are compared unequal, don't insert them. // Otherwise, it was known to cause performance issue: @@ -95,13 +154,29 @@ rmm::device_uvector contains(table_view const& haystack, : haystack_nullable_columns.front().null_mask(); // Insert only rows that do not have any null at any level. - map.insert_if(haystack_it, - haystack_it + haystack.num_rows(), - thrust::counting_iterator(0), // stencil - row_is_valid{row_bitmask_ptr}, - stream.value()); - } else { - map.insert(haystack_it, haystack_it + haystack.num_rows(), stream.value()); + auto const insert_map = [&](auto const value_comp) { + auto const d_eqcomp = strong_index_self_comparator_adapter{ + comparator.equal_to(nullate::DYNAMIC{haystack_has_nulls}, compare_nulls, value_comp)}; + map.insert_if(haystack_it, + haystack_it + haystack.num_rows(), + thrust::counting_iterator(0), // stencil + row_is_valid{row_bitmask_ptr}, + d_hasher, + d_eqcomp, + stream.value()); + }; + + dispatch_nan_comparator(compare_nans, insert_map); + + } else { // haystack_doesn't_have_nulls || compare_nulls == null_equality::EQUAL + auto const insert_map = [&](auto const value_comp) { + auto const d_eqcomp = strong_index_self_comparator_adapter{ + comparator.equal_to(nullate::DYNAMIC{haystack_has_nulls}, compare_nulls, value_comp)}; + map.insert( + haystack_it, haystack_it + haystack.num_rows(), d_hasher, d_eqcomp, stream.value()); + }; + + dispatch_nan_comparator(compare_nans, insert_map); } } @@ -110,36 +185,29 @@ rmm::device_uvector contains(table_view const& haystack, // Check existence for each row of the needles table in the haystack table. { - auto const hasher = cudf::experimental::row::hash::row_hasher(needles, stream); - auto const d_hasher = hasher.device_hasher(nullate::DYNAMIC{has_any_nulls}); + auto const needles_it = cudf::detail::make_counting_transform_iterator( + size_type{0}, [] __device__(auto const idx) { return rhs_index_type{idx}; }); + + auto const hasher = cudf::experimental::row::hash::row_hasher(needles, stream); + auto const d_hasher = + remap_sentinel_hasher(hasher.device_hasher(nullate::DYNAMIC{has_any_nulls}), + static_cast(map.get_empty_key_sentinel())); auto const comparator = cudf::experimental::row::equality::two_table_comparator(haystack, needles, stream); - using make_pair_fn = make_pair_function; - - auto const needles_it = cudf::detail::make_counting_transform_iterator( - size_type{0}, make_pair_fn{d_hasher, map.get_empty_key_sentinel()}); - auto const check_contains = [&](auto const value_comp) { auto const d_eqcomp = comparator.equal_to(nullate::DYNAMIC{has_any_nulls}, compare_nulls, value_comp); - map.pair_contains(needles_it, - needles_it + needles.num_rows(), - contained.begin(), - pair_equality{d_eqcomp}, - stream.value()); + map.contains(needles_it, + needles_it + needles.num_rows(), + contained.begin(), + d_hasher, + d_eqcomp, + stream.value()); }; - if (compare_nans == nan_equality::ALL_EQUAL) { - using nan_equal_comparator = - cudf::experimental::row::equality::nan_equal_physical_equality_comparator; - check_contains(nan_equal_comparator{}); - } else { - using nan_unequal_comparator = - cudf::experimental::row::equality::physical_equality_comparator; - check_contains(nan_unequal_comparator{}); - } + dispatch_nan_comparator(compare_nans, check_contains); } return contained; From e2f0a3a12ed37580915abdaa2a82a42a9517556b Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Thu, 21 Jul 2022 11:49:04 -0700 Subject: [PATCH 2/5] Fix type in `numeric_limits` --- cpp/src/search/contains_table.cu | 11 ++++++----- 1 file changed, 6 insertions(+), 5 deletions(-) diff --git a/cpp/src/search/contains_table.cu b/cpp/src/search/contains_table.cu index fa823e3bf29..384cdc606c6 100644 --- a/cpp/src/search/contains_table.cu +++ b/cpp/src/search/contains_table.cu @@ -110,11 +110,12 @@ rmm::device_uvector contains(table_view const& haystack, cuda::thread_scope_device, rmm::mr::stream_allocator_adaptor>>; - auto map = static_map(compute_hash_table_size(haystack.num_rows()), - cuco::sentinel::empty_key{std::numeric_limits::max()}, - cuco::sentinel::empty_value{detail::JoinNoneValue}, - detail::hash_table_allocator_type{default_allocator{}, stream}, - stream.value()); + auto map = + static_map(compute_hash_table_size(haystack.num_rows()), + cuco::sentinel::empty_key{lhs_index_type{std::numeric_limits::max()}}, + cuco::sentinel::empty_value{detail::JoinNoneValue}, + detail::hash_table_allocator_type{default_allocator{}, stream}, + stream.value()); auto const haystack_has_nulls = has_nested_nulls(haystack); auto const needles_has_nulls = has_nested_nulls(needles); From 0b7f6bfe2bebcf53c4c8ffaf083a55826964d952 Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Thu, 21 Jul 2022 12:14:55 -0700 Subject: [PATCH 3/5] Remove remap sentinel --- cpp/src/search/contains_table.cu | 22 +++++++--------------- 1 file changed, 7 insertions(+), 15 deletions(-) diff --git a/cpp/src/search/contains_table.cu b/cpp/src/search/contains_table.cu index 384cdc606c6..d1ec714aebd 100644 --- a/cpp/src/search/contains_table.cu +++ b/cpp/src/search/contains_table.cu @@ -37,30 +37,24 @@ using cudf::experimental::row::lhs_index_type; using cudf::experimental::row::rhs_index_type; /** - * @brief Functor to remap hash values to new values if they are equal to the specified sentinel - * value. + * @brief An adapter functor to support strong index types for row hasher. */ template -struct remap_sentinel_hasher { - remap_sentinel_hasher(Hasher&& row_hasher, hash_value_type const sentinel) - : _hasher{std::move(row_hasher)}, _sentinel(sentinel) - { - } +struct strong_index_hasher_adapter { + strong_index_hasher_adapter(Hasher const& hasher) : _hasher{hasher} {} template __device__ inline auto operator()(T const idx) const noexcept { - return remap_sentinel_hash(_hasher(static_cast(idx)), _sentinel); + return _hasher(static_cast(idx)); } private: Hasher _hasher; - hash_value_type const _sentinel; }; /** - * @brief A comparator adapter so that the underlying self comparator can work with strong index - * types. + * @brief An adapter functor to support strong index types for table self comparator. */ template struct strong_index_self_comparator_adapter { @@ -129,8 +123,7 @@ rmm::device_uvector contains(table_view const& haystack, auto const hasher = cudf::experimental::row::hash::row_hasher(haystack, stream); auto const d_hasher = - remap_sentinel_hasher(hasher.device_hasher(nullate::DYNAMIC{has_any_nulls}), - static_cast(map.get_empty_key_sentinel())); + strong_index_hasher_adapter{hasher.device_hasher(nullate::DYNAMIC{has_any_nulls})}; auto const comparator = cudf::experimental::row::equality::self_comparator(haystack, stream); @@ -191,8 +184,7 @@ rmm::device_uvector contains(table_view const& haystack, auto const hasher = cudf::experimental::row::hash::row_hasher(needles, stream); auto const d_hasher = - remap_sentinel_hasher(hasher.device_hasher(nullate::DYNAMIC{has_any_nulls}), - static_cast(map.get_empty_key_sentinel())); + strong_index_hasher_adapter{hasher.device_hasher(nullate::DYNAMIC{has_any_nulls})}; auto const comparator = cudf::experimental::row::equality::two_table_comparator(haystack, needles, stream); From 132bb1cad101d241d1ef5d0fb2c8e3c1f36244a9 Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Thu, 21 Jul 2022 12:23:22 -0700 Subject: [PATCH 4/5] Change headers order --- cpp/src/search/contains_table.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/src/search/contains_table.cu b/cpp/src/search/contains_table.cu index d1ec714aebd..ab4aa55fa49 100644 --- a/cpp/src/search/contains_table.cu +++ b/cpp/src/search/contains_table.cu @@ -22,11 +22,11 @@ #include #include -#include - #include #include +#include + #include namespace cudf::detail { From d57211f8128b247bba2b37ce0a202375ae6a81b1 Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Thu, 21 Jul 2022 20:42:40 -0700 Subject: [PATCH 5/5] Fix comment --- cpp/src/search/contains_table.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/src/search/contains_table.cu b/cpp/src/search/contains_table.cu index ab4aa55fa49..9ef7a8d9c0d 100644 --- a/cpp/src/search/contains_table.cu +++ b/cpp/src/search/contains_table.cu @@ -115,7 +115,7 @@ rmm::device_uvector contains(table_view const& haystack, auto const needles_has_nulls = has_nested_nulls(needles); auto const has_any_nulls = haystack_has_nulls || needles_has_nulls; - // Insert all row hash values and indices of the haystack table. + // Insert row indices of the haystack table as map keys. { auto const haystack_it = cudf::detail::make_counting_transform_iterator( size_type{0},