From e0f526ac3ce48df90edb02d3716898f91b3feea0 Mon Sep 17 00:00:00 2001 From: niranda perera Date: Tue, 11 Aug 2026 18:12:03 -0700 Subject: [PATCH 1/4] making preprocesses table accept mr Signed-off-by: niranda perera --- .../cudf/detail/row_operator/equality.cuh | 11 +- .../cudf/detail/row_operator/hashing.cuh | 7 +- .../row_operator/preprocessed_table.cuh | 5 +- cpp/include/cudf/dictionary/detail/encode.hpp | 8 +- cpp/include/cudf/dictionary/encode.hpp | 14 +- cpp/include/cudf_test/column_wrapper.hpp | 8 +- .../binaryop/compiled/struct_binary_ops.cuh | 5 +- cpp/src/dictionary/decode.cu | 15 +- cpp/src/dictionary/detail/concatenate.cu | 5 +- cpp/src/dictionary/encode.cu | 35 +- cpp/src/dictionary/match_keys.cu | 5 +- cpp/src/groupby/hash/groupby.cu | 9 +- cpp/src/groupby/sort/group_nunique.cu | 5 +- cpp/src/groupby/sort/group_rank_scan.cu | 23 +- .../sort/sort_helper_group_offsets.cuh | 11 +- cpp/src/groupby/streaming_groupby/insert.cuh | 7 +- cpp/src/hash/murmurhash3_x86_32.cu | 5 +- cpp/src/hash/xxhash_32.cu | 7 +- cpp/src/hash/xxhash_64.cu | 7 +- cpp/src/join/distinct_hash_join.cu | 26 +- cpp/src/join/filtered_join/filtered_join.cu | 10 +- cpp/src/join/hash_join/hash_join.cu | 5 +- cpp/src/join/hash_join/match_context.cu | 11 +- .../hash_join/partitioned_join_retrieve.cu | 9 +- cpp/src/join/hash_join/retrieve_impl.cuh | 4 +- cpp/src/join/hash_join/size_impl.cuh | 10 +- cpp/src/join/key_remapping.cu | 37 +- cpp/src/join/mark_join.cu | 8 +- cpp/src/join/mixed_join_semi.cu | 14 +- cpp/src/lists/contains.cu | 4 +- cpp/src/partitioning/partitioning.cu | 22 +- cpp/src/reductions/approx_distinct_count.cu | 11 +- cpp/src/reductions/distinct_count.cu | 13 +- cpp/src/reductions/histogram.cu | 16 +- cpp/src/reductions/scan/rank_scan.cu | 9 +- cpp/src/reductions/segmented/nunique.cu | 14 +- cpp/src/reductions/unique_count.cu | 18 +- cpp/src/reductions/unique_count_column.cu | 7 +- cpp/src/row_operator/row_operators.cu | 18 +- cpp/src/search/contains_scalar.cu | 27 +- cpp/src/search/contains_table.cu | 7 +- cpp/src/sort/rank.cu | 7 +- cpp/src/stream_compaction/distinct.cu | 4 +- cpp/src/stream_compaction/unique.cu | 11 +- cpp/src/table/table_equal.cu | 8 +- cpp/src/transform/one_hot_encode.cu | 12 +- cpp/tests/row_operator/row_operator_tests.cu | 357 ++++++++++++------ .../row_operator_tests_utilities.hpp | 18 +- .../row_operator/self_comparison_utilities.cu | 27 +- .../two_table_comparison_utilities.cu | 53 ++- .../two_table_equality_utilities.cu | 30 +- cpp/tests/utilities/column_utilities.cu | 6 +- .../utilities_tests/column_wrapper_tests.cpp | 4 - 53 files changed, 607 insertions(+), 422 deletions(-) diff --git a/cpp/include/cudf/detail/row_operator/equality.cuh b/cpp/include/cudf/detail/row_operator/equality.cuh index f38acc8b7d78..9a5fd68b6ceb 100644 --- a/cpp/include/cudf/detail/row_operator/equality.cuh +++ b/cpp/include/cudf/detail/row_operator/equality.cuh @@ -407,9 +407,12 @@ class self_comparator { * @param t The table to compare * @param stream The stream to construct this object on. Not the stream that will be used for * comparisons using this object. + * @param temp_mr Device memory resource used for temporary allocations */ - self_comparator(table_view const& t, rmm::cuda_stream_view stream) - : d_t(preprocessed_table::create(t, stream)) + self_comparator(table_view const& t, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref temp_mr) + : d_t(preprocessed_table::create(t, stream, temp_mr)) { } @@ -515,10 +518,12 @@ class two_table_comparator { * @param right The right table to compare. * @param stream The stream to construct this object on. Not the stream that will be used for * comparisons using this object. + * @param temp_mr Device memory resource used for temporary allocations */ two_table_comparator(table_view const& left, table_view const& right, - rmm::cuda_stream_view stream); + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref temp_mr); /** * @brief Construct an owning object for performing equality comparisons between two rows from two diff --git a/cpp/include/cudf/detail/row_operator/hashing.cuh b/cpp/include/cudf/detail/row_operator/hashing.cuh index e71fc5213483..5bb77b125d7c 100644 --- a/cpp/include/cudf/detail/row_operator/hashing.cuh +++ b/cpp/include/cudf/detail/row_operator/hashing.cuh @@ -240,9 +240,12 @@ class row_hasher { * @param t The table containing rows to hash * @param stream The stream to construct this object on. Not the stream that will be used for * comparisons using this object. + * @param temp_mr Device memory resource used for temporary allocations */ - row_hasher(table_view const& t, rmm::cuda_stream_view stream) - : d_t(preprocessed_table::create(t, stream)) + row_hasher(table_view const& t, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref temp_mr) + : d_t(preprocessed_table::create(t, stream, temp_mr)) { } diff --git a/cpp/include/cudf/detail/row_operator/preprocessed_table.cuh b/cpp/include/cudf/detail/row_operator/preprocessed_table.cuh index 233294201ccd..08949be1f0b7 100644 --- a/cpp/include/cudf/detail/row_operator/preprocessed_table.cuh +++ b/cpp/include/cudf/detail/row_operator/preprocessed_table.cuh @@ -10,6 +10,7 @@ #include #include #include +#include #include #include @@ -47,10 +48,12 @@ struct preprocessed_table { * * @param table The table to preprocess * @param stream The cuda stream to use while preprocessing. + * @param temp_mr Device memory resource used for temporary allocations * @return A preprocessed table as shared pointer */ static std::shared_ptr create(table_view const& table, - rmm::cuda_stream_view stream); + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref temp_mr); /** * @brief Implicit conversion operator to a `table_device_view` of the preprocessed table. diff --git a/cpp/include/cudf/dictionary/detail/encode.hpp b/cpp/include/cudf/dictionary/detail/encode.hpp index c5c0860871c9..4de1b0c7e151 100644 --- a/cpp/include/cudf/dictionary/detail/encode.hpp +++ b/cpp/include/cudf/dictionary/detail/encode.hpp @@ -37,13 +37,13 @@ namespace dictionary::detail { * @param column The column to dictionary encode. * @param indices_type The integer type to use for the indices. * @param stream CUDA stream used for device memory operations and kernel launches. - * @param mr Device memory resource used to allocate the returned column's device memory. + * @param mr Memory resources used for temporary allocations and the returned column. * @return Returns a dictionary column. */ std::unique_ptr encode(column_view const& column, data_type indices_type, rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr); + cudf::memory_resources mr); /** * @brief Create a column by gathering the keys from the provided @@ -57,12 +57,12 @@ std::unique_ptr encode(column_view const& column, * * @param dictionary_column Existing dictionary column. * @param stream CUDA stream used for device memory operations and kernel launches. - * @param mr Device memory resource used to allocate the returned column's device memory. + * @param mr Memory resources used for temporary allocations and the returned column. * @return New column with type matching the dictionary_column's keys. */ std::unique_ptr decode(dictionary_column_view const& dictionary_column, rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr); + cudf::memory_resources mr); /** * @brief Return minimal integer type for the given number of elements. diff --git a/cpp/include/cudf/dictionary/encode.hpp b/cpp/include/cudf/dictionary/encode.hpp index 555f9d85c376..16fc5ba56f5f 100644 --- a/cpp/include/cudf/dictionary/encode.hpp +++ b/cpp/include/cudf/dictionary/encode.hpp @@ -47,14 +47,14 @@ namespace dictionary { * @param column The column to dictionary encode * @param indices_type The integer type to use for the indices * @param stream CUDA stream used for device memory operations and kernel launches - * @param mr Device memory resource used to allocate the returned column's device memory + * @param mr Memory resources used for temporary allocations and the returned column * @return Returns a dictionary column */ std::unique_ptr encode( column_view const& column, - data_type indices_type = data_type{type_id::INT32}, - rmm::cuda_stream_view stream = cudf::get_default_stream(), - rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); + data_type indices_type = data_type{type_id::INT32}, + rmm::cuda_stream_view stream = cudf::get_default_stream(), + cudf::memory_resources mr = cudf::get_current_device_resource_ref()); /** * @brief Create a column by gathering the keys from the provided @@ -68,13 +68,13 @@ std::unique_ptr encode( * * @param dictionary_column Existing dictionary column * @param stream CUDA stream used for device memory operations and kernel launches - * @param mr Device memory resource used to allocate the returned column's device memory + * @param mr Memory resources used for temporary allocations and the returned column * @return New column with type matching the dictionary_column's keys */ std::unique_ptr decode( dictionary_column_view const& dictionary_column, - rmm::cuda_stream_view stream = cudf::get_default_stream(), - rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); + rmm::cuda_stream_view stream = cudf::get_default_stream(), + cudf::memory_resources mr = cudf::get_current_device_resource_ref()); /** @} */ // end of group } // namespace dictionary diff --git a/cpp/include/cudf_test/column_wrapper.hpp b/cpp/include/cudf_test/column_wrapper.hpp index 54acc5c5e7a9..77b8874861c0 100644 --- a/cpp/include/cudf_test/column_wrapper.hpp +++ b/cpp/include/cudf_test/column_wrapper.hpp @@ -1076,7 +1076,7 @@ class dictionary_column_wrapper : public detail::column_wrapper { begin, end, stream, mr.get_temporary_mr()), cudf::data_type{type_id::INT32}, stream, - mr.get_output_mr()); + mr); } /** @@ -1118,7 +1118,7 @@ class dictionary_column_wrapper : public detail::column_wrapper { begin, end, v, stream, mr.get_temporary_mr()), cudf::data_type{type_id::INT32}, stream, - mr.get_output_mr()); + mr); } /** @@ -1308,7 +1308,7 @@ class dictionary_column_wrapper : public detail::column_wrapper { cudf::dictionary::encode(strings_column_wrapper(begin, end, stream, mr.get_temporary_mr()), cudf::data_type{type_id::INT32}, stream, - mr.get_output_mr()); + mr); } /** @@ -1353,7 +1353,7 @@ class dictionary_column_wrapper : public detail::column_wrapper { cudf::dictionary::encode(strings_column_wrapper(begin, end, v, stream, mr.get_temporary_mr()), cudf::data_type{type_id::INT32}, stream, - mr.get_output_mr()); + mr); } /** diff --git a/cpp/src/binaryop/compiled/struct_binary_ops.cuh b/cpp/src/binaryop/compiled/struct_binary_ops.cuh index 18f4a193b4cd..1126c12a8365 100644 --- a/cpp/src/binaryop/compiled/struct_binary_ops.cuh +++ b/cpp/src/binaryop/compiled/struct_binary_ops.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -146,7 +146,8 @@ void apply_struct_equality_op(mutable_column_view& out, auto tlhs = table_view{{lhs}}; auto trhs = table_view{{rhs}}; - auto table_comparator = cudf::detail::row::equality::two_table_comparator{tlhs, trhs, stream}; + auto table_comparator = cudf::detail::row::equality::two_table_comparator{ + tlhs, trhs, stream, cudf::get_current_device_resource_ref()}; auto outd = column_device_view::create(out, stream); auto optional_iter = diff --git a/cpp/src/dictionary/decode.cu b/cpp/src/dictionary/decode.cu index 6265815b7487..419b3f24ac00 100644 --- a/cpp/src/dictionary/decode.cu +++ b/cpp/src/dictionary/decode.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2024, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -36,13 +36,16 @@ struct indices_handler_fn { */ std::unique_ptr decode(dictionary_column_view const& source, rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr) + cudf::memory_resources mr) { if (source.is_empty()) return make_empty_column(type_id::EMPTY); + auto const output_mr = mr.get_output_mr(); + auto const temp_mr = mr.get_temporary_mr(); + // annotated indices include the offset, size and bitmask from it's parent auto const indices = source.get_indices_annotated(); - auto const d_indices = column_device_view::create(indices, stream); + auto const d_indices = column_device_view::create(indices, stream, temp_mr); auto const d_iterator = cudf::detail::indexalator_factory::make_input_iterator(indices); auto const indices_begin = cudf::detail::make_counting_transform_iterator( 0, indices_handler_fn{d_iterator, *d_indices, source.keys().size()}); @@ -52,12 +55,12 @@ std::unique_ptr decode(dictionary_column_view const& source, indices_begin + source.size(), cudf::out_of_bounds_policy::NULLIFY, stream, - mr) + output_mr) ->release(); auto output_column = std::unique_ptr(std::move(table_column.front())); // apply any nulls to the output column - output_column->set_null_mask(cudf::detail::copy_bitmask(source.parent(), stream, mr), + output_column->set_null_mask(cudf::detail::copy_bitmask(source.parent(), stream, output_mr), source.null_count()); return output_column; @@ -67,7 +70,7 @@ std::unique_ptr decode(dictionary_column_view const& source, std::unique_ptr decode(dictionary_column_view const& source, rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr) + cudf::memory_resources mr) { CUDF_FUNC_RANGE(); return detail::decode(source, stream, mr); diff --git a/cpp/src/dictionary/detail/concatenate.cu b/cpp/src/dictionary/detail/concatenate.cu index 9b2b275aa9c9..d374a847cdbb 100644 --- a/cpp/src/dictionary/detail/concatenate.cu +++ b/cpp/src/dictionary/detail/concatenate.cu @@ -187,8 +187,9 @@ std::unique_ptr concatenate(host_span columns, cudf::detail::row::hash::device_row_hasher>; auto const tv = cudf::table_view({all_keys->view()}); - auto const row_hash = cudf::detail::row::hash::row_hasher(tv, stream); - auto const row_equal = cudf::detail::row::equality::self_comparator(tv, stream); + auto const temp_mr = cudf::get_current_device_resource_ref(); + auto const row_hash = cudf::detail::row::hash::row_hasher(tv, stream, temp_mr); + auto const row_equal = cudf::detail::row::equality::self_comparator(tv, stream, temp_mr); auto const comparator = cudf::detail::row::equality::nan_equal_physical_equality_comparator{}; auto const d_equal = row_equal.equal_to(cudf::nullate::NO{}, null_equality::EQUAL, comparator); diff --git a/cpp/src/dictionary/encode.cu b/cpp/src/dictionary/encode.cu index a43d18c080a4..e279cac55fdb 100644 --- a/cpp/src/dictionary/encode.cu +++ b/cpp/src/dictionary/encode.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -56,7 +56,7 @@ struct encode_fn { std::unique_ptr encode(column_view const& input, data_type indices_type, rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr) + cudf::memory_resources mr) { CUDF_EXPECTS(is_signed(indices_type) && is_index_type(indices_type), "indices must be type signed integer", @@ -68,8 +68,11 @@ std::unique_ptr encode(column_view const& input, "encoding nested types not supported", std::invalid_argument); + auto const output_mr = mr.get_output_mr(); + auto const temp_mr = mr.get_temporary_mr(); + auto indices_column = cudf::make_numeric_column( - indices_type, input.size(), cudf::mask_state::UNALLOCATED, stream, mr); + indices_type, input.size(), cudf::mask_state::UNALLOCATED, stream, output_mr); if (input.is_empty()) { return make_dictionary_column( make_empty_column(input.type()), std::move(indices_column), rmm::device_buffer{}, 0); @@ -82,13 +85,13 @@ std::unique_ptr encode(column_view const& input, auto const has_nulls = nullate::DYNAMIC{input.has_nulls()}; auto const tv = cudf::table_view({input}); - auto const row_hash = cudf::detail::row::hash::row_hasher(tv, stream); - auto const row_equal = cudf::detail::row::equality::self_comparator(tv, stream); + auto const row_hash = cudf::detail::row::hash::row_hasher(tv, stream, temp_mr); + auto const row_equal = cudf::detail::row::equality::self_comparator(tv, stream, temp_mr); auto const comparator = cudf::detail::row::equality::nan_equal_physical_equality_comparator{}; auto const d_equal = row_equal.equal_to(has_nulls, null_equality::EQUAL, comparator); auto const empty_key = cuco::empty_key{cudf::detail::CUDF_SIZE_TYPE_SENTINEL}; auto probe = encode_probe_t{row_hash.device_hasher(has_nulls)}; - auto allocator = rmm::mr::polymorphic_allocator{}; + auto allocator = rmm::mr::polymorphic_allocator{temp_mr}; auto set = cuco::static_set{ input.size(), 0.5, empty_key, d_equal, probe, {}, {}, allocator, stream.value()}; auto set_ref = set.ref(cuco::insert_and_find); @@ -96,35 +99,33 @@ std::unique_ptr encode(column_view const& input, // build a static_set of the input values // and keep track of the indices of the unique values - auto d_indices = rmm::device_uvector(input.size(), stream); - auto d_input = column_device_view::create(input, stream); - thrust::transform(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + auto d_indices = rmm::device_uvector(input.size(), stream, temp_mr); + auto d_input = column_device_view::create(input, stream, temp_mr); + thrust::transform(rmm::exec_policy_nosync(stream, temp_mr), cuda::counting_iterator{0}, cuda::counting_iterator{input.size()}, d_indices.begin(), encode_fn{set_ref, *d_input}); - auto keys_indices = rmm::device_uvector(input.size(), stream); + auto keys_indices = rmm::device_uvector(input.size(), stream, temp_mr); auto keys_end = set.retrieve_all(keys_indices.begin(), stream.value()); keys_indices.resize(cuda::std::distance(keys_indices.begin(), keys_end), stream); // sort the keys_indices so we can use lower-bound on them - thrust::sort(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), - keys_indices.begin(), - keys_indices.end()); + thrust::sort(rmm::exec_policy_nosync(stream, temp_mr), keys_indices.begin(), keys_indices.end()); // use keys_indices to retrieve the keys auto const oob_policy = cudf::out_of_bounds_policy::DONT_CHECK; auto const index_policy = cudf::negative_index_policy::NOT_ALLOWED; auto keys_column = - std::move(cudf::detail::gather(tv, keys_indices, oob_policy, index_policy, stream, mr) + std::move(cudf::detail::gather(tv, keys_indices, oob_policy, index_policy, stream, output_mr) ->release() .front()); // call lower-bound with keys_indices and d_indices to get the output indices_column auto d_result = cudf::detail::indexalator_factory::make_output_iterator(indices_column->mutable_view()); - thrust::lower_bound(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + thrust::lower_bound(rmm::exec_policy_nosync(stream, temp_mr), keys_indices.begin(), keys_indices.end(), d_indices.begin(), @@ -134,7 +135,7 @@ std::unique_ptr encode(column_view const& input, // create column with keys_column and indices_column return make_dictionary_column(std::move(keys_column), std::move(indices_column), - cudf::detail::copy_bitmask(input, stream, mr), + cudf::detail::copy_bitmask(input, stream, output_mr), input.null_count()); } @@ -155,7 +156,7 @@ data_type get_indices_type_for_size(size_type keys_size) std::unique_ptr encode(column_view const& input_column, data_type indices_type, rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr) + cudf::memory_resources mr) { CUDF_FUNC_RANGE(); return detail::encode(input_column, indices_type, stream, mr); diff --git a/cpp/src/dictionary/match_keys.cu b/cpp/src/dictionary/match_keys.cu index 3637d231d41d..f0b9bfbc099e 100644 --- a/cpp/src/dictionary/match_keys.cu +++ b/cpp/src/dictionary/match_keys.cu @@ -50,8 +50,9 @@ struct unique_keys_dispatch_fn { auto const has_nulls = nullate::DYNAMIC{false}; auto const keys_tv = table_view({all_keys}); - auto const row_hash = cudf::detail::row::hash::row_hasher(keys_tv, stream); - auto const row_equal = cudf::detail::row::equality::self_comparator(keys_tv, stream); + auto const temp_mr = cudf::get_current_device_resource_ref(); + auto const row_hash = cudf::detail::row::hash::row_hasher(keys_tv, stream, temp_mr); + auto const row_equal = cudf::detail::row::equality::self_comparator(keys_tv, stream, temp_mr); auto const comparator = cudf::detail::row::equality::nan_equal_physical_equality_comparator{}; auto const d_equal = row_equal.equal_to(has_nulls, null_equality::EQUAL, comparator); auto const empty_key = cuco::empty_key{cudf::detail::CUDF_SIZE_TYPE_SENTINEL}; diff --git a/cpp/src/groupby/hash/groupby.cu b/cpp/src/groupby/hash/groupby.cu index 161b0384537b..f78669799a84 100644 --- a/cpp/src/groupby/hash/groupby.cu +++ b/cpp/src/groupby/hash/groupby.cu @@ -42,10 +42,11 @@ std::unique_ptr dispatch_groupby(table_view const& keys, auto const has_null = nullate::DYNAMIC{cudf::has_nested_nulls(keys)}; auto const skip_rows_with_nulls = keys_have_nulls and include_null_keys == null_policy::EXCLUDE; - auto preprocessed_keys = cudf::detail::row::hash::preprocessed_table::create(keys, stream); - auto const comparator = cudf::detail::row::equality::self_comparator{preprocessed_keys}; - auto const row_hash = cudf::detail::row::hash::row_hasher{std::move(preprocessed_keys)}; - auto const d_row_hash = row_hash.device_hasher(has_null); + auto preprocessed_keys = cudf::detail::row::hash::preprocessed_table::create( + keys, stream, cudf::get_current_device_resource_ref()); + auto const comparator = cudf::detail::row::equality::self_comparator{preprocessed_keys}; + auto const row_hash = cudf::detail::row::hash::row_hasher{std::move(preprocessed_keys)}; + auto const d_row_hash = row_hash.device_hasher(has_null); if (cudf::detail::has_nested_columns(keys)) { auto const d_row_equal = comparator.equal_to(has_null, null_keys_are_equal); diff --git a/cpp/src/groupby/sort/group_nunique.cu b/cpp/src/groupby/sort/group_nunique.cu index 96aaa4cace1c..2df7c9cf3205 100644 --- a/cpp/src/groupby/sort/group_nunique.cu +++ b/cpp/src/groupby/sort/group_nunique.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -79,7 +79,8 @@ std::unique_ptr group_nunique(column_view const& values, if (num_groups == 0) { return result; } auto const values_view = table_view{{values}}; - auto const comparator = cudf::detail::row::equality::self_comparator{values_view, stream}; + auto const comparator = cudf::detail::row::equality::self_comparator{ + values_view, stream, cudf::get_current_device_resource_ref()}; auto const d_values_view = column_device_view::create(values, stream); diff --git a/cpp/src/groupby/sort/group_rank_scan.cu b/cpp/src/groupby/sort/group_rank_scan.cu index 31ed09ff8cd7..2c96e45fc79f 100644 --- a/cpp/src/groupby/sort/group_rank_scan.cu +++ b/cpp/src/groupby/sort/group_rank_scan.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -94,7 +94,9 @@ std::unique_ptr rank_generator(column_view const& grouped_values, rmm::device_async_resource_ref mr) { auto const grouped_values_view = table_view{{grouped_values}}; - auto const comparator = cudf::detail::row::equality::self_comparator{grouped_values_view, stream}; + auto const temp_mr = cudf::get_current_device_resource_ref(); + auto const comparator = + cudf::detail::row::equality::self_comparator{grouped_values_view, stream, temp_mr}; auto ranks = make_fixed_width_column( data_type{type_to_id()}, grouped_values.size(), mask_state::UNALLOCATED, stream, mr); @@ -104,7 +106,7 @@ std::unique_ptr rank_generator(column_view const& grouped_values, auto const permuted_equal = permuted_row_equality_comparator(d_equal, value_order.begin()); - thrust::transform(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + thrust::transform(rmm::exec_policy_nosync(stream, temp_mr), cuda::counting_iterator(0), cuda::counting_iterator(grouped_values.size()), mutable_ranks.begin(), @@ -130,14 +132,13 @@ std::unique_ptr rank_generator(column_view const& grouped_values, cuda::std::reverse_iterator(mutable_ranks.end())}; } }(); - thrust::inclusive_scan_by_key( - rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), - group_labels_begin, - group_labels_begin + group_labels.size(), - mutable_rank_begin, - mutable_rank_begin, - cuda::std::equal_to{}, - scan_op); + thrust::inclusive_scan_by_key(rmm::exec_policy_nosync(stream, temp_mr), + group_labels_begin, + group_labels_begin + group_labels.size(), + mutable_rank_begin, + mutable_rank_begin, + cuda::std::equal_to{}, + scan_op); return ranks; } } // namespace diff --git a/cpp/src/groupby/sort/sort_helper_group_offsets.cuh b/cpp/src/groupby/sort/sort_helper_group_offsets.cuh index 958d424e2bf4..ddd2084c4d9b 100644 --- a/cpp/src/groupby/sort/sort_helper_group_offsets.cuh +++ b/cpp/src/groupby/sort/sort_helper_group_offsets.cuh @@ -38,22 +38,19 @@ size_type compute_group_offsets(table_view const& keys, rmm::device_uvector& group_offsets, rmm::cuda_stream_view stream) { - auto const comparator = cudf::detail::row::equality::self_comparator{keys, stream}; + auto const temp_mr = cudf::get_current_device_resource_ref(); + auto const comparator = cudf::detail::row::equality::self_comparator{keys, stream, temp_mr}; auto const d_key_equal = comparator.equal_to( cudf::nullate::DYNAMIC{cudf::has_nested_nulls(keys)}, null_equality::EQUAL); // Using a temporary buffer for intermediate transform results from the iterator containing // the comparator speeds up compile-time significantly without much degradation in // runtime performance over using the comparator directly in thrust::unique_copy. - auto result = rmm::device_uvector(size, stream); + auto result = rmm::device_uvector(size, stream, temp_mr); auto const itr = cuda::counting_iterator{0}; auto const row_eq = permuted_row_equality_comparator(d_key_equal, sorted_order); auto const ufn = cudf::detail::unique_copy_fn{ itr, duplicate_keep_option::KEEP_FIRST, row_eq, size - 1}; - thrust::transform(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), - itr, - itr + size, - result.begin(), - ufn); + thrust::transform(rmm::exec_policy_nosync(stream, temp_mr), itr, itr + size, result.begin(), ufn); auto const result_end = cudf::detail::copy_if( itr, itr + size, result.begin(), group_offsets.begin(), cuda::std::identity{}, stream); return cuda::std::distance(group_offsets.begin(), result_end); diff --git a/cpp/src/groupby/streaming_groupby/insert.cuh b/cpp/src/groupby/streaming_groupby/insert.cuh index 39be36b4af7e..4c51f0d9eefd 100644 --- a/cpp/src/groupby/streaming_groupby/insert.cuh +++ b/cpp/src/groupby/streaming_groupby/insert.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -34,7 +34,8 @@ streaming_groupby::impl::batch_insert_result streaming_groupby::impl::probe_and_ auto const has_null = cudf::nullate::DYNAMIC{_has_nullable_keys}; // Preprocess batch for row operators. - auto preprocessed_batch = cudf::detail::row::hash::preprocessed_table::create(batch_keys, stream); + auto preprocessed_batch = + cudf::detail::row::hash::preprocessed_table::create(batch_keys, stream, temp_mr); auto const batch_hasher_obj = cudf::detail::row::hash::row_hasher{preprocessed_batch}; auto const d_batch_hash = batch_hasher_obj.device_hasher(has_null); @@ -110,7 +111,7 @@ streaming_groupby::impl::batch_insert_result streaming_groupby::impl::probe_and_ temp_mr); auto preprocessed_compacted = - cudf::detail::row::hash::preprocessed_table::create(compacted->view(), stream); + cudf::detail::row::hash::preprocessed_table::create(compacted->view(), stream, temp_mr); // Store the compacted batch. auto const new_batch_id = static_cast(_compacted_batches.size()); diff --git a/cpp/src/hash/murmurhash3_x86_32.cu b/cpp/src/hash/murmurhash3_x86_32.cu index f82d552456d7..4e1f5ced36d5 100644 --- a/cpp/src/hash/murmurhash3_x86_32.cu +++ b/cpp/src/hash/murmurhash3_x86_32.cu @@ -15,6 +15,7 @@ #include #include #include +#include #include #include @@ -66,8 +67,8 @@ std::unique_ptr murmurhash3_x86_32(table_view const& input, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) { - auto const preprocessed_input = - cudf::detail::row::hash::preprocessed_table::create(input, stream); + auto const preprocessed_input = cudf::detail::row::hash::preprocessed_table::create( + input, stream, cudf::get_current_device_resource_ref()); return murmurhash3_x86_32_impl( preprocessed_input, input.num_rows(), seed, nullate::DYNAMIC{has_nulls(input)}, stream, mr); } diff --git a/cpp/src/hash/xxhash_32.cu b/cpp/src/hash/xxhash_32.cu index 759a491b193b..0efd6f02ce12 100644 --- a/cpp/src/hash/xxhash_32.cu +++ b/cpp/src/hash/xxhash_32.cu @@ -30,9 +30,10 @@ std::unique_ptr xxhash_32(table_view const& input, if (input.num_rows() == 0) { return output; } - bool const nullable = has_nulls(input); - auto const row_hasher = cudf::detail::row::hash::row_hasher(input, stream); - auto output_view = output->mutable_view(); + bool const nullable = has_nulls(input); + auto const row_hasher = + cudf::detail::row::hash::row_hasher(input, stream, cudf::get_current_device_resource_ref()); + auto output_view = output->mutable_view(); // Compute the hash value for each row auto const output_begin = output_view.begin(); diff --git a/cpp/src/hash/xxhash_64.cu b/cpp/src/hash/xxhash_64.cu index fcf7009bd128..69c97724bbd6 100644 --- a/cpp/src/hash/xxhash_64.cu +++ b/cpp/src/hash/xxhash_64.cu @@ -32,9 +32,10 @@ std::unique_ptr xxhash_64(table_view const& input, if (input.num_rows() == 0) { return output; } - bool const nullable = has_nulls(input); - auto const row_hasher = cudf::detail::row::hash::row_hasher(input, stream); - auto output_view = output->mutable_view(); + bool const nullable = has_nulls(input); + auto const row_hasher = + cudf::detail::row::hash::row_hasher(input, stream, cudf::get_current_device_resource_ref()); + auto output_view = output->mutable_view(); // Compute the hash value for each row auto const output_begin = output_view.begin(); diff --git a/cpp/src/join/distinct_hash_join.cu b/cpp/src/join/distinct_hash_join.cu index 8a5cf7b2a279..d08e3da7406c 100644 --- a/cpp/src/join/distinct_hash_join.cu +++ b/cpp/src/join/distinct_hash_join.cu @@ -128,9 +128,9 @@ void find_matches_in_hash_table(HashTableType const& hash_table, hash_table.find_async( iter, iter + left_table_num_rows, d_equal, hasher, found_begin, stream.value()); } else { - auto stencil = cuda::counting_iterator{0}; - auto const row_bitmask = - cudf::detail::bitmask_and(left, stream, cudf::get_current_device_resource_ref()).first; + auto const temp_mr = cudf::get_current_device_resource_ref(); + auto stencil = cuda::counting_iterator{0}; + auto const row_bitmask = cudf::detail::bitmask_and(left, stream, temp_mr).first; auto const pred = cudf::detail::row_is_valid{reinterpret_cast(row_bitmask.data())}; @@ -163,7 +163,8 @@ distinct_hash_join::distinct_hash_join(cudf::table_view const& right, : _has_nested_columns{cudf::has_nested_columns(right)}, _nulls_equal{compare_nulls}, _right{right}, - _preprocessed_right{cudf::detail::row::equality::preprocessed_table::create(_right, stream)}, + _preprocessed_right{cudf::detail::row::equality::preprocessed_table::create( + _right, stream, cudf::get_current_device_resource_ref())}, _hash_table{cuco::extent{static_cast(right.num_rows())}, checked_load_factor(load_factor), cuco::empty_key{cuco::pair{std::numeric_limits::max(), @@ -181,13 +182,14 @@ distinct_hash_join::distinct_hash_join(cudf::table_view const& right, if (right_table_num_rows == 0) { return; } + auto const temp_mr = cudf::get_current_device_resource_ref(); + auto const build_hash_table = [&](auto iter) { if (this->_nulls_equal == cudf::null_equality::EQUAL or (not cudf::nullable(right))) { this->_hash_table.insert_async(iter, iter + right_table_num_rows, stream.value()); } else { - auto stencil = cuda::counting_iterator{0}; - auto const row_bitmask = - cudf::detail::bitmask_and(_right, stream, cudf::get_current_device_resource_ref()).first; + auto stencil = cuda::counting_iterator{0}; + auto const row_bitmask = cudf::detail::bitmask_and(_right, stream, temp_mr).first; auto const pred = cudf::detail::row_is_valid{reinterpret_cast(row_bitmask.data())}; @@ -240,7 +242,9 @@ distinct_hash_join::inner_join(cudf::table_view const& left, auto found_indices = rmm::device_uvector(left_table_num_rows, stream); auto const found_begin = cuda::make_transform_output_iterator(found_indices.begin(), output_fn{}); - auto preprocessed_left = cudf::detail::row::equality::preprocessed_table::create(left, stream); + auto const temp_mr = cudf::get_current_device_resource_ref(); + auto preprocessed_left = + cudf::detail::row::equality::preprocessed_table::create(left, stream, temp_mr); if (cudf::detail::is_primitive_row_op_compatible(_right)) { auto const d_hasher = cudf::detail::row::primitive::row_hasher{nullate::DYNAMIC{has_nulls}, preprocessed_left}; @@ -333,7 +337,9 @@ std::unique_ptr> distinct_hash_join::left_join( auto const output_begin = cuda::make_transform_output_iterator(right_indices->begin(), output_fn{}); - auto preprocessed_left = cudf::detail::row::equality::preprocessed_table::create(left, stream); + auto const temp_mr = cudf::get_current_device_resource_ref(); + auto preprocessed_left = + cudf::detail::row::equality::preprocessed_table::create(left, stream, temp_mr); if (cudf::detail::is_primitive_row_op_compatible(_right)) { auto const d_hasher = @@ -355,7 +361,7 @@ std::unique_ptr> distinct_hash_join::left_join( } else { // If right table is empty, return left table if (this->_right.num_rows() == 0) { - thrust::fill(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + thrust::fill(rmm::exec_policy_nosync(stream, temp_mr), right_indices->begin(), right_indices->end(), cudf::JoinNoMatch); diff --git a/cpp/src/join/filtered_join/filtered_join.cu b/cpp/src/join/filtered_join/filtered_join.cu index d1766f0ffab3..024d000be34f 100644 --- a/cpp/src/join/filtered_join/filtered_join.cu +++ b/cpp/src/join/filtered_join/filtered_join.cu @@ -103,11 +103,12 @@ filtered_join::filtered_join(cudf::table_view const& right, : _right_mode{select_row_operator_mode(right)}, _bucket_storage{cuco::extent{compute_bucket_storage_size( right.num_rows(), checked_load_factor(load_factor), _right_mode)}, - rmm::mr::polymorphic_allocator{std::move(mr)}, + rmm::mr::polymorphic_allocator{mr}, stream.value()}, _right{right}, _nulls_equal{compare_nulls}, - _preprocessed_right{cudf::detail::row::equality::preprocessed_table::create(_right, stream)} + _preprocessed_right{cudf::detail::row::equality::preprocessed_table::create( + _right, stream, cudf::get_current_device_resource_ref())} { cudf::scoped_range range{"filtered_join::filtered_join"}; if (_right.num_rows() == 0) return; @@ -129,9 +130,10 @@ std::unique_ptr> filtered_join::semi_anti_j { cudf::scoped_range range{"filtered_join::semi_anti_join"}; - auto const preprocessed_left = [&left, stream] { + auto const preprocessed_left = [&left, stream, mr] { cudf::scoped_range range{"filtered_join::semi_anti_join::preprocessed_left"}; - return cudf::detail::row::equality::preprocessed_table::create(left, stream); + return cudf::detail::row::equality::preprocessed_table::create( + left, stream, cudf::get_current_device_resource_ref()); }(); auto contains_map = rmm::device_uvector(left.num_rows(), stream); diff --git a/cpp/src/join/hash_join/hash_join.cu b/cpp/src/join/hash_join/hash_join.cu index b699e04fff42..e0f7b8d2d87b 100644 --- a/cpp/src/join/hash_join/hash_join.cu +++ b/cpp/src/join/hash_join/hash_join.cu @@ -127,10 +127,11 @@ hash_join::hash_join(cudf::table_view const& right, {}, {}, {}, - rmm::mr::polymorphic_allocator{std::move(mr)}, + rmm::mr::polymorphic_allocator{mr}, stream.value()}})}, _right{right}, - _preprocessed_right{cudf::detail::row::equality::preprocessed_table::create(_right, stream)} + _preprocessed_right{cudf::detail::row::equality::preprocessed_table::create( + _right, stream, cudf::get_current_device_resource_ref())} { CUDF_FUNC_RANGE(); CUDF_EXPECTS(0 != right.num_columns(), "Hash join right table is empty", std::invalid_argument); diff --git a/cpp/src/join/hash_join/match_context.cu b/cpp/src/join/hash_join/match_context.cu index 5fc2dd5ba9cf..dbe44e01649d 100644 --- a/cpp/src/join/hash_join/match_context.cu +++ b/cpp/src/join/hash_join/match_context.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -32,9 +32,10 @@ std::unique_ptr> make_join_match_counts( rmm::device_async_resource_ref mr) { auto match_counts = std::make_unique>(left.num_rows(), stream, mr); + auto const temp_mr = cudf::get_current_device_resource_ref(); if (is_empty) { - thrust::fill(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + thrust::fill(rmm::exec_policy_nosync(stream, temp_mr), match_counts->begin(), match_counts->end(), join == join_kind::INNER_JOIN ? 0 : 1); @@ -46,14 +47,14 @@ std::unique_ptr> make_join_match_counts( std::invalid_argument); auto const preprocessed_left = - cudf::detail::row::equality::preprocessed_table::create(left, stream); + cudf::detail::row::equality::preprocessed_table::create(left, stream, temp_mr); auto const left_table_num_rows = left.num_rows(); auto count_matches = [&](auto equality, auto d_hasher) { // Precompute left keys: {hash(row_idx), row_idx} for each left row. auto const n = static_cast(left_table_num_rows); - rmm::device_uvector left_keys(n, stream); - thrust::transform(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + rmm::device_uvector left_keys(n, stream, temp_mr); + thrust::transform(rmm::exec_policy_nosync(stream, temp_mr), cuda::counting_iterator(0), cuda::counting_iterator(left_table_num_rows), left_keys.begin(), diff --git a/cpp/src/join/hash_join/partitioned_join_retrieve.cu b/cpp/src/join/hash_join/partitioned_join_retrieve.cu index 77dcae1e9325..af66590c39a8 100644 --- a/cpp/src/join/hash_join/partitioned_join_retrieve.cu +++ b/cpp/src/join/hash_join/partitioned_join_retrieve.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -106,8 +106,9 @@ hash_join::partitioned_join_retrieve(join_kind join, validate_hash_join_probe(_right, left_partition_view, _has_nulls); + auto const temp_mr = cudf::get_current_device_resource_ref(); auto const preprocessed_left = - cudf::detail::row::equality::preprocessed_table::create(left_partition_view, stream); + cudf::detail::row::equality::preprocessed_table::create(left_partition_view, stream, temp_mr); // For FULL_JOIN, probe with LEFT_JOIN semantics (no complement here) bool const is_outer = (join != join_kind::INNER_JOIN); @@ -123,8 +124,8 @@ hash_join::partitioned_join_retrieve(join_kind join, auto retrieve_partition = [&](auto equality, auto d_hasher) { // Precompute left keys for this partition slice. - rmm::device_uvector left_keys(n, stream); - thrust::transform(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + rmm::device_uvector left_keys(n, stream, temp_mr); + thrust::transform(rmm::exec_policy_nosync(stream, temp_mr), cuda::counting_iterator(0), cuda::counting_iterator(partition_size), left_keys.begin(), diff --git a/cpp/src/join/hash_join/retrieve_impl.cuh b/cpp/src/join/hash_join/retrieve_impl.cuh index 5efe69afe850..5c7fefeb1cc0 100644 --- a/cpp/src/join/hash_join/retrieve_impl.cuh +++ b/cpp/src/join/hash_join/retrieve_impl.cuh @@ -173,8 +173,8 @@ hash_join::join_retrieve(cudf::table_view const& left, } } - auto const preprocessed_left = - cudf::detail::row::equality::preprocessed_table::create(left, stream); + auto const preprocessed_left = cudf::detail::row::equality::preprocessed_table::create( + left, stream, cudf::get_current_device_resource_ref()); auto join_indices = cudf::detail::probe_join_hash_table(_right, left, diff --git a/cpp/src/join/hash_join/size_impl.cuh b/cpp/src/join/hash_join/size_impl.cuh index 3e20ebc7367e..48218accc3f0 100644 --- a/cpp/src/join/hash_join/size_impl.cuh +++ b/cpp/src/join/hash_join/size_impl.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -80,8 +80,8 @@ std::size_t hash_join::join_size(cudf::table_view const& left, "Left table has nulls while right table was not hashed with null check.", std::invalid_argument); - auto const preprocessed_left = - cudf::detail::row::equality::preprocessed_table::create(left, stream); + auto const preprocessed_left = cudf::detail::row::equality::preprocessed_table::create( + left, stream, cudf::get_current_device_resource_ref()); return cudf::detail::compute_join_output_size(_right, left, @@ -109,8 +109,8 @@ std::size_t hash_join::join_size(cudf::table_view const& left, "Left table has nulls while right table was not hashed with null check.", std::invalid_argument); - auto const preprocessed_left = - cudf::detail::row::equality::preprocessed_table::create(left, stream); + auto const preprocessed_left = cudf::detail::row::equality::preprocessed_table::create( + left, stream, cudf::get_current_device_resource_ref()); return cudf::detail::get_full_join_size(_right, left, diff --git a/cpp/src/join/key_remapping.cu b/cpp/src/join/key_remapping.cu index ef8540c261f4..a1ea75b7f3f2 100644 --- a/cpp/src/join/key_remapping.cu +++ b/cpp/src/join/key_remapping.cu @@ -305,6 +305,8 @@ class key_remap_table : public key_remap_table_interface { cudf::size_type const right_num_rows{_right.num_rows()}; if (right_num_rows == 0) { return; } + auto const temp_mr = cudf::get_current_device_resource_ref(); + auto const key_iter = cudf::detail::make_counting_transform_iterator( 0, make_key_pair{row_hasher}); @@ -312,9 +314,7 @@ class key_remap_table : public key_remap_table_interface { (_compare_nulls == cudf::null_equality::UNEQUAL) && cudf::nullable(right); auto const row_bitmask = - skip_nulls - ? cudf::detail::bitmask_and(_right, stream, cudf::get_current_device_resource_ref()).first - : rmm::device_buffer{}; + skip_nulls ? cudf::detail::bitmask_and(_right, stream, temp_mr).first : rmm::device_buffer{}; auto const bitmask_ptr = skip_nulls ? reinterpret_cast(row_bitmask.data()) : nullptr; @@ -325,7 +325,7 @@ class key_remap_table : public key_remap_table_interface { // No metrics - simple insert auto set_ref = _hash_table.ref(cuco::op::insert); thrust::for_each_n( - rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + rmm::exec_policy_nosync(stream, temp_mr), cuda::counting_iterator{0}, right_num_rows, insert_only_fn{set_ref, key_iter, bitmask_ptr}); @@ -339,14 +339,11 @@ class key_remap_table : public key_remap_table_interface { cudf::bitmask_type const* bitmask_ptr, rmm::cuda_stream_view stream) { + auto const temp_mr = cudf::get_current_device_resource_ref(); rmm::device_uvector counts(right_num_rows, stream); - thrust::fill(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), - counts.begin(), - counts.end(), - 0); + thrust::fill(rmm::exec_policy_nosync(stream, temp_mr), counts.begin(), counts.end(), 0); - cudf::detail::device_scalar d_distinct_count{ - 0, stream, cudf::get_current_device_resource_ref()}; + cudf::detail::device_scalar d_distinct_count{0, stream, temp_mr}; auto set_ref = _hash_table.ref(cuco::op::insert_and_find); @@ -358,12 +355,11 @@ class key_remap_table : public key_remap_table_interface { _distinct_count = d_distinct_count.value(stream); - _max_duplicate_count = - thrust::reduce(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), - counts.begin(), - counts.end(), - cudf::size_type{0}, - cuda::maximum{}); + _max_duplicate_count = thrust::reduce(rmm::exec_policy_nosync(stream, temp_mr), + counts.begin(), + counts.end(), + cudf::size_type{0}, + cuda::maximum{}); } public: @@ -380,10 +376,12 @@ class key_remap_table : public key_remap_table_interface { return std::make_unique>(0, stream, mr); } + auto const temp_mr = cudf::get_current_device_resource_ref(); + if (this->_right.num_rows() == 0) { auto result = std::make_unique>(left_num_rows, stream, mr); - thrust::fill(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + thrust::fill(rmm::exec_policy_nosync(stream, temp_mr), result->begin(), result->end(), cudf::JoinNoMatch); @@ -395,7 +393,7 @@ class key_remap_table : public key_remap_table_interface { cuda::make_transform_output_iterator(result->begin(), extract_index{}); auto preprocessed_left = - cudf::detail::row::equality::preprocessed_table::create(left_keys, stream); + cudf::detail::row::equality::preprocessed_table::create(left_keys, stream, temp_mr); if (cudf::detail::is_primitive_row_op_compatible(_right)) { auto const d_hasher = cudf::detail::row::primitive::row_hasher{ @@ -500,7 +498,8 @@ std::unique_ptr create_key_remap_table( if (right.num_rows() == 0 || right.num_columns() == 0) { return nullptr; } - auto preprocessed_right = cudf::detail::row::equality::preprocessed_table::create(right, stream); + auto preprocessed_right = cudf::detail::row::equality::preprocessed_table::create( + right, stream, cudf::get_current_device_resource_ref()); if (cudf::detail::is_primitive_row_op_compatible(right)) { auto const d_hasher = cudf::detail::row::primitive::row_hasher{ diff --git a/cpp/src/join/mark_join.cu b/cpp/src/join/mark_join.cu index c050170e3853..94d531a2fb50 100644 --- a/cpp/src/join/mark_join.cu +++ b/cpp/src/join/mark_join.cu @@ -597,7 +597,8 @@ mark_join::mark_join(cudf::table_view const& left, _left{left}, _nulls_equal{compare_nulls}, _prefilter{prefilter}, - _preprocessed_left{cudf::detail::row::equality::preprocessed_table::create(left, stream)}, + _preprocessed_left{cudf::detail::row::equality::preprocessed_table::create( + left, stream, cudf::get_current_device_resource_ref())}, _bucket_storage{ cuco::extent{compute_mark_join_capacity(left, checked_load_factor(load_factor))}, rmm::mr::polymorphic_allocator{mr}, @@ -742,9 +743,10 @@ std::unique_ptr> mark_join::semi_anti_join( { clear_marks(stream); - auto const preprocessed_right = [&right, stream] { + auto const preprocessed_right = [&right, stream, mr] { cudf::scoped_range range{"mark_join::semi_anti_join::preprocessed_right"}; - return cudf::detail::row::equality::preprocessed_table::create(right, stream); + return cudf::detail::row::equality::preprocessed_table::create( + right, stream, cudf::get_current_device_resource_ref()); }(); if (is_primitive_row_op_compatible(_left)) { diff --git a/cpp/src/join/mixed_join_semi.cu b/cpp/src/join/mixed_join_semi.cu index 8021cdab21bc..0b44b2b6c3ed 100644 --- a/cpp/src/join/mixed_join_semi.cu +++ b/cpp/src/join/mixed_join_semi.cu @@ -105,10 +105,11 @@ std::unique_ptr> mixed_join_semi( auto left_conditional_view = table_device_view::create(left_conditional, stream); auto right_conditional_view = table_device_view::create(right_conditional, stream); + auto const temp_mr = cudf::get_current_device_resource_ref(); auto const preprocessed_right = - cudf::detail::row::equality::preprocessed_table::create(right, stream); + cudf::detail::row::equality::preprocessed_table::create(right, stream, temp_mr); auto const preprocessed_left = - cudf::detail::row::equality::preprocessed_table::create(left, stream); + cudf::detail::row::equality::preprocessed_table::create(left, stream, temp_mr); auto const row_comparator = cudf::detail::row::equality::two_table_comparator{preprocessed_left, preprocessed_right}; auto const equality_left = row_comparator.equal_to(has_nulls, compare_nulls); @@ -134,7 +135,7 @@ std::unique_ptr> mixed_join_semi( auto const equality_right_equality = row_comparator_right.equal_to(right_nulls, compare_nulls); auto const preprocessed_right_condtional = - cudf::detail::row::equality::preprocessed_table::create(right_conditional, stream); + cudf::detail::row::equality::preprocessed_table::create(right_conditional, stream, temp_mr); auto const row_comparator_conditional_right = cudf::detail::row::equality::two_table_comparator{ preprocessed_right_condtional, preprocessed_right_condtional}; auto const equality_right_conditional = @@ -147,7 +148,7 @@ std::unique_ptr> mixed_join_semi( {row_hash_right.device_hasher(right_nulls)}, {}, {}, - rmm::mr::polymorphic_allocator{}, + rmm::mr::polymorphic_allocator{temp_mr}, {stream.value()}}; auto iter = cuda::counting_iterator{0}; @@ -157,8 +158,7 @@ std::unique_ptr> mixed_join_semi( row_set.insert_async(iter, iter + right_num_rows, stream.value()); } else { cuda::counting_iterator stencil(0); - auto const [row_bitmask, _] = - cudf::detail::bitmask_and(right, stream, cudf::get_current_device_resource_ref()); + auto const [row_bitmask, _] = cudf::detail::bitmask_and(right, stream, temp_mr); row_is_valid pred{static_cast(row_bitmask.data())}; // insert valid rows @@ -176,7 +176,7 @@ std::unique_ptr> mixed_join_semi( hash_set_ref_type const row_set_ref = row_set.ref(cuco::contains).rebind_hash_function(hash_left); // Vector used to indicate indices from the left table which are present in output - auto left_table_keep_mask = rmm::device_uvector(left.num_rows(), stream); + auto left_table_keep_mask = rmm::device_uvector(left.num_rows(), stream, temp_mr); launch_mixed_join_semi(has_nulls, *left_conditional_view, diff --git a/cpp/src/lists/contains.cu b/cpp/src/lists/contains.cu index 58251a141111..6dc114ac0468 100644 --- a/cpp/src/lists/contains.cu +++ b/cpp/src/lists/contains.cu @@ -210,8 +210,8 @@ std::unique_ptr dispatch_index_of(lists_column_view const& lists, auto const keys_tview = cudf::table_view{{search_keys}}; auto const child_tview = cudf::table_view{{child}}; auto const has_nulls = has_nested_nulls(child_tview) || has_nested_nulls(keys_tview); - auto const comparator = - cudf::detail::row::equality::two_table_comparator(child_tview, keys_tview, stream); + auto const comparator = cudf::detail::row::equality::two_table_comparator( + child_tview, keys_tview, stream, cudf::get_current_device_resource_ref()); if (cudf::is_nested(search_keys.type())) { auto const d_comp = comparator.equal_to(nullate::DYNAMIC{has_nulls}); index_of(input_it, num_rows, output_it, child, search_keys, find_option, d_comp, stream); diff --git a/cpp/src/partitioning/partitioning.cu b/cpp/src/partitioning/partitioning.cu index d9bf0d99864f..48aa34d38289 100644 --- a/cpp/src/partitioning/partitioning.cu +++ b/cpp/src/partitioning/partitioning.cu @@ -575,8 +575,9 @@ std::pair, std::vector> hash_partition_table( rmm::device_async_resource_ref mr) { auto const num_rows = table_to_hash.num_rows(); + auto const temp_mr = cudf::get_current_device_resource_ref(); - auto const row_hasher = detail::row::hash::row_hasher(table_to_hash, stream); + auto const row_hasher = detail::row::hash::row_hasher(table_to_hash, stream, temp_mr); auto const hasher = row_hasher.device_hasher(nullate::DYNAMIC{hash_has_nulls}, seed); @@ -600,7 +601,7 @@ std::pair, std::vector> hash_partition_table( std::size_t const grid_size = util::div_rounding_up_safe(num_rows, rows_per_block); // Allocate array to hold which partition each row belongs to - auto row_partition_numbers = rmm::device_uvector(num_rows, stream); + auto row_partition_numbers = rmm::device_uvector(num_rows, stream, temp_mr); // Array to hold the size of each partition computed by each block // i.e., { {block0 partition0 size, block1 partition0 size, ...}, @@ -608,17 +609,18 @@ std::pair, std::vector> hash_partition_table( // ... // {block0 partition(num_partitions-1) size, block1 // partition(num_partitions -1) size, ...} } - auto block_partition_sizes = rmm::device_uvector(grid_size * num_partitions, stream); + auto block_partition_sizes = + rmm::device_uvector(grid_size * num_partitions, stream, temp_mr); auto scanned_block_partition_sizes = - rmm::device_uvector(grid_size * num_partitions, stream); + rmm::device_uvector(grid_size * num_partitions, stream, temp_mr); // Holds the total number of rows in each partition - auto global_partition_sizes = cudf::detail::make_zeroed_device_uvector_async( - num_partitions, stream, cudf::get_current_device_resource_ref()); + auto global_partition_sizes = + cudf::detail::make_zeroed_device_uvector_async(num_partitions, stream, temp_mr); - auto row_partition_offset = cudf::detail::make_zeroed_device_uvector_async( - num_rows, stream, cudf::get_current_device_resource_ref()); + auto row_partition_offset = + cudf::detail::make_zeroed_device_uvector_async(num_rows, stream, temp_mr); // If the number of partitions is a power of two, we can compute the partition // number of each row more efficiently with bitwise operations @@ -668,7 +670,7 @@ std::pair, std::vector> hash_partition_table( // Compute exclusive scan of all blocks' partition sizes in-place to determine // the starting point for each blocks portion of each partition in the output - thrust::exclusive_scan(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + thrust::exclusive_scan(rmm::exec_policy_nosync(stream, temp_mr), block_partition_sizes.begin(), block_partition_sizes.end(), scanned_block_partition_sizes.data()); @@ -676,7 +678,7 @@ std::pair, std::vector> hash_partition_table( // Compute exclusive scan of size of each partition to determine offset // location of each partition in final output. // TODO This can be done independently on a separate stream - thrust::exclusive_scan(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + thrust::exclusive_scan(rmm::exec_policy_nosync(stream, temp_mr), global_partition_sizes.begin(), global_partition_sizes.end(), global_partition_sizes.begin()); diff --git a/cpp/src/reductions/approx_distinct_count.cu b/cpp/src/reductions/approx_distinct_count.cu index 70b1c46ff58b..064f2eaa1b17 100644 --- a/cpp/src/reductions/approx_distinct_count.cu +++ b/cpp/src/reductions/approx_distinct_count.cu @@ -222,8 +222,9 @@ void approx_distinct_count::add(table_view const& input, rmm::cuda_strea typename approx_distinct_count::hll_ref_type ref{sketch(), cuda::std::identity{}}; auto const has_nulls = nullate::DYNAMIC{cudf::has_nested_nulls(input)}; + auto const temp_mr = cudf::get_current_device_resource_ref(); auto const preprocessed_input = - cudf::detail::row::hash::preprocessed_table::create(input, stream); + cudf::detail::row::hash::preprocessed_table::create(input, stream, temp_mr); auto const row_hasher = cudf::detail::row::hash::row_hasher(preprocessed_input); auto const hash_key = row_hasher.device_hasher(has_nulls); @@ -245,9 +246,8 @@ void approx_distinct_count::add(table_view const& input, rmm::cuda_strea if (!has_nulls) { ref.add_async(hash_iter, hash_iter + num_rows, stream); } else { - auto const row_bitmask = - cudf::detail::bitmask_and(input, stream, cudf::get_current_device_resource_ref()).first; - auto const pred = row_is_valid{static_cast(row_bitmask.data())}; + auto const row_bitmask = cudf::detail::bitmask_and(input, stream, temp_mr).first; + auto const pred = row_is_valid{static_cast(row_bitmask.data())}; ref.add_if_async(hash_iter, hash_iter + num_rows, stencil, pred, stream); } } else { @@ -256,8 +256,7 @@ void approx_distinct_count::add(table_view const& input, rmm::cuda_strea auto const pred = check_nans_predicate{*d_table, nullptr}; ref.add_if_async(hash_iter, hash_iter + num_rows, stencil, pred, stream); } else { - auto const row_bitmask = - cudf::detail::bitmask_and(input, stream, cudf::get_current_device_resource_ref()).first; + auto const row_bitmask = cudf::detail::bitmask_and(input, stream, temp_mr).first; auto const bitmask_ptr = static_cast(row_bitmask.data()); auto const pred = check_nans_predicate{*d_table, bitmask_ptr}; ref.add_if_async(hash_iter, hash_iter + num_rows, stencil, pred, stream); diff --git a/cpp/src/reductions/distinct_count.cu b/cpp/src/reductions/distinct_count.cu index bd2571e2dd0b..0a91ea95c356 100644 --- a/cpp/src/reductions/distinct_count.cu +++ b/cpp/src/reductions/distinct_count.cu @@ -135,11 +135,13 @@ cudf::size_type distinct_count(table_view const& keys, auto const num_rows = keys.num_rows(); if (num_rows == 0) { return 0; } // early exit for empty input auto const has_nulls = nullate::DYNAMIC{cudf::has_nested_nulls(keys)}; + auto const temp_mr = cudf::get_current_device_resource_ref(); - auto const preprocessed_input = cudf::detail::row::hash::preprocessed_table::create(keys, stream); - auto const row_hasher = cudf::detail::row::hash::row_hasher(preprocessed_input); - auto const hash_key = row_hasher.device_hasher(has_nulls); - auto const row_comp = cudf::detail::row::equality::self_comparator(preprocessed_input); + auto const preprocessed_input = + cudf::detail::row::hash::preprocessed_table::create(keys, stream, temp_mr); + auto const row_hasher = cudf::detail::row::hash::row_hasher(preprocessed_input); + auto const hash_key = row_hasher.device_hasher(has_nulls); + auto const row_comp = cudf::detail::row::equality::self_comparator(preprocessed_input); auto const comparator_helper = [&](auto const row_equal) { using hasher_type = decltype(hash_key); @@ -160,8 +162,7 @@ cudf::size_type distinct_count(table_view const& keys, cuda::counting_iterator stencil(0); // We must consider a row if any of its column entries is valid, // hence OR together the validities of the columns. - auto const [row_bitmask, null_count] = - cudf::detail::bitmask_or(keys, stream, cudf::get_current_device_resource_ref()); + auto const [row_bitmask, null_count] = cudf::detail::bitmask_or(keys, stream, temp_mr); // Unless all columns have a null mask, row_bitmask will be // null, and null_count will be zero. Equally, unless there is diff --git a/cpp/src/reductions/histogram.cu b/cpp/src/reductions/histogram.cu index e0408d64ee4f..661b1aaf2e24 100644 --- a/cpp/src/reductions/histogram.cu +++ b/cpp/src/reductions/histogram.cu @@ -111,8 +111,9 @@ compute_row_frequencies(table_view const& input, "Nested types are not yet supported in histogram aggregation.", std::invalid_argument); + auto const temp_mr = cudf::get_current_device_resource_ref(); auto const preprocessed_input = - cudf::detail::row::hash::preprocessed_table::create(input, stream); + cudf::detail::row::hash::preprocessed_table::create(input, stream, temp_mr); auto const has_nulls = nullate::DYNAMIC{cudf::has_nested_nulls(input)}; auto const row_hasher = cudf::detail::row::hash::row_hasher(preprocessed_input); @@ -131,12 +132,11 @@ compute_row_frequencies(table_view const& input, std::size_t const num_rows = input.num_rows(); // Construct a vector to store reduced counts and init to zero - rmm::device_uvector reduction_results(num_rows, stream, mr); - thrust::uninitialized_fill( - rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), - reduction_results.begin(), - reduction_results.end(), - histogram_count_type{0}); + rmm::device_uvector reduction_results(num_rows, stream, temp_mr); + thrust::uninitialized_fill(rmm::exec_policy_nosync(stream, temp_mr), + reduction_results.begin(), + reduction_results.end(), + histogram_count_type{0}); // Construct a hash set auto row_set = @@ -156,7 +156,7 @@ compute_row_frequencies(table_view const& input, // Compute frequencies (aka distinct counts) for the input rows. // Note that we consider null and NaNs as always equal. thrust::for_each( - rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + rmm::exec_policy_nosync(stream, temp_mr), cuda::counting_iterator{0}, cuda::counting_iterator{num_rows}, [set_ref = row_set_ref, diff --git a/cpp/src/reductions/scan/rank_scan.cu b/cpp/src/reductions/scan/rank_scan.cu index 4d99f4199224..3ed92dfd88a1 100644 --- a/cpp/src/reductions/scan/rank_scan.cu +++ b/cpp/src/reductions/scan/rank_scan.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -60,14 +60,15 @@ std::unique_ptr rank_generator(column_view const& order_by, rmm::device_async_resource_ref mr) { auto const order_by_tview = table_view{{order_by}}; - auto comp = cudf::detail::row::equality::self_comparator(order_by_tview, stream); + auto const temp_mr = cudf::get_current_device_resource_ref(); + auto comp = cudf::detail::row::equality::self_comparator(order_by_tview, stream, temp_mr); auto ranks = make_fixed_width_column( data_type{type_to_id()}, order_by.size(), mask_state::UNALLOCATED, stream, mr); auto mutable_ranks = ranks->mutable_view(); auto const comparator_helper = [&](auto const device_comparator) { - thrust::transform(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + thrust::transform(rmm::exec_policy_nosync(stream, temp_mr), cuda::counting_iterator(0), cuda::counting_iterator(order_by.size()), mutable_ranks.begin(), @@ -85,7 +86,7 @@ std::unique_ptr rank_generator(column_view const& order_by, comparator_helper(device_comparator); } - thrust::inclusive_scan(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + thrust::inclusive_scan(rmm::exec_policy_nosync(stream, temp_mr), mutable_ranks.begin(), mutable_ranks.end(), mutable_ranks.begin(), diff --git a/cpp/src/reductions/segmented/nunique.cu b/cpp/src/reductions/segmented/nunique.cu index 62730cbb78c2..74f6caa814cb 100644 --- a/cpp/src/reductions/segmented/nunique.cu +++ b/cpp/src/reductions/segmented/nunique.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -52,19 +52,21 @@ std::unique_ptr segmented_nunique(column_view const& col, // compute the unique identifiers within each segment auto const identifiers = [&] { - auto const d_col = column_device_view::create(col, stream); - auto const comparator = cudf::detail::row::equality::self_comparator{table_view({col}), stream}; + auto const temp_mr = cudf::get_current_device_resource_ref(); + auto const d_col = column_device_view::create(col, stream); + auto const comparator = + cudf::detail::row::equality::self_comparator{table_view({col}), stream, temp_mr}; auto const row_equal = comparator.equal_to(cudf::nullate::DYNAMIC{col.has_nulls()}, null_equality::EQUAL); - auto labels = rmm::device_uvector(col.size(), stream); + auto labels = rmm::device_uvector(col.size(), stream, temp_mr); cudf::detail::label_segments( offsets.begin(), offsets.end(), labels.begin(), labels.end(), stream); auto fn = is_unique_fn{ *d_col, row_equal, null_handling, offsets.data(), labels.data()}; - auto identifiers = rmm::device_uvector(col.size(), stream); - thrust::transform(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + auto identifiers = rmm::device_uvector(col.size(), stream, temp_mr); + thrust::transform(rmm::exec_policy_nosync(stream, temp_mr), cuda::counting_iterator{0}, cuda::counting_iterator{col.size()}, identifiers.begin(), diff --git a/cpp/src/reductions/unique_count.cu b/cpp/src/reductions/unique_count.cu index 278a12391e26..443d67466489 100644 --- a/cpp/src/reductions/unique_count.cu +++ b/cpp/src/reductions/unique_count.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -25,32 +25,30 @@ cudf::size_type unique_count(table_view const& keys, null_equality nulls_equal, rmm::cuda_stream_view stream) { - auto const row_comp = cudf::detail::row::equality::self_comparator(keys, stream); + auto const temp_mr = cudf::get_current_device_resource_ref(); + auto const row_comp = cudf::detail::row::equality::self_comparator(keys, stream, temp_mr); if (cudf::detail::has_nested_columns(keys)) { auto const comp = row_comp.equal_to(nullate::DYNAMIC{has_nested_nulls(keys)}, nulls_equal); // Using a temporary buffer for intermediate transform results from the lambda containing // the comparator speeds up compile-time significantly without much degradation in // runtime performance over using the comparator directly in thrust::count_if. - auto d_results = rmm::device_uvector(keys.num_rows(), stream); - thrust::transform(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + auto d_results = rmm::device_uvector(keys.num_rows(), stream, temp_mr); + thrust::transform(rmm::exec_policy_nosync(stream, temp_mr), cuda::counting_iterator{0}, cuda::counting_iterator{keys.num_rows()}, d_results.begin(), [comp] __device__(auto i) { return (i == 0 or not comp(i, i - 1)); }); - return static_cast( - thrust::count(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), - d_results.begin(), - d_results.end(), - true)); + return static_cast(thrust::count( + rmm::exec_policy_nosync(stream, temp_mr), d_results.begin(), d_results.end(), true)); } else { auto const comp = row_comp.equal_to(nullate::DYNAMIC{has_nested_nulls(keys)}, nulls_equal); // Using thrust::copy_if with the comparator directly will compile more slowly but // improves runtime by up to 2x over the transform/count approach above. return thrust::count_if( - rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + rmm::exec_policy_nosync(stream, temp_mr), cuda::counting_iterator{0}, cuda::counting_iterator{keys.num_rows()}, [comp] __device__(cudf::size_type i) { return (i == 0 or not comp(i, i - 1)); }); diff --git a/cpp/src/reductions/unique_count_column.cu b/cpp/src/reductions/unique_count_column.cu index effdd3a5a323..48e62a0f7010 100644 --- a/cpp/src/reductions/unique_count_column.cu +++ b/cpp/src/reductions/unique_count_column.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -61,8 +61,9 @@ cudf::size_type unique_count(column_view const& input, auto device_view = *input_device_view; auto input_table_view = table_view{{input}}; - auto const comparator = cudf::detail::row::equality::self_comparator{input_table_view, stream}; - auto const comp = comparator.equal_to( + auto const comparator = cudf::detail::row::equality::self_comparator{ + input_table_view, stream, cudf::get_current_device_resource_ref()}; + auto const comp = comparator.equal_to( nullate::DYNAMIC{cudf::has_nulls(input_table_view)}, null_equality::EQUAL, cudf::detail::row::equality::nan_equal_physical_equality_comparator{}); diff --git a/cpp/src/row_operator/row_operators.cu b/cpp/src/row_operator/row_operators.cu index 698b184abef8..d4652cbbb919 100644 --- a/cpp/src/row_operator/row_operators.cu +++ b/cpp/src/row_operator/row_operators.cu @@ -842,27 +842,27 @@ two_table_comparator::two_table_comparator(table_view const& left, namespace equality { -std::shared_ptr preprocessed_table::create(table_view const& t, - rmm::cuda_stream_view stream) +std::shared_ptr preprocessed_table::create( + table_view const& t, rmm::cuda_stream_view stream, rmm::device_async_resource_ref temp_mr) { check_eq_compatibility(t); - auto [null_pushed_table, nullable_data] = - structs::detail::push_down_nulls(t, stream, cudf::get_current_device_resource_ref()); - auto struct_offset_removed_table = remove_struct_child_offsets(null_pushed_table); + auto [null_pushed_table, nullable_data] = structs::detail::push_down_nulls(t, stream, temp_mr); + auto struct_offset_removed_table = remove_struct_child_offsets(null_pushed_table); auto verticalized_t = std::get<0>(decompose_structs(struct_offset_removed_table, decompose_lists_column::YES)); - auto d_t = table_device_view_owner(table_device_view::create(verticalized_t, stream)); + auto d_t = table_device_view_owner(table_device_view::create(verticalized_t, stream, temp_mr)); return std::shared_ptr(new preprocessed_table( std::move(d_t), std::move(nullable_data.new_null_masks), std::move(nullable_data.new_columns))); } two_table_comparator::two_table_comparator(table_view const& left, table_view const& right, - rmm::cuda_stream_view stream) - : d_left_table{preprocessed_table::create(left, stream)}, - d_right_table{preprocessed_table::create(right, stream)} + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref temp_mr) + : d_left_table{preprocessed_table::create(left, stream, temp_mr)}, + d_right_table{preprocessed_table::create(right, stream, temp_mr)} { check_shape_compatibility(left, right); } diff --git a/cpp/src/search/contains_scalar.cu b/cpp/src/search/contains_scalar.cu index 3f70f0260f69..7bbc3a95a5df 100644 --- a/cpp/src/search/contains_scalar.cu +++ b/cpp/src/search/contains_scalar.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -93,30 +93,30 @@ struct contains_scalar_dispatch { // In addition, haystack and needle structure compatibility will be checked later on by // constructor of the table comparator. - auto const haystack_tv = table_view{{haystack}}; - auto const needle_as_col = - make_column_from_scalar(needle, 1, stream, cudf::get_current_device_resource_ref()); - auto const needle_tv = table_view{{needle_as_col->view()}}; - auto const has_nulls = has_nested_nulls(haystack_tv) || has_nested_nulls(needle_tv); + auto const haystack_tv = table_view{{haystack}}; + auto const temp_mr = cudf::get_current_device_resource_ref(); + auto const needle_as_col = make_column_from_scalar(needle, 1, stream, temp_mr); + auto const needle_tv = table_view{{needle_as_col->view()}}; + auto const has_nulls = has_nested_nulls(haystack_tv) || has_nested_nulls(needle_tv); auto const comparator = - cudf::detail::row::equality::two_table_comparator(haystack_tv, needle_tv, stream); + cudf::detail::row::equality::two_table_comparator(haystack_tv, needle_tv, stream, temp_mr); auto const begin = cudf::detail::row::lhs_iterator(0); auto const end = begin + haystack.size(); using cudf::detail::row::rhs_index_type; auto const check_nulls = haystack.has_nulls(); - auto const haystack_cdv_ptr = column_device_view::create(haystack, stream); + auto const haystack_cdv_ptr = column_device_view::create(haystack, stream, temp_mr); auto const d_comp = comparator.equal_to(nullate::DYNAMIC{has_nulls}); // Using a temporary buffer for intermediate transform results from the lambda containing // the comparator speeds up compile-time significantly without much degradation in // runtime performance over using the comparator in a transform iterator with thrust::count_if. - auto d_results = rmm::device_uvector(haystack.size(), stream); + auto d_results = rmm::device_uvector(haystack.size(), stream, temp_mr); thrust::transform( - rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + rmm::exec_policy_nosync(stream, temp_mr), begin, end, d_results.begin(), @@ -127,10 +127,9 @@ struct contains_scalar_dispatch { return d_comp(idx, rhs_index_type{0}); // compare haystack[idx] == needle[0]. }); - return thrust::count(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), - d_results.begin(), - d_results.end(), - true) > 0; + return thrust::count( + rmm::exec_policy_nosync(stream, temp_mr), d_results.begin(), d_results.end(), true) > + 0; } }; diff --git a/cpp/src/search/contains_table.cu b/cpp/src/search/contains_table.cu index c3ff93adf4a3..ec6cad7dac24 100644 --- a/cpp/src/search/contains_table.cu +++ b/cpp/src/search/contains_table.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -35,10 +35,11 @@ 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; + auto const temp_mr = cudf::get_current_device_resource_ref(); auto const preprocessed_needles = - cudf::detail::row::equality::preprocessed_table::create(needles, stream); + cudf::detail::row::equality::preprocessed_table::create(needles, stream, temp_mr); auto const preprocessed_haystack = - cudf::detail::row::equality::preprocessed_table::create(haystack, stream); + cudf::detail::row::equality::preprocessed_table::create(haystack, stream, temp_mr); // The output vector. auto contained = rmm::device_uvector(needles.num_rows(), stream, mr); diff --git a/cpp/src/sort/rank.cu b/cpp/src/sort/rank.cu index 27e9f3596845..833b7d58205b 100644 --- a/cpp/src/sort/rank.cu +++ b/cpp/src/sort/rank.cu @@ -61,7 +61,8 @@ rmm::device_uvector sorted_dense_rank(column_view input_col, rmm::cuda_stream_view stream) { auto const t_input = table_view{{input_col}}; - auto const comparator = cudf::detail::row::equality::self_comparator{t_input, stream}; + auto const temp_mr = cudf::get_current_device_resource_ref(); + auto const comparator = cudf::detail::row::equality::self_comparator{t_input, stream, temp_mr}; auto const sorted_index_order = cuda::make_permutation_iterator( sorted_order_view.begin(), cuda::counting_iterator{0}); @@ -70,7 +71,7 @@ rmm::device_uvector sorted_dense_rank(column_view input_col, rmm::device_uvector dense_rank_sorted(input_size, stream); auto const comparator_helper = [&](auto const device_comparator) { - thrust::transform(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + thrust::transform(rmm::exec_policy_nosync(stream, temp_mr), cuda::counting_iterator{0}, cuda::counting_iterator{input_size}, dense_rank_sorted.data(), @@ -88,7 +89,7 @@ rmm::device_uvector sorted_dense_rank(column_view input_col, comparator_helper(device_comparator); } - thrust::inclusive_scan(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + thrust::inclusive_scan(rmm::exec_policy_nosync(stream, temp_mr), dense_rank_sorted.begin(), dense_rank_sorted.end(), dense_rank_sorted.data()); diff --git a/cpp/src/stream_compaction/distinct.cu b/cpp/src/stream_compaction/distinct.cu index 9db6bbd76c8c..25a01a1cb8dc 100644 --- a/cpp/src/stream_compaction/distinct.cu +++ b/cpp/src/stream_compaction/distinct.cu @@ -86,8 +86,8 @@ rmm::device_uvector distinct_indices(table_view const& input, return rmm::device_uvector(0, stream, mr); } - auto const preprocessed_input = - cudf::detail::row::hash::preprocessed_table::create(input, stream); + auto const preprocessed_input = cudf::detail::row::hash::preprocessed_table::create( + input, stream, cudf::get_current_device_resource_ref()); auto const has_nulls = nullate::DYNAMIC{cudf::has_nested_nulls(input)}; auto const has_nested_columns = cudf::detail::has_nested_columns(input); diff --git a/cpp/src/stream_compaction/unique.cu b/cpp/src/stream_compaction/unique.cu index 12ccfcb97129..6ea7e3d63c53 100644 --- a/cpp/src/stream_compaction/unique.cu +++ b/cpp/src/stream_compaction/unique.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -50,12 +50,13 @@ std::unique_ptr
unique(table_view const& input, auto const num_rows = input.num_rows(); if (num_rows == 0 or input.num_columns() == 0 or keys.empty()) { return empty_like(input); } + auto const temp_mr = cudf::get_current_device_resource_ref(); auto unique_indices = make_numeric_column( - data_type{type_to_id()}, num_rows, mask_state::UNALLOCATED, stream, mr); + data_type{type_to_id()}, num_rows, mask_state::UNALLOCATED, stream, temp_mr); auto mutable_view = mutable_column_device_view::create(*unique_indices, stream); auto keys_view = input.select(keys); - auto comp = cudf::detail::row::equality::self_comparator(keys_view, stream); + auto comp = cudf::detail::row::equality::self_comparator(keys_view, stream, temp_mr); size_type const unique_size = [&] { if (cudf::detail::has_nested_columns(keys_view)) { @@ -64,10 +65,10 @@ std::unique_ptr
unique(table_view const& input, // runtime performance over using the comparator directly in thrust::unique_copy. auto row_equal = comp.equal_to(nullate::DYNAMIC{has_nested_nulls(keys_view)}, nulls_equal); - auto d_results = rmm::device_uvector(num_rows, stream); + auto d_results = rmm::device_uvector(num_rows, stream, temp_mr); auto itr = cuda::counting_iterator{0}; thrust::transform( - rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + rmm::exec_policy_nosync(stream, temp_mr), itr, itr + num_rows, d_results.begin(), diff --git a/cpp/src/table/table_equal.cu b/cpp/src/table/table_equal.cu index 0cc97e0da260..8033d2d408a5 100644 --- a/cpp/src/table/table_equal.cu +++ b/cpp/src/table/table_equal.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -31,11 +31,11 @@ template null_equality nulls_equal, rmm::cuda_stream_view stream) { - auto const comparator = detail::row::equality::two_table_comparator{left, right, stream}; + auto const temp_mr = cudf::get_current_device_resource_ref(); + auto const comparator = detail::row::equality::two_table_comparator{left, right, stream, temp_mr}; auto const rows_equal = comparator.equal_to( nullate::DYNAMIC{has_nested_nulls(left) or has_nested_nulls(right)}, nulls_equal); - rmm::device_uvector eq_rows{ - static_cast(left.num_rows()), stream, cudf::get_current_device_resource_ref()}; + rmm::device_uvector eq_rows{static_cast(left.num_rows()), stream, temp_mr}; CUDF_CUDA_TRY(cub::DeviceTransform::Transform( cuda::counting_iterator{0}, eq_rows.begin(), diff --git a/cpp/src/transform/one_hot_encode.cu b/cpp/src/transform/one_hot_encode.cu index 098618c7d96c..d35f1f0eaaf7 100644 --- a/cpp/src/transform/one_hot_encode.cu +++ b/cpp/src/transform/one_hot_encode.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -68,12 +68,14 @@ std::pair, table_view> one_hot_encode(column_view const& auto all_encodings = make_numeric_column(data_type{type_id::BOOL8}, total_size, mask_state::UNALLOCATED, stream, mr); - auto const t_lhs = table_view{{input}}; - auto const t_rhs = table_view{{categories}}; - auto const comparator = cudf::detail::row::equality::two_table_comparator{t_lhs, t_rhs, stream}; + auto const temp_mr = cudf::get_current_device_resource_ref(); + auto const t_lhs = table_view{{input}}; + auto const t_rhs = table_view{{categories}}; + auto const comparator = + cudf::detail::row::equality::two_table_comparator{t_lhs, t_rhs, stream, temp_mr}; auto const comparator_helper = [&](auto const d_equal) { - thrust::transform(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + thrust::transform(rmm::exec_policy_nosync(stream, temp_mr), cuda::counting_iterator{0}, cuda::counting_iterator{total_size}, all_encodings->mutable_view().begin(), diff --git a/cpp/tests/row_operator/row_operator_tests.cu b/cpp/tests/row_operator/row_operator_tests.cu index 08046c5e3128..785a33952a45 100644 --- a/cpp/tests/row_operator/row_operator_tests.cu +++ b/cpp/tests/row_operator/row_operator_tests.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -24,7 +24,7 @@ #include template -struct TypedTableViewTest : public cudf::test::BaseFixture {}; +struct TypedTableViewTest : public cudf::test::BaseFixtureWithHarness {}; using NumericTypesNotBool = cudf::test::Concat; @@ -33,66 +33,101 @@ TYPED_TEST_SUITE(TypedTableViewTest, NumericTypesNotBool); template std::unique_ptr self_comparison(cudf::table_view input, std::vector const& column_order, - PhysicalElementComparator comparator); + PhysicalElementComparator comparator, + rmm::cuda_stream_view stream, + cudf::memory_resources mr); template std::unique_ptr two_table_comparison(cudf::table_view lhs, cudf::table_view rhs, std::vector const& column_order, - PhysicalElementComparator comparator); + PhysicalElementComparator comparator, + rmm::cuda_stream_view stream, + cudf::memory_resources mr); template std::unique_ptr two_table_equality(cudf::table_view lhs, cudf::table_view rhs, std::vector const& column_order, - PhysicalElementComparator comparator); + PhysicalElementComparator comparator, + rmm::cuda_stream_view stream, + cudf::memory_resources mr); template std::unique_ptr sorted_order( std::shared_ptr preprocessed_input, cudf::size_type num_rows, bool has_nested, PhysicalElementComparator comparator, - rmm::cuda_stream_view stream); + rmm::cuda_stream_view stream, + cudf::memory_resources mr); TYPED_TEST(TypedTableViewTest, TestLexicographicalComparatorTwoTables) { using T = TypeParam; - auto const col1 = cudf::test::fixed_width_column_wrapper{{1, 2, 3, 4}}; - auto const col2 = cudf::test::fixed_width_column_wrapper{{0, 1, 4, 3}}; + // TODO: lexicographic row operators still allocate from the current device resource. + this->enable_current_device_resource_use(); + + auto const stream = this->stream(); + auto const mr = this->resources(); + + auto const col1 = cudf::test::fixed_width_column_wrapper{{1, 2, 3, 4}, stream, mr}; + auto const col2 = cudf::test::fixed_width_column_wrapper{{0, 1, 4, 3}, stream, mr}; auto const column_order = std::vector{cudf::order::DESCENDING}; auto const lhs = cudf::table_view{{col1}}; auto const rhs = cudf::table_view{{col2}}; - auto const expected = cudf::test::fixed_width_column_wrapper{{1, 1, 0, 1}}; - auto const got = two_table_comparison( - lhs, rhs, column_order, cudf::detail::row::lexicographic::physical_element_comparator{}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); + auto const expected = cudf::test::fixed_width_column_wrapper{{1, 1, 0, 1}, stream, mr}; + auto const got = + two_table_comparison(lhs, + rhs, + column_order, + cudf::detail::row::lexicographic::physical_element_comparator{}, + stream, + mr); + CUDF_TEST_EXPECT_COLUMNS_EQUAL( + expected, got->view(), cudf::test::debug_output_level::FIRST_ERROR, stream, mr); auto const sorting_got = two_table_comparison(lhs, rhs, column_order, - cudf::detail::row::lexicographic::sorting_physical_element_comparator{}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, sorting_got->view()); + cudf::detail::row::lexicographic::sorting_physical_element_comparator{}, + stream, + mr); + CUDF_TEST_EXPECT_COLUMNS_EQUAL( + expected, sorting_got->view(), cudf::test::debug_output_level::FIRST_ERROR, stream, mr); } TYPED_TEST(TypedTableViewTest, TestLexicographicalComparatorSameTable) { using T = TypeParam; - auto const col1 = cudf::test::fixed_width_column_wrapper{{1, 2, 3, 4}}; + // TODO: lexicographic row operators still allocate from the current device resource. + this->enable_current_device_resource_use(); + + auto const stream = this->stream(); + auto const mr = this->resources(); + + auto const col1 = cudf::test::fixed_width_column_wrapper{{1, 2, 3, 4}, stream, mr}; auto const column_order = std::vector{cudf::order::DESCENDING}; auto const input_table = cudf::table_view{{col1}}; - auto const expected = cudf::test::fixed_width_column_wrapper{{0, 0, 0, 0}}; - auto const got = self_comparison( - input_table, column_order, cudf::detail::row::lexicographic::physical_element_comparator{}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); + auto const expected = cudf::test::fixed_width_column_wrapper{{0, 0, 0, 0}, stream, mr}; + auto const got = self_comparison(input_table, + column_order, + cudf::detail::row::lexicographic::physical_element_comparator{}, + stream, + mr); + CUDF_TEST_EXPECT_COLUMNS_EQUAL( + expected, got->view(), cudf::test::debug_output_level::FIRST_ERROR, stream, mr); auto const sorting_got = self_comparison(input_table, column_order, - cudf::detail::row::lexicographic::sorting_physical_element_comparator{}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, sorting_got->view()); + cudf::detail::row::lexicographic::sorting_physical_element_comparator{}, + stream, + mr); + CUDF_TEST_EXPECT_COLUMNS_EQUAL( + expected, sorting_got->view(), cudf::test::debug_output_level::FIRST_ERROR, stream, mr); } TYPED_TEST(TypedTableViewTest, TestSortSameTableFromTwoTables) @@ -100,23 +135,29 @@ TYPED_TEST(TypedTableViewTest, TestSortSameTableFromTwoTables) using data_col = cudf::test::fixed_width_column_wrapper; using int32s_col = cudf::test::fixed_width_column_wrapper; - auto const col1 = data_col{5, 2, 7, 1, 3}; + // TODO: lexicographic row operators still allocate from the current device resource. + this->enable_current_device_resource_use(); + + auto const stream = this->stream(); + auto const mr = this->resources(); + + auto const col1 = data_col{{5, 2, 7, 1, 3}, stream, mr}; auto const col2 = data_col{}; // empty auto const lhs = cudf::table_view{{col1}}; auto const empty_rhs = cudf::table_view{{col2}}; - auto const stream = cudf::get_default_stream(); auto const test_sort = - [stream]( + [stream, mr]( auto const& preprocessed, auto const& input, auto const& comparator, auto const& expected) { auto const order = sorted_order( - preprocessed, input.num_rows(), cudf::has_nested_columns(input), comparator, stream); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, order->view()); + preprocessed, input.num_rows(), cudf::has_nested_columns(input), comparator, stream, mr); + CUDF_TEST_EXPECT_COLUMNS_EQUAL( + expected, order->view(), cudf::test::debug_output_level::FIRST_ERROR, stream, mr); }; auto const test_sort_two_tables = [&](auto const& preprocessed_lhs, auto const& preprocessed_empty_rhs) { - auto const expected_lhs = int32s_col{3, 1, 4, 0, 2}; + auto const expected_lhs = int32s_col{{3, 1, 4, 0, 2}, stream, mr}; test_sort(preprocessed_lhs, lhs, cudf::detail::row::lexicographic::physical_element_comparator{}, @@ -161,17 +202,23 @@ TYPED_TEST(TypedTableViewTest, TestSortSameTableFromTwoTablesWithListsOfStructs) using strings_col = cudf::test::strings_column_wrapper; using structs_col = cudf::test::structs_column_wrapper; - auto const col1 = [] { - auto const get_structs = [] { - auto child0 = data_col{0, 3, 0, 2}; - auto child1 = strings_col{"a", "c", "a", "b"}; - return structs_col{{child0, child1}}; + // TODO: lexicographic row operators still allocate from the current device resource. + this->enable_current_device_resource_use(); + + auto const stream = this->stream(); + auto const mr = this->resources(); + + auto const col1 = [&] { + auto const get_structs = [&] { + auto child0 = data_col{{0, 3, 0, 2}, stream, mr}; + auto child1 = strings_col{{"a", "c", "a", "b"}, stream, mr}; + return structs_col{{child0, child1}, {}, stream, mr}; }; return cudf::make_lists_column( - 2, int32s_col{0, 2, 4}.release(), get_structs().release(), 0, {}); + 2, int32s_col{{0, 2, 4}, stream, mr}.release(), get_structs().release(), 0, {}); }(); - auto const col2 = [] { - auto const get_structs = [] { + auto const col2 = [&] { + auto const get_structs = [&] { auto child0 = data_col{}; auto child1 = strings_col{}; return structs_col{{child0, child1}}; @@ -183,18 +230,18 @@ TYPED_TEST(TypedTableViewTest, TestSortSameTableFromTwoTablesWithListsOfStructs) auto const lhs = cudf::table_view{{*col1}}; auto const empty_rhs = cudf::table_view{{*col2}}; - auto const stream = cudf::get_default_stream(); auto const test_sort = - [stream]( + [stream, mr]( auto const& preprocessed, auto const& input, auto const& comparator, auto const& expected) { auto const order = sorted_order( - preprocessed, input.num_rows(), cudf::has_nested_columns(input), comparator, stream); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, order->view()); + preprocessed, input.num_rows(), cudf::has_nested_columns(input), comparator, stream, mr); + CUDF_TEST_EXPECT_COLUMNS_EQUAL( + expected, order->view(), cudf::test::debug_output_level::FIRST_ERROR, stream, mr); }; auto const test_sort_two_tables = [&](auto const& preprocessed_lhs, auto const& preprocessed_empty_rhs) { - auto const expected_lhs = int32s_col{1, 0}; + auto const expected_lhs = int32s_col{{1, 0}, stream, mr}; test_sort(preprocessed_lhs, lhs, cudf::detail::row::lexicographic::sorting_physical_element_comparator{}, @@ -236,7 +283,7 @@ TYPED_TEST(TypedTableViewTest, TestSortSameTableFromTwoTablesWithListsOfStructs) } template -struct NaNTableViewTest : public cudf::test::BaseFixture {}; +struct NaNTableViewTest : public cudf::test::BaseFixtureWithHarness {}; TYPED_TEST_SUITE(NaNTableViewTest, cudf::test::FloatingPointTypes); @@ -244,65 +291,104 @@ TYPED_TEST(NaNTableViewTest, TestLexicographicalComparatorTwoTableNaNCase) { using T = TypeParam; - auto const col1 = cudf::test::fixed_width_column_wrapper{{T(NAN), T(NAN), T(1), T(1)}}; - auto const col2 = cudf::test::fixed_width_column_wrapper{{T(NAN), T(1), T(NAN), T(1)}}; + // TODO: lexicographic row operators still allocate from the current device resource. + this->enable_current_device_resource_use(); + + auto const stream = this->stream(); + auto const mr = this->resources(); + + auto const col1 = + cudf::test::fixed_width_column_wrapper{{T(NAN), T(NAN), T(1), T(1)}, stream, mr}; + auto const col2 = + cudf::test::fixed_width_column_wrapper{{T(NAN), T(1), T(NAN), T(1)}, stream, mr}; auto const column_order = std::vector{cudf::order::DESCENDING}; auto const lhs = cudf::table_view{{col1}}; auto const rhs = cudf::table_view{{col2}}; - auto const expected = cudf::test::fixed_width_column_wrapper{{0, 0, 0, 0}}; - auto const got = two_table_comparison( - lhs, rhs, column_order, cudf::detail::row::lexicographic::physical_element_comparator{}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); - - auto const sorting_expected = cudf::test::fixed_width_column_wrapper{{0, 1, 0, 0}}; + auto const expected = cudf::test::fixed_width_column_wrapper{{0, 0, 0, 0}, stream, mr}; + auto const got = + two_table_comparison(lhs, + rhs, + column_order, + cudf::detail::row::lexicographic::physical_element_comparator{}, + stream, + mr); + CUDF_TEST_EXPECT_COLUMNS_EQUAL( + expected, got->view(), cudf::test::debug_output_level::FIRST_ERROR, stream, mr); + + auto const sorting_expected = + cudf::test::fixed_width_column_wrapper{{0, 1, 0, 0}, stream, mr}; auto const sorting_got = two_table_comparison(lhs, rhs, column_order, - cudf::detail::row::lexicographic::sorting_physical_element_comparator{}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(sorting_expected, sorting_got->view()); + cudf::detail::row::lexicographic::sorting_physical_element_comparator{}, + stream, + mr); + CUDF_TEST_EXPECT_COLUMNS_EQUAL( + sorting_expected, sorting_got->view(), cudf::test::debug_output_level::FIRST_ERROR, stream, mr); } TYPED_TEST(NaNTableViewTest, TestEqualityComparatorTwoTableNaNCase) { using T = TypeParam; - auto const col1 = cudf::test::fixed_width_column_wrapper{{T(NAN), T(NAN), T(1), T(1)}}; - auto const col2 = cudf::test::fixed_width_column_wrapper{{T(NAN), T(1), T(NAN), T(1)}}; + auto const stream = this->stream(); + auto const mr = this->resources(); + + auto const col1 = + cudf::test::fixed_width_column_wrapper{{T(NAN), T(NAN), T(1), T(1)}, stream, mr}; + auto const col2 = + cudf::test::fixed_width_column_wrapper{{T(NAN), T(1), T(NAN), T(1)}, stream, mr}; auto const column_order = std::vector{cudf::order::DESCENDING}; auto const lhs = cudf::table_view{{col1}}; auto const rhs = cudf::table_view{{col2}}; - auto const expected = cudf::test::fixed_width_column_wrapper{{0, 0, 0, 1}}; - auto const got = two_table_equality( - lhs, rhs, column_order, cudf::detail::row::equality::physical_equality_comparator{}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); - - auto const nan_equal_expected = cudf::test::fixed_width_column_wrapper{{1, 0, 0, 1}}; - auto const nan_equal_got = two_table_equality( - lhs, rhs, column_order, cudf::detail::row::equality::nan_equal_physical_equality_comparator{}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(nan_equal_expected, nan_equal_got->view()); + auto const expected = cudf::test::fixed_width_column_wrapper{{0, 0, 0, 1}, stream, mr}; + auto const got = two_table_equality(lhs, + rhs, + column_order, + cudf::detail::row::equality::physical_equality_comparator{}, + stream, + mr); + CUDF_TEST_EXPECT_COLUMNS_EQUAL( + expected, got->view(), cudf::test::debug_output_level::FIRST_ERROR, stream, mr); + + auto const nan_equal_expected = + cudf::test::fixed_width_column_wrapper{{1, 0, 0, 1}, stream, mr}; + auto const nan_equal_got = + two_table_equality(lhs, + rhs, + column_order, + cudf::detail::row::equality::nan_equal_physical_equality_comparator{}, + stream, + mr); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(nan_equal_expected, + nan_equal_got->view(), + cudf::test::debug_output_level::FIRST_ERROR, + stream, + mr); } -struct RowOperatorTest : public cudf::test::BaseFixture {}; +struct RowOperatorTest : public cudf::test::BaseFixtureWithHarness {}; TEST_F(RowOperatorTest, TestTwoTableComparatorColumnCountCheck) { - rmm::cuda_stream_view stream{cudf::get_default_stream()}; + auto const stream = this->stream(); + auto const mr = this->resources(); - auto left_col1 = cudf::test::fixed_width_column_wrapper{{1, 2}}; - auto left_col2 = cudf::test::fixed_width_column_wrapper{{3, 4}}; + auto left_col1 = cudf::test::fixed_width_column_wrapper{{1, 2}, stream, mr}; + auto left_col2 = cudf::test::fixed_width_column_wrapper{{3, 4}, stream, mr}; auto const left_table = cudf::table_view{{left_col1, left_col2}}; - auto right_col = cudf::test::fixed_width_column_wrapper{{1, 2}}; + auto right_col = cudf::test::fixed_width_column_wrapper{{1, 2}, stream, mr}; auto const right_table = cudf::table_view{{right_col}}; - auto left_preprocessed = - cudf::detail::row::equality::preprocessed_table::create(left_table, stream); - auto right_preprocessed = - cudf::detail::row::equality::preprocessed_table::create(right_table, stream); + auto left_preprocessed = cudf::detail::row::equality::preprocessed_table::create( + left_table, stream, mr.get_temporary_mr()); + auto right_preprocessed = cudf::detail::row::equality::preprocessed_table::create( + right_table, stream, mr.get_temporary_mr()); EXPECT_THROW( cudf::detail::row::equality::two_table_comparator(left_preprocessed, right_preprocessed), @@ -311,48 +397,54 @@ TEST_F(RowOperatorTest, TestTwoTableComparatorColumnCountCheck) TEST_F(RowOperatorTest, TestCheckShapeCompatibility) { - rmm::cuda_stream_view stream{cudf::get_default_stream()}; + auto const stream = this->stream(); + auto const mr = this->resources(); - auto left_col1_2 = cudf::test::fixed_width_column_wrapper{{1, 2}}; - auto left_col2_2 = cudf::test::fixed_width_column_wrapper{{3, 4}}; + auto left_col1_2 = cudf::test::fixed_width_column_wrapper{{1, 2}, stream, mr}; + auto left_col2_2 = cudf::test::fixed_width_column_wrapper{{3, 4}, stream, mr}; auto const left_table = cudf::table_view{{left_col1_2, left_col2_2}}; - auto right_col_2 = cudf::test::fixed_width_column_wrapper{{1, 2}}; + auto right_col_2 = cudf::test::fixed_width_column_wrapper{{1, 2}, stream, mr}; auto const right_table = cudf::table_view{{right_col_2}}; - EXPECT_THROW(cudf::detail::row::equality::two_table_comparator(left_table, right_table, stream), + EXPECT_THROW(cudf::detail::row::equality::two_table_comparator( + left_table, right_table, stream, mr.get_temporary_mr()), std::invalid_argument); - auto int_col = cudf::test::fixed_width_column_wrapper{{1, 2}}; + auto int_col = cudf::test::fixed_width_column_wrapper{{1, 2}, stream, mr}; auto const int_table = cudf::table_view{{int_col}}; - auto float_col = cudf::test::fixed_width_column_wrapper{{1.0f, 2.0f}}; + auto float_col = cudf::test::fixed_width_column_wrapper{{1.0f, 2.0f}, stream, mr}; auto const float_table = cudf::table_view{{float_col}}; - EXPECT_THROW(cudf::detail::row::equality::two_table_comparator(int_table, float_table, stream), + EXPECT_THROW(cudf::detail::row::equality::two_table_comparator( + int_table, float_table, stream, mr.get_temporary_mr()), std::invalid_argument); - auto str_col = cudf::test::strings_column_wrapper({"hello", "world"}); + auto str_col = cudf::test::strings_column_wrapper({"hello", "world"}, stream, mr); auto const string_table = cudf::table_view{{str_col}}; - auto num_col = cudf::test::fixed_width_column_wrapper({1, 2}); + auto num_col = cudf::test::fixed_width_column_wrapper({1, 2}, stream, mr); auto const numeric_table = cudf::table_view{{num_col}}; - EXPECT_THROW( - cudf::detail::row::equality::two_table_comparator(string_table, numeric_table, stream), - std::invalid_argument); + EXPECT_THROW(cudf::detail::row::equality::two_table_comparator( + string_table, numeric_table, stream, mr.get_temporary_mr()), + std::invalid_argument); } TEST_F(RowOperatorTest, TestRowHasher64BitHash) { - auto const col = cudf::test::fixed_width_column_wrapper{{0, 42, 123456789}}; + auto const stream = this->stream(); + auto const mr = this->resources(); + + auto const col = cudf::test::fixed_width_column_wrapper{{0, 42, 123456789}, stream, mr}; auto const input = cudf::table_view{{col}}; - auto const stream = cudf::get_default_stream(); - auto const preprocessed = cudf::detail::row::hash::preprocessed_table::create(input, stream); - auto const row_hasher = cudf::detail::row::hash::row_hasher{preprocessed}; + auto const preprocessed = + cudf::detail::row::hash::preprocessed_table::create(input, stream, mr.get_temporary_mr()); + auto const row_hasher = cudf::detail::row::hash::row_hasher{preprocessed}; auto const hasher = row_hasher.device_hasher(cudf::nullate::DYNAMIC{false}); - auto results = cudf::test::fixed_width_column_wrapper{{0, 0, 0}}; - thrust::transform(rmm::exec_policy_nosync(stream), + auto results = cudf::test::fixed_width_column_wrapper{{0, 0, 0}, stream, mr}; + thrust::transform(rmm::exec_policy_nosync(stream, mr.get_temporary_mr()), cuda::counting_iterator{0}, cuda::counting_iterator{3}, cudf::mutable_column_view{results}.begin(), @@ -361,24 +453,27 @@ TEST_F(RowOperatorTest, TestRowHasher64BitHash) // Expected values match cuCollections xxhash_64 reference implementation // https://github.com/NVIDIA/cuCollections/blob/4f03dcccb3a944594c693aa8cebc89302bbd8e20/tests/utility/hash_test.cu#L134-L137 auto const expected = cudf::test::fixed_width_column_wrapper{ - {4246796580750024372ul, 15516826743637085169ul, 9462334144942111946ul}}; - CUDF_TEST_EXPECT_COLUMNS_EQUAL(results, expected); + {4246796580750024372ul, 15516826743637085169ul, 9462334144942111946ul}, stream, mr}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL( + results, expected, cudf::test::debug_output_level::FIRST_ERROR, stream, mr); } TEST_F(RowOperatorTest, TestPrimitiveRowHasher64BitHash) { - auto const col = cudf::test::fixed_width_column_wrapper{{0, 42, 123456789}}; + auto const stream = this->stream(); + auto const mr = this->resources(); + + auto const col = cudf::test::fixed_width_column_wrapper{{0, 42, 123456789}, stream, mr}; auto const input = cudf::table_view{{col}}; - auto const stream = cudf::get_default_stream(); - auto const d_input = cudf::table_device_view::create(input, stream); + auto const d_input = cudf::table_device_view::create(input, stream, mr.get_temporary_mr()); auto const hasher = cudf::detail::row::primitive::row_hasher( cudf::nullate::DYNAMIC{false}, *d_input, static_cast(cudf::DEFAULT_HASH_SEED)); - auto results = cudf::test::fixed_width_column_wrapper{{0, 0, 0}}; + auto results = cudf::test::fixed_width_column_wrapper{{0, 0, 0}, stream, mr}; - thrust::transform(rmm::exec_policy_nosync(stream), + thrust::transform(rmm::exec_policy_nosync(stream, mr.get_temporary_mr()), cuda::counting_iterator{0}, cuda::counting_iterator{3}, cudf::mutable_column_view{results}.begin(), @@ -387,77 +482,89 @@ TEST_F(RowOperatorTest, TestPrimitiveRowHasher64BitHash) // Expected values match cuCollections xxhash_64 reference implementation // https://github.com/NVIDIA/cuCollections/blob/4f03dcccb3a944594c693aa8cebc89302bbd8e20/tests/utility/hash_test.cu#L134-L137 auto const expected = cudf::test::fixed_width_column_wrapper{ - {4246796580750024372ul, 15516826743637085169ul, 9462334144942111946ul}}; - CUDF_TEST_EXPECT_COLUMNS_EQUAL(results, expected); + {4246796580750024372ul, 15516826743637085169ul, 9462334144942111946ul}, stream, mr}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL( + results, expected, cudf::test::debug_output_level::FIRST_ERROR, stream, mr); } TEST_F(RowOperatorTest, TestRowHasherDictionaryColumn) { + auto const stream = this->stream(); + auto const mr = this->resources(); + // Dictionary and equivalent string column should produce identical hashes. // This also verifies same logical values get same hashes (e.g., "baz" at rows 0 and 2). - auto const dict_col = - cudf::test::dictionary_column_wrapper({"baz", "foo", "baz", "bar", "foo"}); - auto const str_col = cudf::test::strings_column_wrapper({"baz", "foo", "baz", "bar", "foo"}); + auto const dict_col = cudf::test::dictionary_column_wrapper( + {"baz", "foo", "baz", "bar", "foo"}, stream, mr); + auto const str_col = + cudf::test::strings_column_wrapper({"baz", "foo", "baz", "bar", "foo"}, stream, mr); - auto const stream = cudf::get_default_stream(); - auto const dict_row_hasher = - cudf::detail::row::hash::row_hasher(cudf::table_view{{dict_col}}, stream); + auto const dict_row_hasher = cudf::detail::row::hash::row_hasher( + cudf::table_view{{dict_col}}, stream, mr.get_temporary_mr()); auto const str_row_hasher = - cudf::detail::row::hash::row_hasher(cudf::table_view{{str_col}}, stream); + cudf::detail::row::hash::row_hasher(cudf::table_view{{str_col}}, stream, mr.get_temporary_mr()); auto const dict_hasher = dict_row_hasher.device_hasher(cudf::nullate::DYNAMIC{false}); auto const str_hasher = str_row_hasher.device_hasher(cudf::nullate::DYNAMIC{false}); - auto dict_results = cudf::test::fixed_width_column_wrapper{{0, 0, 0, 0, 0}}; - auto str_results = cudf::test::fixed_width_column_wrapper{{0, 0, 0, 0, 0}}; + auto dict_results = + cudf::test::fixed_width_column_wrapper{{0, 0, 0, 0, 0}, stream, mr}; + auto str_results = + cudf::test::fixed_width_column_wrapper{{0, 0, 0, 0, 0}, stream, mr}; - thrust::transform(rmm::exec_policy_nosync(stream), + thrust::transform(rmm::exec_policy_nosync(stream, mr.get_temporary_mr()), cuda::counting_iterator{0}, cuda::counting_iterator{5}, cudf::mutable_column_view{dict_results}.begin(), dict_hasher); - thrust::transform(rmm::exec_policy_nosync(stream), + thrust::transform(rmm::exec_policy_nosync(stream, mr.get_temporary_mr()), cuda::counting_iterator{0}, cuda::counting_iterator{5}, cudf::mutable_column_view{str_results}.begin(), str_hasher); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(dict_results, str_results); + CUDF_TEST_EXPECT_COLUMNS_EQUAL( + dict_results, str_results, cudf::test::debug_output_level::FIRST_ERROR, stream, mr); } TEST_F(RowOperatorTest, TestRowHasherDictionaryColumnWithNulls) { - auto const dict_col = - cudf::test::dictionary_column_wrapper({100, 200, 300, 100, 200}, {1, 0, 1, 0, 1}); - auto const int_col = - cudf::test::fixed_width_column_wrapper({100, 200, 300, 100, 200}, {1, 0, 1, 0, 1}); - - auto const stream = cudf::get_default_stream(); - auto const dict_row_hasher = - cudf::detail::row::hash::row_hasher(cudf::table_view{{dict_col}}, stream); + auto const stream = this->stream(); + auto const mr = this->resources(); + + auto const dict_col = cudf::test::dictionary_column_wrapper( + {100, 200, 300, 100, 200}, {1, 0, 1, 0, 1}, stream, mr); + auto const int_col = cudf::test::fixed_width_column_wrapper( + {100, 200, 300, 100, 200}, {1, 0, 1, 0, 1}, stream, mr); + + auto const dict_row_hasher = cudf::detail::row::hash::row_hasher( + cudf::table_view{{dict_col}}, stream, mr.get_temporary_mr()); auto const int_row_hasher = - cudf::detail::row::hash::row_hasher(cudf::table_view{{int_col}}, stream); + cudf::detail::row::hash::row_hasher(cudf::table_view{{int_col}}, stream, mr.get_temporary_mr()); auto const dict_hasher = dict_row_hasher.device_hasher(cudf::nullate::DYNAMIC{true}); auto const int_hasher = int_row_hasher.device_hasher(cudf::nullate::DYNAMIC{true}); - auto dict_results = cudf::test::fixed_width_column_wrapper{{0, 0, 0, 0, 0}}; - auto int_results = cudf::test::fixed_width_column_wrapper{{0, 0, 0, 0, 0}}; + auto dict_results = + cudf::test::fixed_width_column_wrapper{{0, 0, 0, 0, 0}, stream, mr}; + auto int_results = + cudf::test::fixed_width_column_wrapper{{0, 0, 0, 0, 0}, stream, mr}; - thrust::transform(rmm::exec_policy_nosync(stream), + thrust::transform(rmm::exec_policy_nosync(stream, mr.get_temporary_mr()), cuda::counting_iterator{0}, cuda::counting_iterator{5}, cudf::mutable_column_view{dict_results}.begin(), dict_hasher); - thrust::transform(rmm::exec_policy_nosync(stream), + thrust::transform(rmm::exec_policy_nosync(stream, mr.get_temporary_mr()), cuda::counting_iterator{0}, cuda::counting_iterator{5}, cudf::mutable_column_view{int_results}.begin(), int_hasher); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(dict_results, int_results); + CUDF_TEST_EXPECT_COLUMNS_EQUAL( + dict_results, int_results, cudf::test::debug_output_level::FIRST_ERROR, stream, mr); } diff --git a/cpp/tests/row_operator/row_operator_tests_utilities.hpp b/cpp/tests/row_operator/row_operator_tests_utilities.hpp index 3841187d14d2..5b8f35558b40 100644 --- a/cpp/tests/row_operator/row_operator_tests_utilities.hpp +++ b/cpp/tests/row_operator/row_operator_tests_utilities.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -7,6 +7,7 @@ #include #include #include +#include #include @@ -20,21 +21,28 @@ using nan_equality_t = cudf::detail::row::equality::nan_equal_physical_eq template std::unique_ptr self_comparison(cudf::table_view input, std::vector const& column_order, - PhysicalElementComparator comparator); + PhysicalElementComparator comparator, + rmm::cuda_stream_view stream, + cudf::memory_resources mr); template std::unique_ptr two_table_comparison(cudf::table_view lhs, cudf::table_view rhs, std::vector const& column_order, - PhysicalElementComparator comparator); + PhysicalElementComparator comparator, + rmm::cuda_stream_view stream, + cudf::memory_resources mr); template std::unique_ptr two_table_equality(cudf::table_view lhs, cudf::table_view rhs, std::vector const& column_order, - PhysicalElementComparator comparator); + PhysicalElementComparator comparator, + rmm::cuda_stream_view stream, + cudf::memory_resources mr); template std::unique_ptr sorted_order( std::shared_ptr preprocessed_input, cudf::size_type num_rows, bool has_nested, PhysicalElementComparator comparator, - rmm::cuda_stream_view stream); + rmm::cuda_stream_view stream, + cudf::memory_resources mr); diff --git a/cpp/tests/row_operator/self_comparison_utilities.cu b/cpp/tests/row_operator/self_comparison_utilities.cu index 6e43af6dc7bd..898601ef62ea 100644 --- a/cpp/tests/row_operator/self_comparison_utilities.cu +++ b/cpp/tests/row_operator/self_comparison_utilities.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -20,25 +20,28 @@ template std::unique_ptr self_comparison(cudf::table_view input, std::vector const& column_order, - PhysicalElementComparator comparator) + PhysicalElementComparator comparator, + rmm::cuda_stream_view stream, + cudf::memory_resources mr) { - rmm::cuda_stream_view stream{cudf::get_default_stream()}; - auto const table_comparator = cudf::detail::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); + auto output = cudf::make_numeric_column(cudf::data_type(cudf::type_id::BOOL8), + input.num_rows(), + cudf::mask_state::UNALLOCATED, + stream, + mr.get_output_mr()); if (cudf::has_nested_columns(input)) { - thrust::transform(rmm::exec_policy_nosync(stream), + thrust::transform(rmm::exec_policy_nosync(stream, mr.get_temporary_mr()), cuda::counting_iterator{0}, cuda::counting_iterator{input.num_rows()}, cuda::counting_iterator{0}, output->mutable_view().data(), table_comparator.less(cudf::nullate::NO{}, comparator)); } else { - thrust::transform(rmm::exec_policy_nosync(stream), + thrust::transform(rmm::exec_policy_nosync(stream, mr.get_temporary_mr()), cuda::counting_iterator{0}, cuda::counting_iterator{input.num_rows()}, cuda::counting_iterator{0}, @@ -51,8 +54,12 @@ std::unique_ptr self_comparison(cudf::table_view input, template std::unique_ptr self_comparison( cudf::table_view input, std::vector const& column_order, - physical_comparator_t comparator); + physical_comparator_t comparator, + rmm::cuda_stream_view stream, + cudf::memory_resources mr); template std::unique_ptr self_comparison( cudf::table_view input, std::vector const& column_order, - sorting_comparator_t comparator); + sorting_comparator_t comparator, + rmm::cuda_stream_view stream, + cudf::memory_resources mr); diff --git a/cpp/tests/row_operator/two_table_comparison_utilities.cu b/cpp/tests/row_operator/two_table_comparison_utilities.cu index f42bec2a4f9a..55eb526eb809 100644 --- a/cpp/tests/row_operator/two_table_comparison_utilities.cu +++ b/cpp/tests/row_operator/two_table_comparison_utilities.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -21,27 +21,31 @@ template std::unique_ptr two_table_comparison(cudf::table_view lhs, cudf::table_view rhs, std::vector const& column_order, - PhysicalElementComparator comparator) + PhysicalElementComparator comparator, + rmm::cuda_stream_view stream, + cudf::memory_resources mr) { - rmm::cuda_stream_view stream{cudf::get_default_stream()}; - + // TODO: lexicographic::two_table_comparator still allocates from the current device resource. auto const table_comparator = cudf::detail::row::lexicographic::two_table_comparator{lhs, rhs, column_order, {}, stream}; auto const lhs_it = cudf::detail::row::lhs_iterator(0); auto const rhs_it = cudf::detail::row::rhs_iterator(0); - auto output = cudf::make_numeric_column( - cudf::data_type(cudf::type_id::BOOL8), lhs.num_rows(), cudf::mask_state::UNALLOCATED); + auto output = cudf::make_numeric_column(cudf::data_type(cudf::type_id::BOOL8), + lhs.num_rows(), + cudf::mask_state::UNALLOCATED, + stream, + mr.get_output_mr()); if (cudf::has_nested_columns(lhs) || cudf::has_nested_columns(rhs)) { - thrust::transform(rmm::exec_policy_nosync(stream), + thrust::transform(rmm::exec_policy_nosync(stream, mr.get_temporary_mr()), lhs_it, lhs_it + lhs.num_rows(), rhs_it, output->mutable_view().data(), table_comparator.less(cudf::nullate::NO{}, comparator)); } else { - thrust::transform(rmm::exec_policy_nosync(stream), + thrust::transform(rmm::exec_policy_nosync(stream, mr.get_temporary_mr()), lhs_it, lhs_it + lhs.num_rows(), rhs_it, @@ -55,12 +59,16 @@ template std::unique_ptr two_table_comparison const& column_order, - physical_comparator_t comparator); + physical_comparator_t comparator, + rmm::cuda_stream_view stream, + cudf::memory_resources mr); template std::unique_ptr two_table_comparison( cudf::table_view lhs, cudf::table_view rhs, std::vector const& column_order, - sorting_comparator_t comparator); + sorting_comparator_t comparator, + rmm::cuda_stream_view stream, + cudf::memory_resources mr); template std::unique_ptr sorted_order( @@ -68,23 +76,32 @@ std::unique_ptr sorted_order( cudf::size_type num_rows, bool has_nested, PhysicalElementComparator comparator, - rmm::cuda_stream_view stream) + rmm::cuda_stream_view stream, + cudf::memory_resources mr) { auto output = cudf::make_numeric_column(cudf::data_type(cudf::type_to_id()), num_rows, cudf::mask_state::UNALLOCATED, - stream); + stream, + mr.get_output_mr()); auto const out_begin = output->mutable_view().begin(); - thrust::sequence(rmm::exec_policy_nosync(stream), out_begin, out_begin + num_rows, 0); + thrust::sequence( + rmm::exec_policy_nosync(stream, mr.get_temporary_mr()), out_begin, out_begin + num_rows, 0); auto const table_comparator = cudf::detail::row::lexicographic::self_comparator{preprocessed_input}; if (has_nested) { auto const comp = table_comparator.less(cudf::nullate::NO{}, comparator); - thrust::stable_sort(rmm::exec_policy_nosync(stream), out_begin, out_begin + num_rows, comp); + thrust::stable_sort(rmm::exec_policy_nosync(stream, mr.get_temporary_mr()), + out_begin, + out_begin + num_rows, + comp); } else { auto const comp = table_comparator.less(cudf::nullate::NO{}, comparator); - thrust::stable_sort(rmm::exec_policy_nosync(stream), out_begin, out_begin + num_rows, comp); + thrust::stable_sort(rmm::exec_policy_nosync(stream, mr.get_temporary_mr()), + out_begin, + out_begin + num_rows, + comp); } return output; @@ -95,10 +112,12 @@ template std::unique_ptr sorted_order( cudf::size_type num_rows, bool has_nested, physical_comparator_t comparator, - rmm::cuda_stream_view stream); + rmm::cuda_stream_view stream, + cudf::memory_resources mr); template std::unique_ptr sorted_order( std::shared_ptr preprocessed_input, cudf::size_type num_rows, bool has_nested, sorting_comparator_t comparator, - rmm::cuda_stream_view stream); + rmm::cuda_stream_view stream, + cudf::memory_resources mr); diff --git a/cpp/tests/row_operator/two_table_equality_utilities.cu b/cpp/tests/row_operator/two_table_equality_utilities.cu index 4d167f911b0d..94deaf7e19e7 100644 --- a/cpp/tests/row_operator/two_table_equality_utilities.cu +++ b/cpp/tests/row_operator/two_table_equality_utilities.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -17,23 +17,27 @@ template std::unique_ptr two_table_equality(cudf::table_view lhs, cudf::table_view rhs, std::vector const& column_order, - PhysicalElementComparator comparator) + PhysicalElementComparator comparator, + rmm::cuda_stream_view stream, + cudf::memory_resources mr) { - rmm::cuda_stream_view stream{cudf::get_default_stream()}; - - auto const table_comparator = cudf::detail::row::equality::two_table_comparator{lhs, rhs, stream}; + auto const table_comparator = + cudf::detail::row::equality::two_table_comparator{lhs, rhs, stream, mr.get_temporary_mr()}; auto const lhs_it = cudf::detail::row::lhs_iterator(0); auto const rhs_it = cudf::detail::row::rhs_iterator(0); - auto output = cudf::make_numeric_column( - cudf::data_type(cudf::type_id::BOOL8), lhs.num_rows(), cudf::mask_state::UNALLOCATED); + auto output = cudf::make_numeric_column(cudf::data_type(cudf::type_id::BOOL8), + lhs.num_rows(), + cudf::mask_state::UNALLOCATED, + stream, + mr.get_output_mr()); if (cudf::has_nested_columns(lhs) or cudf::has_nested_columns(rhs)) { auto const equal_comparator = table_comparator.equal_to(cudf::nullate::NO{}, cudf::null_equality::EQUAL, comparator); - thrust::transform(rmm::exec_policy_nosync(stream), + thrust::transform(rmm::exec_policy_nosync(stream, mr.get_temporary_mr()), lhs_it, lhs_it + lhs.num_rows(), rhs_it, @@ -43,7 +47,7 @@ std::unique_ptr two_table_equality(cudf::table_view lhs, auto const equal_comparator = table_comparator.equal_to(cudf::nullate::NO{}, cudf::null_equality::EQUAL, comparator); - thrust::transform(rmm::exec_policy_nosync(stream), + thrust::transform(rmm::exec_policy_nosync(stream, mr.get_temporary_mr()), lhs_it, lhs_it + lhs.num_rows(), rhs_it, @@ -57,9 +61,13 @@ template std::unique_ptr two_table_equality( cudf::table_view lhs, cudf::table_view rhs, std::vector const& column_order, - physical_equality_t comparator); + physical_equality_t comparator, + rmm::cuda_stream_view stream, + cudf::memory_resources mr); template std::unique_ptr two_table_equality( cudf::table_view lhs, cudf::table_view rhs, std::vector const& column_order, - nan_equality_t comparator); + nan_equality_t comparator, + rmm::cuda_stream_view stream, + cudf::memory_resources mr); diff --git a/cpp/tests/utilities/column_utilities.cu b/cpp/tests/utilities/column_utilities.cu index 3d8d247b1fa5..4a0f35b4208c 100644 --- a/cpp/tests/utilities/column_utilities.cu +++ b/cpp/tests/utilities/column_utilities.cu @@ -541,10 +541,8 @@ struct column_comparator_impl { auto lhs_tview = table_view{{lhs}}; auto rhs_tview = table_view{{rhs}}; - // TODO: Pass `mr` once two_table_comparator / equality preprocessed_table::create accept - // memory_resources instead of allocating from the current device resource. - auto const comparator = - cudf::detail::row::equality::two_table_comparator{lhs_tview, rhs_tview, stream}; + auto const comparator = cudf::detail::row::equality::two_table_comparator{ + lhs_tview, rhs_tview, stream, mr.get_temporary_mr()}; auto const has_nulls = cudf::has_nulls(lhs_tview) or cudf::has_nulls(rhs_tview); auto const device_comparator = comparator.equal_to(cudf::nullate::DYNAMIC{has_nulls}); diff --git a/cpp/tests/utilities_tests/column_wrapper_tests.cpp b/cpp/tests/utilities_tests/column_wrapper_tests.cpp index fb779516d994..95035d7507e7 100644 --- a/cpp/tests/utilities_tests/column_wrapper_tests.cpp +++ b/cpp/tests/utilities_tests/column_wrapper_tests.cpp @@ -410,8 +410,6 @@ TYPED_TEST(FixedWidthColumnWrapperTest, NullablePairListConstructorAllNullMatch) this->resources()); cudf::column_view view = col; - // TODO: Remove once equality row preprocessing uses the supplied memory resources. - this->enable_current_device_resource_use(); CUDF_TEST_EXPECT_COLUMNS_EQUAL(view, match_view, cudf::test::debug_output_level::FIRST_ERROR, @@ -517,8 +515,6 @@ TYPED_TEST(StringsColumnWrapperTest, NullablePairListConstructorAllNullMatch) this->resources()); cudf::column_view view = col; - // TODO: Remove once equality row preprocessing uses the supplied memory resources. - this->enable_current_device_resource_use(); CUDF_TEST_EXPECT_COLUMNS_EQUAL(view, match_view, cudf::test::debug_output_level::FIRST_ERROR, From 1e1ecfe2f8ab4f97284d5539748ff3aabdc71518 Mon Sep 17 00:00:00 2001 From: niranda perera Date: Tue, 11 Aug 2026 19:07:14 -0700 Subject: [PATCH 2/4] revert encode changes Signed-off-by: niranda perera --- cpp/include/cudf/dictionary/detail/encode.hpp | 10 +++++----- cpp/include/cudf/dictionary/encode.hpp | 14 +++++++------- cpp/include/cudf_test/column_wrapper.hpp | 8 ++++---- cpp/src/dictionary/decode.cu | 16 +++++++--------- cpp/src/dictionary/encode.cu | 14 ++++++-------- cpp/tests/row_operator/row_operator_tests.cu | 8 ++++++++ 6 files changed, 37 insertions(+), 33 deletions(-) diff --git a/cpp/include/cudf/dictionary/detail/encode.hpp b/cpp/include/cudf/dictionary/detail/encode.hpp index 4de1b0c7e151..045df1d448a5 100644 --- a/cpp/include/cudf/dictionary/detail/encode.hpp +++ b/cpp/include/cudf/dictionary/detail/encode.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -37,13 +37,13 @@ namespace dictionary::detail { * @param column The column to dictionary encode. * @param indices_type The integer type to use for the indices. * @param stream CUDA stream used for device memory operations and kernel launches. - * @param mr Memory resources used for temporary allocations and the returned column. + * @param mr Device memory resource used to allocate the returned column's device memory. * @return Returns a dictionary column. */ std::unique_ptr encode(column_view const& column, data_type indices_type, rmm::cuda_stream_view stream, - cudf::memory_resources mr); + rmm::device_async_resource_ref mr); /** * @brief Create a column by gathering the keys from the provided @@ -57,12 +57,12 @@ std::unique_ptr encode(column_view const& column, * * @param dictionary_column Existing dictionary column. * @param stream CUDA stream used for device memory operations and kernel launches. - * @param mr Memory resources used for temporary allocations and the returned column. + * @param mr Device memory resource used to allocate the returned column's device memory. * @return New column with type matching the dictionary_column's keys. */ std::unique_ptr decode(dictionary_column_view const& dictionary_column, rmm::cuda_stream_view stream, - cudf::memory_resources mr); + rmm::device_async_resource_ref mr); /** * @brief Return minimal integer type for the given number of elements. diff --git a/cpp/include/cudf/dictionary/encode.hpp b/cpp/include/cudf/dictionary/encode.hpp index 16fc5ba56f5f..555f9d85c376 100644 --- a/cpp/include/cudf/dictionary/encode.hpp +++ b/cpp/include/cudf/dictionary/encode.hpp @@ -47,14 +47,14 @@ namespace dictionary { * @param column The column to dictionary encode * @param indices_type The integer type to use for the indices * @param stream CUDA stream used for device memory operations and kernel launches - * @param mr Memory resources used for temporary allocations and the returned column + * @param mr Device memory resource used to allocate the returned column's device memory * @return Returns a dictionary column */ std::unique_ptr encode( column_view const& column, - data_type indices_type = data_type{type_id::INT32}, - rmm::cuda_stream_view stream = cudf::get_default_stream(), - cudf::memory_resources mr = cudf::get_current_device_resource_ref()); + data_type indices_type = data_type{type_id::INT32}, + rmm::cuda_stream_view stream = cudf::get_default_stream(), + rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** * @brief Create a column by gathering the keys from the provided @@ -68,13 +68,13 @@ std::unique_ptr encode( * * @param dictionary_column Existing dictionary column * @param stream CUDA stream used for device memory operations and kernel launches - * @param mr Memory resources used for temporary allocations and the returned column + * @param mr Device memory resource used to allocate the returned column's device memory * @return New column with type matching the dictionary_column's keys */ std::unique_ptr decode( dictionary_column_view const& dictionary_column, - rmm::cuda_stream_view stream = cudf::get_default_stream(), - cudf::memory_resources mr = cudf::get_current_device_resource_ref()); + rmm::cuda_stream_view stream = cudf::get_default_stream(), + rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @} */ // end of group } // namespace dictionary diff --git a/cpp/include/cudf_test/column_wrapper.hpp b/cpp/include/cudf_test/column_wrapper.hpp index 77b8874861c0..54acc5c5e7a9 100644 --- a/cpp/include/cudf_test/column_wrapper.hpp +++ b/cpp/include/cudf_test/column_wrapper.hpp @@ -1076,7 +1076,7 @@ class dictionary_column_wrapper : public detail::column_wrapper { begin, end, stream, mr.get_temporary_mr()), cudf::data_type{type_id::INT32}, stream, - mr); + mr.get_output_mr()); } /** @@ -1118,7 +1118,7 @@ class dictionary_column_wrapper : public detail::column_wrapper { begin, end, v, stream, mr.get_temporary_mr()), cudf::data_type{type_id::INT32}, stream, - mr); + mr.get_output_mr()); } /** @@ -1308,7 +1308,7 @@ class dictionary_column_wrapper : public detail::column_wrapper { cudf::dictionary::encode(strings_column_wrapper(begin, end, stream, mr.get_temporary_mr()), cudf::data_type{type_id::INT32}, stream, - mr); + mr.get_output_mr()); } /** @@ -1353,7 +1353,7 @@ class dictionary_column_wrapper : public detail::column_wrapper { cudf::dictionary::encode(strings_column_wrapper(begin, end, v, stream, mr.get_temporary_mr()), cudf::data_type{type_id::INT32}, stream, - mr); + mr.get_output_mr()); } /** diff --git a/cpp/src/dictionary/decode.cu b/cpp/src/dictionary/decode.cu index 419b3f24ac00..1623dde5de19 100644 --- a/cpp/src/dictionary/decode.cu +++ b/cpp/src/dictionary/decode.cu @@ -36,16 +36,14 @@ struct indices_handler_fn { */ std::unique_ptr decode(dictionary_column_view const& source, rmm::cuda_stream_view stream, - cudf::memory_resources mr) + rmm::device_async_resource_ref mr) { if (source.is_empty()) return make_empty_column(type_id::EMPTY); - auto const output_mr = mr.get_output_mr(); - auto const temp_mr = mr.get_temporary_mr(); - // annotated indices include the offset, size and bitmask from it's parent - auto const indices = source.get_indices_annotated(); - auto const d_indices = column_device_view::create(indices, stream, temp_mr); + auto const indices = source.get_indices_annotated(); + auto const d_indices = + column_device_view::create(indices, stream, cudf::get_current_device_resource_ref()); auto const d_iterator = cudf::detail::indexalator_factory::make_input_iterator(indices); auto const indices_begin = cudf::detail::make_counting_transform_iterator( 0, indices_handler_fn{d_iterator, *d_indices, source.keys().size()}); @@ -55,12 +53,12 @@ std::unique_ptr decode(dictionary_column_view const& source, indices_begin + source.size(), cudf::out_of_bounds_policy::NULLIFY, stream, - output_mr) + mr) ->release(); auto output_column = std::unique_ptr(std::move(table_column.front())); // apply any nulls to the output column - output_column->set_null_mask(cudf::detail::copy_bitmask(source.parent(), stream, output_mr), + output_column->set_null_mask(cudf::detail::copy_bitmask(source.parent(), stream, mr), source.null_count()); return output_column; @@ -70,7 +68,7 @@ std::unique_ptr decode(dictionary_column_view const& source, std::unique_ptr decode(dictionary_column_view const& source, rmm::cuda_stream_view stream, - cudf::memory_resources mr) + rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); return detail::decode(source, stream, mr); diff --git a/cpp/src/dictionary/encode.cu b/cpp/src/dictionary/encode.cu index e279cac55fdb..7e1496e228bc 100644 --- a/cpp/src/dictionary/encode.cu +++ b/cpp/src/dictionary/encode.cu @@ -56,7 +56,7 @@ struct encode_fn { std::unique_ptr encode(column_view const& input, data_type indices_type, rmm::cuda_stream_view stream, - cudf::memory_resources mr) + rmm::device_async_resource_ref mr) { CUDF_EXPECTS(is_signed(indices_type) && is_index_type(indices_type), "indices must be type signed integer", @@ -68,11 +68,8 @@ std::unique_ptr encode(column_view const& input, "encoding nested types not supported", std::invalid_argument); - auto const output_mr = mr.get_output_mr(); - auto const temp_mr = mr.get_temporary_mr(); - auto indices_column = cudf::make_numeric_column( - indices_type, input.size(), cudf::mask_state::UNALLOCATED, stream, output_mr); + indices_type, input.size(), cudf::mask_state::UNALLOCATED, stream, mr); if (input.is_empty()) { return make_dictionary_column( make_empty_column(input.type()), std::move(indices_column), rmm::device_buffer{}, 0); @@ -85,6 +82,7 @@ std::unique_ptr encode(column_view const& input, auto const has_nulls = nullate::DYNAMIC{input.has_nulls()}; auto const tv = cudf::table_view({input}); + auto const temp_mr = cudf::get_current_device_resource_ref(); auto const row_hash = cudf::detail::row::hash::row_hasher(tv, stream, temp_mr); auto const row_equal = cudf::detail::row::equality::self_comparator(tv, stream, temp_mr); auto const comparator = cudf::detail::row::equality::nan_equal_physical_equality_comparator{}; @@ -118,7 +116,7 @@ std::unique_ptr encode(column_view const& input, auto const oob_policy = cudf::out_of_bounds_policy::DONT_CHECK; auto const index_policy = cudf::negative_index_policy::NOT_ALLOWED; auto keys_column = - std::move(cudf::detail::gather(tv, keys_indices, oob_policy, index_policy, stream, output_mr) + std::move(cudf::detail::gather(tv, keys_indices, oob_policy, index_policy, stream, mr) ->release() .front()); @@ -135,7 +133,7 @@ std::unique_ptr encode(column_view const& input, // create column with keys_column and indices_column return make_dictionary_column(std::move(keys_column), std::move(indices_column), - cudf::detail::copy_bitmask(input, stream, output_mr), + cudf::detail::copy_bitmask(input, stream, mr), input.null_count()); } @@ -156,7 +154,7 @@ data_type get_indices_type_for_size(size_type keys_size) std::unique_ptr encode(column_view const& input_column, data_type indices_type, rmm::cuda_stream_view stream, - cudf::memory_resources mr) + rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); return detail::encode(input_column, indices_type, stream, mr); diff --git a/cpp/tests/row_operator/row_operator_tests.cu b/cpp/tests/row_operator/row_operator_tests.cu index 785a33952a45..0bd1afd38e19 100644 --- a/cpp/tests/row_operator/row_operator_tests.cu +++ b/cpp/tests/row_operator/row_operator_tests.cu @@ -489,6 +489,10 @@ TEST_F(RowOperatorTest, TestPrimitiveRowHasher64BitHash) TEST_F(RowOperatorTest, TestRowHasherDictionaryColumn) { + // TODO: dictionary encoding gathers the keys, and gather still allocates temporaries from the + // current device resource. + this->disable_current_device_resource_use(); + auto const stream = this->stream(); auto const mr = this->resources(); @@ -531,6 +535,10 @@ TEST_F(RowOperatorTest, TestRowHasherDictionaryColumn) TEST_F(RowOperatorTest, TestRowHasherDictionaryColumnWithNulls) { + // TODO: dictionary encoding gathers the keys, and gather still allocates temporaries from the + // current device resource. + this->disable_current_device_resource_use(); + auto const stream = this->stream(); auto const mr = this->resources(); From 14b5abda82ef02b3ccf09d372192b70c8ebdb988 Mon Sep 17 00:00:00 2001 From: niranda perera Date: Tue, 11 Aug 2026 19:32:50 -0700 Subject: [PATCH 3/4] disable test case Signed-off-by: niranda perera --- cpp/tests/utilities_tests/column_wrapper_tests.cpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/cpp/tests/utilities_tests/column_wrapper_tests.cpp b/cpp/tests/utilities_tests/column_wrapper_tests.cpp index 95035d7507e7..60253ff0e09b 100644 --- a/cpp/tests/utilities_tests/column_wrapper_tests.cpp +++ b/cpp/tests/utilities_tests/column_wrapper_tests.cpp @@ -515,6 +515,9 @@ TYPED_TEST(StringsColumnWrapperTest, NullablePairListConstructorAllNullMatch) this->resources()); cudf::column_view view = col; + // TODO: has_nonempty_nulls (via count_if/transform_reduce) still allocates temporaries from the + // current device resource for strings columns. + this->disable_current_device_resource_use(); CUDF_TEST_EXPECT_COLUMNS_EQUAL(view, match_view, cudf::test::debug_output_level::FIRST_ERROR, From da6ce695ca1db494566a54a23354323b23ab94a2 Mon Sep 17 00:00:00 2001 From: niranda perera Date: Tue, 11 Aug 2026 19:34:22 -0700 Subject: [PATCH 4/4] rename util Signed-off-by: niranda perera --- cpp/tests/row_operator/row_operator_tests.cu | 4 ++-- cpp/tests/utilities_tests/column_wrapper_tests.cpp | 2 +- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/cpp/tests/row_operator/row_operator_tests.cu b/cpp/tests/row_operator/row_operator_tests.cu index 0bd1afd38e19..92dfbe9eee5c 100644 --- a/cpp/tests/row_operator/row_operator_tests.cu +++ b/cpp/tests/row_operator/row_operator_tests.cu @@ -491,7 +491,7 @@ TEST_F(RowOperatorTest, TestRowHasherDictionaryColumn) { // TODO: dictionary encoding gathers the keys, and gather still allocates temporaries from the // current device resource. - this->disable_current_device_resource_use(); + this->enable_current_device_resource_use(); auto const stream = this->stream(); auto const mr = this->resources(); @@ -537,7 +537,7 @@ TEST_F(RowOperatorTest, TestRowHasherDictionaryColumnWithNulls) { // TODO: dictionary encoding gathers the keys, and gather still allocates temporaries from the // current device resource. - this->disable_current_device_resource_use(); + this->enable_current_device_resource_use(); auto const stream = this->stream(); auto const mr = this->resources(); diff --git a/cpp/tests/utilities_tests/column_wrapper_tests.cpp b/cpp/tests/utilities_tests/column_wrapper_tests.cpp index 60253ff0e09b..d193bbaa5075 100644 --- a/cpp/tests/utilities_tests/column_wrapper_tests.cpp +++ b/cpp/tests/utilities_tests/column_wrapper_tests.cpp @@ -517,7 +517,7 @@ TYPED_TEST(StringsColumnWrapperTest, NullablePairListConstructorAllNullMatch) // TODO: has_nonempty_nulls (via count_if/transform_reduce) still allocates temporaries from the // current device resource for strings columns. - this->disable_current_device_resource_use(); + this->enable_current_device_resource_use(); CUDF_TEST_EXPECT_COLUMNS_EQUAL(view, match_view, cudf::test::debug_output_level::FIRST_ERROR,