[BUG] Parquet writer emits a value for every row of a nullable struct's child, shifting the data

[mhaseeb123]

Describe the bug

When writing a child of a nullable struct, the cuDF Parquet writer stores a value for every row, including rows where the parent struct is null.

Parquet requires a value to be stored only where definition_level == max_definition_level. The definition levels cuDF writes are correct; only the value stream is wrong. Because readers consume stored values positionally against the definition levels, they take the first num_non_null values, so every value after the first null parent is shifted and the trailing
values are lost.

This is silent data corruption on a plain write→read round trip. It affects any reader (reproduced below with both cuDF and pyarrow), and it is not specific to required children — see "Also affects nullable children" below.

message schema {
  required group schema {
    optional group s {        # nullable
      required int32 a;       # non-nullable child
    }
  }
}

Steps/Code to reproduce bug

16 rows, parent struct null at rows 3 and 11, child values 100..115.

Self-contained C++ reproducer

#include <cudf/column/column_factories.hpp>
#include <cudf/io/parquet.hpp>
#include <cudf/null_mask.hpp>
#include <cudf/table/table_view.hpp>
#include <cudf/utilities/default_stream.hpp>

#include <rmm/device_buffer.hpp>

#include <cuda_runtime_api.h>

#include <cstdint>
#include <memory>
#include <vector>

namespace {
constexpr cudf::size_type num_rows = 16;
bool parent_is_null(cudf::size_type i) { return i == 3 || i == 11; }
}  // namespace

int main()
{
  // child: 100, 101, ... 115 (distinct, so a shift is obvious); no null mask of its own
  std::vector<int32_t> host(num_rows);
  for (cudf::size_type i = 0; i < num_rows; ++i) { host[i] = 100 + i; }
  auto child = cudf::make_fixed_width_column(
    cudf::data_type{cudf::type_id::INT32}, num_rows, cudf::mask_state::UNALLOCATED);
  cudaMemcpy(child->mutable_view().data<int32_t>(),
             host.data(), host.size() * sizeof(int32_t), cudaMemcpyHostToDevice);

  // parent struct: null at rows 3 and 11
  auto mask = cudf::create_null_mask(num_rows, cudf::mask_state::ALL_VALID);
  cudf::size_type null_count = 0;
  for (cudf::size_type i = 0; i < num_rows; ++i) {
    if (parent_is_null(i)) {
      cudf::set_null_mask(static_cast<cudf::bitmask_type*>(mask.data()), i, i + 1, false);
      ++null_count;
    }
  }
  cudf::get_default_stream().synchronize();

  std::vector<std::unique_ptr<cudf::column>> children;
  children.push_back(std::move(child));
  auto parent = cudf::create_structs_hierarchy(
    num_rows, std::move(children), null_count, std::move(mask));

  auto const input = cudf::table_view{{parent->view()}};
  cudf::io::table_input_metadata meta(input);
  meta.column_metadata[0].set_name("s");
  meta.column_metadata[0].child(0).set_name("a").set_nullability(false);  // required child

  cudf::io::write_parquet(
    cudf::io::parquet_writer_options::builder(cudf::io::sink_info{"out.parquet"}, input)
      .metadata(std::move(meta))
      .dictionary_policy(cudf::io::dictionary_policy::NEVER)
      .compression(cudf::io::compression_type::NONE)
      .build());
  return 0;
}

Reading out.parquet back (pyarrow shown; cuDF gives the same answer):

import pyarrow.parquet as pq
f = pq.ParquetFile("out.parquet")
print(f.schema)
c = f.metadata.row_group(0).column(0)
print("num_values =", c.num_values, " uncompressed =", c.total_uncompressed_size)
print([None if x is None else x["a"] for x in f.read().column(0).to_pylist()])

required group field_id=-1 schema {
  optional group field_id=-1 s {
    required int32 field_id=-1 a;
  }
}
num_values = 16  uncompressed = 90
[100, 101, 102, None, 103, 104, 105, 106, 107, 108, 109, None, 110, 111, 112, 113]

Two things to note:

1. The data is shifted. Row 4 should be 104; it reads back 103, which is the value that sat at the null row 3. Rows 4-10 and 12-15 are all wrong, and the last two input values (114, 115) are lost entirely.
2. The page holds too many values. total_uncompressed_size = 90 bytes == 16 * 4 + 26 (levels + header). With the required 14 values it would be 14 * 4 = 56 plus overhead. So 16 values were encoded where 14 are permitted.

Same behaviour with a required binary child, at both PLAIN and DELTA_LENGTH_BYTE_ARRAY, and it scales:

child type	rows	null structs	encoding	total_uncompressed_size	num_rows * (len+4)	num_non_null * (len+4)
int32	16	2	PLAIN	90	64 (+26) ✔	56
binary (36 B)	16	3	PLAIN	666	640 (+26) ✔	520
binary (36 B)	2000	286	PLAIN	80280	80000 (+280) ✔	68560
binary (36 B)	2000	286	DELTA_LENGTH_BYTE_ARRAY	72369	—	—

For the 2000-row cases, reading the file back with cuDF and inspecting the child's chars_size:

encoding	cuDF chars_size	correct value
PLAIN	61704	61704 ✔
DELTA_LENGTH_BYTE_ARRAY	72000 (== 2000 * 36)	61704 ✘

Also affects nullable children

The same corruption occurs for optional group s { optional int32 a; } whenever the child's own null mask does not already have the parent's nulls pushed into it. Re-running the reproducer with the child given an all-valid mask (instead of set_nullability(false)):

required group field_id=-1 schema {
  optional group field_id=-1 s {
    optional int32 field_id=-1 a;      # note: OPTIONAL this time
  }
}
num_values = 16  uncompressed = 92
[100, 101, 102, None, 103, 104, 105, 106, 107, 108, 109, None, 110, 111, 112, 113]

Identical shift. This matters because cudf::create_structs_hierarchy explicitly documents that the struct's null mask "is orthogonal to the null values of individual child columns", so a struct built through the public API and then written out is corrupted. Columns that happen to come from an op that sanitizes children round-trip fine, which is probably why this went
unnoticed.

Expected behavior

Only the values at rows where the whole path is defined should be encoded — 14 values in the example above, 100, 101, 102, 104, ..., 110, 112, ..., 115 — so that a write→read round trip returns the original data.

Root cause

Value emission and definition levels are computed from two different sources and are never reconciled.

Values are gated on the leaf column's own validity —
gpuEncodePages, cpp/src/io/parquet/page_enc.cu:1668-1686:

size_type const val_idx_in_leaf_col = s->page_start_val + val_idx_in_block;

is_valid = (val_idx_in_leaf_col < s->col.leaf_column->size() &&
            val_idx_in_block < s->page.num_leaf_values)
             ? s->col.leaf_column->is_valid(val_idx_in_leaf_col)
             : 0;

with the emit gate at :1694. There is no reference to def_level / max_def_level in that loop.

Definition levels walk the whole ancestor path —
gpuEncodePageLevels, cpp/src/io/parquet/page_enc.cu:1382-1407:

auto col = *s->col.parent_column;          // starts at the root, not the leaf
do {
  if (s->col.nullability[l]) {
    if (col.is_valid(row)) { ++def; } else { break; }   // shallowest null wins
  }
  ...
} while (is_col_struct);

Nothing superimposes ancestor nulls onto the leaf before encoding:
create_leaf_column_device_views (cpp/src/io/utilities/column_utils.cuh:56-73) walks down to the leaf and stores it verbatim, and parquet_column_view (cpp/src/io/parquet/writer_impl.cu:1008-1136) preserves each level's own mask without OR-ing parents in. grep -rn "superimpose\|push_down_nulls" cpp/src/io/ has no hits in the Parquet writer.

The same leaf_column->is_valid(...) gate is repeated in the fragment sizing and in every other encoder, which is why total_uncompressed_size is budgeted for all rows too: page_enc.cu:167-170 (calculate_frag_size), :340-353 (byte-stream-split), :1919-1921 (gpuEncodeDictPages), :2161-2169 (DELTA_LENGTH_BYTE_ARRAY),
:2306-2320 (DELTA_BYTE_ARRAY), and cpp/src/io/parquet/chunk_dict.cu:136-141 — so values under null parents also get inserted into the dictionary.

The invariant the writer needs is emit iff def_level == max_def_level; it currently emits iff leaf_column->is_valid(i). Those agree only when the leaf's mask already has all ancestor nulls pushed down.

Additional context

• The DELTA_LENGTH_BYTE_ARRAY row in the second table above is a knock-on effect on the read side: compute_delta_length_page_string_sizes (cpp/src/io/parquet/page_string_decode.cu:692) derives a page's string size as page_data_size - delta_block_size. Because the extra unreachable payload bytes sit inside the page, cuDF computes a size that is too large and the column's terminating offset comes out wrong when reading back its own file (72000 vs 61704). That shortcut is correct for a conformant file, so fixing the writer should fix this too, but it is worth a regression test.
• Discovered while working on [BUG] Parquet reader can use an uninitialized required BINARY length under a null ancestor #23655. The two are independent.

Environment overview

• Environment location: Docker (rapidsai devcontainer, rapids-cudf-conda-26.10-cuda13.3-conda)
• Method of cuDF install: from source

Environment details

• cuDF 26.10.00, main @ f588e0c838b
• CUDA 13.3 (V13.3.73), driver 580.173.02
• NVIDIA RTX 5880 Ada Generation
• Linux-7.0.0-28-generic-x86_64-with-glibc2.39