Fix race in ORC string dictionary creation #13214
Merged
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
Signed-off-by: Robert (Bobby) Evans <[email protected]>
revans2
added
bug
abellina
reviewed
Apr 25, 2023
| @@ -72,6 +72,10 @@ static __device__ void LoadNonNullIndices(volatile dictinit_state_s* s, | |||
| Storage& temp_storage) | |||
| { | |||
| if (t == 0) { s->nnz = 0; } | |||
| if (s->chunk.num_rows <= 0) { | |||
There was a problem hiding this comment.
copyrights need to be updated
There was a problem hiding this comment.
@nvdbaranec Is there anything similar in parquet code?
|
Nice catch :) |
ttnghia
approved these changes
Apr 25, 2023
vuule
approved these changes
Apr 25, 2023
|
/merge |
|
Thank you @revans2 for investigating this - excellent work |
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
Unfortunately this is really hard to reproduce. For whatever reason I had to try and reproduce this on a relatively small data set with at least 140,001 rows or more, where one column is a LIST but all of the lists are empty lists and another column is a STRUCT column with two STRING child columns where all of the STRINGS are empty. I also had to sort and partition the data before doing the write, and it had to be in a very specific environment with T4 GPUs. I don't know why all of those were needed to make the race happen regularly, but it did.
Because of this complexity in reproducing it I have not added in any unit tests.
The problem was essentially a race when trying to calculate dictionary duplication for strings in ORC. As a part of this a function
LoadNonNullIndiceswas being called that was supposed to set a valuennzin a shared memory locations. In the normal case a loop was taken where__syncthreads()was called, but if there were no rows in the column (the LIST column) then the loop was not taken and it was a race to see ifnnzwhich was set to 0 by thread 0 showed up in all of the threads or not.What made this crash is that this
nnzvalue is used to determine what happens in the rest of the kernel to see if it reads data, or writes to temp memory (which is not allocated if previous processing shows that there is no need for it), or any of that. Ifnnzis non-zero then it tries to do all of those things and bad stuff starts to happen.Checklist
As a side note I am not a C++ or CUDA expert so I am happy to any suggestions.