Enable chunked row masks in hybrid scan reader - #22716
Conversation
|
Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually. Contributors can view more details about this message here. |
| cudf::host_span<cudf::size_type const>(filtered_row_group_indices); | ||
|
|
||
| // Build a row mask for the filtered row groups using page index stats, or all true if no filter | ||
| auto row_mask = |
There was a problem hiding this comment.
Build a single row mask column across all row groups.
|
|
||
| // Helper to materialize filter and payload columns for a row group pass | ||
| std::size_t rows_materialized = 0; | ||
| auto const materialize_pass = [&](cudf::host_span<cudf::size_type const> row_group_indices) { |
There was a problem hiding this comment.
Create mutable views out of the global row_mask above at current pass' row group boundaries to reuse the same memory. Alternatively, one can just create a per-pass row_mask here with the row_group_indices parameter (row groups in this pass) and discard it once we are done materializing it.
| "Total number of rows exceeds cudf::size_type's limit"); | ||
|
|
||
| return static_cast<size_type>(total_rows); | ||
| return std::accumulate( |
There was a problem hiding this comment.
No logical differences here except using std::accumulate instead of for_each.
| total_rows += pfm.row_groups[row_group_idx].num_rows; | ||
| } | ||
| }); | ||
| CUDF_EXPECTS(std::cmp_less_equal(total_rows, std::numeric_limits<size_type>::max()), |
There was a problem hiding this comment.
Removed this as we now allow more than 2B rows
mhaseeb123
added
bug
mhaseeb123
added
3 - Ready for Review
📝 WalkthroughSummary by CodeRabbit
WalkthroughThis PR refactors hybrid scan's row-mask offset tracking from a global ChangesHybrid scan row-mask offset refactoring
Estimated code review effort🎯 3 (Moderate) | ⏱️ ~25 minutes Suggested labels
Suggested reviewers
🚥 Pre-merge checks | ✅ 4 | ❌ 1❌ Failed checks (1 warning)
✅ Passed checks (4 passed)
✏️ Tip: You can configure your own custom pre-merge checks in the settings. ✨ Finishing Touches🧪 Generate unit tests (beta)
Warning There were issues while running some tools. Please review the errors and either fix the tool's configuration or disable the tool if it's a critical failure. 🔧 Infer (1.2.0)cpp/tests/io/experimental/hybrid_scan_test.cppIn file included from cpp/tests/io/experimental/hybrid_scan_test.cpp:6: ... [truncated 2200 characters] ... ed from ClangFrontend__CFrontend_errors.protect in file "src/clang/cFrontend_errors.ml", line 48, characters 6-141 Comment |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
Actionable comments posted: 2
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/experimental/page_index_filter.cu (1)
999-1004:⚠️ Potential issue | 🟠 Major | ⚡ Quick winDon't run
all_ofover nullable row-mask data.This fast path still touches
row_mask.begin<bool>()when a mutable filter mask has mixed null and non-null values. Those null entries are only normalized later at Lines 1087-1097, so this branch can read undefined BOOL8 values and derive the wrong page mask.As per coding guidelines, Do not read null values of fixed-width columns; null values are undefined and must not be accessed.Suggested fix
- // Return an empty vector if all rows are invalid or all rows are required - if (std::cmp_equal(row_mask.null_count(row_mask_offset, row_mask_offset + total_rows, stream), - total_rows) or - cudf::detail::all_of(row_mask.template begin<bool>() + row_mask_offset, - row_mask.template begin<bool>() + row_mask_offset + total_rows, - cuda::std::identity{}, - stream)) { + auto const null_count = + row_mask.null_count(row_mask_offset, row_mask_offset + total_rows, stream); + + // Return an empty vector if all rows are unknown or all rows are required. + // Avoid reading nullable BOOL8 data before nulls are sanitized below. + if (std::cmp_equal(null_count, total_rows) or + (std::cmp_equal(null_count, 0) and + cudf::detail::all_of(row_mask.template begin<bool>() + row_mask_offset, + row_mask.template begin<bool>() + row_mask_offset + total_rows, + cuda::std::identity{}, + stream))) { return thrust::host_vector<bool>(0, stream); }🤖 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/experimental/page_index_filter.cu` around lines 999 - 1004, The current conditional calls row_mask.template begin<bool>() in the all_of call even when there are nulls, which can read undefined BOOL8 values; fix by first computing the null count (using row_mask.null_count(row_mask_offset, row_mask_offset + total_rows, stream)) into a variable and only invoke cudf::detail::all_of over row_mask.template begin<bool>() + row_mask_offset when that null count is zero (i.e., no nulls present); update the conditional around the existing check so the all_of branch is skipped if any nulls exist, referencing row_mask.null_count, row_mask.template begin<bool>(), row_mask_offset, total_rows and the current if condition.
♻️ Duplicate comments (1)
cpp/tests/io/experimental/hybrid_scan_composer.cpp (1)
256-260:⚠️ Potential issue | 🟠 Major | ⚡ Quick winBuild the row mask per pass instead of once for the full filtered set.
This still materializes a single
row_maskacrosscurrent_row_group_indices, so the helper keeps the same one-columncudf::size_typeceiling that this PR is trying to remove. Therows_materializedslicing is just the downstream symptom: later passes still depend on representing the full mask as one column first. Build the mask insidematerialize_pass()for that pass'srow_group_indices(or feed in pre-chunked mask views) and keep the pass-local view at offset0.Suggested direction
- auto row_mask = - options.get_filter().has_value() - ? reader->build_row_mask_with_page_index_stats(current_row_group_indices, options, stream, mr) - : reader->build_all_true_row_mask(current_row_group_indices, stream, mr); - auto filter_tables = std::vector<std::unique_ptr<cudf::table>>{}; auto payload_tables = std::vector<std::unique_ptr<cudf::table>>{}; - std::size_t rows_materialized = 0; auto const materialize_pass = [&](cudf::host_span<cudf::size_type const> row_group_indices) { - auto const rows_in_pass = reader->total_rows_in_row_groups(row_group_indices); - auto* null_mask = row_mask->nullable() ? row_mask->mutable_view().null_mask() : nullptr; - auto const slice_null_count = - cudf::null_count(null_mask, rows_materialized, rows_materialized + rows_in_pass, stream); - auto row_mask_view = cudf::mutable_column_view(..., rows_materialized); + auto row_mask = + options.get_filter().has_value() + ? reader->build_row_mask_with_page_index_stats(row_group_indices, options, stream, mr) + : reader->build_all_true_row_mask(row_group_indices, stream, mr); + auto row_mask_view = row_mask->mutable_view(); ... - - rows_materialized += rows_in_pass; };Also applies to: 266-278, 326-327
🤖 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/tests/io/experimental/hybrid_scan_composer.cpp` around lines 256 - 260, The code currently builds a single row_mask for the entire current_row_group_indices (via options.get_filter() ? reader->build_row_mask_with_page_index_stats(...) : reader->build_all_true_row_mask(...)), which preserves a one-column cudf::size_type ceiling; instead, move the mask construction into materialize_pass() so each pass builds a pass-local mask for that pass's row_group_indices (or accept and use pre-chunked mask views passed into materialize_pass()), and ensure the pass-local mask view is created at offset 0 (so downstream rows_materialized slicing sees a mask starting at column 0). Apply the same change for the other build_row_mask usages (the similar calls near the later occurrences) so no pass depends on a full-set single mask.
🧹 Nitpick comments (1)
cpp/tests/io/experimental/hybrid_scan_composer.cpp (1)
53-60: ⚡ Quick winDocument the
readerparameter in the Doxygen block.The updated comment is still missing
@param reader, so the function documentation is incomplete. As per coding guidelines, C++/CUDA code must include proper Doxygen documentation.🤖 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/tests/io/experimental/hybrid_scan_composer.cpp` around lines 53 - 60, The Doxygen block for "Apply hybrid scan row group filters" is missing documentation for the reader parameter; update the comment above the corresponding function (the Apply hybrid scan row group filters docblock) to include an `@param` reader line that briefly describes what reader represents and how it's used (e.g., the input reader providing row-group metadata/IO for filtering), ensuring the `@param` name exactly matches the function signature's parameter name.
🤖 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.
Inline comments:
In `@cpp/src/io/parquet/experimental/hybrid_scan_helpers.cpp`:
- Around line 191-198: The lambda that sums
file_metadata.row_groups[row_group_idx].num_rows into an unsigned accumulator
can accept negative num_rows and overflow; update the lambda (used with sum and
row_group_idx) to first assert that
file_metadata.row_groups[row_group_idx].num_rows is >= 0 (using CUDF_EXPECTS
with context including src_idx and row_group_idx), then convert that value to
the accumulator's unsigned type and check that sum + converted_num_rows does not
exceed std::numeric_limits<size_type>::max() (or otherwise detect overflow) and
fail early (CUDF_EXPECTS/CUDF_FAIL) if it would overflow; only then return the
safe sum.
In `@cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp`:
- Around line 386-393: The guard in hybrid_scan_impl.cpp incorrectly rejects the
boundary case by checking num_rows <
std::numeric_limits<cudf::size_type>::max(); update the CUDF_EXPECTS so it only
fails when num_rows is strictly greater than the max (i.e., allow equality) —
locate the check around the total_rows_in_row_groups(row_group_indices)
computation and change the comparison to permit num_rows ==
std::numeric_limits<cudf::size_type>::max(), leaving the error message and
exception type intact; no other logic (true_scalar creation or
make_column_from_scalar call) needs to change.
---
Outside diff comments:
In `@cpp/src/io/parquet/experimental/page_index_filter.cu`:
- Around line 999-1004: The current conditional calls row_mask.template
begin<bool>() in the all_of call even when there are nulls, which can read
undefined BOOL8 values; fix by first computing the null count (using
row_mask.null_count(row_mask_offset, row_mask_offset + total_rows, stream)) into
a variable and only invoke cudf::detail::all_of over row_mask.template
begin<bool>() + row_mask_offset when that null count is zero (i.e., no nulls
present); update the conditional around the existing check so the all_of branch
is skipped if any nulls exist, referencing row_mask.null_count,
row_mask.template begin<bool>(), row_mask_offset, total_rows and the current if
condition.
---
Duplicate comments:
In `@cpp/tests/io/experimental/hybrid_scan_composer.cpp`:
- Around line 256-260: The code currently builds a single row_mask for the
entire current_row_group_indices (via options.get_filter() ?
reader->build_row_mask_with_page_index_stats(...) :
reader->build_all_true_row_mask(...)), which preserves a one-column
cudf::size_type ceiling; instead, move the mask construction into
materialize_pass() so each pass builds a pass-local mask for that pass's
row_group_indices (or accept and use pre-chunked mask views passed into
materialize_pass()), and ensure the pass-local mask view is created at offset 0
(so downstream rows_materialized slicing sees a mask starting at column 0).
Apply the same change for the other build_row_mask usages (the similar calls
near the later occurrences) so no pass depends on a full-set single mask.
---
Nitpick comments:
In `@cpp/tests/io/experimental/hybrid_scan_composer.cpp`:
- Around line 53-60: The Doxygen block for "Apply hybrid scan row group filters"
is missing documentation for the reader parameter; update the comment above the
corresponding function (the Apply hybrid scan row group filters docblock) to
include an `@param` reader line that briefly describes what reader represents and
how it's used (e.g., the input reader providing row-group metadata/IO for
filtering), ensuring the `@param` name exactly matches the function signature's
parameter name.
🪄 Autofix (Beta)
Fix all unresolved CodeRabbit comments on this PR:
- Push a commit to this branch (recommended)
- Create a new PR with the fixes
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: ae2a7766-73c9-4d55-b5f0-e44e8efa5159
📒 Files selected for processing (6)
cpp/src/io/parquet/experimental/hybrid_scan_helpers.cppcpp/src/io/parquet/experimental/hybrid_scan_helpers.hppcpp/src/io/parquet/experimental/hybrid_scan_impl.cppcpp/src/io/parquet/experimental/hybrid_scan_impl.hppcpp/src/io/parquet/experimental/page_index_filter.cucpp/tests/io/experimental/hybrid_scan_composer.cpp
| [&](auto sum, auto const row_group_idx) { | ||
| CUDF_EXPECTS( | ||
| std::cmp_greater_equal(row_group_idx, size_type{0}) and | ||
| std::cmp_less(row_group_idx, file_metadata.row_groups.size()), | ||
| "Encountered out-of-bounds row group index for data source. Row group index: " + | ||
| std::to_string(row_group_idx) + ", Source index: " + std::to_string(src_idx) + | ||
| ", Number of row groups: " + std::to_string(file_metadata.row_groups.size())); | ||
| return sum + file_metadata.row_groups[row_group_idx].num_rows; |
There was a problem hiding this comment.
Validate row counts before adding them to the unsigned accumulator.
num_rows is metadata-derived and signed. sum + file_metadata.row_groups[row_group_idx].num_rows will wrap on negative metadata and still has no guard for std::size_t overflow, so this can return a bogus total and mis-size later row-mask or buffer work instead of failing fast.
🩹 Suggested guard
[&](auto sum, auto const row_group_idx) {
CUDF_EXPECTS(
std::cmp_greater_equal(row_group_idx, size_type{0}) and
std::cmp_less(row_group_idx, file_metadata.row_groups.size()),
"Encountered out-of-bounds row group index for data source. Row group index: " +
std::to_string(row_group_idx) + ", Source index: " + std::to_string(src_idx) +
", Number of row groups: " + std::to_string(file_metadata.row_groups.size()));
- return sum + file_metadata.row_groups[row_group_idx].num_rows;
+ auto const num_rows = file_metadata.row_groups[row_group_idx].num_rows;
+ CUDF_EXPECTS(num_rows >= 0,
+ "Encountered negative row count in row group metadata",
+ std::invalid_argument);
+ auto const rows = static_cast<std::size_t>(num_rows);
+ CUDF_EXPECTS(sum <= std::numeric_limits<std::size_t>::max() - rows,
+ "Total number of rows exceeds std::size_t",
+ std::overflow_error);
+ return sum + rows;
});🤖 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/experimental/hybrid_scan_helpers.cpp` around lines 191 -
198, The lambda that sums file_metadata.row_groups[row_group_idx].num_rows into
an unsigned accumulator can accept negative num_rows and overflow; update the
lambda (used with sum and row_group_idx) to first assert that
file_metadata.row_groups[row_group_idx].num_rows is >= 0 (using CUDF_EXPECTS
with context including src_idx and row_group_idx), then convert that value to
the accumulator's unsigned type and check that sum + converted_num_rows does not
exceed std::numeric_limits<size_type>::max() (or otherwise detect overflow) and
fail early (CUDF_EXPECTS/CUDF_FAIL) if it would overflow; only then return the
safe sum.
| auto const num_rows = total_rows_in_row_groups(row_group_indices); | ||
| CUDF_EXPECTS(num_rows < std::numeric_limits<cudf::size_type>::max(), | ||
| "Total rows in row groups exceed the cudf's column size limit. Retry with a smaller " | ||
| "set of row groups", | ||
| std::invalid_argument); | ||
| auto true_scalar = | ||
| cudf::numeric_scalar<bool>(true, true, stream, cudf::get_current_device_resource_ref()); | ||
| return cudf::make_column_from_scalar(true_scalar, num_rows, stream, mr); |
There was a problem hiding this comment.
Allow the max-sized pass through this guard.
Line 387 rejects num_rows == std::numeric_limits<cudf::size_type>::max(), even though that is still a valid column size. That turns the exact boundary case into a false overflow and needlessly breaks a max-sized pass.
Suggested fix
- CUDF_EXPECTS(num_rows < std::numeric_limits<cudf::size_type>::max(),
+ CUDF_EXPECTS(num_rows <= static_cast<std::size_t>(std::numeric_limits<cudf::size_type>::max()),
"Total rows in row groups exceed the cudf's column size limit. Retry with a smaller "
"set of row groups",
std::invalid_argument);
@@
- return cudf::make_column_from_scalar(true_scalar, num_rows, stream, mr);
+ return cudf::make_column_from_scalar(
+ true_scalar, static_cast<cudf::size_type>(num_rows), stream, mr);📝 Committable suggestion
‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.
| auto const num_rows = total_rows_in_row_groups(row_group_indices); | |
| CUDF_EXPECTS(num_rows < std::numeric_limits<cudf::size_type>::max(), | |
| "Total rows in row groups exceed the cudf's column size limit. Retry with a smaller " | |
| "set of row groups", | |
| std::invalid_argument); | |
| auto true_scalar = | |
| cudf::numeric_scalar<bool>(true, true, stream, cudf::get_current_device_resource_ref()); | |
| return cudf::make_column_from_scalar(true_scalar, num_rows, stream, mr); | |
| auto const num_rows = total_rows_in_row_groups(row_group_indices); | |
| CUDF_EXPECTS(num_rows <= static_cast<std::size_t>(std::numeric_limits<cudf::size_type>::max()), | |
| "Total rows in row groups exceed the cudf's column size limit. Retry with a smaller " | |
| "set of row groups", | |
| std::invalid_argument); | |
| auto true_scalar = | |
| cudf::numeric_scalar<bool>(true, true, stream, cudf::get_current_device_resource_ref()); | |
| return cudf::make_column_from_scalar( | |
| true_scalar, static_cast<cudf::size_type>(num_rows), stream, mr); |
🤖 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/experimental/hybrid_scan_impl.cpp` around lines 386 - 393,
The guard in hybrid_scan_impl.cpp incorrectly rejects the boundary case by
checking num_rows < std::numeric_limits<cudf::size_type>::max(); update the
CUDF_EXPECTS so it only fails when num_rows is strictly greater than the max
(i.e., allow equality) — locate the check around the
total_rows_in_row_groups(row_group_indices) computation and change the
comparison to permit num_rows == std::numeric_limits<cudf::size_type>::max(),
leaving the error message and exception type intact; no other logic (true_scalar
creation or make_column_from_scalar call) needs to change.
qbacpey
left a comment
qbacpey
left a comment
There was a problem hiding this comment.
LGTM. Optional: would it be useful to add a test that builds the row mask per row-group pass, to show the intended chunked row-mask workflow more directly?
Perhaps a separate test that does that. It was like that until my last commit but it looked a bit inefficient. |
mhaseeb123
added
4 - Needs Review
There was a problem hiding this comment.
🧹 Nitpick comments (2)
cpp/tests/io/experimental/hybrid_scan_test.cpp (2)
953-958: 💤 Low valueConsider asserting that page index is present to ensure focused test coverage.
The page index setup is currently optional. If page index is not present, the test would still run but might not exercise the intended page-level filtering code path (since
build_row_mask_with_page_index_statsat line 968 relies on page index).To make the test's intent clearer and ensure it exercises page-level filtering, consider adding an assertion after line 958:
ASSERT_FALSE(page_index_byte_range.is_empty()) << "Page index required for this test";This would make explicit that the test expects and requires page index, preventing silent test degradation if the parquet writing behavior changes.
🤖 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/tests/io/experimental/hybrid_scan_test.cpp` around lines 953 - 958, The test optionally sets up a page index but later calls build_row_mask_with_page_index_stats, so add an explicit assertion that the page index exists: after obtaining page_index_byte_range from reader->page_index_byte_range() (and before calling fetch_page_index_to_host / reader->setup_page_index), assert that page_index_byte_range.is_empty() is false (e.g., ASSERT_FALSE(page_index_byte_range.is_empty()) with a message like "Page index required for this test") so the test fails loudly if fetch_page_index_to_host / setup_page_index would be skipped and the page-level filtering code path (build_row_mask_with_page_index_stats) would not be exercised.
926-1033: ⚡ Quick winConsider adding edge case: single row group (single pass).
The test currently requires at least 2 row groups (line 1012:
ASSERT_GT(row_group_span.size(), 1)). To improve edge case coverage, consider adding a separate test case or relaxing the assertion to also validate the per-pass row mask workflow when there is only one pass (single row group). This would ensure that the chunking logic degrades gracefully and that per-pass row masks work correctly even without actual chunking.As per coding guidelines, test functions should cover boundary sizes and edge cases.
🤖 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/tests/io/experimental/hybrid_scan_test.cpp` around lines 926 - 1033, The test ChunkedReadRowMaskPerPass currently asserts multiple row groups via ASSERT_GT(row_group_span.size(), 1) and thus misses the single-row-group (single-pass) edge case; update the test to also exercise a single-pass path by either adding a new test or relaxing the assertion and running materialize_pass once when row_group_span.size() == 1, ensuring build_row_mask_with_page_index_stats, materialize_filter_columns and materialize_payload_columns are invoked with a single row_group_span.subspan and that concatenation/expectation logic still validates against expected from read_parquet; reference functions/vars: ChunkedReadRowMaskPerPass, row_group_span, materialize_pass, build_row_mask_with_page_index_stats, materialize_filter_columns, materialize_payload_columns, and ASSERT_GT.
🤖 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.
Nitpick comments:
In `@cpp/tests/io/experimental/hybrid_scan_test.cpp`:
- Around line 953-958: The test optionally sets up a page index but later calls
build_row_mask_with_page_index_stats, so add an explicit assertion that the page
index exists: after obtaining page_index_byte_range from
reader->page_index_byte_range() (and before calling fetch_page_index_to_host /
reader->setup_page_index), assert that page_index_byte_range.is_empty() is false
(e.g., ASSERT_FALSE(page_index_byte_range.is_empty()) with a message like "Page
index required for this test") so the test fails loudly if
fetch_page_index_to_host / setup_page_index would be skipped and the page-level
filtering code path (build_row_mask_with_page_index_stats) would not be
exercised.
- Around line 926-1033: The test ChunkedReadRowMaskPerPass currently asserts
multiple row groups via ASSERT_GT(row_group_span.size(), 1) and thus misses the
single-row-group (single-pass) edge case; update the test to also exercise a
single-pass path by either adding a new test or relaxing the assertion and
running materialize_pass once when row_group_span.size() == 1, ensuring
build_row_mask_with_page_index_stats, materialize_filter_columns and
materialize_payload_columns are invoked with a single row_group_span.subspan and
that concatenation/expectation logic still validates against expected from
read_parquet; reference functions/vars: ChunkedReadRowMaskPerPass,
row_group_span, materialize_pass, build_row_mask_with_page_index_stats,
materialize_filter_columns, materialize_payload_columns, and ASSERT_GT.
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 7d328131-3eda-44b3-86be-7533c67a5b47
📒 Files selected for processing (1)
cpp/tests/io/experimental/hybrid_scan_test.cpp
|
/merge |
mhaseeb123
added
5 - Ready to Merge
3 participants
Description
Closes #22672
This PR enables the hybrid scan reader to accept and use chunked row masks to support two-step materialization of more than
std::numeric_limits<cudf::size_type>::max()rows. Row mask columns (or their views) must be chunked across the same row boundaries as row group passes being materialized.Checklist