diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index d553359b0769..f743bc022f8e 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -374,6 +374,13 @@ add_library( src/reductions/scan/scan.cpp src/reductions/scan/scan_exclusive.cu src/reductions/scan/scan_inclusive.cu + src/reductions/segmented_all.cu + src/reductions/segmented_any.cu + src/reductions/segmented_max.cu + src/reductions/segmented_min.cu + src/reductions/segmented_product.cu + src/reductions/segmented_reductions.cpp + src/reductions/segmented_sum.cu src/reductions/std.cu src/reductions/sum.cu src/reductions/sum_of_squares.cu diff --git a/cpp/benchmarks/CMakeLists.txt b/cpp/benchmarks/CMakeLists.txt index 054410c32656..67c5ba0b229b 100644 --- a/cpp/benchmarks/CMakeLists.txt +++ b/cpp/benchmarks/CMakeLists.txt @@ -24,7 +24,7 @@ target_compile_options( target_link_libraries( cudf_datagen PUBLIC GTest::gmock GTest::gtest GTest::gmock_main GTest::gtest_main - benchmark::benchmark nvbench::nvbench Threads::Threads cudf + benchmark::benchmark nvbench::nvbench Threads::Threads cudf cudftestutil ) target_include_directories( @@ -175,9 +175,10 @@ ConfigureBench(TYPE_DISPATCHER_BENCH type_dispatcher/type_dispatcher.cu) # ################################################################################################## # * reduction benchmark --------------------------------------------------------------------------- ConfigureBench( - REDUCTION_BENCH reduction/anyall.cpp reduction/dictionary.cpp reduction/reduce.cpp - reduction/scan.cpp reduction/minmax.cpp + REDUCTION_BENCH reduction/anyall.cpp reduction/dictionary.cpp reduction/minmax.cpp + reduction/reduce.cpp reduction/scan.cpp ) +ConfigureNVBench(REDUCTION_NVBENCH reduction/segment_reduce.cu) # ################################################################################################## # * reduction benchmark --------------------------------------------------------------------------- diff --git a/cpp/benchmarks/reduction/segment_reduce.cu b/cpp/benchmarks/reduction/segment_reduce.cu new file mode 100644 index 000000000000..47e47943d36b --- /dev/null +++ b/cpp/benchmarks/reduction/segment_reduce.cu @@ -0,0 +1,134 @@ +/* + * 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. + */ + +#include +#include + +#include +#include + +#include +#include +#include +#include +#include +#include + +#include + +#include + +#include +#include +#include + +namespace cudf { + +bool constexpr is_boolean_output_agg(segmented_reduce_aggregation::Kind kind) +{ + return kind == segmented_reduce_aggregation::ALL || kind == segmented_reduce_aggregation::ANY; +} + +template +std::unique_ptr make_simple_aggregation() +{ + switch (kind) { + case segmented_reduce_aggregation::SUM: + return make_sum_aggregation(); + case segmented_reduce_aggregation::PRODUCT: + return make_product_aggregation(); + case segmented_reduce_aggregation::MIN: + return make_min_aggregation(); + case segmented_reduce_aggregation::MAX: + return make_max_aggregation(); + case segmented_reduce_aggregation::ALL: + return make_all_aggregation(); + case segmented_reduce_aggregation::ANY: + return make_any_aggregation(); + default: CUDF_FAIL("Unsupported simple segmented aggregation"); + } +} + +template +std::pair, thrust::device_vector> make_test_data( + nvbench::state& state) +{ + auto const column_size{size_type(state.get_int64("column_size"))}; + auto const num_segments{size_type(state.get_int64("num_segments"))}; + + auto segment_length = column_size / num_segments; + + test::UniformRandomGenerator rand_gen(0, 100); + auto data_it = detail::make_counting_transform_iterator( + 0, [&rand_gen](auto i) { return rand_gen.generate(); }); + + auto offset_it = + detail::make_counting_transform_iterator(0, [&column_size, &segment_length](auto i) { + return column_size < i * segment_length ? column_size : i * segment_length; + }); + + test::fixed_width_column_wrapper input(data_it, data_it + column_size); + std::vector h_offsets(offset_it, offset_it + num_segments + 1); + thrust::device_vector d_offsets(h_offsets); + + return std::make_pair(input.release(), d_offsets); +} + +template +std::enable_if_t, void> +BM_Simple_Segmented_Reduction(nvbench::state& state, + nvbench::type_list>) +{ + // TODO: to be replaced by nvbench fixture once it's ready + cudf::rmm_pool_raii rmm_pool; + + auto const column_size{size_type(state.get_int64("column_size"))}; + auto [input, offsets] = make_test_data(state); + auto agg = make_simple_aggregation(); + + state.add_element_count(column_size); + state.add_global_memory_reads(column_size); + state.add_global_memory_writes(column_size); + + state.exec( + nvbench::exec_tag::sync, + [input_view = input->view(), offset_span = device_span{offsets}, &agg]( + nvbench::launch& launch) { + segmented_reduce( + input_view, offset_span, *agg, data_type{type_to_id()}, null_policy::INCLUDE); + }); +} + +template +std::enable_if_t, void> +BM_Simple_Segmented_Reduction(nvbench::state& state, + nvbench::type_list>) +{ + state.skip("Invalid combination of dtype and aggregation type."); +} + +using Types = nvbench::type_list; +// Skip benchmarking MAX/ANY since they are covered by MIN/ALL respectively. +using AggKinds = nvbench:: + enum_type_list; + +NVBENCH_BENCH_TYPES(BM_Simple_Segmented_Reduction, NVBENCH_TYPE_AXES(Types, Types, AggKinds)) + .set_name("segmented_reduction_simple") + .set_type_axes_names({"InputType", "OutputType", "AggregationKinds"}) + .add_int64_axis("column_size", {100'000, 1'000'000, 10'000'000, 100'000'000}) + .add_int64_axis("num_segments", {1'000, 10'000, 100'000}); + +} // namespace cudf diff --git a/cpp/include/cudf/aggregation.hpp b/cpp/include/cudf/aggregation.hpp index fb5b968671fc..2a4b86588cb8 100644 --- a/cpp/include/cudf/aggregation.hpp +++ b/cpp/include/cudf/aggregation.hpp @@ -148,6 +148,17 @@ class groupby_scan_aggregation : public virtual aggregation { groupby_scan_aggregation() {} }; +/** + * @brief Derived class intended for segmented reduction usage. + */ +class segmented_reduce_aggregation : public virtual aggregation { + public: + ~segmented_reduce_aggregation() override = default; + + protected: + segmented_reduce_aggregation() {} +}; + enum class udf_type : bool { CUDA, PTX }; enum class correlation_type : int32_t { PEARSON, KENDALL, SPEARMAN }; diff --git a/cpp/include/cudf/detail/aggregation/aggregation.hpp b/cpp/include/cudf/detail/aggregation/aggregation.hpp index 08ab9ad61314..cb30a54fe843 100644 --- a/cpp/include/cudf/detail/aggregation/aggregation.hpp +++ b/cpp/include/cudf/detail/aggregation/aggregation.hpp @@ -147,7 +147,8 @@ class aggregation_finalizer { // Declares the interface for the finalizer */ class sum_aggregation final : public rolling_aggregation, public groupby_aggregation, - public groupby_scan_aggregation { + public groupby_scan_aggregation, + public segmented_reduce_aggregation { public: sum_aggregation() : aggregation(SUM) {} @@ -166,7 +167,7 @@ class sum_aggregation final : public rolling_aggregation, /** * @brief Derived class for specifying a product aggregation */ -class product_aggregation final : public groupby_aggregation { +class product_aggregation final : public groupby_aggregation, public segmented_reduce_aggregation { public: product_aggregation() : aggregation(PRODUCT) {} @@ -187,7 +188,8 @@ class product_aggregation final : public groupby_aggregation { */ class min_aggregation final : public rolling_aggregation, public groupby_aggregation, - public groupby_scan_aggregation { + public groupby_scan_aggregation, + public segmented_reduce_aggregation { public: min_aggregation() : aggregation(MIN) {} @@ -208,7 +210,8 @@ class min_aggregation final : public rolling_aggregation, */ class max_aggregation final : public rolling_aggregation, public groupby_aggregation, - public groupby_scan_aggregation { + public groupby_scan_aggregation, + public segmented_reduce_aggregation { public: max_aggregation() : aggregation(MAX) {} @@ -248,7 +251,7 @@ class count_aggregation final : public rolling_aggregation, /** * @brief Derived class for specifying an any aggregation */ -class any_aggregation final : public aggregation { +class any_aggregation final : public segmented_reduce_aggregation { public: any_aggregation() : aggregation(ANY) {} @@ -267,7 +270,7 @@ class any_aggregation final : public aggregation { /** * @brief Derived class for specifying an all aggregation */ -class all_aggregation final : public aggregation { +class all_aggregation final : public segmented_reduce_aggregation { public: all_aggregation() : aggregation(ALL) {} diff --git a/cpp/include/cudf/detail/null_mask.cuh b/cpp/include/cudf/detail/null_mask.cuh index df06ad9e4f37..78eaa4f24483 100644 --- a/cpp/include/cudf/detail/null_mask.cuh +++ b/cpp/include/cudf/detail/null_mask.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021, NVIDIA CORPORATION. + * Copyright (c) 2021-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. @@ -18,6 +18,7 @@ #include #include #include +#include #include #include #include @@ -32,6 +33,7 @@ #include #include #include +#include #include #include @@ -279,7 +281,8 @@ rmm::device_uvector segmented_count_bits(bitmask_type const* bitmask, OffsetIterator first_bit_indices_end, OffsetIterator last_bit_indices_begin, count_bits_policy count_bits, - rmm::cuda_stream_view stream) + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) { auto const num_ranges = static_cast(std::distance(first_bit_indices_begin, first_bit_indices_end)); @@ -329,14 +332,15 @@ rmm::device_uvector segmented_count_bits(bitmask_type const* bitmask, // set bits from the length of the segment. auto segments_begin = thrust::make_zip_iterator(first_bit_indices_begin, last_bit_indices_begin); - auto segments_size = thrust::transform_iterator(segments_begin, [] __device__(auto segment) { - auto const begin = thrust::get<0>(segment); - auto const end = thrust::get<1>(segment); - return end - begin; - }); + auto segment_length_iterator = + thrust::transform_iterator(segments_begin, [] __device__(auto const& segment) { + auto const begin = thrust::get<0>(segment); + auto const end = thrust::get<1>(segment); + return end - begin; + }); thrust::transform(rmm::exec_policy(stream), - segments_size, - segments_size + num_ranges, + segment_length_iterator, + segment_length_iterator + num_ranges, d_bit_counts.data(), d_bit_counts.data(), [] __device__(auto segment_size, auto segment_bit_count) { @@ -438,7 +442,8 @@ std::vector segmented_count_bits(bitmask_type const* bitmask, first_bit_indices_end, last_bit_indices_begin, count_bits, - stream); + stream, + rmm::mr::get_current_device_resource()); // Copy the results back to the host. return make_std_vector_sync(d_bit_counts, stream); @@ -501,6 +506,80 @@ std::vector segmented_null_count(bitmask_type const* bitmask, return detail::segmented_count_unset_bits(bitmask, indices_begin, indices_end, stream); } +/** + * @brief Create an output null mask whose validity is determined by the + * validity of any/all elements of segments of an input null mask. + * + * @tparam OffsetIterator Random-access input iterator type. + * @param bitmask Null mask residing in device memory whose segments will be + * reduced into a new mask. + * @param first_bit_indices_begin Random-access input iterator to the beginning + * of a sequence of indices of the first bit in each segment (inclusive). + * @param first_bit_indices_end Random-access input iterator to the end of a + * sequence of indices of the first bit in each segment (inclusive). + * @param last_bit_indices_begin Random-access input iterator to the beginning + * of a sequence of indices of the last bit in each segment (exclusive). + * @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 reduction is valid if any element in the segment is valid. + * @param stream CUDA stream used for device memory operations and kernel launches. + * @param mr Device memory resource used to allocate the returned buffer's device memory. + * @return A pair containing the reduced null mask and number of nulls. + */ +template +std::pair segmented_null_mask_reduction( + bitmask_type const* bitmask, + OffsetIterator first_bit_indices_begin, + OffsetIterator first_bit_indices_end, + OffsetIterator last_bit_indices_begin, + null_policy null_handling, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + auto const segments_begin = + thrust::make_zip_iterator(first_bit_indices_begin, last_bit_indices_begin); + auto const segment_length_iterator = + thrust::make_transform_iterator(segments_begin, [] __device__(auto const& segment) { + auto const begin = thrust::get<0>(segment); + auto const end = thrust::get<1>(segment); + return end - begin; + }); + + auto const num_segments = + static_cast(std::distance(first_bit_indices_begin, first_bit_indices_end)); + + if (bitmask == nullptr) { + return cudf::detail::valid_if( + segment_length_iterator, + segment_length_iterator + num_segments, + [] __device__(auto const& length) { return length > 0; }, + stream, + mr); + } + + auto const segment_valid_counts = + cudf::detail::segmented_count_bits(bitmask, + first_bit_indices_begin, + first_bit_indices_end, + last_bit_indices_begin, + cudf::detail::count_bits_policy::SET_BITS, + stream, + rmm::mr::get_current_device_resource()); + auto const length_and_valid_count = + thrust::make_zip_iterator(segment_length_iterator, segment_valid_counts.begin()); + return cudf::detail::valid_if( + length_and_valid_count, + length_and_valid_count + num_segments, + [null_handling] __device__(auto const& length_and_valid_count) { + auto const length = thrust::get<0>(length_and_valid_count); + auto const valid_count = thrust::get<1>(length_and_valid_count); + return (length > 0) and + ((null_handling == null_policy::EXCLUDE) ? valid_count > 0 : valid_count == length); + }, + stream, + mr); +} + } // namespace detail } // namespace cudf diff --git a/cpp/include/cudf/detail/reduction.cuh b/cpp/include/cudf/detail/reduction.cuh index e176529ed6db..76afbf7e4b8f 100644 --- a/cpp/include/cudf/detail/reduction.cuh +++ b/cpp/include/cudf/detail/reduction.cuh @@ -18,6 +18,7 @@ #include "reduction_operators.cuh" +#include #include #include @@ -26,6 +27,7 @@ #include #include +#include #include #include @@ -40,6 +42,8 @@ namespace detail { * @param[in] num_items the number of items * @param[in] op 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 scalar's device + * memory * @returns Output scalar in device memory * * @tparam Op the reduction operator with device binary operator @@ -102,7 +106,6 @@ std::unique_ptr reduce(InputIterator d_in, CUDF_FAIL( "This function should never be called. fixed_point reduce should always go through the reduce " "for the corresponding device_storage_type_t"); - ; } // @brief string_view specialization of simple reduction @@ -157,6 +160,8 @@ std::unique_ptr reduce(InputIterator d_in, * @param[in] valid_count the intermediate operator argument 1 * @param[in] ddof the intermediate operator argument 2 * @param[in] stream CUDA stream used for device memory operations and kernel launches. + * @param[in] mr Device memory resource used to allocate the returned scalar's device + * memory * @returns Output scalar in device memory * * The reduction operator must have `intermediate::compute_result()` method. @@ -218,6 +223,92 @@ std::unique_ptr reduce(InputIterator d_in, return std::unique_ptr(result); } +/** + * @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 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] 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> +std::unique_ptr segmented_reduce(InputIterator d_in, + OffsetIterator d_offset, + cudf::size_type num_segments, + op::simple_op sop, + 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(); + + // 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 dev_result; +} + +template ::type, + typename std::enable_if_t() || + is_fixed_point()>* = nullptr> +std::unique_ptr segmented_reduce(InputIterator, + OffsetIterator, + cudf::size_type, + op::simple_op, + rmm::cuda_stream_view, + rmm::mr::device_memory_resource*) +{ + CUDF_FAIL( + "Unsupported data types called on segmented_reduce. Only numeric and chrono types are " + "supported."); +} + } // namespace detail } // namespace reduction } // namespace cudf diff --git a/cpp/include/cudf/detail/reduction_functions.hpp b/cpp/include/cudf/detail/reduction_functions.hpp index d8f23e8d7cb3..ccec4bf8a6c5 100644 --- a/cpp/include/cudf/detail/reduction_functions.hpp +++ b/cpp/include/cudf/detail/reduction_functions.hpp @@ -30,7 +30,7 @@ namespace reduction { * If all elements in input column are null, output scalar is null. * * @throw cudf::logic_error if input column type is not convertible to `output_dtype` - * @throw cudf::logic_error if `output_dtype` is not arithmetic point type + * @throw cudf::logic_error if `output_dtype` is not an arithmetic type * * @param col input column to compute sum * @param output_dtype data type of return type and typecast elements of input column @@ -128,7 +128,7 @@ std::unique_ptr all( * If all elements in input column are null, output scalar is null. * * @throw cudf::logic_error if input column type is not convertible to `output_dtype` - * @throw cudf::logic_error if `output_dtype` is not arithmetic point type + * @throw cudf::logic_error if `output_dtype` is not an arithmetic type * * @param col input column to compute product. * @param output_dtype data type of return type and typecast elements of input column @@ -148,7 +148,7 @@ std::unique_ptr product( * If all elements in input column are null, output scalar is null. * * @throw cudf::logic_error if input column type is not convertible to `output_dtype` - * @throw cudf::logic_error if `output_dtype` is not arithmetic point type + * @throw cudf::logic_error if `output_dtype` is not an arithmetic type * * @param col input column to compute sum of squares. * @param output_dtype data type of return type and typecast elements of input column @@ -245,7 +245,7 @@ std::unique_ptr standard_deviation( * @param n index of element to get * @param null_handling Indicates if null values will be counted while indexing. * @param stream CUDA stream used for device memory operations and kernel launches. - * @param mr Device memory resource used to allocate the returned scalar's device memory + * @param mr Device memory resource used to allocate the returned scalar's device memory. * @return nth element as scalar */ std::unique_ptr nth_element( @@ -319,5 +319,159 @@ std::unique_ptr merge_sets( rmm::cuda_stream_view stream = rmm::cuda_stream_default, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +/** + * @brief Compute sum of each segment in input column. + * + * If an input segment is empty, the segment result is null. + * + * @throw cudf::logic_error if input column type is not convertible to `output_dtype`. + * @throw cudf::logic_error if `output_dtype` is not an arithmetic type. + * + * @param col Input column to compute sum. + * @param offsets Indices to identify segment boundaries. + * @param output_dtype Data type of return type and typecast elements of input column. + * @param null_handling If `INCLUDE`, the reduction is valid if all elements in + * a segment are valid, otherwise null. If `EXCLUDE`, the reduction is valid if + * any element in the segment is valid, otherwise null. + * @param stream CUDA stream used for device memory operations and kernel launches. + * @param mr Device memory resource used to allocate the returned column's device memory. + * @return Sums of segments in type `output_dtype`. + */ +std::unique_ptr segmented_sum( + column_view const& col, + device_span offsets, + data_type const output_dtype, + null_policy null_handling, + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + +/** + * @brief Computes product of each segment in input column. + * + * If an input segment is empty, the segment result is null. + * + * @throw cudf::logic_error if input column type is not convertible to `output_dtype`. + * @throw cudf::logic_error if `output_dtype` is not an arithmetic type. + * + * @param col Input column to compute product. + * @param offsets Indices to identify segment boundaries. + * @param output_dtype data type of return type and typecast elements of input column. + * @param null_handling If `INCLUDE`, the reduction is valid if all elements in + * a segment are valid, otherwise null. If `EXCLUDE`, the reduction is valid if + * any element in the segment is valid, otherwise null. + * @param stream CUDA stream used for device memory operations and kernel launches. + * @param mr Device memory resource used to allocate the returned scalar's device memory. + * @return Product as scalar of type `output_dtype`. + */ +std::unique_ptr segmented_product( + column_view const& col, + device_span offsets, + data_type const output_dtype, + null_policy null_handling, + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + +/** + * @brief Compute minimum of each segment in input column. + * + * If an input segment is empty, the segment result is null. + * + * @throw cudf::logic_error if input column type is convertible to `output_dtype`. + * + * @param col Input column to compute minimum. + * @param offsets Indices to identify segment boundaries. + * @param output_dtype Data type of return type and typecast elements of input column. + * @param null_handling If `INCLUDE`, the reduction is valid if all elements in + * a segment are valid, otherwise null. If `EXCLUDE`, the reduction is valid if + * any element in the segment is valid, otherwise null. + * @param stream CUDA stream used for device memory operations and kernel launches. + * @param mr Device memory resource used to allocate the returned scalar's device memory. + * @return Minimums of segments in type `output_dtype`. + */ +std::unique_ptr segmented_min( + column_view const& col, + device_span offsets, + data_type const output_dtype, + null_policy null_handling, + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + +/** + * @brief Compute maximum of each segment in input column. + * + * If an input segment is empty, the segment result is null. + * + * @throw cudf::logic_error if input column type is convertible to `output_dtype`. + * + * @param col Input column to compute maximum. + * @param offsets Indices to identify segment boundaries. + * @param output_dtype Data type of return type and typecast elements of input column. + * @param null_handling If `INCLUDE`, the reduction is valid if all elements in + * a segment are valid, otherwise null. If `EXCLUDE`, the reduction is valid if + * any element in the segment is valid, otherwise null. + * @param stream CUDA stream used for device memory operations and kernel launches. + * @param mr Device memory resource used to allocate the returned scalar's device memory. + * @return Maximums of segments in type `output_dtype`. + */ +std::unique_ptr segmented_max( + column_view const& col, + device_span offsets, + data_type const output_dtype, + null_policy null_handling, + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + +/** + * @brief Compute if any of the values in the segment are true when typecasted to bool. + * + * If an input segment is empty, the segment result is null. + * + * @throw cudf::logic_error if input column type is not convertible to bool. + * @throw cudf::logic_error if `output_dtype` is not bool8. + * + * @param col Input column to compute any_of. + * @param offsets Indices to identify segment boundaries. + * @param output_dtype Data type of return type and typecast elements of input column. + * @param null_handling If `INCLUDE`, the reduction is valid if all elements in + * a segment are valid, otherwise null. If `EXCLUDE`, the reduction is valid if + * any element in the segment is valid, otherwise null. + * @param stream CUDA stream used for device memory operations and kernel launches. + * @param mr Device memory resource used to allocate the returned scalar's device memory. + * @return Column of bool8 for the results of the segments. + */ +std::unique_ptr segmented_any( + column_view const& col, + device_span offsets, + data_type const output_dtype, + null_policy null_handling, + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + +/** + * @brief Compute if all of the values in the segment are true when typecasted to bool. + * + * If an input segment is empty, the segment result is null. + * + * @throw cudf::logic_error if input column type is not convertible to bool. + * @throw cudf::logic_error if `output_dtype` is not bool8. + * + * @param col Input column to compute all_of. + * @param offsets Indices to identify segment boundaries. + * @param output_dtype Data type of return type and typecast elements of input column. + * @param null_handling If `INCLUDE`, the reduction is valid if all elements in + * a segment are valid, otherwise null. If `EXCLUDE`, the reduction is valid if + * any element in the segment is valid, otherwise null. + * @param stream CUDA stream used for device memory operations and kernel launches. + * @param mr Device memory resource used to allocate the returned scalar's device memory. + * @return Column of bool8 for the results of the segments. + */ +std::unique_ptr segmented_all( + column_view const& col, + device_span offsets, + data_type const output_dtype, + null_policy null_handling, + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + } // namespace reduction } // namespace cudf diff --git a/cpp/include/cudf/detail/valid_if.cuh b/cpp/include/cudf/detail/valid_if.cuh index 4a7e9b89c802..ee9e4b2c6871 100644 --- a/cpp/include/cudf/detail/valid_if.cuh +++ b/cpp/include/cudf/detail/valid_if.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * Copyright (c) 2019-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. @@ -64,7 +64,7 @@ __global__ void valid_if_kernel( size_type block_count = single_lane_block_sum_reduce(warp_valid_count); if (threadIdx.x == 0) { atomicAdd(valid_count, block_count); } -} // namespace detail +} /** * @brief Generate a bitmask where every bit is set for which a predicate is diff --git a/cpp/include/cudf/reduction.hpp b/cpp/include/cudf/reduction.hpp index d094118293b0..367814cda8e1 100644 --- a/cpp/include/cudf/reduction.hpp +++ b/cpp/include/cudf/reduction.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * Copyright (c) 2019-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. @@ -48,18 +48,21 @@ enum class scan_type : bool { INCLUSIVE, EXCLUSIVE }; * output data type. * @throw cudf::logic_error if `min` or `max` reduction is called and the * output type does not match the input column data type. + * @throw cudf::logic_error if `any` or `all` reduction is called and the + * output type is not bool8. + * @throw cudf::logic_error if `mean`, `var`, or `std` reduction is called and + * the output type is not floating point. * * If the input column has arithmetic type, output_dtype can be any arithmetic - * type. For `mean`, `var` and `std` ops, a floating point output type must be - * specified. If the input column has non-arithmetic type - * eg.(timestamp, string...), the same type must be specified. + * type. If the input column has non-arithmetic type, e.g. timestamp or string, + * the same output type must be specified. * * If the reduction fails, the member is_valid of the output scalar * will contain `false`. * * @param col Input column view * @param agg Aggregation operator applied by the reduction - * @param output_dtype The computation and output precision. + * @param output_dtype The computation and output precision. * @param mr Device memory resource used to allocate the returned scalar's device memory * @returns Output scalar with reduce result. */ @@ -69,6 +72,56 @@ std::unique_ptr reduce( data_type output_dtype, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +/** + * @brief Compute reduction of each segment in the input column + * + * This function does not detect overflows in reductions. When given integral and + * floating point inputs, their values are promoted to `int64_t` and `double` + * respectively to compute, and casted to @p output_dtype before returning. + * + * Null values are treated as identities during reduction. + * + * If the segment is empty, the row corresponding to the result of the + * segment is null. + * + * If any index in @p offsets is out of bound of @p segmented_values , the behavior + * is undefined. + * + * @note If the input column has arithmetic type, output_dtype can be any arithmetic + * type. If the input column has non-arithmetic type, e.g. timestamp, the same + * output type must be specified. + * + * @note If input is not empty, the result is always nullable. + * + * @throw cudf::logic_error if reduction is called for non-arithmetic output + * type and operator other than `min` and `max`. + * @throw cudf::logic_error if input column data type is not convertible to + * output data type. + * @throw cudf::logic_error if `min` or `max` reduction is called and the + * output type does not match the input column data type. + * @throw cudf::logic_error if `any` or `all` reduction is called and the + * output type is not bool8. + * + * @param segmented_values Column view of segmented inputs. + * @param offsets Each segment's offset of @p segmented_values. A list of offsets + * with size `num_segments + 1`. The size of `i`th segment is `offsets[i+1] - + * offsets[i]`. + * @param agg Aggregation operator applied by the reduction. + * @param output_dtype The output precision. + * @param null_handling If `INCLUDE`, the reduction is valid if all elements in + * a segment are valid, otherwise null. If `EXCLUDE`, the reduction is valid if + * any element in the segment is valid, otherwise null. + * @param mr Device memory resource used to allocate the returned scalar's device memory + * @returns Output column with results of segmented reduction. + */ +std::unique_ptr segmented_reduce( + column_view const& segmented_values, + device_span offsets, + segmented_reduce_aggregation const& agg, + data_type output_dtype, + null_policy null_handling, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + /** * @brief Computes the scan of a column. * diff --git a/cpp/src/aggregation/aggregation.cpp b/cpp/src/aggregation/aggregation.cpp index 8e2a167f7b2f..68fe6d86103e 100644 --- a/cpp/src/aggregation/aggregation.cpp +++ b/cpp/src/aggregation/aggregation.cpp @@ -417,6 +417,8 @@ template std::unique_ptr make_sum_aggregation(); template std::unique_ptr make_sum_aggregation(); template std::unique_ptr make_sum_aggregation(); template std::unique_ptr make_sum_aggregation(); +template std::unique_ptr +make_sum_aggregation(); /// Factory to create a PRODUCT aggregation template @@ -426,6 +428,8 @@ std::unique_ptr make_product_aggregation() } template std::unique_ptr make_product_aggregation(); template std::unique_ptr make_product_aggregation(); +template std::unique_ptr +make_product_aggregation(); /// Factory to create a MIN aggregation template @@ -437,6 +441,8 @@ template std::unique_ptr make_min_aggregation(); template std::unique_ptr make_min_aggregation(); template std::unique_ptr make_min_aggregation(); template std::unique_ptr make_min_aggregation(); +template std::unique_ptr +make_min_aggregation(); /// Factory to create a MAX aggregation template @@ -448,6 +454,8 @@ template std::unique_ptr make_max_aggregation(); template std::unique_ptr make_max_aggregation(); template std::unique_ptr make_max_aggregation(); template std::unique_ptr make_max_aggregation(); +template std::unique_ptr +make_max_aggregation(); /// Factory to create a COUNT aggregation template @@ -473,6 +481,8 @@ std::unique_ptr make_any_aggregation() return std::make_unique(); } template std::unique_ptr make_any_aggregation(); +template std::unique_ptr +make_any_aggregation(); /// Factory to create a ALL aggregation template @@ -481,6 +491,8 @@ std::unique_ptr make_all_aggregation() return std::make_unique(); } template std::unique_ptr make_all_aggregation(); +template std::unique_ptr +make_all_aggregation(); /// Factory to create a SUM_OF_SQUARES aggregation template diff --git a/cpp/src/reductions/segmented_all.cu b/cpp/src/reductions/segmented_all.cu new file mode 100644 index 000000000000..a04da1ac2fa3 --- /dev/null +++ b/cpp/src/reductions/segmented_all.cu @@ -0,0 +1,46 @@ +/* + * 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. + */ + +#include "simple_segmented.cuh" + +#include + +namespace cudf { +namespace reduction { + +std::unique_ptr segmented_all(column_view const& col, + device_span offsets, + cudf::data_type const output_dtype, + null_policy null_handling, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + CUDF_EXPECTS(output_dtype == cudf::data_type(cudf::type_id::BOOL8), + "segmented_all() operation requires output type `BOOL8`"); + + // A minimum over bool types is used to implement all() + return cudf::type_dispatcher( + col.type(), + simple::detail::bool_result_column_dispatcher{}, + col, + offsets, + null_handling, + stream, + mr); +} + +} // namespace reduction +} // namespace cudf diff --git a/cpp/src/reductions/segmented_any.cu b/cpp/src/reductions/segmented_any.cu new file mode 100644 index 000000000000..ad44289175b0 --- /dev/null +++ b/cpp/src/reductions/segmented_any.cu @@ -0,0 +1,46 @@ +/* + * 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. + */ + +#include "simple_segmented.cuh" + +#include + +namespace cudf { +namespace reduction { + +std::unique_ptr segmented_any(column_view const& col, + device_span offsets, + cudf::data_type const output_dtype, + null_policy null_handling, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + CUDF_EXPECTS(output_dtype == cudf::data_type(cudf::type_id::BOOL8), + "segmented_any() operation requires output type `BOOL8`"); + + // A maximum over bool types is used to implement any() + return cudf::type_dispatcher( + col.type(), + simple::detail::bool_result_column_dispatcher{}, + col, + offsets, + null_handling, + stream, + mr); +} + +} // namespace reduction +} // namespace cudf diff --git a/cpp/src/reductions/segmented_max.cu b/cpp/src/reductions/segmented_max.cu new file mode 100644 index 000000000000..198960643434 --- /dev/null +++ b/cpp/src/reductions/segmented_max.cu @@ -0,0 +1,44 @@ +/* + * 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. + */ + +#include "simple_segmented.cuh" + +#include + +namespace cudf { +namespace reduction { + +std::unique_ptr segmented_max(column_view const& col, + device_span offsets, + cudf::data_type const output_dtype, + null_policy null_handling, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + CUDF_EXPECTS(col.type() == output_dtype, + "segmented_max() operation requires matching output type"); + return cudf::type_dispatcher( + col.type(), + simple::detail::same_column_type_dispatcher{}, + col, + offsets, + null_handling, + stream, + mr); +} + +} // namespace reduction +} // namespace cudf diff --git a/cpp/src/reductions/segmented_min.cu b/cpp/src/reductions/segmented_min.cu new file mode 100644 index 000000000000..5c880f45bf84 --- /dev/null +++ b/cpp/src/reductions/segmented_min.cu @@ -0,0 +1,44 @@ +/* + * 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. + */ + +#include "simple_segmented.cuh" + +#include + +namespace cudf { +namespace reduction { + +std::unique_ptr segmented_min(column_view const& col, + device_span offsets, + data_type const output_dtype, + null_policy null_handling, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + CUDF_EXPECTS(col.type() == output_dtype, + "segmented_min() operation requires matching output type"); + return cudf::type_dispatcher( + col.type(), + simple::detail::same_column_type_dispatcher{}, + col, + offsets, + null_handling, + stream, + mr); +} + +} // namespace reduction +} // namespace cudf diff --git a/cpp/src/reductions/segmented_product.cu b/cpp/src/reductions/segmented_product.cu new file mode 100644 index 000000000000..1b8528708208 --- /dev/null +++ b/cpp/src/reductions/segmented_product.cu @@ -0,0 +1,43 @@ +/* + * 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. + */ + +#include "simple_segmented.cuh" + +#include + +namespace cudf { +namespace reduction { + +std::unique_ptr segmented_product(column_view const& col, + device_span offsets, + cudf::data_type const output_dtype, + null_policy null_handling, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + return cudf::type_dispatcher( + col.type(), + simple::detail::column_type_dispatcher{}, + col, + offsets, + output_dtype, + null_handling, + stream, + mr); +} + +} // namespace reduction +} // namespace cudf diff --git a/cpp/src/reductions/segmented_reductions.cpp b/cpp/src/reductions/segmented_reductions.cpp new file mode 100644 index 000000000000..415f5ae488ec --- /dev/null +++ b/cpp/src/reductions/segmented_reductions.cpp @@ -0,0 +1,105 @@ +/* + * 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. + */ + +#include +#include +#include +#include +#include +#include +#include + +#include + +namespace cudf { +namespace detail { +struct segmented_reduce_dispatch_functor { + column_view const& col; + device_span offsets; + data_type output_dtype; + null_policy null_handling; + rmm::mr::device_memory_resource* mr; + rmm::cuda_stream_view stream; + + segmented_reduce_dispatch_functor(column_view const& segmented_values, + device_span offsets, + data_type output_dtype, + null_policy null_handling, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) + : col(segmented_values), + offsets(offsets), + output_dtype(output_dtype), + null_handling(null_handling), + mr(mr), + stream(stream) + { + } + + template + std::unique_ptr operator()() + { + switch (k) { + case segmented_reduce_aggregation::SUM: + return reduction::segmented_sum(col, offsets, output_dtype, null_handling, stream, mr); + case segmented_reduce_aggregation::PRODUCT: + return reduction::segmented_product(col, offsets, output_dtype, null_handling, stream, mr); + case segmented_reduce_aggregation::MIN: + return reduction::segmented_min(col, offsets, output_dtype, null_handling, stream, mr); + case segmented_reduce_aggregation::MAX: + return reduction::segmented_max(col, offsets, output_dtype, null_handling, stream, mr); + case segmented_reduce_aggregation::ANY: + return reduction::segmented_any(col, offsets, output_dtype, null_handling, stream, mr); + case segmented_reduce_aggregation::ALL: + return reduction::segmented_all(col, offsets, output_dtype, null_handling, stream, mr); + default: + CUDF_FAIL("Unsupported aggregation type."); + // TODO: Add support for compound_ops + } + } +}; + +std::unique_ptr segmented_reduce(column_view const& segmented_values, + device_span offsets, + segmented_reduce_aggregation const& agg, + data_type output_dtype, + null_policy null_handling, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + CUDF_EXPECTS(offsets.size() > 0, "`offsets` should have at least 1 element."); + if (segmented_values.is_empty()) { return empty_like(segmented_values); } + + return aggregation_dispatcher( + agg.kind, + segmented_reduce_dispatch_functor{ + segmented_values, offsets, output_dtype, null_handling, stream, mr}); +} +} // namespace detail + +std::unique_ptr segmented_reduce(column_view const& segmented_values, + device_span offsets, + segmented_reduce_aggregation const& agg, + data_type output_dtype, + null_policy null_handling, + rmm::mr::device_memory_resource* mr) +{ + CUDF_FUNC_RANGE(); + return detail::segmented_reduce( + segmented_values, offsets, agg, output_dtype, null_handling, rmm::cuda_stream_default, mr); +} + +} // namespace cudf diff --git a/cpp/src/reductions/segmented_sum.cu b/cpp/src/reductions/segmented_sum.cu new file mode 100644 index 000000000000..f2deeddbcbb7 --- /dev/null +++ b/cpp/src/reductions/segmented_sum.cu @@ -0,0 +1,42 @@ +/* + * 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. + */ + +#include "simple_segmented.cuh" + +#include + +namespace cudf { +namespace reduction { + +std::unique_ptr segmented_sum(column_view const& col, + device_span offsets, + cudf::data_type const output_dtype, + null_policy null_handling, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + return cudf::type_dispatcher(col.type(), + simple::detail::column_type_dispatcher{}, + col, + offsets, + output_dtype, + null_handling, + stream, + mr); +} + +} // namespace reduction +} // namespace cudf diff --git a/cpp/src/reductions/simple_segmented.cuh b/cpp/src/reductions/simple_segmented.cuh new file mode 100644 index 000000000000..99837e67398a --- /dev/null +++ b/cpp/src/reductions/simple_segmented.cuh @@ -0,0 +1,274 @@ +/* + * 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 + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +#include +#include + +namespace cudf { +namespace reduction { +namespace simple { +namespace detail { + +/** + * @brief Segment reduction for 'sum', 'product', 'min', 'max', 'sum of squares' + * which directly compute the reduction by a single step reduction call. + * + * @tparam InputType the input column data-type + * @tparam ResultType the output 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::unique_ptr 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) +{ + // TODO: Rewrites this function to accept a pair of iterators for start/end indices + // to enable `2N` type offset input. + // reduction by iterator + auto dcol = cudf::column_device_view::create(col, stream); + auto simple_op = Op{}; + size_type num_segments = offsets.size() - 1; + + // 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); + } 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); + } + }(); + + // Compute the output null mask + auto const bitmask = col.null_mask(); + auto const first_bit_indices_begin = offsets.begin(); + auto const first_bit_indices_end = offsets.end() - 1; + auto const last_bit_indices_begin = first_bit_indices_begin + 1; + auto const [output_null_mask, output_null_count] = + cudf::detail::segmented_null_mask_reduction(bitmask, + first_bit_indices_begin, + first_bit_indices_end, + last_bit_indices_begin, + null_handling, + stream, + mr); + result->set_null_mask(output_null_mask, output_null_count, stream); + + return result; +} + +/** + * @brief Call reduce and return a column of type bool. + * + * This is used by operations `any()` and `all()`. + * + * @tparam Op The reduce operation to execute on the column. + */ +template +struct bool_result_column_dispatcher { + template ()>* = nullptr> + 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 simple_segmented_reduction( + col, offsets, null_handling, stream, mr); + } + + template ()>* = nullptr> + std::unique_ptr operator()(column_view const&, + device_span, + null_policy, + rmm::cuda_stream_view, + rmm::mr::device_memory_resource*) + { + CUDF_FAIL("Reduction operator not supported for this type"); + } +}; + +/** + * @brief Call reduce and return a column of type matching the input column. + * + * This is used by operations `min()` and `max()`. + * + * @tparam Op The reduce operation to execute on the column. + */ +template +struct same_column_type_dispatcher { + private: + template + static constexpr bool is_supported() + { + return !(cudf::is_fixed_point() || cudf::is_dictionary() || + std::is_same_v || + std::is_same_v); + } + + public: + template ()>* = nullptr> + 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 simple_segmented_reduction( + col, offsets, null_handling, stream, mr); + } + + template ()>* = nullptr> + std::unique_ptr operator()(column_view const&, + device_span, + null_policy, + rmm::cuda_stream_view, + rmm::mr::device_memory_resource*) + { + CUDF_FAIL("Reduction operator not supported for this type"); + } +}; + +/** + * @brief Call reduce and return a column of the type specified. + * + * This is used by operations sum(), product(), and sum_of_squares(). + * It only supports numeric types. If the output type is not the + * same as the input type, an extra cast operation may occur. + * + * @tparam Op The reduce operation to execute on the column. + */ +template +struct column_type_dispatcher { + /** + * @brief Specialization for reducing floating-point column types to any output type. + */ + template ::value>* = nullptr> + std::unique_ptr reduce_numeric(column_view const& col, + device_span offsets, + data_type const output_type, + null_policy null_handling, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) + { + // TODO: per gh-9988, we should change the compute precision to `output_type`. + auto result = + simple_segmented_reduction(col, offsets, null_handling, stream, mr); + if (output_type == result->type()) { return result; } + return cudf::detail::cast(*result, output_type, stream, mr); + } + + /** + * @brief Specialization for reducing integer column types to any output type. + */ + template ::value>* = nullptr> + std::unique_ptr reduce_numeric(column_view const& col, + device_span offsets, + data_type const output_type, + null_policy null_handling, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) + { + // TODO: per gh-9988, we should change the compute precision to `output_type`. + auto result = + simple_segmented_reduction(col, offsets, null_handling, stream, mr); + if (output_type == result->type()) { return result; } + return cudf::detail::cast(*result, output_type, stream, mr); + } + + /** + * @brief Called by the type-dispatcher to reduce the input column `col` using + * the `Op` operation. + * + * @tparam ElementType The input column type or key type. + * @param col Input column (must be numeric) + * @param offsets Indices to segment boundaries + * @param output_type Requested type of the scalar result + * @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 CUDA stream used for device memory operations and kernel launches. + * @param mr Device memory resource used to allocate the returned scalar's device memory + */ + template ()>* = nullptr> + std::unique_ptr operator()(column_view const& col, + device_span offsets, + data_type const output_type, + null_policy null_handling, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) + { + if (output_type.id() == cudf::type_to_id()) { + return simple_segmented_reduction( + col, offsets, null_handling, stream, mr); + } + // reduce and map to output type + return reduce_numeric(col, offsets, output_type, null_handling, stream, mr); + } + + template ()>* = nullptr> + std::unique_ptr operator()(column_view const&, + device_span, + data_type const, + null_policy, + rmm::cuda_stream_view, + rmm::mr::device_memory_resource*) + { + CUDF_FAIL("Reduction operator not supported for this type"); + } +}; + +} // namespace detail +} // namespace simple +} // namespace reduction +} // namespace cudf diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index b8ea8b301c70..05b90095562f 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -151,7 +151,7 @@ ConfigureTest( # * reduction tests ------------------------------------------------------------------------------- ConfigureTest( REDUCTION_TEST reductions/collect_ops_tests.cpp reductions/rank_tests.cpp - reductions/reduction_tests.cpp reductions/scan_tests.cpp + reductions/reduction_tests.cpp reductions/scan_tests.cpp reductions/segmented_reduction_tests.cpp ) # ################################################################################################## diff --git a/cpp/tests/reductions/segmented_reduction_tests.cpp b/cpp/tests/reductions/segmented_reduction_tests.cpp new file mode 100644 index 000000000000..3a432cce801e --- /dev/null +++ b/cpp/tests/reductions/segmented_reduction_tests.cpp @@ -0,0 +1,391 @@ +/* + * 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. + */ + +#include +#include +#include + +#include +#include +#include + +#include + +namespace cudf { +namespace test { + +#define XXX 0 // null placeholder + +template +struct SegmentedReductionTest : public cudf::test::BaseFixture { +}; + +struct SegmentedReductionTestUntyped : public cudf::test::BaseFixture { +}; + +TYPED_TEST_CASE(SegmentedReductionTest, NumericTypes); + +TYPED_TEST(SegmentedReductionTest, SumExcludeNulls) +{ + // [1, 2, 3], [1, null, 3], [1], [null], [null, null], [] + // values: {1, 2, 3, 1, XXX, 3, 1, XXX, XXX, XXX} + // offsets: {0, 3, 6, 7, 8, 10, 10} + // nullmask: {1, 1, 1, 1, 0, 1, 1, 0, 0, 0} + // outputs: {6, 4, 1, XXX, XXX, XXX} + // output nullmask: {1, 1, 1, 0, 0, 0} + auto input = fixed_width_column_wrapper{{1, 2, 3, 1, XXX, 3, 1, XXX, XXX, XXX}, + {1, 1, 1, 1, 0, 1, 1, 0, 0, 0}}; + auto offsets = std::vector{0, 3, 6, 7, 8, 10, 10}; + auto d_offsets = thrust::device_vector(offsets); + auto expect = fixed_width_column_wrapper{{6, 4, 1, XXX, XXX, XXX}, {1, 1, 1, 0, 0, 0}}; + + auto res = segmented_reduce(input, + d_offsets, + *make_sum_aggregation(), + data_type{type_to_id()}, + null_policy::EXCLUDE); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*res, expect); +} + +TYPED_TEST(SegmentedReductionTest, ProductExcludeNulls) +{ + // [1, 3, 5], [null, 3, 5], [1], [null], [null, null], [] + // values: {1, 3, 5, XXX, 3, 5, 1, XXX, XXX, XXX} + // offsets: {0, 3, 6, 7, 8, 10, 10} + // nullmask: {1, 1, 1, 0, 1, 1, 1, 0, 0, 0} + // outputs: {15, 15, 1, XXX, XXX, XXX} + // output nullmask: {1, 1, 1, 0, 0, 0} + auto input = fixed_width_column_wrapper{{1, 3, 5, XXX, 3, 5, 1, XXX, XXX, XXX}, + {1, 1, 1, 0, 1, 1, 1, 0, 0, 0}}; + auto offsets = std::vector{0, 3, 6, 7, 8, 10, 10}; + auto d_offsets = thrust::device_vector(offsets); + auto expect = + fixed_width_column_wrapper{{15, 15, 1, XXX, XXX, XXX}, {1, 1, 1, 0, 0, 0}}; + + auto res = segmented_reduce(input, + d_offsets, + *make_product_aggregation(), + data_type{type_to_id()}, + null_policy::EXCLUDE); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*res, expect); +} + +TYPED_TEST(SegmentedReductionTest, MaxExcludeNulls) +{ + // [1, 2, 3], [1, null, 3], [1], [null], [null, null], [] + // values: {1, 2, 3, 1, XXX, 3, 1, XXX, XXX, XXX} + // offsets: {0, 3, 6, 7, 8, 10, 10} + // nullmask: {1, 1, 1, 1, 0, 1, 1, 0, 0, 0} + // outputs: {3, 3, 1, XXX, XXX, XXX} + // output nullmask: {1, 1, 1, 0, 0, 0} + auto input = fixed_width_column_wrapper{{1, 2, 3, 1, XXX, 3, 1, XXX, XXX, XXX}, + {1, 1, 1, 1, 0, 1, 1, 0, 0, 0}}; + auto offsets = std::vector{0, 3, 6, 7, 8, 10, 10}; + auto d_offsets = thrust::device_vector(offsets); + auto expect = fixed_width_column_wrapper{{3, 3, 1, XXX, XXX, XXX}, {1, 1, 1, 0, 0, 0}}; + + auto res = segmented_reduce(input, + d_offsets, + *make_max_aggregation(), + data_type{type_to_id()}, + null_policy::EXCLUDE); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*res, expect); +} + +TYPED_TEST(SegmentedReductionTest, MinExcludeNulls) +{ + // [1, 2, 3], [1, null, 3], [1], [null], [null, null], [] + // values: {1, 2, 3, 1, XXX, 3, 1, XXX, XXX, XXX} + // offsets: {0, 3, 6, 7, 8, 10, 10} + // nullmask: {1, 1, 1, 1, 0, 1, 1, 0, 0, 0} + // outputs: {1, 1, 1, XXX, XXX, XXX} + // output nullmask: {1, 1, 1, 0, 0, 0} + auto input = fixed_width_column_wrapper{{1, 2, 3, 1, XXX, 3, 1, XXX, XXX, XXX}, + {1, 1, 1, 1, 0, 1, 1, 0, 0, 0}}; + auto offsets = std::vector{0, 3, 6, 7, 8, 10, 10}; + auto d_offsets = thrust::device_vector(offsets); + auto expect = fixed_width_column_wrapper{{1, 1, 1, XXX, XXX, XXX}, {1, 1, 1, 0, 0, 0}}; + + auto res = segmented_reduce(input, + d_offsets, + *make_min_aggregation(), + data_type{type_to_id()}, + null_policy::EXCLUDE); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*res, expect); +} + +TYPED_TEST(SegmentedReductionTest, AnyExcludeNulls) +{ + // [0, 0, 0], [0, null, 0], [0, 1, 0], [1, null, 0], [], [0], [1], [null], [null, null] + // values: {0, 0, 0, 0, XXX, 0, 0, 1, 0, 1, XXX, 0, 0, 1, XXX, XXX, XXX} + // offsets: {0, 3, 6, 9, 12, 12, 13, 14, 15, 17} + // nullmask:{1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 0, 1, 1, 1, 0, 0, 0} + // outputs: {0, 0, 1, 1, XXX, 0, 1, XXX, XXX} + // output nullmask: {1, 1, 1, 1, 0, 1, 1, 0, 0} + auto input = fixed_width_column_wrapper{ + {0, 0, 0, 0, XXX, 0, 0, 1, 0, 1, XXX, 0, 0, 1, XXX, XXX, XXX}, + {1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 0, 1, 1, 1, 0, 0, 0}}; + auto offsets = std::vector{0, 3, 6, 9, 12, 12, 13, 14, 15, 17}; + auto d_offsets = thrust::device_vector(offsets); + auto expect = fixed_width_column_wrapper{ + {false, false, true, true, bool{XXX}, false, true, bool{XXX}, bool{XXX}}, + {true, true, true, true, false, true, true, false, false}}; + + auto res = segmented_reduce(input, + d_offsets, + *make_any_aggregation(), + data_type{type_id::BOOL8}, + null_policy::EXCLUDE); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*res, expect); +} + +TYPED_TEST(SegmentedReductionTest, AllExcludeNulls) +{ + // [1, 2, 3], [1, null, 3], [], [1], [null], [null, null], [1, 0, 3], [1, null, 0], [0] + // values: {1, 2, 3, 1, XXX, 3, 1, XXX, XXX, XXX, 1, 0, 3, 1, XXX, 0, 0} + // offsets: {0, 3, 6, 6, 7, 8, 10, 13, 16, 17} + // nullmask: {1, 1, 1, 1, 0, 1, 1, 0, 0, 0, 1, 1, 1, 1, 0, 1, 1} + // outputs: {true, true, XXX, true, XXX, XXX, false, false, false} + // output nullmask: {1, 1, 0, 1, 0, 0, 1, 1, 1} + auto input = fixed_width_column_wrapper{ + {1, 2, 3, 1, XXX, 3, 1, XXX, XXX, XXX, 1, 0, 3, 1, XXX, 0, 0}, + {1, 1, 1, 1, 0, 1, 1, 0, 0, 0, 1, 1, 1, 1, 0, 1, 1}}; + auto offsets = std::vector{0, 3, 6, 6, 7, 8, 10, 13, 16, 17}; + auto d_offsets = thrust::device_vector(offsets); + auto expect = fixed_width_column_wrapper{ + {true, true, bool{XXX}, true, bool{XXX}, bool{XXX}, false, false, false}, + {true, true, false, true, false, false, true, true, true}}; + + auto res = segmented_reduce(input, + d_offsets, + *make_all_aggregation(), + data_type{type_id::BOOL8}, + null_policy::EXCLUDE); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*res, expect); +} + +TYPED_TEST(SegmentedReductionTest, SumIncludeNulls) +{ + // [1, 2, 3], [1, null, 3], [1], [null], [null, null], [] + // values: {1, 2, 3, 1, XXX, 3, 1, XXX, XXX, XXX} + // offsets: {0, 3, 6, 7, 8, 10, 10} + // nullmask: {1, 1, 1, 1, 0, 1, 1, 0, 0, 0} + // outputs: {6, XXX, 1, XXX, XXX, XXX} + // output nullmask: {1, 0, 1, 0, 0, 0} + auto input = fixed_width_column_wrapper{{1, 2, 3, 1, XXX, 3, 1, XXX, XXX, XXX}, + {1, 1, 1, 1, 0, 1, 1, 0, 0, 0}}; + auto offsets = std::vector{0, 3, 6, 7, 8, 10, 10}; + auto d_offsets = thrust::device_vector(offsets); + auto expect = + fixed_width_column_wrapper{{6, XXX, 1, XXX, XXX, XXX}, {1, 0, 1, 0, 0, 0}}; + + auto res = segmented_reduce(input, + d_offsets, + *make_sum_aggregation(), + data_type{type_to_id()}, + null_policy::INCLUDE); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*res, expect); +} + +TYPED_TEST(SegmentedReductionTest, ProductIncludeNulls) +{ + // [1, 3, 5], [null, 3, 5], [1], [null], [null, null], [] + // values: {1, 3, 5, XXX, 3, 5, 1, XXX, XXX, XXX} + // offsets: {0, 3, 6, 7, 8, 10, 10} + // nullmask: {1, 1, 1, 0, 1, 1, 1, 0, 0, 0} + // outputs: {15, XXX, 1, XXX, XXX, XXX} + // output nullmask: {1, 0, 1, 0, 0, 0} + auto input = fixed_width_column_wrapper{{1, 3, 5, XXX, 3, 5, 1, XXX, XXX, XXX}, + {1, 1, 1, 0, 1, 1, 1, 0, 0, 0}}; + auto offsets = std::vector{0, 3, 6, 7, 8, 10, 10}; + auto d_offsets = thrust::device_vector(offsets); + auto expect = + fixed_width_column_wrapper{{15, XXX, 1, XXX, XXX, XXX}, {1, 0, 1, 0, 0, 0}}; + + auto res = segmented_reduce(input, + d_offsets, + *make_product_aggregation(), + data_type{type_to_id()}, + null_policy::INCLUDE); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*res, expect); +} + +TYPED_TEST(SegmentedReductionTest, MaxIncludeNulls) +{ + // [1, 2, 3], [1, null, 3], [1], [null], [null, null], [] + // values: {1, 2, 3, 1, XXX, 3, 1, XXX, XXX, XXX} + // offsets: {0, 3, 6, 7, 8, 10, 10} + // nullmask: {1, 1, 1, 1, 0, 1, 1, 0, 0, 0} + // outputs: {3, XXX, 1, XXX, XXX, XXX} + // output nullmask: {1, 0, 1, 0, 0, 0} + auto input = fixed_width_column_wrapper{{1, 2, 3, 1, XXX, 3, 1, XXX, XXX, XXX}, + {1, 1, 1, 1, 0, 1, 1, 0, 0, 0}}; + auto offsets = std::vector{0, 3, 6, 7, 8, 10, 10}; + auto d_offsets = thrust::device_vector(offsets); + auto expect = + fixed_width_column_wrapper{{3, XXX, 1, XXX, XXX, XXX}, {1, 0, 1, 0, 0, 0}}; + + auto res = segmented_reduce(input, + d_offsets, + *make_max_aggregation(), + data_type{type_to_id()}, + null_policy::INCLUDE); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*res, expect); +} + +TYPED_TEST(SegmentedReductionTest, MinIncludeNulls) +{ + // [1, 2, 3], [1, null, 3], [1], [null], [null, null], [] + // values: {1, 2, 3, 1, XXX, 3, 1, XXX, XXX} + // offsets: {0, 3, 6, 7, 8, 10, 10} + // nullmask: {1, 1, 1, 1, 0, 1, 1, 0, 0} + // outputs: {1, XXX, 1, XXX, XXX, XXX} + // output nullmask: {1, 0, 1, 0, 0, 0} + auto input = fixed_width_column_wrapper{{1, 2, 3, 1, XXX, 3, 1, XXX, XXX, XXX}, + {1, 1, 1, 1, 0, 1, 1, 0, 0}}; + auto offsets = std::vector{0, 3, 6, 7, 8, 10, 10}; + auto d_offsets = thrust::device_vector(offsets); + auto expect = + fixed_width_column_wrapper{{1, XXX, 1, XXX, XXX, XXX}, {1, 0, 1, 0, 0, 0}}; + + auto res = segmented_reduce(input, + d_offsets, + *make_min_aggregation(), + data_type{type_to_id()}, + null_policy::INCLUDE); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*res, expect); +} + +TYPED_TEST(SegmentedReductionTest, AnyIncludeNulls) +{ + // [0, 0, 0], [0, null, 0], [0, 1, 0], [1, null, 0], [], [0], [1], [null], [null, null] + // values: {0, 0, 0, 0, XXX, 0, 0, 1, 0, 1, XXX, 0, 0, 1, XXX, XXX, XXX} + // offsets: {0, 3, 6, 9, 12, 12, 13, 14, 15, 17} + // nullmask:{1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 0, 1, 1, 1, 0, 0, 0} + // outputs: {0, XXX, 1, XXX, XXX, 0, 1, XXX, XXX} + // output nullmask: {1, 0, 1, 0, 0, 1, 1, 0, 0} + auto input = fixed_width_column_wrapper{ + {0, 0, 0, 0, XXX, 0, 0, 1, 0, 1, XXX, 0, 0, 1, XXX, XXX, XXX}, + {1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 0, 1, 1, 1, 0, 0, 0}}; + auto offsets = std::vector{0, 3, 6, 9, 12, 12, 13, 14, 15, 17}; + auto d_offsets = thrust::device_vector(offsets); + auto expect = fixed_width_column_wrapper{ + {false, bool{XXX}, true, bool{XXX}, bool{XXX}, false, true, bool{XXX}, bool{XXX}}, + {true, false, true, false, false, true, true, false, false}}; + + auto res = segmented_reduce(input, + d_offsets, + *make_any_aggregation(), + data_type{type_id::BOOL8}, + null_policy::INCLUDE); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*res, expect); +} + +TYPED_TEST(SegmentedReductionTest, AllIncludeNulls) +{ + // [1, 2, 3], [1, null, 3], [], [1], [null], [null, null], [1, 0, 3], [1, null, 0], [0] + // values: {1, 2, 3, 1, XXX, 3, 1, XXX, XXX, XXX, 1, 0, 3, 1, XXX, 0, 0} + // offsets: {0, 3, 6, 6, 7, 8, 10, 13, 16, 17} + // nullmask: {1, 1, 1, 1, 0, 1, 1, 0, 0, 0, 1, 1, 1, 1, 0, 1, 1} + // outputs: {true, XXX, XXX, true, XXX, XXX, false, XXX, false} + // output nullmask: {1, 0, 0, 1, 0, 0, 1, 0, 1} + auto input = fixed_width_column_wrapper{ + {1, 2, 3, 1, XXX, 3, 1, XXX, XXX, XXX, 1, 0, 3, 1, XXX, 0, 0}, + {1, 1, 1, 1, 0, 1, 1, 0, 0, 0, 1, 1, 1, 1, 0, 1, 1}}; + auto offsets = std::vector{0, 3, 6, 6, 7, 8, 10, 13, 16, 17}; + auto d_offsets = thrust::device_vector(offsets); + auto expect = fixed_width_column_wrapper{ + {true, bool{XXX}, bool{XXX}, true, bool{XXX}, bool{XXX}, false, bool{XXX}, false}, + {true, false, false, true, false, false, true, false, true}}; + + auto res = segmented_reduce(input, + d_offsets, + *make_all_aggregation(), + data_type{type_id::BOOL8}, + null_policy::INCLUDE); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*res, expect); +} + +TEST_F(SegmentedReductionTestUntyped, PartialSegmentReudction) +{ + // Segmented reduction allows offsets only specify part of the input columns. + // [1], [2, 3], [4] + // values: {1, 2, 3, 4, 5, 6, 7} + // offsets: {0, 1, 3, 4} + // nullmask: {1, 1, 1, 1, 1, 1, 1} + // outputs: {1, 5, 4} + // output nullmask: {1, 1, 1} + + auto input = fixed_width_column_wrapper{{1, 2, 3, 4, 5, 6, 7}, + {true, true, true, true, true, true, true}}; + auto offsets = std::vector{0, 1, 3, 4}; + auto d_offsets = thrust::device_vector(offsets); + auto expect = fixed_width_column_wrapper{{1, 5, 4}, {true, true, true}}; + + auto res = segmented_reduce(input, + d_offsets, + *make_sum_aggregation(), + data_type{type_id::INT32}, + null_policy::INCLUDE); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*res, expect); +} + +TEST_F(SegmentedReductionTestUntyped, NonNullableInput) +{ + // Segmented reduction allows offsets only specify part of the input columns. + // [1], [], [2, 3], [4, 5, 6, 7] + // values: {1, 2, 3, 4, 5, 6, 7} + // offsets: {0, 1, 1, 3, 7} + // nullmask: nullptr + // outputs: {1, 5, 4} + // output nullmask: {1, 1, 1} + + auto input = fixed_width_column_wrapper{1, 2, 3, 4, 5, 6, 7}; + auto offsets = std::vector{0, 1, 1, 3, 7}; + auto d_offsets = thrust::device_vector(offsets); + auto expect = fixed_width_column_wrapper{{1, XXX, 5, 22}, {true, false, true, true}}; + + auto res = segmented_reduce(input, + d_offsets, + *make_sum_aggregation(), + data_type{type_id::INT32}, + null_policy::INCLUDE); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*res, expect); +} + +TEST_F(SegmentedReductionTestUntyped, ReduceEmptyColumn) +{ + auto input = fixed_width_column_wrapper{}; + auto offsets = std::vector{0}; + auto d_offsets = thrust::device_vector(offsets); + auto expect = fixed_width_column_wrapper{}; + + auto res = segmented_reduce(input, + d_offsets, + *make_sum_aggregation(), + data_type{type_to_id()}, + null_policy::EXCLUDE); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*res, expect); +} + +#undef XXX + +} // namespace test +} // namespace cudf