Zero-fill non-nullable string offsets with a nullable ancestor

cpp/src/io/parquet/page_decode.cuh

@@ -1458,10 +1458,15 @@ 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 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
+ * may not be available to this leaf.
+ *
+ * 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
@@ -1481,7 +1486,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;

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

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<cudf::device_span<cudf::bitmask_type>> 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<cudf::device_span<cuda::std::byte>>(
+    _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.push_back({static_cast<cuda::std::byte*>(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<size_type>(_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<size_type>(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.push_back({static_cast<cuda::std::byte*>(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<cudf::device_span<cudf::bitmask_type> const>{nullmask_bufs}, _stream);
   cudf::detail::batched_memset<cudf::bitmask_type>(
     pinned_nullmask_bufs, std::numeric_limits<cudf::bitmask_type>::max(), _stream);
+
+  // Need to zero non-nullable string lengths with nullable ancestors
+  if (not unwritten_bufs.empty()) {
+    cudf::detail::batched_memset<cuda::std::byte>(
+      cudf::host_span<cudf::device_span<cuda::std::byte> const>{unwritten_bufs},
+      static_cast<cuda::std::byte>(0),
+      _stream);
+  }
 }
 
 void reader_impl::fill_pruned_offsets(size_t skip_rows,

cpp/tests/io/parquet_reader_test.cpp

@@ -17,6 +17,7 @@
 #include <cudf_test/table_utilities.hpp>
 
 #include <cudf/column/column.hpp>
+#include <cudf/column/column_factories.hpp>
 #include <cudf/copying.hpp>
 #include <cudf/io/parquet.hpp>
 #include <cudf/io/parquet_metadata.hpp>
@@ -39,6 +40,7 @@
 #include <memory>
 #include <optional>
 #include <stdexcept>
+#include <string_view>
 #include <utility>
 
 using ParquetDecompressionTest = DecompressionTest<ParquetReaderTest>;
@@ -6337,3 +6339,149 @@ 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<cudf::column> make_optional_struct(
+  std::vector<std::unique_ptr<cudf::column>>&& 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<std::unique_ptr<cudf::column>> 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<std::unique_ptr<cudf::column>> 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<std::unique_ptr<cudf::column>> 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<std::unique_ptr<cudf::column>> 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<cudf::size_type>(i * list_size); });
+  column_wrapper<cudf::size_type> 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("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());
+
+  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());
+}