[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,
@@ -1302,25 +1301,17 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size_t, 8)
     valid_count = next_valid_count;
   }
 
-  // Zero-fill null positions after decoding valid values
-  if constexpr (has_strings_t || has_lists_t || is_dict_int32_t) {
-    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;
-      }();
-      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 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<decode_block_size_t>(s, dtype_len, init_valid_map_offset, t);
   }
 
   if constexpr (has_strings_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

@@ -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<bitmask_type>();
   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 <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/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<LogicalType> logical_type{};  // logical type
   int32_t ts_clock_rate{};  // output timestamp clock frequency (0=default, 1000=ms, 1000000000=ns)