From 76f0fe36cf2f9ff592c16fa66fe36acfbb6bc1cf Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Wed, 15 Mar 2023 11:57:36 -0700 Subject: [PATCH 01/13] extract --- cpp/include/cudf/lists/detail/extract.hpp | 20 +++++++++----------- 1 file changed, 9 insertions(+), 11 deletions(-) diff --git a/cpp/include/cudf/lists/detail/extract.hpp b/cpp/include/cudf/lists/detail/extract.hpp index 44c31c9ddb2..013f9b491dd 100644 --- a/cpp/include/cudf/lists/detail/extract.hpp +++ b/cpp/include/cudf/lists/detail/extract.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022, NVIDIA CORPORATION. + * Copyright (c) 2022-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -27,22 +27,20 @@ namespace detail { * rmm::mr::device_memory_resource*) * @param stream CUDA stream used for device memory operations and kernel launches. */ -std::unique_ptr extract_list_element( - lists_column_view lists_column, - size_type const index, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr extract_list_element(lists_column_view lists_column, + size_type const index, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @copydoc cudf::lists::extract_list_element(lists_column_view, column_view const&, * rmm::mr::device_memory_resource*) * @param stream CUDA stream used for device memory operations and kernel launches. */ -std::unique_ptr extract_list_element( - lists_column_view lists_column, - column_view const& indices, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr extract_list_element(lists_column_view lists_column, + column_view const& indices, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); } // namespace detail } // namespace lists From af51dd5661ecd66ebdc178e72725f1689b872408 Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Wed, 15 Mar 2023 12:08:44 -0700 Subject: [PATCH 02/13] interleave_columns --- cpp/include/cudf/lists/detail/interleave_columns.hpp | 11 +++++------ 1 file changed, 5 insertions(+), 6 deletions(-) diff --git a/cpp/include/cudf/lists/detail/interleave_columns.hpp b/cpp/include/cudf/lists/detail/interleave_columns.hpp index 7ae90779fdc..a5cf67c95b9 100644 --- a/cpp/include/cudf/lists/detail/interleave_columns.hpp +++ b/cpp/include/cudf/lists/detail/interleave_columns.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021, NVIDIA CORPORATION. + * Copyright (c) 2021-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -44,11 +44,10 @@ namespace detail { * @param mr Device memory resource used to allocate the returned column's device memory. * @return The interleaved columns as a single column. */ -std::unique_ptr interleave_columns( - table_view const& input, - bool has_null_mask, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr interleave_columns(table_view const& input, + bool has_null_mask, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); } // namespace detail } // namespace lists From a35d134b2d79df81ac657785a9bbd81e23301eae Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Wed, 15 Mar 2023 12:11:41 -0700 Subject: [PATCH 03/13] reverse --- cpp/include/cudf/lists/detail/reverse.hpp | 9 ++++----- 1 file changed, 4 insertions(+), 5 deletions(-) diff --git a/cpp/include/cudf/lists/detail/reverse.hpp b/cpp/include/cudf/lists/detail/reverse.hpp index d467a9ac70e..6e3b952a3b0 100644 --- a/cpp/include/cudf/lists/detail/reverse.hpp +++ b/cpp/include/cudf/lists/detail/reverse.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022, NVIDIA CORPORATION. + * Copyright (c) 2022-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -23,9 +23,8 @@ namespace cudf::lists::detail { * @copydoc cudf::lists::reverse * @param stream CUDA stream used for device memory operations and kernel launches */ -std::unique_ptr reverse( - lists_column_view const& input, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr reverse(lists_column_view const& input, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); } // namespace cudf::lists::detail From 45fcebfb070e6dbe46f87291ca6de1c43d36c740 Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Wed, 15 Mar 2023 12:13:56 -0700 Subject: [PATCH 04/13] scatter --- cpp/include/cudf/lists/detail/scatter.cuh | 43 +++++++++++------------ 1 file changed, 20 insertions(+), 23 deletions(-) diff --git a/cpp/include/cudf/lists/detail/scatter.cuh b/cpp/include/cudf/lists/detail/scatter.cuh index c2b4778aac8..856914b445e 100644 --- a/cpp/include/cudf/lists/detail/scatter.cuh +++ b/cpp/include/cudf/lists/detail/scatter.cuh @@ -89,15 +89,14 @@ rmm::device_uvector list_vector_from_column( * @return New lists column. */ template -std::unique_ptr scatter_impl( - rmm::device_uvector const& source_vector, - rmm::device_uvector& target_vector, - MapIterator scatter_map_begin, - MapIterator scatter_map_end, - column_view const& source, - column_view const& target, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) +std::unique_ptr scatter_impl(rmm::device_uvector const& source_vector, + rmm::device_uvector& target_vector, + MapIterator scatter_map_begin, + MapIterator scatter_map_end, + column_view const& source, + column_view const& target, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) { CUDF_EXPECTS(column_types_equal(source, target), "Mismatched column types."); @@ -170,13 +169,12 @@ std::unique_ptr scatter_impl( * @return New lists column. */ template -std::unique_ptr scatter( - column_view const& source, - MapIterator scatter_map_begin, - MapIterator scatter_map_end, - column_view const& target, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) +std::unique_ptr scatter(column_view const& source, + MapIterator scatter_map_begin, + MapIterator scatter_map_end, + column_view const& target, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) { auto const num_rows = target.size(); if (num_rows == 0) { return cudf::empty_like(target); } @@ -227,13 +225,12 @@ std::unique_ptr scatter( * @return New lists column. */ template -std::unique_ptr scatter( - scalar const& slr, - MapIterator scatter_map_begin, - MapIterator scatter_map_end, - column_view const& target, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) +std::unique_ptr scatter(scalar const& slr, + MapIterator scatter_map_begin, + MapIterator scatter_map_end, + column_view const& target, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) { auto const num_rows = target.size(); if (num_rows == 0) { return cudf::empty_like(target); } From 6db24efa93261a563076bdfc9c41930f0064bb46 Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Wed, 15 Mar 2023 12:16:14 -0700 Subject: [PATCH 05/13] sorting --- cpp/include/cudf/lists/detail/sorting.hpp | 24 +++++++++++------------ 1 file changed, 11 insertions(+), 13 deletions(-) diff --git a/cpp/include/cudf/lists/detail/sorting.hpp b/cpp/include/cudf/lists/detail/sorting.hpp index 1068a4c4b69..c378ca8cf06 100644 --- a/cpp/include/cudf/lists/detail/sorting.hpp +++ b/cpp/include/cudf/lists/detail/sorting.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021, NVIDIA CORPORATION. + * Copyright (c) 2021-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -28,24 +28,22 @@ namespace detail { * * @param stream CUDA stream used for device memory operations and kernel launches. */ -std::unique_ptr sort_lists( - lists_column_view const& input, - order column_order, - null_order null_precedence, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr sort_lists(lists_column_view const& input, + order column_order, + null_order null_precedence, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @copydoc cudf::lists::stable_sort_lists * * @param stream CUDA stream used for device memory operations and kernel launches. */ -std::unique_ptr stable_sort_lists( - lists_column_view const& input, - order column_order, - null_order null_precedence, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr stable_sort_lists(lists_column_view const& input, + order column_order, + null_order null_precedence, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); } // namespace detail } // namespace lists From 06667802418dca1bbba89a763bdc75fbade97013 Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Wed, 15 Mar 2023 12:17:53 -0700 Subject: [PATCH 06/13] stream_compaction --- .../cudf/lists/detail/stream_compaction.hpp | 22 +++++++++---------- 1 file changed, 10 insertions(+), 12 deletions(-) diff --git a/cpp/include/cudf/lists/detail/stream_compaction.hpp b/cpp/include/cudf/lists/detail/stream_compaction.hpp index ba3dbb6594b..7ab9cf9a343 100644 --- a/cpp/include/cudf/lists/detail/stream_compaction.hpp +++ b/cpp/include/cudf/lists/detail/stream_compaction.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022, NVIDIA CORPORATION. + * Copyright (c) 2022-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -28,22 +28,20 @@ namespace cudf::lists::detail { * * @param stream CUDA stream used for device memory operations and kernel launches */ -std::unique_ptr apply_boolean_mask( - lists_column_view const& input, - lists_column_view const& boolean_mask, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr apply_boolean_mask(lists_column_view const& input, + lists_column_view const& boolean_mask, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @copydoc cudf::list::distinct * * @param stream CUDA stream used for device memory operations and kernel launches. */ -std::unique_ptr distinct( - lists_column_view const& input, - null_equality nulls_equal, - nan_equality nans_equal, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr distinct(lists_column_view const& input, + null_equality nulls_equal, + nan_equality nans_equal, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); } // namespace cudf::lists::detail From 8f719033fd1176914b804076b766fcaead884c7e Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Wed, 15 Mar 2023 12:23:01 -0700 Subject: [PATCH 07/13] lists concatenate --- cpp/include/cudf/lists/detail/concatenate.hpp | 9 ++++----- 1 file changed, 4 insertions(+), 5 deletions(-) diff --git a/cpp/include/cudf/lists/detail/concatenate.hpp b/cpp/include/cudf/lists/detail/concatenate.hpp index 5a8b4bc3bf3..a1f149d4ccf 100644 --- a/cpp/include/cudf/lists/detail/concatenate.hpp +++ b/cpp/include/cudf/lists/detail/concatenate.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2022, NVIDIA CORPORATION. + * Copyright (c) 2020-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -43,10 +43,9 @@ namespace detail { * @param mr Device memory resource used to allocate the returned column's device memory. * @return New column with concatenated results. */ -std::unique_ptr concatenate( - host_span columns, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr concatenate(host_span columns, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); } // namespace detail } // namespace lists From e2a32f4e599e72c7019730e05a635820cfddc471 Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Wed, 15 Mar 2023 12:24:23 -0700 Subject: [PATCH 08/13] dict concatenate --- cpp/src/dictionary/detail/concatenate.cu | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/cpp/src/dictionary/detail/concatenate.cu b/cpp/src/dictionary/detail/concatenate.cu index d4f3a9ca495..4706c2f9ef3 100644 --- a/cpp/src/dictionary/detail/concatenate.cu +++ b/cpp/src/dictionary/detail/concatenate.cu @@ -114,7 +114,8 @@ struct compute_children_offsets_fn { [](auto lhs, auto rhs) { return offsets_pair{lhs.first + rhs.first, lhs.second + rhs.second}; }); - return cudf::detail::make_device_uvector_sync(offsets, stream); + return cudf::detail::make_device_uvector_sync( + offsets, stream, rmm::mr::get_current_device_resource()); } private: From 3418568c85dcbaa7b295c7a11cc761039fa7e6a9 Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Wed, 15 Mar 2023 13:03:18 -0700 Subject: [PATCH 09/13] null mask --- cpp/include/cudf/detail/null_mask.hpp | 54 +++++++++------------ cpp/src/groupby/hash/groupby.cu | 4 +- cpp/src/groupby/sort/group_scan_util.cuh | 6 +-- cpp/src/groupby/sort/sort_helper.cu | 3 +- cpp/src/join/hash_join.cu | 48 +++++++++++++----- cpp/src/join/mixed_join.cu | 6 ++- cpp/src/join/mixed_join_semi.cu | 6 ++- cpp/src/scalar/scalar.cpp | 3 +- cpp/src/search/contains_table.cu | 5 +- cpp/src/stream_compaction/distinct_count.cu | 3 +- cpp/src/unary/cast_ops.cu | 4 +- 11 files changed, 87 insertions(+), 55 deletions(-) diff --git a/cpp/include/cudf/detail/null_mask.hpp b/cpp/include/cudf/detail/null_mask.hpp index a0e04d7b215..7f1b15893c5 100644 --- a/cpp/include/cudf/detail/null_mask.hpp +++ b/cpp/include/cudf/detail/null_mask.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2022, NVIDIA CORPORATION. + * Copyright (c) 2019-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -31,11 +31,10 @@ namespace detail { * * @param stream CUDA stream used for device memory operations and kernel launches. */ -rmm::device_buffer create_null_mask( - size_type size, - mask_state state, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +rmm::device_buffer create_null_mask(size_type size, + mask_state state, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @copydoc cudf::set_null_mask(bitmask_type*, size_type, size_type, bool) @@ -209,22 +208,20 @@ std::vector segmented_null_count(bitmask_type const* bitmask, * * @param stream CUDA stream used for device memory operations and kernel launches. */ -rmm::device_buffer copy_bitmask( - bitmask_type const* mask, - size_type begin_bit, - size_type end_bit, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +rmm::device_buffer copy_bitmask(bitmask_type const* mask, + size_type begin_bit, + size_type end_bit, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @copydoc cudf::copy_bitmask(column_view const& view, rmm::mr::device_memory_resource*) * * @param stream CUDA stream used for device memory operations and kernel launches. */ -rmm::device_buffer copy_bitmask( - column_view const& view, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +rmm::device_buffer copy_bitmask(column_view const& view, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @copydoc bitmask_and(host_span, host_span const, @@ -232,32 +229,29 @@ rmm::device_buffer copy_bitmask( * * @param stream CUDA stream used for device memory operations and kernel launches */ -std::pair bitmask_and( - host_span masks, - host_span masks_begin_bits, - size_type mask_size_bits, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::pair bitmask_and(host_span masks, + host_span masks_begin_bits, + size_type mask_size_bits, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @copydoc cudf::bitmask_and * * @param[in] stream CUDA stream used for device memory operations and kernel launches. */ -std::pair bitmask_and( - table_view const& view, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::pair bitmask_and(table_view const& view, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @copydoc cudf::bitmask_or * * @param[in] stream CUDA stream used for device memory operations and kernel launches. */ -std::pair bitmask_or( - table_view const& view, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::pair bitmask_or(table_view const& view, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @brief Performs a bitwise AND of the specified bitmasks, diff --git a/cpp/src/groupby/hash/groupby.cu b/cpp/src/groupby/hash/groupby.cu index 72ac6255549..93a8d19a3c8 100644 --- a/cpp/src/groupby/hash/groupby.cu +++ b/cpp/src/groupby/hash/groupby.cu @@ -486,7 +486,9 @@ void compute_single_pass_aggs(table_view const& keys, keys_have_nulls and include_null_keys == null_policy::EXCLUDE; auto row_bitmask = - skip_key_rows_with_nulls ? cudf::detail::bitmask_and(keys, stream).first : rmm::device_buffer{}; + skip_key_rows_with_nulls + ? cudf::detail::bitmask_and(keys, stream, rmm::mr::get_current_device_resource()).first + : rmm::device_buffer{}; thrust::for_each_n(rmm::exec_policy(stream), thrust::make_counting_iterator(0), diff --git a/cpp/src/groupby/sort/group_scan_util.cuh b/cpp/src/groupby/sort/group_scan_util.cuh index cb954e614f2..f12efd3cd24 100644 --- a/cpp/src/groupby/sort/group_scan_util.cuh +++ b/cpp/src/groupby/sort/group_scan_util.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2022, NVIDIA CORPORATION. + * Copyright (c) 2021-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -124,7 +124,7 @@ struct group_scan_functor() make_null_replacement_iterator(*values_view, OpType::template identity()), thrust::identity{}); do_scan(input, result_view->begin(), OpType{}); - result->set_null_mask(cudf::detail::copy_bitmask(values, stream)); + result->set_null_mask(cudf::detail::copy_bitmask(values, stream, mr)); } else { auto input = thrust::make_transform_iterator(values_view->begin(), thrust::identity{}); @@ -175,7 +175,7 @@ struct group_scan_functorset_null_mask(cudf::detail::copy_bitmask(values, stream), values.null_count()); + results->set_null_mask(cudf::detail::copy_bitmask(values, stream, mr), values.null_count()); return results; } }; diff --git a/cpp/src/groupby/sort/sort_helper.cu b/cpp/src/groupby/sort/sort_helper.cu index ebafcd75e6d..6e992f2f53b 100644 --- a/cpp/src/groupby/sort/sort_helper.cu +++ b/cpp/src/groupby/sort/sort_helper.cu @@ -223,7 +223,8 @@ column_view sort_groupby_helper::keys_bitmask_column(rmm::cuda_stream_view strea { if (_keys_bitmask_column) return _keys_bitmask_column->view(); - auto [row_bitmask, null_count] = cudf::detail::bitmask_and(_keys, stream); + auto [row_bitmask, null_count] = + cudf::detail::bitmask_and(_keys, stream, rmm::mr::get_current_device_resource()); _keys_bitmask_column = make_numeric_column( data_type(type_id::INT8), _keys.num_rows(), std::move(row_bitmask), null_count, stream); diff --git a/cpp/src/join/hash_join.cu b/cpp/src/join/hash_join.cu index 7fb35e179e9..d0bdad73614 100644 --- a/cpp/src/join/hash_join.cu +++ b/cpp/src/join/hash_join.cu @@ -284,7 +284,8 @@ hash_join::hash_join(cudf::table_view const& build, cudf::null_equality compare_nulls, rmm::cuda_stream_view stream) : _is_empty{build.num_rows() == 0}, - _composite_bitmask{cudf::detail::bitmask_and(build, stream).first}, + _composite_bitmask{ + cudf::detail::bitmask_and(build, stream, rmm::mr::get_current_device_resource()).first}, _nulls_equal{compare_nulls}, _hash_table{compute_hash_table_size(build.num_rows()), cuco::empty_key{std::numeric_limits::max()}, @@ -298,8 +299,13 @@ hash_join::hash_join(cudf::table_view const& build, "Build column size is too big for hash join"); // need to store off the owning structures for some of the views in _build - _flattened_build_table = structs::detail::flatten_nested_columns( - build, {}, {}, structs::detail::column_nullability::FORCE, stream); + _flattened_build_table = + structs::detail::flatten_nested_columns(build, + {}, + {}, + structs::detail::column_nullability::FORCE, + stream, + rmm::mr::get_current_device_resource()); _build = _flattened_build_table->flattened_columns(); if (_is_empty) { return; } @@ -356,8 +362,13 @@ std::size_t hash_join::inner_join_size(cudf::table_view const& probe, // Return directly if build table is empty if (_is_empty) { return 0; } - auto flattened_probe = structs::detail::flatten_nested_columns( - probe, {}, {}, structs::detail::column_nullability::FORCE, stream); + auto flattened_probe = + structs::detail::flatten_nested_columns(probe, + {}, + {}, + structs::detail::column_nullability::FORCE, + stream, + rmm::mr::get_current_device_resource()); auto const flattened_probe_table = flattened_probe->flattened_columns(); auto build_table_ptr = cudf::table_device_view::create(_build, stream); @@ -381,8 +392,13 @@ std::size_t hash_join::left_join_size(cudf::table_view const& probe, // Trivial left join case - exit early if (_is_empty) { return probe.num_rows(); } - auto flattened_probe = structs::detail::flatten_nested_columns( - probe, {}, {}, structs::detail::column_nullability::FORCE, stream); + auto flattened_probe = + structs::detail::flatten_nested_columns(probe, + {}, + {}, + structs::detail::column_nullability::FORCE, + stream, + rmm::mr::get_current_device_resource()); auto const flattened_probe_table = flattened_probe->flattened_columns(); auto build_table_ptr = cudf::table_device_view::create(_build, stream); @@ -407,8 +423,13 @@ std::size_t hash_join::full_join_size(cudf::table_view const& probe, // Trivial left join case - exit early if (_is_empty) { return probe.num_rows(); } - auto flattened_probe = structs::detail::flatten_nested_columns( - probe, {}, {}, structs::detail::column_nullability::FORCE, stream); + auto flattened_probe = + structs::detail::flatten_nested_columns(probe, + {}, + {}, + structs::detail::column_nullability::FORCE, + stream, + rmm::mr::get_current_device_resource()); auto const flattened_probe_table = flattened_probe->flattened_columns(); auto build_table_ptr = cudf::table_device_view::create(_build, stream); @@ -474,8 +495,13 @@ hash_join::compute_hash_join(cudf::table_view const& probe, CUDF_EXPECTS(probe.num_rows() < cudf::detail::MAX_JOIN_SIZE, "Probe column size is too big for hash join"); - auto flattened_probe = structs::detail::flatten_nested_columns( - probe, {}, {}, structs::detail::column_nullability::FORCE, stream); + auto flattened_probe = + structs::detail::flatten_nested_columns(probe, + {}, + {}, + structs::detail::column_nullability::FORCE, + stream, + rmm::mr::get_current_device_resource()); auto const flattened_probe_table = flattened_probe->flattened_columns(); CUDF_EXPECTS(_build.num_columns() == flattened_probe_table.num_columns(), diff --git a/cpp/src/join/mixed_join.cu b/cpp/src/join/mixed_join.cu index 46e337a3363..fb403ae5a7e 100644 --- a/cpp/src/join/mixed_join.cu +++ b/cpp/src/join/mixed_join.cu @@ -136,7 +136,8 @@ mixed_join( // TODO: To add support for nested columns we will need to flatten in many // places. However, this probably isn't worth adding any time soon since we // won't be able to support AST conditions for those types anyway. - auto const row_bitmask = cudf::detail::bitmask_and(build, stream).first; + auto const row_bitmask = + cudf::detail::bitmask_and(build, stream, rmm::mr::get_current_device_resource()).first; build_join_hash_table( build, hash_table, compare_nulls, static_cast(row_bitmask.data()), stream); auto hash_table_view = hash_table.get_device_view(); @@ -384,7 +385,8 @@ compute_mixed_join_output_size(table_view const& left_equality, // TODO: To add support for nested columns we will need to flatten in many // places. However, this probably isn't worth adding any time soon since we // won't be able to support AST conditions for those types anyway. - auto const row_bitmask = cudf::detail::bitmask_and(build, stream).first; + auto const row_bitmask = + cudf::detail::bitmask_and(build, stream, rmm::mr::get_current_device_resource()).first; build_join_hash_table( build, hash_table, compare_nulls, static_cast(row_bitmask.data()), stream); auto hash_table_view = hash_table.get_device_view(); diff --git a/cpp/src/join/mixed_join_semi.cu b/cpp/src/join/mixed_join_semi.cu index b32df9316e2..164206ee56d 100644 --- a/cpp/src/join/mixed_join_semi.cu +++ b/cpp/src/join/mixed_join_semi.cu @@ -193,7 +193,8 @@ std::unique_ptr> mixed_join_semi( hash_table.insert(iter, iter + right_num_rows, hash_build, equality_build, stream.value()); } else { thrust::counting_iterator stencil(0); - auto const [row_bitmask, _] = cudf::detail::bitmask_and(build, stream); + auto const [row_bitmask, _] = + cudf::detail::bitmask_and(build, stream, rmm::mr::get_current_device_resource()); row_is_valid pred{static_cast(row_bitmask.data())}; // insert valid rows @@ -431,7 +432,8 @@ compute_mixed_join_output_size_semi(table_view const& left_equality, hash_table.insert(iter, iter + right_num_rows, hash_build, equality_build, stream.value()); } else { thrust::counting_iterator stencil(0); - auto const [row_bitmask, _] = cudf::detail::bitmask_and(build, stream); + auto const [row_bitmask, _] = + cudf::detail::bitmask_and(build, stream, rmm::mr::get_current_device_resource()); row_is_valid pred{static_cast(row_bitmask.data())}; // insert valid rows diff --git a/cpp/src/scalar/scalar.cpp b/cpp/src/scalar/scalar.cpp index 046bfee9e41..403dc8c9189 100644 --- a/cpp/src/scalar/scalar.cpp +++ b/cpp/src/scalar/scalar.cpp @@ -583,7 +583,8 @@ table struct_scalar::init_data(table&& data, auto data_cols = data.release(); // push validity mask down - auto const validity = cudf::detail::create_null_mask(1, mask_state::ALL_NULL, stream); + auto const validity = cudf::detail::create_null_mask( + 1, mask_state::ALL_NULL, stream, rmm::mr::get_current_device_resource()); for (auto& col : data_cols) { col = cudf::structs::detail::superimpose_nulls( static_cast(validity.data()), 1, std::move(col), stream, mr); diff --git a/cpp/src/search/contains_table.cu b/cpp/src/search/contains_table.cu index f770b4598cf..8f82f426aea 100644 --- a/cpp/src/search/contains_table.cu +++ b/cpp/src/search/contains_table.cu @@ -125,7 +125,10 @@ std::pair build_row_bitmask(table_view // If there are more than one nullable column, we compute `bitmask_and` of their null masks. // Otherwise, we have only one nullable column and can use its null mask directly. if (nullable_columns.size() > 1) { - auto row_bitmask = cudf::detail::bitmask_and(table_view{nullable_columns}, stream).first; + auto row_bitmask = + cudf::detail::bitmask_and( + table_view{nullable_columns}, stream, rmm::mr::get_current_device_resource()) + .first; auto const row_bitmask_ptr = static_cast(row_bitmask.data()); return std::pair(std::move(row_bitmask), row_bitmask_ptr); } diff --git a/cpp/src/stream_compaction/distinct_count.cu b/cpp/src/stream_compaction/distinct_count.cu index 0dae26c18a9..8c50f8d29e8 100644 --- a/cpp/src/stream_compaction/distinct_count.cu +++ b/cpp/src/stream_compaction/distinct_count.cu @@ -151,7 +151,8 @@ cudf::size_type distinct_count(table_view const& keys, // when nulls are equal, insert non-null rows only to improve efficiency if (nulls_equal == null_equality::EQUAL and has_nulls) { thrust::counting_iterator stencil(0); - auto const [row_bitmask, null_count] = cudf::detail::bitmask_or(keys, stream); + auto const [row_bitmask, null_count] = + cudf::detail::bitmask_or(keys, stream, rmm::mr::get_current_device_resource()); row_validity pred{static_cast(row_bitmask.data())}; key_map.insert_if(iter, iter + num_rows, stencil, pred, hash_key, row_equal, stream.value()); diff --git a/cpp/src/unary/cast_ops.cu b/cpp/src/unary/cast_ops.cu index b569ce04c31..6e19fc2ca3f 100644 --- a/cpp/src/unary/cast_ops.cu +++ b/cpp/src/unary/cast_ops.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2022, NVIDIA CORPORATION. + * Copyright (c) 2019-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -328,7 +328,7 @@ struct dispatch_unary_cast_to { auto output = std::make_unique(cudf::data_type{type.id(), input.type().scale()}, size, rmm::device_buffer{size * cudf::size_of(type), stream}, - copy_bitmask(input, stream), + copy_bitmask(input, stream, mr), input.null_count()); mutable_column_view output_mutable = *output; From 763048e3447824d3ea01466ff793a83b5512771d Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Wed, 15 Mar 2023 13:24:18 -0700 Subject: [PATCH 10/13] struct utilities and associated files --- cpp/include/cudf/detail/structs/utilities.hpp | 10 +- cpp/src/reductions/struct_minmax_util.cuh | 3 +- cpp/src/search/contains_table.cu | 14 +-- cpp/src/table/row_operators.cu | 8 +- cpp/tests/join/join_tests.cpp | 4 +- cpp/tests/structs/utilities_tests.cpp | 110 ++++++++++++------ 6 files changed, 95 insertions(+), 54 deletions(-) diff --git a/cpp/include/cudf/detail/structs/utilities.hpp b/cpp/include/cudf/detail/structs/utilities.hpp index 4a708d2fb51..5fcc331a382 100644 --- a/cpp/include/cudf/detail/structs/utilities.hpp +++ b/cpp/include/cudf/detail/structs/utilities.hpp @@ -175,7 +175,7 @@ class flattened_table { std::vector const& null_precedence, column_nullability nullability, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::mr::device_memory_resource* mr); /** * @brief Superimpose nulls from a given null mask into the input column, using bitwise AND. @@ -222,9 +222,7 @@ class flattened_table { * to be kept alive. */ [[nodiscard]] std::pair push_down_nulls( - column_view const& input, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + column_view const& input, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr); /** * @brief Push down nulls from columns of the input table into their children columns, using @@ -251,9 +249,7 @@ class flattened_table { * to be kept alive. */ [[nodiscard]] std::pair push_down_nulls( - table_view const& input, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + table_view const& input, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr); /** * @brief Checks if a column or any of its children is a struct column with structs that are null. diff --git a/cpp/src/reductions/struct_minmax_util.cuh b/cpp/src/reductions/struct_minmax_util.cuh index 796d10a3477..be6b5645739 100644 --- a/cpp/src/reductions/struct_minmax_util.cuh +++ b/cpp/src/reductions/struct_minmax_util.cuh @@ -102,7 +102,8 @@ class comparison_binop_generator { {}, std::vector{DEFAULT_NULL_ORDER}, cudf::structs::detail::column_nullability::MATCH_INCOMING, - stream)}, + stream, + rmm::mr::get_current_device_resource())}, d_flattened_input_ptr{ table_device_view::create(flattened_input->flattened_columns(), stream)}, is_min_op(is_min_op), diff --git a/cpp/src/search/contains_table.cu b/cpp/src/search/contains_table.cu index 8f82f426aea..1a2f242ef87 100644 --- a/cpp/src/search/contains_table.cu +++ b/cpp/src/search/contains_table.cu @@ -325,13 +325,13 @@ rmm::device_uvector contains_without_lists_or_nans(table_view const& hayst auto const has_any_nulls = haystack_has_nulls || needles_has_nulls; // Flatten the input tables. - auto const flatten_nullability = has_any_nulls - ? structs::detail::column_nullability::FORCE - : structs::detail::column_nullability::MATCH_INCOMING; - auto const haystack_flattened_tables = - structs::detail::flatten_nested_columns(haystack, {}, {}, flatten_nullability, stream); - auto const needles_flattened_tables = - structs::detail::flatten_nested_columns(needles, {}, {}, flatten_nullability, stream); + auto const flatten_nullability = has_any_nulls + ? structs::detail::column_nullability::FORCE + : structs::detail::column_nullability::MATCH_INCOMING; + auto const haystack_flattened_tables = structs::detail::flatten_nested_columns( + haystack, {}, {}, flatten_nullability, stream, rmm::mr::get_current_device_resource()); + auto const needles_flattened_tables = structs::detail::flatten_nested_columns( + needles, {}, {}, flatten_nullability, stream, rmm::mr::get_current_device_resource()); auto const haystack_flattened = haystack_flattened_tables->flattened_columns(); auto const needles_flattened = needles_flattened_tables->flattened_columns(); auto const haystack_tdv_ptr = table_device_view::create(haystack_flattened, stream); diff --git a/cpp/src/table/row_operators.cu b/cpp/src/table/row_operators.cu index 8a63a6f6411..73f6d85b941 100644 --- a/cpp/src/table/row_operators.cu +++ b/cpp/src/table/row_operators.cu @@ -25,6 +25,7 @@ #include #include +#include #include namespace cudf { @@ -397,9 +398,10 @@ std::shared_ptr preprocessed_table::create(table_view const& { check_eq_compatibility(t); - auto [null_pushed_table, nullable_data] = structs::detail::push_down_nulls(t, stream); - auto struct_offset_removed_table = remove_struct_child_offsets(null_pushed_table); - auto verticalized_t = std::get<0>(decompose_structs(struct_offset_removed_table)); + auto [null_pushed_table, nullable_data] = + structs::detail::push_down_nulls(t, stream, rmm::mr::get_current_device_resource()); + auto struct_offset_removed_table = remove_struct_child_offsets(null_pushed_table); + auto verticalized_t = std::get<0>(decompose_structs(struct_offset_removed_table)); auto d_t = table_device_view_owner(table_device_view::create(verticalized_t, stream)); return std::shared_ptr(new preprocessed_table( diff --git a/cpp/tests/join/join_tests.cpp b/cpp/tests/join/join_tests.cpp index 31500319592..404ff7d8380 100644 --- a/cpp/tests/join/join_tests.cpp +++ b/cpp/tests/join/join_tests.cpp @@ -1865,8 +1865,8 @@ TEST_F(JoinTest, Repro_StructsWithoutNullsPushedDown) // Note: Join result might not have nulls pushed down, since it's an output of gather(). // Must superimpose parent nulls before comparisons. - auto [superimposed_results, _] = - cudf::structs::detail::push_down_nulls(*result, cudf::get_default_stream()); + auto [superimposed_results, _] = cudf::structs::detail::push_down_nulls( + *result, cudf::get_default_stream(), rmm::mr::get_current_device_resource()); auto const expected = [] { auto fact_ints = ints{0}; diff --git a/cpp/tests/structs/utilities_tests.cpp b/cpp/tests/structs/utilities_tests.cpp index e92b96553c0..327fede6126 100644 --- a/cpp/tests/structs/utilities_tests.cpp +++ b/cpp/tests/structs/utilities_tests.cpp @@ -53,9 +53,14 @@ TYPED_TEST(TypedStructUtilitiesTest, ListsAtTopLevel) auto lists_col = lists{{0, 1}, {22, 33}, {44, 55, 66}}; auto nums_col = nums{{0, 1, 2}, cudf::test::iterators::null_at(6)}; - auto table = cudf::table_view{{lists_col, nums_col}}; - auto flattened_table = cudf::structs::detail::flatten_nested_columns( - table, {}, {}, cudf::structs::detail::column_nullability::FORCE, cudf::get_default_stream()); + auto table = cudf::table_view{{lists_col, nums_col}}; + auto flattened_table = + cudf::structs::detail::flatten_nested_columns(table, + {}, + {}, + cudf::structs::detail::column_nullability::FORCE, + cudf::get_default_stream(), + rmm::mr::get_current_device_resource()); CUDF_TEST_EXPECT_TABLES_EQUAL(table, flattened_table->flattened_columns()); } @@ -76,7 +81,8 @@ TYPED_TEST(TypedStructUtilitiesTest, NestedListsUnsupported) {}, {}, cudf::structs::detail::column_nullability::FORCE, - cudf::get_default_stream()), + cudf::get_default_stream(), + rmm::mr::get_current_device_resource()), cudf::logic_error); } @@ -90,9 +96,14 @@ TYPED_TEST(TypedStructUtilitiesTest, NoStructs) {"", "1", "22", "333", "4444", "55555", "666666"}, cudf::test::iterators::null_at(1)}; auto nuther_nums_col = nums{{0, 1, 2, 3, 4, 5, 6}, cudf::test::iterators::null_at(6)}; - auto table = cudf::table_view{{nums_col, strings_col, nuther_nums_col}}; - auto flattened_table = cudf::structs::detail::flatten_nested_columns( - table, {}, {}, cudf::structs::detail::column_nullability::FORCE, cudf::get_default_stream()); + auto table = cudf::table_view{{nums_col, strings_col, nuther_nums_col}}; + auto flattened_table = + cudf::structs::detail::flatten_nested_columns(table, + {}, + {}, + cudf::structs::detail::column_nullability::FORCE, + cudf::get_default_stream(), + rmm::mr::get_current_device_resource()); CUDF_TEST_EXPECT_TABLES_EQUAL(table, flattened_table->flattened_columns()); } @@ -118,8 +129,13 @@ TYPED_TEST(TypedStructUtilitiesTest, SingleLevelStruct) auto expected = cudf::table_view{ {expected_nums_col_1, expected_structs_col, expected_nums_col_2, expected_strings_col}}; - auto flattened_table = cudf::structs::detail::flatten_nested_columns( - table, {}, {}, cudf::structs::detail::column_nullability::FORCE, cudf::get_default_stream()); + auto flattened_table = + cudf::structs::detail::flatten_nested_columns(table, + {}, + {}, + cudf::structs::detail::column_nullability::FORCE, + cudf::get_default_stream(), + rmm::mr::get_current_device_resource()); CUDF_TEST_EXPECT_TABLES_EQUAL(expected, flattened_table->flattened_columns()); } @@ -146,8 +162,13 @@ TYPED_TEST(TypedStructUtilitiesTest, SingleLevelStructWithNulls) auto expected = cudf::table_view{ {expected_nums_col_1, expected_structs_col, expected_nums_col_2, expected_strings_col}}; - auto flattened_table = cudf::structs::detail::flatten_nested_columns( - table, {}, {}, cudf::structs::detail::column_nullability::FORCE, cudf::get_default_stream()); + auto flattened_table = + cudf::structs::detail::flatten_nested_columns(table, + {}, + {}, + cudf::structs::detail::column_nullability::FORCE, + cudf::get_default_stream(), + rmm::mr::get_current_device_resource()); CUDF_TEST_EXPECT_TABLES_EQUAL(expected, flattened_table->flattened_columns()); } @@ -185,8 +206,13 @@ TYPED_TEST(TypedStructUtilitiesTest, StructOfStruct) expected_nums_col_3, expected_strings_col}}; - auto flattened_table = cudf::structs::detail::flatten_nested_columns( - table, {}, {}, cudf::structs::detail::column_nullability::FORCE, cudf::get_default_stream()); + auto flattened_table = + cudf::structs::detail::flatten_nested_columns(table, + {}, + {}, + cudf::structs::detail::column_nullability::FORCE, + cudf::get_default_stream(), + rmm::mr::get_current_device_resource()); CUDF_TEST_EXPECT_TABLES_EQUAL(expected, flattened_table->flattened_columns()); } @@ -225,8 +251,13 @@ TYPED_TEST(TypedStructUtilitiesTest, StructOfStructWithNullsAtLeafLevel) expected_nums_col_3, expected_strings_col}}; - auto flattened_table = cudf::structs::detail::flatten_nested_columns( - table, {}, {}, cudf::structs::detail::column_nullability::FORCE, cudf::get_default_stream()); + auto flattened_table = + cudf::structs::detail::flatten_nested_columns(table, + {}, + {}, + cudf::structs::detail::column_nullability::FORCE, + cudf::get_default_stream(), + rmm::mr::get_current_device_resource()); CUDF_TEST_EXPECT_TABLES_EQUAL(expected, flattened_table->flattened_columns()); } @@ -266,8 +297,13 @@ TYPED_TEST(TypedStructUtilitiesTest, StructOfStructWithNullsAtTopLevel) expected_nums_col_3, expected_strings_col}}; - auto flattened_table = cudf::structs::detail::flatten_nested_columns( - table, {}, {}, cudf::structs::detail::column_nullability::FORCE, cudf::get_default_stream()); + auto flattened_table = + cudf::structs::detail::flatten_nested_columns(table, + {}, + {}, + cudf::structs::detail::column_nullability::FORCE, + cudf::get_default_stream(), + rmm::mr::get_current_device_resource()); CUDF_TEST_EXPECT_TABLES_EQUAL(expected, flattened_table->flattened_columns()); } @@ -307,8 +343,13 @@ TYPED_TEST(TypedStructUtilitiesTest, StructOfStructWithNullsAtAllLevels) expected_nums_col_3, expected_strings_col}}; - auto flattened_table = cudf::structs::detail::flatten_nested_columns( - table, {}, {}, cudf::structs::detail::column_nullability::FORCE, cudf::get_default_stream()); + auto flattened_table = + cudf::structs::detail::flatten_nested_columns(table, + {}, + {}, + cudf::structs::detail::column_nullability::FORCE, + cudf::get_default_stream(), + rmm::mr::get_current_device_resource()); CUDF_TEST_EXPECT_TABLES_EQUAL(expected, flattened_table->flattened_columns()); } @@ -330,7 +371,8 @@ TYPED_TEST(TypedStructUtilitiesTest, ListsAreUnsupported) {}, {}, cudf::structs::detail::column_nullability::FORCE, - cudf::get_default_stream()), + cudf::get_default_stream(), + rmm::mr::get_current_device_resource()), cudf::logic_error); } @@ -346,8 +388,8 @@ TYPED_TEST_SUITE(TypedSuperimposeTest, cudf::test::FixedWidthTypes); void test_non_struct_columns(cudf::column_view const& input) { // push_down_nulls() on non-struct columns should return the input column, unchanged. - auto [superimposed, backing_data] = - cudf::structs::detail::push_down_nulls(input, cudf::get_default_stream()); + auto [superimposed, backing_data] = cudf::structs::detail::push_down_nulls( + input, cudf::get_default_stream(), rmm::mr::get_current_device_resource()); CUDF_TEST_EXPECT_COLUMNS_EQUAL(input, superimposed); EXPECT_TRUE(backing_data.new_null_masks.empty()); @@ -410,8 +452,8 @@ TYPED_TEST(TypedSuperimposeTest, BasicStruct) CUDF_TEST_EXPECT_COLUMNS_EQUAL(structs_view.child(1), make_lists_member(cudf::test::iterators::nulls_at({4, 5}))); - auto [output, backing_data] = - cudf::structs::detail::push_down_nulls(structs_view, cudf::get_default_stream()); + auto [output, backing_data] = cudf::structs::detail::push_down_nulls( + structs_view, cudf::get_default_stream(), rmm::mr::get_current_device_resource()); // After push_down_nulls(), the struct nulls (i.e. at index-0) should have been pushed // down to the children. All members should have nulls at row-index 0. @@ -436,8 +478,8 @@ TYPED_TEST(TypedSuperimposeTest, NonNullableParentStruct) cudf::test::iterators::no_nulls()} .release(); - auto [output, backing_data] = - cudf::structs::detail::push_down_nulls(structs_input->view(), cudf::get_default_stream()); + auto [output, backing_data] = cudf::structs::detail::push_down_nulls( + structs_input->view(), cudf::get_default_stream(), rmm::mr::get_current_device_resource()); // After push_down_nulls(), none of the child structs should have changed, // because the parent had no nulls to begin with. @@ -471,8 +513,8 @@ TYPED_TEST(TypedSuperimposeTest, NestedStruct_ChildNullable_ParentNonNullable) auto structs_of_structs = cudf::test::structs_column_wrapper{std::move(outer_struct_members)}.release(); - auto [output, backing_data] = - cudf::structs::detail::push_down_nulls(structs_of_structs->view(), cudf::get_default_stream()); + auto [output, backing_data] = cudf::structs::detail::push_down_nulls( + structs_of_structs->view(), cudf::get_default_stream(), rmm::mr::get_current_device_resource()); // After push_down_nulls(), outer-struct column should not have pushed nulls to child // structs. But the child struct column must push its nulls to its own children. @@ -514,8 +556,8 @@ TYPED_TEST(TypedSuperimposeTest, NestedStruct_ChildNullable_ParentNullable) cudf::detail::set_null_mask( structs_of_structs_view.null_mask(), 1, 2, false, cudf::get_default_stream()); - auto [output, backing_data] = - cudf::structs::detail::push_down_nulls(structs_of_structs->view(), cudf::get_default_stream()); + auto [output, backing_data] = cudf::structs::detail::push_down_nulls( + structs_of_structs->view(), cudf::get_default_stream(), rmm::mr::get_current_device_resource()); // After push_down_nulls(), outer-struct column should not have pushed nulls to child // structs. But the child struct column must push its nulls to its own children. @@ -570,8 +612,8 @@ TYPED_TEST(TypedSuperimposeTest, Struct_Sliced) // nums_member: 11011 // lists_member: 00111 - auto [output, backing_data] = - cudf::structs::detail::push_down_nulls(sliced_structs, cudf::get_default_stream()); + auto [output, backing_data] = cudf::structs::detail::push_down_nulls( + sliced_structs, cudf::get_default_stream(), rmm::mr::get_current_device_resource()); // After push_down_nulls(), the null masks should be: // STRUCT: 11110 @@ -623,8 +665,8 @@ TYPED_TEST(TypedSuperimposeTest, NestedStruct_Sliced) // nums_member: 11010 // lists_member: 00110 - auto [output, backing_data] = - cudf::structs::detail::push_down_nulls(sliced_structs, cudf::get_default_stream()); + auto [output, backing_data] = cudf::structs::detail::push_down_nulls( + sliced_structs, cudf::get_default_stream(), rmm::mr::get_current_device_resource()); // After push_down_nulls(), the null masks will be: // STRUCT: 11101 From 9906ca926f4af603eff67163cc2ebed5af9aa8ed Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Fri, 17 Mar 2023 12:20:54 -0700 Subject: [PATCH 11/13] Fix copyrights. --- cpp/src/dictionary/detail/concatenate.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/src/dictionary/detail/concatenate.cu b/cpp/src/dictionary/detail/concatenate.cu index 4706c2f9ef3..bc54f65bbd3 100644 --- a/cpp/src/dictionary/detail/concatenate.cu +++ b/cpp/src/dictionary/detail/concatenate.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2022, NVIDIA CORPORATION. + * Copyright (c) 2020-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. From b46554bb35ce80b3fdafef2bf6f06762a1d1f2bb Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Fri, 17 Mar 2023 12:22:58 -0700 Subject: [PATCH 12/13] Group includes. --- cpp/src/table/row_operators.cu | 1 + 1 file changed, 1 insertion(+) diff --git a/cpp/src/table/row_operators.cu b/cpp/src/table/row_operators.cu index 73f6d85b941..7b41392810a 100644 --- a/cpp/src/table/row_operators.cu +++ b/cpp/src/table/row_operators.cu @@ -26,6 +26,7 @@ #include #include + #include namespace cudf { From 32a91918a2a5053fa46236073c460a43a99c2ce9 Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Fri, 17 Mar 2023 15:42:36 -0700 Subject: [PATCH 13/13] Fix java --- java/src/main/native/src/ColumnVectorJni.cpp | 6 ++++-- java/src/main/native/src/ColumnViewJni.cu | 10 ++++++---- 2 files changed, 10 insertions(+), 6 deletions(-) diff --git a/java/src/main/native/src/ColumnVectorJni.cpp b/java/src/main/native/src/ColumnVectorJni.cpp index 6dc7de13560..1d22d8a5d79 100644 --- a/java/src/main/native/src/ColumnVectorJni.cpp +++ b/java/src/main/native/src/ColumnVectorJni.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2022, NVIDIA CORPORATION. + * Copyright (c) 2019-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -32,6 +32,7 @@ #include #include #include +#include #include "cudf_jni_apis.hpp" #include "dtype_utils.hpp" @@ -296,7 +297,8 @@ JNIEXPORT jlong JNICALL Java_ai_rapids_cudf_ColumnVector_concatenate(JNIEnv *env cudf::jni::native_jpointerArray{env, column_handles}.get_dereferenced(); auto const is_lists_column = columns[0].type().id() == cudf::type_id::LIST; return release_as_jlong( - is_lists_column ? cudf::lists::detail::concatenate(columns, cudf::get_default_stream()) : + is_lists_column ? cudf::lists::detail::concatenate(columns, cudf::get_default_stream(), + rmm::mr::get_current_device_resource()) : cudf::concatenate(columns)); } CATCH_STD(env, 0); diff --git a/java/src/main/native/src/ColumnViewJni.cu b/java/src/main/native/src/ColumnViewJni.cu index 8a2c0b2b411..fd8d4c5b9ae 100644 --- a/java/src/main/native/src/ColumnViewJni.cu +++ b/java/src/main/native/src/ColumnViewJni.cu @@ -163,7 +163,8 @@ void post_process_list_overlap(cudf::column_view const &lhs, cudf::column_view c auto [null_mask, null_count] = cudf::detail::bitmask_and( std::vector{ overlap_cv.null_mask(), static_cast(new_null_mask.data())}, - std::vector{0, 0}, overlap_cv.size(), stream); + std::vector{0, 0}, overlap_cv.size(), stream, + rmm::mr::get_current_device_resource()); overlap_result->set_null_mask(std::move(null_mask), null_count); } else { // Just set the output nullmask as the new nullmask. @@ -213,9 +214,10 @@ std::unique_ptr lists_distinct_by_key(cudf::lists_column_view cons cudf::detail::labels_to_offsets(labels_begin, labels_begin + out_labels.size(), offsets_begin, offsets_begin + out_offsets->size(), stream); - return cudf::make_lists_column(input.size(), std::move(out_offsets), std::move(out_structs), - input.null_count(), - cudf::detail::copy_bitmask(input.parent(), stream), stream); + return cudf::make_lists_column( + input.size(), std::move(out_offsets), std::move(out_structs), input.null_count(), + cudf::detail::copy_bitmask(input.parent(), stream, rmm::mr::get_current_device_resource()), + stream); } } // namespace cudf::jni