Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Fix memory read error in get_dremel_data in page_enc.cu #8995

Merged

Conversation

davidwendt
Copy link
Contributor

@davidwendt davidwendt commented Aug 9, 2021

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.

@davidwendt davidwendt added bug Something isn't working 2 - In Progress Currently a work in progress libcudf Affects libcudf (C++/CUDA) code. cuIO cuIO issue non-breaking Non-breaking change labels Aug 9, 2021
@davidwendt davidwendt self-assigned this Aug 9, 2021
@codecov
Copy link

codecov bot commented Aug 9, 2021

Codecov Report

❗ No coverage uploaded for pull request base (branch-21.10@2c5a2ad). Click here to learn what that means.
The diff coverage is n/a.

Impacted file tree graph

@@               Coverage Diff               @@
##             branch-21.10    #8995   +/-   ##
===============================================
  Coverage                ?   10.60%           
===============================================
  Files                   ?      116           
  Lines                   ?    19064           
  Branches                ?        0           
===============================================
  Hits                    ?     2022           
  Misses                  ?    17042           
  Partials                ?        0           

Continue to review full report at Codecov.

Legend - Click here to learn more
Δ = absolute <relative> (impact), ø = not affected, ? = missing data
Powered by Codecov. Last update 2c5a2ad...df219e9. Read the comment docs.

@davidwendt davidwendt marked this pull request as ready for review August 9, 2021 20:15
@davidwendt davidwendt requested a review from a team as a code owner August 9, 2021 20:15
@davidwendt davidwendt added 3 - Ready for Review Ready for review by team and removed 2 - In Progress Currently a work in progress labels Aug 9, 2021
Copy link
Contributor

@devavret devavret left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I thought hard but there's really no other easy way to achieve this.

@davidwendt
Copy link
Contributor Author

@gpucibot merge

@rapids-bot rapids-bot bot merged commit 77baf09 into rapidsai:branch-21.10 Aug 13, 2021
@davidwendt davidwendt deleted the bug-memcheck-parquet-page-enc branch August 13, 2021 22:01
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
3 - Ready for Review Ready for review by team bug Something isn't working cuIO cuIO issue libcudf Affects libcudf (C++/CUDA) code. non-breaking Non-breaking change
Projects
None yet
Development

Successfully merging this pull request may close these issues.

3 participants