diff --git a/cpp/include/cudf/detail/cuco_helpers.hpp b/cpp/include/cudf/detail/cuco_helpers.hpp index 926df921715..27cb3e5d543 100644 --- a/cpp/include/cudf/detail/cuco_helpers.hpp +++ b/cpp/include/cudf/detail/cuco_helpers.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024, NVIDIA CORPORATION. + * Copyright (c) 2024-2025, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -29,17 +29,4 @@ static cudf::size_type constexpr CUDF_SIZE_TYPE_SENTINEL = -1; /// Default load factor for cuco data structures static double constexpr CUCO_DESIRED_LOAD_FACTOR = 0.5; -/** - * @brief Stream-ordered allocator adaptor used for cuco data structures - * - * The stream-ordered `rmm::mr::polymorphic_allocator` cannot be used in `cuco` directly since the - * later expects a standard C++ `Allocator` interface. This allocator helper provides a simple way - * to handle cuco memory allocation/deallocation with the given `stream` and the rmm default memory - * resource. - * - * @tparam T The allocator's value type. - */ -template -using cuco_allocator = rmm::mr::stream_allocator_adaptor>; - } // namespace cudf::detail diff --git a/cpp/include/cudf/detail/join/distinct_hash_join.cuh b/cpp/include/cudf/detail/join/distinct_hash_join.cuh index 620be9b35ee..74e9c170c9a 100644 --- a/cpp/include/cudf/detail/join/distinct_hash_join.cuh +++ b/cpp/include/cudf/detail/join/distinct_hash_join.cuh @@ -164,7 +164,7 @@ class distinct_hash_join { cuda::thread_scope_device, always_not_equal, probing_scheme_type, - cudf::detail::cuco_allocator, + rmm::mr::polymorphic_allocator, cuco_storage_type>; bool _has_nested_columns; ///< True if nested columns are present in build and probe tables diff --git a/cpp/include/cudf/detail/join/filtered_join.cuh b/cpp/include/cudf/detail/join/filtered_join.cuh index 315a9b1a47e..cfd0cc13aa8 100644 --- a/cpp/include/cudf/detail/join/filtered_join.cuh +++ b/cpp/include/cudf/detail/join/filtered_join.cuh @@ -175,7 +175,7 @@ class filtered_join { cuco::bucket_storage, - cudf::detail::cuco_allocator>; + rmm::mr::polymorphic_allocator>; // Hasher for primitive row types using primitive_row_hasher = diff --git a/cpp/include/cudf/detail/join/hash_join.cuh b/cpp/include/cudf/detail/join/hash_join.cuh index 99f11126999..772b19a4791 100644 --- a/cpp/include/cudf/detail/join/hash_join.cuh +++ b/cpp/include/cudf/detail/join/hash_join.cuh @@ -93,7 +93,7 @@ struct hash_join { cuda::thread_scope_device, always_not_equal, cuco::double_hashing, - cudf::detail::cuco_allocator, + rmm::mr::polymorphic_allocator, cuco::storage<2>>; hash_join() = delete; diff --git a/cpp/src/groupby/hash/compute_groupby.cu b/cpp/src/groupby/hash/compute_groupby.cu index c8d92621764..5d64676db80 100644 --- a/cpp/src/groupby/hash/compute_groupby.cu +++ b/cpp/src/groupby/hash/compute_groupby.cu @@ -109,16 +109,16 @@ std::unique_ptr compute_groupby(table_view const& keys, return hashes; }(); - auto set = cuco::static_set{ - cuco::extent{static_cast(num_keys)}, - cudf::detail::CUCO_DESIRED_LOAD_FACTOR, // 50% load factor - cuco::empty_key{cudf::detail::CUDF_SIZE_TYPE_SENTINEL}, - d_row_equal, - probing_scheme_t{row_hasher_with_cache_t{d_row_hash, cached_hashes.data()}}, - cuco::thread_scope_device, - cuco::storage{}, - cudf::detail::cuco_allocator{rmm::mr::polymorphic_allocator{}, stream}, - stream.value()}; + auto set = + cuco::static_set{cuco::extent{static_cast(num_keys)}, + cudf::detail::CUCO_DESIRED_LOAD_FACTOR, // 50% load factor + cuco::empty_key{cudf::detail::CUDF_SIZE_TYPE_SENTINEL}, + d_row_equal, + probing_scheme_t{row_hasher_with_cache_t{d_row_hash, cached_hashes.data()}}, + cuco::thread_scope_device, + cuco::storage{}, + rmm::mr::polymorphic_allocator{}, + stream.value()}; auto const gather_keys = [&](auto const& gather_map) { return cudf::detail::gather(keys, diff --git a/cpp/src/groupby/hash/helpers.cuh b/cpp/src/groupby/hash/helpers.cuh index ffffaec0680..c9d9210949c 100644 --- a/cpp/src/groupby/hash/helpers.cuh +++ b/cpp/src/groupby/hash/helpers.cuh @@ -98,7 +98,7 @@ using global_set_t = cuco::static_set, + rmm::mr::polymorphic_allocator, cuco::storage>; using nullable_global_set_t = cuco::static_set, + rmm::mr::polymorphic_allocator, cuco::storage>; template diff --git a/cpp/src/io/json/json_tree.cu b/cpp/src/io/json/json_tree.cu index 15797f4a800..64ee8bf0df9 100644 --- a/cpp/src/io/json/json_tree.cu +++ b/cpp/src/io/json/json_tree.cu @@ -553,15 +553,14 @@ std::pair> remapped_field_nodes_after_uni using hasher_type = decltype(d_hasher); constexpr size_type empty_node_index_sentinel = -1; - auto key_set = cuco::static_set{ - cuco::extent{compute_hash_table_size(num_keys)}, - cuco::empty_key{empty_node_index_sentinel}, - d_equal, - cuco::linear_probing<1, hasher_type>{d_hasher}, - {}, - {}, - cudf::detail::cuco_allocator{rmm::mr::polymorphic_allocator{}, stream}, - stream.value()}; + auto key_set = cuco::static_set{cuco::extent{compute_hash_table_size(num_keys)}, + cuco::empty_key{empty_node_index_sentinel}, + d_equal, + cuco::linear_probing<1, hasher_type>{d_hasher}, + {}, + {}, + rmm::mr::polymorphic_allocator{}, + stream.value()}; auto const counting_iter = thrust::make_counting_iterator(0); rmm::device_uvector found_keys(num_keys, stream); key_set.insert_and_find_async(counting_iter, @@ -629,15 +628,15 @@ rmm::device_uvector hash_node_type_with_field_name(device_span{d_hasher}, - {}, - {}, - cudf::detail::cuco_allocator{rmm::mr::polymorphic_allocator{}, stream}, - stream.value()}; + auto key_set = + cuco::static_set{cuco::extent{compute_hash_table_size(num_fields, 40)}, // 40% occupancy + cuco::empty_key{empty_node_index_sentinel}, + d_equal, + cuco::linear_probing<1, hasher_type>{d_hasher}, + {}, + {}, + rmm::mr::polymorphic_allocator{}, + stream.value()}; key_set.insert_if_async(counting_iter, counting_iter + num_nodes, thrust::counting_iterator(0), // stencil @@ -659,7 +658,7 @@ rmm::device_uvector hash_node_type_with_field_name(device_span{hasher_type3{}}, {}, {}, - cudf::detail::cuco_allocator{rmm::mr::polymorphic_allocator{}, stream}, + rmm::mr::polymorphic_allocator{}, stream.value()}; }; if (!is_enabled_experimental) { return std::pair{false, make_map(0)}; } @@ -864,15 +863,14 @@ std::pair, rmm::device_uvector> hash_n constexpr size_type empty_node_index_sentinel = -1; using hasher_type = decltype(d_hashed_cache); - auto key_set = cuco::static_set{ - cuco::extent{compute_hash_table_size(num_nodes)}, - cuco::empty_key{empty_node_index_sentinel}, - d_equal, - cuco::linear_probing<1, hasher_type>{d_hashed_cache}, - {}, - {}, - cudf::detail::cuco_allocator{rmm::mr::polymorphic_allocator{}, stream}, - stream.value()}; + auto key_set = cuco::static_set{cuco::extent{compute_hash_table_size(num_nodes)}, + cuco::empty_key{empty_node_index_sentinel}, + d_equal, + cuco::linear_probing<1, hasher_type>{d_hashed_cache}, + {}, + {}, + rmm::mr::polymorphic_allocator{}, + stream.value()}; // insert and convert node ids to unique set ids auto nodes_itr = thrust::make_counting_iterator(0); diff --git a/cpp/src/io/orc/orc_gpu.hpp b/cpp/src/io/orc/orc_gpu.hpp index 3a478740b83..7206f2c146d 100644 --- a/cpp/src/io/orc/orc_gpu.hpp +++ b/cpp/src/io/orc/orc_gpu.hpp @@ -50,7 +50,7 @@ auto constexpr occupancy_factor = 1.43f; ///< cuCollections suggests using a ha using storage_type = cuco::bucket_storage, - cudf::detail::cuco_allocator>; + rmm::mr::polymorphic_allocator>; using storage_ref_type = typename storage_type::ref_type; auto constexpr KEY_SENTINEL = size_type{-1}; diff --git a/cpp/src/io/orc/writer_impl.cu b/cpp/src/io/orc/writer_impl.cu index b34ed1fd2dd..b8189cf5db3 100644 --- a/cpp/src/io/orc/writer_impl.cu +++ b/cpp/src/io/orc/writer_impl.cu @@ -2166,8 +2166,7 @@ stripe_dictionaries build_dictionaries(orc_table_view& orc_table, // Create a single bulk storage to use for all sub-dictionaries auto map_storage = std::make_unique( - total_map_storage_size, - cudf::detail::cuco_allocator{rmm::mr::polymorphic_allocator{}, stream}); + total_map_storage_size, rmm::mr::polymorphic_allocator{}, stream.value()); // Initialize stripe dictionaries for (auto col_idx : orc_table.string_column_indices) { diff --git a/cpp/src/io/parquet/experimental/deletion_vectors.cu b/cpp/src/io/parquet/experimental/deletion_vectors.cu index c40fd31021e..c9992cb6fe5 100644 --- a/cpp/src/io/parquet/experimental/deletion_vectors.cu +++ b/cpp/src/io/parquet/experimental/deletion_vectors.cu @@ -14,7 +14,6 @@ * limitations under the License. */ -#include #include #include #include @@ -25,6 +24,7 @@ #include #include #include +#include #include #include @@ -40,7 +40,7 @@ namespace cudf::io::parquet::experimental { // Type alias for the cuco 64-bit roaring bitmap using roaring_bitmap_type = - cuco::experimental::roaring_bitmap>; + cuco::experimental::roaring_bitmap>; namespace { @@ -299,9 +299,7 @@ chunked_parquet_reader::chunked_parquet_reader( if (not serialized_roaring64.empty()) { _deletion_vector = std::make_unique( - serialized_roaring64.data(), - cudf::detail::cuco_allocator{rmm::mr::polymorphic_allocator{}, _stream.value()}, - _stream); + serialized_roaring64.data(), rmm::mr::polymorphic_allocator{}, _stream); } } @@ -311,7 +309,7 @@ chunked_parquet_reader::chunked_parquet_reader( struct chunked_parquet_reader::roaring_bitmap_impl { roaring_bitmap_type roaring_bitmap; roaring_bitmap_impl(cuda::std::byte const* const serialized_roaring64_data, - cudf::detail::cuco_allocator const& allocator, + rmm::mr::polymorphic_allocator const& allocator, rmm::cuda_stream_view stream) : roaring_bitmap(serialized_roaring64_data, allocator, stream) { @@ -445,10 +443,8 @@ table_with_metadata read_parquet(parquet_reader_options const& options, } // Filter the table using the deletion vector - auto deletion_vector = roaring_bitmap_type( - serialized_roaring64.data(), - cudf::detail::cuco_allocator{rmm::mr::polymorphic_allocator{}, stream.value()}, - stream); + auto deletion_vector = roaring_bitmap_type{ + serialized_roaring64.data(), rmm::mr::polymorphic_allocator{}, stream}; auto row_mask = build_row_mask_column(table_with_index->get_column(0).view(), deletion_vector, num_rows, diff --git a/cpp/src/io/parquet/experimental/dictionary_page_filter.cu b/cpp/src/io/parquet/experimental/dictionary_page_filter.cu index 944f7e69333..2cb0e1b6163 100644 --- a/cpp/src/io/parquet/experimental/dictionary_page_filter.cu +++ b/cpp/src/io/parquet/experimental/dictionary_page_filter.cu @@ -97,7 +97,7 @@ using hasher_type = cudf::hashing::detail::MurmurHash3_x86_32; using storage_type = cuco::bucket_storage, - cudf::detail::cuco_allocator>; + rmm::mr::polymorphic_allocator>; using storage_ref_type = typename storage_type::ref_type; /** @@ -1105,9 +1105,8 @@ struct dictionary_caster { auto const total_num_literals = static_cast(literals.size()); // Create a single bulk storage used by all cuco hash sets - auto set_storage = storage_type{ - total_set_storage_size, - cudf::detail::cuco_allocator{rmm::mr::polymorphic_allocator{}, stream}}; + auto set_storage = + storage_type{total_set_storage_size, rmm::mr::polymorphic_allocator{}, stream.value()}; // Initialize storage with the empty key sentinel set_storage.initialize_async(EMPTY_KEY_SENTINEL, {stream.value()}); diff --git a/cpp/src/io/parquet/parquet_gpu.cuh b/cpp/src/io/parquet/parquet_gpu.cuh index d916a9e0c55..2fa9f5be018 100644 --- a/cpp/src/io/parquet/parquet_gpu.cuh +++ b/cpp/src/io/parquet/parquet_gpu.cuh @@ -46,7 +46,7 @@ auto constexpr SCOPE = cuda::thread_scope_block; using storage_type = cuco::bucket_storage, - cudf::detail::cuco_allocator>; + rmm::mr::polymorphic_allocator>; using storage_ref_type = typename storage_type::ref_type; /** diff --git a/cpp/src/io/parquet/writer_impl.cu b/cpp/src/io/parquet/writer_impl.cu index 621eadb0e5b..48f420a2257 100644 --- a/cpp/src/io/parquet/writer_impl.cu +++ b/cpp/src/io/parquet/writer_impl.cu @@ -1346,9 +1346,8 @@ build_chunk_dictionaries(hostdevice_2dvector& chunks, if (total_map_storage_size == 0) { return {std::move(dict_data), std::move(dict_index)}; } // Create a single bulk storage used by all sub-dictionaries - auto map_storage = storage_type{ - total_map_storage_size, - cudf::detail::cuco_allocator{rmm::mr::polymorphic_allocator{}, stream}}; + auto map_storage = + storage_type{total_map_storage_size, rmm::mr::polymorphic_allocator{}, stream.value()}; // Create a span of non-const map_storage as map_storage_ref takes in a non-const pointer. device_span const map_storage_data{map_storage.data(), total_map_storage_size}; diff --git a/cpp/src/join/distinct_hash_join.cu b/cpp/src/join/distinct_hash_join.cu index 80972d32d1d..6613012e2a0 100644 --- a/cpp/src/join/distinct_hash_join.cu +++ b/cpp/src/join/distinct_hash_join.cu @@ -177,7 +177,7 @@ distinct_hash_join::distinct_hash_join(cudf::table_view const& build, {}, cuco::thread_scope_device, cuco_storage_type{}, - cudf::detail::cuco_allocator{rmm::mr::polymorphic_allocator{}, stream}, + rmm::mr::polymorphic_allocator{}, stream.value()} { CUDF_FUNC_RANGE(); diff --git a/cpp/src/join/filtered_join.cu b/cpp/src/join/filtered_join.cu index d830aaf21c4..9a14d223b84 100644 --- a/cpp/src/join/filtered_join.cu +++ b/cpp/src/join/filtered_join.cu @@ -226,7 +226,8 @@ filtered_join::filtered_join(cudf::table_view const& build, _build{build}, _preprocessed_build{cudf::detail::row::equality::preprocessed_table::create(_build, stream)}, _bucket_storage{cuco::extent{compute_bucket_storage_size(build, load_factor)}, - cuco_allocator{rmm::mr::polymorphic_allocator{}, stream.value()}} + rmm::mr::polymorphic_allocator{}, + stream.value()} { _bucket_storage.initialize(empty_sentinel_key, stream); } diff --git a/cpp/src/join/hash_join.cu b/cpp/src/join/hash_join.cu index 793ddd84f2d..e4148a877a7 100644 --- a/cpp/src/join/hash_join.cu +++ b/cpp/src/join/hash_join.cu @@ -516,17 +516,16 @@ hash_join::hash_join(cudf::table_view const& build, : _has_nulls(has_nulls), _is_empty{build.num_rows() == 0}, _nulls_equal{compare_nulls}, - _hash_table{ - cuco::extent{static_cast(build.num_rows())}, - load_factor, - cuco::empty_key{ - cuco::pair{std::numeric_limits::max(), cudf::detail::JoinNoneValue}}, - {}, - {}, - {}, - {}, - cudf::detail::cuco_allocator{rmm::mr::polymorphic_allocator{}, stream.value()}, - stream.value()}, + _hash_table{cuco::extent{static_cast(build.num_rows())}, + load_factor, + cuco::empty_key{cuco::pair{std::numeric_limits::max(), + cudf::detail::JoinNoneValue}}, + {}, + {}, + {}, + {}, + rmm::mr::polymorphic_allocator{}, + stream.value()}, _build{build}, _preprocessed_build{cudf::detail::row::equality::preprocessed_table::create(_build, stream)} { diff --git a/cpp/src/join/join_common_utils.hpp b/cpp/src/join/join_common_utils.hpp index 94b31449ba5..96f800834ed 100644 --- a/cpp/src/join/join_common_utils.hpp +++ b/cpp/src/join/join_common_utils.hpp @@ -85,7 +85,7 @@ using mixed_multiset_type = cuda::thread_scope_device, mixed_join_always_not_equal, cuco::double_hashing<1, mixed_join_hasher1, mixed_join_hasher2>, - cudf::detail::cuco_allocator, + rmm::mr::polymorphic_allocator, cuco::storage<2>>; bool is_trivial_join(table_view const& left, table_view const& right, join_kind join_type); diff --git a/cpp/src/join/mixed_join.cu b/cpp/src/join/mixed_join.cu index 75419cefe44..6fb42fd89da 100644 --- a/cpp/src/join/mixed_join.cu +++ b/cpp/src/join/mixed_join.cu @@ -189,7 +189,7 @@ mixed_join_setup_data setup_mixed_join_common(table_view const& left_equality, {}, {}, {}, - cudf::detail::cuco_allocator{rmm::mr::polymorphic_allocator{}, stream}, + rmm::mr::polymorphic_allocator{}, stream.value()}; // TODO: To add support for nested columns we will need to flatten in many diff --git a/cpp/src/join/mixed_join_common_utils.cuh b/cpp/src/join/mixed_join_common_utils.cuh index ead6814dbe9..0c7f4681ee8 100644 --- a/cpp/src/join/mixed_join_common_utils.cuh +++ b/cpp/src/join/mixed_join_common_utils.cuh @@ -166,7 +166,7 @@ using hash_set_type = cuda::thread_scope_device, double_row_equality_comparator, cuco::linear_probing, - cudf::detail::cuco_allocator, + rmm::mr::polymorphic_allocator, cuco::storage<1>>; // The hash_set_ref_type used by mixed_semi_join kernels for probing. diff --git a/cpp/src/join/mixed_join_semi.cu b/cpp/src/join/mixed_join_semi.cu index 06b5e9d4c7a..ce652fb271b 100644 --- a/cpp/src/join/mixed_join_semi.cu +++ b/cpp/src/join/mixed_join_semi.cu @@ -153,15 +153,14 @@ std::unique_ptr> mixed_join_semi( auto const equality_build_conditional = row_comparator_conditional_build.equal_to(build_nulls, compare_nulls); - hash_set_type row_set{ - {compute_hash_table_size(build.num_rows())}, - cuco::empty_key{JoinNoneValue}, - {equality_build_equality, equality_build_conditional}, - {row_hash_build.device_hasher(build_nulls)}, - {}, - {}, - cudf::detail::cuco_allocator{rmm::mr::polymorphic_allocator{}, stream}, - {stream.value()}}; + hash_set_type row_set{{compute_hash_table_size(build.num_rows())}, + cuco::empty_key{JoinNoneValue}, + {equality_build_equality, equality_build_conditional}, + {row_hash_build.device_hasher(build_nulls)}, + {}, + {}, + rmm::mr::polymorphic_allocator{}, + {stream.value()}}; auto iter = thrust::make_counting_iterator(0); diff --git a/cpp/src/reductions/histogram.cu b/cpp/src/reductions/histogram.cu index aea895db788..f459cb402ba 100644 --- a/cpp/src/reductions/histogram.cu +++ b/cpp/src/reductions/histogram.cu @@ -148,16 +148,16 @@ compute_row_frequencies(table_view const& input, histogram_count_type{0}); // Construct a hash set - auto row_set = cuco::static_set{ - cuco::extent{num_rows}, - cudf::detail::CUCO_DESIRED_LOAD_FACTOR, - cuco::empty_key{-1}, - key_equal, - cuco::linear_probing{key_hasher}, - {}, // thread scope - {}, // storage - cudf::detail::cuco_allocator{rmm::mr::polymorphic_allocator{}, stream}, - stream.value()}; + auto row_set = + cuco::static_set{cuco::extent{num_rows}, + cudf::detail::CUCO_DESIRED_LOAD_FACTOR, + cuco::empty_key{-1}, + key_equal, + cuco::linear_probing{key_hasher}, + {}, // thread scope + {}, // storage + rmm::mr::polymorphic_allocator{}, + stream.value()}; // Device-accessible reference to the hash set with `insert_and_find` operator auto row_set_ref = row_set.ref(cuco::op::insert_and_find); diff --git a/cpp/src/search/contains_table_impl.cuh b/cpp/src/search/contains_table_impl.cuh index ca75a223370..f305be084a8 100644 --- a/cpp/src/search/contains_table_impl.cuh +++ b/cpp/src/search/contains_table_impl.cuh @@ -143,15 +143,14 @@ void perform_contains(table_view const& haystack, return lhs_index_type{idx}; })); - auto set = cuco::static_set{ - cuco::extent{compute_hash_table_size(haystack.num_rows())}, - cuco::empty_key{rhs_index_type{-1}}, - d_equal, - probing_scheme, - {}, - {}, - cudf::detail::cuco_allocator{rmm::mr::polymorphic_allocator{}, stream}, - stream.value()}; + auto set = cuco::static_set{cuco::extent{compute_hash_table_size(haystack.num_rows())}, + cuco::empty_key{rhs_index_type{-1}}, + d_equal, + probing_scheme, + {}, + {}, + rmm::mr::polymorphic_allocator{}, + stream.value()}; if (haystack_has_nulls && compare_nulls == null_equality::UNEQUAL) { auto const bitmask_buffer_and_ptr = build_row_bitmask(haystack, stream); diff --git a/cpp/src/stream_compaction/distinct.cu b/cpp/src/stream_compaction/distinct.cu index aa5746c739c..e13b63e736a 100644 --- a/cpp/src/stream_compaction/distinct.cu +++ b/cpp/src/stream_compaction/distinct.cu @@ -97,16 +97,15 @@ rmm::device_uvector distinct_indices(table_view const& input, auto const helper_func = [&](auto const& d_equal) { using RowEqual = std::decay_t; - auto set = distinct_set_t{ - num_rows, - 0.5, // desired load factor - cuco::empty_key{cudf::detail::CUDF_SIZE_TYPE_SENTINEL}, - d_equal, - {row_hash.device_hasher(has_nulls)}, - {}, - {}, - cudf::detail::cuco_allocator{rmm::mr::polymorphic_allocator{}, stream}, - stream.value()}; + auto set = distinct_set_t{num_rows, + 0.5, // desired load factor + cuco::empty_key{cudf::detail::CUDF_SIZE_TYPE_SENTINEL}, + d_equal, + {row_hash.device_hasher(has_nulls)}, + {}, + {}, + rmm::mr::polymorphic_allocator{}, + stream.value()}; return detail::reduce_by_row(set, num_rows, keep, stream, mr); }; diff --git a/cpp/src/stream_compaction/distinct_count.cu b/cpp/src/stream_compaction/distinct_count.cu index 79ac9033a63..c823376fb6a 100644 --- a/cpp/src/stream_compaction/distinct_count.cu +++ b/cpp/src/stream_compaction/distinct_count.cu @@ -144,15 +144,14 @@ cudf::size_type distinct_count(table_view const& keys, auto const comparator_helper = [&](auto const row_equal) { using hasher_type = decltype(hash_key); - auto key_set = cuco::static_set{ - cuco::extent{compute_hash_table_size(num_rows)}, - cuco::empty_key{-1}, - row_equal, - cuco::linear_probing<1, hasher_type>{hash_key}, - {}, - {}, - cudf::detail::cuco_allocator{rmm::mr::polymorphic_allocator{}, stream}, - stream.value()}; + auto key_set = cuco::static_set{cuco::extent{compute_hash_table_size(num_rows)}, + cuco::empty_key{-1}, + row_equal, + cuco::linear_probing<1, hasher_type>{hash_key}, + {}, + {}, + rmm::mr::polymorphic_allocator{}, + stream.value()}; auto const iter = thrust::counting_iterator(0); // when nulls are equal, we skip hashing any row that has a null diff --git a/cpp/src/stream_compaction/distinct_helpers.hpp b/cpp/src/stream_compaction/distinct_helpers.hpp index 13270e5393c..587db410ebb 100644 --- a/cpp/src/stream_compaction/distinct_helpers.hpp +++ b/cpp/src/stream_compaction/distinct_helpers.hpp @@ -60,7 +60,7 @@ using distinct_set_t = 1, cudf::detail::row::hash::device_row_hasher>, - cudf::detail::cuco_allocator, + rmm::mr::polymorphic_allocator, cuco::storage<1>>; /** diff --git a/cpp/src/text/bpe/byte_pair_encoding.cuh b/cpp/src/text/bpe/byte_pair_encoding.cuh index 1307b694c3e..155c1c82b09 100644 --- a/cpp/src/text/bpe/byte_pair_encoding.cuh +++ b/cpp/src/text/bpe/byte_pair_encoding.cuh @@ -106,7 +106,7 @@ using merge_pairs_map_type = cuco::static_map, + rmm::mr::polymorphic_allocator, cuco_storage>; /** @@ -164,7 +164,7 @@ using mp_table_map_type = cuco::static_map, + rmm::mr::polymorphic_allocator, cuco_storage>; } // namespace detail diff --git a/cpp/src/text/bpe/load_merge_pairs.cu b/cpp/src/text/bpe/load_merge_pairs.cu index bdf36d43103..1b040748b8e 100644 --- a/cpp/src/text/bpe/load_merge_pairs.cu +++ b/cpp/src/text/bpe/load_merge_pairs.cu @@ -44,18 +44,18 @@ namespace { std::unique_ptr initialize_merge_pairs_map( cudf::column_device_view const& input, rmm::cuda_stream_view stream) { - auto const elements = input.size() / 2; - auto merge_pairs_map = std::make_unique( - static_cast(elements), - cudf::detail::CUCO_DESIRED_LOAD_FACTOR, - cuco::empty_key{-1}, - cuco::empty_value{-1}, - bpe_equal{input}, - bpe_probe_scheme{bpe_hasher{input}}, - cuco::thread_scope_device, - cuco_storage{}, - cudf::detail::cuco_allocator{rmm::mr::polymorphic_allocator{}, stream}, - stream.value()); + auto const elements = input.size() / 2; + auto merge_pairs_map = + std::make_unique(static_cast(elements), + cudf::detail::CUCO_DESIRED_LOAD_FACTOR, + cuco::empty_key{-1}, + cuco::empty_value{-1}, + bpe_equal{input}, + bpe_probe_scheme{bpe_hasher{input}}, + cuco::thread_scope_device, + cuco_storage{}, + rmm::mr::polymorphic_allocator{}, + stream.value()); auto iter = cudf::detail::make_counting_transform_iterator( 0, @@ -70,17 +70,16 @@ std::unique_ptr initialize_merge_pairs_map( std::unique_ptr initialize_mp_table_map( cudf::column_device_view const& input, rmm::cuda_stream_view stream) { - auto mp_table_map = std::make_unique( - static_cast(input.size()), - cudf::detail::CUCO_DESIRED_LOAD_FACTOR, - cuco::empty_key{-1}, - cuco::empty_value{-1}, - mp_equal{input}, - mp_probe_scheme{mp_hasher{input}}, - cuco::thread_scope_device, - cuco_storage{}, - cudf::detail::cuco_allocator{rmm::mr::polymorphic_allocator{}, stream}, - stream.value()); + auto mp_table_map = std::make_unique(static_cast(input.size()), + cudf::detail::CUCO_DESIRED_LOAD_FACTOR, + cuco::empty_key{-1}, + cuco::empty_value{-1}, + mp_equal{input}, + mp_probe_scheme{mp_hasher{input}}, + cuco::thread_scope_device, + cuco_storage{}, + rmm::mr::polymorphic_allocator{}, + stream.value()); auto iter = cudf::detail::make_counting_transform_iterator( 0, diff --git a/cpp/src/text/vocabulary_tokenize.cu b/cpp/src/text/vocabulary_tokenize.cu index 15069a8f01b..bca3b8e2bf9 100644 --- a/cpp/src/text/vocabulary_tokenize.cu +++ b/cpp/src/text/vocabulary_tokenize.cu @@ -102,7 +102,7 @@ using vocabulary_map_type = cuco::static_map, + rmm::mr::polymorphic_allocator, cuco_storage>; } // namespace } // namespace detail @@ -154,7 +154,7 @@ tokenize_vocabulary::tokenize_vocabulary(cudf::strings_column_view const& input, detail::probe_scheme{detail::vocab_hasher{*d_vocabulary}}, cuco::thread_scope_device, detail::cuco_storage{}, - cudf::detail::cuco_allocator{rmm::mr::polymorphic_allocator{}, stream}, + rmm::mr::polymorphic_allocator{}, stream.value()); // the row index is the token id (value for each key in the map) diff --git a/cpp/src/text/wordpiece_tokenize.cu b/cpp/src/text/wordpiece_tokenize.cu index b27bd379356..1e40bc4ca05 100644 --- a/cpp/src/text/wordpiece_tokenize.cu +++ b/cpp/src/text/wordpiece_tokenize.cu @@ -93,7 +93,7 @@ using vocabulary_map_type = cuco::static_map, + rmm::mr::polymorphic_allocator, cuco_storage>; /** @@ -137,7 +137,7 @@ using sub_vocabulary_map_type = cuco::static_map, + rmm::mr::polymorphic_allocator, cuco_storage>; } // namespace } // namespace detail @@ -238,7 +238,7 @@ wordpiece_vocabulary::wordpiece_vocabulary(cudf::strings_column_view const& inpu detail::probe_scheme{detail::vocab_hasher{*d_vocabulary}}, cuco::thread_scope_thread, detail::cuco_storage{}, - cudf::detail::cuco_allocator{rmm::mr::polymorphic_allocator{}, stream}, + rmm::mr::polymorphic_allocator{}, stream.value()); // the row index is the token id (data value for each key in the map) auto iter = cudf::detail::make_counting_transform_iterator(0, key_pair{}); @@ -264,7 +264,7 @@ wordpiece_vocabulary::wordpiece_vocabulary(cudf::strings_column_view const& inpu detail::sub_probe_scheme{detail::sub_vocab_hasher{*d_vocabulary}}, cuco::thread_scope_thread, detail::cuco_storage{}, - cudf::detail::cuco_allocator{rmm::mr::polymorphic_allocator{}, stream}, + rmm::mr::polymorphic_allocator{}, stream.value()); // insert them without the '##' prefix since that is how they will be looked up auto iter_sub = thrust::make_transform_iterator(sub_map_indices.begin(), key_pair{}); diff --git a/cpp/tests/io/parquet_bloom_filter_test.cu b/cpp/tests/io/parquet_bloom_filter_test.cu index d858f58fa56..4a6ce09bdb1 100644 --- a/cpp/tests/io/parquet_bloom_filter_test.cu +++ b/cpp/tests/io/parquet_bloom_filter_test.cu @@ -59,11 +59,11 @@ TEST_F(ParquetBloomFilterTest, TestStrings) cuco::extent, cuda::thread_scope_device, policy_type, - cudf::detail::cuco_allocator> + rmm::mr::polymorphic_allocator> filter{num_filter_blocks, cuco::thread_scope_device, {{cudf::DEFAULT_HASH_SEED}}, - cudf::detail::cuco_allocator{rmm::mr::polymorphic_allocator{}, stream}, + rmm::mr::polymorphic_allocator{}, stream}; // Add strings to the bloom filter