From c07329896c914ae69e8e03621fec55530727869c Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Wed, 10 Nov 2021 09:56:29 -0800 Subject: [PATCH 01/12] Initial version of semi/anti join using cuco that compiles successfully for non-nullable types. --- cpp/src/join/hash_join.cu | 16 ----- cpp/src/join/hash_join.cuh | 35 ---------- cpp/src/join/join_common_utils.cuh | 53 +++++++++++++++ cpp/src/join/semi_join.cu | 101 ++++++++++++++++++----------- 4 files changed, 115 insertions(+), 90 deletions(-) diff --git a/cpp/src/join/hash_join.cu b/cpp/src/join/hash_join.cu index e4bd1938ecc..a0c56d5dc25 100644 --- a/cpp/src/join/hash_join.cu +++ b/cpp/src/join/hash_join.cu @@ -36,22 +36,6 @@ namespace detail { namespace { -/** - * @brief Device functor to determine if a row is valid. - */ -class row_is_valid { - public: - row_is_valid(bitmask_type const* row_bitmask) : _row_bitmask{row_bitmask} {} - - __device__ __inline__ bool operator()(const size_type& i) const noexcept - { - return cudf::bit_is_set(_row_bitmask, i); - } - - private: - bitmask_type const* _row_bitmask; -}; - } // anonymous namespace std::pair, std::unique_ptr> get_empty_joined_table( diff --git a/cpp/src/join/hash_join.cuh b/cpp/src/join/hash_join.cuh index 10b0e420ef6..eceb2be787c 100644 --- a/cpp/src/join/hash_join.cuh +++ b/cpp/src/join/hash_join.cuh @@ -39,41 +39,6 @@ namespace cudf { namespace detail { -/** - * @brief Remaps a hash value to a new value if it is equal to the specified sentinel value. - * - * @param hash The hash value to potentially remap - * @param sentinel The reserved value - */ -template -constexpr auto remap_sentinel_hash(H hash, S sentinel) -{ - // Arbitrarily choose hash - 1 - return (hash == sentinel) ? (hash - 1) : hash; -} - -/** - * @brief Device functor to create a pair of hash value and index for a given row. - */ -class make_pair_function { - public: - make_pair_function(row_hash const& hash, hash_value_type const empty_key_sentinel) - : _hash{hash}, _empty_key_sentinel{empty_key_sentinel} - { - } - - __device__ __forceinline__ cudf::detail::pair_type operator()(size_type i) const noexcept - { - // Compute the hash value of row `i` - auto row_hash_value = remap_sentinel_hash(_hash(i), _empty_key_sentinel); - return cuco::make_pair(std::move(row_hash_value), std::move(i)); - } - - private: - row_hash _hash; - hash_value_type const _empty_key_sentinel; -}; - /** * @brief Calculates the exact size of the join output produced when * joining two tables together. diff --git a/cpp/src/join/join_common_utils.cuh b/cpp/src/join/join_common_utils.cuh index cec633765c7..9abd83896f4 100644 --- a/cpp/src/join/join_common_utils.cuh +++ b/cpp/src/join/join_common_utils.cuh @@ -24,9 +24,62 @@ #include +// TODO: Figure out what to include for cuco::make_pair. + namespace cudf { namespace detail { +/** + * @brief Remaps a hash value to a new value if it is equal to the specified sentinel value. + * + * @param hash The hash value to potentially remap + * @param sentinel The reserved value + */ +template +constexpr auto remap_sentinel_hash(H hash, S sentinel) +{ + // Arbitrarily choose hash - 1 + return (hash == sentinel) ? (hash - 1) : hash; +} + +/** + * @brief Device functor to create a pair of hash value and index for a given row. + */ +class make_pair_function { + public: + make_pair_function(row_hash const& hash, hash_value_type const empty_key_sentinel) + : _hash{hash}, _empty_key_sentinel{empty_key_sentinel} + { + } + + __device__ __forceinline__ cudf::detail::pair_type operator()(size_type i) const noexcept + { + // Compute the hash value of row `i` + auto row_hash_value = remap_sentinel_hash(_hash(i), _empty_key_sentinel); + return cuco::make_pair(std::move(row_hash_value), std::move(i)); + } + + private: + row_hash _hash; + hash_value_type const _empty_key_sentinel; +}; + +/** + * @brief Device functor to determine if a row is valid. + */ +class row_is_valid { + public: + row_is_valid(bitmask_type const* row_bitmask) : _row_bitmask{row_bitmask} {} + + __device__ __inline__ bool operator()(const size_type& i) const noexcept + { + return cudf::bit_is_set(_row_bitmask, i); + } + + private: + bitmask_type const* _row_bitmask; +}; + /** * @brief Device functor to determine if two pairs are identical. */ diff --git a/cpp/src/join/semi_join.cu b/cpp/src/join/semi_join.cu index 5b5dd418a97..0a8807c28a0 100644 --- a/cpp/src/join/semi_join.cu +++ b/cpp/src/join/semi_join.cu @@ -15,10 +15,12 @@ */ #include +#include #include #include #include +#include #include #include #include @@ -34,10 +36,30 @@ #include #include #include +#include + +#include namespace cudf { namespace detail { +namespace { +// Use the default cuco hasher. +using Hash = cuco::detail::MurmurHash3_32; + +/** + * @brief Device functor to create a pair of hash value and index for a given row. + */ +struct new_pair_function { + __device__ __forceinline__ cudf::detail::pair_type operator()(size_type i) const noexcept + { + // Compute the hash value of row `i` + return cuco::make_pair(i, true); + } +}; + +} // namespace + std::unique_ptr> left_semi_anti_join( join_kind const kind, cudf::table_view const& left_keys, @@ -71,65 +93,66 @@ std::unique_ptr> left_semi_anti_join( auto right_flattened_keys = right_flattened_tables.flattened_columns(); auto left_flattened_keys = left_flattened_tables.flattened_columns(); - // Only care about existence, so we'll use an unordered map (other joins need a multimap) - using hash_table_type = concurrent_unordered_map; + // Create hash table. + auto hash_table = cuco:: + static_map>{ + compute_hash_table_size(right_num_rows), + std::numeric_limits::max(), + cudf::detail::JoinNoneValue, + }; // Create hash table containing all keys found in right table - auto right_rows_d = table_device_view::create(right_flattened_keys, stream); - size_t const hash_table_size = compute_hash_table_size(right_num_rows); - row_hash hash_build{*right_rows_d}; + auto right_rows_d = table_device_view::create(right_flattened_keys, stream); + row_hash const hash_build{*right_rows_d}; row_equality equality_build{*right_rows_d, *right_rows_d, compare_nulls == null_equality::EQUAL}; + Hash hash{}; // Need to create a hash because we need a custom equality comparator and the APIs + // are ordered (..., Hash, Equality) + new_pair_function pair_func_build{}; - // Going to join it with left table - auto left_rows_d = table_device_view::create(left_flattened_keys, stream); - row_hash hash_probe{*left_rows_d}; - row_equality equality_probe{*left_rows_d, *right_rows_d, compare_nulls == null_equality::EQUAL}; + auto iter = cudf::detail::make_counting_transform_iterator(0, pair_func_build); - auto hash_table_ptr = hash_table_type::create(hash_table_size, - stream, - std::numeric_limits::max(), - std::numeric_limits::max(), - hash_build, - equality_build); - auto hash_table = *hash_table_ptr; - - // if compare_nulls == UNEQUAL, we can simply ignore any rows that - // contain a NULL in any column as they will never compare to equal. - auto const row_bitmask = (compare_nulls == null_equality::EQUAL) - ? rmm::device_buffer{} - : cudf::detail::bitmask_and(right_flattened_keys, stream).first; // skip rows that are null here. - thrust::for_each_n( - rmm::exec_policy(stream), - thrust::make_counting_iterator(0), - right_num_rows, - [hash_table, row_bitmask = static_cast(row_bitmask.data())] __device__( - size_type idx) mutable { - if (!row_bitmask || cudf::bit_is_set(row_bitmask, idx)) { - hash_table.insert(thrust::make_pair(idx, true)); - } - }); + if ((compare_nulls == null_equality::EQUAL) or (not nullable(right_keys))) { + // TODO: Need to enable a stream argument in cuco. + hash_table.insert(iter, iter + right_num_rows, hash, equality_build); + // hash_table.insert(iter, iter + right_num_rows, hash, pair_equality, stream.value()); + } else { + CUDF_FAIL("Not supported yet!"); + // thrust::counting_iterator stencil(0); + // auto const row_bitmask = cudf::detail::bitmask_and(right_flattened_keys, stream);; + // row_is_valid pred{static_cast(row_bitmask.data())}; + // + //// insert valid rows + //// TODO: This needs to be implemented still for static_map, it only exists for multimap. + // hash_table.insert_if(iter, iter + right_num_rows, stencil, pred, hash, pair_equality, + // stream.value()); + } - // // Now we have a hash table, we need to iterate over the rows of the left table // and check to see if they are contained in the hash table - // + auto left_rows_d = table_device_view::create(left_flattened_keys, stream); + row_hash hash_probe{*left_rows_d}; + row_equality equality_probe{*left_rows_d, *right_rows_d, compare_nulls == null_equality::EQUAL}; // For semi join we want contains to be true, for anti join we want contains to be false bool const join_type_boolean = (kind == join_kind::LEFT_SEMI_JOIN); + auto hash_table_view = hash_table.get_device_view(); + auto gather_map = std::make_unique>(left_num_rows, stream, mr); // gather_map_end will be the end of valid data in gather_map auto gather_map_end = thrust::copy_if( rmm::exec_policy(stream), - thrust::make_counting_iterator(0), - thrust::make_counting_iterator(left_num_rows), + thrust::make_counting_iterator(0), + thrust::make_counting_iterator(left_num_rows), gather_map->begin(), - [hash_table, join_type_boolean, hash_probe, equality_probe] __device__(size_type idx) { - auto pos = hash_table.find(idx, hash_probe, equality_probe); - return (pos != hash_table.end()) == join_type_boolean; + [hash_table_view, join_type_boolean, hash_probe, equality_probe] __device__( + size_type const idx) { + // Look up this row. The hash function used here needs to map a (left) row index to the hash + // of the row, so it's a row hash. The equality check needs to verify + return hash_table_view.contains(idx, hash_probe, equality_probe) == join_type_boolean; }); auto join_size = thrust::distance(gather_map->begin(), gather_map_end); From e6e6365c26278cac0ad75c16100c6a1b7d210b20 Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Wed, 10 Nov 2021 18:29:14 -0800 Subject: [PATCH 02/12] Add a basic semi join test to verify that the new code works for non-nullable cases and fix a hashing bug. --- cpp/src/join/semi_join.cu | 11 +++++------ cpp/tests/join/semi_anti_join_tests.cpp | 15 +++++++++++++++ 2 files changed, 20 insertions(+), 6 deletions(-) diff --git a/cpp/src/join/semi_join.cu b/cpp/src/join/semi_join.cu index 0a8807c28a0..26fe38fdcf2 100644 --- a/cpp/src/join/semi_join.cu +++ b/cpp/src/join/semi_join.cu @@ -105,8 +105,6 @@ std::unique_ptr> left_semi_anti_join( auto right_rows_d = table_device_view::create(right_flattened_keys, stream); row_hash const hash_build{*right_rows_d}; row_equality equality_build{*right_rows_d, *right_rows_d, compare_nulls == null_equality::EQUAL}; - Hash hash{}; // Need to create a hash because we need a custom equality comparator and the APIs - // are ordered (..., Hash, Equality) new_pair_function pair_func_build{}; auto iter = cudf::detail::make_counting_transform_iterator(0, pair_func_build); @@ -114,8 +112,8 @@ std::unique_ptr> left_semi_anti_join( // skip rows that are null here. if ((compare_nulls == null_equality::EQUAL) or (not nullable(right_keys))) { // TODO: Need to enable a stream argument in cuco. - hash_table.insert(iter, iter + right_num_rows, hash, equality_build); - // hash_table.insert(iter, iter + right_num_rows, hash, pair_equality, stream.value()); + hash_table.insert(iter, iter + right_num_rows, hash_build, equality_build); + // hash_table.insert(iter, iter + right_num_rows, hash_build, equality_build, stream.value()); } else { CUDF_FAIL("Not supported yet!"); // thrust::counting_iterator stencil(0); @@ -124,7 +122,7 @@ std::unique_ptr> left_semi_anti_join( // //// insert valid rows //// TODO: This needs to be implemented still for static_map, it only exists for multimap. - // hash_table.insert_if(iter, iter + right_num_rows, stencil, pred, hash, pair_equality, + // hash_table.insert_if(iter, iter + right_num_rows, stencil, pred, hash_build, equality_build, // stream.value()); } @@ -132,7 +130,8 @@ std::unique_ptr> left_semi_anti_join( // and check to see if they are contained in the hash table auto left_rows_d = table_device_view::create(left_flattened_keys, stream); row_hash hash_probe{*left_rows_d}; - row_equality equality_probe{*left_rows_d, *right_rows_d, compare_nulls == null_equality::EQUAL}; + // TODO: Make sure this order is correct. + row_equality equality_probe{*right_rows_d, *left_rows_d, compare_nulls == null_equality::EQUAL}; // For semi join we want contains to be true, for anti join we want contains to be false bool const join_type_boolean = (kind == join_kind::LEFT_SEMI_JOIN); diff --git a/cpp/tests/join/semi_anti_join_tests.cpp b/cpp/tests/join/semi_anti_join_tests.cpp index 5b38bafb122..ff4270058cd 100644 --- a/cpp/tests/join/semi_anti_join_tests.cpp +++ b/cpp/tests/join/semi_anti_join_tests.cpp @@ -39,6 +39,21 @@ using Table = cudf::table; struct JoinTest : public cudf::test::BaseFixture { }; +TEST_F(JoinTest, TestSimple) +{ + column_wrapper left_col0{0, 1, 2}; + column_wrapper right_col0{0, 1, 3}; + + auto left = cudf::table_view{{left_col0}}; + auto right = cudf::table_view{{right_col0}}; + + auto result = cudf::left_semi_join(left, right); + auto result_cv = cudf::column_view( + cudf::data_type{cudf::type_to_id()}, result->size(), result->data()); + column_wrapper expected{0, 1}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result_cv); +}; + std::pair, std::unique_ptr> get_saj_tables( std::vector const& left_is_human_nulls, std::vector const& right_is_human_nulls) { From e0d947ced3349bdeac121e7fcba544deee4c7984 Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Thu, 11 Nov 2021 12:38:22 -0800 Subject: [PATCH 03/12] Fully functional implementation using cuco. --- cpp/src/join/semi_join.cu | 15 ++++++++------- 1 file changed, 8 insertions(+), 7 deletions(-) diff --git a/cpp/src/join/semi_join.cu b/cpp/src/join/semi_join.cu index 26fe38fdcf2..c7180dfc949 100644 --- a/cpp/src/join/semi_join.cu +++ b/cpp/src/join/semi_join.cu @@ -115,13 +115,14 @@ std::unique_ptr> left_semi_anti_join( hash_table.insert(iter, iter + right_num_rows, hash_build, equality_build); // hash_table.insert(iter, iter + right_num_rows, hash_build, equality_build, stream.value()); } else { - CUDF_FAIL("Not supported yet!"); - // thrust::counting_iterator stencil(0); - // auto const row_bitmask = cudf::detail::bitmask_and(right_flattened_keys, stream);; - // row_is_valid pred{static_cast(row_bitmask.data())}; - // - //// insert valid rows - //// TODO: This needs to be implemented still for static_map, it only exists for multimap. + thrust::counting_iterator stencil(0); + auto const row_bitmask = cudf::detail::bitmask_and(right_flattened_keys, stream); + ; + row_is_valid pred{static_cast(row_bitmask.data())}; + + // insert valid rows + // TODO: Need to enable a stream argument in cuco. + hash_table.insert_if(iter, iter + right_num_rows, stencil, pred, hash_build, equality_build); // hash_table.insert_if(iter, iter + right_num_rows, stencil, pred, hash_build, equality_build, // stream.value()); } From f7e478c498872641af8cadfbe7e3e0e5e6fd7cd5 Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Thu, 11 Nov 2021 14:29:43 -0800 Subject: [PATCH 04/12] Add some more illuminating comments. --- cpp/src/join/semi_join.cu | 11 ++++++++--- 1 file changed, 8 insertions(+), 3 deletions(-) diff --git a/cpp/src/join/semi_join.cu b/cpp/src/join/semi_join.cu index c7180dfc949..037b9d89429 100644 --- a/cpp/src/join/semi_join.cu +++ b/cpp/src/join/semi_join.cu @@ -111,7 +111,8 @@ std::unique_ptr> left_semi_anti_join( // skip rows that are null here. if ((compare_nulls == null_equality::EQUAL) or (not nullable(right_keys))) { - // TODO: Need to enable a stream argument in cuco. + // TODO: Need to enable a stream argument in cuco, see + // https://github.com/NVIDIA/cuCollections/pull/113 hash_table.insert(iter, iter + right_num_rows, hash_build, equality_build); // hash_table.insert(iter, iter + right_num_rows, hash_build, equality_build, stream.value()); } else { @@ -121,7 +122,8 @@ std::unique_ptr> left_semi_anti_join( row_is_valid pred{static_cast(row_bitmask.data())}; // insert valid rows - // TODO: Need to enable a stream argument in cuco. + // TODO: Need to enable a stream argument in cuco, see + // https://github.com/NVIDIA/cuCollections/pull/113 hash_table.insert_if(iter, iter + right_num_rows, stencil, pred, hash_build, equality_build); // hash_table.insert_if(iter, iter + right_num_rows, stencil, pred, hash_build, equality_build, // stream.value()); @@ -131,7 +133,10 @@ std::unique_ptr> left_semi_anti_join( // and check to see if they are contained in the hash table auto left_rows_d = table_device_view::create(left_flattened_keys, stream); row_hash hash_probe{*left_rows_d}; - // TODO: Make sure this order is correct. + // Note: This equality comparator violates symmetry of equality and is + // therefore relying on the implementation detail of the order in which its + // operator is invoked. If cuco makes no promises about the order of + // invocation this seems a bit unsafe. row_equality equality_probe{*right_rows_d, *left_rows_d, compare_nulls == null_equality::EQUAL}; // For semi join we want contains to be true, for anti join we want contains to be false From 2ab97023c70135a6374b6179dfbf2f6b580781ff Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Thu, 11 Nov 2021 15:42:23 -0800 Subject: [PATCH 05/12] Remove now unnecessary header. --- cpp/src/join/semi_join.cu | 1 - 1 file changed, 1 deletion(-) diff --git a/cpp/src/join/semi_join.cu b/cpp/src/join/semi_join.cu index 037b9d89429..ac3734037b0 100644 --- a/cpp/src/join/semi_join.cu +++ b/cpp/src/join/semi_join.cu @@ -14,7 +14,6 @@ * limitations under the License. */ -#include #include #include From 603105130d1f1de4f8d3ae5f41a8734f649adc88 Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Thu, 11 Nov 2021 16:02:19 -0800 Subject: [PATCH 06/12] Update for current new bitmask_and API. --- cpp/src/join/semi_join.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/src/join/semi_join.cu b/cpp/src/join/semi_join.cu index ac3734037b0..0e00e51d965 100644 --- a/cpp/src/join/semi_join.cu +++ b/cpp/src/join/semi_join.cu @@ -116,8 +116,8 @@ std::unique_ptr> left_semi_anti_join( // hash_table.insert(iter, iter + right_num_rows, hash_build, equality_build, stream.value()); } else { thrust::counting_iterator stencil(0); - auto const row_bitmask = cudf::detail::bitmask_and(right_flattened_keys, stream); - ; + [[maybe_unused]] auto const [row_bitmask, null_count] = + cudf::detail::bitmask_and(right_flattened_keys, stream); row_is_valid pred{static_cast(row_bitmask.data())}; // insert valid rows From ad4e29506bf4ff8ff36e8d79350f52f361ed0e77 Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Thu, 11 Nov 2021 17:58:07 -0800 Subject: [PATCH 07/12] Remove cruft from initial impls. --- cpp/src/join/hash_join.cuh | 35 ++++++++++++++++++++++++++++ cpp/src/join/join_common_utils.cuh | 37 ------------------------------ cpp/src/join/semi_join.cu | 5 ++-- 3 files changed, 37 insertions(+), 40 deletions(-) diff --git a/cpp/src/join/hash_join.cuh b/cpp/src/join/hash_join.cuh index eceb2be787c..10b0e420ef6 100644 --- a/cpp/src/join/hash_join.cuh +++ b/cpp/src/join/hash_join.cuh @@ -39,6 +39,41 @@ namespace cudf { namespace detail { +/** + * @brief Remaps a hash value to a new value if it is equal to the specified sentinel value. + * + * @param hash The hash value to potentially remap + * @param sentinel The reserved value + */ +template +constexpr auto remap_sentinel_hash(H hash, S sentinel) +{ + // Arbitrarily choose hash - 1 + return (hash == sentinel) ? (hash - 1) : hash; +} + +/** + * @brief Device functor to create a pair of hash value and index for a given row. + */ +class make_pair_function { + public: + make_pair_function(row_hash const& hash, hash_value_type const empty_key_sentinel) + : _hash{hash}, _empty_key_sentinel{empty_key_sentinel} + { + } + + __device__ __forceinline__ cudf::detail::pair_type operator()(size_type i) const noexcept + { + // Compute the hash value of row `i` + auto row_hash_value = remap_sentinel_hash(_hash(i), _empty_key_sentinel); + return cuco::make_pair(std::move(row_hash_value), std::move(i)); + } + + private: + row_hash _hash; + hash_value_type const _empty_key_sentinel; +}; + /** * @brief Calculates the exact size of the join output produced when * joining two tables together. diff --git a/cpp/src/join/join_common_utils.cuh b/cpp/src/join/join_common_utils.cuh index 9abd83896f4..73470793d5d 100644 --- a/cpp/src/join/join_common_utils.cuh +++ b/cpp/src/join/join_common_utils.cuh @@ -24,46 +24,9 @@ #include -// TODO: Figure out what to include for cuco::make_pair. - namespace cudf { namespace detail { -/** - * @brief Remaps a hash value to a new value if it is equal to the specified sentinel value. - * - * @param hash The hash value to potentially remap - * @param sentinel The reserved value - */ -template -constexpr auto remap_sentinel_hash(H hash, S sentinel) -{ - // Arbitrarily choose hash - 1 - return (hash == sentinel) ? (hash - 1) : hash; -} - -/** - * @brief Device functor to create a pair of hash value and index for a given row. - */ -class make_pair_function { - public: - make_pair_function(row_hash const& hash, hash_value_type const empty_key_sentinel) - : _hash{hash}, _empty_key_sentinel{empty_key_sentinel} - { - } - - __device__ __forceinline__ cudf::detail::pair_type operator()(size_type i) const noexcept - { - // Compute the hash value of row `i` - auto row_hash_value = remap_sentinel_hash(_hash(i), _empty_key_sentinel); - return cuco::make_pair(std::move(row_hash_value), std::move(i)); - } - - private: - row_hash _hash; - hash_value_type const _empty_key_sentinel; -}; - /** * @brief Device functor to determine if a row is valid. */ diff --git a/cpp/src/join/semi_join.cu b/cpp/src/join/semi_join.cu index 0e00e51d965..a60778b430f 100644 --- a/cpp/src/join/semi_join.cu +++ b/cpp/src/join/semi_join.cu @@ -49,10 +49,9 @@ using Hash = cuco::detail::MurmurHash3_32; /** * @brief Device functor to create a pair of hash value and index for a given row. */ -struct new_pair_function { +struct make_pair_function { __device__ __forceinline__ cudf::detail::pair_type operator()(size_type i) const noexcept { - // Compute the hash value of row `i` return cuco::make_pair(i, true); } }; @@ -104,7 +103,7 @@ std::unique_ptr> left_semi_anti_join( auto right_rows_d = table_device_view::create(right_flattened_keys, stream); row_hash const hash_build{*right_rows_d}; row_equality equality_build{*right_rows_d, *right_rows_d, compare_nulls == null_equality::EQUAL}; - new_pair_function pair_func_build{}; + make_pair_function pair_func_build{}; auto iter = cudf::detail::make_counting_transform_iterator(0, pair_func_build); From e2afdc5aa4e48032a88cbdae8dca0f90eb39e5f4 Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Thu, 11 Nov 2021 17:59:38 -0800 Subject: [PATCH 08/12] Remove unnecessary typedef. --- cpp/src/join/semi_join.cu | 3 --- 1 file changed, 3 deletions(-) diff --git a/cpp/src/join/semi_join.cu b/cpp/src/join/semi_join.cu index a60778b430f..f5f55915626 100644 --- a/cpp/src/join/semi_join.cu +++ b/cpp/src/join/semi_join.cu @@ -43,9 +43,6 @@ namespace cudf { namespace detail { namespace { -// Use the default cuco hasher. -using Hash = cuco::detail::MurmurHash3_32; - /** * @brief Device functor to create a pair of hash value and index for a given row. */ From 69d2d0773ab9a0ef7e39eb1ba3ca5214e0ed7b89 Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Thu, 18 Nov 2021 13:08:56 -0800 Subject: [PATCH 09/12] Remove comment now that streams are enabled. --- cpp/src/join/semi_join.cu | 7 ++----- 1 file changed, 2 insertions(+), 5 deletions(-) diff --git a/cpp/src/join/semi_join.cu b/cpp/src/join/semi_join.cu index f5f55915626..744705c0ddd 100644 --- a/cpp/src/join/semi_join.cu +++ b/cpp/src/join/semi_join.cu @@ -117,11 +117,8 @@ std::unique_ptr> left_semi_anti_join( row_is_valid pred{static_cast(row_bitmask.data())}; // insert valid rows - // TODO: Need to enable a stream argument in cuco, see - // https://github.com/NVIDIA/cuCollections/pull/113 - hash_table.insert_if(iter, iter + right_num_rows, stencil, pred, hash_build, equality_build); - // hash_table.insert_if(iter, iter + right_num_rows, stencil, pred, hash_build, equality_build, - // stream.value()); + hash_table.insert_if( + iter, iter + right_num_rows, stencil, pred, hash_build, equality_build, stream.value()); } // Now we have a hash table, we need to iterate over the rows of the left table From dab5739bf8c1b4017d166c3f83ecc7a1c221b5e5 Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Tue, 23 Nov 2021 09:24:46 -0800 Subject: [PATCH 10/12] Address PR comments. --- cpp/src/join/semi_join.cu | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/cpp/src/join/semi_join.cu b/cpp/src/join/semi_join.cu index 744705c0ddd..f692ef91490 100644 --- a/cpp/src/join/semi_join.cu +++ b/cpp/src/join/semi_join.cu @@ -49,7 +49,9 @@ namespace { struct make_pair_function { __device__ __forceinline__ cudf::detail::pair_type operator()(size_type i) const noexcept { - return cuco::make_pair(i, true); + // The value is irrelevant since we only ever use the hash map to check for + // membership of a particular row index. + return cuco::make_pair(i, 0); } }; @@ -112,8 +114,7 @@ std::unique_ptr> left_semi_anti_join( // hash_table.insert(iter, iter + right_num_rows, hash_build, equality_build, stream.value()); } else { thrust::counting_iterator stencil(0); - [[maybe_unused]] auto const [row_bitmask, null_count] = - cudf::detail::bitmask_and(right_flattened_keys, stream); + auto const [row_bitmask, _] = cudf::detail::bitmask_and(right_flattened_keys, stream); row_is_valid pred{static_cast(row_bitmask.data())}; // insert valid rows From ec8185ee8f39b3598baae65de9de3ad13d6b35ce Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Tue, 23 Nov 2021 09:25:41 -0800 Subject: [PATCH 11/12] Update cuco commit hash. --- cpp/cmake/thirdparty/get_cucollections.cmake | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/cmake/thirdparty/get_cucollections.cmake b/cpp/cmake/thirdparty/get_cucollections.cmake index b6cb9757ae8..cadf87ae78d 100644 --- a/cpp/cmake/thirdparty/get_cucollections.cmake +++ b/cpp/cmake/thirdparty/get_cucollections.cmake @@ -21,7 +21,7 @@ function(find_and_configure_cucollections) cuco 0.0 GLOBAL_TARGETS cuco::cuco CPM_ARGS GITHUB_REPOSITORY NVIDIA/cuCollections - GIT_TAG f0eecb203590f1f4ac4a9f1700229f4434ac64dc + GIT_TAG bb8c34ddeaf6b526c0e9a6571912d3f4514c0753 OPTIONS "BUILD_TESTS OFF" "BUILD_BENCHMARKS OFF" "BUILD_EXAMPLES OFF" ) From 3360100cfce2dc0b7ad684c2755b08d5f5820275 Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Thu, 16 Dec 2021 10:35:05 -0800 Subject: [PATCH 12/12] Address outstanding issues with recently added cuco features. --- cpp/src/join/semi_join.cu | 10 ++++------ 1 file changed, 4 insertions(+), 6 deletions(-) diff --git a/cpp/src/join/semi_join.cu b/cpp/src/join/semi_join.cu index 203e1c4c549..9daf5b52693 100644 --- a/cpp/src/join/semi_join.cu +++ b/cpp/src/join/semi_join.cu @@ -92,11 +92,12 @@ std::unique_ptr> left_semi_anti_join( // Create hash table. auto hash_table = cuco:: - static_map>{ + static_map{ compute_hash_table_size(right_num_rows), std::numeric_limits::max(), cudf::detail::JoinNoneValue, - }; + hash_table_allocator_type{default_allocator{}, stream}, + stream.value()}; // Create hash table containing all keys found in right table auto right_rows_d = table_device_view::create(right_flattened_keys, stream); @@ -108,10 +109,7 @@ std::unique_ptr> left_semi_anti_join( // skip rows that are null here. if ((compare_nulls == null_equality::EQUAL) or (not nullable(right_keys))) { - // TODO: Need to enable a stream argument in cuco, see - // https://github.com/NVIDIA/cuCollections/pull/113 - hash_table.insert(iter, iter + right_num_rows, hash_build, equality_build); - // hash_table.insert(iter, iter + right_num_rows, hash_build, equality_build, stream.value()); + hash_table.insert(iter, iter + right_num_rows, hash_build, equality_build, stream.value()); } else { thrust::counting_iterator stencil(0); auto const [row_bitmask, _] = cudf::detail::bitmask_and(right_flattened_keys, stream);