Skip to content

Fix racecheck reported in parquet decode_delta_length_byte_array_kernel #24034

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.

Sign up for GitHub

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

Merged
rapids-bot merged 2 commits into from
Sep 9, 2026
+4 −2
Merged

Fix racecheck reported in parquet decode_delta_length_byte_array_kernel #24034

Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
af4f5e5
Fix racecheck reported in parquet decode_delta_length_byte_array_kernel
davidwendt Sep 8, 2026
d7d4077
Merge branch 'release/26.10' into racecheck-parquet-reader
davidwendt Sep 9, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 4 additions & 2 deletions cpp/src/io/parquet/page_delta_decode.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -830,8 +830,6 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size)

int const leaf_level_index = s->setup.col.max_nesting_depth - 1;

// db->init_binary_block below resets db->values_per_mb
block.sync();
// if this is a bounds page, then we need to decode up to the first mini-block
// that has a value we need, and set string_offset to the position of the first value in the
// string data block.
Expand All @@ -848,6 +846,10 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size)
: min(db->values_per_mb, static_cast<uint32_t>(delta_max_batch_size));
uint32_t const passes_per_batch = batch_size / cudf::detail::warp_size;

// db->init_binary_block below resets db->values_per_mb, so make sure every thread has read it
// for batch_size above before warp 0 re-initializes the decoder
block.sync();

if (is_skip_resume) {
if (warp.meta_group_rank() == 0) {
// string_off is only valid on thread 0
Expand Down
Loading