[BUG] Parquet reader can use an uninitialized required BINARY length under a null ancestor

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<decode_block_size_t,
@@ -1310,16 +1309,7 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size_t, 8)
         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<decode_block_size_t>(
-        s, dtype_len, init_valid_map_offset, num_values, t);
+      zero_fill_null_positions<decode_block_size_t>(s, dtype_len, init_valid_map_offset, t);
     }
   }
 

cpp/src/io/parquet/page_data.cu

@@ -103,8 +103,7 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size)
   auto* const rep    = reinterpret_cast<level_t*>(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<decode_block_size>(s,
-                                                         s->output_cvt.dtype_len,
-                                                         init_valid_map_offset,
-                                                         num_values,
-                                                         static_cast<int>(block.thread_rank()));
-    }
+    zero_fill_null_positions<decode_block_size>(
+      s, s->output_cvt.dtype_len, init_valid_map_offset, static_cast<int>(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<decode_block_size>(s,
-                                                         s->output_cvt.dtype_len,
-                                                         init_valid_map_offset,
-                                                         num_values,
-                                                         static_cast<int>(block.thread_rank()));
-    }
+    zero_fill_null_positions<decode_block_size>(
+      s, s->output_cvt.dtype_len, init_valid_map_offset, static_cast<int>(block.thread_rank()));
   }
 
   if (s->setup.error != 0) {

cpp/src/io/parquet/page_decode.cuh

@@ -1455,6 +1455,49 @@ inline __device__ bool setup_local_page_info(auto* const s,
   return true;
 }
 
+/**
+ * @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.
+ *
+ * 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 nesting level whose validity bitmap should be used
+ */
+__device__ auto& get_null_fill_nesting_info(auto* s)
+{
+  int const leaf_level_index = s->setup.col.max_nesting_depth - 1;
+  auto& leaf_ni              = s->nesting.nesting_info[leaf_level_index];
+  if (leaf_ni.valid_map != nullptr) { return leaf_ni; }
+  if (s->setup.col.max_level[level_type::REPETITION] != 0) { return leaf_ni; }
+  for (int idx = leaf_level_index - 1; idx >= 0; --idx) {
+    auto& ancestor_ni = s->nesting.nesting_info[idx];
+    if (ancestor_ni.valid_map != nullptr) { return ancestor_ni; }
+  }
+  return leaf_ni;
+}
+
+/**
+ * @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 +1522,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];
+  auto const& leaf_ni        = s->nesting.nesting_info[leaf_level_index];
+  auto const& ni             = get_null_fill_nesting_info(s);
 
-  // Check if we have nulls to fill
-  if ((ni.valid_map == nullptr) || (num_values == 0)) { return; }
-
-  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<bitmask_type>();
   using cudf::detail::warp_size;
@@ -1564,4 +1605,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 <int block_size>
+__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<block_size>(s, dtype_len, init_valid_map_offset, num_values, t);
+}
+
 }  // namespace cudf::io::parquet::detail

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<decode_block_size>(s,
-                                                         s->output_cvt.dtype_len,
-                                                         init_valid_map_offset,
-                                                         num_values,
-                                                         static_cast<int>(block.thread_rank()));
-    }
+    zero_fill_null_positions<decode_block_size>(
+      s, s->output_cvt.dtype_len, init_valid_map_offset, static_cast<int>(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<decode_block_size>(s,
-                                                       sizeof(size_type),
-                                                       init_valid_map_offset,
-                                                       num_values,
-                                                       static_cast<int>(block.thread_rank()));
-  }
+  zero_fill_null_positions<decode_block_size>(
+    s, sizeof(size_type), init_valid_map_offset, static_cast<int>(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<decode_block_size>(s,
-                                                       sizeof(size_type),
-                                                       init_valid_map_offset,
-                                                       num_values,
-                                                       static_cast<int>(block.thread_rank()));
-  }
+  zero_fill_null_positions<decode_block_size>(
+    s, sizeof(size_type), init_valid_map_offset, static_cast<int>(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.

cpp/tests/io/parquet_reader_test.cpp

@@ -1759,6 +1759,57 @@ 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<int32_t>(i); });
+  column_wrapper<int32_t> child_col{
+    child_values, child_values + num_rows, cudf::test::iterators::no_nulls()};
+
+  std::vector<bool> 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<int32_t>(i); });
+  column_wrapper<int32_t> 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, NestingOptimizationTest)
 {
   // test nesting levels > cudf::io::parquet::detail::max_cacheable_nesting_decode_info deep.