diff --git a/cpp/src/io/parquet/decode_fixed.cu b/cpp/src/io/parquet/decode_fixed.cu index 2d262342ea9f..eafe418a67d2 100644 --- a/cpp/src/io/parquet/decode_fixed.cu +++ b/cpp/src/io/parquet/decode_fixed.cu @@ -1187,11 +1187,10 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size_t, 8) // - valid_count: number of non-null values we have decoded so far. In each iteration of the // loop below, we look at the number of valid items (which could be all for non-nullable), // and valid_count is that running count. - int processed_count = 0; - int valid_count = 0; - size_t string_output_offset = 0; - int const init_valid_map_offset = - s->nesting.nesting_info[s->setup.col.max_nesting_depth - 1].valid_map_offset; + int processed_count = 0; + int valid_count = 0; + size_t string_output_offset = 0; + int const init_valid_map_offset = init_null_fill_valid_map_offset(s); // Skip ahead in the decoding so that we don't repeat work skip_ahead_in_decoding uint32_t { - if constexpr (is_dict_int32_t) { return sizeof(int32_t); } - if constexpr (has_strings_t) { return sizeof(cudf::size_type); } - return s->output_cvt.dtype_len; - }(); - int const num_values = [&]() { - if constexpr (has_lists_t) { - auto const& ni = s->nesting.nesting_info[s->setup.col.max_nesting_depth - 1]; - return ni.valid_map_offset - init_valid_map_offset; - } else { - return s->setup.num_rows; - } - }(); - zero_fill_null_positions_shared( - s, dtype_len, init_valid_map_offset, num_values, t); - } + // Zero-fill null positions after decoding valid values. zero_fill_null_positions() itself is a + // no-op when there is no ancestor validity buffer to resolve, so this must not be gated on + // has_strings_t/has_lists_t/is_dict_int32_t alone: plain fixed-width (non-dict, non-string, + // non-list) leaves under an optional non-list ancestor need the same zero-fill. + if (process_nulls) { + uint32_t const dtype_len = [&]() -> uint32_t { + if constexpr (is_dict_int32_t) { return sizeof(int32_t); } + if constexpr (has_strings_t) { return sizeof(cudf::size_type); } + return s->output_cvt.dtype_len; + }(); + zero_fill_null_positions(s, dtype_len, init_valid_map_offset, t); } if constexpr (has_strings_t) { diff --git a/cpp/src/io/parquet/page_data.cu b/cpp/src/io/parquet/page_data.cu index 54f43b9585d3..5c1a8dd4a3f8 100644 --- a/cpp/src/io/parquet/page_data.cu +++ b/cpp/src/io/parquet/page_data.cu @@ -103,8 +103,7 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size) auto* const rep = reinterpret_cast(pp->lvl_decode_buf[level_type::REPETITION]); // Capture initial valid_map_offset before any processing that might modify it - int const init_valid_map_offset = - s->nesting.nesting_info[s->setup.col.max_nesting_depth - 1].valid_map_offset; + int const init_valid_map_offset = init_null_fill_valid_map_offset(s); // skipped_leaf_values will always be 0 for flat hierarchies. uint32_t skipped_leaf_values = s->setup.page.skipped_leaf_values; @@ -223,16 +222,8 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size) // Zero-fill null positions after decoding valid values if (has_repetition) { - int const leaf_level_index = s->setup.col.max_nesting_depth - 1; - auto const& ni = s->nesting.nesting_info[leaf_level_index]; - 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( + s, s->output_cvt.dtype_len, init_valid_map_offset, static_cast(block.thread_rank())); } if (s->setup.error != 0) { @@ -299,8 +290,7 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size) PageNestingDecodeInfo* nesting_info_base = s->nesting.nesting_info; // Capture initial valid_map_offset before any processing that might modify it - int const init_valid_map_offset = - s->nesting.nesting_info[s->setup.col.max_nesting_depth - 1].valid_map_offset; + int const init_valid_map_offset = init_null_fill_valid_map_offset(s); if (s->stream.dict_base) { out_warp_id = (s->stream.dict_bits > 0) ? 2 : 1; @@ -475,15 +465,8 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size) auto const is_string = ((dtype == Type::BYTE_ARRAY) && !is_decimal) || (dtype == Type::FIXED_LEN_BYTE_ARRAY); if (is_string || has_repetition) { - 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( + s, s->output_cvt.dtype_len, init_valid_map_offset, static_cast(block.thread_rank())); } if (s->setup.error != 0) { diff --git a/cpp/src/io/parquet/page_decode.cuh b/cpp/src/io/parquet/page_decode.cuh index 3c24912095aa..853f8a5b74f0 100644 --- a/cpp/src/io/parquet/page_decode.cuh +++ b/cpp/src/io/parquet/page_decode.cuh @@ -872,6 +872,12 @@ __device__ void gpuUpdateValidityOffsetsAndRowIndices(int32_t target_input_value nesting_info->valid_map_offset += warp_valid_mask_bit_count; nesting_info->null_count += warp_valid_mask_bit_count - __popc(warp_output_valid_mask); } + // this column does not own (write) the validity bitmap at this level, but still needs to + // track how far its read-only view has conceptually advanced, in lockstep with the + // owning column's valid_map_offset above, so it can later resolve null positions there. + if (nesting_info->null_fill_valid_map != nullptr && warp_valid_mask_bit_count > 0) { + nesting_info->null_fill_valid_map_offset += warp_valid_mask_bit_count; + } nesting_info->valid_count += warp_valid_count; nesting_info->value_count += warp_value_count; } @@ -1312,6 +1318,11 @@ inline __device__ bool setup_local_page_info(auto* const s, nesting_info->valid_map += output_offset >> 5; nesting_info->valid_map_offset = (int32_t)(output_offset & 0x1f); } + nesting_info->null_fill_valid_map = s->setup.col.null_fill_valid_map_base[idx]; + if (nesting_info->null_fill_valid_map != nullptr) { + nesting_info->null_fill_valid_map += output_offset >> 5; + nesting_info->null_fill_valid_map_offset = (int32_t)(output_offset & 0x1f); + } } } } @@ -1455,6 +1466,69 @@ inline __device__ bool setup_local_page_info(auto* const s, return true; } +/** + * @brief A validity bitmap and its current bit offset, used to identify null positions when + * zero-filling a leaf's output data. May be a level's own (owned/written) `valid_map`, or, for a + * required leaf whose ancestors do not own the shared bitmap, the read-only `null_fill_valid_map` + * view of it. + */ +struct null_fill_source { + bitmask_type* valid_map; + int32_t valid_map_offset; +}; + +/** + * @brief Determine which nesting level's validity bitmap should be used to identify null + * positions when zero-filling the leaf's output data. + * + * A required Parquet leaf can be absent from its own nesting level's validity buffer because + * one of its ancestors is optional; in that case cudf does not allocate a validity buffer for + * the (non-nullable) leaf, so the nearest optional ancestor's bitmap is used instead. This + * substitution is only correct when there is no repetition (no list ancestor), since only then + * does a bit position in the ancestor's bitmap map 1:1 to a row position in the leaf's output. + * + * An ancestor level shared by multiple children (e.g. two required children of the same nullable + * struct) is only ever written by one "owning" child (see reader_impl.cpp); the others only have + * a read-only `null_fill_valid_map` view of it, which is what is returned here for them. + * + * This resolution only depends on state that is fixed for the page (validity buffer pointers + * and the column's repetition level), not on decode progress, so it is safe to call both before + * and after decoding a batch and get a consistent answer. + * + * @param s Page state containing all necessary information + * @return The validity bitmap and offset that should be used + */ +__device__ null_fill_source get_null_fill_nesting_info(auto* s) +{ + int const leaf_level_index = s->setup.col.max_nesting_depth - 1; + auto const& leaf_ni = s->nesting.nesting_info[leaf_level_index]; + if (leaf_ni.valid_map != nullptr) { return {leaf_ni.valid_map, leaf_ni.valid_map_offset}; } + if (s->setup.col.max_level[level_type::REPETITION] != 0) { return {nullptr, 0}; } + for (int idx = leaf_level_index - 1; idx >= 0; --idx) { + auto const& ancestor_ni = s->nesting.nesting_info[idx]; + if (ancestor_ni.valid_map != nullptr) { + return {ancestor_ni.valid_map, ancestor_ni.valid_map_offset}; + } + if (ancestor_ni.null_fill_valid_map != nullptr) { + return {ancestor_ni.null_fill_valid_map, ancestor_ni.null_fill_valid_map_offset}; + } + } + return {nullptr, 0}; +} + +/** + * @brief Capture the validity-map bit offset of the nesting level that will later be used to + * zero-fill null positions in the leaf output (see get_null_fill_nesting_info()), before that + * level's offset advances further while decoding the current batch. + * + * @param s Page state containing all necessary information + * @return The bit offset to later pass to zero_fill_null_positions() as init_valid_map_offset + */ +__device__ int init_null_fill_valid_map_offset(auto* s) +{ + return get_null_fill_nesting_info(s).valid_map_offset; +} + /** * @brief Zero-fill null positions in output data using parallel per-validity-block processing * @@ -1479,12 +1553,10 @@ __device__ void zero_fill_null_positions_shared( // nesting level that is storing actual leaf values 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 - if ((ni.valid_map == nullptr) || (num_values == 0)) { return; } + auto const& leaf_ni = s->nesting.nesting_info[leaf_level_index]; + auto const ni = get_null_fill_nesting_info(s); - auto const data_out = ni.data_out; + auto const data_out = leaf_ni.data_out; constexpr int bits_per_mask = cudf::detail::size_in_bits(); using cudf::detail::warp_size; @@ -1564,4 +1636,33 @@ __device__ void zero_fill_null_positions_shared( __syncthreads(); } +/** + * @brief Zero-fill null positions in the leaf's output data for values decoded since + * @p init_valid_map_offset was captured by init_null_fill_valid_map_offset(). + * + * This resolves the correct validity bitmap via get_null_fill_nesting_info() (the leaf's own, + * or the nearest optional ancestor's if the leaf itself is required and therefore has no + * validity buffer of its own) and derives the number of newly-decoded values from how far that + * bitmap's offset has advanced since @p init_valid_map_offset was captured, so that callers do + * not need to duplicate this resolution logic themselves. + * + * @tparam block_size CUDA block size for the kernel + * @param s Page state containing all necessary information + * @param dtype_len Size of each data element in bytes + * @param init_valid_map_offset Bit offset captured before decoding the current batch, via + * init_null_fill_valid_map_offset() + * @param t Thread index within the block + */ +template +__device__ void zero_fill_null_positions(auto* s, + uint32_t dtype_len, + int init_valid_map_offset, + int t) +{ + auto const ni = get_null_fill_nesting_info(s); + int const num_values = ni.valid_map_offset - init_valid_map_offset; + if (ni.valid_map == nullptr || num_values == 0) { return; } + zero_fill_null_positions_shared(s, dtype_len, init_valid_map_offset, num_values, t); +} + } // namespace cudf::io::parquet::detail diff --git a/cpp/src/io/parquet/page_delta_decode.cu b/cpp/src/io/parquet/page_delta_decode.cu index 72425abc59de..b94600a47f3d 100644 --- a/cpp/src/io/parquet/page_delta_decode.cu +++ b/cpp/src/io/parquet/page_delta_decode.cu @@ -370,8 +370,7 @@ CUDF_KERNEL void __launch_bounds__(decode_delta_binary_block_size) bool const process_nulls = should_process_nulls(s); // Capture initial valid_map_offset before any processing that might modify it - int const init_valid_map_offset = - s->nesting.nesting_info[s->setup.col.max_nesting_depth - 1].valid_map_offset; + int const init_valid_map_offset = init_null_fill_valid_map_offset(s); // copying logic from gpuDecodePageData. PageNestingDecodeInfo const* nesting_info_base = s->nesting.nesting_info; @@ -480,15 +479,8 @@ CUDF_KERNEL void __launch_bounds__(decode_delta_binary_block_size) if (has_repetition) { // Zero-fill null positions after decoding valid values - 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( + s, s->output_cvt.dtype_len, init_valid_map_offset, static_cast(block.thread_rank())); } if (block.thread_rank() == 0 and s->setup.error != 0) { set_error(s->setup.error, error_code); } @@ -549,8 +541,7 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size) bool const process_nulls = should_process_nulls(s); // Capture initial valid_map_offset before any processing that might modify it - int const init_valid_map_offset = - s->nesting.nesting_info[s->setup.col.max_nesting_depth - 1].valid_map_offset; + int const init_valid_map_offset = init_null_fill_valid_map_offset(s); // choose a character parallel string copy when the average string is longer than a warp auto const use_char_ll = @@ -689,15 +680,8 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size) } // Zero-fill null positions after decoding valid values - auto const& ni = s->nesting.nesting_info[leaf_level_index]; - if (ni.valid_map != nullptr) { - int const num_values = ni.valid_map_offset - init_valid_map_offset; - zero_fill_null_positions_shared(s, - sizeof(size_type), - init_valid_map_offset, - num_values, - static_cast(block.thread_rank())); - } + zero_fill_null_positions( + s, sizeof(size_type), init_valid_map_offset, static_cast(block.thread_rank())); // For large strings, update the initial string buffer offset to be used during large string // column construction. Otherwise, convert string sizes to final offsets. @@ -772,8 +756,7 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size) bool const process_nulls = should_process_nulls(s); // Capture initial valid_map_offset before any processing that might modify it - int const init_valid_map_offset = - s->nesting.nesting_info[s->setup.col.max_nesting_depth - 1].valid_map_offset; + int const init_valid_map_offset = init_null_fill_valid_map_offset(s); // copying logic from gpuDecodePageData. PageNestingDecodeInfo const* nesting_info_base = s->nesting.nesting_info; @@ -906,15 +889,8 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size) } // Zero-fill null positions after decoding valid values - auto const& ni = nesting_info_base[leaf_level_index]; - if (ni.valid_map != nullptr) { - int const num_values = ni.valid_map_offset - init_valid_map_offset; - zero_fill_null_positions_shared(s, - sizeof(size_type), - init_valid_map_offset, - num_values, - static_cast(block.thread_rank())); - } + zero_fill_null_positions( + s, sizeof(size_type), init_valid_map_offset, static_cast(block.thread_rank())); // For large strings, update the initial string buffer offset to be used during large string // column construction. Otherwise, convert string sizes to final offsets. diff --git a/cpp/src/io/parquet/parquet_gpu.hpp b/cpp/src/io/parquet/parquet_gpu.hpp index f0399fcbc646..baae69f57d8d 100644 --- a/cpp/src/io/parquet/parquet_gpu.hpp +++ b/cpp/src/io/parquet/parquet_gpu.hpp @@ -303,6 +303,14 @@ struct PageNestingDecodeInfo { uint8_t* data_out; uint8_t* string_out; bitmask_type* valid_map; + + // read-only view of this level's validity bitmap, populated even when this column does not + // own (write) the bitmap at this nesting level (i.e. a sibling column under the same nullable + // struct owns and writes it). Used to resolve null positions for zero-filling a required leaf + // whose own nesting level has no validity buffer of its own; never written to, so it carries no + // risk of the concurrent-writer race that `valid_map` ownership avoids. + int32_t null_fill_valid_map_offset; + bitmask_type* null_fill_valid_map; }; // Use up to 512 bytes of shared memory as a cache for nesting information. @@ -493,12 +501,17 @@ struct ColumnChunkDesc { int32_t num_data_pages{}; // number of data pages int32_t num_dict_pages{}; // number of dictionary pages PageInfo const* dict_page{}; - string_index_pair* str_dict_index{}; // index for string dictionary - bitmask_type** valid_map_base{}; // base pointers of valid bit map for this column - void** column_data_base{}; // base pointers of column data - void** column_string_base{}; // base pointers of column string data - uint32_t* column_string_offset_base{}; // base pointer of column string offset data - Compression codec{}; // compressed codec enum + string_index_pair* str_dict_index{}; // index for string dictionary + bitmask_type** valid_map_base{}; // base pointers of valid bit map for this column + // base pointers to a read-only view of each nesting level's validity bitmap, populated for + // every column sharing a nesting level (not just the owning column that writes valid_map). + // Used only to resolve null positions when zero-filling a required leaf's output; never written + // to. + bitmask_type** null_fill_valid_map_base{}; + void** column_data_base{}; // base pointers of column data + void** column_string_base{}; // base pointers of column string data + uint32_t* column_string_offset_base{}; // base pointer of column string offset data + Compression codec{}; // compressed codec enum cuda::std::optional logical_type{}; // logical type int32_t ts_clock_rate{}; // output timestamp clock frequency (0=default, 1000=ms, 1000000000=ns) diff --git a/cpp/src/io/parquet/reader_impl.cpp b/cpp/src/io/parquet/reader_impl.cpp index f9f69d270b31..b4b959b6fbcf 100644 --- a/cpp/src/io/parquet/reader_impl.cpp +++ b/cpp/src/io/parquet/reader_impl.cpp @@ -91,6 +91,12 @@ void reader_impl::decode_page_data(read_mode mode, size_t skip_rows, size_t num_ // offset into `chunk_nested_data`/`chunk_nested_valids` for the array of pointers for chunk `i` auto chunk_nested_valids = cudf::detail::hostdevice_vector(sum_max_depths, _stream); + // read-only view of each nesting level's validity bitmap, populated for every column sharing a + // level (unlike chunk_nested_valids, which only the owning column gets). Used to zero-fill null + // positions for a required leaf under a nullable ancestor when this column doesn't own (write) + // that ancestor's validity buffer. + auto chunk_nested_null_fill_valids = + cudf::detail::hostdevice_vector(sum_max_depths, _stream); auto chunk_nested_data = cudf::detail::hostdevice_vector(sum_max_depths, _stream); auto chunk_offsets = std::vector(); auto chunk_nested_str_data = @@ -110,6 +116,9 @@ void reader_impl::decode_page_data(read_mode mode, size_t skip_rows, size_t num_ auto valids = chunk_nested_valids.host_ptr(chunk_off); pass.chunks[c].valid_map_base = chunk_nested_valids.device_ptr(chunk_off); + auto null_fill_valids = chunk_nested_null_fill_valids.host_ptr(chunk_off); + pass.chunks[c].null_fill_valid_map_base = chunk_nested_null_fill_valids.device_ptr(chunk_off); + // get a slice of size `nesting depth` from `chunk_nested_data` to store an array of pointers to // out data auto data = chunk_nested_data.host_ptr(chunk_off); @@ -157,6 +166,11 @@ void reader_impl::decode_page_data(read_mode mode, size_t skip_rows, size_t num_ auto& out_buf = (*cols)[input_col.nesting[idx]]; cols = &out_buf.children; + // every column sharing this nesting level gets a read-only view of its validity bitmap + // (used only to resolve null positions for zero-filling a required leaf's output data), even + // though only the owning column below gets a writable pointer to it. + null_fill_valids[idx] = out_buf.null_mask(); + int const owning_schema = out_buf.user_data & PARQUET_COLUMN_BUFFER_SCHEMA_MASK; if (owning_schema == 0 || owning_schema == input_col.schema_idx) { valids[idx] = out_buf.null_mask(); @@ -184,6 +198,7 @@ void reader_impl::decode_page_data(read_mode mode, size_t skip_rows, size_t num_ pass.chunks.host_to_device_async(_stream); chunk_nested_valids.host_to_device_async(_stream); + chunk_nested_null_fill_valids.host_to_device_async(_stream); chunk_nested_data.host_to_device_async(_stream); if (has_strings) { // Host vector to initialize the initial string offsets diff --git a/cpp/tests/io/parquet_reader_test.cpp b/cpp/tests/io/parquet_reader_test.cpp index 6e55ccbd5ca0..0c91d4184a41 100644 --- a/cpp/tests/io/parquet_reader_test.cpp +++ b/cpp/tests/io/parquet_reader_test.cpp @@ -1759,6 +1759,115 @@ TEST_F(ParquetReaderTest, StructByteArray) CUDF_TEST_EXPECT_TABLES_EQUAL(expected, result.tbl->view()); } +TEST_F(ParquetReaderTest, StructRequiredChildNullGaps) +{ + // A REQUIRED leaf column nested under an OPTIONAL struct ancestor has no validity buffer of + // its own in cudf, so when the struct is null the leaf has no encoded value for that row. + // The reader must zero-fill the resulting output gaps using the struct's own validity bitmap + // rather than leaving them as garbage/stale data. Use enough rows to span multiple 32-bit + // validity blocks and warps in the zero-fill kernel, with nulls at irregular offsets so gaps + // straddle block/warp boundaries. + constexpr auto num_rows = 200; + + auto child_values = cudf::detail::make_counting_transform_iterator( + 0, [](auto i) { return static_cast(i); }); + column_wrapper child_col{ + child_values, child_values + num_rows, cudf::test::iterators::no_nulls()}; + + std::vector struct_validity(num_rows); + for (int i = 0; i < num_rows; ++i) { + struct_validity[i] = (i % 7) != 0; + } + auto struct_col = cudf::test::structs_column_wrapper{{child_col}, struct_validity}; + + auto const written = table_view{{struct_col}}; + cudf::io::table_input_metadata output_metadata(written); + output_metadata.column_metadata[0].child(0).set_nullability(false); + + auto filepath = temp_env->get_temp_filepath("StructRequiredChildNullGaps.parquet"); + cudf::io::parquet_writer_options out_opts = + cudf::io::parquet_writer_options::builder(cudf::io::sink_info{filepath}, written) + .metadata(std::move(output_metadata)); + cudf::io::write_parquet(out_opts); + + cudf::io::parquet_reader_options in_opts = + cudf::io::parquet_reader_options::builder(cudf::io::source_info{filepath}); + auto result = cudf::io::read_parquet(in_opts); + + // the child column has no validity buffer of its own; its raw values at struct-null rows must + // be zero-filled, so comparing raw values (not just the struct's own null mask) exercises the + // fix + ASSERT_FALSE(result.tbl->view().column(0).child(0).nullable()); + + auto expected_child_values = cudf::detail::make_counting_transform_iterator( + 0, [](auto i) { return (i % 7) == 0 ? 0 : static_cast(i); }); + column_wrapper expected_child_col{ + expected_child_values, expected_child_values + num_rows, cudf::test::iterators::no_nulls()}; + auto expected_struct_col = + cudf::test::structs_column_wrapper{{expected_child_col}, struct_validity}; + auto const expected = table_view{{expected_struct_col}}; + + CUDF_TEST_EXPECT_TABLES_EQUAL(expected, result.tbl->view()); +} + +TEST_F(ParquetReaderTest, StructTwoRequiredChildrenNullGaps) +{ + // Only the first child column of a nullable struct owns (writes) the struct's validity bitmap + // during decode; other children only get a read-only view of it. Both required children must + // still have their output gaps zero-filled at struct-null rows, not just the owning one. + constexpr auto num_rows = 200; + + auto make_values = [](int multiplier) { + return cudf::detail::make_counting_transform_iterator( + 0, [multiplier](auto i) { return static_cast(i) * multiplier; }); + }; + + auto values_a = make_values(1); + auto values_b = make_values(3); + column_wrapper child_col_a{ + values_a, values_a + num_rows, cudf::test::iterators::no_nulls()}; + column_wrapper child_col_b{ + values_b, values_b + num_rows, cudf::test::iterators::no_nulls()}; + + std::vector struct_validity(num_rows); + for (int i = 0; i < num_rows; ++i) { + struct_validity[i] = (i % 7) != 0; + } + auto struct_col = cudf::test::structs_column_wrapper{{child_col_a, child_col_b}, struct_validity}; + + auto const written = table_view{{struct_col}}; + cudf::io::table_input_metadata output_metadata(written); + output_metadata.column_metadata[0].child(0).set_nullability(false); + output_metadata.column_metadata[0].child(1).set_nullability(false); + + auto filepath = temp_env->get_temp_filepath("StructTwoRequiredChildrenNullGaps.parquet"); + cudf::io::parquet_writer_options out_opts = + cudf::io::parquet_writer_options::builder(cudf::io::sink_info{filepath}, written) + .metadata(std::move(output_metadata)); + cudf::io::write_parquet(out_opts); + + cudf::io::parquet_reader_options in_opts = + cudf::io::parquet_reader_options::builder(cudf::io::source_info{filepath}); + auto result = cudf::io::read_parquet(in_opts); + + ASSERT_FALSE(result.tbl->view().column(0).child(0).nullable()); + ASSERT_FALSE(result.tbl->view().column(0).child(1).nullable()); + + auto expected_values_a = cudf::detail::make_counting_transform_iterator( + 0, [](auto i) { return (i % 7) == 0 ? 0 : static_cast(i) * 1; }); + auto expected_values_b = cudf::detail::make_counting_transform_iterator( + 0, [](auto i) { return (i % 7) == 0 ? 0 : static_cast(i) * 3; }); + column_wrapper expected_child_col_a{ + expected_values_a, expected_values_a + num_rows, cudf::test::iterators::no_nulls()}; + column_wrapper expected_child_col_b{ + expected_values_b, expected_values_b + num_rows, cudf::test::iterators::no_nulls()}; + auto expected_struct_col = cudf::test::structs_column_wrapper{ + {expected_child_col_a, expected_child_col_b}, struct_validity}; + auto const expected = table_view{{expected_struct_col}}; + + CUDF_TEST_EXPECT_TABLES_EQUAL(expected, result.tbl->view()); +} + TEST_F(ParquetReaderTest, NestingOptimizationTest) { // test nesting levels > cudf::io::parquet::detail::max_cacheable_nesting_decode_info deep.