From 3ae4319021539ead8d2b46e2395442d6f4585160 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Fri, 15 May 2026 15:32:19 -0700 Subject: [PATCH] Fix potential malformed headers in parquet delta decoder (#22275) This PR checks for potential malformed headers in Parquet delta decoder and propagates it in decoders. Authors: - Muhammad Haseeb (https://github.com/mhaseeb123) Approvers: - Paul Mattione (https://github.com/pmattione-nvidia) - Bradley Dice (https://github.com/bdice) - Nghia Truong (https://github.com/ttnghia) URL: https://github.com/rapidsai/cudf/pull/22275 --- cpp/src/io/parquet/delta_binary.cuh | 17 +++++++++++- cpp/src/io/parquet/page_delta_decode.cu | 35 ++++++++++++++++++------- 2 files changed, 42 insertions(+), 10 deletions(-) diff --git a/cpp/src/io/parquet/delta_binary.cuh b/cpp/src/io/parquet/delta_binary.cuh index 513c65515bf7..6dac50ce80bd 100644 --- a/cpp/src/io/parquet/delta_binary.cuh +++ b/cpp/src/io/parquet/delta_binary.cuh @@ -91,6 +91,7 @@ struct delta_binary_decoder { uint32_t cur_mb; // index of the current mini-block within the block uint8_t const* cur_mb_start; // pointer to the start of the current mini-block data uint8_t const* cur_bitwidths; // pointer to the bitwidth array in the block + bool error; // flag to catch malformed headers zigzag128_t value[delta_rolling_buf_size]; // circular buffer of delta values @@ -148,7 +149,21 @@ struct delta_binary_decoder { last_value = first_value; current_value_idx = 0; - values_per_mb = block_size / mini_block_count; + error = false; + + // Validate header against the DELTA_BINARY_PACKED spec invariants + if (mini_block_count == 0 or block_size == 0 or (block_size % mini_block_count) != 0) { + error = true; + value_count = 0; + values_per_mb = 1; + block_start = d_end; + cur_mb = 0; + cur_mb_start = d_end; + cur_bitwidths = d_end; + return; + } + + values_per_mb = block_size / mini_block_count; // init the first mini-block block_start = d_start; diff --git a/cpp/src/io/parquet/page_delta_decode.cu b/cpp/src/io/parquet/page_delta_decode.cu index 0900e73cd6dd..82ac47391811 100644 --- a/cpp/src/io/parquet/page_delta_decode.cu +++ b/cpp/src/io/parquet/page_delta_decode.cu @@ -366,9 +366,11 @@ CUDF_KERNEL void __launch_bounds__(decode_delta_binary_block_size) block.sync(); auto const batch_size = db->values_per_mb; - if (batch_size > max_delta_mini_block_size) { - set_error(static_cast(decode_error::DELTA_PARAMS_UNSUPPORTED), - error_code); + if (db->error or batch_size > max_delta_mini_block_size) { + if (block.thread_rank() == 0) { + set_error(static_cast(decode_error::DELTA_PARAMS_UNSUPPORTED), + error_code); + } return; } @@ -546,6 +548,15 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size) } block.sync(); + // Propagate malformed-header errors from either underlying DELTA_BINARY_PACKED decoder. + if (prefix_db->error or suffix_db->error) { + if (block.thread_rank() == 0) { + set_error(static_cast(decode_error::DELTA_PARAMS_UNSUPPORTED), + error_code); + } + return; + } + // assert that prefix and suffix have same mini-block size if (prefix_db->values_per_mb != suffix_db->values_per_mb or prefix_db->block_size != suffix_db->block_size or @@ -562,8 +573,10 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size) // sanity check to make sure we can process this page auto const batch_size = prefix_db->values_per_mb; if (batch_size > max_delta_mini_block_size) { - set_error(static_cast(decode_error::DELTA_PARAMS_UNSUPPORTED), - error_code); + if (block.thread_rank() == 0) { + set_error(static_cast(decode_error::DELTA_PARAMS_UNSUPPORTED), + error_code); + } return; } @@ -759,14 +772,18 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size) } block.sync(); - int const leaf_level_index = s->col.max_nesting_depth - 1; - // sanity check to make sure we can process this page auto const batch_size = db->values_per_mb; - if (batch_size > max_delta_mini_block_size) { - set_error(static_cast(decode_error::DELTA_PARAMS_UNSUPPORTED), error_code); + if (db->error or batch_size > max_delta_mini_block_size) { + if (block.thread_rank() == 0) { + set_error(static_cast(decode_error::DELTA_PARAMS_UNSUPPORTED), + error_code); + } return; } + + int const leaf_level_index = s->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