refactor(parquet): introduce level_state_scan struct for preprocess_levels_kernel

[vyasr]

Description

This PR builds on #23471 to create a minimal struct for the preprocess_levels kernel. The primary addition in this PR is page_decode_stream_state struct encapsulating the level stream decoding components of page_state_s, which now has one of those structs as a member. preprocess_levels now uses a minimal level_scan_state struct, which only contains the two sub-structs required by that kernel. The net result is a savings of 560 bytes of shared memory in the kernel. Interestingly, on a couple of architectures (sm75 and sm86) we also saw a drop register usage by 8 (80->72 on sm75, 54->46 on sm86). That is hopefully indicative of the ancillary benefits of reducing shared memory usage as we will be able to incrementally reduce the complexity of the kernels and make it easier for the compiler to produce better optimized code (in this case presumably with better register reuse by eliding intermediate shared memory accesses that made live range counting harder).

Checklist

• I am familiar with the Contributing Guidelines.
• New or existing tests cover these changes.
• The documentation is up to date with these changes.

[coderabbitai]

📝 Walkthrough

Summary by CodeRabbit

• Bug Fixes
  • Improved Parquet decoding reliability across dictionary, boolean, string, timestamp, nested-level, and byte-stream-split data.
  • Corrected handling of page data boundaries, dictionary metadata, and level information during decoding.
  • Improved processing for DELTA-encoded and fixed-width values, including string offsets and null handling.

Walkthrough

Changes

The Parquet decoder now composes page setup and stream state, updates level preprocessing to use the composed state, and migrates fixed-width, dictionary, DELTA, and string decoding paths to stream-backed pointers and metadata.

Parquet stream state refactor

Layer / File(s)	Summary
Composed page state and setup wiring cpp/src/io/parquet/page_decode.cuh, cpp/src/io/parquet/page_state_composed.cuh	Adds page_decode_stream_state, embeds it in page_state_s, and updates page setup, level initialization, dictionary state, nesting metadata, and decoding helpers.
Composed level-scan state cpp/src/io/parquet/decode_preprocess.cu	Uses level_scan_state and stream-backed repetition and definition level ranges in preprocessing.
Fixed-width and page value decoding cpp/src/io/parquet/decode_fixed.cu, cpp/src/io/parquet/page_data.cu, cpp/src/io/parquet/page_data.cuh	Switches fixed-width, boolean, split-page, timestamp, byte-array, and dictionary value reads to s->stream fields.
Delta and string stream consumers cpp/src/io/parquet/page_delta_decode.cu, cpp/src/io/parquet/page_string_decode.cu, cpp/src/io/parquet/page_string_utils.cuh	Updates DELTA decoding, string sizing, offset preprocessing, and string materialization to use stream-backed data and dictionary metadata.

Estimated code review effort: 4 (Complex) | ~45 minutes

Possibly related PRs

• rapidsai/cudf#22532: Overlaps with dictionary-index initialization and dispatch in the page decoding path.
• rapidsai/cudf#23090: Modifies level preprocessing and RLE decoder initialization in the same kernel.
• rapidsai/cudf#23471: Refactors overlapping page_state_s fields and page decoder state access.

Suggested labels: cuIO, Performance

Suggested reviewers: vuule, pmattione-nvidia, bdice, mhaseeb123, qbacpey

🚥 Pre-merge checks | ✅ 5

✅ Passed checks (5 passed)

Check name	Status	Explanation
Title check	✅ Passed	The title matches the main refactor and kernel state simplification, though the struct name appears slightly off.
Description check	✅ Passed	The description accurately explains the Parquet preprocess-level refactor and its shared-memory reduction.
Docstring Coverage	✅ Passed	No functions found in the changed files to evaluate docstring coverage. Skipping docstring coverage check.
Linked Issues check	✅ Passed	Check skipped because no linked issues were found for this pull request.
Out of Scope Changes check	✅ Passed	Check skipped because no linked issues were found for this pull request.

✨ Finishing Touches

🧪 Generate unit tests (beta)

• Create PR with unit tests

---

_{Comment @coderabbitai help to get the list of available commands.}

[coderabbitai]

Caution

Some comments are outside the diff and can’t be posted inline due to platform limitations.

⚠️ Outside diff range comments (1)

cpp/src/io/parquet/decode_preprocess.cu (1)

446-460: 🩺 Stability & Availability | 🟠 Major | ⚡ Quick win

Guard decoder init when s->setup.error is set. Malformed pages can set an error while InitLevelSection leaves s->stream.abs_lvl_start[lvl] unset, and setup_local_page_info still reaches both decoders[...].init(...) calls. Skip the init/decode path or return early so the RLE decoder never sees an indeterminate level-stream pointer.

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/src/io/parquet/decode_preprocess.cu` around lines 446 - 460, Guard the
repetition-level decoder setup in the shown preprocessing flow with
s->setup.error before invoking decoders[level_type::REPETITION].init or
decode_next. When an error is set, skip the barrier/init/decode path or return
early, ensuring no decoder receives an unset s->stream.abs_lvl_start pointer;
preserve normal processing when no error exists.

🧹 Nitpick comments (3)

cpp/src/io/parquet/page_state_composed.cuh (2)

29-33: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

Add an explicit constexpr default constructor, mirroring page_state_s.

level_scan_state's members carry NSDMIs (page_decode_setup_state setup{}, arrays in page_decode_stream_state), so the __shared__ level_scan_state state_g; in decode_preprocess.cu is only accepted because of the #pragma nv_diag_suppress static_var_with_dynamic_init above the kernel. page_state_s deliberately declares CUDF_HOST_DEVICE constexpr page_state_s() noexcept {} for this reason; doing the same here keeps the shared declaration valid without relying on the suppression, and avoids the (dead) shared-memory zero-init intent.

♻️ Proposed change

 struct level_scan_state {
+  CUDF_HOST_DEVICE constexpr level_scan_state() noexcept {}
   page_decode_setup_state setup;
   page_decode_stream_state stream;
   CUDF_PARQUET_PAGE_STATE_ERROR_METHODS
 };

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/src/io/parquet/page_state_composed.cuh` around lines 29 - 33, Update the
level_scan_state struct with an explicit CUDF_HOST_DEVICE constexpr noexcept
default constructor, matching page_state_s, while preserving its existing
members and error methods. This should make its default construction explicit
for shared-memory use without relying on the diagnostic suppression.

---

26-33: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low value

Use doxygen comments for the new struct.

level_scan_state is a new public-in-namespace type documented with plain // comments; a /** @brief ... */ block keeps it consistent with the doxygen tooling used for the rest of the C++/CUDA sources.

As per coding guidelines: "Use doxygen as a documentation generator and linter for C++ and CUDA code."

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/src/io/parquet/page_state_composed.cuh` around lines 26 - 33, Replace the
plain `//` documentation above `level_scan_state` with a Doxygen block comment
using the project’s standard `/** `@brief` ... */` style, while preserving the
existing description of the struct’s setup and stream state.

Source: Coding guidelines

cpp/src/io/parquet/page_decode.cuh (1)

128-133: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low value

Consider constraining the auto* parameters.

Switching to auto* makes these helpers usable with both page_state_s and level_scan_state, but it also silently accepts anything with a setup member and moves diagnostics to the point of use. A small concept (e.g. requires { s->setup.col; s->stream.initial_rle_run; }) or a named template parameter would document the contract these helpers actually rely on.

Also applies to: 147-163, 171-174

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/src/io/parquet/page_decode.cuh` around lines 128 - 133, Constrain the
auto* parameters in is_nullable and the related helpers at lines 147-163 and
171-174 to the intended page_state_s/level_scan_state contract. Add a shared
concept or named template requirement covering the setup.col and
stream.initial_rle_run members those helpers use, so unrelated types are
rejected at definition time.

🤖 Prompt for all review comments with AI agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

Outside diff comments:
In `@cpp/src/io/parquet/decode_preprocess.cu`:
- Around line 446-460: Guard the repetition-level decoder setup in the shown
preprocessing flow with s->setup.error before invoking
decoders[level_type::REPETITION].init or decode_next. When an error is set, skip
the barrier/init/decode path or return early, ensuring no decoder receives an
unset s->stream.abs_lvl_start pointer; preserve normal processing when no error
exists.

---

Nitpick comments:
In `@cpp/src/io/parquet/page_decode.cuh`:
- Around line 128-133: Constrain the auto* parameters in is_nullable and the
related helpers at lines 147-163 and 171-174 to the intended
page_state_s/level_scan_state contract. Add a shared concept or named template
requirement covering the setup.col and stream.initial_rle_run members those
helpers use, so unrelated types are rejected at definition time.

In `@cpp/src/io/parquet/page_state_composed.cuh`:
- Around line 29-33: Update the level_scan_state struct with an explicit
CUDF_HOST_DEVICE constexpr noexcept default constructor, matching page_state_s,
while preserving its existing members and error methods. This should make its
default construction explicit for shared-memory use without relying on the
diagnostic suppression.
- Around line 26-33: Replace the plain `//` documentation above
`level_scan_state` with a Doxygen block comment using the project’s standard
`/** `@brief` ... */` style, while preserving the existing description of the
struct’s setup and stream state.

---

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 3579ebe1-a151-442b-954c-dd74403c1978

📥 Commits

Reviewing files that changed from the base of the PR and between bd36a1e and 5a5e241.

📒 Files selected for processing (9)

• cpp/src/io/parquet/decode_fixed.cu
• cpp/src/io/parquet/decode_preprocess.cu
• cpp/src/io/parquet/page_data.cu
• cpp/src/io/parquet/page_data.cuh
• cpp/src/io/parquet/page_decode.cuh
• cpp/src/io/parquet/page_delta_decode.cu
• cpp/src/io/parquet/page_state_composed.cuh
• cpp/src/io/parquet/page_string_decode.cu
• cpp/src/io/parquet/page_string_utils.cuh

[vyasr]

/merge