From af4f5e5b237cc32cd0c650891693fdf4d3273bab Mon Sep 17 00:00:00 2001 From: David Wendt Date: Tue, 8 Sep 2026 12:44:29 -0400 Subject: [PATCH] Fix racecheck reported in parquet decode_delta_length_byte_array_kernel --- cpp/src/io/parquet/page_delta_decode.cu | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/cpp/src/io/parquet/page_delta_decode.cu b/cpp/src/io/parquet/page_delta_decode.cu index a49babac3bf0..65328dfecd7d 100644 --- a/cpp/src/io/parquet/page_delta_decode.cu +++ b/cpp/src/io/parquet/page_delta_decode.cu @@ -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. @@ -848,6 +846,10 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size) : min(db->values_per_mb, static_cast(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