From ee5e89533241f1c221bfffeb1fc964414c12a869 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Tue, 23 Nov 2021 17:13:43 -0500 Subject: [PATCH 1/8] Use CTAD --- cpp/include/cudf_test/column_wrapper.hpp | 2 +- cpp/src/groupby/sort/group_merge_m2.cu | 2 +- cpp/src/groupby/sort/group_single_pass_reduction_util.cuh | 2 +- cpp/src/join/hash_join.cu | 2 +- cpp/src/join/join_utils.cu | 2 +- cpp/src/lists/combine/concatenate_list_elements.cu | 2 +- cpp/src/lists/contains.cu | 2 +- cpp/src/lists/interleave_columns.cu | 6 +++--- cpp/src/reductions/scan/scan_inclusive.cu | 2 +- cpp/src/sort/rank.cu | 8 ++++---- cpp/tests/stream_compaction/apply_boolean_mask_tests.cpp | 4 ++-- cpp/tests/transform/row_bit_count_test.cu | 2 +- 12 files changed, 18 insertions(+), 18 deletions(-) diff --git a/cpp/include/cudf_test/column_wrapper.hpp b/cpp/include/cudf_test/column_wrapper.hpp index f291b04776a..e81a849b8e7 100644 --- a/cpp/include/cudf_test/column_wrapper.hpp +++ b/cpp/include/cudf_test/column_wrapper.hpp @@ -1503,7 +1503,7 @@ class lists_column_wrapper : public detail::column_wrapper { std::cend(cols), valids, // stencil std::back_inserter(children), - thrust::identity{}); + thrust::identity{}); auto data = children.empty() ? cudf::empty_like(expected_hierarchy) : concatenate(children); diff --git a/cpp/src/groupby/sort/group_merge_m2.cu b/cpp/src/groupby/sort/group_merge_m2.cu index 4e2a5b68abc..59b2f6d0989 100644 --- a/cpp/src/groupby/sort/group_merge_m2.cu +++ b/cpp/src/groupby/sort/group_merge_m2.cu @@ -174,7 +174,7 @@ std::unique_ptr group_merge_m2(column_view const& values, // Generate bitmask for the output. // Only mean and M2 values can be nullable. Count column must be non-nullable. auto [null_mask, null_count] = cudf::detail::valid_if( - validities.begin(), validities.end(), thrust::identity{}, stream, mr); + validities.begin(), validities.end(), thrust::identity{}, stream, mr); if (null_count > 0) { result_means->set_null_mask(null_mask, null_count); // copy null_mask result_M2s->set_null_mask(std::move(null_mask), null_count); // take over null_mask diff --git a/cpp/src/groupby/sort/group_single_pass_reduction_util.cuh b/cpp/src/groupby/sort/group_single_pass_reduction_util.cuh index decb127b264..fe24f022a5e 100644 --- a/cpp/src/groupby/sort/group_single_pass_reduction_util.cuh +++ b/cpp/src/groupby/sort/group_single_pass_reduction_util.cuh @@ -218,7 +218,7 @@ struct group_reduction_functor{}); auto [null_mask, null_count] = cudf::detail::valid_if( - validity.begin(), validity.end(), thrust::identity{}, stream, mr); + validity.begin(), validity.end(), thrust::identity{}, stream, mr); result->set_null_mask(std::move(null_mask), null_count); } return result; diff --git a/cpp/src/join/hash_join.cu b/cpp/src/join/hash_join.cu index e4bd1938ecc..c5b680f129e 100644 --- a/cpp/src/join/hash_join.cu +++ b/cpp/src/join/hash_join.cu @@ -266,7 +266,7 @@ std::size_t get_full_join_size(cudf::table_device_view build_table, left_join_complement_size = thrust::count_if(rmm::exec_policy(stream), invalid_index_map->begin(), invalid_index_map->end(), - thrust::identity()); + thrust::identity()); } return join_size + left_join_complement_size; } diff --git a/cpp/src/join/join_utils.cu b/cpp/src/join/join_utils.cu index 4aca4b4a9cf..9e98f87e7f0 100644 --- a/cpp/src/join/join_utils.cu +++ b/cpp/src/join/join_utils.cu @@ -136,7 +136,7 @@ get_left_join_indices_complement(std::unique_ptr> thrust::make_counting_iterator(end_counter), invalid_index_map->begin(), right_indices_complement->begin(), - thrust::identity()) - + thrust::identity{}) - right_indices_complement->begin(); right_indices_complement->resize(indices_count, stream); } diff --git a/cpp/src/lists/combine/concatenate_list_elements.cu b/cpp/src/lists/combine/concatenate_list_elements.cu index 4bef312b396..2ddede97ce4 100644 --- a/cpp/src/lists/combine/concatenate_list_elements.cu +++ b/cpp/src/lists/combine/concatenate_list_elements.cu @@ -225,7 +225,7 @@ std::unique_ptr concatenate_lists_nullifying_rows(column_view const& inp auto list_entries = gather_list_entries(input, offsets_view, num_rows, num_output_entries, stream, mr); auto [null_mask, null_count] = cudf::detail::valid_if( - list_validities.begin(), list_validities.end(), thrust::identity{}, stream, mr); + list_validities.begin(), list_validities.end(), thrust::identity{}, stream, mr); return make_lists_column(num_rows, std::move(list_offsets), diff --git a/cpp/src/lists/contains.cu b/cpp/src/lists/contains.cu index bdbc9ae013c..3d0473a3ad9 100644 --- a/cpp/src/lists/contains.cu +++ b/cpp/src/lists/contains.cu @@ -76,7 +76,7 @@ struct lookup_functor { } else { return cudf::detail::valid_if(result_validity.begin(), result_validity.end(), - thrust::identity{}, + thrust::identity{}, stream, mr); } diff --git a/cpp/src/lists/interleave_columns.cu b/cpp/src/lists/interleave_columns.cu index b9b73d98ed2..7aa504b224a 100644 --- a/cpp/src/lists/interleave_columns.cu +++ b/cpp/src/lists/interleave_columns.cu @@ -229,7 +229,7 @@ struct interleave_list_entries_impl{}, stream, mr); + validities.begin(), validities.end(), thrust::identity{}, stream, mr); return make_strings_column(num_output_entries, std::move(offsets_column), @@ -306,7 +306,7 @@ struct interleave_list_entries_impl( if (data_has_null_mask) { auto [null_mask, null_count] = cudf::detail::valid_if( - validities.begin(), validities.end(), thrust::identity{}, stream, mr); + validities.begin(), validities.end(), thrust::identity{}, stream, mr); if (null_count > 0) { output->set_null_mask(null_mask, null_count); } } @@ -405,7 +405,7 @@ std::unique_ptr interleave_columns(table_view const& input, } auto [null_mask, null_count] = cudf::detail::valid_if( - list_validities.begin(), list_validities.end(), thrust::identity{}, stream, mr); + list_validities.begin(), list_validities.end(), thrust::identity{}, stream, mr); return make_lists_column(num_output_lists, std::move(list_offsets), std::move(list_entries), diff --git a/cpp/src/reductions/scan/scan_inclusive.cu b/cpp/src/reductions/scan/scan_inclusive.cu index 02ecd6df4d9..52788cdeff2 100644 --- a/cpp/src/reductions/scan/scan_inclusive.cu +++ b/cpp/src/reductions/scan/scan_inclusive.cu @@ -49,7 +49,7 @@ rmm::device_buffer mask_scan(column_view const& input_view, size_type const first_null = thrust::find_if_not(rmm::exec_policy(stream), valid_itr, valid_itr + input_view.size(), - thrust::identity{}) - + thrust::identity{}) - valid_itr; size_type const exclusive_offset = (inclusive == scan_type::EXCLUSIVE) ? 1 : 0; return std::min(input_view.size(), first_null + exclusive_offset); diff --git a/cpp/src/sort/rank.cu b/cpp/src/sort/rank.cu index c8a908e44cd..cbc2d0a0730 100644 --- a/cpp/src/sort/rank.cu +++ b/cpp/src/sort/rank.cu @@ -171,8 +171,8 @@ void rank_min(cudf::device_span group_keys, thrust::make_counting_iterator(1), sorted_order_view, rank_mutable_view.begin(), - thrust::minimum{}, - thrust::identity{}, + thrust::minimum{}, + thrust::identity{}, stream); } @@ -189,8 +189,8 @@ void rank_max(cudf::device_span group_keys, thrust::make_counting_iterator(1), sorted_order_view, rank_mutable_view.begin(), - thrust::maximum{}, - thrust::identity{}, + thrust::maximum{}, + thrust::identity{}, stream); } diff --git a/cpp/tests/stream_compaction/apply_boolean_mask_tests.cpp b/cpp/tests/stream_compaction/apply_boolean_mask_tests.cpp index 813cceb0861..c80a8fba55c 100644 --- a/cpp/tests/stream_compaction/apply_boolean_mask_tests.cpp +++ b/cpp/tests/stream_compaction/apply_boolean_mask_tests.cpp @@ -204,13 +204,13 @@ TEST_F(ApplyBooleanMask, FixedPointLargeColumnTest) dec32_data.cend(), mask_data.cbegin(), std::back_inserter(expect_dec32_data), - thrust::identity()); + thrust::identity{}); thrust::copy_if(thrust::seq, dec64_data.cbegin(), dec64_data.cend(), mask_data.cbegin(), std::back_inserter(expect_dec64_data), - thrust::identity()); + thrust::identity{}); decimal32_wrapper expect_col32( expect_dec32_data.begin(), expect_dec32_data.end(), numeric::scale_type{-3}); diff --git a/cpp/tests/transform/row_bit_count_test.cu b/cpp/tests/transform/row_bit_count_test.cu index 7fb7326f221..77edd20ea06 100644 --- a/cpp/tests/transform/row_bit_count_test.cu +++ b/cpp/tests/transform/row_bit_count_test.cu @@ -242,7 +242,7 @@ TEST_F(RowBitCount, StructsWithLists_RowsExceedingASingleBlock) thrust::tabulate(thrust::device, ints_view.begin(), ints_view.end(), - thrust::identity()); + thrust::identity{}); // List offsets = {0, 2, 4, 6, 8, ..., num_rows*2}; auto list_offsets = make_numeric_column(data_type{type_id::INT32}, num_rows + 1); From be0634da4fe024e666204892eb88db21d3d8ae4f Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Tue, 23 Nov 2021 17:33:29 -0500 Subject: [PATCH 2/8] Clang-format --- cpp/tests/transform/row_bit_count_test.cu | 6 ++---- 1 file changed, 2 insertions(+), 4 deletions(-) diff --git a/cpp/tests/transform/row_bit_count_test.cu b/cpp/tests/transform/row_bit_count_test.cu index 77edd20ea06..43d63c9fd22 100644 --- a/cpp/tests/transform/row_bit_count_test.cu +++ b/cpp/tests/transform/row_bit_count_test.cu @@ -239,10 +239,8 @@ TEST_F(RowBitCount, StructsWithLists_RowsExceedingASingleBlock) // List child column = {0, 1, 2, 3, 4, ..., 2*num_rows}; auto ints = make_numeric_column(data_type{type_id::INT32}, num_rows * 2); auto ints_view = ints->mutable_view(); - thrust::tabulate(thrust::device, - ints_view.begin(), - ints_view.end(), - thrust::identity{}); + thrust::tabulate( + thrust::device, ints_view.begin(), ints_view.end(), thrust::identity{}); // List offsets = {0, 2, 4, 6, 8, ..., num_rows*2}; auto list_offsets = make_numeric_column(data_type{type_id::INT32}, num_rows + 1); From 9d664f71c10415e00612b40052357f144af0d04e Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Tue, 23 Nov 2021 17:37:46 -0500 Subject: [PATCH 3/8] Clang-format --- cpp/src/lists/contains.cu | 7 ++----- cpp/src/lists/interleave_columns.cu | 4 ++-- cpp/src/reductions/scan/scan_inclusive.cu | 9 ++++----- cpp/tests/strings/fixed_point_tests.cpp | 2 +- 4 files changed, 9 insertions(+), 13 deletions(-) diff --git a/cpp/src/lists/contains.cu b/cpp/src/lists/contains.cu index 3d0473a3ad9..b48982d205a 100644 --- a/cpp/src/lists/contains.cu +++ b/cpp/src/lists/contains.cu @@ -74,11 +74,8 @@ struct lookup_functor { if (!search_keys_have_nulls && !input_lists.has_nulls() && !input_lists.child().has_nulls()) { return {rmm::device_buffer{0, stream, mr}, size_type{0}}; } else { - return cudf::detail::valid_if(result_validity.begin(), - result_validity.end(), - thrust::identity{}, - stream, - mr); + return cudf::detail::valid_if( + result_validity.begin(), result_validity.end(), thrust::identity{}, stream, mr); } } diff --git a/cpp/src/lists/interleave_columns.cu b/cpp/src/lists/interleave_columns.cu index 7aa504b224a..220cb25a942 100644 --- a/cpp/src/lists/interleave_columns.cu +++ b/cpp/src/lists/interleave_columns.cu @@ -228,8 +228,8 @@ struct interleave_list_entries_impl(c), "1.70141183460469231731687303715884105727"); -} \ No newline at end of file +} From 20e2e8146a988f207396943b7b8a65866636d814 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Tue, 23 Nov 2021 18:14:01 -0500 Subject: [PATCH 4/8] Clang-format --- cpp/tests/transform/row_bit_count_test.cu | 14 +------------- 1 file changed, 1 insertion(+), 13 deletions(-) diff --git a/cpp/tests/transform/row_bit_count_test.cu b/cpp/tests/transform/row_bit_count_test.cu index 43d63c9fd22..52924fb813c 100644 --- a/cpp/tests/transform/row_bit_count_test.cu +++ b/cpp/tests/transform/row_bit_count_test.cu @@ -211,18 +211,6 @@ TEST_F(RowBitCount, StringsWithNulls) CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, *result); } -namespace { - -/** - * @brief __device__ functor to multiply input by 2, defined out of line because __device__ lambdas - * cannot be defined in a TEST_F(). - */ -struct times_2 { - int32_t __device__ operator()(int32_t i) const { return i * 2; } -}; - -} // namespace - TEST_F(RowBitCount, StructsWithLists_RowsExceedingASingleBlock) { // Tests that `row_bit_count()` can handle struct> with more @@ -248,7 +236,7 @@ TEST_F(RowBitCount, StructsWithLists_RowsExceedingASingleBlock) thrust::tabulate(thrust::device, list_offsets_view.begin(), list_offsets_view.end(), - times_2{}); + [] __device__ (auto e) { return e * 2; }); // List = {{0,1}, {2,3}, {4,5}, ..., {2*(num_rows-1), 2*num_rows-1}}; auto lists_column = make_lists_column(num_rows, std::move(list_offsets), std::move(ints), 0, {}); From fbce0d4aaf3eaaf6adc9a10fc450b9aa0b66ce7b Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Tue, 23 Nov 2021 21:05:25 -0500 Subject: [PATCH 5/8] Clang-format --- cpp/src/groupby/sort/group_merge_m2.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/src/groupby/sort/group_merge_m2.cu b/cpp/src/groupby/sort/group_merge_m2.cu index 59b2f6d0989..bde7c985df1 100644 --- a/cpp/src/groupby/sort/group_merge_m2.cu +++ b/cpp/src/groupby/sort/group_merge_m2.cu @@ -173,8 +173,8 @@ std::unique_ptr group_merge_m2(column_view const& values, // Generate bitmask for the output. // Only mean and M2 values can be nullable. Count column must be non-nullable. - auto [null_mask, null_count] = cudf::detail::valid_if( - validities.begin(), validities.end(), thrust::identity{}, stream, mr); + auto [null_mask, null_count] = + cudf::detail::valid_if(validities.begin(), validities.end(), thrust::identity{}, stream, mr); if (null_count > 0) { result_means->set_null_mask(null_mask, null_count); // copy null_mask result_M2s->set_null_mask(std::move(null_mask), null_count); // take over null_mask From b2179a534d7220370cf489d642171f5baa6aa6ba Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Tue, 23 Nov 2021 21:19:18 -0500 Subject: [PATCH 6/8] Revert change --- cpp/tests/transform/row_bit_count_test.cu | 14 +++++++++++++- 1 file changed, 13 insertions(+), 1 deletion(-) diff --git a/cpp/tests/transform/row_bit_count_test.cu b/cpp/tests/transform/row_bit_count_test.cu index 52924fb813c..43d63c9fd22 100644 --- a/cpp/tests/transform/row_bit_count_test.cu +++ b/cpp/tests/transform/row_bit_count_test.cu @@ -211,6 +211,18 @@ TEST_F(RowBitCount, StringsWithNulls) CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, *result); } +namespace { + +/** + * @brief __device__ functor to multiply input by 2, defined out of line because __device__ lambdas + * cannot be defined in a TEST_F(). + */ +struct times_2 { + int32_t __device__ operator()(int32_t i) const { return i * 2; } +}; + +} // namespace + TEST_F(RowBitCount, StructsWithLists_RowsExceedingASingleBlock) { // Tests that `row_bit_count()` can handle struct> with more @@ -236,7 +248,7 @@ TEST_F(RowBitCount, StructsWithLists_RowsExceedingASingleBlock) thrust::tabulate(thrust::device, list_offsets_view.begin(), list_offsets_view.end(), - [] __device__ (auto e) { return e * 2; }); + times_2{}); // List = {{0,1}, {2,3}, {4,5}, ..., {2*(num_rows-1), 2*num_rows-1}}; auto lists_column = make_lists_column(num_rows, std::move(list_offsets), std::move(ints), 0, {}); From dbc8b6db6fa799e6a1f555cf48d5b9d3a3f4a35a Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Tue, 23 Nov 2021 21:22:57 -0500 Subject: [PATCH 7/8] Clang-format --- cpp/src/groupby/sort/group_single_pass_reduction_util.cuh | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/src/groupby/sort/group_single_pass_reduction_util.cuh b/cpp/src/groupby/sort/group_single_pass_reduction_util.cuh index fe24f022a5e..c41fab144a8 100644 --- a/cpp/src/groupby/sort/group_single_pass_reduction_util.cuh +++ b/cpp/src/groupby/sort/group_single_pass_reduction_util.cuh @@ -217,8 +217,8 @@ struct group_reduction_functor{}); - auto [null_mask, null_count] = cudf::detail::valid_if( - validity.begin(), validity.end(), thrust::identity{}, stream, mr); + auto [null_mask, null_count] = + cudf::detail::valid_if(validity.begin(), validity.end(), thrust::identity{}, stream, mr); result->set_null_mask(std::move(null_mask), null_count); } return result; From 7239a6df59009750c2308e33b76f6fa810ee0205 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Tue, 23 Nov 2021 23:04:04 -0500 Subject: [PATCH 8/8] More CTAD --- cpp/include/cudf/strings/detail/gather.cuh | 2 +- cpp/include/cudf_test/column_wrapper.hpp | 7 ++----- cpp/src/copying/concatenate.cu | 2 +- cpp/src/groupby/sort/group_rank_scan.cu | 2 +- cpp/src/groupby/sort/group_scan_util.cuh | 6 +++--- .../sort/group_single_pass_reduction_util.cuh | 12 ++++++------ cpp/src/groupby/sort/group_tdigest.cu | 10 +++++----- cpp/src/quantiles/tdigest/tdigest.cu | 7 ++----- cpp/src/rolling/grouped_rolling.cu | 6 +++--- cpp/src/rolling/rolling_collect_list.cu | 2 +- cpp/src/sort/rank.cu | 2 +- cpp/src/strings/copying/concatenate.cu | 2 +- cpp/src/strings/findall.cu | 7 ++----- cpp/src/strings/repeat_strings.cu | 2 +- cpp/src/strings/split/split.cu | 14 ++++---------- cpp/tests/iterator/iterator_tests.cuh | 11 +++-------- 16 files changed, 37 insertions(+), 57 deletions(-) diff --git a/cpp/include/cudf/strings/detail/gather.cuh b/cpp/include/cudf/strings/detail/gather.cuh index ec4a88a0e46..eb7258830ce 100644 --- a/cpp/include/cudf/strings/detail/gather.cuh +++ b/cpp/include/cudf/strings/detail/gather.cuh @@ -315,7 +315,7 @@ std::unique_ptr gather( d_out_offsets + output_count, [] __device__(auto size) { return static_cast(size); }, size_t{0}, - thrust::plus{}); + thrust::plus{}); CUDF_EXPECTS(total_bytes < static_cast(std::numeric_limits::max()), "total size of output strings is too large for a cudf column"); diff --git a/cpp/include/cudf_test/column_wrapper.hpp b/cpp/include/cudf_test/column_wrapper.hpp index e81a849b8e7..4d0c68f9ecd 100644 --- a/cpp/include/cudf_test/column_wrapper.hpp +++ b/cpp/include/cudf_test/column_wrapper.hpp @@ -1499,11 +1499,8 @@ class lists_column_wrapper : public detail::column_wrapper { // concatenate them together, skipping children that are null. std::vector children; - thrust::copy_if(std::cbegin(cols), - std::cend(cols), - valids, // stencil - std::back_inserter(children), - thrust::identity{}); + thrust::copy_if( + std::cbegin(cols), std::cend(cols), valids, std::back_inserter(children), thrust::identity{}); auto data = children.empty() ? cudf::empty_like(expected_hierarchy) : concatenate(children); diff --git a/cpp/src/copying/concatenate.cu b/cpp/src/copying/concatenate.cu index f4b6a8bf5fd..34c0cea683e 100644 --- a/cpp/src/copying/concatenate.cu +++ b/cpp/src/copying/concatenate.cu @@ -79,7 +79,7 @@ auto create_device_views(host_span views, rmm::cuda_stream_vi device_views.cend(), std::next(offsets.begin()), [](auto const& col) { return col.size(); }, - thrust::plus{}); + thrust::plus{}); auto d_offsets = make_device_uvector_async(offsets, stream); auto const output_size = offsets.back(); diff --git a/cpp/src/groupby/sort/group_rank_scan.cu b/cpp/src/groupby/sort/group_rank_scan.cu index 935ef9554a9..f36bdc0a660 100644 --- a/cpp/src/groupby/sort/group_rank_scan.cu +++ b/cpp/src/groupby/sort/group_rank_scan.cu @@ -79,7 +79,7 @@ std::unique_ptr rank_generator(column_view const& order_by, group_labels.end(), mutable_ranks.begin(), mutable_ranks.begin(), - thrust::equal_to{}, + thrust::equal_to{}, scan_op); return ranks; diff --git a/cpp/src/groupby/sort/group_scan_util.cuh b/cpp/src/groupby/sort/group_scan_util.cuh index b565e8dc6d8..15afe158af1 100644 --- a/cpp/src/groupby/sort/group_scan_util.cuh +++ b/cpp/src/groupby/sort/group_scan_util.cuh @@ -115,7 +115,7 @@ struct group_scan_functor() group_labels.end(), inp_iter, out_iter, - thrust::equal_to{}, + thrust::equal_to{}, binop); }; @@ -160,7 +160,7 @@ struct group_scan_functor{}, + thrust::equal_to{}, binop); }; @@ -214,7 +214,7 @@ struct group_scan_functor{}, + thrust::equal_to{}, binop); }; diff --git a/cpp/src/groupby/sort/group_single_pass_reduction_util.cuh b/cpp/src/groupby/sort/group_single_pass_reduction_util.cuh index c41fab144a8..95a36f40e57 100644 --- a/cpp/src/groupby/sort/group_single_pass_reduction_util.cuh +++ b/cpp/src/groupby/sort/group_single_pass_reduction_util.cuh @@ -191,7 +191,7 @@ struct group_reduction_functor{}, + thrust::equal_to{}, binop); }; @@ -215,7 +215,7 @@ struct group_reduction_functor validity(num_groups, stream); do_reduction(cudf::detail::make_validity_iterator(*d_values_ptr), validity.begin(), - thrust::logical_or{}); + thrust::logical_or{}); auto [null_mask, null_count] = cudf::detail::valid_if(validity.begin(), validity.end(), thrust::identity{}, stream, mr); @@ -264,7 +264,7 @@ struct group_reduction_functor< inp_iter, thrust::make_discard_iterator(), out_iter, - thrust::equal_to{}, + thrust::equal_to{}, binop); }; @@ -283,10 +283,10 @@ struct group_reduction_functor< auto validity = rmm::device_uvector(num_groups, stream); do_reduction(cudf::detail::make_validity_iterator(*d_values_ptr), validity.begin(), - thrust::logical_or{}); + thrust::logical_or{}); - auto [null_mask, null_count] = cudf::detail::valid_if( - validity.begin(), validity.end(), thrust::identity{}, stream, mr); + auto [null_mask, null_count] = + cudf::detail::valid_if(validity.begin(), validity.end(), thrust::identity{}, stream, mr); result->set_null_mask(std::move(null_mask), null_count); } else { auto const binop = diff --git a/cpp/src/groupby/sort/group_tdigest.cu b/cpp/src/groupby/sort/group_tdigest.cu index 146a6a8c31c..551eb128231 100644 --- a/cpp/src/groupby/sort/group_tdigest.cu +++ b/cpp/src/groupby/sort/group_tdigest.cu @@ -625,7 +625,7 @@ std::unique_ptr compute_tdigests(int delta, centroids_begin, // values thrust::make_discard_iterator(), // key output output, // output - thrust::equal_to{}, // key equality check + thrust::equal_to{}, // key equality check merge_centroids{}); // create final tdigest column @@ -850,8 +850,8 @@ std::unique_ptr group_merge_tdigest(column_view const& input, min_iter, thrust::make_discard_iterator(), merged_min_col->mutable_view().begin(), - thrust::equal_to{}, // key equality check - thrust::minimum{}); + thrust::equal_to{}, // key equality check + thrust::minimum{}); auto merged_max_col = cudf::make_numeric_column( data_type{type_id::FLOAT64}, num_groups, mask_state::UNALLOCATED, stream, mr); @@ -864,8 +864,8 @@ std::unique_ptr group_merge_tdigest(column_view const& input, max_iter, thrust::make_discard_iterator(), merged_max_col->mutable_view().begin(), - thrust::equal_to{}, // key equality check - thrust::maximum{}); + thrust::equal_to{}, // key equality check + thrust::maximum{}); // for any empty groups, set the min and max to be 0. not technically necessary but it makes // testing simpler. diff --git a/cpp/src/quantiles/tdigest/tdigest.cu b/cpp/src/quantiles/tdigest/tdigest.cu index 57c221b15ed..18e7d02d086 100644 --- a/cpp/src/quantiles/tdigest/tdigest.cu +++ b/cpp/src/quantiles/tdigest/tdigest.cu @@ -348,11 +348,8 @@ std::unique_ptr percentile_approx(tdigest_column_view const& input, if (null_count == 0) { return std::pair{rmm::device_buffer{}, null_count}; } - return cudf::detail::valid_if(tdigest_is_empty, - tdigest_is_empty + tdv.size(), - thrust::logical_not{}, - stream, - mr); + return cudf::detail::valid_if( + tdigest_is_empty, tdigest_is_empty + tdv.size(), thrust::logical_not{}, stream, mr); }(); return cudf::make_lists_column( diff --git a/cpp/src/rolling/grouped_rolling.cu b/cpp/src/rolling/grouped_rolling.cu index 509f67bb5c6..5a7f15148d8 100644 --- a/cpp/src/rolling/grouped_rolling.cu +++ b/cpp/src/rolling/grouped_rolling.cu @@ -142,8 +142,8 @@ std::unique_ptr grouped_rolling_window(table_view const& group_keys, preceding_window] __device__(size_type idx) { auto group_label = d_group_labels[idx]; auto group_start = d_group_offsets[group_label]; - return thrust::minimum{}(preceding_window, - idx - group_start + 1); // Preceding includes current row. + return thrust::minimum{}(preceding_window, + idx - group_start + 1); // Preceding includes current row. }; auto following_calculator = [d_group_offsets = group_offsets.data(), @@ -152,7 +152,7 @@ std::unique_ptr grouped_rolling_window(table_view const& group_keys, auto group_label = d_group_labels[idx]; auto group_end = d_group_offsets[group_label + 1]; // Cannot fall off the end, since offsets // is capped with `input.size()`. - return thrust::minimum{}(following_window, (group_end - 1) - idx); + return thrust::minimum{}(following_window, (group_end - 1) - idx); }; if (aggr.kind == aggregation::CUDA || aggr.kind == aggregation::PTX) { diff --git a/cpp/src/rolling/rolling_collect_list.cu b/cpp/src/rolling/rolling_collect_list.cu index ecef90dc8e1..30c39bde7d2 100644 --- a/cpp/src/rolling/rolling_collect_list.cu +++ b/cpp/src/rolling/rolling_collect_list.cu @@ -75,7 +75,7 @@ std::unique_ptr get_list_child_to_list_row_mapping(cudf::column_view con per_row_mapping_begin, per_row_mapping_begin + num_child_rows, per_row_mapping_begin, - thrust::maximum{}); + thrust::maximum{}); return per_row_mapping; } diff --git a/cpp/src/sort/rank.cu b/cpp/src/sort/rank.cu index cbc2d0a0730..e9589e6c4b3 100644 --- a/cpp/src/sort/rank.cu +++ b/cpp/src/sort/rank.cu @@ -117,7 +117,7 @@ void tie_break_ranks_transform(cudf::device_span dense_rank_sor tie_iter, thrust::make_discard_iterator(), tie_sorted.begin(), - thrust::equal_to{}, + thrust::equal_to{}, tie_breaker); auto sorted_tied_rank = thrust::make_transform_iterator( dense_rank_sorted.begin(), diff --git a/cpp/src/strings/copying/concatenate.cu b/cpp/src/strings/copying/concatenate.cu index db8b37a9592..3822fa8bf5a 100644 --- a/cpp/src/strings/copying/concatenate.cu +++ b/cpp/src/strings/copying/concatenate.cu @@ -96,7 +96,7 @@ auto create_strings_device_views(host_span views, rmm::cuda_s device_views_ptr + views.size(), std::next(d_partition_offsets.begin()), chars_size_transform{}, - thrust::plus{}); + thrust::plus{}); auto const output_chars_size = d_partition_offsets.back_element(stream); stream.synchronize(); // ensure copy of output_chars_size is complete before returning diff --git a/cpp/src/strings/findall.cu b/cpp/src/strings/findall.cu index 3ab5b55020c..8d96f0de415 100644 --- a/cpp/src/strings/findall.cu +++ b/cpp/src/strings/findall.cu @@ -153,11 +153,8 @@ std::unique_ptr findall_re( std::vector> results; - size_type const columns = thrust::reduce(rmm::exec_policy(stream), - find_counts.begin(), - find_counts.end(), - 0, - thrust::maximum{}); + size_type const columns = thrust::reduce( + rmm::exec_policy(stream), find_counts.begin(), find_counts.end(), 0, thrust::maximum{}); // boundary case: if no columns, return all nulls column (issue #119) if (columns == 0) results.emplace_back(std::make_unique( diff --git a/cpp/src/strings/repeat_strings.cu b/cpp/src/strings/repeat_strings.cu index 458f3ed885c..7820e0064a6 100644 --- a/cpp/src/strings/repeat_strings.cu +++ b/cpp/src/strings/repeat_strings.cu @@ -369,7 +369,7 @@ std::pair, int64_t> repeat_strings_output_sizes( thrust::make_counting_iterator(strings_count), fn, int64_t{0}, - thrust::plus{}); + thrust::plus{}); return std::make_pair(std::move(output_sizes), total_bytes); } diff --git a/cpp/src/strings/split/split.cu b/cpp/src/strings/split/split.cu index 5113b418501..c6e52a79059 100644 --- a/cpp/src/strings/split/split.cu +++ b/cpp/src/strings/split/split.cu @@ -490,11 +490,8 @@ std::unique_ptr
split_fn(strings_column_view const& strings_column, }); // the columns_count is the maximum number of tokens for any string - auto const columns_count = thrust::reduce(rmm::exec_policy(stream), - token_counts.begin(), - token_counts.end(), - 0, - thrust::maximum{}); + auto const columns_count = thrust::reduce( + rmm::exec_policy(stream), token_counts.begin(), token_counts.end(), 0, thrust::maximum{}); // boundary case: if no columns, return one null column (custrings issue #119) if (columns_count == 0) { results.push_back(std::make_unique( @@ -748,11 +745,8 @@ std::unique_ptr
whitespace_split_fn(size_type strings_count, [tokenizer] __device__(size_type idx) { return tokenizer.count_tokens(idx); }); // column count is the maximum number of tokens for any string - size_type const columns_count = thrust::reduce(rmm::exec_policy(stream), - token_counts.begin(), - token_counts.end(), - 0, - thrust::maximum{}); + size_type const columns_count = thrust::reduce( + rmm::exec_policy(stream), token_counts.begin(), token_counts.end(), 0, thrust::maximum{}); std::vector> results; // boundary case: if no columns, return one null column (issue #119) diff --git a/cpp/tests/iterator/iterator_tests.cuh b/cpp/tests/iterator/iterator_tests.cuh index 4ec347c4bc1..0cf627407e9 100644 --- a/cpp/tests/iterator/iterator_tests.cuh +++ b/cpp/tests/iterator/iterator_tests.cuh @@ -50,13 +50,8 @@ struct IteratorTest : public cudf::test::BaseFixture { // Get temporary storage size size_t temp_storage_bytes = 0; - cub::DeviceReduce::Reduce(nullptr, - temp_storage_bytes, - d_in, - dev_result.begin(), - num_items, - thrust::minimum{}, - init); + cub::DeviceReduce::Reduce( + nullptr, temp_storage_bytes, d_in, dev_result.begin(), num_items, thrust::minimum{}, init); // Allocate temporary storage rmm::device_buffer d_temp_storage(temp_storage_bytes, rmm::cuda_stream_default); @@ -67,7 +62,7 @@ struct IteratorTest : public cudf::test::BaseFixture { d_in, dev_result.begin(), num_items, - thrust::minimum{}, + thrust::minimum{}, init); evaluate(expected, dev_result, "cub test");