From 05368fe0e3ba8bdd53d83452c841edbfd55c650e Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Fri, 28 Aug 2026 19:44:52 +0000 Subject: [PATCH 01/10] Pre-zero required string offsets under a nullable ancestor --- cpp/src/io/parquet/page_decode.cuh | 14 +- cpp/src/io/parquet/page_delta_decode.cu | 11 +- cpp/src/io/parquet/reader_impl_preprocess.cu | 44 +++- cpp/tests/io/parquet_reader_test.cpp | 203 +++++++++++++++++++ 4 files changed, 261 insertions(+), 11 deletions(-) diff --git a/cpp/src/io/parquet/page_decode.cuh b/cpp/src/io/parquet/page_decode.cuh index 3c24912095aa..0a90987f5611 100644 --- a/cpp/src/io/parquet/page_decode.cuh +++ b/cpp/src/io/parquet/page_decode.cuh @@ -1458,10 +1458,14 @@ inline __device__ bool setup_local_page_info(auto* const s, /** * @brief Zero-fill null positions in output data using parallel per-validity-block processing * - * This function processes the validity bitmap and zero-fills all positions in the output - * data that correspond to null values. It uses a parallel approach where each thread - * handles one 32-bit validity block at a time, looping only over the zero bits (null positions) - * within that block. + * Each thread handles one 32-bit validity block and zero-fills only its null positions. + * + * @note This handles only nulls in a leaf's own bitmap. Nulls inherited from `optional` + * ancestors are zero-filled by `reader_impl::allocate_columns` because an ancestor's validity map + * may not be available to this leaf. + * + * Callers use this for structural outputs: string lengths, list offsets, and dictionary indices. + * Fixed-width null values need no initialization because they are masked. * * @tparam block_size CUDA block size for the kernel * @param s Page state containing all necessary information @@ -1481,7 +1485,7 @@ __device__ void zero_fill_null_positions_shared( int const leaf_level_index = s->setup.col.max_nesting_depth - 1; auto const& ni = s->nesting.nesting_info[leaf_level_index]; - // Check if we have nulls to fill + // Check if this leaf has a validity map to zero out nulls if ((ni.valid_map == nullptr) || (num_values == 0)) { return; } auto const data_out = ni.data_out; diff --git a/cpp/src/io/parquet/page_delta_decode.cu b/cpp/src/io/parquet/page_delta_decode.cu index 72425abc59de..0041e40bd44c 100644 --- a/cpp/src/io/parquet/page_delta_decode.cu +++ b/cpp/src/io/parquet/page_delta_decode.cu @@ -483,11 +483,12 @@ CUDF_KERNEL void __launch_bounds__(decode_delta_binary_block_size) auto const& ni = s->nesting.nesting_info[s->setup.col.max_nesting_depth - 1]; if (ni.valid_map != nullptr) { int const num_values = ni.valid_map_offset - init_valid_map_offset; - zero_fill_null_positions_shared(s, - s->output_cvt.dtype_len, - init_valid_map_offset, - num_values, - static_cast(block.thread_rank())); + zero_fill_null_positions_shared( + s, + s->output_cvt.dtype_len, + init_valid_map_offset, + num_values, + static_cast(block.thread_rank())); } } diff --git a/cpp/src/io/parquet/reader_impl_preprocess.cu b/cpp/src/io/parquet/reader_impl_preprocess.cu index 6637d54bc0cd..16ba74aba2b2 100644 --- a/cpp/src/io/parquet/reader_impl_preprocess.cu +++ b/cpp/src/io/parquet/reader_impl_preprocess.cu @@ -948,14 +948,31 @@ void reader_impl::allocate_columns(read_mode mode, size_t skip_rows, size_t num_ // Validity Buffer is a uint32_t pointer std::vector> nullmask_bufs; + // An optional ancestor leaves unwritten output slots until the next repeated level. + // Pre-zero non-nullable STRING buffers as their lengths are converted to offsets. + // Not handling needed here for nullable strings (zero-filled by decoder using their own validity + // bitmap),fixed-width (masked), LIST offsets (never have gaps), and dictionary indices + // (have no ancestors). + auto const compute_has_unwritten_slots = [](auto const& out_buf, bool has_nullable_ancestor) { + return has_nullable_ancestor and out_buf.type.id() == type_id::STRING and + not out_buf.is_nullable; + }; + auto unwritten_bufs = cudf::detail::make_empty_pinned_vector>( + _input_columns.size(), _stream); + for (auto const& input_col : _input_columns) { size_t const max_depth = input_col.nesting_depth(); - auto* cols = &_output_buffers; + auto* cols = &_output_buffers; + bool has_nullable_ancestor = false; for (size_t l_idx = 0; l_idx < max_depth; l_idx++) { auto& out_buf = (*cols)[input_col.nesting[l_idx]]; cols = &out_buf.children; + auto const has_unwritten_slots = compute_has_unwritten_slots(out_buf, has_nullable_ancestor); + has_nullable_ancestor = + out_buf.type.id() == type_id::LIST ? false : (has_nullable_ancestor or out_buf.is_nullable); + // if this has a list parent, we have to get column sizes from the // data computed during compute_page_sizes if (out_buf.user_data & PARQUET_COLUMN_BUFFER_FLAG_HAS_LIST_PARENT) { @@ -976,6 +993,10 @@ void reader_impl::allocate_columns(read_mode mode, size_t skip_rows, size_t num_ out_buf.null_mask(), cudf::util::round_up_safe(out_buf.null_mask_size(), sizeof(cudf::bitmask_type)) / sizeof(cudf::bitmask_type)); + if (has_unwritten_slots and out_buf.data() != nullptr) { + unwritten_bufs.emplace_back(static_cast(out_buf.data()), + out_buf.data_size()); + } } } } @@ -1068,10 +1089,19 @@ void reader_impl::allocate_columns(read_mode mode, size_t skip_rows, size_t num_ for (size_type idx = 0; idx < static_cast(_input_columns.size()); idx++) { auto const& input_col = _input_columns[idx]; auto* cols = &_output_buffers; + // See the identically named variable in the non-list allocation loop above + bool has_nullable_ancestor = false; for (size_type l_idx = 0; l_idx < static_cast(input_col.nesting_depth()); l_idx++) { auto& out_buf = (*cols)[input_col.nesting[l_idx]]; cols = &out_buf.children; + + auto const has_unwritten_slots = + compute_has_unwritten_slots(out_buf, has_nullable_ancestor); + has_nullable_ancestor = out_buf.type.id() == type_id::LIST + ? false + : (has_nullable_ancestor or out_buf.is_nullable); + // if this buffer is part of a list hierarchy, we need to determine it's // final size and allocate it here. // @@ -1095,6 +1125,10 @@ void reader_impl::allocate_columns(read_mode mode, size_t skip_rows, size_t num_ out_buf.null_mask(), cudf::util::round_up_safe(out_buf.null_mask_size(), sizeof(cudf::bitmask_type)) / sizeof(cudf::bitmask_type)); + if (has_unwritten_slots and out_buf.data() != nullptr) { + unwritten_bufs.emplace_back(static_cast(out_buf.data()), + out_buf.data_size()); + } } } } @@ -1105,6 +1139,14 @@ void reader_impl::allocate_columns(read_mode mode, size_t skip_rows, size_t num_ cudf::host_span const>{nullmask_bufs}, _stream); cudf::detail::batched_memset( pinned_nullmask_bufs, std::numeric_limits::max(), _stream); + + // Need to zero non-nullable string lengths with nullable ancestors + if (not unwritten_bufs.empty()) { + cudf::detail::batched_memset( + cudf::host_span const>{unwritten_bufs}, + static_cast(0), + _stream); + } } void reader_impl::fill_pruned_offsets(size_t skip_rows, diff --git a/cpp/tests/io/parquet_reader_test.cpp b/cpp/tests/io/parquet_reader_test.cpp index a77ad1fd1de9..f4f3e87b2a5c 100644 --- a/cpp/tests/io/parquet_reader_test.cpp +++ b/cpp/tests/io/parquet_reader_test.cpp @@ -17,6 +17,7 @@ #include #include +#include #include #include #include @@ -35,10 +36,12 @@ #include #include #include +#include #include #include #include #include +#include #include using ParquetDecompressionTest = DecompressionTest; @@ -6337,3 +6340,203 @@ TEST_F(ParquetReaderTest, NestedMismatchedSchemaColumnValidation) EXPECT_THROW(cudf::io::read_parquet(opts), std::invalid_argument); } } +namespace { + +/** + * @brief Create an optional struct with required children + * + * @param children Child columns without null masks + * @param num_rows Number of rows + * @return Struct column with every seventh row null + */ +std::unique_ptr make_optional_struct( + std::vector>&& children, cudf::size_type num_rows) +{ + auto validity = + cudf::detail::make_counting_transform_iterator(0, [](auto i) { return (i % 7) != 0; }); + auto [null_mask, null_count] = cudf::test::detail::make_null_mask(validity, validity + num_rows); + return cudf::create_structs_hierarchy( + num_rows, std::move(children), null_count, std::move(null_mask)); +} + +} // namespace + +TEST_F(ParquetReaderTest, StructTwoRequiredChildrenNullGaps) +{ + // Two required string children below one nullable struct. + constexpr cudf::size_type num_rows = 2000; + constexpr auto const value = std::string_view{"fixed_width_payload"}; + + auto const values = cuda::make_constant_iterator(value); + cudf::test::strings_column_wrapper a_col{values, values + num_rows}; + cudf::test::strings_column_wrapper b_col{values, values + num_rows}; + + std::vector> children; + children.push_back(a_col.release()); + children.push_back(b_col.release()); + auto struct_col = make_optional_struct(std::move(children), num_rows); + auto expected = cudf::purge_nonempty_nulls(struct_col->view()); + + auto const written = table_view{{struct_col->view()}}; + cudf::io::table_input_metadata input_metadata(written); + input_metadata.column_metadata[0].set_name("s"); + input_metadata.column_metadata[0].child(0).set_name("a").set_nullability(false); + input_metadata.column_metadata[0].child(1).set_name("b").set_nullability(false); + + auto const filepath = temp_env->get_temp_filepath("StructTwoRequiredChildrenNullGaps.parquet"); + cudf::io::write_parquet( + cudf::io::parquet_writer_options::builder(cudf::io::sink_info{filepath}, written) + .metadata(std::move(input_metadata)) + .dictionary_policy(cudf::io::dictionary_policy::NEVER) + .compression(cudf::io::compression_type::NONE) + .build()); + + auto const result = cudf::io::read_parquet( + cudf::io::parquet_reader_options::builder(cudf::io::source_info{filepath}).build()); + CUDF_TEST_EXPECT_TABLES_EQUAL(table_view{{expected->view()}}, result.tbl->view()); +} + +TEST_F(ParquetReaderTest, NestedStructRequiredStringChildNullGaps) +{ + // The nullable ancestor is separated from the required string by a required struct. + constexpr cudf::size_type num_rows = 2000; + constexpr auto const value = std::string_view{"fixed_width_payload"}; + + auto const values = cuda::make_constant_iterator(value); + cudf::test::strings_column_wrapper child_col{values, values + num_rows}; + + std::vector> inner_children; + inner_children.push_back(child_col.release()); + auto inner_struct = + cudf::create_structs_hierarchy(num_rows, std::move(inner_children), 0, rmm::device_buffer{}); + + std::vector> outer_children; + outer_children.push_back(std::move(inner_struct)); + auto outer_struct = make_optional_struct(std::move(outer_children), num_rows); + auto expected = cudf::purge_nonempty_nulls(outer_struct->view()); + + auto const written = table_view{{outer_struct->view()}}; + cudf::io::table_input_metadata input_metadata(written); + input_metadata.column_metadata[0] + .set_name("outer") + .child(0) + .set_name("inner") + .set_nullability(false) + .child(0) + .set_name("value") + .set_nullability(false); + + auto const filepath = + temp_env->get_temp_filepath("NestedStructRequiredStringChildNullGaps.parquet"); + cudf::io::write_parquet( + cudf::io::parquet_writer_options::builder(cudf::io::sink_info{filepath}, written) + .metadata(std::move(input_metadata)) + .dictionary_policy(cudf::io::dictionary_policy::NEVER) + .compression(cudf::io::compression_type::NONE) + .build()); + + auto const result = cudf::io::read_parquet( + cudf::io::parquet_reader_options::builder(cudf::io::source_info{filepath}).build()); + CUDF_TEST_EXPECT_TABLES_EQUAL(table_view{{expected->view()}}, result.tbl->view()); +} + +TEST_F(ParquetReaderTest, ListOfStructRequiredStringChildNullGaps) +{ + // Exercise list allocation with null structs inside the list. + constexpr cudf::size_type num_lists = 500; + constexpr cudf::size_type list_size = 4; + constexpr cudf::size_type num_elements = num_lists * list_size; + constexpr auto const value = std::string_view{"fixed_width_payload"}; + + auto const values = cuda::make_constant_iterator(value); + cudf::test::strings_column_wrapper child_col{values, values + num_elements}; + + std::vector> children; + children.push_back(child_col.release()); + auto struct_col = make_optional_struct(std::move(children), num_elements); + + auto offsets = cudf::detail::make_counting_transform_iterator( + 0, [](auto i) { return static_cast(i * list_size); }); + column_wrapper offsets_col(offsets, offsets + num_lists + 1); + + auto list_col = cudf::make_lists_column( + num_lists, offsets_col.release(), std::move(struct_col), 0, rmm::device_buffer{}); + auto expected = cudf::purge_nonempty_nulls(list_col->view()); + + auto const written = table_view{{list_col->view()}}; + cudf::io::table_input_metadata input_metadata(written); + input_metadata.column_metadata[0] + .set_name("l") + .child(1) + .set_name("s") + .child(0) + .set_name("a") + .set_nullability(false); + + auto const filepath = + temp_env->get_temp_filepath("ListOfStructRequiredStringChildNullGaps.parquet"); + cudf::io::write_parquet( + cudf::io::parquet_writer_options::builder(cudf::io::sink_info{filepath}, written) + .metadata(std::move(input_metadata)) + .dictionary_policy(cudf::io::dictionary_policy::NEVER) + .compression(cudf::io::compression_type::NONE) + .build()); + + auto const result = cudf::io::read_parquet( + cudf::io::parquet_reader_options::builder(cudf::io::source_info{filepath}).build()); + CUDF_TEST_EXPECT_TABLES_EQUAL(table_view{{expected->view()}}, result.tbl->view()); +} + +TEST_F(ParquetReaderTest, StructRequiredListChildNullGaps) +{ + // `optional struct { required list }` writes LIST offsets even when the struct is null, + // so no gap needs zero-filling. + constexpr cudf::size_type num_rows = 2000; + constexpr cudf::size_type list_size = 4; + constexpr cudf::size_type num_elements = num_rows * list_size; + + constexpr int32_t null_gap_value = 0x5a5a5a5a; + + auto values = cudf::detail::make_counting_transform_iterator( + 0, [null_gap_value](auto) { return null_gap_value; }); + column_wrapper leaf_col(values, values + num_elements); + + auto offsets = cudf::detail::make_counting_transform_iterator( + 0, [](auto i) { return static_cast(i * list_size); }); + column_wrapper offsets_col(offsets, offsets + num_rows + 1); + + auto list_col = cudf::make_lists_column( + num_rows, offsets_col.release(), leaf_col.release(), 0, rmm::device_buffer{}); + + std::vector> children; + children.push_back(std::move(list_col)); + auto struct_col = make_optional_struct(std::move(children), num_rows); + + auto const written = table_view{{struct_col->view()}}; + cudf::io::table_input_metadata input_metadata(written); + input_metadata.column_metadata[0].set_name("s"); + input_metadata.column_metadata[0].child(0).set_name("l").set_nullability(false); + + auto const filepath = temp_env->get_temp_filepath("StructRequiredListChildNullGaps.parquet"); + cudf::io::write_parquet( + cudf::io::parquet_writer_options::builder(cudf::io::sink_info{filepath}, written) + .metadata(std::move(input_metadata)) + .dictionary_policy(cudf::io::dictionary_policy::NEVER) + .compression(cudf::io::compression_type::NONE) + .build()); + + auto const result = cudf::io::read_parquet( + cudf::io::parquet_reader_options::builder(cudf::io::source_info{filepath}).build()); + auto const list_child = result.tbl->view().column(0).child(0); + auto const lcv = cudf::lists_column_view{list_child}; + auto const output_offsets = cudf::test::to_host(lcv.offsets()).first; + ASSERT_EQ(output_offsets.size(), static_cast(num_rows) + 1); + ASSERT_EQ(output_offsets.front(), 0); + for (std::size_t i = 1; i < output_offsets.size(); ++i) { + ASSERT_GE(output_offsets[i], output_offsets[i - 1]) + << std::format("non-monotonic list offset at {}", i); + ASSERT_LE(output_offsets[i] - output_offsets[i - 1], list_size) + << std::format("oversized list at {}", i - 1); + } + ASSERT_EQ(output_offsets.back(), lcv.child().size()); +} From 67fd4641659fa1849438c3ffbeb904b196d3d391 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Fri, 28 Aug 2026 19:53:44 +0000 Subject: [PATCH 02/10] Remove unnecessary test --- cpp/tests/io/parquet_reader_test.cpp | 53 ---------------------------- 1 file changed, 53 deletions(-) diff --git a/cpp/tests/io/parquet_reader_test.cpp b/cpp/tests/io/parquet_reader_test.cpp index f4f3e87b2a5c..614d69c91ea8 100644 --- a/cpp/tests/io/parquet_reader_test.cpp +++ b/cpp/tests/io/parquet_reader_test.cpp @@ -36,7 +36,6 @@ #include #include #include -#include #include #include #include @@ -6487,56 +6486,4 @@ TEST_F(ParquetReaderTest, ListOfStructRequiredStringChildNullGaps) CUDF_TEST_EXPECT_TABLES_EQUAL(table_view{{expected->view()}}, result.tbl->view()); } -TEST_F(ParquetReaderTest, StructRequiredListChildNullGaps) -{ - // `optional struct { required list }` writes LIST offsets even when the struct is null, - // so no gap needs zero-filling. - constexpr cudf::size_type num_rows = 2000; - constexpr cudf::size_type list_size = 4; - constexpr cudf::size_type num_elements = num_rows * list_size; - - constexpr int32_t null_gap_value = 0x5a5a5a5a; - - auto values = cudf::detail::make_counting_transform_iterator( - 0, [null_gap_value](auto) { return null_gap_value; }); - column_wrapper leaf_col(values, values + num_elements); - - auto offsets = cudf::detail::make_counting_transform_iterator( - 0, [](auto i) { return static_cast(i * list_size); }); - column_wrapper offsets_col(offsets, offsets + num_rows + 1); - - auto list_col = cudf::make_lists_column( - num_rows, offsets_col.release(), leaf_col.release(), 0, rmm::device_buffer{}); - - std::vector> children; - children.push_back(std::move(list_col)); - auto struct_col = make_optional_struct(std::move(children), num_rows); - - auto const written = table_view{{struct_col->view()}}; - cudf::io::table_input_metadata input_metadata(written); - input_metadata.column_metadata[0].set_name("s"); - input_metadata.column_metadata[0].child(0).set_name("l").set_nullability(false); - auto const filepath = temp_env->get_temp_filepath("StructRequiredListChildNullGaps.parquet"); - cudf::io::write_parquet( - cudf::io::parquet_writer_options::builder(cudf::io::sink_info{filepath}, written) - .metadata(std::move(input_metadata)) - .dictionary_policy(cudf::io::dictionary_policy::NEVER) - .compression(cudf::io::compression_type::NONE) - .build()); - - auto const result = cudf::io::read_parquet( - cudf::io::parquet_reader_options::builder(cudf::io::source_info{filepath}).build()); - auto const list_child = result.tbl->view().column(0).child(0); - auto const lcv = cudf::lists_column_view{list_child}; - auto const output_offsets = cudf::test::to_host(lcv.offsets()).first; - ASSERT_EQ(output_offsets.size(), static_cast(num_rows) + 1); - ASSERT_EQ(output_offsets.front(), 0); - for (std::size_t i = 1; i < output_offsets.size(); ++i) { - ASSERT_GE(output_offsets[i], output_offsets[i - 1]) - << std::format("non-monotonic list offset at {}", i); - ASSERT_LE(output_offsets[i] - output_offsets[i - 1], list_size) - << std::format("oversized list at {}", i - 1); - } - ASSERT_EQ(output_offsets.back(), lcv.child().size()); -} From 60ef3cd724a02c8d3f38122b1f05460c80977ca1 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Fri, 28 Aug 2026 19:56:31 +0000 Subject: [PATCH 03/10] minor --- cpp/src/io/parquet/page_decode.cuh | 4 ++-- cpp/src/io/parquet/reader_impl_preprocess.cu | 8 ++++---- cpp/tests/io/parquet_reader_test.cpp | 6 +++--- 3 files changed, 9 insertions(+), 9 deletions(-) diff --git a/cpp/src/io/parquet/page_decode.cuh b/cpp/src/io/parquet/page_decode.cuh index 0a90987f5611..f973584bcc65 100644 --- a/cpp/src/io/parquet/page_decode.cuh +++ b/cpp/src/io/parquet/page_decode.cuh @@ -1464,8 +1464,8 @@ inline __device__ bool setup_local_page_info(auto* const s, * ancestors are zero-filled by `reader_impl::allocate_columns` because an ancestor's validity map * may not be available to this leaf. * - * Callers use this for structural outputs: string lengths, list offsets, and dictionary indices. - * Fixed-width null values need no initialization because they are masked. + * Callers use this for structural outputs: nullable string lengths, list offsets, and dictionary + * indices. Fixed-width null values need no initialization because they are masked. * * @tparam block_size CUDA block size for the kernel * @param s Page state containing all necessary information diff --git a/cpp/src/io/parquet/reader_impl_preprocess.cu b/cpp/src/io/parquet/reader_impl_preprocess.cu index 42806315b552..c97c8b4e6837 100644 --- a/cpp/src/io/parquet/reader_impl_preprocess.cu +++ b/cpp/src/io/parquet/reader_impl_preprocess.cu @@ -994,8 +994,8 @@ void reader_impl::allocate_columns(read_mode mode, size_t skip_rows, size_t num_ cudf::util::round_up_safe(out_buf.null_mask_size(), sizeof(cudf::bitmask_type)) / sizeof(cudf::bitmask_type)); if (has_unwritten_slots and out_buf.data() != nullptr) { - unwritten_bufs.emplace_back(static_cast(out_buf.data()), - out_buf.data_size()); + unwritten_bufs.push_back(static_cast(out_buf.data()), + out_buf.data_size()); } } } @@ -1126,8 +1126,8 @@ void reader_impl::allocate_columns(read_mode mode, size_t skip_rows, size_t num_ cudf::util::round_up_safe(out_buf.null_mask_size(), sizeof(cudf::bitmask_type)) / sizeof(cudf::bitmask_type)); if (has_unwritten_slots and out_buf.data() != nullptr) { - unwritten_bufs.emplace_back(static_cast(out_buf.data()), - out_buf.data_size()); + unwritten_bufs.push_back(static_cast(out_buf.data()), + out_buf.data_size()); } } } diff --git a/cpp/tests/io/parquet_reader_test.cpp b/cpp/tests/io/parquet_reader_test.cpp index 614d69c91ea8..f2efb1268fae 100644 --- a/cpp/tests/io/parquet_reader_test.cpp +++ b/cpp/tests/io/parquet_reader_test.cpp @@ -6465,11 +6465,11 @@ TEST_F(ParquetReaderTest, ListOfStructRequiredStringChildNullGaps) auto const written = table_view{{list_col->view()}}; cudf::io::table_input_metadata input_metadata(written); input_metadata.column_metadata[0] - .set_name("l") + .set_name("outer") .child(1) - .set_name("s") + .set_name("inner") .child(0) - .set_name("a") + .set_name("value") .set_nullability(false); auto const filepath = From 782058010b0119fb9662acf4a1a59a8130fd115d Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Fri, 28 Aug 2026 20:00:01 +0000 Subject: [PATCH 04/10] minor docs update --- cpp/src/io/parquet/page_decode.cuh | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/cpp/src/io/parquet/page_decode.cuh b/cpp/src/io/parquet/page_decode.cuh index f973584bcc65..0e6a4dbd8a58 100644 --- a/cpp/src/io/parquet/page_decode.cuh +++ b/cpp/src/io/parquet/page_decode.cuh @@ -1458,7 +1458,8 @@ inline __device__ bool setup_local_page_info(auto* const s, /** * @brief Zero-fill null positions in output data using parallel per-validity-block processing * - * Each thread handles one 32-bit validity block and zero-fills only its null positions. + * Each warp handles one 32-bit validity block, with each lane zero-filling a single null position. + * Remaining blocks that exceed the warp count are handled `process_block_sequential`. * * @note This handles only nulls in a leaf's own bitmap. Nulls inherited from `optional` * ancestors are zero-filled by `reader_impl::allocate_columns` because an ancestor's validity map From a1fe6f24ba4c19665f0c77e3f7e80d580d415ed9 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Fri, 28 Aug 2026 20:01:18 +0000 Subject: [PATCH 05/10] style --- cpp/tests/io/parquet_reader_test.cpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/cpp/tests/io/parquet_reader_test.cpp b/cpp/tests/io/parquet_reader_test.cpp index f2efb1268fae..02c7d9a68176 100644 --- a/cpp/tests/io/parquet_reader_test.cpp +++ b/cpp/tests/io/parquet_reader_test.cpp @@ -6485,5 +6485,3 @@ TEST_F(ParquetReaderTest, ListOfStructRequiredStringChildNullGaps) cudf::io::parquet_reader_options::builder(cudf::io::source_info{filepath}).build()); CUDF_TEST_EXPECT_TABLES_EQUAL(table_view{{expected->view()}}, result.tbl->view()); } - - From 665bee9e456c9a9698e2372cbb1c1da24b42248e Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Fri, 28 Aug 2026 20:14:35 +0000 Subject: [PATCH 06/10] minor syntax fix --- cpp/src/io/parquet/reader_impl_preprocess.cu | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/cpp/src/io/parquet/reader_impl_preprocess.cu b/cpp/src/io/parquet/reader_impl_preprocess.cu index c97c8b4e6837..64948f30b542 100644 --- a/cpp/src/io/parquet/reader_impl_preprocess.cu +++ b/cpp/src/io/parquet/reader_impl_preprocess.cu @@ -994,8 +994,8 @@ void reader_impl::allocate_columns(read_mode mode, size_t skip_rows, size_t num_ cudf::util::round_up_safe(out_buf.null_mask_size(), sizeof(cudf::bitmask_type)) / sizeof(cudf::bitmask_type)); if (has_unwritten_slots and out_buf.data() != nullptr) { - unwritten_bufs.push_back(static_cast(out_buf.data()), - out_buf.data_size()); + unwritten_bufs.push_back({static_cast(out_buf.data()), + out_buf.data_size()}); } } } @@ -1126,8 +1126,8 @@ void reader_impl::allocate_columns(read_mode mode, size_t skip_rows, size_t num_ cudf::util::round_up_safe(out_buf.null_mask_size(), sizeof(cudf::bitmask_type)) / sizeof(cudf::bitmask_type)); if (has_unwritten_slots and out_buf.data() != nullptr) { - unwritten_bufs.push_back(static_cast(out_buf.data()), - out_buf.data_size()); + unwritten_bufs.push_back({static_cast(out_buf.data()), + out_buf.data_size()}); } } } From 5b7f7cf3dc6604e6be8a28d4145a016042d49ede Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Fri, 28 Aug 2026 22:06:55 +0000 Subject: [PATCH 07/10] clang-format for the trillionth time --- cpp/src/io/parquet/reader_impl_preprocess.cu | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/cpp/src/io/parquet/reader_impl_preprocess.cu b/cpp/src/io/parquet/reader_impl_preprocess.cu index 64948f30b542..92cfc57fcfad 100644 --- a/cpp/src/io/parquet/reader_impl_preprocess.cu +++ b/cpp/src/io/parquet/reader_impl_preprocess.cu @@ -994,8 +994,8 @@ void reader_impl::allocate_columns(read_mode mode, size_t skip_rows, size_t num_ cudf::util::round_up_safe(out_buf.null_mask_size(), sizeof(cudf::bitmask_type)) / sizeof(cudf::bitmask_type)); if (has_unwritten_slots and out_buf.data() != nullptr) { - unwritten_bufs.push_back({static_cast(out_buf.data()), - out_buf.data_size()}); + unwritten_bufs.push_back( + {static_cast(out_buf.data()), out_buf.data_size()}); } } } @@ -1126,8 +1126,8 @@ void reader_impl::allocate_columns(read_mode mode, size_t skip_rows, size_t num_ cudf::util::round_up_safe(out_buf.null_mask_size(), sizeof(cudf::bitmask_type)) / sizeof(cudf::bitmask_type)); if (has_unwritten_slots and out_buf.data() != nullptr) { - unwritten_bufs.push_back({static_cast(out_buf.data()), - out_buf.data_size()}); + unwritten_bufs.push_back( + {static_cast(out_buf.data()), out_buf.data_size()}); } } } From b06df3b2cae40f7d959d99605ee51689625c1a61 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Fri, 28 Aug 2026 22:15:11 +0000 Subject: [PATCH 08/10] Improve doc --- cpp/src/io/parquet/reader_impl_preprocess.cu | 12 +++++++----- 1 file changed, 7 insertions(+), 5 deletions(-) diff --git a/cpp/src/io/parquet/reader_impl_preprocess.cu b/cpp/src/io/parquet/reader_impl_preprocess.cu index 92cfc57fcfad..adf97205f45d 100644 --- a/cpp/src/io/parquet/reader_impl_preprocess.cu +++ b/cpp/src/io/parquet/reader_impl_preprocess.cu @@ -948,11 +948,13 @@ void reader_impl::allocate_columns(read_mode mode, size_t skip_rows, size_t num_ // Validity Buffer is a uint32_t pointer std::vector> nullmask_bufs; - // An optional ancestor leaves unwritten output slots until the next repeated level. - // Pre-zero non-nullable STRING buffers as their lengths are converted to offsets. - // Not handling needed here for nullable strings (zero-filled by decoder using their own validity - // bitmap),fixed-width (masked), LIST offsets (never have gaps), and dictionary indices - // (have no ancestors). + // An optional ancestor leaves unwritten output slots until the next repeated level. So, for a + // non-nullable STRING (FIELD) with a nullable ancestor, the column is nullable and not all rows + // will be decoded. The decoder may not detect this because it may not have a validity map from + // the ancestor. To avoid this, zero-fill such STRING buffers here as their uninitialized lengths + // are converted to offsets. No handling needed here for nullable strings (zero-filled by decoder + // using their own validity bitmap), fixed-width (masked), LIST offsets (never have gaps), and + // dictionary indices (have no ancestors). auto const compute_has_unwritten_slots = [](auto const& out_buf, bool has_nullable_ancestor) { return has_nullable_ancestor and out_buf.type.id() == type_id::STRING and not out_buf.is_nullable; From 67431e051d86a0d0c0cc1818809ea5e7eca7bbb1 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Sat, 29 Aug 2026 06:04:54 +0000 Subject: [PATCH 09/10] Fix tests --- cpp/tests/io/parquet_reader_test.cpp | 185 ++++++++++++++++++--------- 1 file changed, 121 insertions(+), 64 deletions(-) diff --git a/cpp/tests/io/parquet_reader_test.cpp b/cpp/tests/io/parquet_reader_test.cpp index 02c7d9a68176..64c6d6da7348 100644 --- a/cpp/tests/io/parquet_reader_test.cpp +++ b/cpp/tests/io/parquet_reader_test.cpp @@ -6346,23 +6346,31 @@ namespace { * * @param children Child columns without null masks * @param num_rows Number of rows - * @return Struct column with every seventh row null + * @param should_propagate_nulls Whether to push the struct's nulls down into its children + * @return Optional struct column with every seventh row null */ std::unique_ptr make_optional_struct( - std::vector>&& children, cudf::size_type num_rows) + std::vector>&& children, + cudf::size_type num_rows, + bool should_propagate_nulls) { - auto validity = + auto const validity = cudf::detail::make_counting_transform_iterator(0, [](auto i) { return (i % 7) != 0; }); auto [null_mask, null_count] = cudf::test::detail::make_null_mask(validity, validity + num_rows); - return cudf::create_structs_hierarchy( - num_rows, std::move(children), null_count, std::move(null_mask)); + return should_propagate_nulls + ? cudf::make_structs_column( + num_rows, std::move(children), null_count, std::move(null_mask)) + : cudf::create_structs_hierarchy( + num_rows, std::move(children), null_count, std::move(null_mask)); } } // namespace -TEST_F(ParquetReaderTest, StructTwoRequiredChildrenNullGaps) +TEST_F(ParquetReaderTest, TwoRequiredStringLeavesWithNullableAncestor) { - // Two required string children below one nullable struct. + // Build a table with an optional struct { required string, required string } column + // to test the case with two required strings sharing an immediate nullable ancestor. + constexpr cudf::size_type num_rows = 2000; constexpr auto const value = std::string_view{"fixed_width_payload"}; @@ -6373,31 +6381,45 @@ TEST_F(ParquetReaderTest, StructTwoRequiredChildrenNullGaps) std::vector> children; children.push_back(a_col.release()); children.push_back(b_col.release()); - auto struct_col = make_optional_struct(std::move(children), num_rows); - auto expected = cudf::purge_nonempty_nulls(struct_col->view()); + auto struct_col = make_optional_struct(std::move(children), num_rows, false); - auto const written = table_view{{struct_col->view()}}; - cudf::io::table_input_metadata input_metadata(written); - input_metadata.column_metadata[0].set_name("s"); - input_metadata.column_metadata[0].child(0).set_name("a").set_nullability(false); - input_metadata.column_metadata[0].child(1).set_name("b").set_nullability(false); + auto const filepath = + temp_env->get_temp_filepath("TwoRequiredStringLeavesWithNullableAncestor.parquet"); - auto const filepath = temp_env->get_temp_filepath("StructTwoRequiredChildrenNullGaps.parquet"); - cudf::io::write_parquet( - cudf::io::parquet_writer_options::builder(cudf::io::sink_info{filepath}, written) - .metadata(std::move(input_metadata)) - .dictionary_policy(cudf::io::dictionary_policy::NEVER) - .compression(cudf::io::compression_type::NONE) - .build()); + // Write the table to Parquet + { + auto const written = table_view{{struct_col->view()}}; + cudf::io::table_input_metadata input_metadata(written); + input_metadata.column_metadata[0].set_name("s"); + input_metadata.column_metadata[0].child(0).set_name("a").set_nullability(false); + input_metadata.column_metadata[0].child(1).set_name("b").set_nullability(false); + + cudf::io::write_parquet( + cudf::io::parquet_writer_options::builder(cudf::io::sink_info{filepath}, written) + .metadata(std::move(input_metadata)) + .dictionary_policy(cudf::io::dictionary_policy::NEVER) + .compression(cudf::io::compression_type::NONE) + .build()); + } + // Build expected table from written struct's children + auto const expected = + make_optional_struct(std::move(struct_col->release().children), num_rows, true); + + // Read the table from Parquet auto const result = cudf::io::read_parquet( cudf::io::parquet_reader_options::builder(cudf::io::source_info{filepath}).build()); + + // Compare CUDF_TEST_EXPECT_TABLES_EQUAL(table_view{{expected->view()}}, result.tbl->view()); } -TEST_F(ParquetReaderTest, NestedStructRequiredStringChildNullGaps) +TEST_F(ParquetReaderTest, RequiredStringLeafWithSeparatedNullableAncestor) { - // The nullable ancestor is separated from the required string by a required struct. + // Build a table with an optional struct { required struct { required string } } column + // to test the case with an nullable ancestor of a required string separated by a required + // struct. + constexpr cudf::size_type num_rows = 2000; constexpr auto const value = std::string_view{"fixed_width_payload"}; @@ -6411,37 +6433,48 @@ TEST_F(ParquetReaderTest, NestedStructRequiredStringChildNullGaps) std::vector> outer_children; outer_children.push_back(std::move(inner_struct)); - auto outer_struct = make_optional_struct(std::move(outer_children), num_rows); - auto expected = cudf::purge_nonempty_nulls(outer_struct->view()); - - auto const written = table_view{{outer_struct->view()}}; - cudf::io::table_input_metadata input_metadata(written); - input_metadata.column_metadata[0] - .set_name("outer") - .child(0) - .set_name("inner") - .set_nullability(false) - .child(0) - .set_name("value") - .set_nullability(false); + auto outer_struct = make_optional_struct(std::move(outer_children), num_rows, false); auto const filepath = - temp_env->get_temp_filepath("NestedStructRequiredStringChildNullGaps.parquet"); - cudf::io::write_parquet( - cudf::io::parquet_writer_options::builder(cudf::io::sink_info{filepath}, written) - .metadata(std::move(input_metadata)) - .dictionary_policy(cudf::io::dictionary_policy::NEVER) - .compression(cudf::io::compression_type::NONE) - .build()); + temp_env->get_temp_filepath("RequiredStringLeafWithSeparatedNullableAncestor.parquet"); + + // Write the table to Parquet + { + auto const written = table_view{{outer_struct->view()}}; + cudf::io::table_input_metadata input_metadata(written); + input_metadata.column_metadata[0] + .set_name("outer") + .child(0) + .set_name("inner") + .set_nullability(false) + .child(0) + .set_name("value") + .set_nullability(false); + cudf::io::write_parquet( + cudf::io::parquet_writer_options::builder(cudf::io::sink_info{filepath}, written) + .metadata(std::move(input_metadata)) + .dictionary_policy(cudf::io::dictionary_policy::NEVER) + .compression(cudf::io::compression_type::NONE) + .build()); + } + + // Build expected table from written struct's children + auto const expected = + make_optional_struct(std::move(outer_struct->release().children), num_rows, true); + // Read the table from Parquet auto const result = cudf::io::read_parquet( cudf::io::parquet_reader_options::builder(cudf::io::source_info{filepath}).build()); + + // Compare CUDF_TEST_EXPECT_TABLES_EQUAL(table_view{{expected->view()}}, result.tbl->view()); } -TEST_F(ParquetReaderTest, ListOfStructRequiredStringChildNullGaps) +TEST_F(ParquetReaderTest, RequiredStringLeafWithNullableAncestorUnderList) { - // Exercise list allocation with null structs inside the list. + // Build a table with a required list column to test the + // case with a nullable ancestor of a required string inside a list + constexpr cudf::size_type num_lists = 500; constexpr cudf::size_type list_size = 4; constexpr cudf::size_type num_elements = num_lists * list_size; @@ -6452,7 +6485,7 @@ TEST_F(ParquetReaderTest, ListOfStructRequiredStringChildNullGaps) std::vector> children; children.push_back(child_col.release()); - auto struct_col = make_optional_struct(std::move(children), num_elements); + auto struct_col = make_optional_struct(std::move(children), num_elements, false); auto offsets = cudf::detail::make_counting_transform_iterator( 0, [](auto i) { return static_cast(i * list_size); }); @@ -6460,28 +6493,52 @@ TEST_F(ParquetReaderTest, ListOfStructRequiredStringChildNullGaps) auto list_col = cudf::make_lists_column( num_lists, offsets_col.release(), std::move(struct_col), 0, rmm::device_buffer{}); - auto expected = cudf::purge_nonempty_nulls(list_col->view()); - - auto const written = table_view{{list_col->view()}}; - cudf::io::table_input_metadata input_metadata(written); - input_metadata.column_metadata[0] - .set_name("outer") - .child(1) - .set_name("inner") - .child(0) - .set_name("value") - .set_nullability(false); auto const filepath = - temp_env->get_temp_filepath("ListOfStructRequiredStringChildNullGaps.parquet"); - cudf::io::write_parquet( - cudf::io::parquet_writer_options::builder(cudf::io::sink_info{filepath}, written) - .metadata(std::move(input_metadata)) - .dictionary_policy(cudf::io::dictionary_policy::NEVER) - .compression(cudf::io::compression_type::NONE) - .build()); + temp_env->get_temp_filepath("RequiredStringLeafWithNullableAncestorUnderList.parquet"); + // Write the table to Parquet + { + auto const written = table_view{{list_col->view()}}; + cudf::io::table_input_metadata input_metadata(written); + input_metadata.column_metadata[0] + .set_name("outer") + .child(1) + .set_name("inner") + .child(0) + .set_name("value") + .set_nullability(false); + + cudf::io::write_parquet( + cudf::io::parquet_writer_options::builder(cudf::io::sink_info{filepath}, written) + .metadata(std::move(input_metadata)) + .dictionary_policy(cudf::io::dictionary_policy::NEVER) + .compression(cudf::io::compression_type::NONE) + .build()); + } + + // Build expected table from written offsets and the leaf strings + auto written_contents = list_col->release(); + auto exp_offsets = + std::move(written_contents.children[cudf::lists_column_view::offsets_column_index]); + auto const written_struct = + written_contents.children[cudf::lists_column_view::child_column_index]->view(); + + auto const struct_validity = cudf::is_valid(written_struct); + auto exp_leaf = + cudf::copy_if_else(written_struct.child(0), cudf::string_scalar{""}, struct_validity->view()); + + std::vector> exp_children; + exp_children.push_back(std::move(exp_leaf)); + auto exp_struct = make_optional_struct(std::move(exp_children), num_elements, false); + + auto const expected = cudf::make_lists_column( + num_lists, std::move(exp_offsets), std::move(exp_struct), 0, rmm::device_buffer{}); + + // Read the table from Parquet auto const result = cudf::io::read_parquet( cudf::io::parquet_reader_options::builder(cudf::io::source_info{filepath}).build()); + + // Compare CUDF_TEST_EXPECT_TABLES_EQUAL(table_view{{expected->view()}}, result.tbl->view()); } From 3d8a179fccc9ac24c6f8aff012720016f6a642eb Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Sat, 29 Aug 2026 06:10:19 +0000 Subject: [PATCH 10/10] Minor --- cpp/tests/io/parquet_reader_test.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/cpp/tests/io/parquet_reader_test.cpp b/cpp/tests/io/parquet_reader_test.cpp index 64c6d6da7348..043218125dde 100644 --- a/cpp/tests/io/parquet_reader_test.cpp +++ b/cpp/tests/io/parquet_reader_test.cpp @@ -17,7 +17,6 @@ #include #include -#include #include #include #include