From 6a32d883683bcd0c5750d46c2ff1bd7badd1ef6c Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Thu, 23 Apr 2026 20:29:38 +0000 Subject: [PATCH 1/6] Check malformed block headers in parquet delta decoder --- cpp/src/io/parquet/delta_binary.cuh | 16 ++++++++++++- cpp/src/io/parquet/page_delta_decode.cu | 30 +++++++++++++++++++++++++ 2 files changed, 45 insertions(+), 1 deletion(-) diff --git a/cpp/src/io/parquet/delta_binary.cuh b/cpp/src/io/parquet/delta_binary.cuh index 513c65515bf7..873a6210e66f 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; // whether to catch malformed headers zigzag128_t value[delta_rolling_buf_size]; // circular buffer of delta values @@ -148,7 +149,20 @@ struct delta_binary_decoder { last_value = first_value; current_value_idx = 0; - values_per_mb = block_size / mini_block_count; + error = false; + + if (mini_block_count == 0 || block_size == 0 || (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..14abd30ce2ac 100644 --- a/cpp/src/io/parquet/page_delta_decode.cu +++ b/cpp/src/io/parquet/page_delta_decode.cu @@ -365,6 +365,13 @@ CUDF_KERNEL void __launch_bounds__(decode_delta_binary_block_size) if (block.thread_rank() == 0) { db->init_binary_block(s->data_start, s->data_end); } block.sync(); + // Exit if the DELTA_BINARY_PACKED header is malformed + if (db->error) { + set_error(static_cast(decode_error::DELTA_PARAMS_UNSUPPORTED), + error_code); + return; + } + 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), @@ -378,6 +385,9 @@ CUDF_KERNEL void __launch_bounds__(decode_delta_binary_block_size) while (s->error == 0 && (s->input_value_count < s->num_input_values || s->src_pos < s->nz_count)) { + int const prev_input_value_count = s->input_value_count; + uint32_t const prev_src_pos = s->src_pos; + uint32_t target_pos; uint32_t const src_pos = s->src_pos; @@ -432,6 +442,13 @@ CUDF_KERNEL void __launch_bounds__(decode_delta_binary_block_size) } block.sync(); + + // Ensure we advanced position and input value count in this iteration. + if (s->input_value_count == prev_input_value_count and s->src_pos == prev_src_pos) { + cg::invoke_one(block, [&] { s->set_error_code(decode_error::DELTA_PARAMS_UNSUPPORTED); }); + block.sync(); + break; + } } if (has_repetition) { @@ -546,6 +563,13 @@ 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 || suffix_db->error) { + 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 @@ -759,6 +783,12 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size) } block.sync(); + // Propagate malformed-header errors from the underlying DELTA_BINARY_PACKED decoder + if (db->error) { + set_error(static_cast(decode_error::DELTA_PARAMS_UNSUPPORTED), error_code); + return; + } + int const leaf_level_index = s->col.max_nesting_depth - 1; // sanity check to make sure we can process this page From 23aade2c777437df794ddf404775a5daf7d39e2b Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Fri, 24 Apr 2026 13:00:58 -0700 Subject: [PATCH 2/6] Apply suggestion from @mhaseeb123 --- cpp/src/io/parquet/delta_binary.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/src/io/parquet/delta_binary.cuh b/cpp/src/io/parquet/delta_binary.cuh index 873a6210e66f..1bbcbad329c1 100644 --- a/cpp/src/io/parquet/delta_binary.cuh +++ b/cpp/src/io/parquet/delta_binary.cuh @@ -91,7 +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; // whether to catch malformed headers + bool error; // flag to catch malformed headers zigzag128_t value[delta_rolling_buf_size]; // circular buffer of delta values From 1a52a6dce759bfa83deadbbb89060e481907745d Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Fri, 24 Apr 2026 13:02:03 -0700 Subject: [PATCH 3/6] Apply suggestions from code review Co-authored-by: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> --- cpp/src/io/parquet/delta_binary.cuh | 2 +- cpp/src/io/parquet/page_delta_decode.cu | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/src/io/parquet/delta_binary.cuh b/cpp/src/io/parquet/delta_binary.cuh index 1bbcbad329c1..374e53a24abe 100644 --- a/cpp/src/io/parquet/delta_binary.cuh +++ b/cpp/src/io/parquet/delta_binary.cuh @@ -151,7 +151,7 @@ struct delta_binary_decoder { current_value_idx = 0; error = false; - if (mini_block_count == 0 || block_size == 0 || (block_size % mini_block_count) != 0) { + 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; diff --git a/cpp/src/io/parquet/page_delta_decode.cu b/cpp/src/io/parquet/page_delta_decode.cu index 14abd30ce2ac..5067d3e25cf1 100644 --- a/cpp/src/io/parquet/page_delta_decode.cu +++ b/cpp/src/io/parquet/page_delta_decode.cu @@ -564,7 +564,7 @@ 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 || suffix_db->error) { + if (prefix_db->error or suffix_db->error) { set_error(static_cast(decode_error::DELTA_PARAMS_UNSUPPORTED), error_code); return; From 66df3a303f437c2ba60173f9749f6b6395228d74 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb Date: Thu, 14 May 2026 21:36:52 +0000 Subject: [PATCH 4/6] Cleanup --- cpp/src/io/parquet/delta_binary.cuh | 1 + cpp/src/io/parquet/page_delta_decode.cu | 41 ++++++++++--------------- 2 files changed, 17 insertions(+), 25 deletions(-) diff --git a/cpp/src/io/parquet/delta_binary.cuh b/cpp/src/io/parquet/delta_binary.cuh index 374e53a24abe..6dac50ce80bd 100644 --- a/cpp/src/io/parquet/delta_binary.cuh +++ b/cpp/src/io/parquet/delta_binary.cuh @@ -151,6 +151,7 @@ struct delta_binary_decoder { current_value_idx = 0; 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; diff --git a/cpp/src/io/parquet/page_delta_decode.cu b/cpp/src/io/parquet/page_delta_decode.cu index 5067d3e25cf1..6bcb69357ee0 100644 --- a/cpp/src/io/parquet/page_delta_decode.cu +++ b/cpp/src/io/parquet/page_delta_decode.cu @@ -365,17 +365,12 @@ CUDF_KERNEL void __launch_bounds__(decode_delta_binary_block_size) if (block.thread_rank() == 0) { db->init_binary_block(s->data_start, s->data_end); } block.sync(); - // Exit if the DELTA_BINARY_PACKED header is malformed - if (db->error) { - set_error(static_cast(decode_error::DELTA_PARAMS_UNSUPPORTED), - error_code); - return; - } - 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; } @@ -385,9 +380,6 @@ CUDF_KERNEL void __launch_bounds__(decode_delta_binary_block_size) while (s->error == 0 && (s->input_value_count < s->num_input_values || s->src_pos < s->nz_count)) { - int const prev_input_value_count = s->input_value_count; - uint32_t const prev_src_pos = s->src_pos; - uint32_t target_pos; uint32_t const src_pos = s->src_pos; @@ -442,13 +434,6 @@ CUDF_KERNEL void __launch_bounds__(decode_delta_binary_block_size) } block.sync(); - - // Ensure we advanced position and input value count in this iteration. - if (s->input_value_count == prev_input_value_count and s->src_pos == prev_src_pos) { - cg::invoke_one(block, [&] { s->set_error_code(decode_error::DELTA_PARAMS_UNSUPPORTED); }); - block.sync(); - break; - } } if (has_repetition) { @@ -565,8 +550,10 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size) // Propagate malformed-header errors from either underlying DELTA_BINARY_PACKED decoder. if (prefix_db->error or suffix_db->error) { - 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; } @@ -586,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; } @@ -785,7 +774,9 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size) // Propagate malformed-header errors from the underlying DELTA_BINARY_PACKED decoder if (db->error) { - 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; } From 95f3d8ac85a2361ff20ced0c7664c7c10b59bb41 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb Date: Thu, 14 May 2026 21:46:39 +0000 Subject: [PATCH 5/6] simplify check --- cpp/src/io/parquet/page_delta_decode.cu | 11 +++-------- 1 file changed, 3 insertions(+), 8 deletions(-) diff --git a/cpp/src/io/parquet/page_delta_decode.cu b/cpp/src/io/parquet/page_delta_decode.cu index 6bcb69357ee0..316e0417993d 100644 --- a/cpp/src/io/parquet/page_delta_decode.cu +++ b/cpp/src/io/parquet/page_delta_decode.cu @@ -772,8 +772,9 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size) } block.sync(); - // Propagate malformed-header errors from the underlying DELTA_BINARY_PACKED decoder - if (db->error) { + // sanity check to make sure we can process this page + auto const batch_size = db->values_per_mb; + 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); } @@ -782,12 +783,6 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size) 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); - return; - } // 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 From c839f1b87ac5a85ab60a62bcfc8d1fb073161340 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb Date: Thu, 14 May 2026 21:48:35 +0000 Subject: [PATCH 6/6] Fix nit --- cpp/src/io/parquet/page_delta_decode.cu | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/cpp/src/io/parquet/page_delta_decode.cu b/cpp/src/io/parquet/page_delta_decode.cu index 316e0417993d..82ac47391811 100644 --- a/cpp/src/io/parquet/page_delta_decode.cu +++ b/cpp/src/io/parquet/page_delta_decode.cu @@ -776,7 +776,8 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size) auto const batch_size = db->values_per_mb; 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); + set_error(static_cast(decode_error::DELTA_PARAMS_UNSUPPORTED), + error_code); } return; }