From a2d2726f7f38fd06130074e7d174f5f8e2e2f21c Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Wed, 2 Dec 2020 15:37:26 -0800 Subject: [PATCH] Remove bounds check for `cudf::gather`(#6875) Closes #6478 `cudf::gather` now will not run a pre-pass to check for index validity. For `out_of_bounds_policy`, remove `FAIL`, while exposing `NULLIFY` and `DONT_CHECK` to user. `NULLIFY` checks out-of-bounds indices and sets them to null rows, while `DONT_CHECK` skips all checks. Using `DONT_CHECK` should yield higher performance, given `gather_map` contains only valid indices. Note that the negative index (wrap-arounds) policy is unchanged. When gather map dtype is `signed`, wrap-around is applied. A new Cython binding to `cudf::minmax`, used for Cython `gather` bound checking is added. Will also close #6731 Authors: - Michael Wang - Michael Wang Approvers: - null - Devavret Makkar - Ashwin Srinath - Keith Kraus - Jake Hemstad URL: https://github.com/rapidsai/cudf/pull/6875 --- CHANGELOG.md | 1 + cpp/include/cudf/copying.hpp | 26 ++++++++++++---- cpp/include/cudf/detail/copy_if.cuh | 8 +++-- cpp/include/cudf/detail/gather.cuh | 32 ++++++++++++-------- cpp/include/cudf/detail/gather.hpp | 16 +++++----- cpp/src/copying/gather.cu | 29 +++--------------- cpp/src/copying/sample.cu | 4 +-- cpp/src/dictionary/add_keys.cu | 2 +- cpp/src/dictionary/decode.cu | 2 +- cpp/src/dictionary/remove_keys.cu | 3 +- cpp/src/filling/repeat.cu | 5 +-- cpp/src/groupby/groupby.cu | 2 +- cpp/src/groupby/hash/groupby.cu | 21 +++++++++---- cpp/src/groupby/sort/group_argmax.cu | 4 +-- cpp/src/groupby/sort/group_argmin.cu | 4 +-- cpp/src/groupby/sort/group_nth_element.cu | 9 ++++-- cpp/src/groupby/sort/groupby.cu | 8 ++--- cpp/src/groupby/sort/sort_helper.cu | 15 ++++++--- cpp/src/join/hash_join.cu | 14 +++++---- cpp/src/join/semi_join.cu | 8 +++-- cpp/src/lists/extract.cu | 3 +- cpp/src/partitioning/round_robin.cu | 7 +++-- cpp/src/quantiles/quantiles.cu | 7 ++++- cpp/src/reshape/tile.cu | 3 +- cpp/src/rolling/rolling.cu | 2 +- cpp/src/sort/sort.cu | 2 +- cpp/src/stream_compaction/drop_duplicates.cu | 2 +- cpp/src/strings/copying/copying.cu | 2 +- cpp/src/strings/sorting/sorting.cu | 2 +- cpp/src/transform/encode.cu | 2 +- cpp/tests/copying/detail_gather_tests.cu | 2 +- cpp/tests/copying/gather_list_tests.cu | 2 +- cpp/tests/copying/gather_str_tests.cu | 17 +++++------ python/cudf/cudf/_lib/copying.pyx | 18 +++++++++-- python/cudf/cudf/_lib/cpp/copying.pxd | 6 +++- python/cudf/cudf/_lib/cpp/reduce.pxd | 6 ++++ python/cudf/cudf/_lib/reduce.pyx | 29 ++++++++++++++++-- 37 files changed, 207 insertions(+), 118 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 55cd081470e..b0f23c99896 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -53,6 +53,7 @@ - PR #6471 Replace index type-dispatch call with indexalator in cudf::strings::substring - PR #6485 Add File IO to cuIO benchmarks - PR #6504 Update Java bindings version to 0.17-SNAPSHOT +- PR #6875 Remove bounds check for `cudf::gather` - PR #6489 Add `AVRO` fuzz tests with varying function parameters - PR #6540 Add dictionary support to `cudf::unary_operation` - PR #6537 Refactor ORC timezone diff --git a/cpp/include/cudf/copying.hpp b/cpp/include/cudf/copying.hpp index 7b7f4c03afb..5e3b3673053 100644 --- a/cpp/include/cudf/copying.hpp +++ b/cpp/include/cudf/copying.hpp @@ -33,6 +33,18 @@ namespace cudf { * @brief Column APIs for gather, scatter, split, slice, etc. */ +/** + * @brief Policy to account for possible out-of-bounds indices + * + * `NULLIFY` means to nullify output values corresponding to out-of-bounds gather_map values. + * `DONT_CHECK` means do not check whether the indices are out-of-bounds, for better performance. + */ + +enum class out_of_bounds_policy : bool { + NULLIFY, /// Output values corresponding to out-of-bounds indices are null + DONT_CHECK /// No bounds checking is performed, better performance +}; + /** * @brief Gathers the specified rows (including null values) of a set of columns. * @@ -49,22 +61,24 @@ namespace cudf { * For dictionary columns, the keys column component is copied and not trimmed * if the gather results in abandoned key elements. * - * @throws cudf::logic_error if `check_bounds == true` and an index exists in - * `gather_map` outside the range `[-n, n)`, where `n` is the number of rows in - * the source table. If `check_bounds == false`, the behavior is undefined. + * @throws cudf::logic_error if gather_map contains null values. * * @param[in] source_table The input columns whose rows will be gathered * @param[in] gather_map View into a non-nullable column of integral indices that maps the * rows in the source columns to rows in the destination columns. - * @param[in] check_bounds Optionally perform bounds checking on the values - * of `gather_map` and throw an error if any of its values are out of bounds. + * @param[in] bounds_policy Policy to apply to account for possible out-of-bounds indices + * `DONT_CHECK` skips all bounds checking for gather map values. `NULLIFY` coerces rows that + * corresponds to out-of-bounds indices in the gather map to be null elements. Callers should + * use `DONT_CHECK` when they are certain that the gather_map contains only valid indices for + * better performance. If `policy` is set to `DONT_CHECK` and there are out-of-bounds indices + * in the gather map, the behavior is undefined. Defaults to `DONT_CHECK`. * @param[in] mr Device memory resource used to allocate the returned table's device memory * @return std::unique_ptr Result of the gather */ std::unique_ptr
gather( table_view const& source_table, column_view const& gather_map, - bool check_bounds = false, + out_of_bounds_policy bounds_policy = out_of_bounds_policy::DONT_CHECK, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** diff --git a/cpp/include/cudf/detail/copy_if.cuh b/cpp/include/cudf/detail/copy_if.cuh index b6c40a3a86e..0b960deb826 100644 --- a/cpp/include/cudf/detail/copy_if.cuh +++ b/cpp/include/cudf/detail/copy_if.cuh @@ -268,8 +268,12 @@ struct scatter_gather_functor { indices.begin(), filter); - auto output_table = cudf::detail::gather( - cudf::table_view{{input}}, indices.begin(), indices.end(), false, stream, mr); + auto output_table = cudf::detail::gather(cudf::table_view{{input}}, + indices.begin(), + indices.end(), + cudf::out_of_bounds_policy::DONT_CHECK, + stream, + mr); // There will be only one column return std::make_unique(std::move(output_table->get_column(0))); diff --git a/cpp/include/cudf/detail/gather.cuh b/cpp/include/cudf/detail/gather.cuh index 8e586b231bc..4bc99a15850 100644 --- a/cpp/include/cudf/detail/gather.cuh +++ b/cpp/include/cudf/detail/gather.cuh @@ -610,7 +610,12 @@ struct column_gatherer_impl { * the source columns to rows in the destination columns * @param[in] gather_map_end End of iterator range of integer indices that map the rows in the * source columns to rows in the destination columns - * @param[in] nullify_out_of_bounds Nullify values in `gather_map` that are out of bounds. + * @param[in] bounds_policy Policy to apply to account for possible out-of-bound indices + * `DONT_CHECK` skips all bound checking for gather map values. `NULLIFY` coerces rows that + * corresponds to out-of-bound indices in the gather map to be null elements. Callers should + * use `DONT_CHECK` when they are certain that the gather_map contains only valid indices for + * better performance. In case there are out-of-bound indices in the gather map, the behavior + * is undefined. Defaults to `DONT_CHECK`. * @param[in] mr Device memory resource used to allocate the returned table's device memory * @param[in] stream CUDA stream used for device memory operations and kernel launches. * @return cudf::table Result of the gather @@ -620,7 +625,7 @@ std::unique_ptr
gather( table_view const& source_table, MapIterator gather_map_begin, MapIterator gather_map_end, - bool nullify_out_of_bounds = false, + out_of_bounds_policy bounds_policy = out_of_bounds_policy::DONT_CHECK, rmm::cuda_stream_view stream = rmm::cuda_stream_default, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) { @@ -630,18 +635,21 @@ std::unique_ptr
gather( for (auto const& source_column : source_table) { // The data gather for n columns will be put on the first n streams - destination_columns.push_back(cudf::type_dispatcher(source_column.type(), - column_gatherer{}, - source_column, - gather_map_begin, - gather_map_end, - nullify_out_of_bounds, - stream, - mr)); + destination_columns.push_back( + cudf::type_dispatcher(source_column.type(), + column_gatherer{}, + source_column, + gather_map_begin, + gather_map_end, + bounds_policy == out_of_bounds_policy::NULLIFY, + stream, + mr)); } - auto const op = - nullify_out_of_bounds ? gather_bitmask_op::NULLIFY : gather_bitmask_op::DONT_CHECK; + gather_bitmask_op const op = bounds_policy == out_of_bounds_policy::NULLIFY + ? gather_bitmask_op::NULLIFY + : gather_bitmask_op::DONT_CHECK; + gather_bitmask(source_table, gather_map_begin, destination_columns, op, stream, mr); return std::make_unique
(std::move(destination_columns)); diff --git a/cpp/include/cudf/detail/gather.hpp b/cpp/include/cudf/detail/gather.hpp index adace7e27f8..268c4878444 100644 --- a/cpp/include/cudf/detail/gather.hpp +++ b/cpp/include/cudf/detail/gather.hpp @@ -18,6 +18,7 @@ #include #include +#include #include #include @@ -25,9 +26,8 @@ #include namespace cudf { -namespace detail { -enum class out_of_bounds_policy : int8_t { FAIL, NULLIFY, IGNORE }; +namespace detail { enum class negative_index_policy : bool { ALLOWED, NOT_ALLOWED }; @@ -49,11 +49,11 @@ enum class negative_index_policy : bool { ALLOWED, NOT_ALLOWED }; * @param[in] source_table The input columns whose rows will be gathered * @param[in] gather_map View into a non-nullable column of integral indices that maps the * rows in the source columns to rows in the destination columns. - * @param[in] out_of_bounds_policy How to treat out of bounds indices. FAIL: check `gather_map` - * values and throw an exception if any are out of bounds. `NULLIFY` means to nullify output values - * corresponding to out-of-bounds gather_map values. `IGNORE` means to ignore values in - * `gather_map` that are out of bounds. `IGNORE` is incompatible with `negative_index_policy == - * ALLOW`. + * @param[in] bounds_policy How to treat out-of-bounds indices. `NULLIFY` coerces rows that + * correspond to out-of-bounds indices in the gather map to be null elements. For better + * performance, use `DONT_CHECK` when the `gather_map` is known to contain only valid + * indices. If `policy` is set to `DONT_CHECK` and there are out-of-bounds indices in `gather_map`, + * the behavior is undefined. * @param[in] negative_index_policy Interpret each negative index `i` in the * gathermap as the positive index `i+num_source_rows`. * @param[in] mr Device memory resource used to allocate the returned table's device memory @@ -63,7 +63,7 @@ enum class negative_index_policy : bool { ALLOWED, NOT_ALLOWED }; std::unique_ptr
gather( table_view const& source_table, column_view const& gather_map, - out_of_bounds_policy bounds, + out_of_bounds_policy bounds_policy, negative_index_policy neg_indices, rmm::cuda_stream_view stream = rmm::cuda_stream_default, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); diff --git a/cpp/src/copying/gather.cu b/cpp/src/copying/gather.cu index 4e186c00ac3..87721184536 100644 --- a/cpp/src/copying/gather.cu +++ b/cpp/src/copying/gather.cu @@ -32,7 +32,7 @@ namespace detail { std::unique_ptr
gather(table_view const& source_table, column_view const& gather_map, - out_of_bounds_policy bounds, + out_of_bounds_policy bounds_policy, negative_index_policy neg_indices, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) @@ -43,19 +43,6 @@ std::unique_ptr
gather(table_view const& source_table, auto map_begin = indexalator_factory::make_input_iterator(gather_map); auto map_end = map_begin + gather_map.size(); - if (bounds == out_of_bounds_policy::FAIL) { - cudf::size_type begin = - neg_indices == negative_index_policy::ALLOWED ? -source_table.num_rows() : 0; - cudf::size_type end = source_table.num_rows(); - CUDF_EXPECTS(gather_map.size() == thrust::count_if(rmm::exec_policy(stream)->on(stream.value()), - map_begin, - map_end, - [begin, end] __device__(size_type index) { - return ((index >= begin) && (index < end)); - }), - "Index out of bounds."); - } - if (neg_indices == negative_index_policy::ALLOWED) { cudf::size_type n_rows = source_table.num_rows(); auto idx_converter = [n_rows] __device__(size_type in) { @@ -64,19 +51,18 @@ std::unique_ptr
gather(table_view const& source_table, return gather(source_table, thrust::make_transform_iterator(map_begin, idx_converter), thrust::make_transform_iterator(map_end, idx_converter), - bounds == out_of_bounds_policy::IGNORE, + bounds_policy, stream, mr); } - return gather( - source_table, map_begin, map_end, bounds == out_of_bounds_policy::IGNORE, stream, mr); + return gather(source_table, map_begin, map_end, bounds_policy, stream, mr); } } // namespace detail std::unique_ptr
gather(table_view const& source_table, column_view const& gather_map, - bool check_bounds, + out_of_bounds_policy bounds_policy, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); @@ -85,12 +71,7 @@ std::unique_ptr
gather(table_view const& source_table, : detail::negative_index_policy::ALLOWED; return detail::gather( - source_table, - gather_map, - check_bounds ? detail::out_of_bounds_policy::FAIL : detail::out_of_bounds_policy::NULLIFY, - index_policy, - rmm::cuda_stream_default, - mr); + source_table, gather_map, bounds_policy, index_policy, rmm::cuda_stream_default, mr); } } // namespace cudf diff --git a/cpp/src/copying/sample.cu b/cpp/src/copying/sample.cu index 15dc3565da4..fc82989a7c2 100644 --- a/cpp/src/copying/sample.cu +++ b/cpp/src/copying/sample.cu @@ -60,7 +60,7 @@ std::unique_ptr
sample(table_view const& input, thrust::make_transform_iterator(thrust::counting_iterator(0), RandomGen); auto end = thrust::make_transform_iterator(thrust::counting_iterator(n), RandomGen); - return detail::gather(input, begin, end, false, stream, mr); + return detail::gather(input, begin, end, out_of_bounds_policy::DONT_CHECK, stream, mr); } else { auto gather_map = make_numeric_column( data_type{type_id::INT32}, num_rows, mask_state::UNALLOCATED, stream.value()); @@ -77,7 +77,7 @@ std::unique_ptr
sample(table_view const& input, return detail::gather(input, gather_map_view.begin(), gather_map_view.end(), - false, + out_of_bounds_policy::DONT_CHECK, stream, mr); } diff --git a/cpp/src/dictionary/add_keys.cu b/cpp/src/dictionary/add_keys.cu index daf1bb76916..50bdf219a85 100644 --- a/cpp/src/dictionary/add_keys.cu +++ b/cpp/src/dictionary/add_keys.cu @@ -90,7 +90,7 @@ std::unique_ptr add_keys( // and the corresponding index is therefore invalid/undefined auto table_indices = cudf::detail::gather(table_view{{map_indices->view()}}, indices_view, - cudf::detail::out_of_bounds_policy::IGNORE, + cudf::out_of_bounds_policy::NULLIFY, cudf::detail::negative_index_policy::NOT_ALLOWED, stream, mr) diff --git a/cpp/src/dictionary/decode.cu b/cpp/src/dictionary/decode.cu index 3822edfc9ef..d701fb689cb 100644 --- a/cpp/src/dictionary/decode.cu +++ b/cpp/src/dictionary/decode.cu @@ -47,7 +47,7 @@ std::unique_ptr decode(dictionary_column_view const& source, // use gather to create the output column -- use ignore_out_of_bounds=true auto table_column = cudf::detail::gather(table_view{{source.keys()}}, indices, - cudf::detail::out_of_bounds_policy::IGNORE, + cudf::out_of_bounds_policy::NULLIFY, cudf::detail::negative_index_policy::NOT_ALLOWED, stream, mr) diff --git a/cpp/src/dictionary/remove_keys.cu b/cpp/src/dictionary/remove_keys.cu index b36b110b13f..76d51878d3b 100644 --- a/cpp/src/dictionary/remove_keys.cu +++ b/cpp/src/dictionary/remove_keys.cu @@ -114,12 +114,13 @@ std::unique_ptr remove_keys_fn( // Example: gather([0,max,1,max,2],[4,0,3,1,2,2,2,4,0]) => [2,0,max,max,1,1,1,2,0] auto table_indices = cudf::detail::gather(table_view{{map_indices->view()}}, indices_view, - cudf::detail::out_of_bounds_policy::NULLIFY, + cudf::out_of_bounds_policy::NULLIFY, cudf::detail::negative_index_policy::NOT_ALLOWED, stream, mr) ->release(); std::unique_ptr indices_column(std::move(table_indices.front())); + indices_column->set_null_mask(rmm::device_buffer{}, 0); // compute new nulls -- merge the existing nulls with the newly created ones (value<0) auto const offset = dictionary_column.offset(); diff --git a/cpp/src/filling/repeat.cu b/cpp/src/filling/repeat.cu index 46b16ac0949..30b05dfc3df 100644 --- a/cpp/src/filling/repeat.cu +++ b/cpp/src/filling/repeat.cu @@ -135,7 +135,8 @@ std::unique_ptr
repeat(table_view const& input_table, thrust::make_counting_iterator(output_size), indices.begin()); - return gather(input_table, indices.begin(), indices.end(), false, stream, mr); + return gather( + input_table, indices.begin(), indices.end(), out_of_bounds_policy::DONT_CHECK, stream, mr); } std::unique_ptr
repeat(table_view const& input_table, @@ -155,7 +156,7 @@ std::unique_ptr
repeat(table_view const& input_table, thrust::make_counting_iterator(0), [count] __device__(auto i) { return i / count; }); auto map_end = map_begin + output_size; - return gather(input_table, map_begin, map_end, false, stream, mr); + return gather(input_table, map_begin, map_end, out_of_bounds_policy::DONT_CHECK, stream, mr); } } // namespace detail diff --git a/cpp/src/groupby/groupby.cu b/cpp/src/groupby/groupby.cu index 3df6e0ece85..862cbcf5e82 100644 --- a/cpp/src/groupby/groupby.cu +++ b/cpp/src/groupby/groupby.cu @@ -149,7 +149,7 @@ groupby::groups groupby::get_groups(table_view values, rmm::mr::device_memory_re if (values.num_columns()) { grouped_values = cudf::detail::gather(values, helper().key_sort_order(), - cudf::detail::out_of_bounds_policy::NULLIFY, + cudf::out_of_bounds_policy::DONT_CHECK, cudf::detail::negative_index_policy::NOT_ALLOWED, rmm::cuda_stream_default, mr); diff --git a/cpp/src/groupby/hash/groupby.cu b/cpp/src/groupby/hash/groupby.cu index 44bed810059..597e58b211b 100644 --- a/cpp/src/groupby/hash/groupby.cu +++ b/cpp/src/groupby/hash/groupby.cu @@ -21,6 +21,7 @@ #include #include #include +#include #include #include #include @@ -139,8 +140,12 @@ class hash_compound_agg_finalizer final : public cudf::detail::aggregation_final auto to_dense_agg_result(cudf::aggregation const& agg) { auto s = sparse_results->get_result(col_idx, agg); - auto dense_result_table = cudf::detail::gather( - table_view({s}), gather_map.begin(), gather_map.begin() + map_size, false, stream, mr); + auto dense_result_table = cudf::detail::gather(table_view({s}), + gather_map.begin(), + gather_map.begin() + map_size, + out_of_bounds_policy::DONT_CHECK, + stream, + mr); return std::move(dense_result_table->release()[0]); } @@ -161,8 +166,8 @@ class hash_compound_agg_finalizer final : public cudf::detail::aggregation_final auto gather_argminmax = cudf::detail::gather(table_view({col}), null_removed_map, - arg_result->nullable() ? cudf::detail::out_of_bounds_policy::IGNORE - : cudf::detail::out_of_bounds_policy::NULLIFY, + arg_result->nullable() ? cudf::out_of_bounds_policy::NULLIFY + : cudf::out_of_bounds_policy::DONT_CHECK, cudf::detail::negative_index_policy::NOT_ALLOWED, stream, mr); @@ -538,8 +543,12 @@ std::unique_ptr
groupby_null_templated(table_view const& keys, stream, mr); - return cudf::detail::gather( - keys, gather_map.begin(), gather_map.begin() + map_size, false, stream, mr); + return cudf::detail::gather(keys, + gather_map.begin(), + gather_map.begin() + map_size, + out_of_bounds_policy::DONT_CHECK, + stream, + mr); } } // namespace diff --git a/cpp/src/groupby/sort/group_argmax.cu b/cpp/src/groupby/sort/group_argmax.cu index a22e7619694..94b5d3817a7 100644 --- a/cpp/src/groupby/sort/group_argmax.cu +++ b/cpp/src/groupby/sort/group_argmax.cu @@ -54,8 +54,8 @@ std::unique_ptr group_argmax(column_view const& values, auto result_table = cudf::detail::gather(table_view({key_sort_order}), null_removed_indices, - indices->nullable() ? cudf::detail::out_of_bounds_policy::IGNORE - : cudf::detail::out_of_bounds_policy::NULLIFY, + indices->nullable() ? cudf::out_of_bounds_policy::NULLIFY + : cudf::out_of_bounds_policy::DONT_CHECK, cudf::detail::negative_index_policy::NOT_ALLOWED, stream, mr); diff --git a/cpp/src/groupby/sort/group_argmin.cu b/cpp/src/groupby/sort/group_argmin.cu index 6cdcd7cd94a..11a350ae1c4 100644 --- a/cpp/src/groupby/sort/group_argmin.cu +++ b/cpp/src/groupby/sort/group_argmin.cu @@ -54,8 +54,8 @@ std::unique_ptr group_argmin(column_view const& values, auto result_table = cudf::detail::gather(table_view({key_sort_order}), null_removed_indices, - indices->nullable() ? cudf::detail::out_of_bounds_policy::IGNORE - : cudf::detail::out_of_bounds_policy::NULLIFY, + indices->nullable() ? cudf::out_of_bounds_policy::NULLIFY + : cudf::out_of_bounds_policy::DONT_CHECK, cudf::detail::negative_index_policy::NOT_ALLOWED, stream, mr); diff --git a/cpp/src/groupby/sort/group_nth_element.cu b/cpp/src/groupby/sort/group_nth_element.cu index e7e947b65fc..48cbde535dc 100644 --- a/cpp/src/groupby/sort/group_nth_element.cu +++ b/cpp/src/groupby/sort/group_nth_element.cu @@ -18,6 +18,7 @@ #include #include #include +#include #include #include #include @@ -105,8 +106,12 @@ std::unique_ptr group_nth_element(column_view const &values, return (bitmask_iterator[i] && intra_group_index[i] == nth); }); } - auto output_table = cudf::detail::gather( - table_view{{values}}, nth_index.begin(), nth_index.end(), true, stream, mr); + auto output_table = cudf::detail::gather(table_view{{values}}, + nth_index.begin(), + nth_index.end(), + out_of_bounds_policy::NULLIFY, + stream, + mr); if (!output_table->get_column(0).has_nulls()) output_table->get_column(0).set_null_mask({}, 0); return std::make_unique(std::move(output_table->get_column(0))); } diff --git a/cpp/src/groupby/sort/groupby.cu b/cpp/src/groupby/sort/groupby.cu index 964d973dcc5..a88e45c4c7f 100644 --- a/cpp/src/groupby/sort/groupby.cu +++ b/cpp/src/groupby/sort/groupby.cu @@ -203,8 +203,8 @@ void store_result_functor::operator()(aggregation const& agg) auto transformed_result = cudf::detail::gather(table_view({values}), null_removed_map, - argmin_result.nullable() ? cudf::detail::out_of_bounds_policy::IGNORE - : cudf::detail::out_of_bounds_policy::NULLIFY, + argmin_result.nullable() ? cudf::out_of_bounds_policy::NULLIFY + : cudf::out_of_bounds_policy::DONT_CHECK, cudf::detail::negative_index_policy::NOT_ALLOWED, stream, mr); @@ -240,8 +240,8 @@ void store_result_functor::operator()(aggregation const& agg) auto transformed_result = cudf::detail::gather(table_view({values}), null_removed_map, - argmax_result.nullable() ? cudf::detail::out_of_bounds_policy::IGNORE - : cudf::detail::out_of_bounds_policy::NULLIFY, + argmax_result.nullable() ? cudf::out_of_bounds_policy::NULLIFY + : cudf::out_of_bounds_policy::DONT_CHECK, cudf::detail::negative_index_policy::NOT_ALLOWED, stream, mr); diff --git a/cpp/src/groupby/sort/sort_helper.cu b/cpp/src/groupby/sort/sort_helper.cu index 595efc8198d..7788a6d867b 100644 --- a/cpp/src/groupby/sort/sort_helper.cu +++ b/cpp/src/groupby/sort/sort_helper.cu @@ -15,6 +15,7 @@ */ #include +#include #include #include #include @@ -276,7 +277,7 @@ sort_groupby_helper::column_ptr sort_groupby_helper::sorted_values( auto sorted_values_table = cudf::detail::gather(table_view({values}), gather_map, - cudf::detail::out_of_bounds_policy::NULLIFY, + cudf::out_of_bounds_policy::DONT_CHECK, cudf::detail::negative_index_policy::NOT_ALLOWED, stream, mr); @@ -291,7 +292,7 @@ sort_groupby_helper::column_ptr sort_groupby_helper::grouped_values( auto grouped_values_table = cudf::detail::gather(table_view({values}), gather_map, - cudf::detail::out_of_bounds_policy::NULLIFY, + cudf::out_of_bounds_policy::DONT_CHECK, cudf::detail::negative_index_policy::NOT_ALLOWED, stream, mr); @@ -307,8 +308,12 @@ std::unique_ptr
sort_groupby_helper::unique_keys(rmm::cuda_stream_view st auto gather_map_it = thrust::make_transform_iterator( group_offsets().begin(), [idx_data] __device__(size_type i) { return idx_data[i]; }); - return cudf::detail::gather( - _keys, gather_map_it, gather_map_it + num_groups(), false, stream, mr); + return cudf::detail::gather(_keys, + gather_map_it, + gather_map_it + num_groups(), + out_of_bounds_policy::DONT_CHECK, + stream, + mr); } std::unique_ptr
sort_groupby_helper::sorted_keys(rmm::cuda_stream_view stream, @@ -316,7 +321,7 @@ std::unique_ptr
sort_groupby_helper::sorted_keys(rmm::cuda_stream_view st { return cudf::detail::gather(_keys, key_sort_order(), - cudf::detail::out_of_bounds_policy::NULLIFY, + cudf::out_of_bounds_policy::DONT_CHECK, cudf::detail::negative_index_policy::NOT_ALLOWED, stream, mr); diff --git a/cpp/src/join/hash_join.cu b/cpp/src/join/hash_join.cu index 67b9d3436d8..a88432fe88d 100644 --- a/cpp/src/join/hash_join.cu +++ b/cpp/src/join/hash_join.cu @@ -407,7 +407,9 @@ std::pair, std::unique_ptr
> construct_join_output_ std::vector build_noncommon_col = non_common_column_indices(build.num_columns(), build_common_col); - bool const nullify_out_of_bounds{JoinKind != join_kind::INNER_JOIN}; + out_of_bounds_policy const bounds_policy = JoinKind != join_kind::INNER_JOIN + ? out_of_bounds_policy::NULLIFY + : out_of_bounds_policy::DONT_CHECK; std::unique_ptr
common_table = std::make_unique
(); // Construct the joined columns @@ -418,13 +420,13 @@ std::pair, std::unique_ptr
> construct_join_output_ auto common_from_build = detail::gather(build.select(build_common_col), complement_indices.second.begin(), complement_indices.second.end(), - nullify_out_of_bounds, + bounds_policy, stream, rmm::mr::get_current_device_resource()); auto common_from_probe = detail::gather(probe.select(probe_common_col), joined_indices.first.begin(), joined_indices.first.end(), - nullify_out_of_bounds, + bounds_policy, stream, rmm::mr::get_current_device_resource()); common_table = cudf::detail::concatenate( @@ -436,7 +438,7 @@ std::pair, std::unique_ptr
> construct_join_output_ common_table = detail::gather(probe.select(probe_common_col), joined_indices.first.begin(), joined_indices.first.end(), - nullify_out_of_bounds, + bounds_policy, stream, mr); } @@ -446,14 +448,14 @@ std::pair, std::unique_ptr
> construct_join_output_ std::unique_ptr
probe_table = detail::gather(probe.select(probe_noncommon_col), joined_indices.first.begin(), joined_indices.first.end(), - nullify_out_of_bounds, + bounds_policy, stream, mr); std::unique_ptr
build_table = detail::gather(build.select(build_noncommon_col), joined_indices.second.begin(), joined_indices.second.end(), - nullify_out_of_bounds, + bounds_policy, stream, mr); diff --git a/cpp/src/join/semi_join.cu b/cpp/src/join/semi_join.cu index 6df329243ed..c3618afc0c2 100644 --- a/cpp/src/join/semi_join.cu +++ b/cpp/src/join/semi_join.cu @@ -158,8 +158,12 @@ std::unique_ptr left_semi_anti_join( // rebuild left table for call to gather auto const left_updated = scatter_columns(left_selected, left_on, left); - return cudf::detail::gather( - left_updated.select(return_columns), gather_map.begin(), gather_map_end, false, stream, mr); + return cudf::detail::gather(left_updated.select(return_columns), + gather_map.begin(), + gather_map_end, + out_of_bounds_policy::DONT_CHECK, + stream, + mr); } } // namespace detail diff --git a/cpp/src/lists/extract.cu b/cpp/src/lists/extract.cu index 342bd006ea2..3ba02ac555f 100644 --- a/cpp/src/lists/extract.cu +++ b/cpp/src/lists/extract.cu @@ -15,6 +15,7 @@ */ #include #include +#include #include #include @@ -98,7 +99,7 @@ std::unique_ptr extract_list_element(lists_column_view lists_column, auto result = cudf::detail::gather(table_view({child_column}), d_gather_map, d_gather_map + gather_map->size(), - true, // nullify-out-of-bounds + out_of_bounds_policy::NULLIFY, // nullify-out-of-bounds stream, mr) ->release(); diff --git a/cpp/src/partitioning/round_robin.cu b/cpp/src/partitioning/round_robin.cu index 6367293a9d3..a28978a294d 100644 --- a/cpp/src/partitioning/round_robin.cu +++ b/cpp/src/partitioning/round_robin.cu @@ -100,7 +100,7 @@ std::pair, std::vector> degenerate auto uniq_tbl = cudf::detail::gather(input, rotated_iter_begin, rotated_iter_begin + nrows, // map - false, + cudf::out_of_bounds_policy::DONT_CHECK, stream, mr); @@ -135,7 +135,7 @@ std::pair, std::vector> degenerate auto uniq_tbl = cudf::detail::gather(input, d_row_indices.begin(), d_row_indices.end(), // map - false, + cudf::out_of_bounds_policy::DONT_CHECK, stream, mr); @@ -253,7 +253,8 @@ std::pair, std::vector> round_robin_part return num_partitions * index_within_partition + partition_index; }); - auto uniq_tbl = cudf::detail::gather(input, iter_begin, iter_begin + nrows, false, stream, mr); + auto uniq_tbl = cudf::detail::gather( + input, iter_begin, iter_begin + nrows, cudf::out_of_bounds_policy::DONT_CHECK, stream, mr); auto ret_pair = std::make_pair(std::move(uniq_tbl), std::vector(num_partitions)); // this has the effect of rotating the set of partition sizes diff --git a/cpp/src/quantiles/quantiles.cu b/cpp/src/quantiles/quantiles.cu index 51e71104cac..715c8c211d7 100644 --- a/cpp/src/quantiles/quantiles.cu +++ b/cpp/src/quantiles/quantiles.cu @@ -47,7 +47,12 @@ std::unique_ptr
quantiles(table_view const& input, auto quantile_idx_iter = thrust::make_transform_iterator(q_device.begin(), quantile_idx_lookup); - return detail::gather(input, quantile_idx_iter, quantile_idx_iter + q.size(), false, stream, mr); + return detail::gather(input, + quantile_idx_iter, + quantile_idx_iter + q.size(), + out_of_bounds_policy::DONT_CHECK, + stream, + mr); } } // namespace detail diff --git a/cpp/src/reshape/tile.cu b/cpp/src/reshape/tile.cu index e1c665cf8dd..b4b9ac42739 100644 --- a/cpp/src/reshape/tile.cu +++ b/cpp/src/reshape/tile.cu @@ -54,7 +54,8 @@ std::unique_ptr
tile(const table_view &in, auto counting_it = thrust::make_counting_iterator(0); auto tiled_it = thrust::make_transform_iterator(counting_it, tile_functor{in_num_rows}); - return detail::gather(in, tiled_it, tiled_it + out_num_rows, false, stream, mr); + return detail::gather( + in, tiled_it, tiled_it + out_num_rows, out_of_bounds_policy::DONT_CHECK, stream, mr); } } // namespace detail diff --git a/cpp/src/rolling/rolling.cu b/cpp/src/rolling/rolling.cu index 8856b64428e..19d3b70f6c1 100644 --- a/cpp/src/rolling/rolling.cu +++ b/cpp/src/rolling/rolling.cu @@ -701,7 +701,7 @@ struct rolling_window_launcher { // and that's why nullify_out_of_bounds/ignore_out_of_bounds is true. auto output_table = detail::gather(table_view{{input}}, output->view(), - detail::out_of_bounds_policy::IGNORE, + cudf::out_of_bounds_policy::NULLIFY, detail::negative_index_policy::NOT_ALLOWED, stream, mr); diff --git a/cpp/src/sort/sort.cu b/cpp/src/sort/sort.cu index 028796d59cb..d94f09daf52 100644 --- a/cpp/src/sort/sort.cu +++ b/cpp/src/sort/sort.cu @@ -49,7 +49,7 @@ std::unique_ptr
sort_by_key(table_view const& values, return detail::gather(values, sorted_order->view(), - detail::out_of_bounds_policy::NULLIFY, + out_of_bounds_policy::DONT_CHECK, detail::negative_index_policy::NOT_ALLOWED, stream, mr); diff --git a/cpp/src/stream_compaction/drop_duplicates.cu b/cpp/src/stream_compaction/drop_duplicates.cu index 3e08b1d1b0c..f8f3445ca7c 100644 --- a/cpp/src/stream_compaction/drop_duplicates.cu +++ b/cpp/src/stream_compaction/drop_duplicates.cu @@ -198,7 +198,7 @@ std::unique_ptr
drop_duplicates(table_view const& input, // run gather operation to establish new order return detail::gather(input, unique_indices_view, - detail::out_of_bounds_policy::NULLIFY, + out_of_bounds_policy::DONT_CHECK, detail::negative_index_policy::NOT_ALLOWED, stream, mr); diff --git a/cpp/src/strings/copying/copying.cu b/cpp/src/strings/copying/copying.cu index 58bcdccf578..ae8af63cb05 100644 --- a/cpp/src/strings/copying/copying.cu +++ b/cpp/src/strings/copying/copying.cu @@ -75,7 +75,7 @@ std::unique_ptr copy_slice(strings_column_view const& strings, // build a new strings column from the indices auto sliced_table = cudf::detail::gather(table_view{{strings.parent()}}, indices_view, - cudf::detail::out_of_bounds_policy::NULLIFY, + cudf::out_of_bounds_policy::DONT_CHECK, cudf::detail::negative_index_policy::NOT_ALLOWED, stream, mr) diff --git a/cpp/src/strings/sorting/sorting.cu b/cpp/src/strings/sorting/sorting.cu index 3d78024064e..0d59ad053e5 100644 --- a/cpp/src/strings/sorting/sorting.cu +++ b/cpp/src/strings/sorting/sorting.cu @@ -67,7 +67,7 @@ std::unique_ptr sort(strings_column_view strings, // now build a new strings column from the indices auto table_sorted = cudf::detail::gather(table_view{{strings.parent()}}, indices_view, - cudf::detail::out_of_bounds_policy::NULLIFY, + cudf::out_of_bounds_policy::DONT_CHECK, cudf::detail::negative_index_policy::NOT_ALLOWED, stream, mr) diff --git a/cpp/src/transform/encode.cu b/cpp/src/transform/encode.cu index 5abe4e1aaf7..8ddb5471cf9 100644 --- a/cpp/src/transform/encode.cu +++ b/cpp/src/transform/encode.cu @@ -76,7 +76,7 @@ std::pair, std::unique_ptr> encode( keys_table = cudf::detail::gather(keys_table->view(), gather_map_column, - cudf::detail::out_of_bounds_policy::FAIL, + cudf::out_of_bounds_policy::DONT_CHECK, cudf::detail::negative_index_policy::NOT_ALLOWED, stream, mr); diff --git a/cpp/tests/copying/detail_gather_tests.cu b/cpp/tests/copying/detail_gather_tests.cu index 97c0498abc4..d14036516d6 100644 --- a/cpp/tests/copying/detail_gather_tests.cu +++ b/cpp/tests/copying/detail_gather_tests.cu @@ -86,7 +86,7 @@ TYPED_TEST(GatherTest, GatherDetailInvalidIndexTest) std::unique_ptr result = cudf::detail::gather(source_table, gather_map, - cudf::detail::out_of_bounds_policy::IGNORE, + cudf::out_of_bounds_policy::NULLIFY, cudf::detail::negative_index_policy::NOT_ALLOWED); auto expect_data = diff --git a/cpp/tests/copying/gather_list_tests.cu b/cpp/tests/copying/gather_list_tests.cu index 03cfef3783e..3369472687e 100644 --- a/cpp/tests/copying/gather_list_tests.cu +++ b/cpp/tests/copying/gather_list_tests.cu @@ -265,7 +265,7 @@ TYPED_TEST(GatherTestListTyped, GatherDetailInvalidIndex) cudf::table_view source_table({list}); auto results = cudf::detail::gather(source_table, gather_map, - cudf::detail::out_of_bounds_policy::IGNORE, + cudf::out_of_bounds_policy::NULLIFY, cudf::detail::negative_index_policy::NOT_ALLOWED); std::vector expected_validity{1, 0, 0, 1}; diff --git a/cpp/tests/copying/gather_str_tests.cu b/cpp/tests/copying/gather_str_tests.cu index a8b8acb030a..75cea81c950 100644 --- a/cpp/tests/copying/gather_str_tests.cu +++ b/cpp/tests/copying/gather_str_tests.cu @@ -58,7 +58,7 @@ TEST_F(GatherTestStr, Gather) cudf::test::fixed_width_column_wrapper gather_map(h_map.begin(), h_map.end()); auto results = cudf::detail::gather(source_table, gather_map, - cudf::detail::out_of_bounds_policy::IGNORE, + cudf::out_of_bounds_policy::NULLIFY, cudf::detail::negative_index_policy::NOT_ALLOWED); std::vector h_expected; @@ -78,7 +78,7 @@ TEST_F(GatherTestStr, Gather) CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view().column(0), expected); } -TEST_F(GatherTestStr, GatherIgnoreOutOfBounds) +TEST_F(GatherTestStr, GatherDontCheckOutOfBounds) { std::vector h_strings{"eee", "bb", "", "aa", "bbb", "ééé"}; cudf::test::strings_column_wrapper strings(h_strings.begin(), h_strings.end()); @@ -88,17 +88,14 @@ TEST_F(GatherTestStr, GatherIgnoreOutOfBounds) cudf::test::fixed_width_column_wrapper gather_map(h_map.begin(), h_map.end()); auto results = cudf::detail::gather(source_table, gather_map, - cudf::detail::out_of_bounds_policy::IGNORE, + cudf::out_of_bounds_policy::DONT_CHECK, cudf::detail::negative_index_policy::NOT_ALLOWED); std::vector h_expected; - std::vector expected_validity; for (auto itr = h_map.begin(); itr != h_map.end(); ++itr) { h_expected.push_back(h_strings[*itr]); - expected_validity.push_back(1); } - cudf::test::strings_column_wrapper expected( - h_expected.begin(), h_expected.end(), expected_validity.begin()); + cudf::test::strings_column_wrapper expected(h_expected.begin(), h_expected.end()); CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view().column(0), expected); } @@ -107,7 +104,9 @@ TEST_F(GatherTestStr, GatherZeroSizeStringsColumn) cudf::column_view zero_size_strings_column( cudf::data_type{cudf::type_id::STRING}, 0, nullptr, nullptr, 0); rmm::device_vector gather_map{}; - auto results = cudf::detail::gather( - cudf::table_view({zero_size_strings_column}), gather_map.begin(), gather_map.end(), true); + auto results = cudf::detail::gather(cudf::table_view({zero_size_strings_column}), + gather_map.begin(), + gather_map.end(), + cudf::out_of_bounds_policy::NULLIFY); cudf::test::expect_strings_empty(results->get_column(0).view()); } diff --git a/python/cudf/cudf/_lib/copying.pyx b/python/cudf/cudf/_lib/copying.pyx index f35b5d39501..ad798a73ed2 100644 --- a/python/cudf/cudf/_lib/copying.pyx +++ b/python/cudf/cudf/_lib/copying.pyx @@ -12,6 +12,7 @@ from cudf._lib.column cimport Column from cudf._lib.scalar import as_device_scalar from cudf._lib.scalar cimport DeviceScalar from cudf._lib.table cimport Table +from cudf._lib.reduce import minmax from cudf._lib.cpp.column.column cimport column from cudf._lib.cpp.column.column_view cimport ( @@ -130,7 +131,16 @@ def copy_range(Column input_column, def gather(Table source_table, Column gather_map, bool keep_index=True): - assert pd.api.types.is_integer_dtype(gather_map.dtype) + if not pd.api.types.is_integer_dtype(gather_map.dtype): + raise ValueError("Gather map is not integer dtype.") + + if len(gather_map) > 0: + gm_min, gm_max = minmax(gather_map) + if gm_min < -len(source_table) or gm_max >= len(source_table): + raise IndexError(f"Gather map index with min {gm_min}," + f" max {gm_max} is out of bounds in" + f" {type(source_table)} with {len(source_table)}" + f" rows.") cdef unique_ptr[table] c_result cdef table_view source_table_view @@ -139,14 +149,16 @@ def gather(Table source_table, Column gather_map, bool keep_index=True): else: source_table_view = source_table.data_view() cdef column_view gather_map_view = gather_map.view() - cdef bool c_bounds_check = True + cdef cpp_copying.out_of_bounds_policy policy = ( + cpp_copying.out_of_bounds_policy.DONT_CHECK + ) with nogil: c_result = move( cpp_copying.gather( source_table_view, gather_map_view, - c_bounds_check + policy ) ) diff --git a/python/cudf/cudf/_lib/cpp/copying.pxd b/python/cudf/cudf/_lib/cpp/copying.pxd index f499aeb4482..55cbc3880ac 100644 --- a/python/cudf/cudf/_lib/cpp/copying.pxd +++ b/python/cudf/cudf/_lib/cpp/copying.pxd @@ -21,10 +21,14 @@ from cudf._lib.cpp.types cimport size_type ctypedef const scalar constscalar cdef extern from "cudf/copying.hpp" namespace "cudf" nogil: + ctypedef enum out_of_bounds_policy: + NULLIFY 'cudf::out_of_bounds_policy::NULLIFY' + DONT_CHECK 'cudf::out_of_bounds_policy::DONT_CHECK' + cdef unique_ptr[table] gather ( const table_view& source_table, const column_view& gather_map, - bool bounds_check + out_of_bounds_policy policy ) except + cdef unique_ptr[column] shift( diff --git a/python/cudf/cudf/_lib/cpp/reduce.pxd b/python/cudf/cudf/_lib/cpp/reduce.pxd index 817b72ea3bd..dfe1ffd3669 100644 --- a/python/cudf/cudf/_lib/cpp/reduce.pxd +++ b/python/cudf/cudf/_lib/cpp/reduce.pxd @@ -7,6 +7,7 @@ from cudf._lib.cpp.column.column cimport column from cudf._lib.scalar cimport DeviceScalar from cudf._lib.aggregation cimport aggregation from libcpp.memory cimport unique_ptr +from libcpp.utility cimport pair cdef extern from "cudf/reduction.hpp" namespace "cudf" nogil: @@ -25,3 +26,8 @@ cdef extern from "cudf/reduction.hpp" namespace "cudf" nogil: const unique_ptr[aggregation] agg, scan_type inclusive ) except + + + cdef pair[unique_ptr[scalar], + unique_ptr[scalar]] cpp_minmax "cudf::minmax" ( + column_view col + ) except + diff --git a/python/cudf/cudf/_lib/reduce.pyx b/python/cudf/cudf/_lib/reduce.pyx index 4399a1311e9..7b455dd574b 100644 --- a/python/cudf/cudf/_lib/reduce.pyx +++ b/python/cudf/cudf/_lib/reduce.pyx @@ -1,7 +1,7 @@ # Copyright (c) 2020, NVIDIA CORPORATION. import cudf -from cudf._lib.cpp.reduce cimport cpp_reduce, cpp_scan, scan_type +from cudf._lib.cpp.reduce cimport cpp_reduce, cpp_scan, scan_type, cpp_minmax from cudf._lib.cpp.scalar.scalar cimport scalar from cudf._lib.cpp.types cimport data_type, type_id from cudf._lib.cpp.column.column_view cimport column_view @@ -12,7 +12,7 @@ from cudf._lib.types import np_to_cudf_types from cudf._lib.types cimport underlying_type_t_type_id from cudf._lib.aggregation cimport make_aggregation, aggregation from libcpp.memory cimport unique_ptr -from libcpp.utility cimport move +from libcpp.utility cimport move, pair import numpy as np @@ -105,3 +105,28 @@ def scan(scan_op, Column incol, inclusive, **kwargs): py_result = Column.from_unique_ptr(move(c_result)) return py_result + + +def minmax(Column incol): + """ + Top level Cython minmax function wrapping libcudf++ minmax. + + Parameters + ---------- + incol : Column + A cuDF Column object + + Returns + ------- + A pair of ``(min, max)`` values of ``incol`` + """ + cdef column_view c_incol_view = incol.view() + cdef pair[unique_ptr[scalar], unique_ptr[scalar]] c_result + + with nogil: + c_result = move(cpp_minmax(c_incol_view)) + + py_result_min = DeviceScalar.from_unique_ptr(move(c_result.first)) + py_result_max = DeviceScalar.from_unique_ptr(move(c_result.second)) + + return cudf.Scalar(py_result_min), cudf.Scalar(py_result_max)