Skip to content

Commit

Permalink
Fix synchronization bug in bool parquet mukernels (rapidsai#17302)
Browse files Browse the repository at this point in the history
This fixes a synchronization bug in the parquet microkernels for plain-decoding bools.  This closes [several](NVIDIA/spark-rapids#11715) timing [issues](NVIDIA/spark-rapids#11716) found during testing of spark-rapids.

Authors:
  - Paul Mattione (https://github.com/pmattione-nvidia)

Approvers:
  - Bradley Dice (https://github.com/bdice)
  - Vukasin Milovanovic (https://github.com/vuule)

URL: rapidsai#17302
  • Loading branch information
pmattione-nvidia authored Nov 13, 2024
1 parent 6acd33d commit 5e40691
Showing 1 changed file with 1 addition and 0 deletions.
1 change: 1 addition & 0 deletions cpp/src/io/parquet/decode_fixed.cu
Original file line number Diff line number Diff line change
Expand Up @@ -848,6 +848,7 @@ inline __device__ void bool_plain_decode(page_state_s* s, state_buf* sb, int t,
{
int pos = s->dict_pos;
int const target_pos = pos + to_decode;
__syncthreads(); // Make sure all threads have read dict_pos before it changes at the end.

while (pos < target_pos) {
int const batch_len = min(target_pos - pos, decode_block_size_t);
Expand Down

0 comments on commit 5e40691

Please sign in to comment.