-
Notifications
You must be signed in to change notification settings - Fork 915
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 read out of bounds in string concatenate #13838
Fix read out of bounds in string concatenate #13838
Conversation
xrefing #13771 (comment) as I was unable to generate tests. For completeness, here's what I wrote there referring missing tests:
|
cd3c128
to
9e4b456
Compare
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
This appears to be the same fix as #10344. Do we need any other explicit casts for bounds checking, like that PR had? I did not see anything immediately where that would be necessary but may be worth another inspection.
I'm assuming you're referring specifically to this line. AFAIU, this is already covered here, but please correct me if I'm wrong. |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
LGTM
@pentschev My question above seems fine / already covered. Thanks! |
I managed to write a proper test in 25712ec. This is the output if the fix from this PR is reverted:
It's a bit of a long running test though (~3.5s), I hope this is ok. |
@@ -121,7 +121,7 @@ __global__ void fused_concatenate_string_offset_kernel(column_device_view const* | |||
bitmask_type* output_mask, | |||
size_type* out_valid_count) | |||
{ | |||
size_type output_index = threadIdx.x + blockIdx.x * blockDim.x; | |||
int64_t output_index = threadIdx.x + blockIdx.x * blockDim.x; |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Would you please use the cudf::thread_index_type
alias? Thank you again for diagnosing this issue.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Thanks @GregoryKimball , addressed that.
Reviews/questions were addressed and tests are passing. Anything else needed here or could we get it merged? |
/merge |
@pentschev If you have write access, feel free to merge the PR on your own after getting enough approvals. |
Description
If data is sufficiently large,
fused_concatenate_string_chars_kernel
will attempt to read out of bounds and ultimately cause CUDA to raisecudaErrorIllegalAddress
. Details on how the issue was encountered are in #13771, although this was an already known problem.Fixes #13771 .
Checklist