Skip to content

Commit

Permalink
Fix memory read error in get_dremel_data in page_enc.cu (#8995)
Browse files Browse the repository at this point in the history
Reference issue #8883 and depends on PR #8884

The `PARQUET_TEST` fails with cuda-memcheck when called with the `rmm_mode=cuda` parameter. The 4-byte read error is caused by this code logic:

```
// Scan to get distance by which each offset value is shifted due to the insertion of empties
auto scan_it = cudf::detail::make_counting_transform_iterator( column_offsets[level],
  [off = lcv.offsets().data<size_type>()] __device__(auto i) -> int { return off[i] == off[i + 1]; });
rmm::device_uvector<size_type> scan_out(offset_size_at_level, stream);
thrust::exclusive_scan(rmm::exec_policy(stream), scan_it, scan_it + offset_size_at_level, scan_out.begin());
```

The `scan_it` lambda will read one offset value passed the end due to the `off[i + 1]` statement. The `exclusive_scan` does not actually use the last element so the code was modified to just return `false` if the index, `i` is greater than or equal to the size of the offsets column.

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

Approvers:
  - Robert Maynard (https://github.com/robertmaynard)
  - Devavret Makkar (https://github.com/devavret)

URL: #8995
  • Loading branch information
davidwendt authored Aug 13, 2021
1 parent 233943d commit 77baf09
Showing 1 changed file with 6 additions and 6 deletions.
12 changes: 6 additions & 6 deletions cpp/src/io/parquet/page_enc.cu
Original file line number Diff line number Diff line change
Expand Up @@ -1966,9 +1966,9 @@ dremel_data get_dremel_data(column_view h_col,

// Scan to get distance by which each offset value is shifted due to the insertion of empties
auto scan_it = cudf::detail::make_counting_transform_iterator(
column_offsets[level], [off = lcv.offsets().data<size_type>()] __device__(auto i) -> int {
return off[i] == off[i + 1];
});
column_offsets[level],
[off = lcv.offsets().data<size_type>(), size = lcv.offsets().size()] __device__(
auto i) -> int { return (i + 1 < size) && (off[i] == off[i + 1]); });
rmm::device_uvector<size_type> scan_out(offset_size_at_level, stream);
thrust::exclusive_scan(
rmm::exec_policy(stream), scan_it, scan_it + offset_size_at_level, scan_out.begin());
Expand Down Expand Up @@ -2053,9 +2053,9 @@ dremel_data get_dremel_data(column_view h_col,
// Scan to get distance by which each offset value is shifted due to the insertion of dremel
// level value fof an empty list
auto scan_it = cudf::detail::make_counting_transform_iterator(
column_offsets[level], [off = lcv.offsets().data<size_type>()] __device__(auto i) -> int {
return off[i] == off[i + 1];
});
column_offsets[level],
[off = lcv.offsets().data<size_type>(), size = lcv.offsets().size()] __device__(
auto i) -> int { return (i + 1 < size) && (off[i] == off[i + 1]); });
rmm::device_uvector<size_type> scan_out(offset_size_at_level, stream);
thrust::exclusive_scan(
rmm::exec_policy(stream), scan_it, scan_it + offset_size_at_level, scan_out.begin());
Expand Down

0 comments on commit 77baf09

Please sign in to comment.