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

[nartal1]

Describe the bug

A required Parquet leaf can be absent because one of its ancestors is optional. For the minimal physical shape below, payload has no leaf validity map of its own, but there is no payload value when s is null:

optional group s {
  required binary payload;
}

It looks like that the reader can leave the variable-width length slot unwritten for that inherited-null row. With an RMM pool, the slot can retain data from a prior allocation. Scanning the lengths can then produce invalid offsets and an invalid allocation size. The issue is not specific to variant type though.

Steps/Code to reproduce bug

Copy the program below outside the source tree as n1_self_contained_cpp_repro.cpp. It creates 100 rows with one null parent at row 21 and a 32-byte payload in every other row, then writes one uncompressed, non-dictionary row group with V1 data pages.

The writer metadata explicitly creates this physical schema:

required group schema {
  optional group s {
    required binary payload;
  }
}

Self-contained C++ reproducer

#include <cudf/column/column_factories.hpp>
#include <cudf/io/parquet.hpp>
#include <cudf/null_mask.hpp>
#include <cudf/strings/strings_column_view.hpp>
#include <cudf/utilities/default_stream.hpp>
#include <cudf/utilities/memory_resource.hpp>

#include <rmm/device_buffer.hpp>
#include <rmm/mr/cuda_memory_resource.hpp>
#include <rmm/mr/pool_memory_resource.hpp>

#include <cuda_runtime_api.h>

#include <cstdint>
#include <cstdlib>
#include <iostream>
#include <memory>
#include <stdexcept>
#include <string>
#include <tuple>
#include <utility>
#include <vector>

namespace {

constexpr cudf::size_type rows          = 100;
constexpr cudf::size_type null_row      = 21;
constexpr cudf::size_type bytes_per_row = 32;
constexpr int64_t expected_chars        = (rows - 1) * bytes_per_row;

void check_cuda(cudaError_t status)
{
  if (status != cudaSuccess) { throw std::runtime_error{cudaGetErrorString(status)}; }
}

void make_parquet(std::string const& path)
{
  std::vector<cudf::size_type> host_offsets(rows + 1);
  for (cudf::size_type row = 0; row < rows; ++row) {
    host_offsets[row] = row * bytes_per_row;
  }
  host_offsets[rows] = rows * bytes_per_row;

  auto offsets = cudf::make_fixed_width_column(
    cudf::data_type{cudf::type_id::INT32}, rows + 1, cudf::mask_state::UNALLOCATED);
  auto bytes = cudf::make_fixed_width_column(
    cudf::data_type{cudf::type_id::UINT8}, rows * bytes_per_row, cudf::mask_state::UNALLOCATED);
  check_cuda(cudaMemcpy(offsets->mutable_view().data<cudf::size_type>(),
                        host_offsets.data(),
                        host_offsets.size() * sizeof(cudf::size_type),
                        cudaMemcpyHostToDevice));
  check_cuda(cudaMemset(bytes->mutable_view().data<uint8_t>(), 1, rows * bytes_per_row));

  auto payload = cudf::make_lists_column(
    rows, std::move(offsets), std::move(bytes), 0, rmm::device_buffer{});
  auto parent_mask = cudf::create_null_mask(rows, cudf::mask_state::ALL_VALID);
  cudf::set_null_mask(static_cast<cudf::bitmask_type*>(parent_mask.data()),
                      null_row,
                      null_row + 1,
                      false);
  cudf::get_default_stream().synchronize();

  std::vector<std::unique_ptr<cudf::column>> children;
  children.push_back(std::move(payload));
  auto parent = cudf::create_structs_hierarchy(
    rows, std::move(children), 1, std::move(parent_mask));
  auto const input = cudf::table_view{{parent->view()}};

  cudf::io::table_input_metadata metadata(input);
  metadata.column_metadata[0].set_name("s").set_nullability(true);
  metadata.column_metadata[0]
    .child(0)
    .set_name("payload")
    .set_nullability(false)
    .set_output_as_binary(true)
    .set_encoding(cudf::io::column_encoding::PLAIN);

  auto options = cudf::io::parquet_writer_options::builder(cudf::io::sink_info{path}, input)
                   .metadata(std::move(metadata))
                   .dictionary_policy(cudf::io::dictionary_policy::NEVER)
                   .compression(cudf::io::compression_type::NONE)
                   .write_v2_headers(false)
                   .row_group_size_rows(rows)
                   .max_page_size_rows(rows)
                   .max_page_fragment_size(rows)
                   .build();
  cudf::io::write_parquet(options);
}

void validate(cudf::table_view table, int read)
{
  if (table.num_columns() != 1 || table.num_rows() != rows) {
    throw std::runtime_error{"read " + std::to_string(read) + ": unexpected table shape"};
  }
  auto const root = table.column(0);
  auto const payload = cudf::strings_column_view{root.child(0)};
  auto const chars_size = payload.chars_size(cudf::get_default_stream());
  if (chars_size != expected_chars) {
    throw std::runtime_error{"read " + std::to_string(read) + ": payload chars expected " +
                             std::to_string(expected_chars) + " but found " +
                             std::to_string(chars_size)};
  }
}

}  // namespace

int main(int argc, char** argv)
{
  auto const path  = argc > 1 ? std::string{argv[1]} : "/tmp/n1-required-binary.parquet";
  auto const reads = argc > 2 ? std::stoi(argv[2]) : 3;

  try {
    rmm::mr::pool_memory_resource pool{rmm::mr::cuda_memory_resource{}, 1ULL << 30};
    auto previous = cudf::set_current_device_resource(pool);
    try {
      // Generate the exact OPTIONAL STRUCT -> REQUIRED BINARY shape without PyArrow.
      make_parquet(path);

      auto const options = cudf::io::parquet_reader_options::builder(
                             cudf::io::source_info{path})
                             .build();
      for (int read = 0; read < reads; ++read) {
        auto result = cudf::io::read_parquet(options);
        validate(result.tbl->view(), read);
        auto const payload = cudf::strings_column_view{result.tbl->view().column(0).child(0)};
        std::cout << "PASS read=" << read << " rows=" << result.tbl->num_rows()
                  << " chars=" << payload.chars_size(cudf::get_default_stream()) << '\n';
      }
    } catch (...) {
      std::ignore = cudf::set_current_device_resource(std::move(previous));
      throw;
    }
    std::ignore = cudf::set_current_device_resource(std::move(previous));
  } catch (std::exception const& e) {
    std::cerr << "FAIL " << e.what() << '\n';
    return EXIT_FAILURE;
  }
  return EXIT_SUCCESS;
}

Use this minimal CMakeLists.txt in a sibling directory named n1-self-contained-cpp:

cmake_minimum_required(VERSION 3.30.4)
project(n1_self_contained_cpp_repro LANGUAGES CXX)

find_package(cudf REQUIRED)
get_filename_component(CUDF_INSTALL_PREFIX "${cudf_DIR}/../../.." ABSOLUTE)

add_executable(n1_self_contained_cpp_repro ../n1_self_contained_cpp_repro.cpp)
target_compile_features(n1_self_contained_cpp_repro PRIVATE cxx_std_20)
target_link_directories(n1_self_contained_cpp_repro PRIVATE "${CUDF_INSTALL_PREFIX}/lib")
target_link_libraries(n1_self_contained_cpp_repro PRIVATE cudf::cudf)

Build and run it against a cuDF main-branch build/install:

cmake -S n1-self-contained-cpp -B n1-build \
  -DCMAKE_PREFIX_PATH=<cudf-install> \
  -Dcudf_DIR=<cudf-install>/lib/cmake/cudf
cmake --build n1-build --parallel
./n1-build/n1_self_contained_cpp_repro

The program uses one 1 GiB RMM pool for file generation and all reads. The failing assertion checks the decoded nested STRING character count; no artificial allocator poisoning is used.

Observed result:

The first read is correct. The second read deterministically truncates 672 bytes (21 × 32) from the payload:

PASS read=0 rows=100 chars=3168
FAIL read 1: payload chars expected 3168 but found 2496

This result occurred in five of five fresh native processes on the unfixed build. With the proposed initialization change, the exact same self-generating C++ program passed 20 consecutive reads with the expected 3,168 characters.

The 100-row shape is sufficient and stable; it is not claimed to be the theoretical minimum. A larger real-data run also propagated a corrupt negative child size into an approximately 16 EiB RMM allocation request. That exact size depends on prior device-memory contents and is not required to reproduce the correctness bug. The corresponding native stack reached:

strings::detail::gather
cudf::detail::gather
purge_nonempty_nulls
structs::superimpose_and_sanitize_nulls
Parquet reader

Expected behavior

An inherited-null row must contribute zero length. Final offsets and null sanitization must be valid and independent of allocator reuse or prior device-memory contents.

Additional context

The explicit writer nullability settings in the reproducer are significant: the child must remain REQUIRED. Using structs_column_wrapper instead of create_structs_hierarchy superimposes the parent mask on the child and changes the schema, so it does not exercise this bug.