Use cudf::thread_index_type in cuIO to prevent overflow in row indexing
#13910
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
vuule
added
bug
…bug-prevent-overflow-tidx
…cudf into bug-prevent-overflow-tidx
vuule
commented
Aug 18, 2023
Comment on lines
219
to
-220
| while (cur_row < end_row) { | ||
| auto const is_valid = cur_row < col.size() and col.is_valid(cur_row); |
There was a problem hiding this comment.
redundant check with cur_row < end_row
davidwendt
reviewed
Aug 18, 2023
cpp/src/io/csv/csv_gpu.cu
Outdated
| @@ -317,8 +316,8 @@ __global__ void __launch_bounds__(csvparse_block_dim) | |||
| auto const raw_csv = data.data(); | |||
| // thread IDs range per block, so also need the block id. | |||
| // this is entry into the field array - tid is an elements within the num_entries array | |||
| long const rec_id = threadIdx.x + (blockDim.x * blockIdx.x); | |||
| long const rec_id_next = rec_id + 1; | |||
| auto const rec_id = cudf::thread_index_type{threadIdx.x} + (blockDim.x * blockIdx.x); | |||
There was a problem hiding this comment.
Promoting just the threadId.x does not appear to be enough to promote the multiplication
https://godbolt.org/z/coqPKqfWe
There was a problem hiding this comment.
moved the promotion to the multiplication
PointKernel
reviewed
Aug 22, 2023
…bug-prevent-overflow-tidx
harrism
requested changes
Aug 22, 2023
…bug-prevent-overflow-tidx
harrism
approved these changes
Aug 23, 2023
|
/merge |
3 tasks
rapids-bot bot
pushed a commit
that referenced
this pull request
Sep 1, 2023
This PR adds `grid_1d::grid_stride()` and uses it in a handful of kernels. Follow-up to #13910, which added a `grid_1d::global_thread_id()`. We'll need to do a later PR that catches any missing instances where this should be used, since there are a large number of PRs in flight touching thread indexing code in various files. See #10368. Authors: - Bradley Dice (https://github.com/bdice) Approvers: - Yunsong Wang (https://github.com/PointKernel) - Vyas Ramasubramani (https://github.com/vyasr) URL: #13996
3 tasks
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Description
Use wider type for indexing when rows are indexed in pattern
or
Overflow can happen when the number of rows is so close to
max<size_type>that adding block size pushes the index over the max.Also sprinkled auto where increased size is not needed.
Also removed a few redundant conditions.
Checklist