Skip to content

Commit

Permalink
Fix memcheck error in groupby-tdigest get_scalar_minmax (#9339)
Browse files Browse the repository at this point in the history
The groupby tdigest logic in the `get_scalar_minmax` would access incorrect device memory when encountering an empty group. The index value calculated for a column element in an empty group would be off by -1. This was found running a cuda-memcheck (compute-sanitizer) with the `AllNulls` gtest. Here the first group is empty resulting an out-of-bounds read at `col.element<T>(-1)`.

This PR adds a check for the empty group (`valid_count==0`) to prevent reading the incorrect column row.

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

Approvers:
  - https://github.com/nvdbaranec
  - MithunR (https://github.com/mythrocks)

URL: #9339
  • Loading branch information
davidwendt authored Oct 1, 2021
1 parent 8e371cd commit 4090c45
Showing 1 changed file with 6 additions and 4 deletions.
10 changes: 6 additions & 4 deletions cpp/src/groupby/sort/group_tdigest.cu
Original file line number Diff line number Diff line change
Expand Up @@ -535,10 +535,12 @@ struct get_scalar_minmax {

__device__ thrust::tuple<double, double> operator()(size_type group_index)
{
// note: .element<T>() is taking care of fixed-point conversions for us.
return {static_cast<double>(col.element<T>(group_offsets[group_index])),
static_cast<double>(
col.element<T>(group_offsets[group_index] + (group_valid_counts[group_index] - 1)))};
auto const valid_count = group_valid_counts[group_index];
return valid_count > 0
? thrust::make_tuple(
static_cast<double>(col.element<T>(group_offsets[group_index])),
static_cast<double>(col.element<T>(group_offsets[group_index] + valid_count - 1)))
: thrust::make_tuple(0.0, 0.0);
}
};

Expand Down

0 comments on commit 4090c45

Please sign in to comment.