From b2a65e89ba68f1954bc36ada61267cfec40e806a Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Tue, 21 Mar 2023 22:03:13 -0400 Subject: [PATCH] Remove default detail mrs: part3 (#12966) This is the third PR in a sequence removing default mr parameters in detail APIs. Contributes to https://github.com/rapidsai/cudf/issues/12944. Authors: - Vyas Ramasubramani (https://github.com/vyasr) Approvers: - Nghia Truong (https://github.com/ttnghia) URL: https://github.com/rapidsai/cudf/pull/12966 --- cpp/benchmarks/common/generate_input.cu | 35 ++++++++--- cpp/benchmarks/join/join_common.hpp | 7 ++- cpp/include/cudf/detail/round.hpp | 13 ++-- cpp/include/cudf/detail/scatter.cuh | 15 +++-- cpp/include/cudf/detail/scatter.hpp | 48 +++++++------- cpp/include/cudf/detail/search.hpp | 15 +++-- cpp/include/cudf/detail/sequence.hpp | 33 +++++----- cpp/include/cudf/detail/transform.hpp | 63 ++++++++----------- cpp/include/cudf/detail/transpose.hpp | 9 ++- cpp/include/cudf/detail/unary.hpp | 52 +++++++-------- cpp/include/cudf/detail/valid_if.cuh | 13 ++-- .../cudf/io/text/detail/tile_state.hpp | 4 +- cpp/include/cudf/io/text/detail/trie.hpp | 6 +- cpp/include/cudf/lists/detail/contains.hpp | 42 ++++++------- cpp/src/groupby/sort/scan.cpp | 6 +- cpp/src/interop/to_arrow.cu | 2 +- cpp/src/io/text/multibyte_split.cu | 8 ++- cpp/src/join/semi_join.cu | 8 ++- cpp/src/lists/combine/concatenate_rows.cu | 8 ++- cpp/src/lists/extract.cu | 6 +- cpp/src/lists/set_operations.cu | 12 ++-- cpp/tests/bitmask/valid_if_tests.cu | 17 +++-- java/src/main/native/src/ColumnViewJni.cu | 7 ++- java/src/main/native/src/maps_column_view.cu | 8 ++- java/src/main/native/src/row_conversion.cu | 2 +- 25 files changed, 222 insertions(+), 217 deletions(-) diff --git a/cpp/benchmarks/common/generate_input.cu b/cpp/benchmarks/common/generate_input.cu index 545028260b8..762e9640d12 100644 --- a/cpp/benchmarks/common/generate_input.cu +++ b/cpp/benchmarks/common/generate_input.cu @@ -430,8 +430,12 @@ std::unique_ptr create_random_column(data_profile const& profile, null_mask.begin()); } - auto [result_bitmask, null_count] = cudf::detail::valid_if( - null_mask.begin(), null_mask.end(), thrust::identity{}, cudf::get_default_stream()); + auto [result_bitmask, null_count] = + cudf::detail::valid_if(null_mask.begin(), + null_mask.end(), + thrust::identity{}, + cudf::get_default_stream(), + rmm::mr::get_current_device_resource()); return std::make_unique( dtype, @@ -509,8 +513,12 @@ std::unique_ptr create_random_utf8_string_column(data_profile cons thrust::make_zip_iterator(offsets.begin(), offsets.begin() + 1), num_rows, string_generator{chars.data(), engine}); - auto [result_bitmask, null_count] = cudf::detail::valid_if( - null_mask.begin(), null_mask.end() - 1, thrust::identity{}, cudf::get_default_stream()); + auto [result_bitmask, null_count] = + cudf::detail::valid_if(null_mask.begin(), + null_mask.end() - 1, + thrust::identity{}, + cudf::get_default_stream(), + rmm::mr::get_current_device_resource()); return cudf::make_strings_column( num_rows, std::move(offsets), @@ -628,8 +636,11 @@ std::unique_ptr create_random_column(data_profi auto [null_mask, null_count] = [&]() { if (profile.get_null_probability().has_value()) { auto valids = valid_dist(engine, num_rows); - return cudf::detail::valid_if( - valids.begin(), valids.end(), thrust::identity{}, cudf::get_default_stream()); + return cudf::detail::valid_if(valids.begin(), + valids.end(), + thrust::identity{}, + cudf::get_default_stream(), + rmm::mr::get_current_device_resource()); } return std::pair{}; }(); @@ -712,9 +723,12 @@ std::unique_ptr create_random_column(data_profile auto offsets_column = std::make_unique( cudf::data_type{cudf::type_id::INT32}, num_rows + 1, offsets.release()); - auto [null_mask, null_count] = cudf::detail::valid_if( - valids.begin(), valids.end(), thrust::identity{}, cudf::get_default_stream()); - list_column = cudf::make_lists_column( + auto [null_mask, null_count] = cudf::detail::valid_if(valids.begin(), + valids.end(), + thrust::identity{}, + cudf::get_default_stream(), + rmm::mr::get_current_device_resource()); + list_column = cudf::make_lists_column( num_rows, std::move(offsets_column), std::move(current_child_column), @@ -840,7 +854,8 @@ std::pair create_random_null_mask( return cudf::detail::valid_if(thrust::make_counting_iterator(0), thrust::make_counting_iterator(size), bool_generator{seed, 1.0 - *null_probability}, - cudf::get_default_stream()); + cudf::get_default_stream(), + rmm::mr::get_current_device_resource()); } } diff --git a/cpp/benchmarks/join/join_common.hpp b/cpp/benchmarks/join/join_common.hpp index e37a4ca1193..70036a95377 100644 --- a/cpp/benchmarks/join/join_common.hpp +++ b/cpp/benchmarks/join/join_common.hpp @@ -104,8 +104,11 @@ void BM_join(state_type& state, Join JoinFunc) // roughly 75% nulls auto validity = thrust::make_transform_iterator(thrust::make_counting_iterator(0), null75_generator{}); - return cudf::detail::valid_if( - validity, validity + size, thrust::identity{}, cudf::get_default_stream()) + return cudf::detail::valid_if(validity, + validity + size, + thrust::identity{}, + cudf::get_default_stream(), + rmm::mr::get_current_device_resource()) .first; }; diff --git a/cpp/include/cudf/detail/round.hpp b/cpp/include/cudf/detail/round.hpp index 1e5612919f4..cdfc7caef37 100644 --- a/cpp/include/cudf/detail/round.hpp +++ b/cpp/include/cudf/detail/round.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. @@ -31,12 +31,11 @@ namespace detail { * * @param stream CUDA stream used for device memory operations and kernel launches. */ -std::unique_ptr round( - column_view const& input, - int32_t decimal_places, - rounding_method method, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr round(column_view const& input, + int32_t decimal_places, + rounding_method method, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); } // namespace detail } // namespace cudf diff --git a/cpp/include/cudf/detail/scatter.cuh b/cpp/include/cudf/detail/scatter.cuh index c8b17e22df2..dbf7bfa9527 100644 --- a/cpp/include/cudf/detail/scatter.cuh +++ b/cpp/include/cudf/detail/scatter.cuh @@ -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. @@ -386,13 +386,12 @@ struct column_scatterer_impl { * @return Result of scattering values from source to target */ template -std::unique_ptr scatter( - table_view const& source, - MapIterator scatter_map_begin, - MapIterator scatter_map_end, - table_view const& target, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) +std::unique_ptr
scatter(table_view const& source, + MapIterator scatter_map_begin, + MapIterator scatter_map_end, + table_view const& target, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/include/cudf/detail/scatter.hpp b/cpp/include/cudf/detail/scatter.hpp index 7c4b04537ea..39ae4fe1944 100644 --- a/cpp/include/cudf/detail/scatter.hpp +++ b/cpp/include/cudf/detail/scatter.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. @@ -59,12 +59,11 @@ namespace detail { * @param mr Device memory resource used to allocate the returned table's device memory * @return Result of scattering values from source to target */ -std::unique_ptr
scatter( - table_view const& source, - column_view const& scatter_map, - table_view const& target, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr
scatter(table_view const& source, + column_view const& scatter_map, + table_view const& target, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @copydoc cudf::detail::scatter(table_view const&,column_view const&,table_view @@ -72,12 +71,11 @@ std::unique_ptr
scatter( * * @throws cudf::logic_error if `scatter_map` span size is larger than max of `size_type`. */ -std::unique_ptr
scatter( - table_view const& source, - device_span const scatter_map, - table_view const& target, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr
scatter(table_view const& source, + device_span const scatter_map, + table_view const& target, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @brief Scatters a row of scalar values into a copy of the target table @@ -108,12 +106,11 @@ std::unique_ptr
scatter( * @param mr Device memory resource used to allocate the returned table's device memory * @return Result of scattering values from source to target */ -std::unique_ptr
scatter( - std::vector> const& source, - column_view const& indices, - table_view const& target, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr
scatter(std::vector> const& source, + column_view const& indices, + table_view const& target, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @copydoc cudf::boolean_mask_scatter( @@ -123,12 +120,11 @@ std::unique_ptr
scatter( * * @param stream CUDA stream used for device memory operations and kernel launches. */ -std::unique_ptr
boolean_mask_scatter( - table_view const& source, - table_view const& target, - 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
boolean_mask_scatter(table_view const& source, + table_view const& target, + column_view const& boolean_mask, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @copydoc cudf::boolean_mask_scatter( @@ -144,7 +140,7 @@ std::unique_ptr
boolean_mask_scatter( table_view const& target, column_view const& boolean_mask, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::mr::device_memory_resource* mr); } // namespace detail } // namespace cudf diff --git a/cpp/include/cudf/detail/search.hpp b/cpp/include/cudf/detail/search.hpp index 56d41fd635c..4c4ad7834f4 100644 --- a/cpp/include/cudf/detail/search.hpp +++ b/cpp/include/cudf/detail/search.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. @@ -89,12 +89,11 @@ std::unique_ptr contains(column_view const& haystack, * @param mr Device memory resource used to allocate the returned vector * @return A vector of bools indicating if each row in `needles` has matching rows in `haystack` */ -rmm::device_uvector contains( - table_view const& haystack, - table_view const& needles, - null_equality compare_nulls, - nan_equality compare_nans, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +rmm::device_uvector contains(table_view const& haystack, + table_view const& needles, + null_equality compare_nulls, + nan_equality compare_nans, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); } // namespace cudf::detail diff --git a/cpp/include/cudf/detail/sequence.hpp b/cpp/include/cudf/detail/sequence.hpp index 4a9bf5c74e1..3c3d1d0ed9e 100644 --- a/cpp/include/cudf/detail/sequence.hpp +++ b/cpp/include/cudf/detail/sequence.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. @@ -32,12 +32,11 @@ namespace detail { * * @param stream CUDA stream used for device memory operations and kernel launches. */ -std::unique_ptr sequence( - size_type size, - scalar const& init, - scalar const& step, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr sequence(size_type size, + scalar const& init, + scalar const& step, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @copydoc cudf::sequence(size_type size, scalar const& init, @@ -46,11 +45,10 @@ std::unique_ptr sequence( * * @param stream CUDA stream used for device memory operations and kernel launches. */ -std::unique_ptr sequence( - size_type size, - scalar const& init, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr sequence(size_type size, + scalar const& init, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @copydoc cudf::calendrical_month_sequence(size_type size, @@ -60,12 +58,11 @@ std::unique_ptr sequence( * * @param stream CUDA stream used for device memory operations and kernel launches. */ -std::unique_ptr calendrical_month_sequence( - size_type size, - scalar const& init, - size_type months, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr calendrical_month_sequence(size_type size, + scalar const& init, + size_type months, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); } // namespace detail } // namespace cudf diff --git a/cpp/include/cudf/detail/transform.hpp b/cpp/include/cudf/detail/transform.hpp index 8e19ebb8da7..5b64f61f11a 100644 --- a/cpp/include/cudf/detail/transform.hpp +++ b/cpp/include/cudf/detail/transform.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. @@ -29,24 +29,22 @@ namespace detail { * * @param stream CUDA stream used for device memory operations and kernel launches. */ -std::unique_ptr transform( - column_view const& input, - std::string const& unary_udf, - data_type output_type, - bool is_ptx, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr transform(column_view const& input, + std::string const& unary_udf, + data_type output_type, + bool is_ptx, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @copydoc cudf::compute_column * * @param stream CUDA stream used for device memory operations and kernel launches. */ -std::unique_ptr compute_column( - table_view const table, - ast::operation const& expr, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr compute_column(table_view const table, + ast::operation const& expr, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @copydoc cudf::nans_to_nulls @@ -54,9 +52,7 @@ std::unique_ptr compute_column( * @param stream CUDA stream used for device memory operations and kernel launches. */ std::pair, size_type> nans_to_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); /** * @copydoc cudf::bools_to_mask @@ -64,9 +60,7 @@ std::pair, size_type> nans_to_nulls( * @param stream CUDA stream used for device memory operations and kernel launches. */ std::pair, cudf::size_type> bools_to_mask( - 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); /** * @copydoc cudf::encode @@ -74,42 +68,37 @@ std::pair, cudf::size_type> bools_to_mask( * @param stream CUDA stream used for device memory operations and kernel launches. */ std::pair, std::unique_ptr> encode( - cudf::table_view const& input, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + cudf::table_view const& input, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr); /** * @copydoc cudf::one_hot_encode * * @param stream CUDA stream used for device memory operations and kernel launches. */ -std::pair, table_view> one_hot_encode( - column_view const& input, - column_view const& categories, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::pair, table_view> one_hot_encode(column_view const& input, + column_view const& categories, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @copydoc cudf::mask_to_bools * * @param stream CUDA stream used for device memory operations and kernel launches. */ -std::unique_ptr mask_to_bools( - bitmask_type const* null_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()); +std::unique_ptr mask_to_bools(bitmask_type const* null_mask, + size_type begin_bit, + size_type end_bit, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @copydoc cudf::row_bit_count * * @param stream CUDA stream used for device memory operations and kernel launches. */ -std::unique_ptr row_bit_count( - table_view const& t, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr row_bit_count(table_view const& t, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); } // namespace detail } // namespace cudf diff --git a/cpp/include/cudf/detail/transpose.hpp b/cpp/include/cudf/detail/transpose.hpp index 0470d625edc..d0be51860b2 100644 --- a/cpp/include/cudf/detail/transpose.hpp +++ b/cpp/include/cudf/detail/transpose.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. @@ -28,10 +28,9 @@ namespace detail { * * @param stream CUDA stream used for device memory operations and kernel launches. */ -std::pair, table_view> transpose( - table_view const& input, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::pair, table_view> transpose(table_view const& input, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); } // namespace detail } // namespace cudf diff --git a/cpp/include/cudf/detail/unary.hpp b/cpp/include/cudf/detail/unary.hpp index b7ecedc1489..3fbdf4a5a8f 100644 --- a/cpp/include/cudf/detail/unary.hpp +++ b/cpp/include/cudf/detail/unary.hpp @@ -45,13 +45,12 @@ namespace detail { */ template -std::unique_ptr true_if( - InputIterator begin, - InputIterator end, - size_type size, - Predicate p, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) +std::unique_ptr true_if(InputIterator begin, + InputIterator end, + size_type size, + Predicate p, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) { auto output = make_numeric_column(data_type(type_id::BOOL8), size, mask_state::UNALLOCATED, stream, mr); @@ -68,52 +67,47 @@ std::unique_ptr true_if( * * @param stream CUDA stream used for device memory operations and kernel launches. */ -std::unique_ptr unary_operation( - cudf::column_view const& input, - cudf::unary_operator op, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr unary_operation(cudf::column_view const& input, + cudf::unary_operator op, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @copydoc cudf::is_valid * * @param stream CUDA stream used for device memory operations and kernel launches. */ -std::unique_ptr is_valid( - cudf::column_view const& input, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr is_valid(cudf::column_view const& input, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @copydoc cudf::cast * * @param stream CUDA stream used for device memory operations and kernel launches. */ -std::unique_ptr cast( - column_view const& input, - data_type type, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr cast(column_view const& input, + data_type type, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @copydoc cudf::is_nan * * @param[in] stream CUDA stream used for device memory operations and kernel launches. */ -std::unique_ptr is_nan( - cudf::column_view const& input, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr is_nan(cudf::column_view const& input, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @copydoc cudf::is_not_nan * * @param[in] stream CUDA stream used for device memory operations and kernel launches. */ -std::unique_ptr is_not_nan( - cudf::column_view const& input, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr is_not_nan(cudf::column_view const& input, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); } // namespace detail } // namespace cudf diff --git a/cpp/include/cudf/detail/valid_if.cuh b/cpp/include/cudf/detail/valid_if.cuh index 04c78bed17d..76d6fd719a4 100644 --- a/cpp/include/cudf/detail/valid_if.cuh +++ b/cpp/include/cudf/detail/valid_if.cuh @@ -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. @@ -86,12 +86,11 @@ __global__ void valid_if_kernel( * null count */ template -std::pair valid_if( - InputIterator begin, - InputIterator end, - Predicate p, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) +std::pair valid_if(InputIterator begin, + InputIterator end, + Predicate p, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) { CUDF_EXPECTS(begin <= end, "Invalid range."); diff --git a/cpp/include/cudf/io/text/detail/tile_state.hpp b/cpp/include/cudf/io/text/detail/tile_state.hpp index bf833d4720c..6ae399fbe75 100644 --- a/cpp/include/cudf/io/text/detail/tile_state.hpp +++ b/cpp/include/cudf/io/text/detail/tile_state.hpp @@ -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. @@ -82,7 +82,7 @@ struct scan_tile_state { scan_tile_state(cudf::size_type num_tiles, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) + rmm::mr::device_memory_resource* mr) : tile_status(rmm::device_uvector>( num_tiles, stream, mr)), tile_state_partial(rmm::device_uvector(num_tiles, stream, mr)), diff --git a/cpp/include/cudf/io/text/detail/trie.hpp b/cpp/include/cudf/io/text/detail/trie.hpp index a908a9fa227..7bb2e4e2ece 100644 --- a/cpp/include/cudf/io/text/detail/trie.hpp +++ b/cpp/include/cudf/io/text/detail/trie.hpp @@ -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. @@ -165,7 +165,7 @@ struct trie { */ static trie create(std::string const& pattern, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) + rmm::mr::device_memory_resource* mr) { return create(std::vector{pattern}, stream, mr); @@ -181,7 +181,7 @@ struct trie { */ static trie create(std::vector const& patterns, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) + rmm::mr::device_memory_resource* mr) { std::vector tokens; std::vector transitions; diff --git a/cpp/include/cudf/lists/detail/contains.hpp b/cpp/include/cudf/lists/detail/contains.hpp index 24318e72e98..58ec18cb9ef 100644 --- a/cpp/include/cudf/lists/detail/contains.hpp +++ b/cpp/include/cudf/lists/detail/contains.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. @@ -29,12 +29,11 @@ namespace detail { * rmm::mr::device_memory_resource*) * @param stream CUDA stream used for device memory operations and kernel launches. */ -std::unique_ptr index_of( - cudf::lists_column_view const& lists, - cudf::scalar const& search_key, - cudf::lists::duplicate_find_option find_option, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr index_of(cudf::lists_column_view const& lists, + cudf::scalar const& search_key, + cudf::lists::duplicate_find_option find_option, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @copydoc cudf::lists::index_of(cudf::lists_column_view const&, @@ -43,12 +42,11 @@ std::unique_ptr index_of( * rmm::mr::device_memory_resource*) * @param stream CUDA stream used for device memory operations and kernel launches. */ -std::unique_ptr index_of( - cudf::lists_column_view const& lists, - cudf::column_view const& search_keys, - cudf::lists::duplicate_find_option find_option, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr index_of(cudf::lists_column_view const& lists, + cudf::column_view const& search_keys, + cudf::lists::duplicate_find_option find_option, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @copydoc cudf::lists::contains(cudf::lists_column_view const&, @@ -56,11 +54,10 @@ std::unique_ptr index_of( * rmm::mr::device_memory_resource*) * @param stream CUDA stream used for device memory operations and kernel launches. */ -std::unique_ptr contains( - cudf::lists_column_view const& lists, - cudf::scalar const& search_key, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr contains(cudf::lists_column_view const& lists, + cudf::scalar const& search_key, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @copydoc cudf::lists::contains(cudf::lists_column_view const&, @@ -68,11 +65,10 @@ std::unique_ptr contains( * rmm::mr::device_memory_resource*) * @param stream CUDA stream used for device memory operations and kernel launches. */ -std::unique_ptr contains( - cudf::lists_column_view const& lists, - cudf::column_view const& search_keys, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr contains(cudf::lists_column_view const& lists, + cudf::column_view const& search_keys, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); } // namespace detail } // namespace lists } // namespace cudf diff --git a/cpp/src/groupby/sort/scan.cpp b/cpp/src/groupby/sort/scan.cpp index 1aaa06750db..820dc8a3077 100644 --- a/cpp/src/groupby/sort/scan.cpp +++ b/cpp/src/groupby/sort/scan.cpp @@ -129,8 +129,10 @@ void scan_result_functor::operator()(aggregation const& agg) auto const group_labels_view = column_view(cudf::device_span(group_labels)); auto const gather_map = [&]() { if (is_presorted()) { // assumes both keys and values are sorted, Spark does this. - return cudf::detail::sequence( - group_labels.size(), *cudf::make_fixed_width_scalar(size_type{0}, stream), stream); + return cudf::detail::sequence(group_labels.size(), + *cudf::make_fixed_width_scalar(size_type{0}, stream), + stream, + rmm::mr::get_current_device_resource()); } else { auto sort_order = (rank_agg._method == rank_method::FIRST ? cudf::detail::stable_sorted_order : cudf::detail::sorted_order); diff --git a/cpp/src/interop/to_arrow.cu b/cpp/src/interop/to_arrow.cu index 861b5b0fba4..7f88019beb2 100644 --- a/cpp/src/interop/to_arrow.cu +++ b/cpp/src/interop/to_arrow.cu @@ -215,7 +215,7 @@ std::shared_ptr dispatch_to_arrow::operator()(column_view in arrow::MemoryPool* ar_mr, rmm::cuda_stream_view stream) { - auto bitmask = bools_to_mask(input, stream); + auto bitmask = bools_to_mask(input, stream, rmm::mr::get_current_device_resource()); auto data_buffer = allocate_arrow_buffer(static_cast(bitmask.first->size()), ar_mr); diff --git a/cpp/src/io/text/multibyte_split.cu b/cpp/src/io/text/multibyte_split.cu index a0ba3e3ee35..afa260e215a 100644 --- a/cpp/src/io/text/multibyte_split.cu +++ b/cpp/src/io/text/multibyte_split.cu @@ -379,9 +379,11 @@ std::unique_ptr multibyte_split(cudf::io::text::data_chunk_source // must be at least 32 when using warp-reduce on partials // must be at least 1 more than max possible concurrent tiles // best when at least 32 more than max possible concurrent tiles, due to rolling `invalid`s - auto num_tile_states = std::max(32, TILES_PER_CHUNK * concurrency + 32); - auto tile_multistates = scan_tile_state(num_tile_states, stream); - auto tile_offsets = scan_tile_state(num_tile_states, stream); + auto num_tile_states = std::max(32, TILES_PER_CHUNK * concurrency + 32); + auto tile_multistates = + scan_tile_state(num_tile_states, stream, rmm::mr::get_current_device_resource()); + auto tile_offsets = + scan_tile_state(num_tile_states, stream, rmm::mr::get_current_device_resource()); multibyte_split_init_kernel<<> left_semi_anti_join( // Previously, the gather map was generated directly without this array but by calling to // `map.contains` inside the `thrust::copy_if` kernel. However, that led to increasing register // usage and reducing performance, as reported here: https://github.com/rapidsai/cudf/pull/10511. - auto const flagged = - cudf::detail::contains(right_keys, left_keys, compare_nulls, nan_equality::ALL_EQUAL, stream); + auto const flagged = cudf::detail::contains(right_keys, + left_keys, + compare_nulls, + nan_equality::ALL_EQUAL, + stream, + rmm::mr::get_current_device_resource()); auto const left_num_rows = left_keys.num_rows(); auto gather_map = diff --git a/cpp/src/lists/combine/concatenate_rows.cu b/cpp/src/lists/combine/concatenate_rows.cu index 8b006548391..b890a0c82a2 100644 --- a/cpp/src/lists/combine/concatenate_rows.cu +++ b/cpp/src/lists/combine/concatenate_rows.cu @@ -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. @@ -246,7 +246,8 @@ std::unique_ptr concatenate_rows(table_view const& input, auto const row_index = i % num_rows; return row_null_counts[row_index] != num_columns; }, - stream); + stream, + rmm::mr::get_current_device_resource()); } // NULLIFY_OUTPUT_ROW. Output row is nullfied if any input row is null return cudf::detail::valid_if( @@ -257,7 +258,8 @@ std::unique_ptr concatenate_rows(table_view const& input, auto const row_index = i % num_rows; return row_null_counts[row_index] == 0; }, - stream); + stream, + rmm::mr::get_current_device_resource()); }(); concat->set_null_mask(std::move(null_mask), null_count); } diff --git a/cpp/src/lists/extract.cu b/cpp/src/lists/extract.cu index bff63871e29..5d4a20d1cb8 100644 --- a/cpp/src/lists/extract.cu +++ b/cpp/src/lists/extract.cu @@ -101,8 +101,10 @@ std::unique_ptr make_index_child(size_type index, */ std::unique_ptr make_index_offsets(size_type num_lists, rmm::cuda_stream_view stream) { - return cudf::detail::sequence( - num_lists + 1, cudf::scalar_type_t(0, true, stream), stream); + return cudf::detail::sequence(num_lists + 1, + cudf::scalar_type_t(0, true, stream), + stream, + rmm::mr::get_current_device_resource()); } } // namespace diff --git a/cpp/src/lists/set_operations.cu b/cpp/src/lists/set_operations.cu index eb3ec5a8236..c05ef2fd644 100644 --- a/cpp/src/lists/set_operations.cu +++ b/cpp/src/lists/set_operations.cu @@ -83,8 +83,8 @@ std::unique_ptr have_overlap(lists_column_view const& lhs, auto const rhs_table = table_view{{rhs_labels->view(), rhs_child}}; // Check existence for each row of the rhs_table in lhs_table. - auto const contained = - cudf::detail::contains(lhs_table, rhs_table, nulls_equal, nans_equal, stream); + auto const contained = cudf::detail::contains( + lhs_table, rhs_table, nulls_equal, nans_equal, stream, rmm::mr::get_current_device_resource()); auto const num_rows = lhs.size(); @@ -151,8 +151,8 @@ std::unique_ptr intersect_distinct(lists_column_view const& lhs, auto const lhs_table = table_view{{lhs_labels->view(), lhs_child}}; auto const rhs_table = table_view{{rhs_labels->view(), rhs_child}}; - auto const contained = - cudf::detail::contains(lhs_table, rhs_table, nulls_equal, nans_equal, stream); + auto const contained = cudf::detail::contains( + lhs_table, rhs_table, nulls_equal, nans_equal, stream, rmm::mr::get_current_device_resource()); auto const intersect_table = cudf::detail::copy_if( rhs_table, @@ -231,8 +231,8 @@ std::unique_ptr difference_distinct(lists_column_view const& lhs, auto const lhs_table = table_view{{lhs_labels->view(), lhs_child}}; auto const rhs_table = table_view{{rhs_labels->view(), rhs_child}}; - auto const contained = - cudf::detail::contains(rhs_table, lhs_table, nulls_equal, nans_equal, stream); + auto const contained = cudf::detail::contains( + rhs_table, lhs_table, nulls_equal, nans_equal, stream, rmm::mr::get_current_device_resource()); auto const difference_table = cudf::detail::copy_if( lhs_table, diff --git a/cpp/tests/bitmask/valid_if_tests.cu b/cpp/tests/bitmask/valid_if_tests.cu index cdc453be8e4..cb086cda179 100644 --- a/cpp/tests/bitmask/valid_if_tests.cu +++ b/cpp/tests/bitmask/valid_if_tests.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. @@ -43,7 +43,8 @@ TEST_F(ValidIfTest, EmptyRange) auto actual = cudf::detail::valid_if(thrust::make_counting_iterator(0), thrust::make_counting_iterator(0), odds_valid{}, - cudf::get_default_stream()); + cudf::get_default_stream(), + rmm::mr::get_current_device_resource()); auto const& buffer = actual.first; EXPECT_EQ(0u, buffer.size()); EXPECT_EQ(nullptr, buffer.data()); @@ -55,7 +56,8 @@ TEST_F(ValidIfTest, InvalidRange) EXPECT_THROW(cudf::detail::valid_if(thrust::make_counting_iterator(1), thrust::make_counting_iterator(0), odds_valid{}, - cudf::get_default_stream()), + cudf::get_default_stream(), + rmm::mr::get_current_device_resource()), cudf::logic_error); } @@ -66,7 +68,8 @@ TEST_F(ValidIfTest, OddsValid) auto actual = cudf::detail::valid_if(thrust::make_counting_iterator(0), thrust::make_counting_iterator(10000), odds_valid{}, - cudf::get_default_stream()); + cudf::get_default_stream(), + rmm::mr::get_current_device_resource()); CUDF_TEST_EXPECT_EQUAL_BUFFERS(expected.data(), actual.first.data(), expected.size()); EXPECT_EQ(5000, actual.second); } @@ -78,7 +81,8 @@ TEST_F(ValidIfTest, AllValid) auto actual = cudf::detail::valid_if(thrust::make_counting_iterator(0), thrust::make_counting_iterator(10000), all_valid{}, - cudf::get_default_stream()); + cudf::get_default_stream(), + rmm::mr::get_current_device_resource()); CUDF_TEST_EXPECT_EQUAL_BUFFERS(expected.data(), actual.first.data(), expected.size()); EXPECT_EQ(0, actual.second); } @@ -90,7 +94,8 @@ TEST_F(ValidIfTest, AllNull) auto actual = cudf::detail::valid_if(thrust::make_counting_iterator(0), thrust::make_counting_iterator(10000), all_null{}, - cudf::get_default_stream()); + cudf::get_default_stream(), + rmm::mr::get_current_device_resource()); CUDF_TEST_EXPECT_EQUAL_BUFFERS(expected.data(), actual.first.data(), expected.size()); EXPECT_EQ(10000, actual.second); } diff --git a/java/src/main/native/src/ColumnViewJni.cu b/java/src/main/native/src/ColumnViewJni.cu index 8a2c0b2b411..86c2add851a 100644 --- a/java/src/main/native/src/ColumnViewJni.cu +++ b/java/src/main/native/src/ColumnViewJni.cu @@ -56,7 +56,7 @@ new_column_with_boolean_column_as_validity(cudf::column_view const &exemplar, auto [null_mask, null_count] = cudf::detail::valid_if( validity_begin, validity_end, [] __device__(auto optional_bool) { return optional_bool.value_or(false); }, - cudf::get_default_stream()); + cudf::get_default_stream(), rmm::mr::get_current_device_resource()); auto const exemplar_without_null_mask = cudf::column_view{ exemplar.type(), exemplar.size(), @@ -153,8 +153,9 @@ void post_process_list_overlap(cudf::column_view const &lhs, cudf::column_view c }); // Create a new nullmask from the validity data. - auto [new_null_mask, new_null_count] = cudf::detail::valid_if( - validity.begin(), validity.end(), thrust::identity{}, cudf::get_default_stream()); + auto [new_null_mask, new_null_count] = + cudf::detail::valid_if(validity.begin(), validity.end(), thrust::identity{}, + cudf::get_default_stream(), rmm::mr::get_current_device_resource()); if (new_null_count > 0) { // If the `overlap_result` column is nullable, perform `bitmask_and` of its nullmask and the diff --git a/java/src/main/native/src/maps_column_view.cu b/java/src/main/native/src/maps_column_view.cu index 23254c0d501..1af7689f972 100644 --- a/java/src/main/native/src/maps_column_view.cu +++ b/java/src/main/native/src/maps_column_view.cu @@ -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. @@ -55,7 +55,8 @@ std::unique_ptr get_values_for_impl(maps_column_view const &maps_view, CUDF_EXPECTS(lookup_keys.type().id() == keys_.child().type().id(), "Lookup keys must have the same type as the keys of the map column."); auto key_indices = - lists::detail::index_of(keys_, lookup_keys, lists::duplicate_find_option::FIND_LAST, stream); + lists::detail::index_of(keys_, lookup_keys, lists::duplicate_find_option::FIND_LAST, stream, + rmm::mr::get_current_device_resource()); auto constexpr absent_offset = size_type{-1}; auto constexpr nullity_offset = std::numeric_limits::min(); thrust::replace(rmm::exec_policy(stream), key_indices->mutable_view().template begin(), @@ -86,7 +87,8 @@ std::unique_ptr contains_impl(maps_column_view const &maps_view, KeyT co auto const keys = maps_view.keys(); CUDF_EXPECTS(lookup_keys.type().id() == keys.child().type().id(), "Lookup keys must have the same type as the keys of the map column."); - auto const contains = lists::detail::contains(keys, lookup_keys, stream); + auto const contains = + lists::detail::contains(keys, lookup_keys, stream, rmm::mr::get_current_device_resource()); // Replace nulls with BOOL8{false}; auto const scalar_false = numeric_scalar{false, true, stream}; return detail::replace_nulls(contains->view(), scalar_false, stream, mr); diff --git a/java/src/main/native/src/row_conversion.cu b/java/src/main/native/src/row_conversion.cu index 747ff24f055..84f84f8b46f 100644 --- a/java/src/main/native/src/row_conversion.cu +++ b/java/src/main/native/src/row_conversion.cu @@ -1257,7 +1257,7 @@ static std::unique_ptr fixed_width_convert_to_rows( // Allocate and set the offsets row for the byte array std::unique_ptr offsets = - cudf::detail::sequence(num_rows + 1, zero, scalar_size_per_row, stream); + cudf::detail::sequence(num_rows + 1, zero, scalar_size_per_row, stream, mr); std::unique_ptr data = make_numeric_column(data_type(type_id::INT8), static_cast(total_allocation),