From 8d17379bd821edc67c3a3f93801d3e79b8e2d97f Mon Sep 17 00:00:00 2001 From: David Wendt <45795991+davidwendt@users.noreply.github.com> Date: Thu, 9 Feb 2023 08:33:46 -0500 Subject: [PATCH] Fix memcheck read error in compound segmented reduce (#12722) 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: https://github.com/rapidsai/cudf/pull/12722 --- cpp/include/cudf/detail/segmented_reduction.cuh | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/cpp/include/cudf/detail/segmented_reduction.cuh b/cpp/include/cudf/detail/segmented_reduction.cuh index 9a49c1abe38..1c39d5eab1e 100644 --- a/cpp/include/cudf/detail/segmented_reduction.cuh +++ b/cpp/include/cudf/detail/segmented_reduction.cuh @@ -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::type; - using IntermediateType = typename thrust::iterator_value::type; - auto num_segments = static_cast(std::distance(d_offset_begin, d_offset_end)); - auto const binary_op = op.get_binary_op(); + using OutputType = typename thrust::iterator_value::type; + using IntermediateType = typename thrust::iterator_value::type; + auto num_segments = static_cast(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(); rmm::device_uvector intermediate_result{static_cast(num_segments),