From 068542a498455f94571ed4b853e0f3e0d4a50e1c Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Wed, 16 Mar 2022 15:30:29 -0700 Subject: [PATCH 01/20] Passes compile, initial layout for how test should look like --- .../reductions/segmented_reduction_tests.cpp | 59 ++++++++++++++++++- 1 file changed, 58 insertions(+), 1 deletion(-) diff --git a/cpp/tests/reductions/segmented_reduction_tests.cpp b/cpp/tests/reductions/segmented_reduction_tests.cpp index 3a432cce801..de8c3256abb 100644 --- a/cpp/tests/reductions/segmented_reduction_tests.cpp +++ b/cpp/tests/reductions/segmented_reduction_tests.cpp @@ -19,9 +19,11 @@ #include #include +#include #include #include +#include #include namespace cudf { @@ -36,7 +38,7 @@ struct SegmentedReductionTest : public cudf::test::BaseFixture { struct SegmentedReductionTestUntyped : public cudf::test::BaseFixture { }; -TYPED_TEST_CASE(SegmentedReductionTest, NumericTypes); +TYPED_TEST_SUITE(SegmentedReductionTest, NumericTypes); TYPED_TEST(SegmentedReductionTest, SumExcludeNulls) { @@ -385,6 +387,61 @@ TEST_F(SegmentedReductionTestUntyped, ReduceEmptyColumn) CUDF_TEST_EXPECT_COLUMNS_EQUAL(*res, expect); } +int32_t pow10(int32_t exponent) { return exponent == 0 ? 1 : 10 * pow10(exponent - 1); } + +template +struct SegmentedReductionFixedPointTest : public cudf::test::BaseFixture { + public: + std::vector scale_list_by_pow10(std::vector input, + int32_t exponent) + { + std::vector result(input.size()); + std::transform(input.begin(), input.end(), result.begin(), [&exponent](auto x) { + return x * pow10(exponent); + }); + return result; + } +}; + +TYPED_TEST_SUITE(SegmentedReductionFixedPointTest, cudf::test::FixedPointTypes); + +TYPED_TEST(SegmentedReductionFixedPointTest, ProductIncludeNullsZeroInputScale) +{ + // [1, 2, 3], [1], [], [2, NULL, 3], [NULL], [NULL, NULL] | scale: 0 + // values: {1, 2, 3, 1, 2, XXX, 3, XXX, XXX, XXX} + // offsets: {0, 3, 4, 4, 7, 8, 10} + // nullmask:{1, 1, 1, 1, 1, 0, 1, 0, 0, 0} + // output_dtype: decimalXX, scale: -1, 0, 1 + // outputs: {6, 1, XXX, XXX, XXX, XXX} + // output nullmask: {1, 1, 0, 0, 0, 0} + + using DecimalXX = TypeParam; + + for (int output_scale : {-1, 0, 1}) { + fixed_point_column_wrapper input{ + {1, 2, 3, 1, 2, XXX, 3, XXX, XXX, XXX}, + {true, true, true, true, true, false, true, false, false, false}, + numeric::scale_type(0)}; + fixed_width_column_wrapper offsets{0, 3, 4, 4, 7, 8, 10}; + + data_type output_dtype{type_to_id(), numeric::scale_type{output_scale}}; + + auto result_rep = this->scale_list_by_pow10({6, 1, XXX, XXX, XXX, XXX}, -output_scale); + fixed_point_column_wrapper expect{ + result_rep.begin(), + result_rep.end(), + {true, true, false, false, false, false}, + numeric::scale_type(output_scale)}; + + auto res = segmented_reduce(input, + column_view(offsets), + *make_product_aggregation(), + output_dtype, + null_policy::INCLUDE); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*res, expect); + } +} + #undef XXX } // namespace test From 66bdb752397fb13560462771d5a39b16ee417d58 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Wed, 16 Mar 2022 15:33:01 -0700 Subject: [PATCH 02/20] remove unused include --- cpp/tests/reductions/segmented_reduction_tests.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/cpp/tests/reductions/segmented_reduction_tests.cpp b/cpp/tests/reductions/segmented_reduction_tests.cpp index de8c3256abb..613ca9b4ab5 100644 --- a/cpp/tests/reductions/segmented_reduction_tests.cpp +++ b/cpp/tests/reductions/segmented_reduction_tests.cpp @@ -23,7 +23,6 @@ #include #include -#include #include namespace cudf { From 9bdcc6591e19de5a6ae179f0e283110b512336b0 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Wed, 16 Mar 2022 15:40:35 -0700 Subject: [PATCH 03/20] Fix exponent bug --- cpp/tests/reductions/segmented_reduction_tests.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/tests/reductions/segmented_reduction_tests.cpp b/cpp/tests/reductions/segmented_reduction_tests.cpp index 613ca9b4ab5..59598d63641 100644 --- a/cpp/tests/reductions/segmented_reduction_tests.cpp +++ b/cpp/tests/reductions/segmented_reduction_tests.cpp @@ -396,7 +396,7 @@ struct SegmentedReductionFixedPointTest : public cudf::test::BaseFixture { { std::vector result(input.size()); std::transform(input.begin(), input.end(), result.begin(), [&exponent](auto x) { - return x * pow10(exponent); + return exponent >= 0 ? x * pow10(exponent) : x / pow10(-exponent); }); return result; } @@ -404,7 +404,7 @@ struct SegmentedReductionFixedPointTest : public cudf::test::BaseFixture { TYPED_TEST_SUITE(SegmentedReductionFixedPointTest, cudf::test::FixedPointTypes); -TYPED_TEST(SegmentedReductionFixedPointTest, ProductIncludeNullsZeroInputScale) +TYPED_TEST(SegmentedReductionFixedPointTest, ProductIncludeNulls) { // [1, 2, 3], [1], [], [2, NULL, 3], [NULL], [NULL, NULL] | scale: 0 // values: {1, 2, 3, 1, 2, XXX, 3, XXX, XXX, XXX} From 3ddad323dc3fe5a752b8e5fcccf5c01170819cd6 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 17 Mar 2022 16:44:43 -0700 Subject: [PATCH 04/20] Fix to test for max aggregation and add string type test --- .../reductions/segmented_reduction_tests.cpp | 42 ++++++++++++++++--- 1 file changed, 37 insertions(+), 5 deletions(-) diff --git a/cpp/tests/reductions/segmented_reduction_tests.cpp b/cpp/tests/reductions/segmented_reduction_tests.cpp index 59598d63641..01d9474ac19 100644 --- a/cpp/tests/reductions/segmented_reduction_tests.cpp +++ b/cpp/tests/reductions/segmented_reduction_tests.cpp @@ -14,6 +14,7 @@ * limitations under the License. */ +#include "cudf_test/column_utilities.hpp" #include #include #include @@ -404,14 +405,14 @@ struct SegmentedReductionFixedPointTest : public cudf::test::BaseFixture { TYPED_TEST_SUITE(SegmentedReductionFixedPointTest, cudf::test::FixedPointTypes); -TYPED_TEST(SegmentedReductionFixedPointTest, ProductIncludeNulls) +TYPED_TEST(SegmentedReductionFixedPointTest, MaxIncludeNullsScaleZero) { // [1, 2, 3], [1], [], [2, NULL, 3], [NULL], [NULL, NULL] | scale: 0 // values: {1, 2, 3, 1, 2, XXX, 3, XXX, XXX, XXX} - // offsets: {0, 3, 4, 4, 7, 8, 10} // nullmask:{1, 1, 1, 1, 1, 0, 1, 0, 0, 0} + // offsets: {0, 3, 4, 4, 7, 8, 10} // output_dtype: decimalXX, scale: -1, 0, 1 - // outputs: {6, 1, XXX, XXX, XXX, XXX} + // outputs: {3, 1, XXX, XXX, XXX, XXX} // output nullmask: {1, 1, 0, 0, 0, 0} using DecimalXX = TypeParam; @@ -425,7 +426,7 @@ TYPED_TEST(SegmentedReductionFixedPointTest, ProductIncludeNulls) data_type output_dtype{type_to_id(), numeric::scale_type{output_scale}}; - auto result_rep = this->scale_list_by_pow10({6, 1, XXX, XXX, XXX, XXX}, -output_scale); + auto result_rep = this->scale_list_by_pow10({3, 1, XXX, XXX, XXX, XXX}, -output_scale); fixed_point_column_wrapper expect{ result_rep.begin(), result_rep.end(), @@ -434,13 +435,44 @@ TYPED_TEST(SegmentedReductionFixedPointTest, ProductIncludeNulls) auto res = segmented_reduce(input, column_view(offsets), - *make_product_aggregation(), + *make_max_aggregation(), output_dtype, null_policy::INCLUDE); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*res, expect); } } +struct SegmentedReductionStringTest : public cudf::test::BaseFixture { +}; + +TEST_F(SegmentedReductionStringTest, MaxIncludeNulls) +{ + // ['world'], ['cudf', NULL, 'cuml'], ['hello', 'rapids', 'ai'], [], [NULL], [NULL, NULL] + // values: {"world", "cudf", XXX, "cuml", "hello", "rapids", "ai", XXX, XXX, XXX} + // nullmask:{1, 1, 0, 1, 1, 1, 1, 0, 0, 0} + // offsets: {0, 1, 4, 7, 7, 8, 10} + // output_dtype: string dtype + // outputs: {"world", XXX, "rapids", XXX, XXX, XXX} + // output nullmask: {1, 0, 1, 0, 0, 0} + + strings_column_wrapper input{ + {"world", "cudf", XXX, "cuml", "hello", "rapids", "ai", XXX, XXX, XXX}, + {true, true, false, true, true, true, true, false, false, false}}; + fixed_width_column_wrapper offsets{0, 1, 4, 7, 7, 8, 10}; + data_type output_dtype{type_id::STRING}; + + strings_column_wrapper expect{{"world", XXX, "rapids", XXX, XXX, XXX}, + {true, false, true, false, false, false}}; + + auto res = segmented_reduce(input, + column_view(offsets), + *make_max_aggregation(), + output_dtype, + null_policy::INCLUDE); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*res, expect); +} +} + #undef XXX } // namespace test From e78ee186d257dbb04c0f82f0ac384814acb195d8 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 24 Mar 2022 20:54:33 -0700 Subject: [PATCH 05/20] Remove excess brace --- cpp/tests/reductions/segmented_reduction_tests.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/cpp/tests/reductions/segmented_reduction_tests.cpp b/cpp/tests/reductions/segmented_reduction_tests.cpp index 01d9474ac19..4b4b98804b8 100644 --- a/cpp/tests/reductions/segmented_reduction_tests.cpp +++ b/cpp/tests/reductions/segmented_reduction_tests.cpp @@ -471,7 +471,6 @@ TEST_F(SegmentedReductionStringTest, MaxIncludeNulls) null_policy::INCLUDE); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*res, expect); } -} #undef XXX From eef7f2141e035703a312f83309a05a8c7b8e6ef2 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Mon, 4 Apr 2022 15:26:27 -0700 Subject: [PATCH 06/20] Refactor `element_minmax_fn` to a shared utility file Co-authored-by: Bradley Dice --- .../detail/utilities/device_operators.cuh | 32 +++++++++++++++++++ .../sort/group_single_pass_reduction_util.cuh | 32 +------------------ 2 files changed, 33 insertions(+), 31 deletions(-) diff --git a/cpp/include/cudf/detail/utilities/device_operators.cuh b/cpp/include/cudf/detail/utilities/device_operators.cuh index 87fef5bc187..9f75f85bcf3 100644 --- a/cpp/include/cudf/detail/utilities/device_operators.cuh +++ b/cpp/include/cudf/detail/utilities/device_operators.cuh @@ -21,6 +21,7 @@ * @file */ +#include #include #include #include @@ -262,4 +263,35 @@ struct DeviceLeadLag { explicit CUDF_HOST_DEVICE inline DeviceLeadLag(size_type offset_) : row_offset(offset_) {} }; +/** + * @brief Binary `argmin`/`argmax` operator + * + * @tparam T Type of the underlying column. Must support '<' operator. + */ +template +struct element_arg_minmax_fn { + column_device_view const d_col; + bool const has_nulls; + bool const arg_min; + + __device__ inline auto operator()(size_type const& lhs_idx, size_type const& rhs_idx) const + { + // The extra bounds checking is due to issue github.com/rapidsai/cudf/9156 and + // github.com/NVIDIA/thrust/issues/1525 + // where invalid random values may be passed here by thrust::reduce_by_key + if (lhs_idx < 0 || lhs_idx >= d_col.size() || (has_nulls && d_col.is_null_nocheck(lhs_idx))) { + return rhs_idx; + } + if (rhs_idx < 0 || rhs_idx >= d_col.size() || (has_nulls && d_col.is_null_nocheck(rhs_idx))) { + return lhs_idx; + } + + // Return `lhs_idx` iff: + // row(lhs_idx) < row(rhs_idx) and finding ArgMin, or + // row(lhs_idx) >= row(rhs_idx) and finding ArgMax. + auto const less = d_col.element(lhs_idx) < d_col.element(rhs_idx); + return less == arg_min ? lhs_idx : rhs_idx; + } +}; + } // namespace cudf 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 8e1463f7964..ee99b20f12d 100644 --- a/cpp/src/groupby/sort/group_single_pass_reduction_util.cuh +++ b/cpp/src/groupby/sort/group_single_pass_reduction_util.cuh @@ -23,6 +23,7 @@ #include #include #include +#include #include #include #include @@ -40,37 +41,6 @@ namespace cudf { namespace groupby { namespace detail { -/** - * @brief Binary operator with index values into the input column. - * - * @tparam T Type of the underlying column. Must support '<' operator. - */ -template -struct element_arg_minmax_fn { - column_device_view const d_col; - bool const has_nulls; - bool const arg_min; - - __device__ inline auto operator()(size_type const& lhs_idx, size_type const& rhs_idx) const - { - // The extra bounds checking is due to issue github.com/rapidsai/cudf/9156 and - // github.com/NVIDIA/thrust/issues/1525 - // where invalid random values may be passed here by thrust::reduce_by_key - if (lhs_idx < 0 || lhs_idx >= d_col.size() || (has_nulls && d_col.is_null_nocheck(lhs_idx))) { - return rhs_idx; - } - if (rhs_idx < 0 || rhs_idx >= d_col.size() || (has_nulls && d_col.is_null_nocheck(rhs_idx))) { - return lhs_idx; - } - - // Return `lhs_idx` iff: - // row(lhs_idx) < row(rhs_idx) and finding ArgMin, or - // row(lhs_idx) >= row(rhs_idx) and finding ArgMax. - auto const less = d_col.element(lhs_idx) < d_col.element(rhs_idx); - return less == arg_min ? lhs_idx : rhs_idx; - } -}; - /** * @brief Value accessor for column which supports dictionary column too. * From a8dc8ea485fe6be4cba591fffa12e2bd2a44dd14 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Mon, 4 Apr 2022 15:28:38 -0700 Subject: [PATCH 07/20] Initial pass on string type min, max support --- cpp/include/cudf/detail/reduction.cuh | 78 ++++++++++++--- cpp/src/reductions/simple_segmented.cuh | 120 +++++++++++++++++++++++- 2 files changed, 182 insertions(+), 16 deletions(-) diff --git a/cpp/include/cudf/detail/reduction.cuh b/cpp/include/cudf/detail/reduction.cuh index 76afbf7e4b8..eb8668ddc3f 100644 --- a/cpp/include/cudf/detail/reduction.cuh +++ b/cpp/include/cudf/detail/reduction.cuh @@ -227,36 +227,36 @@ std::unique_ptr reduce(InputIterator d_in, * @brief Compute the specified simple reduction over each of the segments in the * input range of elements. * - * @tparam Op the reduction operator with device binary operator * @tparam InputIterator the input column iterator * @tparam OffsetIterator the offset column iterator + * @tparam BinaryOp the device binary operator used to reduce * @tparam OutputType the output type of reduction * * @param[in] d_in the begin iterator to input * @param[in] d_offset the begin iterator to offset * @param[in] num_segments the number of segments - * @param[in] sop the reduction operator + * @param[in] binary_op the reduction operator + * @param[in] identity the identity element of the reduction operator * @param[in] stream CUDA stream used for device memory operations and kernel launches. * @param[in] mr Device memory resource used to allocate the returned column's device * memory * @returns Output column in device memory * */ -template ::type, typename std::enable_if_t() && - not cudf::is_fixed_point()>* = nullptr> + !cudf::is_fixed_point()>* = nullptr> std::unique_ptr segmented_reduce(InputIterator d_in, OffsetIterator d_offset, cudf::size_type num_segments, - op::simple_op sop, + BinaryOp binary_op, + OutputType identity, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - auto binary_op = sop.get_binary_op(); - auto identity = sop.template get_identity(); auto dev_result = make_fixed_width_column( data_type{type_to_id()}, num_segments, mask_state::UNALLOCATED, stream, mr); auto dev_result_mview = dev_result->mutable_view(); @@ -291,16 +291,70 @@ std::unique_ptr segmented_reduce(InputIterator d_in, return dev_result; } -template ::type, + typename std::enable_if_t>* = nullptr> +std::unique_ptr segmented_reduce(InputIterator d_in, + OffsetIterator d_offset, + cudf::size_type num_segments, + BinaryOp binary_op, + OutputType identity, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + auto gather_map = make_fixed_width_column( + data_type{type_to_id()}, num_segments, mask_state::UNALLOCATED, stream, mr); + auto dev_result_mview = gather_map->mutable_view(); + + // Allocate temporary storage + rmm::device_buffer d_temp_storage; + size_t temp_storage_bytes = 0; + cub::DeviceSegmentedReduce::Reduce(d_temp_storage.data(), + temp_storage_bytes, + d_in, + dev_result_mview.data(), + num_segments, + d_offset, + d_offset + 1, + binary_op, + identity, + stream.value()); + d_temp_storage = rmm::device_buffer{temp_storage_bytes, stream}; + + // Run reduction + cub::DeviceSegmentedReduce::Reduce(d_temp_storage.data(), + temp_storage_bytes, + d_in, + dev_result_mview.data(), + num_segments, + d_offset, + d_offset + 1, + binary_op, + identity, + stream.value()); + + return gather_map; +} + +template ::type, - typename std::enable_if_t() || + typename std::enable_if_t<(!is_fixed_width() && + !std::is_same_v()) || is_fixed_point()>* = nullptr> std::unique_ptr segmented_reduce(InputIterator, OffsetIterator, cudf::size_type, - op::simple_op, + BinaryOp, + OutputType, rmm::cuda_stream_view, rmm::mr::device_memory_resource*) { diff --git a/cpp/src/reductions/simple_segmented.cuh b/cpp/src/reductions/simple_segmented.cuh index 99837e67398..a19412afc04 100644 --- a/cpp/src/reductions/simple_segmented.cuh +++ b/cpp/src/reductions/simple_segmented.cuh @@ -16,12 +16,15 @@ #pragma once +#include #include +#include #include #include #include #include #include +#include #include #include #include @@ -31,9 +34,12 @@ #include +#include #include #include +#include + namespace cudf { namespace reduction { namespace simple { @@ -70,18 +76,21 @@ std::unique_ptr simple_segmented_reduction(column_view const& col, auto simple_op = Op{}; size_type num_segments = offsets.size() - 1; + auto binary_op = simple_op.get_binary_op(); + auto identity = simple_op.template get_identity(); + // TODO: Explore rewriting null_replacing_element_transformer/element_transformer with nullate auto result = [&] { if (col.has_nulls()) { auto f = simple_op.template get_null_replacing_element_transformer(); auto it = thrust::make_transform_iterator(dcol->pair_begin(), f); return cudf::reduction::detail::segmented_reduce( - it, offsets.begin(), num_segments, simple_op, stream, mr); + it, offsets.begin(), num_segments, binary_op, identity, stream, mr); } else { auto f = simple_op.template get_element_transformer(); auto it = thrust::make_transform_iterator(dcol->begin(), f); return cudf::reduction::detail::segmented_reduce( - it, offsets.begin(), num_segments, simple_op, stream, mr); + it, offsets.begin(), num_segments, binary_op, identity, stream, mr); } }(); @@ -103,6 +112,95 @@ std::unique_ptr simple_segmented_reduction(column_view const& col, return result; } +template || + std::is_same_v)> +std::unique_ptr string_simple_segmented_reduction(column_view const& col, + device_span offsets, + null_policy null_handling, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + // Pass to simple_segmented_reduction, get indices to gather, perform gather here. + auto device_col = cudf::column_device_view::create(col, stream); + + auto it = thrust::make_counting_iterator(0); + auto const num_segments = static_cast(offsets.size()) - 1; + + bool const is_argmin = std::is_same_v; + auto string_comparator = + element_arg_minmax_fn{*device_col, col.has_nulls(), is_argmin}; + auto const identity = is_argmin ? cudf::detail::ARGMIN_SENTINEL : cudf::detail::ARGMAX_SENTINEL; + + auto gather_map = + cudf::reduction::detail::segmented_reduce(it, + offsets.begin(), + num_segments, + string_comparator, + identity, + stream, + rmm::mr::get_current_device_resource()); + auto result = std::move(cudf::detail::gather(table_view{{col}}, + *gather_map, + cudf::out_of_bounds_policy::NULLIFY, + cudf::detail::negative_index_policy::NOT_ALLOWED, + stream, + mr) + ->release()[0]); + auto const [segmented_null_mask, segmented_null_count] = + cudf::detail::segmented_null_mask_reduction(col.null_mask(), + offsets.begin(), + offsets.end() - 1, + offsets.begin() + 1, + null_handling, + stream, + mr); + + // If the segmented null mask contains any null values, the segmented null mask + // must be combined with the result null mask. + if (segmented_null_count > 0) { + if (result->null_count() == 0) { + // The result has no nulls. Use the segmented null mask. + result->set_null_mask(segmented_null_mask, segmented_null_count, stream); + } else { + // Compute the logical AND of the segmented output null mask and the + // result null mask to update the result null mask and null count. + auto result_mview = result->mutable_view(); + std::vector masks{ + static_cast(result_mview.null_mask()), + static_cast(segmented_null_mask.data())}; + std::vector begin_bits{0, 0}; + auto const valid_count = cudf::detail::inplace_bitmask_and( + device_span(static_cast(result_mview.null_mask()), + num_bitmask_words(result->size())), + masks, + begin_bits, + result->size(), + stream, + mr); + result->set_null_count(result->size() - valid_count); + } + } + + return result; +} + +template () && + !std::is_same_v())> +std::unique_ptr string_simple_segmented_reduction(column_view const& col, + device_span offsets, + null_policy null_handling, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + CUDF_FAIL("Segmented reduction on string column only supports min and max reduction."); +} + /** * @brief Call reduce and return a column of type bool. * @@ -153,7 +251,9 @@ struct same_column_type_dispatcher { } public: - template ()>* = nullptr> + template () && + !std::is_same_v)> std::unique_ptr operator()(column_view const& col, device_span offsets, null_policy null_handling, @@ -164,7 +264,19 @@ struct same_column_type_dispatcher { col, offsets, null_handling, stream, mr); } - template ()>* = nullptr> + template () && std::is_same_v)> + std::unique_ptr operator()(column_view const& col, + device_span offsets, + null_policy null_handling, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) + { + return string_simple_segmented_reduction( + col, offsets, null_handling, stream, mr); + } + + template ())> std::unique_ptr operator()(column_view const&, device_span, null_policy, From fe5da2d526cb33b689311e28f77946f96d11e6af Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Mon, 11 Apr 2022 11:46:23 -0700 Subject: [PATCH 08/20] Add docstrings for `string_segmented_reduction` and tparams cleanups, removed unused functions in `reduction.cuh`. Co-authored-by: Bradley Dice --- cpp/include/cudf/detail/reduction.cuh | 57 +------------------------ cpp/src/reductions/simple_segmented.cuh | 44 +++++++++++++------ 2 files changed, 33 insertions(+), 68 deletions(-) diff --git a/cpp/include/cudf/detail/reduction.cuh b/cpp/include/cudf/detail/reduction.cuh index eb8668ddc3f..cf95678f164 100644 --- a/cpp/include/cudf/detail/reduction.cuh +++ b/cpp/include/cudf/detail/reduction.cuh @@ -291,65 +291,12 @@ std::unique_ptr segmented_reduce(InputIterator d_in, return dev_result; } -/** - * @brief string column specialization for segmented reduction. - * - * This specialization creates the gather map from which - */ -template ::type, - typename std::enable_if_t>* = nullptr> -std::unique_ptr segmented_reduce(InputIterator d_in, - OffsetIterator d_offset, - cudf::size_type num_segments, - BinaryOp binary_op, - OutputType identity, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) -{ - auto gather_map = make_fixed_width_column( - data_type{type_to_id()}, num_segments, mask_state::UNALLOCATED, stream, mr); - auto dev_result_mview = gather_map->mutable_view(); - - // Allocate temporary storage - rmm::device_buffer d_temp_storage; - size_t temp_storage_bytes = 0; - cub::DeviceSegmentedReduce::Reduce(d_temp_storage.data(), - temp_storage_bytes, - d_in, - dev_result_mview.data(), - num_segments, - d_offset, - d_offset + 1, - binary_op, - identity, - stream.value()); - d_temp_storage = rmm::device_buffer{temp_storage_bytes, stream}; - - // Run reduction - cub::DeviceSegmentedReduce::Reduce(d_temp_storage.data(), - temp_storage_bytes, - d_in, - dev_result_mview.data(), - num_segments, - d_offset, - d_offset + 1, - binary_op, - identity, - stream.value()); - - return gather_map; -} - template ::type, - typename std::enable_if_t<(!is_fixed_width() && - !std::is_same_v()) || - is_fixed_point()>* = nullptr> + typename std::enable_if_t() && + !cudf::is_fixed_point())>* = nullptr> std::unique_ptr segmented_reduce(InputIterator, OffsetIterator, cudf::size_type, diff --git a/cpp/src/reductions/simple_segmented.cuh b/cpp/src/reductions/simple_segmented.cuh index a19412afc04..94d6f7081f1 100644 --- a/cpp/src/reductions/simple_segmented.cuh +++ b/cpp/src/reductions/simple_segmented.cuh @@ -112,16 +112,35 @@ std::unique_ptr simple_segmented_reduction(column_view const& col, return result; } + +/** + * @brief String segmented reduction for 'min', 'max'. + * + * This algorithm uses argmin/argmax as a custom comparator to build a gather + * map, then builds the output. + * + * @tparam InputType the input column data-type + * @tparam Op the operator of cudf::reduction::op:: + + * @param col Input column of data to reduce. + * @param offsets Indices to segment boundaries. + * @param null_handling If `null_policy::INCLUDE`, all elements in a segment + * must be valid for the reduced value to be valid. If `null_policy::EXCLUDE`, + * the reduced value is valid if any element in the segment is valid. + * @param stream Used for device memory operations and kernel launches. + * @param mr Device memory resource used to allocate the returned column's device memory + * @return Output column in device memory + */ + template || std::is_same_v)> -std::unique_ptr string_simple_segmented_reduction(column_view const& col, - device_span offsets, - null_policy null_handling, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) +std::unique_ptr string_segmented_reduction(column_view const& col, + device_span offsets, + null_policy null_handling, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) { // Pass to simple_segmented_reduction, get indices to gather, perform gather here. auto device_col = cudf::column_device_view::create(col, stream); @@ -188,15 +207,14 @@ std::unique_ptr string_simple_segmented_reduction(column_view const& col } template () && !std::is_same_v())> -std::unique_ptr string_simple_segmented_reduction(column_view const& col, - device_span offsets, - null_policy null_handling, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) +std::unique_ptr string_segmented_reduction(column_view const& col, + device_span offsets, + null_policy null_handling, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) { CUDF_FAIL("Segmented reduction on string column only supports min and max reduction."); } @@ -272,7 +290,7 @@ struct same_column_type_dispatcher { rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - return string_simple_segmented_reduction( + return string_segmented_reduction( col, offsets, null_handling, stream, mr); } From ab801f143fa612c7698737af4125308406c1cff5 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Mon, 11 Apr 2022 12:37:49 -0700 Subject: [PATCH 09/20] Move element argminmax to a separate file. Co-authored-by: Bradley Dice --- .../detail/utilities/device_operators.cuh | 32 --------- .../detail/utilities/element_argminmax.cuh | 66 +++++++++++++++++++ .../sort/group_single_pass_reduction_util.cuh | 4 +- cpp/src/reductions/simple_segmented.cuh | 4 +- 4 files changed, 70 insertions(+), 36 deletions(-) create mode 100644 cpp/include/cudf/detail/utilities/element_argminmax.cuh diff --git a/cpp/include/cudf/detail/utilities/device_operators.cuh b/cpp/include/cudf/detail/utilities/device_operators.cuh index 9f75f85bcf3..87fef5bc187 100644 --- a/cpp/include/cudf/detail/utilities/device_operators.cuh +++ b/cpp/include/cudf/detail/utilities/device_operators.cuh @@ -21,7 +21,6 @@ * @file */ -#include #include #include #include @@ -263,35 +262,4 @@ struct DeviceLeadLag { explicit CUDF_HOST_DEVICE inline DeviceLeadLag(size_type offset_) : row_offset(offset_) {} }; -/** - * @brief Binary `argmin`/`argmax` operator - * - * @tparam T Type of the underlying column. Must support '<' operator. - */ -template -struct element_arg_minmax_fn { - column_device_view const d_col; - bool const has_nulls; - bool const arg_min; - - __device__ inline auto operator()(size_type const& lhs_idx, size_type const& rhs_idx) const - { - // The extra bounds checking is due to issue github.com/rapidsai/cudf/9156 and - // github.com/NVIDIA/thrust/issues/1525 - // where invalid random values may be passed here by thrust::reduce_by_key - if (lhs_idx < 0 || lhs_idx >= d_col.size() || (has_nulls && d_col.is_null_nocheck(lhs_idx))) { - return rhs_idx; - } - if (rhs_idx < 0 || rhs_idx >= d_col.size() || (has_nulls && d_col.is_null_nocheck(rhs_idx))) { - return lhs_idx; - } - - // Return `lhs_idx` iff: - // row(lhs_idx) < row(rhs_idx) and finding ArgMin, or - // row(lhs_idx) >= row(rhs_idx) and finding ArgMax. - auto const less = d_col.element(lhs_idx) < d_col.element(rhs_idx); - return less == arg_min ? lhs_idx : rhs_idx; - } -}; - } // namespace cudf diff --git a/cpp/include/cudf/detail/utilities/element_argminmax.cuh b/cpp/include/cudf/detail/utilities/element_argminmax.cuh new file mode 100644 index 00000000000..e4bbac585d9 --- /dev/null +++ b/cpp/include/cudf/detail/utilities/element_argminmax.cuh @@ -0,0 +1,66 @@ +/* + * Copyright (c) 2022, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +/** + * @brief Definition of the device operators + * @file + */ + +#include +#include +#include +#include + +#include + +namespace cudf { +namespace detail { + +/** + * @brief Binary `argmin`/`argmax` operator + * + * @tparam T Type of the underlying column. Must support '<' operator. + */ +template +struct element_argminmax_fn { + column_device_view const d_col; + bool const has_nulls; + bool const arg_min; + + __device__ inline auto operator()(size_type const& lhs_idx, size_type const& rhs_idx) const + { + // The extra bounds checking is due to issue github.com/rapidsai/cudf/9156 and + // github.com/NVIDIA/thrust/issues/1525 + // where invalid random values may be passed here by thrust::reduce_by_key + if (lhs_idx < 0 || lhs_idx >= d_col.size() || (has_nulls && d_col.is_null_nocheck(lhs_idx))) { + return rhs_idx; + } + if (rhs_idx < 0 || rhs_idx >= d_col.size() || (has_nulls && d_col.is_null_nocheck(rhs_idx))) { + return lhs_idx; + } + + // Return `lhs_idx` iff: + // row(lhs_idx) < row(rhs_idx) and finding ArgMin, or + // row(lhs_idx) >= row(rhs_idx) and finding ArgMax. + auto const less = d_col.element(lhs_idx) < d_col.element(rhs_idx); + return less == arg_min ? lhs_idx : rhs_idx; + } +}; + +} // namespace detail +} // namespace cudf 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 ee99b20f12d..4b84642a491 100644 --- a/cpp/src/groupby/sort/group_single_pass_reduction_util.cuh +++ b/cpp/src/groupby/sort/group_single_pass_reduction_util.cuh @@ -23,7 +23,7 @@ #include #include #include -#include +#include #include #include #include @@ -182,7 +182,7 @@ struct group_reduction_functor(0); auto const binop = - element_arg_minmax_fn{*d_values_ptr, values.has_nulls(), K == aggregation::ARGMIN}; + cudf::detail::element_argminmax_fn{*d_values_ptr, values.has_nulls(), K == aggregation::ARGMIN}; do_reduction(count_iter, result_begin, binop); } else { using OpType = cudf::detail::corresponding_operator_t; diff --git a/cpp/src/reductions/simple_segmented.cuh b/cpp/src/reductions/simple_segmented.cuh index 94d6f7081f1..dd8edb01cd2 100644 --- a/cpp/src/reductions/simple_segmented.cuh +++ b/cpp/src/reductions/simple_segmented.cuh @@ -24,7 +24,7 @@ #include #include #include -#include +#include #include #include #include @@ -150,7 +150,7 @@ std::unique_ptr string_segmented_reduction(column_view const& col, bool const is_argmin = std::is_same_v; auto string_comparator = - element_arg_minmax_fn{*device_col, col.has_nulls(), is_argmin}; + cudf::detail::element_argminmax_fn{*device_col, col.has_nulls(), is_argmin}; auto const identity = is_argmin ? cudf::detail::ARGMIN_SENTINEL : cudf::detail::ARGMAX_SENTINEL; auto gather_map = From 5e7ca891b853dde00d853f2bdac4ac848ae5f01d Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Mon, 11 Apr 2022 12:38:10 -0700 Subject: [PATCH 10/20] Add tests for string segmented reduction. Co-authored-by: Bradley Dice --- .../reductions/segmented_reduction_tests.cpp | 224 ++++++++++++------ 1 file changed, 157 insertions(+), 67 deletions(-) diff --git a/cpp/tests/reductions/segmented_reduction_tests.cpp b/cpp/tests/reductions/segmented_reduction_tests.cpp index 7a89b30fb6e..da2f5c9d2cc 100644 --- a/cpp/tests/reductions/segmented_reduction_tests.cpp +++ b/cpp/tests/reductions/segmented_reduction_tests.cpp @@ -14,13 +14,11 @@ * limitations under the License. */ -#include "cudf_test/column_utilities.hpp" #include #include #include #include -#include #include #include @@ -40,7 +38,7 @@ struct SegmentedReductionTest : public cudf::test::BaseFixture { struct SegmentedReductionTestUntyped : public cudf::test::BaseFixture { }; -TYPED_TEST_SUITE(SegmentedReductionTest, NumericTypes); +TYPED_TEST_CASE(SegmentedReductionTest, NumericTypes); TYPED_TEST(SegmentedReductionTest, SumExcludeNulls) { @@ -389,91 +387,183 @@ TEST_F(SegmentedReductionTestUntyped, ReduceEmptyColumn) CUDF_TEST_EXPECT_COLUMNS_EQUAL(*res, expect); } -int32_t pow10(int32_t exponent) { return exponent == 0 ? 1 : 10 * pow10(exponent - 1); } +// int32_t pow10(int32_t exponent) { return exponent == 0 ? 1 : 10 * pow10(exponent - 1); } + +// template +// struct SegmentedReductionFixedPointTest : public cudf::test::BaseFixture { +// public: +// std::vector scale_list_by_pow10(std::vector input, +// int32_t exponent) +// { +// std::vector result(input.size()); +// std::transform(input.begin(), input.end(), result.begin(), [&exponent](auto x) { +// return exponent >= 0 ? x * pow10(exponent) : x / pow10(-exponent); +// }); +// return result; +// } +// }; + +// TYPED_TEST_SUITE(SegmentedReductionFixedPointTest, cudf::test::FixedPointTypes); + +// TYPED_TEST(SegmentedReductionFixedPointTest, MaxIncludeNullsScaleZero) +// { +// // [1, 2, 3], [1], [], [2, NULL, 3], [NULL], [NULL, NULL] | scale: 0 +// // values: {1, 2, 3, 1, 2, XXX, 3, XXX, XXX, XXX} +// // nullmask:{1, 1, 1, 1, 1, 0, 1, 0, 0, 0} +// // offsets: {0, 3, 4, 4, 7, 8, 10} +// // output_dtype: decimalXX, scale: -1, 0, 1 +// // outputs: {3, 1, XXX, XXX, XXX, XXX} +// // output nullmask: {1, 1, 0, 0, 0, 0} + +// using DecimalXX = TypeParam; + +// for (int output_scale : {-1, 0, 1}) { +// fixed_point_column_wrapper input{ +// {1, 2, 3, 1, 2, XXX, 3, XXX, XXX, XXX}, +// {true, true, true, true, true, false, true, false, false, false}, +// numeric::scale_type(0)}; +// fixed_width_column_wrapper offsets{0, 3, 4, 4, 7, 8, 10}; + +// data_type output_dtype{type_to_id(), numeric::scale_type{output_scale}}; + +// auto result_rep = this->scale_list_by_pow10({3, 1, XXX, XXX, XXX, XXX}, -output_scale); +// fixed_point_column_wrapper expect{ +// result_rep.begin(), +// result_rep.end(), +// {true, true, false, false, false, false}, +// numeric::scale_type(output_scale)}; + +// auto res = segmented_reduce(input, +// column_view(offsets), +// *make_max_aggregation(), +// output_dtype, +// null_policy::INCLUDE); +// CUDF_TEST_EXPECT_COLUMNS_EQUAL(*res, expect); +// } +// } + +// String min/max test grid +// Segment: Length 0, length 1, length 2 +// Element nulls: No nulls, all nulls, some nulls +// String: Empty string, +// Position of the min/max: start of segment, end of segment +// Include null, exclude null -template -struct SegmentedReductionFixedPointTest : public cudf::test::BaseFixture { - public: - std::vector scale_list_by_pow10(std::vector input, - int32_t exponent) - { - std::vector result(input.size()); - std::transform(input.begin(), input.end(), result.begin(), [&exponent](auto x) { - return exponent >= 0 ? x * pow10(exponent) : x / pow10(-exponent); - }); - return result; - } +struct SegmentedReductionStringTest : public cudf::test::BaseFixture { }; -TYPED_TEST_SUITE(SegmentedReductionFixedPointTest, cudf::test::FixedPointTypes); - -TYPED_TEST(SegmentedReductionFixedPointTest, MaxIncludeNullsScaleZero) +TEST_F(SegmentedReductionStringTest, MaxIncludeNulls) { - // [1, 2, 3], [1], [], [2, NULL, 3], [NULL], [NULL, NULL] | scale: 0 - // values: {1, 2, 3, 1, 2, XXX, 3, XXX, XXX, XXX} - // nullmask:{1, 1, 1, 1, 1, 0, 1, 0, 0, 0} - // offsets: {0, 3, 4, 4, 7, 8, 10} - // output_dtype: decimalXX, scale: -1, 0, 1 - // outputs: {3, 1, XXX, XXX, XXX, XXX} - // output nullmask: {1, 1, 0, 0, 0, 0} - - using DecimalXX = TypeParam; - - for (int output_scale : {-1, 0, 1}) { - fixed_point_column_wrapper input{ - {1, 2, 3, 1, 2, XXX, 3, XXX, XXX, XXX}, - {true, true, true, true, true, false, true, false, false, false}, - numeric::scale_type(0)}; - fixed_width_column_wrapper offsets{0, 3, 4, 4, 7, 8, 10}; - - data_type output_dtype{type_to_id(), numeric::scale_type{output_scale}}; - - auto result_rep = this->scale_list_by_pow10({3, 1, XXX, XXX, XXX, XXX}, -output_scale); - fixed_point_column_wrapper expect{ - result_rep.begin(), - result_rep.end(), - {true, true, false, false, false, false}, - numeric::scale_type(output_scale)}; - - auto res = segmented_reduce(input, - column_view(offsets), - *make_max_aggregation(), - output_dtype, - null_policy::INCLUDE); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*res, expect); - } -} + // data: ['world'], ['cudf', NULL, ''], ['rapids', 'i am', 'ai'], ['apples', 'zebras'], + // [], [NULL], [NULL, NULL] + // values: {"world", "cudf", XXX, "", "rapids", "i am", "ai", "apples", "zebras", XXX, XXX, XXX} + // nullmask:{1, 1, 0, 1, 1, 1, 1, 1, 1, 0, 0, 0} + // offsets: {0, 1, 4, 7, 9, 9, 10, 12} + // output_dtype: string dtype + // outputs: {"world", XXX, "rapids", "zebras", XXX, XXX, XXX} + // output nullmask: {1, 0, 1, 1, 0, 0, 0} -struct SegmentedReductionStringTest : public cudf::test::BaseFixture { -}; + strings_column_wrapper input{ + {"world", "cudf", XXX, "", "rapids", "i am", "ai", "apples", "zebras", XXX, XXX, XXX}, + {true, true, false, true, true, true, true, true, true, false, false, false}}; + fixed_width_column_wrapper offsets{0, 1, 4, 7, 9, 9, 10, 12}; + data_type output_dtype{type_id::STRING}; -TEST_F(SegmentedReductionStringTest, MaxIncludeNulls) + strings_column_wrapper expect{{"world", XXX, "rapids", "zebras", XXX, XXX, XXX}, + {true, false, true, true, false, false, false}}; + + auto res = segmented_reduce(input, + column_view(offsets), + *make_max_aggregation(), + output_dtype, + null_policy::INCLUDE); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*res, expect); +} + +TEST_F(SegmentedReductionStringTest, MaxExcludeNulls) { - // ['world'], ['cudf', NULL, 'cuml'], ['hello', 'rapids', 'ai'], [], [NULL], [NULL, NULL] - // values: {"world", "cudf", XXX, "cuml", "hello", "rapids", "ai", XXX, XXX, XXX} - // nullmask:{1, 1, 0, 1, 1, 1, 1, 0, 0, 0} - // offsets: {0, 1, 4, 7, 7, 8, 10} + // data: ['world'], ['cudf', NULL, ''], ['rapids', 'i am', 'ai'], ['apples', 'zebras'], + // [], [NULL], [NULL, NULL] + // values: {"world", "cudf", XXX, "", "rapids", "i am", "ai", "apples", "zebras", XXX, XXX, XXX} + // nullmask:{1, 1, 0, 1, 1, 1, 1, 1, 1, 0, 0, 0} + // offsets: {0, 1, 4, 7, 9, 9, 10, 12} // output_dtype: string dtype - // outputs: {"world", XXX, "rapids", XXX, XXX, XXX} - // output nullmask: {1, 0, 1, 0, 0, 0} + // outputs: {"world", "cudf", "rapids", "zebras", XXX, XXX, XXX} + // output nullmask: {1, 1, 1, 1, 0, 0, 0} strings_column_wrapper input{ - {"world", "cudf", XXX, "cuml", "hello", "rapids", "ai", XXX, XXX, XXX}, - {true, true, false, true, true, true, true, false, false, false}}; - fixed_width_column_wrapper offsets{0, 1, 4, 7, 7, 8, 10}; + {"world", "cudf", XXX, "", "rapids", "i am", "ai", "apples", "zebras", XXX, XXX, XXX}, + {true, true, false, true, true, true, true, true, true, false, false, false}}; + fixed_width_column_wrapper offsets{0, 1, 4, 7, 9, 9, 10, 12}; data_type output_dtype{type_id::STRING}; - strings_column_wrapper expect{{"world", XXX, "rapids", XXX, XXX, XXX}, - {true, false, true, false, false, false}}; + strings_column_wrapper expect{{"world", "cudf", "rapids", "zebras", XXX, XXX, XXX}, + {true, true, true, true, false, false, false}}; auto res = segmented_reduce(input, column_view(offsets), *make_max_aggregation(), output_dtype, + null_policy::EXCLUDE); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*res, expect); +} + +TEST_F(SegmentedReductionStringTest, MinIncludeNulls) +{ + // data: ['world'], ['cudf', NULL, ''], ['rapids', 'i am', 'ai'], ['apples', 'zebras'], + // [], [NULL], [NULL, NULL] + // values: {"world", "cudf", XXX, "", "rapids", "i am", "ai", "apples", "zebras", XXX, XXX, XXX} + // nullmask:{1, 1, 0, 1, 1, 1, 1, 1, 1, 0, 0, 0} + // offsets: {0, 1, 4, 7, 9, 9, 10, 12} + // output_dtype: string dtype + // outputs: {"world", XXX, "ai", "apples", XXX, XXX, XXX} + // output nullmask: {1, 0, 1, 1, 0, 0, 0} + + strings_column_wrapper input{ + {"world", "cudf", XXX, "", "rapids", "i am", "ai", "apples", "zebras", XXX, XXX, XXX}, + {true, true, false, true, true, true, true, true, true, false, false, false}}; + fixed_width_column_wrapper offsets{0, 1, 4, 7, 9, 9, 10, 12}; + data_type output_dtype{type_id::STRING}; + + strings_column_wrapper expect{{"world", XXX, "ai", "apples", XXX, XXX, XXX}, + {true, false, true, true, false, false, false}}; + + auto res = segmented_reduce(input, + column_view(offsets), + *make_min_aggregation(), + output_dtype, null_policy::INCLUDE); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*res, expect); } +TEST_F(SegmentedReductionStringTest, MinExcludeNulls) +{ + // data: ['world'], ['cudf', NULL, ''], ['rapids', 'i am', 'ai'], ['apples', 'zebras'], + // [], [NULL], [NULL, NULL] + // values: {"world", "cudf", XXX, "", "rapids", "i am", "ai", "apples", "zebras", XXX, XXX, XXX} + // nullmask:{1, 1, 0, 1, 1, 1, 1, 1, 1, 0, 0, 0} + // offsets: {0, 1, 4, 7, 9, 9, 10, 12} + // output_dtype: string dtype + // outputs: {"world", "", "ai", "apples", XXX, XXX, XXX} + // output nullmask: {1, 1, 1, 1, 0, 0, 0} + + strings_column_wrapper input{ + {"world", "cudf", XXX, "", "rapids", "i am", "ai", "apples", "zebras", XXX, XXX, XXX}, + {true, true, false, true, true, true, true, true, true, false, false, false}}; + fixed_width_column_wrapper offsets{0, 1, 4, 7, 9, 9, 10, 12}; + data_type output_dtype{type_id::STRING}; + + strings_column_wrapper expect{{"world", "", "ai", "apples", XXX, XXX, XXX}, + {true, true, true, true, false, false, false}}; + + auto res = segmented_reduce(input, + column_view(offsets), + *make_min_aggregation(), + output_dtype, + null_policy::EXCLUDE); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*res, expect); +} + #undef XXX } // namespace test From 680e0fe8f04788942a415e5f80105a07624cb5bb Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Mon, 11 Apr 2022 12:39:27 -0700 Subject: [PATCH 11/20] Remove stale commented out codes. --- .../reductions/segmented_reduction_tests.cpp | 55 ------------------- 1 file changed, 55 deletions(-) diff --git a/cpp/tests/reductions/segmented_reduction_tests.cpp b/cpp/tests/reductions/segmented_reduction_tests.cpp index da2f5c9d2cc..281b6b08dc2 100644 --- a/cpp/tests/reductions/segmented_reduction_tests.cpp +++ b/cpp/tests/reductions/segmented_reduction_tests.cpp @@ -387,61 +387,6 @@ TEST_F(SegmentedReductionTestUntyped, ReduceEmptyColumn) CUDF_TEST_EXPECT_COLUMNS_EQUAL(*res, expect); } -// int32_t pow10(int32_t exponent) { return exponent == 0 ? 1 : 10 * pow10(exponent - 1); } - -// template -// struct SegmentedReductionFixedPointTest : public cudf::test::BaseFixture { -// public: -// std::vector scale_list_by_pow10(std::vector input, -// int32_t exponent) -// { -// std::vector result(input.size()); -// std::transform(input.begin(), input.end(), result.begin(), [&exponent](auto x) { -// return exponent >= 0 ? x * pow10(exponent) : x / pow10(-exponent); -// }); -// return result; -// } -// }; - -// TYPED_TEST_SUITE(SegmentedReductionFixedPointTest, cudf::test::FixedPointTypes); - -// TYPED_TEST(SegmentedReductionFixedPointTest, MaxIncludeNullsScaleZero) -// { -// // [1, 2, 3], [1], [], [2, NULL, 3], [NULL], [NULL, NULL] | scale: 0 -// // values: {1, 2, 3, 1, 2, XXX, 3, XXX, XXX, XXX} -// // nullmask:{1, 1, 1, 1, 1, 0, 1, 0, 0, 0} -// // offsets: {0, 3, 4, 4, 7, 8, 10} -// // output_dtype: decimalXX, scale: -1, 0, 1 -// // outputs: {3, 1, XXX, XXX, XXX, XXX} -// // output nullmask: {1, 1, 0, 0, 0, 0} - -// using DecimalXX = TypeParam; - -// for (int output_scale : {-1, 0, 1}) { -// fixed_point_column_wrapper input{ -// {1, 2, 3, 1, 2, XXX, 3, XXX, XXX, XXX}, -// {true, true, true, true, true, false, true, false, false, false}, -// numeric::scale_type(0)}; -// fixed_width_column_wrapper offsets{0, 3, 4, 4, 7, 8, 10}; - -// data_type output_dtype{type_to_id(), numeric::scale_type{output_scale}}; - -// auto result_rep = this->scale_list_by_pow10({3, 1, XXX, XXX, XXX, XXX}, -output_scale); -// fixed_point_column_wrapper expect{ -// result_rep.begin(), -// result_rep.end(), -// {true, true, false, false, false, false}, -// numeric::scale_type(output_scale)}; - -// auto res = segmented_reduce(input, -// column_view(offsets), -// *make_max_aggregation(), -// output_dtype, -// null_policy::INCLUDE); -// CUDF_TEST_EXPECT_COLUMNS_EQUAL(*res, expect); -// } -// } - // String min/max test grid // Segment: Length 0, length 1, length 2 // Element nulls: No nulls, all nulls, some nulls From 19e5df80927101ffd209ac7a6e7d307d6058b4b1 Mon Sep 17 00:00:00 2001 From: Bradley Dice Date: Mon, 11 Apr 2022 12:39:55 -0700 Subject: [PATCH 12/20] clang-format --- cpp/include/cudf/detail/reduction.cuh | 2 +- cpp/src/groupby/sort/group_single_pass_reduction_util.cuh | 4 ++-- cpp/src/reductions/simple_segmented.cuh | 4 +--- 3 files changed, 4 insertions(+), 6 deletions(-) diff --git a/cpp/include/cudf/detail/reduction.cuh b/cpp/include/cudf/detail/reduction.cuh index cf95678f164..023d83f3c24 100644 --- a/cpp/include/cudf/detail/reduction.cuh +++ b/cpp/include/cudf/detail/reduction.cuh @@ -296,7 +296,7 @@ template ::type, typename std::enable_if_t() && - !cudf::is_fixed_point())>* = nullptr> + !cudf::is_fixed_point())>* = nullptr> std::unique_ptr segmented_reduce(InputIterator, OffsetIterator, cudf::size_type, 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 4b84642a491..93d5e6c032c 100644 --- a/cpp/src/groupby/sort/group_single_pass_reduction_util.cuh +++ b/cpp/src/groupby/sort/group_single_pass_reduction_util.cuh @@ -181,8 +181,8 @@ struct group_reduction_functor(0); - auto const binop = - cudf::detail::element_argminmax_fn{*d_values_ptr, values.has_nulls(), K == aggregation::ARGMIN}; + auto const binop = cudf::detail::element_argminmax_fn{ + *d_values_ptr, values.has_nulls(), K == aggregation::ARGMIN}; do_reduction(count_iter, result_begin, binop); } else { using OpType = cudf::detail::corresponding_operator_t; diff --git a/cpp/src/reductions/simple_segmented.cuh b/cpp/src/reductions/simple_segmented.cuh index dd8edb01cd2..c5347811da2 100644 --- a/cpp/src/reductions/simple_segmented.cuh +++ b/cpp/src/reductions/simple_segmented.cuh @@ -112,7 +112,6 @@ std::unique_ptr simple_segmented_reduction(column_view const& col, return result; } - /** * @brief String segmented reduction for 'min', 'max'. * @@ -290,8 +289,7 @@ struct same_column_type_dispatcher { rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - return string_segmented_reduction( - col, offsets, null_handling, stream, mr); + return string_segmented_reduction(col, offsets, null_handling, stream, mr); } template ())> From d826cf9155cd2c9c924310fc6008d47745cc9c16 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Mon, 11 Apr 2022 16:11:46 -0700 Subject: [PATCH 13/20] Redefine null placeholder. --- cpp/tests/reductions/segmented_reduction_tests.cpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/cpp/tests/reductions/segmented_reduction_tests.cpp b/cpp/tests/reductions/segmented_reduction_tests.cpp index 281b6b08dc2..42bd5064be2 100644 --- a/cpp/tests/reductions/segmented_reduction_tests.cpp +++ b/cpp/tests/reductions/segmented_reduction_tests.cpp @@ -394,6 +394,9 @@ TEST_F(SegmentedReductionTestUntyped, ReduceEmptyColumn) // Position of the min/max: start of segment, end of segment // Include null, exclude null +#undef XXX +#define XXX "" // null placeholder + struct SegmentedReductionStringTest : public cudf::test::BaseFixture { }; From 2d3bd6bcb8e2589d4d6bab692241654d564f76fc Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 14 Apr 2022 17:02:39 -0700 Subject: [PATCH 14/20] Use constexpr --- cpp/src/reductions/simple_segmented.cuh | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/cpp/src/reductions/simple_segmented.cuh b/cpp/src/reductions/simple_segmented.cuh index a19412afc04..ed990f9523c 100644 --- a/cpp/src/reductions/simple_segmented.cuh +++ b/cpp/src/reductions/simple_segmented.cuh @@ -129,10 +129,11 @@ std::unique_ptr string_simple_segmented_reduction(column_view const& col auto it = thrust::make_counting_iterator(0); auto const num_segments = static_cast(offsets.size()) - 1; - bool const is_argmin = std::is_same_v; + bool constexpr is_argmin = std::is_same_v; auto string_comparator = element_arg_minmax_fn{*device_col, col.has_nulls(), is_argmin}; - auto const identity = is_argmin ? cudf::detail::ARGMIN_SENTINEL : cudf::detail::ARGMAX_SENTINEL; + auto constexpr identity = + is_argmin ? cudf::detail::ARGMIN_SENTINEL : cudf::detail::ARGMAX_SENTINEL; auto gather_map = cudf::reduction::detail::segmented_reduce(it, From d4ad909dd7d0c95c946d2369e813df585777f270 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 14 Apr 2022 17:03:50 -0700 Subject: [PATCH 15/20] adopt ctad for vector types --- cpp/src/reductions/simple_segmented.cuh | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/cpp/src/reductions/simple_segmented.cuh b/cpp/src/reductions/simple_segmented.cuh index ed990f9523c..3ae18d92db1 100644 --- a/cpp/src/reductions/simple_segmented.cuh +++ b/cpp/src/reductions/simple_segmented.cuh @@ -169,9 +169,8 @@ std::unique_ptr string_simple_segmented_reduction(column_view const& col // Compute the logical AND of the segmented output null mask and the // result null mask to update the result null mask and null count. auto result_mview = result->mutable_view(); - std::vector masks{ - static_cast(result_mview.null_mask()), - static_cast(segmented_null_mask.data())}; + std::vector masks{static_cast(result_mview.null_mask()), + static_cast(segmented_null_mask.data())}; std::vector begin_bits{0, 0}; auto const valid_count = cudf::detail::inplace_bitmask_and( device_span(static_cast(result_mview.null_mask()), From 4a3a541f97b326f33f3b0bbbad23e201b5f44a35 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 14 Apr 2022 17:21:25 -0700 Subject: [PATCH 16/20] Wrap argminmax idx condition with lambda --- .../cudf/detail/utilities/element_argminmax.cuh | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/cpp/include/cudf/detail/utilities/element_argminmax.cuh b/cpp/include/cudf/detail/utilities/element_argminmax.cuh index e4bbac585d9..a452b40286e 100644 --- a/cpp/include/cudf/detail/utilities/element_argminmax.cuh +++ b/cpp/include/cudf/detail/utilities/element_argminmax.cuh @@ -47,12 +47,12 @@ struct element_argminmax_fn { // The extra bounds checking is due to issue github.com/rapidsai/cudf/9156 and // github.com/NVIDIA/thrust/issues/1525 // where invalid random values may be passed here by thrust::reduce_by_key - if (lhs_idx < 0 || lhs_idx >= d_col.size() || (has_nulls && d_col.is_null_nocheck(lhs_idx))) { - return rhs_idx; - } - if (rhs_idx < 0 || rhs_idx >= d_col.size() || (has_nulls && d_col.is_null_nocheck(rhs_idx))) { - return lhs_idx; - } + auto out_of_bound_or_null = [d_col = &this->d_col, + has_nulls = &this->has_nulls](auto const& idx) { + return idx < 0 || idx >= d_col.size() || (has_nulls && d_col.is_null_nocheck(idx)); + }; + if (out_of_bound_or_null(lhs_idx)) { return rhs_idx; } + if (out_of_bound_or_null(rhs_idx)) { return lhs_idx; } // Return `lhs_idx` iff: // row(lhs_idx) < row(rhs_idx) and finding ArgMin, or From 6910c164cb115a3615eecc238c92d5a647451bab Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 14 Apr 2022 17:57:04 -0700 Subject: [PATCH 17/20] Minor bug fix --- cpp/include/cudf/detail/utilities/element_argminmax.cuh | 6 +++--- cpp/src/reductions/simple_segmented.cuh | 2 +- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/cpp/include/cudf/detail/utilities/element_argminmax.cuh b/cpp/include/cudf/detail/utilities/element_argminmax.cuh index a452b40286e..428b0740af5 100644 --- a/cpp/include/cudf/detail/utilities/element_argminmax.cuh +++ b/cpp/include/cudf/detail/utilities/element_argminmax.cuh @@ -47,9 +47,9 @@ struct element_argminmax_fn { // The extra bounds checking is due to issue github.com/rapidsai/cudf/9156 and // github.com/NVIDIA/thrust/issues/1525 // where invalid random values may be passed here by thrust::reduce_by_key - auto out_of_bound_or_null = [d_col = &this->d_col, - has_nulls = &this->has_nulls](auto const& idx) { - return idx < 0 || idx >= d_col.size() || (has_nulls && d_col.is_null_nocheck(idx)); + auto out_of_bound_or_null = [this] __device__(size_type const& idx) { + return idx < 0 || idx >= this->d_col.size() || + (this->has_nulls && this->d_col.is_null_nocheck(idx)); }; if (out_of_bound_or_null(lhs_idx)) { return rhs_idx; } if (out_of_bound_or_null(rhs_idx)) { return lhs_idx; } diff --git a/cpp/src/reductions/simple_segmented.cuh b/cpp/src/reductions/simple_segmented.cuh index 33d1422ae23..7796794502d 100644 --- a/cpp/src/reductions/simple_segmented.cuh +++ b/cpp/src/reductions/simple_segmented.cuh @@ -149,7 +149,7 @@ std::unique_ptr string_segmented_reduction(column_view const& col, bool constexpr is_argmin = std::is_same_v; auto string_comparator = - element_arg_minmax_fn{*device_col, col.has_nulls(), is_argmin}; + cudf::detail::element_argminmax_fn{*device_col, col.has_nulls(), is_argmin}; auto constexpr identity = is_argmin ? cudf::detail::ARGMIN_SENTINEL : cudf::detail::ARGMAX_SENTINEL; From 1d0b71696fb67e1eff7ef3d4ea1b1e25a12fd011 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 14 Apr 2022 18:07:17 -0700 Subject: [PATCH 18/20] Consolidate input data as fixtures --- .../reductions/segmented_reduction_tests.cpp | 28 ++++++++----------- 1 file changed, 12 insertions(+), 16 deletions(-) diff --git a/cpp/tests/reductions/segmented_reduction_tests.cpp b/cpp/tests/reductions/segmented_reduction_tests.cpp index 42bd5064be2..8a9a8fb549e 100644 --- a/cpp/tests/reductions/segmented_reduction_tests.cpp +++ b/cpp/tests/reductions/segmented_reduction_tests.cpp @@ -398,6 +398,14 @@ TEST_F(SegmentedReductionTestUntyped, ReduceEmptyColumn) #define XXX "" // null placeholder struct SegmentedReductionStringTest : public cudf::test::BaseFixture { + std::pair> input() + { + return std::pair( + strings_column_wrapper{ + {"world", "cudf", XXX, "", "rapids", "i am", "ai", "apples", "zebras", XXX, XXX, XXX}, + {true, true, false, true, true, true, true, true, true, false, false, false}}, + fixed_width_column_wrapper{0, 1, 4, 7, 9, 9, 10, 12}); + } }; TEST_F(SegmentedReductionStringTest, MaxIncludeNulls) @@ -411,10 +419,7 @@ TEST_F(SegmentedReductionStringTest, MaxIncludeNulls) // outputs: {"world", XXX, "rapids", "zebras", XXX, XXX, XXX} // output nullmask: {1, 0, 1, 1, 0, 0, 0} - strings_column_wrapper input{ - {"world", "cudf", XXX, "", "rapids", "i am", "ai", "apples", "zebras", XXX, XXX, XXX}, - {true, true, false, true, true, true, true, true, true, false, false, false}}; - fixed_width_column_wrapper offsets{0, 1, 4, 7, 9, 9, 10, 12}; + auto const [input, offsets] = this->input(); data_type output_dtype{type_id::STRING}; strings_column_wrapper expect{{"world", XXX, "rapids", "zebras", XXX, XXX, XXX}, @@ -439,10 +444,7 @@ TEST_F(SegmentedReductionStringTest, MaxExcludeNulls) // outputs: {"world", "cudf", "rapids", "zebras", XXX, XXX, XXX} // output nullmask: {1, 1, 1, 1, 0, 0, 0} - strings_column_wrapper input{ - {"world", "cudf", XXX, "", "rapids", "i am", "ai", "apples", "zebras", XXX, XXX, XXX}, - {true, true, false, true, true, true, true, true, true, false, false, false}}; - fixed_width_column_wrapper offsets{0, 1, 4, 7, 9, 9, 10, 12}; + auto const [input, offsets] = this->input(); data_type output_dtype{type_id::STRING}; strings_column_wrapper expect{{"world", "cudf", "rapids", "zebras", XXX, XXX, XXX}, @@ -467,10 +469,7 @@ TEST_F(SegmentedReductionStringTest, MinIncludeNulls) // outputs: {"world", XXX, "ai", "apples", XXX, XXX, XXX} // output nullmask: {1, 0, 1, 1, 0, 0, 0} - strings_column_wrapper input{ - {"world", "cudf", XXX, "", "rapids", "i am", "ai", "apples", "zebras", XXX, XXX, XXX}, - {true, true, false, true, true, true, true, true, true, false, false, false}}; - fixed_width_column_wrapper offsets{0, 1, 4, 7, 9, 9, 10, 12}; + auto const [input, offsets] = this->input(); data_type output_dtype{type_id::STRING}; strings_column_wrapper expect{{"world", XXX, "ai", "apples", XXX, XXX, XXX}, @@ -495,10 +494,7 @@ TEST_F(SegmentedReductionStringTest, MinExcludeNulls) // outputs: {"world", "", "ai", "apples", XXX, XXX, XXX} // output nullmask: {1, 1, 1, 1, 0, 0, 0} - strings_column_wrapper input{ - {"world", "cudf", XXX, "", "rapids", "i am", "ai", "apples", "zebras", XXX, XXX, XXX}, - {true, true, false, true, true, true, true, true, true, false, false, false}}; - fixed_width_column_wrapper offsets{0, 1, 4, 7, 9, 9, 10, 12}; + auto const [input, offsets] = this->input(); data_type output_dtype{type_id::STRING}; strings_column_wrapper expect{{"world", "", "ai", "apples", XXX, XXX, XXX}, From 424f540a4a94f309e56d8694d67e88bd7b5e59f0 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Fri, 15 Apr 2022 10:59:56 -0700 Subject: [PATCH 19/20] Update cpp/include/cudf/detail/utilities/element_argminmax.cuh Co-authored-by: David Wendt <45795991+davidwendt@users.noreply.github.com> --- cpp/include/cudf/detail/utilities/element_argminmax.cuh | 4 ---- 1 file changed, 4 deletions(-) diff --git a/cpp/include/cudf/detail/utilities/element_argminmax.cuh b/cpp/include/cudf/detail/utilities/element_argminmax.cuh index 428b0740af5..f98aa9bcaaa 100644 --- a/cpp/include/cudf/detail/utilities/element_argminmax.cuh +++ b/cpp/include/cudf/detail/utilities/element_argminmax.cuh @@ -16,10 +16,6 @@ #pragma once -/** - * @brief Definition of the device operators - * @file - */ #include #include From 8f2a687dd7b4b6b0384906cc3d639fc5684bef78 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Fri, 15 Apr 2022 11:33:56 -0700 Subject: [PATCH 20/20] style fix --- cpp/include/cudf/detail/utilities/element_argminmax.cuh | 1 - 1 file changed, 1 deletion(-) diff --git a/cpp/include/cudf/detail/utilities/element_argminmax.cuh b/cpp/include/cudf/detail/utilities/element_argminmax.cuh index f98aa9bcaaa..45b56278dba 100644 --- a/cpp/include/cudf/detail/utilities/element_argminmax.cuh +++ b/cpp/include/cudf/detail/utilities/element_argminmax.cuh @@ -16,7 +16,6 @@ #pragma once - #include #include #include