diff --git a/cpp/include/cudf/copying.hpp b/cpp/include/cudf/copying.hpp
index f1ab6e30f268..d5ed72451864 100644
--- a/cpp/include/cudf/copying.hpp
+++ b/cpp/include/cudf/copying.hpp
@@ -75,15 +75,14 @@ enum class negative_index_policy : bool {
* better performance. If `policy` is set to `DONT_CHECK` and there are out-of-bounds indices
* in the gather map, the behavior is undefined. Defaults to `DONT_CHECK`.
* @param stream CUDA stream used for device memory operations and kernel launches
- * @param mr Device memory resource used to allocate the returned table's device memory
+ * @param mr Memory resources used for temporary allocations and the returned table
* @return Result of the gather
*/
-std::unique_ptr
gather(
- table_view const& source_table,
- column_view const& gather_map,
- out_of_bounds_policy bounds_policy = out_of_bounds_policy::DONT_CHECK,
- cuda::stream_ref stream = cudf::get_default_stream(),
- rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
+std::unique_ptr gather(table_view const& source_table,
+ column_view const& gather_map,
+ out_of_bounds_policy bounds_policy = out_of_bounds_policy::DONT_CHECK,
+ cuda::stream_ref stream = cudf::get_default_stream(),
+ cudf::memory_resources mr = cudf::get_current_device_resource_ref());
/**
* @brief Gathers the specified rows of a set of columns according to a gather map.
@@ -112,16 +111,15 @@ std::unique_ptr gather(
* @param bounds_policy Interpretation of out-of-bounds indices
* @param neg_indices Interpretation of a negative index `i` in the `gather_map`
* @param stream CUDA stream used for device memory operations and kernel launches.
- * @param mr Device memory resource used to allocate the returned table's device memory
+ * @param mr Memory resources used for temporary allocations and the returned table
* @return Result of the gather
*/
-std::unique_ptr gather(
- table_view const& source_table,
- column_view const& gather_map,
- out_of_bounds_policy bounds_policy,
- negative_index_policy neg_indices,
- cuda::stream_ref stream = cudf::get_default_stream(),
- rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
+std::unique_ptr gather(table_view const& source_table,
+ column_view const& gather_map,
+ out_of_bounds_policy bounds_policy,
+ negative_index_policy neg_indices,
+ cuda::stream_ref stream = cudf::get_default_stream(),
+ cudf::memory_resources mr = cudf::get_current_device_resource_ref());
/**
* @brief Reverses the rows within a table.
diff --git a/cpp/include/cudf/detail/gather.cuh b/cpp/include/cudf/detail/gather.cuh
index e5bb1f9ff575..faf3efcafc39 100644
--- a/cpp/include/cudf/detail/gather.cuh
+++ b/cpp/include/cudf/detail/gather.cuh
@@ -104,6 +104,7 @@ struct gather_bitmask_functor {
* @param gather_map_end End of the gather map
* @param nullify_out_of_bounds True if map values are checked against `source_size`
* @param stream CUDA stream used for kernel launches.
+ * @param temp_mr Device memory resource used for temporary allocations
*/
template
void gather_helper(InputItr source_itr,
@@ -112,11 +113,12 @@ void gather_helper(InputItr source_itr,
MapIterator gather_map_begin,
MapIterator gather_map_end,
bool nullify_out_of_bounds,
- rmm::cuda_stream_view stream)
+ rmm::cuda_stream_view stream,
+ rmm::device_async_resource_ref temp_mr = cudf::get_current_device_resource_ref())
{
using map_type = typename std::iterator_traits::value_type;
if (nullify_out_of_bounds) {
- thrust::gather_if(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()),
+ thrust::gather_if(rmm::exec_policy_nosync(stream, temp_mr),
gather_map_begin,
gather_map_end,
gather_map_begin,
@@ -124,7 +126,7 @@ void gather_helper(InputItr source_itr,
target_itr,
bounds_checker{0, source_size});
} else {
- thrust::gather(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()),
+ thrust::gather(rmm::exec_policy_nosync(stream, temp_mr),
gather_map_begin,
gather_map_end,
source_itr,
@@ -159,7 +161,7 @@ struct column_gatherer {
* @param gather_map_end End of iterator range of integral values representing the gather map
* @param nullify_out_of_bounds Nullify values in `gather_map` that are out of bounds
* @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
*/
template
std::unique_ptr operator()(column_view const& source_column,
@@ -167,7 +169,7 @@ struct column_gatherer {
MapIterator gather_map_end,
bool nullify_out_of_bounds,
rmm::cuda_stream_view stream,
- rmm::device_async_resource_ref mr)
+ cudf::memory_resources mr)
{
column_gatherer_impl gatherer{};
@@ -199,7 +201,7 @@ struct column_gatherer_impl
std::unique_ptr operator()(column_view const& source_column,
@@ -207,11 +209,12 @@ struct column_gatherer_impl(),
source_column.size(),
@@ -219,7 +222,8 @@ struct column_gatherer_impl {
* @param gather_map_end End of iterator range of integral values representing the gather map
* @param nullify_out_of_bounds Nullify values in `gather_map` that are out of bounds
* @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
*/
template
std::unique_ptr operator()(column_view const& source_column,
@@ -252,7 +256,7 @@ struct column_gatherer_impl {
MapItType gather_map_end,
bool nullify_out_of_bounds,
rmm::cuda_stream_view stream,
- rmm::device_async_resource_ref mr)
+ cudf::memory_resources mr)
{
if (true == nullify_out_of_bounds) {
return cudf::strings::detail::gather(
@@ -326,42 +330,46 @@ struct column_gatherer_impl {
MapItRoot gather_map_end,
bool nullify_out_of_bounds,
rmm::cuda_stream_view stream,
- rmm::device_async_resource_ref mr)
+ cudf::memory_resources mr)
{
+ auto const output_mr = mr.get_output_mr();
+
lists_column_view list(column);
auto gather_map_size = std::distance(gather_map_begin, gather_map_end);
// if the gather map is empty, return an empty column
if (gather_map_size == 0) { return empty_like(column); }
// generate gather_data for the next level (N+1)
- lists::detail::gather_data gd = nullify_out_of_bounds
- ? lists::detail::make_gather_data(
- column, gather_map_begin, gather_map_size, stream, mr)
- : lists::detail::make_gather_data(
- column, gather_map_begin, gather_map_size, stream, mr);
+ lists::detail::gather_data gd =
+ nullify_out_of_bounds ? lists::detail::make_gather_data(
+ column, gather_map_begin, gather_map_size, stream, output_mr)
+ : lists::detail::make_gather_data(
+ column, gather_map_begin, gather_map_size, stream, output_mr);
// the nesting case.
if (list.child().type() == cudf::data_type{type_id::LIST}) {
// gather children
- auto child = lists::detail::gather_list_nested(list.get_sliced_child(stream), gd, stream, mr);
+ auto child =
+ lists::detail::gather_list_nested(list.get_sliced_child(stream), gd, stream, output_mr);
// return the final column
return make_lists_column(gather_map_size,
std::move(gd.offsets),
std::move(child),
0,
- rmm::device_buffer{0, stream, mr});
+ rmm::device_buffer{0, stream, output_mr});
}
// it's a leaf. do a regular gather
- auto child = lists::detail::gather_list_leaf(list.get_sliced_child(stream), gd, stream, mr);
+ auto child =
+ lists::detail::gather_list_leaf(list.get_sliced_child(stream), gd, stream, output_mr);
// assemble final column
return make_lists_column(gather_map_size,
std::move(gd.offsets),
std::move(child),
0,
- rmm::device_buffer{0, stream, mr});
+ rmm::device_buffer{0, stream, output_mr});
}
};
@@ -380,7 +388,7 @@ struct column_gatherer_impl {
* @param gather_map_end End of iterator range of integral values representing the gather map
* @param nullify_out_of_bounds Nullify values in `gather_map` that are out of bounds
* @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 dictionary column with gathered rows.
*/
template
@@ -389,8 +397,11 @@ struct column_gatherer_impl {
MapItType gather_map_end,
bool nullify_out_of_bounds,
rmm::cuda_stream_view stream,
- rmm::device_async_resource_ref mr)
+ cudf::memory_resources mr)
{
+ auto const output_mr = mr.get_output_mr();
+ auto const temp_mr = mr.get_temporary_mr();
+
dictionary_column_view dictionary(source_column);
auto output_count = std::distance(gather_map_begin, gather_map_end);
if (output_count == 0) return make_empty_column(type_id::DICTIONARY32);
@@ -401,11 +412,11 @@ struct column_gatherer_impl {
// be relatively smallish.
// Also, there are scenarios where the keys are common with other dictionaries
// and the original intention was to share the keys here.
- auto keys_copy = std::make_unique(dictionary.keys(), stream, mr);
+ auto keys_copy = std::make_unique(dictionary.keys(), stream, output_mr);
// Perform gather on just the indices
column_view indices = dictionary.get_indices_annotated();
- auto new_indices =
- cudf::allocate_like(indices, output_count, cudf::mask_allocation_policy::NEVER, stream, mr);
+ auto new_indices = cudf::allocate_like(
+ indices, output_count, cudf::mask_allocation_policy::NEVER, stream, output_mr);
gather_helper(
cudf::detail::indexalator_factory::make_input_iterator(indices),
indices.size(),
@@ -413,8 +424,9 @@ struct column_gatherer_impl {
gather_map_begin,
gather_map_end,
nullify_out_of_bounds,
- stream);
- return make_dictionary_column(std::move(keys_copy), std::move(new_indices), stream, mr);
+ stream,
+ temp_mr);
+ return make_dictionary_column(std::move(keys_copy), std::move(new_indices), stream, output_mr);
}
};
@@ -426,8 +438,10 @@ struct column_gatherer_impl {
MapItRoot gather_map_end,
bool nullify_out_of_bounds,
rmm::cuda_stream_view stream,
- rmm::device_async_resource_ref mr)
+ cudf::memory_resources mr)
{
+ auto const output_mr = mr.get_output_mr();
+
auto const gather_map_size = std::distance(gather_map_begin, gather_map_end);
if (gather_map_size == 0) { return empty_like(column); }
@@ -477,9 +491,9 @@ struct column_gatherer_impl {
gather_map_size,
std::move(output_struct_members),
0,
- rmm::device_buffer{0, stream, mr}, // Null mask will be fixed up in cudf::gather().
+ rmm::device_buffer{0, stream, output_mr}, // Null mask will be fixed up in cudf::gather().
stream,
- mr);
+ output_mr);
}
};
@@ -532,10 +546,13 @@ void gather_bitmask(table_view const& source,
std::vector>& target,
gather_bitmask_op op,
rmm::cuda_stream_view stream,
- rmm::device_async_resource_ref mr)
+ cudf::memory_resources mr)
{
if (target.empty()) { return; }
+ auto const output_mr = mr.get_output_mr();
+ auto const temp_mr = mr.get_temporary_mr();
+
// Validate that all target columns have the same size
auto const target_rows = target.front()->size();
CUDF_EXPECTS(std::all_of(target.begin(),
@@ -549,7 +566,7 @@ void gather_bitmask(table_view const& source,
not target[i]->nullable()) {
auto const state =
op == gather_bitmask_op::PASSTHROUGH ? mask_state::ALL_VALID : mask_state::UNINITIALIZED;
- auto mask = cudf::create_null_mask(target[i]->size(), state, stream, mr);
+ auto mask = cudf::create_null_mask(target[i]->size(), state, stream, output_mr);
target[i]->set_null_mask(std::move(mask), 0);
}
}
@@ -559,12 +576,10 @@ void gather_bitmask(table_view const& source,
std::transform(target.begin(), target.end(), target_masks.begin(), [](auto const& col) {
return col->mutable_view().null_mask();
});
- auto d_target_masks =
- make_device_uvector_async(target_masks, stream, cudf::get_current_device_resource_ref());
+ auto d_target_masks = make_device_uvector_async(target_masks, stream, temp_mr);
- auto const device_source = table_device_view::create(source, stream);
- auto d_valid_counts = make_zeroed_device_uvector_async(
- target.size(), stream, cudf::get_current_device_resource_ref());
+ auto const device_source = table_device_view::create(source, stream, temp_mr);
+ auto d_valid_counts = make_zeroed_device_uvector_async(target.size(), stream, temp_mr);
// Dispatch operation enum to get implementation
auto const impl = [op]() {
@@ -621,7 +636,7 @@ void gather_bitmask(table_view const& source,
* better performance. In case there are out-of-bound indices in the gather map, the behavior
* is undefined. Defaults to `DONT_CHECK`.
* @param[in] stream CUDA stream used for device memory operations and kernel launches.
- * @param[in] mr Device memory resource used to allocate the returned table's device memory
+ * @param[in] mr Memory resources used for temporary allocations and the returned table
* @return cudf::table Result of the gather
*/
template
@@ -630,8 +645,10 @@ std::unique_ptr gather(table_view const& source_table,
MapIterator gather_map_end,
out_of_bounds_policy bounds_policy,
rmm::cuda_stream_view stream,
- rmm::device_async_resource_ref mr)
+ cudf::memory_resources mr)
{
+ auto const output_mr = mr.get_output_mr();
+
std::vector> destination_columns;
// TODO: Could be beneficial to use streams internally here
@@ -661,7 +678,8 @@ std::unique_ptr gather(table_view const& source_table,
gather_bitmask(source_table, gather_map_begin, destination_columns, op, stream, mr);
} else {
for (size_type i = 0; i < source_table.num_columns(); ++i) {
- set_all_valid_null_masks(source_table.column(i), *destination_columns[i], stream, mr);
+ set_all_valid_null_masks(
+ source_table.column(i), *destination_columns[i], stream, output_mr);
}
}
}
diff --git a/cpp/include/cudf/detail/gather.hpp b/cpp/include/cudf/detail/gather.hpp
index 9190b08e57f2..847f93586a5c 100644
--- a/cpp/include/cudf/detail/gather.hpp
+++ b/cpp/include/cudf/detail/gather.hpp
@@ -21,21 +21,21 @@ namespace cudf {
namespace detail {
/**
- * @copydoc cudf::gather(table_view const&,column_view const&,table_view
- * const&,cudf::out_of_bounds_policy,cudf::negative_index_policy,cuda::stream_ref,
- * rmm::device_async_resource_ref)
+ * @copydoc cudf::gather(table_view const&,column_view const&,out_of_bounds_policy,
+ * negative_index_policy,cuda::stream_ref,rmm::device_async_resource_ref)
+ *
+ * @param mr Memory resources used for temporary allocations and the returned table
*/
std::unique_ptr gather(table_view const& source_table,
column_view const& gather_map,
out_of_bounds_policy bounds_policy,
negative_index_policy neg_indices,
cuda::stream_ref stream,
- rmm::device_async_resource_ref mr);
+ memory_resources mr);
/**
- * @copydoc cudf::detail::gather(table_view const&,column_view const&,table_view
- * const&,cudf::out_of_bounds_policy,cudf::negative_index_policy,cuda::stream_ref,
- * rmm::device_async_resource_ref)
+ * @copydoc cudf::detail::gather(table_view const&,column_view const&,out_of_bounds_policy,
+ * negative_index_policy,cuda::stream_ref,memory_resources)
*
* @throws cudf::logic_error if `gather_map` span size is larger than max of `size_type`.
*/
@@ -44,7 +44,7 @@ std::unique_ptr gather(table_view const& source_table,
out_of_bounds_policy bounds_policy,
negative_index_policy neg_indices,
cuda::stream_ref stream,
- rmm::device_async_resource_ref mr);
+ memory_resources mr);
} // namespace detail
} // namespace cudf
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/detail/sizes_to_offsets_iterator.cuh b/cpp/include/cudf/detail/sizes_to_offsets_iterator.cuh
index ae79c6b683d2..05a643eca279 100644
--- a/cpp/include/cudf/detail/sizes_to_offsets_iterator.cuh
+++ b/cpp/include/cudf/detail/sizes_to_offsets_iterator.cuh
@@ -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
*/
@@ -249,27 +249,29 @@ static sizes_to_offsets_iterator make_sizes_to_offsets_i
* @param result Output iterator for scan result
* @param initial_offset Initial offset to add to scan
* @param stream CUDA stream used for device memory operations and kernel launches
+ * @param temp_mr Device memory resource used for temporary allocations
* @return The last element of the scan
*/
template
-auto sizes_to_offsets(SizesIterator begin,
- SizesIterator end,
- OffsetsIterator result,
- int64_t initial_offset,
- rmm::cuda_stream_view stream)
+auto sizes_to_offsets(
+ SizesIterator begin,
+ SizesIterator end,
+ OffsetsIterator result,
+ int64_t initial_offset,
+ rmm::cuda_stream_view stream,
+ rmm::device_async_resource_ref temp_mr = cudf::get_current_device_resource_ref())
{
using SizeType = cuda::std::iter_value_t;
static_assert(std::is_integral_v,
"Only numeric types are supported by sizes_to_offsets");
- using LastType = std::conditional_t, int64_t, uint64_t>;
- auto last_element =
- cudf::detail::device_scalar(0, stream, cudf::get_current_device_resource_ref());
+ using LastType = std::conditional_t, int64_t, uint64_t>;
+ auto last_element = cudf::detail::device_scalar(0, stream, temp_mr);
auto output_itr =
make_sizes_to_offsets_iterator(result, result + std::distance(begin, end), last_element.data());
// This function uses the type of the initialization parameter as the accumulator type
// when computing the individual scan output elements.
- thrust::exclusive_scan(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()),
+ thrust::exclusive_scan(rmm::exec_policy_nosync(stream, temp_mr),
begin,
end,
output_itr,
@@ -293,21 +295,21 @@ auto sizes_to_offsets(SizesIterator begin,
* @param begin The beginning of the input sequence
* @param end The end of the input sequence
* @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 Offsets column and total elements
*/
template
std::pair, size_type> make_offsets_child_column(
- InputIterator begin,
- InputIterator end,
- rmm::cuda_stream_view stream,
- rmm::device_async_resource_ref mr)
+ InputIterator begin, InputIterator end, rmm::cuda_stream_view stream, cudf::memory_resources mr)
{
auto count = static_cast(std::distance(begin, end));
- auto offsets_column = make_numeric_column(
- data_type{type_to_id()}, count + 1, mask_state::UNALLOCATED, stream, mr);
- auto offsets_view = offsets_column->mutable_view();
- auto d_offsets = offsets_view.template data();
+ auto offsets_column = make_numeric_column(data_type{type_to_id()},
+ count + 1,
+ mask_state::UNALLOCATED,
+ stream,
+ mr.get_output_mr());
+ auto offsets_view = offsets_column->mutable_view();
+ auto d_offsets = offsets_view.template data();
// The number of offsets is count+1 so to build the offsets from the sizes
// using exclusive-scan technically requires count+1 input values even though
@@ -320,7 +322,7 @@ std::pair, size_type> make_offsets_child_column(
auto input_itr = cudf::detail::make_counting_transform_iterator(0, map_fn);
// Use the sizes-to-offsets iterator to compute the total number of elements
auto const total_elements =
- sizes_to_offsets(input_itr, input_itr + count + 1, d_offsets, 0, stream);
+ sizes_to_offsets(input_itr, input_itr + count + 1, d_offsets, 0, stream, mr.get_temporary_mr());
CUDF_EXPECTS(
total_elements <= static_cast(std::numeric_limits::max()),
"Size of output exceeds the column size limit",
diff --git a/cpp/include/cudf/dictionary/detail/encode.hpp b/cpp/include/cudf/dictionary/detail/encode.hpp
index e728acc1e389..9eec766a80d1 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,
cuda::stream_ref 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,
cuda::stream_ref 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 d915927e7c0b..514782812e22 100644
--- a/cpp/include/cudf/dictionary/encode.hpp
+++ b/cpp/include/cudf/dictionary/encode.hpp
@@ -47,14 +47,13 @@ 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},
- cuda::stream_ref stream = cudf::get_default_stream(),
- rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
+std::unique_ptr encode(column_view const& column,
+ data_type indices_type = data_type{type_id::INT32},
+ cuda::stream_ref 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 +67,12 @@ 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,
- cuda::stream_ref stream = cudf::get_default_stream(),
- rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
+std::unique_ptr decode(dictionary_column_view const& dictionary_column,
+ cuda::stream_ref 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/lists/detail/gather.cuh b/cpp/include/cudf/lists/detail/gather.cuh
index 59ffba0c8367..56ef467cc96e 100644
--- a/cpp/include/cudf/lists/detail/gather.cuh
+++ b/cpp/include/cudf/lists/detail/gather.cuh
@@ -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
@@ -290,7 +290,8 @@ std::unique_ptr gather_list_leaf(column_view const& column,
* @copydoc cudf::lists::segmented_gather(lists_column_view const& source_column,
* lists_column_view const& gather_map_list,
* out_of_bounds_policy bounds_policy,
- * rmm::device_async_resource_ref mr)
+ * rmm::cuda_stream_view stream,
+ * cudf::memory_resources mr)
*
* @param stream CUDA stream on which to execute kernels
*/
@@ -298,7 +299,7 @@ std::unique_ptr segmented_gather(lists_column_view const& source_column,
lists_column_view const& gather_map_list,
out_of_bounds_policy bounds_policy,
rmm::cuda_stream_view stream,
- rmm::device_async_resource_ref mr);
+ cudf::memory_resources mr);
} // namespace detail
} // namespace lists
diff --git a/cpp/include/cudf/lists/gather.hpp b/cpp/include/cudf/lists/gather.hpp
index 21925eff38ee..4da1a948ee34 100644
--- a/cpp/include/cudf/lists/gather.hpp
+++ b/cpp/include/cudf/lists/gather.hpp
@@ -59,7 +59,7 @@ namespace lists {
* output list row's element, when the gather index falls outside the range `[-n, n)`,
* where `n` is the number of elements in list row corresponding to the gather-map row.
* @param stream CUDA stream used for device memory operations and kernel launches.
- * @param mr Device memory resource to allocate any returned objects
+ * @param mr Memory resources used for temporary allocations and the returned column
* @return column with elements in list of rows gathered based on `gather_map_list`
*
*/
@@ -68,7 +68,7 @@ std::unique_ptr segmented_gather(
lists_column_view const& gather_map_list,
out_of_bounds_policy bounds_policy = out_of_bounds_policy::DONT_CHECK,
rmm::cuda_stream_view stream = cudf::get_default_stream(),
- rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
+ cudf::memory_resources mr = cudf::get_current_device_resource_ref());
/** @} */ // end of group
} // namespace lists
diff --git a/cpp/include/cudf/strings/detail/gather.cuh b/cpp/include/cudf/strings/detail/gather.cuh
index 759448ac58a7..3af885f9b4f2 100644
--- a/cpp/include/cudf/strings/detail/gather.cuh
+++ b/cpp/include/cudf/strings/detail/gather.cuh
@@ -215,7 +215,7 @@ CUDF_KERNEL void gather_chars_fn_char_parallel(StringIterator strings_begin,
* @param begin Start of index iterator.
* @param end End of index iterator.
* @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 strings column containing the gathered strings.
*/
template
@@ -223,13 +223,16 @@ std::unique_ptr gather(strings_column_view const& strings,
MapIterator begin,
MapIterator end,
rmm::cuda_stream_view stream,
- rmm::device_async_resource_ref mr)
+ cudf::memory_resources mr)
{
+ auto const output_mr = mr.get_output_mr();
+ auto const temp_mr = mr.get_temporary_mr();
+
auto const output_count = std::distance(begin, end);
if (output_count == 0) return make_empty_column(type_id::STRING);
// build offsets column
- auto const d_strings = column_device_view::create(strings.parent(), stream);
+ auto const d_strings = column_device_view::create(strings.parent(), stream, temp_mr);
auto const d_in_offsets = cudf::detail::offsetalator_factory::make_input_iterator(
strings.is_empty() ? make_empty_column(type_id::INT32)->view() : strings.offsets(),
strings.offset());
@@ -252,7 +255,7 @@ std::unique_ptr gather(strings_column_view const& strings,
cudf::prefetch::detail::prefetch(strings.chars_begin(stream), strings.chars_size(stream), stream);
// build output char column
- auto out_chars_data = rmm::device_uvector(out_char_bytes, stream, mr);
+ auto out_chars_data = rmm::device_uvector(out_char_bytes, stream, output_mr);
cudf::prefetch::detail::prefetch(out_chars_data, stream);
auto d_out_chars = out_chars_data.data();
@@ -318,7 +321,7 @@ std::unique_ptr gather(strings_column_view const& strings,
stream.value());
// Allocate temporary storage
- auto d_temp_storage = rmm::device_buffer(temp_storage_bytes, stream, mr);
+ auto d_temp_storage = rmm::device_buffer(temp_storage_bytes, stream, temp_mr);
// Run batched copy algorithm
cub::DeviceMemcpy::Batched(d_temp_storage.data(),
@@ -358,7 +361,7 @@ std::unique_ptr gather(strings_column_view const& strings,
* @param end End of index iterator.
* @param nullify_out_of_bounds If true, indices outside the column's range are nullified.
* @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 strings column containing the gathered strings.
*/
template
@@ -367,7 +370,7 @@ std::unique_ptr gather(strings_column_view const& strings,
MapIterator end,
bool nullify_out_of_bounds,
rmm::cuda_stream_view stream,
- rmm::device_async_resource_ref mr)
+ cudf::memory_resources mr)
{
if (nullify_out_of_bounds) return gather(strings, begin, end, stream, mr);
return gather(strings, begin, end, stream, mr);
diff --git a/cpp/include/cudf/strings/detail/strings_children.cuh b/cpp/include/cudf/strings/detail/strings_children.cuh
index d15e5fa199f1..59f0334be297 100644
--- a/cpp/include/cudf/strings/detail/strings_children.cuh
+++ b/cpp/include/cudf/strings/detail/strings_children.cuh
@@ -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
*/
#pragma once
@@ -115,23 +115,25 @@ rmm::device_uvector make_chars_buffer(column_view const& offsets,
* @param begin The beginning of the input sequence
* @param end The end of the input sequence
* @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 Offsets column and total elements
*/
template
-std::pair, int64_t> make_offsets_child_column(
- InputIterator begin,
- InputIterator end,
- rmm::cuda_stream_view stream,
- rmm::device_async_resource_ref mr)
+std::pair, int64_t> make_offsets_child_column(InputIterator begin,
+ InputIterator end,
+ rmm::cuda_stream_view stream,
+ cudf::memory_resources mr)
{
+ auto const output_mr = mr.get_output_mr();
+ auto const temp_mr = mr.get_temporary_mr();
+
auto constexpr size_type_max = static_cast(std::numeric_limits::max());
auto const lcount = static_cast(std::distance(begin, end));
CUDF_EXPECTS(
lcount <= size_type_max, "Size of output exceeds the column size limit", std::overflow_error);
auto const strings_count = static_cast(lcount);
auto offsets_column = make_numeric_column(
- data_type{type_id::INT32}, strings_count + 1, mask_state::UNALLOCATED, stream, mr);
+ data_type{type_id::INT32}, strings_count + 1, mask_state::UNALLOCATED, stream, output_mr);
auto d_offsets = offsets_column->mutable_view().template data();
// The number of offsets is strings_count+1 so to build the offsets from the sizes
@@ -141,8 +143,8 @@ std::pair, int64_t> make_offsets_child_column(
auto input_itr =
cudf::detail::make_counting_transform_iterator(0, string_offsets_fn{begin, strings_count});
// Use the sizes-to-offsets iterator to compute the total number of elements
- auto const total_bytes =
- cudf::detail::sizes_to_offsets(input_itr, input_itr + strings_count + 1, d_offsets, 0, stream);
+ auto const total_bytes = cudf::detail::sizes_to_offsets(
+ input_itr, input_itr + strings_count + 1, d_offsets, 0, stream, temp_mr);
auto const threshold = cudf::strings::get_offset64_threshold();
CUDF_EXPECTS(cudf::strings::is_large_strings_enabled() || (total_bytes < threshold),
@@ -151,10 +153,10 @@ std::pair, int64_t> make_offsets_child_column(
if (total_bytes >= cudf::strings::get_offset64_threshold()) {
// recompute as int64 offsets when above the threshold
offsets_column = make_numeric_column(
- data_type{type_id::INT64}, strings_count + 1, mask_state::UNALLOCATED, stream, mr);
+ data_type{type_id::INT64}, strings_count + 1, mask_state::UNALLOCATED, stream, output_mr);
auto d_offsets64 = offsets_column->mutable_view().template data();
cudf::detail::sizes_to_offsets(
- input_itr, input_itr + strings_count + 1, d_offsets64, 0, stream);
+ input_itr, input_itr + strings_count + 1, d_offsets64, 0, stream, temp_mr);
}
return std::pair(std::move(offsets_column), total_bytes);
diff --git a/cpp/include/cudf_test/base_fixture.hpp b/cpp/include/cudf_test/base_fixture.hpp
index 003955369001..326bf1ad3b89 100644
--- a/cpp/include/cudf_test/base_fixture.hpp
+++ b/cpp/include/cudf_test/base_fixture.hpp
@@ -45,8 +45,10 @@ class BaseFixture : public ::testing::Test {
/**
* @brief Base fixture that instruments tests with a memory-resource harness.
*
- * Each test instantiates a fresh harness. Tests should construct results with `resources()`.
- * `TearDown` asserts that no output or temporary allocations remain live.
+ * Each test instantiates a fresh harness and installs a failing current-device resource for the
+ * duration of the test, so accidental fallback to the default MR fails. Tests should construct
+ * results with `resources()`. `TearDown` asserts that no output or temporary allocations remain
+ * live; the prior current resource is restored when `_fail_on_current` is destroyed.
*/
struct BaseFixtureWithHarness : public BaseFixture {
/**
@@ -68,6 +70,7 @@ struct BaseFixtureWithHarness : public BaseFixture {
protected:
memory_resource_test_harness _harness{mr()};
+ scoped_current_device_resource _fail_on_current{_harness.fail_on_current_device_resource_use()};
};
/**
diff --git a/cpp/include/cudf_test/column_wrapper.hpp b/cpp/include/cudf_test/column_wrapper.hpp
index 54acc5c5e7a9..34d217196e16 100644
--- a/cpp/include/cudf_test/column_wrapper.hpp
+++ b/cpp/include/cudf_test/column_wrapper.hpp
@@ -36,9 +36,13 @@
#include
#include
+#include
+#include
#include
#include
#include
+#include
+#include
namespace CUDF_EXPORT cudf {
namespace test {
@@ -340,8 +344,155 @@ auto make_chars_and_offsets(StringsIterator begin, StringsIterator end, Validity
}
return std::pair(std::move(chars), std::move(offsets));
};
+
} // namespace detail
+// Forward declaration for lists_column_initializer::build
+template
+class lists_column_wrapper;
+
+/**
+ * @brief Host-side recursive initializer tree for constructing list columns with an
+ * explicit stream and memory resources at every nesting level.
+ *
+ * Prefer this over brace-nested `lists_column_wrapper` constructions that pass
+ * `stream`/`mr` only at the outer level, which leave brace-constructed children on
+ * the default test resources.
+ *
+ * Example:
+ * @code{.cpp}
+ * using Init = cudf::test::lists_column_initializer;
+ * // List: [{1, 2}, {3}]
+ * lists_column_wrapper col{Init{{{1, 2}, {3}}}, stream, mr};
+ * @endcode
+ *
+ * Leaf and nested constructors accept the existing validity iterators
+ * (`valids`, `null_at(...)`, etc.) and materialize them into owned storage.
+ *
+ * @tparam T Host leaf element type (e.g. `int32_t` or `std::string`)
+ */
+template
+class lists_column_initializer {
+ public:
+ /**
+ * @brief Construct an empty leaf. Avoids ambiguity between the leaf and nested
+ * empty `initializer_list` constructors.
+ */
+ lists_column_initializer() = default;
+
+ /**
+ * @brief Construct a leaf from scalar values.
+ *
+ * @param values Leaf element values
+ */
+ lists_column_initializer(std::initializer_list values) : values_{values} {}
+
+ /**
+ * @brief Construct a leaf from scalar values and a validity iterator.
+ *
+ * @tparam ValidityIterator Iterator convertible to `bool`
+ * @param values Leaf element values
+ * @param v Validity iterator over `values.size()` elements
+ */
+ template
+ lists_column_initializer(std::initializer_list values, ValidityIterator v) : values_{values}
+ {
+ value_validity_.reserve(values_.size());
+ for (std::size_t i = 0; i < values_.size(); ++i) {
+ value_validity_.push_back(static_cast(*v++));
+ }
+ }
+
+ /**
+ * @brief Construct a nested node from child initializers.
+ *
+ * This constructor is a template so the non-template leaf
+ * `initializer_list` constructor is preferred for scalar lists such as
+ * `{1, 2, 3}`. Otherwise both overloads are non-templates and constructing
+ * `Init` from an `int` via the nested overload recurses until the stack
+ * overflows.
+ *
+ * @param children Child list initializers
+ */
+ template
+ requires(std::is_same_v)
+ lists_column_initializer(std::initializer_list children)
+ : children_{children.begin(), children.end()}, nested_{true}
+ {
+ }
+
+ /**
+ * @brief Construct a nested node from child initializers and a row-validity iterator.
+ *
+ * @tparam ValidityIterator Iterator convertible to `bool`
+ * @param children Child list initializers
+ * @param v Validity iterator over `children.size()` rows
+ */
+ template
+ requires(std::is_same_v)
+ lists_column_initializer(std::initializer_list children, ValidityIterator v)
+ : nested_{true}
+ {
+ children_.reserve(children.size());
+ for (auto const& child : children) {
+ if (static_cast(*v++)) {
+ children_.push_back(child);
+ } else {
+ children_.emplace_back();
+ children_.back().valid_ = false;
+ }
+ }
+ }
+
+ [[nodiscard]] bool nested() const { return nested_; }
+ [[nodiscard]] bool valid() const { return valid_; }
+ [[nodiscard]] auto const& values() const { return values_; }
+ [[nodiscard]] auto const& value_validity() const { return value_validity_; }
+ [[nodiscard]] auto const& children() const { return children_; }
+
+ /**
+ * @brief Recursively build child list wrappers and row validity for a nested node.
+ *
+ * Each valid child is allocated with the provided `stream` and `mr`. Null children are
+ * represented as default-constructed wrappers (skipped during concatenate).
+ *
+ * @tparam ElementT List wrapper element type
+ * @tparam SourceElementT Source type used by the list wrapper
+ * @param stream CUDA stream used for device memory operations
+ * @param mr Memory resources used to allocate child columns
+ * @return Child wrappers and an empty validity vector when all rows are valid,
+ * otherwise a validity mask matching `children().size()`
+ */
+ template
+ [[nodiscard]] std::pair>,
+ std::vector>
+ build(rmm::cuda_stream_view stream, cudf::memory_resources mr) const
+ {
+ std::vector> children;
+ std::vector validity;
+ children.reserve(children_.size());
+ validity.reserve(children_.size());
+ bool any_null = false;
+ for (auto const& child : children_) {
+ any_null = any_null || !child.valid();
+ validity.push_back(child.valid());
+ if (child.valid()) {
+ children.emplace_back(child, stream, mr);
+ } else {
+ children.emplace_back(); // null rows are skipped during concatenate
+ }
+ }
+ return {std::move(children), any_null ? std::move(validity) : std::vector{}};
+ }
+
+ private:
+ std::vector values_;
+ std::vector value_validity_;
+ std::vector children_;
+ bool nested_{false};
+ bool valid_{true};
+};
+
/**
* @brief `column_wrapper` derived class for wrapping columns of fixed-width
* elements.
@@ -1076,7 +1227,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 +1269,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 +1459,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 +1504,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);
}
/**
@@ -1476,192 +1627,178 @@ class lists_column_wrapper : public detail::column_wrapper {
*/
operator lists_column_view() const { return cudf::lists_column_view{wrapped->view()}; }
+ using host_element_t =
+ std::conditional_t, std::string, SourceElementT>;
+ using leaf_wrapper_t = std::conditional_t,
+ strings_column_wrapper,
+ fixed_width_column_wrapper>;
+
/**
- * @brief Construct a lists column containing a single list of fixed-width
- * type from an initializer list of values.
+ * @brief Construct a lists column containing a single list from an initializer
+ * list of values.
*
* Example:
* @code{.cpp}
- * Creates a LIST column with 1 list composed of 2 total integers
- * [{0, 1}]
+ * // Creates a LIST column with 1 list composed of 2 total integers
+ * // [{0, 1}]
* lists_column_wrapper l{0, 1};
* @endcode
*
+ * These leaf constructors are templates (via `requires`) so that the non-template
+ * nested `initializer_list` constructor is preferred for
+ * ambiguous cases such as `lists_column_wrapper{{}, {}}`.
+ *
* @param elements The list of elements
* @param stream CUDA stream used for device memory operations
* @param mr Memory resources used to allocate the returned column
*/
- template ()>* = nullptr>
+ template
+ requires(cudf::is_fixed_width())
lists_column_wrapper(std::initializer_list elements,
rmm::cuda_stream_view stream = cudf::test::get_default_stream(),
cudf::memory_resources mr = cudf::get_current_device_resource_ref())
: column_wrapper{}
{
build_from_non_nested(
- cudf::test::fixed_width_column_wrapper(elements, stream, mr).release(),
- stream,
- mr);
+ fixed_width_column_wrapper(elements, stream, mr).release(), stream, mr);
}
/**
- * @brief Construct a lists column containing a single list of fixed-width
- * type from an iterator range.
+ * @brief Construct a lists column containing a single list of strings.
*
* Example:
* @code{.cpp}
- * // Creates a LIST column with 1 list composed of 5 total integers
- * auto elements = make_counting_transform_iterator(0, [](auto i){return i*2;});
- * // [{0, 1, 2, 3, 4}]
- * lists_column_wrapper l(elements, elements+5);
+ * // Creates a LIST column with 1 list composed of 2 total strings
+ * // [{"abc", "def"}]
+ * lists_column_wrapper s{"abc", "def"};
* @endcode
*
- * @param begin Beginning of the sequence
- * @param end End of the sequence
+ * @param elements The list of strings
* @param stream CUDA stream used for device memory operations
* @param mr Memory resources used to allocate the returned column
*/
- template ()>* = nullptr>
- lists_column_wrapper(InputIterator begin,
- InputIterator end,
+ template
+ requires(std::is_same_v)
+ lists_column_wrapper(std::initializer_list elements,
rmm::cuda_stream_view stream = cudf::test::get_default_stream(),
cudf::memory_resources mr = cudf::get_current_device_resource_ref())
: column_wrapper{}
{
- build_from_non_nested(
- cudf::test::fixed_width_column_wrapper(begin, end, stream, mr).release(),
- stream,
- mr);
+ build_from_non_nested(strings_column_wrapper(elements, stream, mr).release(), stream, mr);
}
/**
- * @brief Construct a lists column containing a single list of fixed-width
- * type from an initializer list of values and a validity iterator.
+ * @brief Construct a lists column containing a single list from an iterator range.
*
* Example:
* @code{.cpp}
- * // Creates a LIST column with 1 lists composed of 2 total integers
- * auto validity = make_counting_transform_iterator(0, [](auto i){return i%2;});
- * // [{0, NULL}]
- * lists_column_wrapper l{{0, 1}, validity};
+ * // Creates a LIST column with 1 list composed of 5 total integers
+ * auto elements = make_counting_transform_iterator(0, [](auto i){return i*2;});
+ * // [{0, 1, 2, 3, 4}]
+ * lists_column_wrapper l(elements, elements+5);
* @endcode
*
- * @param elements The list of elements
- * @param v The validity iterator
+ * @param begin Beginning of the sequence
+ * @param end End of the sequence
* @param stream CUDA stream used for device memory operations
* @param mr Memory resources used to allocate the returned column
*/
- template ()>* = nullptr>
- lists_column_wrapper(std::initializer_list elements,
- ValidityIterator v,
+ template
+ lists_column_wrapper(InputIterator begin,
+ InputIterator end,
rmm::cuda_stream_view stream = cudf::test::get_default_stream(),
cudf::memory_resources mr = cudf::get_current_device_resource_ref())
: column_wrapper{}
{
- build_from_non_nested(
- cudf::test::fixed_width_column_wrapper(elements, v, stream, mr).release(),
- stream,
- mr);
+ build_from_non_nested(leaf_wrapper_t(begin, end, stream, mr).release(), stream, mr);
}
/**
- * @brief Construct a lists column containing a single list of fixed-width
- * type from an iterator range and a validity iterator.
+ * @brief Construct a lists column containing a single list from an initializer
+ * list of values and a validity iterator.
*
* Example:
* @code{.cpp}
- * // Creates a LIST column with 1 lists composed of 5 total integers
- * auto elements = make_counting_transform_iterator(0, [](auto i){return i*2;});
+ * // Creates a LIST column with 1 list composed of 2 total integers
* auto validity = make_counting_transform_iterator(0, [](auto i){return i%2;});
- * // [{0, NULL, 2, NULL, 4}]
- * lists_column_wrapper l(elements, elements+5, validity);
+ * // [{0, NULL}]
+ * lists_column_wrapper l{{0, 1}, validity};
* @endcode
*
- * @param begin Beginning of the sequence
- * @param end End of the sequence
+ * @param elements The list of elements
* @param v The validity iterator
* @param stream CUDA stream used for device memory operations
* @param mr Memory resources used to allocate the returned column
*/
- template ()>* = nullptr>
- lists_column_wrapper(InputIterator begin,
- InputIterator end,
+ template
+ requires(cudf::is_fixed_width())
+ lists_column_wrapper(std::initializer_list elements,
ValidityIterator v,
rmm::cuda_stream_view stream = cudf::test::get_default_stream(),
cudf::memory_resources mr = cudf::get_current_device_resource_ref())
: column_wrapper{}
{
build_from_non_nested(
- cudf::test::fixed_width_column_wrapper(begin, end, v, stream, mr)
- .release(),
+ fixed_width_column_wrapper(elements, v, stream, mr).release(),
stream,
mr);
}
/**
- * @brief Construct a lists column containing a single list of strings
- * from an initializer list of values.
+ * @brief Construct a lists column containing a single list of strings and a
+ * validity iterator.
*
* Example:
* @code{.cpp}
- * // Creates a LIST column with 1 list composed of 2 total strings
- * // [{"abc", "def"}]
- * lists_column_wrapper l{"abc", "def"};
+ * auto validity = make_counting_transform_iterator(0, [](auto i){return i%2;});
+ * // [{"abc", NULL}]
+ * lists_column_wrapper l{{"abc", "def"}, validity};
* @endcode
*
- * @param elements The list of elements
+ * @param elements The list of strings
+ * @param v The validity iterator
* @param stream CUDA stream used for device memory operations
* @param mr Memory resources used to allocate the returned column
*/
- template >* = nullptr>
+ template
+ requires(std::is_same_v)
lists_column_wrapper(std::initializer_list elements,
+ ValidityIterator v,
rmm::cuda_stream_view stream = cudf::test::get_default_stream(),
cudf::memory_resources mr = cudf::get_current_device_resource_ref())
: column_wrapper{}
{
- build_from_non_nested(
- cudf::test::strings_column_wrapper(elements.begin(), elements.end(), stream, mr).release(),
- stream,
- mr);
+ build_from_non_nested(strings_column_wrapper(elements, v, stream, mr).release(), stream, mr);
}
/**
- * @brief Construct a lists column containing a single list of strings
- * from an initializer list of values and a validity iterator.
+ * @brief Construct a lists column containing a single list from an iterator
+ * range and a validity iterator.
*
* Example:
* @code{.cpp}
- * // Creates a LIST column with 1 list composed of 2 total strings
+ * // Creates a LIST column with 1 list composed of 5 total integers
+ * auto elements = make_counting_transform_iterator(0, [](auto i){return i*2;});
* auto validity = make_counting_transform_iterator(0, [](auto i){return i%2;});
- * // [{"abc", NULL}]
- * lists_column_wrapper l{{"abc", "def"}, validity};
+ * // [{0, NULL, 2, NULL, 4}]
+ * lists_column_wrapper l(elements, elements+5, validity);
* @endcode
*
- * @param elements The list of elements
+ * @param begin Beginning of the sequence
+ * @param end End of the sequence
* @param v The validity iterator
* @param stream CUDA stream used for device memory operations
* @param mr Memory resources used to allocate the returned column
*/
- template >* = nullptr>
- lists_column_wrapper(std::initializer_list elements,
+ template
+ lists_column_wrapper(InputIterator begin,
+ InputIterator end,
ValidityIterator v,
rmm::cuda_stream_view stream = cudf::test::get_default_stream(),
cudf::memory_resources mr = cudf::get_current_device_resource_ref())
: column_wrapper{}
{
- build_from_non_nested(
- cudf::test::strings_column_wrapper(elements.begin(), elements.end(), v, stream, mr).release(),
- stream,
- mr);
+ build_from_non_nested(leaf_wrapper_t(begin, end, v, stream, mr).release(), stream, mr);
}
/**
@@ -1683,6 +1820,11 @@ class lists_column_wrapper : public detail::column_wrapper {
* lists_column_wrapper l{ {{0, 1}, {2, 3}}, {{4, 5}, {6, 7}} };
* @endcode
*
+ * For multi-row (and deeper) columns that should allocate with an explicit stream/mr, use
+ * `lists_column_initializer` so every nesting level receives those arguments:
+ * `using Init = cudf::test::lists_column_initializer;`
+ * `lists_column_wrapper l{Init{{{0, 1}, {2, 3}, {4, 5}}}, stream, mr};`
+ *
* @param elements The list of elements
* @param stream CUDA stream used for device memory operations
* @param mr Memory resources used to allocate the returned column
@@ -1758,6 +1900,46 @@ class lists_column_wrapper : public detail::column_wrapper {
build_from_nested(elements, validity, stream, mr);
}
+ /**
+ * @brief Construct a lists column from a recursive `lists_column_initializer` tree.
+ *
+ * Every nesting level is allocated with the provided `stream` and `mr`. Prefer this over
+ * brace-nested `lists_column_wrapper` constructions that pass resources only at the outer
+ * level.
+ *
+ * Example:
+ * @code{.cpp}
+ * using Init = cudf::test::lists_column_initializer;
+ * // List: [{0, 1}, {2, 3}, {4, 5}]
+ * lists_column_wrapper l{Init{{{0, 1}, {2, 3}, {4, 5}}}, stream, mr};
+ *
+ * // List>: [{{0, 1}, {2}}, {{3}}]
+ * lists_column_wrapper nested{Init{{{{0, 1}, {2}}, {{3}}}}, stream, mr};
+ * @endcode
+ *
+ * @param init Host-side nested values (and optional validity)
+ * @param stream CUDA stream used for device memory operations
+ * @param mr Memory resources used to allocate the returned column
+ */
+ lists_column_wrapper(lists_column_initializer init,
+ rmm::cuda_stream_view stream,
+ cudf::memory_resources mr)
+ : column_wrapper{}
+ {
+ if (!init.nested()) {
+ if (init.value_validity().empty()) {
+ *this = lists_column_wrapper(init.values().begin(), init.values().end(), stream, mr);
+ } else {
+ *this = lists_column_wrapper(
+ init.values().begin(), init.values().end(), init.value_validity().begin(), stream, mr);
+ }
+ return;
+ }
+
+ auto [children, validity] = init.template build(stream, mr);
+ build_from_nested(children, validity, stream, mr);
+ }
+
/**
* @brief Construct a list column containing a single empty, optionally null row.
*
@@ -1821,7 +2003,8 @@ class lists_column_wrapper : public detail::column_wrapper {
* @param mr Memory resources used to allocate the returned column
*
*/
- void build_from_nested(std::initializer_list> elements,
+ template
+ void build_from_nested(ListsRange const& elements,
std::vector const& v,
rmm::cuda_stream_view stream,
cudf::memory_resources mr)
@@ -1981,8 +2164,9 @@ class lists_column_wrapper : public detail::column_wrapper {
cudf::copy_bitmask(col, stream, temp_mr));
}
+ template
std::pair, std::vector>> preprocess_columns(
- std::initializer_list> const& elements,
+ ListsRange const& elements,
column_view& expected_hierarchy,
int expected_depth,
rmm::cuda_stream_view stream,
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/copying/gather.cu b/cpp/src/copying/gather.cu
index a43e662478a0..ff6eae1c6443 100644
--- a/cpp/src/copying/gather.cu
+++ b/cpp/src/copying/gather.cu
@@ -28,7 +28,7 @@ std::unique_ptr gather(table_view const& source_table,
out_of_bounds_policy bounds_policy,
negative_index_policy neg_indices,
cuda::stream_ref stream,
- rmm::device_async_resource_ref mr)
+ cudf::memory_resources mr)
{
CUDF_EXPECTS(not gather_map.has_nulls(), "gather_map contains nulls", std::invalid_argument);
@@ -55,7 +55,7 @@ std::unique_ptr gather(table_view const& source_table,
out_of_bounds_policy bounds_policy,
negative_index_policy neg_indices,
cuda::stream_ref stream,
- rmm::device_async_resource_ref mr)
+ cudf::memory_resources mr)
{
CUDF_EXPECTS(gather_map.size() <= static_cast(std::numeric_limits::max()),
"gather map size exceeds the column size limit",
@@ -74,7 +74,7 @@ std::unique_ptr gather(table_view const& source_table,
column_view const& gather_map,
out_of_bounds_policy bounds_policy,
cuda::stream_ref stream,
- rmm::device_async_resource_ref mr)
+ cudf::memory_resources mr)
{
CUDF_FUNC_RANGE();
@@ -89,7 +89,7 @@ std::unique_ptr gather(table_view const& source_table,
out_of_bounds_policy bounds_policy,
negative_index_policy neg_indices,
cuda::stream_ref stream,
- rmm::device_async_resource_ref mr)
+ cudf::memory_resources mr)
{
CUDF_FUNC_RANGE();
return detail::gather(source_table, gather_map, bounds_policy, neg_indices, stream, mr);
diff --git a/cpp/src/dictionary/decode.cu b/cpp/src/dictionary/decode.cu
index 0338a939ed7d..0249d0bf66cc 100644
--- a/cpp/src/dictionary/decode.cu
+++ b/cpp/src/dictionary/decode.cu
@@ -36,13 +36,13 @@ struct indices_handler_fn {
*/
std::unique_ptr decode(dictionary_column_view const& source,
cuda::stream_ref stream,
- rmm::device_async_resource_ref mr)
+ cudf::memory_resources mr)
{
if (source.is_empty()) return make_empty_column(type_id::EMPTY);
// 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, mr.get_temporary_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()});
@@ -57,8 +57,8 @@ std::unique_ptr decode(dictionary_column_view const& source,
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),
- source.null_count());
+ output_column->set_null_mask(
+ cudf::detail::copy_bitmask(source.parent(), stream, mr.get_output_mr()), source.null_count());
return output_column;
}
@@ -67,7 +67,7 @@ std::unique_ptr decode(dictionary_column_view const& source,
std::unique_ptr decode(dictionary_column_view const& source,
cuda::stream_ref 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 e3e428c14ace..1380dd0e01d1 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 7f43eb337370..eaab38d2bea3 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,
cuda::stream_ref 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.get()};
auto set_ref = set.ref(cuco::insert_and_find);
@@ -96,22 +99,20 @@ 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.get());
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;
@@ -124,7 +125,7 @@ std::unique_ptr encode(column_view const& input,
// 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,
cuda::stream_ref 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 269cb403998c..327a80c4f1ee 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