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

cpp/src/io/parquet/page_decode.cuh

@@ -1479,12 +1479,27 @@ __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];
+
+  // A required Parquet leaf can be absent because one of its ancestors is optional. Since with RMM,
+  // the reader reader can leave the validity map associated w/the ancestor unwritten, this code
+  // zero-fills the gap rows by borrowing the nearest ancestor's validity bitmap instead.
+  auto const& ni = [&]() -> PageNestingDecodeInfo const& {
+    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 const& ancestor_ni = s->nesting.nesting_info[idx];
+      if (ancestor_ni.valid_map != nullptr) { return ancestor_ni; }
+    }
+    return leaf_ni;
+  }();
 
   // Check if we have nulls to fill
   if ((ni.valid_map == nullptr) || (num_values == 0)) { return; }
 
-  auto const data_out = ni.data_out;
+  if (&ni != &leaf_ni) { valid_map_offset = ni.valid_map_offset; }
+
+  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;