diff --git a/conda/recipes/libcudf/meta.yaml b/conda/recipes/libcudf/meta.yaml index 1d660e2cd74..a1953a2d358 100644 --- a/conda/recipes/libcudf/meta.yaml +++ b/conda/recipes/libcudf/meta.yaml @@ -1,4 +1,4 @@ -# Copyright (c) 2018-2020, NVIDIA CORPORATION. +# Copyright (c) 2018-2021, NVIDIA CORPORATION. {% set version = environ.get('GIT_DESCRIBE_TAG', '0.0.0.dev').lstrip('v') + environ.get('VERSION_SUFFIX', '') %} {% set minor_version = version.split('.')[0] + '.' + version.split('.')[1] %} @@ -128,6 +128,7 @@ test: - test -f $PREFIX/include/cudf/lists/contains.hpp - test -f $PREFIX/include/cudf/lists/gather.hpp - test -f $PREFIX/include/cudf/lists/lists_column_view.hpp + - test -f $PREFIX/include/cudf/lists/sorting.hpp - test -f $PREFIX/include/cudf/merge.hpp - test -f $PREFIX/include/cudf/null_mask.hpp - test -f $PREFIX/include/cudf/partitioning.hpp diff --git a/cpp/include/cudf/detail/sorting.hpp b/cpp/include/cudf/detail/sorting.hpp index 0ac20ed3c94..3127a5f89f1 100644 --- a/cpp/include/cudf/detail/sorting.hpp +++ b/cpp/include/cudf/detail/sorting.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * Copyright (c) 2019-2021, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -63,5 +63,32 @@ std::unique_ptr sort_by_key( rmm::cuda_stream_view stream = rmm::cuda_stream_default, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +/** + * @copydoc cudf::segmented_sorted_order + * + * @param[in] stream CUDA stream used for device memory operations and kernel launches. + */ +std::unique_ptr segmented_sorted_order( + table_view const& keys, + column_view const& segment_offsets, + std::vector const& column_order = {}, + std::vector const& null_precedence = {}, + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + +/** + * @copydoc cudf::segmented_sort_by_key + * + * @param[in] stream CUDA stream used for device memory operations and kernel launches. + */ +std::unique_ptr
segmented_sort_by_key( + table_view const& values, + table_view const& keys, + column_view const& segment_offsets, + std::vector const& column_order = {}, + std::vector const& null_precedence = {}, + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + } // namespace detail } // namespace cudf diff --git a/cpp/include/cudf/lists/list_device_view.cuh b/cpp/include/cudf/lists/list_device_view.cuh index 824b10ced83..81109bedaed 100644 --- a/cpp/include/cudf/lists/list_device_view.cuh +++ b/cpp/include/cudf/lists/list_device_view.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020, NVIDIA CORPORATION. + * Copyright (c) 2020-2021, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -190,4 +190,27 @@ class list_device_view { }; }; +/** + * @brief returns size of the list by row index + * + */ +struct list_size_functor { + column_device_view const d_column; + CUDA_HOST_DEVICE_CALLABLE list_size_functor(column_device_view const& d_col) : d_column(d_col) + { +#if defined(__CUDA_ARCH__) + release_assert(d_col.type().id() == type_id::LIST && "Only list type column is supported"); +#else + CUDF_EXPECTS(d_col.type().id() == type_id::LIST, "Only list type column is supported"); +#endif + } + CUDA_DEVICE_CALLABLE size_type operator()(size_type idx) + { + if (d_column.is_null(idx)) return size_type{0}; + auto d_offsets = + d_column.child(lists_column_view::offsets_column_index).data() + d_column.offset(); + return d_offsets[idx + 1] - d_offsets[idx]; + } +}; + } // namespace cudf diff --git a/cpp/include/cudf/lists/sorting.hpp b/cpp/include/cudf/lists/sorting.hpp new file mode 100644 index 00000000000..e27f3d03d86 --- /dev/null +++ b/cpp/include/cudf/lists/sorting.hpp @@ -0,0 +1,59 @@ +/* + * Copyright (c) 2021, 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 + +namespace cudf { +namespace lists { +/** + * @addtogroup lists_sort + * @{ + * @file + */ + +/** + * @brief Segmented sort of the elements within a list in each row of a list column. + * + * `source_column` with depth 1 is only supported. + * + * * @code{.pseudo} + * source_column : [{4, 2, 3, 1}, {1, 2, NULL, 4}, {-10, 10, 0}] + * + * Ascending, Null After : [{1, 2, 3, 4}, {1, 2, 4, NULL}, {-10, 0, 10}] + * Ascending, Null Before : [{1, 2, 3, 4}, {NULL, 1, 2, 4}, {-10, 0, 10}] + * Descending, Null After : [{4, 3, 2, 1}, {NULL, 4, 2, 1}, {10, 0, -10}] + * Descending, Null Before : [{4, 3, 2, 1}, {4, 2, 1, NULL}, {10, 0, -10}] + * @endcode + * + * @param source_column View of the list column of numeric types to sort + * @param column_order The desired sort order + * @param null_precedence The desired order of null compared to other elements in the list + * @param mr Device memory resource to allocate any returned objects + * @return list column with elements in each list sorted. + * + */ +std::unique_ptr sort_lists( + lists_column_view const& source_column, + order column_order, + null_order null_precedence, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + +/** @} */ // end of group +} // namespace lists +} // namespace cudf diff --git a/cpp/include/cudf/sorting.hpp b/cpp/include/cudf/sorting.hpp index 1116b49c892..2454cfe7c7b 100644 --- a/cpp/include/cudf/sorting.hpp +++ b/cpp/include/cudf/sorting.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * Copyright (c) 2019-2021, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -128,7 +128,7 @@ std::unique_ptr
sort( * @param values The table to reorder * @param keys The table that determines the ordering * @param column_order The desired order for each column in `keys`. Size must be - * equal to `input.num_columns()` or empty. If empty, all columns are sorted in + * equal to `keys.num_columns()` or empty. If empty, all columns are sorted in * ascending order. * @param null_precedence The desired order of a null element compared to other * elements for each column in `keys`. Size must be equal to @@ -184,5 +184,62 @@ std::unique_ptr rank( bool percentage, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +/** + * @brief Returns sorted order after sorting each segment in the table. + * + * If segment_offsets contains values larger than number of rows, behaviour is undefined. + * @throws cudf::logic_error if `segment_offsets` is not `size_type` column. + * + * @param keys The table that determines the ordering of elements in each segment + * @param segment_offsets The column of `size_type` type containing start offset index for each + * contiguous segment. + * @param column_order The desired order for each column in `keys`. Size must be + * equal to `keys.num_columns()` or empty. If empty, all columns are sorted in + * ascending order. + * @param null_precedence The desired order of a null element compared to other + * elements for each column in `keys`. Size must be equal to + * `keys.num_columns()` or empty. If empty, all columns will be sorted with + * `null_order::BEFORE`. + * @param mr Device memory resource to allocate any returned objects + * @return sorted order of the segment sorted table . + * + */ +std::unique_ptr segmented_sorted_order( + table_view const& keys, + column_view const& segment_offsets, + std::vector const& column_order = {}, + std::vector const& null_precedence = {}, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + +/** + * @brief Performs a lexicographic segmented sort of a table + * + * If segment_offsets contains values larger than number of rows, behaviour is undefined. + * @throws cudf::logic_error if `values.num_rows() != keys.num_rows()`. + * @throws cudf::logic_error if `segment_offsets` is not `size_type` column. + * + * @param values The table to reorder + * @param keys The table that determines the ordering of elements in each segment + * @param segment_offsets The column of `size_type` type containing start offset index for each + * contiguous segment. + * @param column_order The desired order for each column in `keys`. Size must be + * equal to `keys.num_columns()` or empty. If empty, all columns are sorted in + * ascending order. + * @param null_precedence The desired order of a null element compared to other + * elements for each column in `keys`. Size must be equal to + * `keys.num_columns()` or empty. If empty, all columns will be sorted with + * `null_order::BEFORE`. + * @param mr Device memory resource to allocate any returned objects + * @return table with elements in each segment sorted. + * + */ +std::unique_ptr
segmented_sort_by_key( + table_view const& values, + table_view const& keys, + column_view const& segment_offsets, + std::vector const& column_order = {}, + std::vector const& null_precedence = {}, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + /** @} */ // end of group } // namespace cudf diff --git a/cpp/src/lists/count_elements.cu b/cpp/src/lists/count_elements.cu index 78549152770..ba366b3a020 100644 --- a/cpp/src/lists/count_elements.cu +++ b/cpp/src/lists/count_elements.cu @@ -20,6 +20,7 @@ #include #include #include +#include #include #include @@ -60,13 +61,7 @@ std::unique_ptr count_elements(lists_column_view const& input, thrust::make_counting_iterator(0), thrust::make_counting_iterator(input.size()), output->mutable_view().begin(), - [d_column] __device__(size_type idx) { - if (d_column.is_null(idx)) return size_type{0}; - auto d_offsets = - d_column.child(lists_column_view::offsets_column_index).data() + - d_column.offset(); - return d_offsets[idx + 1] - d_offsets[idx]; - }); + list_size_functor{d_column}); output->set_null_count(input.null_count()); // reset null count return output; diff --git a/cpp/src/lists/segmented_sort.cu b/cpp/src/lists/segmented_sort.cu new file mode 100644 index 00000000000..5681f7767e0 --- /dev/null +++ b/cpp/src/lists/segmented_sort.cu @@ -0,0 +1,270 @@ +/* + * Copyright (c) 2021, 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 +#include +#include +#include + +#include + +#include + +namespace cudf { +namespace lists { +namespace detail { + +struct SegmentedSortColumn { + template + void SortPairsAscending(KeyT const* keys_in, + KeyT* keys_out, + ValueT const* values_in, + ValueT* values_out, + int num_items, + int num_segments, + OffsetIteratorT begin_offsets, + OffsetIteratorT end_offsets, + rmm::cuda_stream_view stream) + { + rmm::device_buffer d_temp_storage; + size_t temp_storage_bytes = 0; + cub::DeviceSegmentedRadixSort::SortPairs(d_temp_storage.data(), + temp_storage_bytes, + keys_in, + keys_out, + values_in, + values_out, + num_items, + num_segments, + begin_offsets, + end_offsets, + 0, + sizeof(KeyT) * 8, + stream.value()); + d_temp_storage = rmm::device_buffer{temp_storage_bytes, stream}; + + cub::DeviceSegmentedRadixSort::SortPairs(d_temp_storage.data(), + temp_storage_bytes, + keys_in, + keys_out, + values_in, + values_out, + num_items, + num_segments, + begin_offsets, + end_offsets, + 0, + sizeof(KeyT) * 8, + stream.value()); + } + + template + void SortPairsDescending(KeyT const* keys_in, + KeyT* keys_out, + ValueT const* values_in, + ValueT* values_out, + int num_items, + int num_segments, + OffsetIteratorT begin_offsets, + OffsetIteratorT end_offsets, + rmm::cuda_stream_view stream) + { + rmm::device_buffer d_temp_storage; + size_t temp_storage_bytes = 0; + cub::DeviceSegmentedRadixSort::SortPairsDescending(d_temp_storage.data(), + temp_storage_bytes, + keys_in, + keys_out, + values_in, + values_out, + num_items, + num_segments, + begin_offsets, + end_offsets, + 0, + sizeof(KeyT) * 8, + stream.value()); + d_temp_storage = rmm::device_buffer{temp_storage_bytes, stream}; + + cub::DeviceSegmentedRadixSort::SortPairsDescending(d_temp_storage.data(), + temp_storage_bytes, + keys_in, + keys_out, + values_in, + values_out, + num_items, + num_segments, + begin_offsets, + end_offsets, + 0, + sizeof(KeyT) * 8, + stream.value()); + } + + template + std::enable_if_t(), std::unique_ptr> operator()( + column_view const& child, + column_view const& segment_offsets, + order column_order, + null_order null_precedence, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) + { + auto child_table = segmented_sort_by_key(table_view{{child}}, + table_view{{child}}, + segment_offsets, + {column_order}, + {null_precedence}, + stream, + mr); + return std::move(child_table->release().front()); + } + + template + std::enable_if_t(), std::unique_ptr> operator()( + column_view const& child, + column_view const& offsets, + order column_order, + null_order null_precedence, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) + { + auto output = + cudf::detail::allocate_like(child, child.size(), mask_allocation_policy::NEVER, stream, mr); + mutable_column_view mutable_output_view = output->mutable_view(); + + auto keys = [&]() { + if (child.nullable()) { + rmm::device_uvector keys(child.size(), stream); + auto const null_replace_T = null_precedence == null_order::AFTER + ? std::numeric_limits::max() + : std::numeric_limits::min(); + auto device_child = column_device_view::create(child, stream); + auto keys_in = + cudf::detail::make_null_replacement_iterator(*device_child, null_replace_T); + thrust::copy_n(rmm::exec_policy(stream), keys_in, child.size(), keys.begin()); + return keys; + } + return rmm::device_uvector{0, stream}; + }(); + + std::unique_ptr sorted_indices = cudf::make_numeric_column( + data_type(type_to_id()), child.size(), mask_state::UNALLOCATED, stream, mr); + mutable_column_view mutable_indices_view = sorted_indices->mutable_view(); + thrust::sequence(rmm::exec_policy(stream), + mutable_indices_view.begin(), + mutable_indices_view.end(), + 0); + + if (column_order == order::ASCENDING) + SortPairsAscending(child.nullable() ? keys.data() : child.begin(), + mutable_output_view.begin(), + mutable_indices_view.begin(), + mutable_indices_view.begin(), + child.size(), + offsets.size() - 1, + offsets.begin(), + offsets.begin() + 1, + stream); + else + SortPairsDescending(child.nullable() ? keys.data() : child.begin(), + mutable_output_view.begin(), + mutable_indices_view.begin(), + mutable_indices_view.begin(), + child.size(), + offsets.size() - 1, + offsets.begin(), + offsets.begin() + 1, + stream); + std::vector> output_cols; + output_cols.push_back(std::move(output)); + // rearrange the null_mask. + cudf::detail::gather_bitmask(cudf::table_view{{child}}, + mutable_indices_view.begin(), + output_cols, + cudf::detail::gather_bitmask_op::DONT_CHECK, + stream, + mr); + return std::move(output_cols.front()); + } +}; + +std::unique_ptr sort_lists(lists_column_view const& input, + order column_order, + null_order null_precedence, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + if (input.is_empty()) return empty_like(input.parent()); + auto segment_offsets = + cudf::detail::slice(input.offsets(), {input.offset(), input.offsets().size()}, stream)[0]; + // Copy list offsets. + auto output_offset = allocate_like(segment_offsets, mask_allocation_policy::RETAIN, mr); + thrust::transform(rmm::exec_policy(stream), + segment_offsets.begin(), + segment_offsets.end(), + output_offset->mutable_view().begin(), + [first = segment_offsets.begin()] __device__(auto offset_index) { + return offset_index - *first; + }); + // for numeric columns, calls Faster segmented radix sort path + // for non-numeric columns, calls segmented_sort_by_key. + auto output_child = type_dispatcher(input.child().type(), + SegmentedSortColumn{}, + input.get_sliced_child(stream), + output_offset->view(), + column_order, + null_precedence, + stream, + mr); + + auto null_mask = cudf::detail::copy_bitmask(input.parent(), stream, mr); + + // Assemble list column & return + return make_lists_column(input.size(), + std::move(output_offset), + std::move(output_child), + input.null_count(), + std::move(null_mask)); +} +} // namespace detail + +std::unique_ptr sort_lists(lists_column_view const& input, + order column_order, + null_order null_precedence, + rmm::mr::device_memory_resource* mr) +{ + CUDF_FUNC_RANGE(); + return detail::sort_lists(input, column_order, null_precedence, rmm::cuda_stream_default, mr); +} + +} // namespace lists +} // namespace cudf diff --git a/cpp/src/sort/segmented_sort.cu b/cpp/src/sort/segmented_sort.cu new file mode 100644 index 00000000000..f8b0e311c9c --- /dev/null +++ b/cpp/src/sort/segmented_sort.cu @@ -0,0 +1,141 @@ +/* + * Copyright (c) 2021, 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 + +#include +#include +#include +#include + +#include +#include +#include +#include + +namespace cudf { +namespace detail { + +// returns segment indices for each element for all segments. +// first segment begin index = 0, last segment end index = num_rows. +rmm::device_uvector get_segment_indices(size_type num_rows, + column_view const& offsets, + rmm::cuda_stream_view stream) +{ + rmm::device_uvector segment_ids(num_rows, stream); + + auto offset_begin = offsets.begin(); // assumes already offset column contains offset. + auto offsets_minus_one = thrust::make_transform_iterator( + offset_begin, [offset_begin] __device__(auto i) { return i - 1; }); + auto counting_iter = thrust::make_counting_iterator(0); + thrust::lower_bound(rmm::exec_policy(stream), + offsets_minus_one, + offsets_minus_one + offsets.size(), + counting_iter, + counting_iter + segment_ids.size(), + segment_ids.begin()); + return std::move(segment_ids); +} + +std::unique_ptr segmented_sorted_order(table_view const& keys, + column_view const& segment_offsets, + std::vector const& column_order, + std::vector const& null_precedence, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + CUDF_EXPECTS(segment_offsets.type() == data_type(type_to_id()), + "segment offsets should be size_type"); + // Get segment id of each element in all segments. + auto segment_ids = get_segment_indices(keys.num_rows(), segment_offsets, stream); + + // insert segment id before all columns. + std::vector keys_with_segid; + keys_with_segid.reserve(keys.num_columns() + 1); + keys_with_segid.push_back( + column_view(data_type(type_to_id()), segment_ids.size(), segment_ids.data())); + keys_with_segid.insert(keys_with_segid.end(), keys.begin(), keys.end()); + auto segid_keys = table_view(keys_with_segid); + + auto prepend_default = [](auto const& vector, auto default_value) { + if (vector.empty()) return vector; + std::remove_cv_t> pre_vector; + pre_vector.reserve(pre_vector.size() + 1); + pre_vector.push_back(default_value); + pre_vector.insert(pre_vector.end(), vector.begin(), vector.end()); + return pre_vector; + }; + auto child_column_order = prepend_default(column_order, order::ASCENDING); + auto child_null_precedence = prepend_default(null_precedence, null_order::AFTER); + // return sorted order of child columns + return detail::sorted_order(segid_keys, child_column_order, child_null_precedence, stream, mr); +} + +std::unique_ptr
segmented_sort_by_key(table_view const& values, + table_view const& keys, + column_view const& segment_offsets, + std::vector const& column_order, + std::vector const& null_precedence, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + CUDF_EXPECTS(values.num_rows() == keys.num_rows(), + "Mismatch in number of rows for values and keys"); + auto sorted_order = segmented_sorted_order(keys, + segment_offsets, + column_order, + null_precedence, + stream, + rmm::mr::get_current_device_resource()); + + // Gather segmented sort of child value columns` + return detail::gather(values, + sorted_order->view(), + out_of_bounds_policy::DONT_CHECK, + detail::negative_index_policy::NOT_ALLOWED, + stream, + mr); +} +} // namespace detail + +std::unique_ptr
segmented_sort_by_key(table_view const& values, + table_view const& keys, + column_view const& segment_offsets, + std::vector const& column_order, + std::vector const& null_precedence, + rmm::mr::device_memory_resource* mr) +{ + CUDF_FUNC_RANGE(); + return detail::segmented_sort_by_key( + values, keys, segment_offsets, column_order, null_precedence, rmm::cuda_stream_default, mr); +} + +} // namespace cudf diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 8395a3cc1f2..ccbc38f3696 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -362,6 +362,7 @@ ConfigureTest(JSON_TEST "${JSON_TEST_SRC}") # - sort tests ------------------------------------------------------------------------------------ set(SORT_TEST_SRC + "${CMAKE_CURRENT_SOURCE_DIR}/sort/segmented_sort_tests.cpp" "${CMAKE_CURRENT_SOURCE_DIR}/sort/sort_test.cpp" "${CMAKE_CURRENT_SOURCE_DIR}/sort/rank_test.cpp") @@ -663,7 +664,8 @@ ConfigureTest(AST_TEST "${AST_TEST_SRC}") set(LISTS_TEST_SRC "${CMAKE_CURRENT_SOURCE_DIR}/lists/contains_tests.cpp" "${CMAKE_CURRENT_SOURCE_DIR}/lists/count_elements_tests.cpp" - "${CMAKE_CURRENT_SOURCE_DIR}/lists/extract_tests.cpp") + "${CMAKE_CURRENT_SOURCE_DIR}/lists/extract_tests.cpp" + "${CMAKE_CURRENT_SOURCE_DIR}/lists/sort_lists_tests.cpp") ConfigureTest(LISTS_TEST "${LISTS_TEST_SRC}") diff --git a/cpp/tests/lists/sort_lists_tests.cpp b/cpp/tests/lists/sort_lists_tests.cpp new file mode 100644 index 00000000000..ac73297f088 --- /dev/null +++ b/cpp/tests/lists/sort_lists_tests.cpp @@ -0,0 +1,182 @@ +/* + * Copyright (c) 2021, 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 +#include + +template +using LCW = cudf::test::lists_column_wrapper; +using cudf::lists_column_view; +using cudf::lists::sort_lists; + +namespace cudf { +namespace test { + +template +struct SortLists : public BaseFixture { +}; + +TYPED_TEST_CASE(SortLists, NumericTypes); +using SortListsInt = SortLists; + +/* +empty case + empty list + single row with empty list + multi row with empty lists +single case + single list with single element + single list with multi element +normal case without nulls +Null cases + null rows + null elements in list. +Error: + depth>1 +*/ +TYPED_TEST(SortLists, NoNull) +{ + using T = TypeParam; + + // List + LCW list{{3, 2, 1, 4}, {5}, {10, 8, 9}, {6, 7}}; + + // Ascending + // LCW order{{2, 1, 0, 3}, {0}, {1, 2, 0}, {0, 1}}; + LCW expected{{1, 2, 3, 4}, {5}, {8, 9, 10}, {6, 7}}; + auto results = sort_lists(lists_column_view{list}, order::ASCENDING, null_order::AFTER); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), expected); + + results = sort_lists(lists_column_view{list}, order::ASCENDING, null_order::BEFORE); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), expected); + + // Descending + // LCW order{{3, 0, 1, 2}, {0}, {0, 1, 2}, {1, 0}}; + LCW expected2{{4, 3, 2, 1}, {5}, {10, 9, 8}, {7, 6}}; + results = sort_lists(lists_column_view{list}, order::DESCENDING, null_order::AFTER); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), expected2); + + results = sort_lists(lists_column_view{list}, order::DESCENDING, null_order::BEFORE); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), expected2); +} + +TYPED_TEST(SortLists, Null) +{ + using T = TypeParam; + if (std::is_same::value) return; + std::vector valids_o{1, 1, 0, 1}; + std::vector valids_a{1, 1, 1, 0}; + std::vector valids_b{0, 1, 1, 1}; + + // List + LCW list{{{3, 2, 4, 1}, valids_o.begin()}, {5}, {10, 8, 9}, {6, 7}}; + // LCW order{{2, 1, 3, 0}, {0}, {1, 2, 0}, {0, 1}}; + LCW expected1{{{1, 2, 3, 4}, valids_a.begin()}, {5}, {8, 9, 10}, {6, 7}}; + LCW expected2{{{4, 1, 2, 3}, valids_b.begin()}, {5}, {8, 9, 10}, {6, 7}}; + auto results = sort_lists(lists_column_view{list}, order::ASCENDING, null_order::AFTER); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), expected1); + + results = sort_lists(lists_column_view{list}, order::ASCENDING, null_order::BEFORE); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), expected2); + + // Descending + // LCW order{{3, 0, 1, 2}, {0}, {0, 1, 2}, {1, 0}}; + LCW expected3{{{4, 3, 2, 1}, valids_b.begin()}, {5}, {10, 9, 8}, {7, 6}}; + LCW expected4{{{3, 2, 1, 4}, valids_a.begin()}, {5}, {10, 9, 8}, {7, 6}}; + results = sort_lists(lists_column_view{list}, order::DESCENDING, null_order::AFTER); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), expected3); + + results = sort_lists(lists_column_view{list}, order::DESCENDING, null_order::BEFORE); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), expected4); +} + +TEST_F(SortListsInt, Empty) +{ + using T = int; + LCW l1{}; + LCW l2{LCW{}}; + LCW l3{LCW{}, LCW{}}; + + auto results = sort_lists(lists_column_view{l1}, {}, {}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), l1); + results = sort_lists(lists_column_view{l2}, {}, {}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), l2); + results = sort_lists(lists_column_view{l3}, {}, {}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), l3); +} + +TEST_F(SortListsInt, Single) +{ + using T = int; + LCW l1{{1}}; + LCW l2{{1, 2, 3}}; + + auto results = sort_lists(lists_column_view{l1}, {}, {}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), l1); + results = sort_lists(lists_column_view{l2}, {}, {}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), l2); +} + +TEST_F(SortListsInt, NullRows) +{ + using T = int; + std::vector valids{0, 1, 0}; + LCW l1{{{1, 2, 3}, {4, 5, 6}, {7}}, valids.begin()}; // offset 0, 0, 3, 3 + + auto results = sort_lists(lists_column_view{l1}, {}, {}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), l1); +} + +/* +// Disabling this test. +// Reason: After this exception "cudaErrorAssert device-side assert triggered", further tests fail +TEST_F(SortListsInt, Depth) +{ + using T = int; + LCW l1{LCW{{1, 2}, {3}}, LCW{{4, 5}}}; + // device exception + EXPECT_THROW(sort_lists(lists_column_view{l1}, {}, {}), std::exception); +} +*/ + +TEST_F(SortListsInt, Sliced) +{ + using T = int; + LCW l1{{1, 2, 3, 4}, {5, 6, 7}, {8, 9}, {10}}; + auto sliced_list = cudf::slice(l1, {1, 4})[0]; + + auto results = sort_lists(lists_column_view{sliced_list}, {}, {}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), sliced_list); +} + +} // namespace test +} // namespace cudf diff --git a/cpp/tests/sort/segmented_sort_tests.cpp b/cpp/tests/sort/segmented_sort_tests.cpp new file mode 100644 index 00000000000..e907212c9e8 --- /dev/null +++ b/cpp/tests/sort/segmented_sort_tests.cpp @@ -0,0 +1,272 @@ +/* + * Copyright (c) 2021, 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 + +template +using column_wrapper = cudf::test::fixed_width_column_wrapper; +using LCWstr = cudf::test::lists_column_wrapper; + +namespace cudf { +namespace test { + +template +struct SegmentedSort : public BaseFixture { +}; + +TYPED_TEST_CASE(SegmentedSort, NumericTypes); +using SegmentedSortInt = SegmentedSort; + +/* Summary of test cases. +empty case + key{}, + value{}, + segment_offset{} +single case + keys{1}, value{1} + segmented_offset{0}, {0, 1} +normal case +{8, 9, 2, 3, 2, 2, 4, 1, 7, 5, 6} +{0, 2, 5, 8 11} + without null + with null +corner case + sliced table, + sliced segment_offsets + non-zero start of segment_offsets without offset + non-zero start of segment_offsets with offset +mismatch sizes + keys, values num_rows + order, null_order + segmented_offsets beyond num_rows +//*/ +TEST_F(SegmentedSortInt, Empty) +{ + using T = int; + column_wrapper col_empty{}; + // clang-format off + column_wrapper col1{{8, 9, 2, 3, 2, 2, 4, 1, 7, 5, 6}}; + column_wrapper segments{{0, 2, 5, 8, 11}}; + // clang-format on + table_view table_empty{{col_empty}}; + table_view table_valid{{col1}}; + + CUDF_EXPECT_NO_THROW(cudf::segmented_sort_by_key(table_valid, table_valid, segments)); + CUDF_EXPECT_NO_THROW(cudf::segmented_sort_by_key(table_valid, table_valid, col_empty)); + CUDF_EXPECT_NO_THROW(cudf::segmented_sort_by_key(table_empty, table_empty, segments)); + CUDF_EXPECT_NO_THROW(cudf::segmented_sort_by_key(table_empty, table_empty, col_empty)); + + CUDF_EXPECT_THROW_MESSAGE(cudf::segmented_sort_by_key(table_empty, table_valid, segments), + "Mismatch in number of rows for values and keys"); + CUDF_EXPECT_THROW_MESSAGE(cudf::segmented_sort_by_key(table_empty, table_valid, col_empty), + "Mismatch in number of rows for values and keys"); + CUDF_EXPECT_THROW_MESSAGE(cudf::segmented_sort_by_key(table_valid, table_empty, segments), + "Mismatch in number of rows for values and keys"); + CUDF_EXPECT_THROW_MESSAGE(cudf::segmented_sort_by_key(table_valid, table_empty, col_empty), + "Mismatch in number of rows for values and keys"); +} + +TEST_F(SegmentedSortInt, Single) +{ + using T = int; + column_wrapper col1{{1}}; + column_wrapper col3{{8, 9, 2}}; + column_wrapper segments1{{0}}; + column_wrapper segments2{{0, 3}}; + table_view table_1elem{{col1}}; + table_view table_1segm{{col3}}; + CUDF_EXPECT_NO_THROW(cudf::segmented_sort_by_key(table_1elem, table_1elem, segments2)); + CUDF_EXPECT_NO_THROW(cudf::segmented_sort_by_key(table_1elem, table_1elem, segments1)); + CUDF_EXPECT_NO_THROW(cudf::segmented_sort_by_key(table_1segm, table_1segm, segments2)); + CUDF_EXPECT_NO_THROW(cudf::segmented_sort_by_key(table_1segm, table_1segm, segments1)); +} + +TYPED_TEST(SegmentedSort, NoNull) +{ + using T = TypeParam; + + // segments {0 1 2} {3 4} {5} {6 7 8 9 10}{11 12}{13}{14 15} + column_wrapper col1{{10, 36, 14, 32, 49, 23, 10, 34, 12, 45, 12, 37, 43, 26, 21, 16}}; + column_wrapper col2{{10, 63, 41, 23, 94, 32, 10, 43, 21, 54, 22, 73, 34, 62, 12, 61}}; + // segment sorted order {0 2 1} {3 4} {5} {6 8 10 7 9}{11 12}{13}{15 16} + column_wrapper segments{0, 3, 5, 5, 5, 6, 11, 13, 14, 16}; + table_view input1{{col1}}; + table_view input2{{col1, col2}}; + + // Ascending + column_wrapper col1_asc{{10, 14, 36, 32, 49, 23, 10, 12, 12, 34, 45, 37, 43, 26, 16, 21}}; + + auto results = cudf::segmented_sort_by_key(input1, input1, segments, {order::ASCENDING}); + CUDF_TEST_EXPECT_TABLES_EQUAL(results->view(), table_view{{col1_asc}}); + + column_wrapper col1_des{{36, 14, 10, 49, 32, 23, 45, 34, 12, 12, 10, 43, 37, 26, 21, 16}}; + results = cudf::segmented_sort_by_key(input1, input1, segments, {order::DESCENDING}); + CUDF_TEST_EXPECT_TABLES_EQUAL(results->view(), table_view{{col1_des}}); + + column_wrapper col1_12_asc{{10, 14, 36, 32, 49, 23, 10, 12, 12, 34, 45, 37, 43, 26, 16, 21}}; + column_wrapper col2_12_asc{{10, 41, 63, 23, 94, 32, 10, 21, 22, 43, 54, 73, 34, 62, 61, 12}}; + column_wrapper col2_12_des{{10, 41, 63, 23, 94, 32, 10, 22, 21, 43, 54, 73, 34, 62, 61, 12}}; + + table_view expected12_aa{{col1_12_asc, col2_12_asc}}; + results = cudf::segmented_sort_by_key(input2, input2, segments, {}); + CUDF_TEST_EXPECT_TABLES_EQUAL(results->view(), expected12_aa); + + table_view expected12_ad{{col1_12_asc, col2_12_des}}; + results = + cudf::segmented_sort_by_key(input2, input2, segments, {order::ASCENDING, order::DESCENDING}); + CUDF_TEST_EXPECT_TABLES_EQUAL(results->view(), expected12_ad); +} + +TYPED_TEST(SegmentedSort, Null) +{ + using T = TypeParam; + if (std::is_same::value) return; + + // segments {0 1 2} {3 4} {5} {6 7 8 9 10}{11 12}{13}{14 15} + column_wrapper col1{{1, 3, 2, 4, 5, 23, 6, 8, 7, 9, 7, 37, 43, 26, 21, 16}, + {1, 1, 0, 1, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1}}; + column_wrapper col2{{0, 0, 0, 1, 1, 4, 5, 5, 21, 5, 22, 6, 6, 7, 8, 8}, + {1, 1, 1, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1}}; + column_wrapper segments{0, 3, 5, 5, 5, 6, 11, 13, 14, 16}; + table_view input1{{col1}}; + table_view input2{{col1, col2}}; + + // Ascending + column_wrapper col1_aa{{1, 3, 2, 4, 5, 23, 6, 7, 7, 8, 9, 37, 43, 26, 16, 21}, + {1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1}}; + column_wrapper col1_ab{{2, 1, 3, 4, 5, 23, 9, 6, 7, 7, 8, 37, 43, 26, 16, 21}, + {0, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 1}}; + + auto results = cudf::segmented_sort_by_key(input1, input1, segments, {}, {null_order::AFTER}); + CUDF_TEST_EXPECT_TABLES_EQUAL(results->view(), table_view{{col1_aa}}); + results = cudf::segmented_sort_by_key(input1, input1, segments, {}, {null_order::BEFORE}); + CUDF_TEST_EXPECT_TABLES_EQUAL(results->view(), table_view{{col1_ab}}); + + // Descending + column_wrapper col1_da{{2, 3, 1, 5, 4, 23, 9, 8, 7, 7, 6, 43, 37, 26, 21, 16}, + {0, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 1}}; + column_wrapper col1_db{{3, 1, 2, 5, 4, 23, 8, 7, 7, 6, 9, 43, 37, 26, 21, 16}, + {1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1}}; + results = + cudf::segmented_sort_by_key(input1, input1, segments, {order::DESCENDING}, {null_order::AFTER}); + CUDF_TEST_EXPECT_TABLES_EQUAL(results->view(), table_view{{col1_da}}); + results = cudf::segmented_sort_by_key( + input1, input1, segments, {order::DESCENDING}, {null_order::BEFORE}); + CUDF_TEST_EXPECT_TABLES_EQUAL(results->view(), table_view{{col1_db}}); + + // second row null order. + column_wrapper col2_12_aa{{0, 0, 0, 1, 1, 4, 5, 22, 21, 5, 5, 6, 6, 7, 8, 8}, + {1, 1, 1, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1}}; + column_wrapper col2_12_ab{{0, 0, 0, 1, 1, 4, 5, 5, 21, 22, 5, 6, 6, 7, 8, 8}, + {1, 1, 1, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1}}; + table_view expected12_aa{{col1_aa, col2_12_aa}}; + table_view expected12_ab{{col1_ab, col2_12_ab}}; + results = cudf::segmented_sort_by_key( + input2, input2, segments, {}, {null_order::AFTER, null_order::AFTER}); + CUDF_TEST_EXPECT_TABLES_EQUAL(results->view(), expected12_aa); + results = cudf::segmented_sort_by_key( + input2, input2, segments, {}, {null_order::BEFORE, null_order::BEFORE}); + CUDF_TEST_EXPECT_TABLES_EQUAL(results->view(), expected12_ab); +} + +TEST_F(SegmentedSortInt, NonZeroSegmentsStart) +{ + using T = int; + // clang-format off + column_wrapper col1{{8, 9, 2, 3, 2, 2, 4, 1, 7, 5, 6}}; + column_wrapper segments1{{0, 2, 5, 8, 11}}; + column_wrapper segments2{{ 2, 5, 8, 11}}; + column_wrapper segments3{{ 6, 8, 11}}; + column_wrapper expected1{{0, 1, 2, 4, 3, 7, 5, 6, 9, 10, 8}}; + column_wrapper expected2{{0, 1, 2, 4, 3, 7, 5, 6, 9, 10, 8}}; + column_wrapper expected3{{2, 4, 5, 3, 0, 1, 7, 6, 9, 10, 8}}; + // clang-format on + table_view input{{col1}}; + auto results = cudf::detail::segmented_sorted_order(input, segments1); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), expected1); + results = cudf::detail::segmented_sorted_order(input, segments2); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), expected2); + results = cudf::detail::segmented_sorted_order(input, segments3); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), expected3); +} + +TEST_F(SegmentedSortInt, Sliced) +{ + using T = int; + // clang-format off + column_wrapper col1{{8, 9, 2, 3, 2, 2, 4, 1, 7, 5, 6}}; + // sliced 2, 2, 4, 1, 7, 5, 6 + column_wrapper segments1{{0, 2, 5}}; + column_wrapper segments2{{-4, 0, 2, 5}}; + column_wrapper segments3{{ 7}}; + column_wrapper expected1{{0, 1, 3, 2, 4, 5, 6}}; + column_wrapper expected2{{0, 1, 3, 2, 4, 5, 6}}; + column_wrapper expected3{{3, 0, 1, 2, 5, 6, 4}}; + // clang-format on + auto slice = cudf::slice(col1, {4, 11})[0]; // 7 elements + table_view input{{slice}}; + auto seg_slice = cudf::slice(segments2, {2, 4})[0]; // 2 elements + + // sliced input + auto results = cudf::detail::segmented_sorted_order(input, segments1); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), expected1); + // sliced input and sliced segment + results = cudf::detail::segmented_sorted_order(input, seg_slice); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), expected2); + // sliced input, segment end. + results = cudf::detail::segmented_sorted_order(input, segments3); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), expected3); +} + +TEST_F(SegmentedSortInt, ErrorsMismatchArgSizes) +{ + using T = int; + column_wrapper col1{{1, 2, 3, 4}}; + column_wrapper col2{{5, 6, 7, 8, 9}}; + table_view input1{{col1}}; + + // Mismatch order sizes + EXPECT_THROW( + cudf::segmented_sort_by_key(input1, input1, col2, {order::ASCENDING, order::ASCENDING}, {}), + logic_error); + // Mismatch null precedence sizes + EXPECT_THROW( + cudf::segmented_sort_by_key(input1, input1, col2, {}, {null_order::AFTER, null_order::AFTER}), + logic_error); + // Both + EXPECT_THROW(cudf::segmented_sort_by_key(input1, + input1, + col2, + {order::ASCENDING, order::ASCENDING}, + {null_order::AFTER, null_order::AFTER}), + logic_error); + // segmented_offsets beyond num_rows - undefined behaviour, no throw. + CUDF_EXPECT_NO_THROW(cudf::segmented_sort_by_key(input1, input1, col2)); +} + +} // namespace test +} // namespace cudf