Skip to content

Commit

Permalink
Fix memcheck read error in compound segmented reduce (#12722)
Browse files Browse the repository at this point in the history
Fixes an out-of-bounds memory read in the compound segmented reduction logic. The number of segments was computed incorrectly causing an extra read passed the end of the valid-counts vector. This was found by running compute-sanitizer test on the reductions gtests as follows:
```
compute-sanitizer --tool memcheck --demangle full gtests/REDUCTIONS_TEST --rmm_mode=cuda
```
The number of segments is 1 less than the number of offsets.

Authors:
  - David Wendt (https://github.com/davidwendt)

Approvers:
  - Bradley Dice (https://github.com/bdice)
  - Vyas Ramasubramani (https://github.com/vyasr)

URL: #12722
  • Loading branch information
davidwendt authored Feb 9, 2023
1 parent b0335f0 commit 8d17379
Showing 1 changed file with 4 additions and 4 deletions.
8 changes: 4 additions & 4 deletions cpp/include/cudf/detail/segmented_reduction.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -145,10 +145,10 @@ void segmented_reduce(InputIterator d_in,
size_type* d_valid_counts,
rmm::cuda_stream_view stream)
{
using OutputType = typename thrust::iterator_value<OutputIterator>::type;
using IntermediateType = typename thrust::iterator_value<InputIterator>::type;
auto num_segments = static_cast<size_type>(std::distance(d_offset_begin, d_offset_end));
auto const binary_op = op.get_binary_op();
using OutputType = typename thrust::iterator_value<OutputIterator>::type;
using IntermediateType = typename thrust::iterator_value<InputIterator>::type;
auto num_segments = static_cast<size_type>(std::distance(d_offset_begin, d_offset_end)) - 1;
auto const binary_op = op.get_binary_op();
auto const initial_value = op.template get_identity<IntermediateType>();

rmm::device_uvector<IntermediateType> intermediate_result{static_cast<std::size_t>(num_segments),
Expand Down

0 comments on commit 8d17379

Please sign in to comment.