-
Notifications
You must be signed in to change notification settings - Fork 915
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
Fix memory write error in get_list_child_to_list_row_mapping utility (#…
…8994) Reference issue #8883 and depends on fixes in PR #8884 The `get_list_child_to_list_row_mapping` builds a map for rolling operation on a lists column. In the `thrust::scatter` call a map value includes the last offset which will always be out-of-bounds to given output vector. This output vector is used to build the resultant output map by calling `thrust::inclusive_scan` but the out-of-bounds offset value is not used -- which is why the utility does not fail. The fix in this PR simply allocates an extra row in the intermediate vector so the `thrust::scatter` will not write to out-of-bounds memory. Since the value is eventually ignored, it does not effect the result. The code in this function was creating many temporary columns incorrectly using the passed in `device_resource_manager` variable `mr`. The code was corrected by changing these to be just `device_uvector's` instead making it more clear that these are internal temporary memory buffers. Further the code calling `get_list_child_to_list_row_mapping` utility is using the output as a temporary column and so this PR fixes the logic to correct the memory resource usage. Authors: - David Wendt (https://github.com/davidwendt) Approvers: - Robert Maynard (https://github.com/robertmaynard) - MithunR (https://github.com/mythrocks) - Karthikeyan (https://github.com/karthikeyann) URL: #8994
- Loading branch information
1 parent
e42464c
commit d4c3f32
Showing
3 changed files
with
176 additions
and
147 deletions.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,157 @@ | ||
| /* | ||
| * 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 <rolling/rolling_collect_list.cuh> | ||
|
|
||
| #include <cudf/detail/get_value.cuh> | ||
| #include <cudf/detail/iterator.cuh> | ||
|
|
||
| #include <rmm/device_uvector.hpp> | ||
|
|
||
| #include <thrust/copy.h> | ||
| #include <thrust/iterator/constant_iterator.h> | ||
| #include <thrust/scan.h> | ||
| #include <thrust/scatter.h> | ||
|
|
||
| namespace cudf { | ||
| namespace detail { | ||
|
|
||
| /** | ||
| * @see cudf::detail::get_list_child_to_list_row_mapping | ||
| */ | ||
| std::unique_ptr<column> get_list_child_to_list_row_mapping(cudf::column_view const& offsets, | ||
| rmm::cuda_stream_view stream) | ||
| { | ||
| // First, scatter the count for each repeated offset (except the first and last), | ||
| // into a column of N `0`s, where N == number of child rows. | ||
| // For example: | ||
| // offsets == [0, 2, 5, 8, 11, 13] | ||
| // scatter result == [0, 0, 1, 0, 0, 1, 0, 0, 1, 0, 0, 1, 0] | ||
| // | ||
| // An example with empty list row at index 2: | ||
| // offsets == [0, 2, 5, 5, 8, 11, 13] | ||
| // scatter result == [0, 0, 1, 0, 0, 2, 0, 0, 1, 0, 0, 1, 0] | ||
| // | ||
| auto const num_child_rows{ | ||
| cudf::detail::get_value<size_type>(offsets, offsets.size() - 1, stream)}; | ||
| auto per_row_mapping = make_fixed_width_column( | ||
| data_type{type_to_id<size_type>()}, num_child_rows, mask_state::UNALLOCATED, stream); | ||
| auto per_row_mapping_begin = per_row_mapping->mutable_view().template begin<size_type>(); | ||
| thrust::fill_n(rmm::exec_policy(stream), per_row_mapping_begin, num_child_rows, 0); | ||
|
|
||
| auto const begin = thrust::make_counting_iterator<size_type>(0); | ||
| thrust::scatter_if(rmm::exec_policy(stream), | ||
| begin, | ||
| begin + offsets.size() - 1, | ||
| offsets.begin<size_type>(), | ||
| begin, // stencil iterator | ||
| per_row_mapping_begin, | ||
| [offset = offsets.begin<size_type>()] __device__(auto i) { | ||
| return offset[i] != offset[i + 1]; | ||
| }); // [0,0,1,0,0,3,...] | ||
|
|
||
| // Next, generate mapping with inclusive_scan(max) on the scatter result. | ||
| // For the example above: | ||
| // scatter result == [0, 0, 1, 0, 0, 2, 0, 0, 3, 0, 0, 4, 0] | ||
| // inclusive_scan == [0, 0, 1, 1, 1, 2, 2, 2, 3, 3, 3, 4, 4] | ||
| // | ||
| // For the case with an empty list at index 2: | ||
| // scatter result == [0, 0, 1, 0, 0, 3, 0, 0, 4, 0, 0, 5, 0] | ||
| // inclusive_scan == [0, 0, 1, 1, 1, 3, 3, 3, 4, 4, 4, 5, 5] | ||
| thrust::inclusive_scan(rmm::exec_policy(stream), | ||
| per_row_mapping_begin, | ||
| per_row_mapping_begin + num_child_rows, | ||
| per_row_mapping_begin, | ||
| thrust::maximum<size_type>{}); | ||
| return per_row_mapping; | ||
| } | ||
|
|
||
| /** | ||
| * @see cudf::detail::count_child_nulls | ||
| */ | ||
| size_type count_child_nulls(column_view const& input, | ||
| std::unique_ptr<column> const& gather_map, | ||
| rmm::cuda_stream_view stream) | ||
| { | ||
| auto input_device_view = column_device_view::create(input, stream); | ||
|
|
||
| auto input_row_is_null = [d_input = *input_device_view] __device__(auto i) { | ||
| return d_input.is_null_nocheck(i); | ||
| }; | ||
|
|
||
| return thrust::count_if(rmm::exec_policy(stream), | ||
| gather_map->view().begin<size_type>(), | ||
| gather_map->view().end<size_type>(), | ||
| input_row_is_null); | ||
| } | ||
|
|
||
| /** | ||
| * @see cudf::detail::rolling_collect_list | ||
| */ | ||
| std::pair<std::unique_ptr<column>, std::unique_ptr<column>> purge_null_entries( | ||
| column_view const& input, | ||
| column_view const& gather_map, | ||
| column_view const& offsets, | ||
| size_type num_child_nulls, | ||
| rmm::cuda_stream_view stream, | ||
| rmm::mr::device_memory_resource* mr) | ||
| { | ||
| auto input_device_view = column_device_view::create(input, stream); | ||
|
|
||
| auto input_row_not_null = [d_input = *input_device_view] __device__(auto i) { | ||
| return d_input.is_valid_nocheck(i); | ||
| }; | ||
|
|
||
| // Purge entries in gather_map that correspond to null input. | ||
| auto new_gather_map = make_fixed_width_column(data_type{type_to_id<size_type>()}, | ||
| gather_map.size() - num_child_nulls, | ||
| mask_state::UNALLOCATED, | ||
| stream); | ||
| thrust::copy_if(rmm::exec_policy(stream), | ||
| gather_map.template begin<size_type>(), | ||
| gather_map.template end<size_type>(), | ||
| new_gather_map->mutable_view().template begin<size_type>(), | ||
| input_row_not_null); | ||
|
|
||
| // Recalculate offsets after null entries are purged. | ||
| auto new_sizes = make_fixed_width_column( | ||
| data_type{type_to_id<size_type>()}, input.size(), mask_state::UNALLOCATED, stream); | ||
|
|
||
| thrust::transform(rmm::exec_policy(stream), | ||
| thrust::make_counting_iterator<size_type>(0), | ||
| thrust::make_counting_iterator<size_type>(input.size()), | ||
| new_sizes->mutable_view().template begin<size_type>(), | ||
| [d_gather_map = gather_map.template begin<size_type>(), | ||
| d_old_offsets = offsets.template begin<size_type>(), | ||
| input_row_not_null] __device__(auto i) { | ||
| return thrust::count_if(thrust::seq, | ||
| d_gather_map + d_old_offsets[i], | ||
| d_gather_map + d_old_offsets[i + 1], | ||
| input_row_not_null); | ||
| }); | ||
|
|
||
| auto new_offsets = | ||
| strings::detail::make_offsets_child_column(new_sizes->view().template begin<size_type>(), | ||
| new_sizes->view().template end<size_type>(), | ||
| stream, | ||
| mr); | ||
|
|
||
| return std::make_pair<std::unique_ptr<column>, std::unique_ptr<column>>(std::move(new_gather_map), | ||
| std::move(new_offsets)); | ||
| } | ||
|
|
||
| } // namespace detail | ||
| } // namespace cudf |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters