From e32363f102be74cb7fc6ed7acbdc40fba763e767 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Tue, 25 Aug 2026 07:53:24 +0000 Subject: [PATCH 01/21] Add hybrid scan page pruning without offset indexes --- .../hybrid_scan_io/hybrid_scan_composer.cpp | 26 +- .../cudf/io/experimental/hybrid_scan.hpp | 6 +- .../io/experimental/hybrid_scan_multifile.hpp | 8 +- .../io/parquet/experimental/hybrid_scan.cpp | 2 +- .../experimental/hybrid_scan_chunking.cu | 13 +- .../experimental/hybrid_scan_helpers.hpp | 10 +- .../parquet/experimental/hybrid_scan_impl.cpp | 90 ++++- .../parquet/experimental/hybrid_scan_impl.hpp | 8 +- .../experimental/hybrid_scan_multifile.cpp | 2 +- .../parquet/experimental/page_index_filter.cu | 356 +---------------- .../experimental/page_index_filter_utils.cu | 357 +++++++++++++++++- .../experimental/page_index_filter_utils.hpp | 33 +- .../experimental/hybrid_scan_filters_test.cpp | 6 +- .../java/ai/rapids/cudf/HybridScanReader.java | 18 +- .../src/HybridScanReaderJniMaterialize.cpp | 2 +- .../ai/rapids/cudf/HybridScanReaderTest.java | 133 ++++++- .../pylibcudf/io/experimental/hybrid_scan.pyx | 4 +- .../pylibcudf/libcudf/io/hybrid_scan.pxd | 2 +- .../tests/io/test_experimental_hybrid_scan.py | 67 ++++ 19 files changed, 703 insertions(+), 440 deletions(-) diff --git a/cpp/examples/hybrid_scan_io/hybrid_scan_composer.cpp b/cpp/examples/hybrid_scan_io/hybrid_scan_composer.cpp index a26839d9ffd6..c9fa43bea8eb 100644 --- a/cpp/examples/hybrid_scan_io/hybrid_scan_composer.cpp +++ b/cpp/examples/hybrid_scan_io/hybrid_scan_composer.cpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -426,22 +426,6 @@ std::unique_ptr hybrid_scan( } } -// Specialization for two-step read without page index -template - requires(not single_step_read and not use_page_index) -std::unique_ptr inline hybrid_scan( - io_source const& io_source, - std::optional filter_expression, - std::unordered_set const& filters, - bool verbose, - rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr) -{ - static_assert(single_step_read or use_page_index, - "Hybrid scan requires parquet page index for two-step parquet read"); - return nullptr; -} - // Instantiations for hybrid_scan template template std::unique_ptr hybrid_scan( @@ -460,6 +444,14 @@ template std::unique_ptr hybrid_scan( rmm::cuda_stream_view, rmm::device_async_resource_ref); +template std::unique_ptr hybrid_scan( + io_source const&, + std::optional, + std::unordered_set const&, + bool, + rmm::cuda_stream_view, + rmm::device_async_resource_ref); + template std::unique_ptr hybrid_scan( io_source const&, std::optional, diff --git a/cpp/include/cudf/io/experimental/hybrid_scan.hpp b/cpp/include/cudf/io/experimental/hybrid_scan.hpp index 772b65e62fc9..333fbfcf2c05 100644 --- a/cpp/include/cudf/io/experimental/hybrid_scan.hpp +++ b/cpp/include/cudf/io/experimental/hybrid_scan.hpp @@ -564,7 +564,7 @@ class hybrid_scan_reader { * * @param row_group_indices Input row groups indices * @param column_chunk_data Device spans of column chunk data of filter columns - * @param[in,out] row_mask Mutable boolean column indicating surviving rows from page pruning + * @param[in,out] row_mask Mutable boolean column indicating surviving rows * @param mask_data_pages Whether to build and use a data page mask using the row mask * @param options Parquet reader options * @param stream CUDA stream used for device memory operations and kernel launches @@ -645,7 +645,7 @@ class hybrid_scan_reader { * @param pass_read_limit Limit on the memory used for reading and decompressing data. `0` if * there is no limit * @param row_group_indices Input row groups indices - * @param row_mask Boolean column indicating which rows need to be read + * @param[in,out] row_mask Mutable boolean column indicating surviving rows * @param mask_data_pages Whether to build and use a data page mask using the row mask * @param column_chunk_data Device spans of column chunk data of filter columns * @param options Parquet reader options @@ -656,7 +656,7 @@ class hybrid_scan_reader { std::size_t chunk_read_limit, std::size_t pass_read_limit, std::span row_group_indices, - cudf::column_view const& row_mask, + cudf::mutable_column_view const& row_mask, use_data_page_mask mask_data_pages, std::span const> column_chunk_data, parquet_reader_options const& options, diff --git a/cpp/include/cudf/io/experimental/hybrid_scan_multifile.hpp b/cpp/include/cudf/io/experimental/hybrid_scan_multifile.hpp index c75fa3d186d3..0cc290270ff5 100644 --- a/cpp/include/cudf/io/experimental/hybrid_scan_multifile.hpp +++ b/cpp/include/cudf/io/experimental/hybrid_scan_multifile.hpp @@ -274,7 +274,7 @@ class hybrid_scan_multifile { * @param column_chunk_data Flattened device spans of filter column chunk data returned in the * same order as `filter_column_chunks_byte_ranges` * @param[in,out] row_mask Mutable boolean column spanning all selected rows across all sources - * and indicating surviving rows from page pruning + * indicating surviving rows * @param mask_data_pages Whether to build and use a data page mask using the row mask * @param options Parquet reader options * @param stream CUDA stream used for device memory operations and kernel launches @@ -389,8 +389,8 @@ class hybrid_scan_multifile { * @param pass_read_limit Limit on the memory used for reading and decompressing data. `0` if * there is no limit * @param row_group_indices Span of vectors of input row group indices, one per source - * @param row_mask Boolean column spanning all selected rows across all sources and indicating - * which rows need to be read + * @param[in,out] row_mask Mutable boolean column spanning all selected rows across all sources + * indicating surviving rows * @param mask_data_pages Whether to build and use a data page mask using the row mask * @param column_chunk_data Flattened device spans of filter column chunk data returned in the * same order as `filter_column_chunks_byte_ranges` @@ -402,7 +402,7 @@ class hybrid_scan_multifile { std::size_t chunk_read_limit, std::size_t pass_read_limit, cudf::host_span const> row_group_indices, - cudf::column_view const& row_mask, + cudf::mutable_column_view const& row_mask, use_data_page_mask mask_data_pages, cudf::host_span const> column_chunk_data, parquet_reader_options const& options, diff --git a/cpp/src/io/parquet/experimental/hybrid_scan.cpp b/cpp/src/io/parquet/experimental/hybrid_scan.cpp index 09faac8c261a..4f98d14ee92c 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan.cpp +++ b/cpp/src/io/parquet/experimental/hybrid_scan.cpp @@ -293,7 +293,7 @@ void hybrid_scan_reader::setup_chunking_for_filter_columns( std::size_t chunk_read_limit, std::size_t pass_read_limit, std::span row_group_indices, - cudf::column_view const& row_mask, + cudf::mutable_column_view const& row_mask, use_data_page_mask mask_data_pages, std::span const> column_chunk_data, parquet_reader_options const& options, diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu b/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu index 970f790da3e3..4918b8b1c532 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu +++ b/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu @@ -126,7 +126,18 @@ void hybrid_scan_reader_impl::setup_next_pass( set_sparse_pass_page_mask(column_chunk_data); } else { setup_compressed_data(column_chunk_data); - set_pass_page_mask(data_page_mask); + // When offset index is absent, compute and use the data page mask using the decoded page + // headers from `setup_compressed_data`. + auto const data_page_mask_pghdr = [&]() { + if (not _has_offset_index and not _row_mask.is_empty()) { + return compute_data_page_mask_with_page_headers(); + } + return thrust::host_vector{}; + }(); + set_pass_page_mask( + data_page_mask_pghdr.empty() + ? data_page_mask + : std::span{data_page_mask_pghdr.data(), data_page_mask_pghdr.size()}); } // detect malformed columns. diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_helpers.hpp b/cpp/src/io/parquet/experimental/hybrid_scan_helpers.hpp index 59591d438913..44e199a52061 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_helpers.hpp +++ b/cpp/src/io/parquet/experimental/hybrid_scan_helpers.hpp @@ -343,24 +343,18 @@ class aggregate_reader_metadata : public aggregate_reader_metadata_base { * Compute a vector of boolean vectors indicating which data pages need to be decoded to * construct each input column based on the row mask, one vector per column * - * @tparam ColumnView Type of the row mask column view - cudf::mutable_column_view for filter - * columns and cudf::column_view for payload columns - * - * @param row_mask Boolean column indicating which rows need to be read after page-pruning + * @param row_mask Non-nullable boolean column view indicating surviving rows * @param row_group_indices Input row groups indices * @param input_columns Input column information - * @param row_mask_offset Offset into the row mask column for the current pass * @param stream CUDA stream used for device memory operations and kernel launches * * @return Boolean vector indicating which data pages need to be decoded to produce * the output table based on the input row mask across all input columns */ - template [[nodiscard]] thrust::host_vector compute_data_page_mask( - ColumnView const& row_mask, + cudf::column_view const& row_mask, std::span const> row_group_indices, std::span input_columns, - cudf::size_type row_mask_offset, cuda::stream_ref stream) const; }; diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp b/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp index 82b15871408f..2b6f703e5464 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp +++ b/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp @@ -8,6 +8,7 @@ #include "cudf/io/text/byte_range_info.hpp" #include "hybrid_scan_helpers.hpp" #include "io/parquet/reader_impl_chunking_utils.cuh" +#include "page_index_filter_utils.hpp" #include #include @@ -594,8 +595,8 @@ hybrid_scan_reader_impl::payload_pages_byte_ranges( // Compute the data page mask auto const mask_size = mask_offsets.back(); - auto data_page_mask = _extended_metadata->compute_data_page_mask( - row_mask, row_group_indices, _input_columns, 0, stream); + auto data_page_mask = + _extended_metadata->compute_data_page_mask(row_mask, row_group_indices, _input_columns, stream); CUDF_EXPECTS(data_page_mask.empty() or data_page_mask.size() == mask_size, "Computed data page mask does not match offset indexes"); @@ -696,8 +697,9 @@ table_with_metadata hybrid_scan_reader_impl::materialize_filter_columns( auto data_page_mask = thrust::host_vector{}; if (mask_data_pages == use_data_page_mask::YES) { + _row_mask = set_nulls_to_true(row_mask, stream); data_page_mask = _extended_metadata->compute_data_page_mask( - row_mask, row_group_indices, _input_columns, _row_mask_offset, stream); + _row_mask, row_group_indices, _input_columns, stream); } prepare_data(read_mode::READ_ALL, row_group_indices, column_chunk_data, data_page_mask); @@ -734,8 +736,9 @@ table_with_metadata hybrid_scan_reader_impl::materialize_payload_columns( auto data_page_mask = thrust::host_vector{}; if (not row_mask.is_empty() and mask_data_pages == use_data_page_mask::YES) { + _row_mask = row_mask; data_page_mask = _extended_metadata->compute_data_page_mask( - row_mask, row_group_indices, _input_columns, _row_mask_offset, stream); + _row_mask, row_group_indices, _input_columns, stream); } prepare_data(read_mode::READ_ALL, row_group_indices, column_chunk_data, data_page_mask); @@ -773,7 +776,7 @@ void hybrid_scan_reader_impl::setup_chunking_for_filter_columns( std::size_t chunk_read_limit, std::size_t pass_read_limit, std::span const> row_group_indices, - cudf::column_view const& row_mask, + cudf::mutable_column_view const& row_mask, use_data_page_mask mask_data_pages, std::span const> column_chunk_data, parquet_reader_options const& options, @@ -805,8 +808,9 @@ void hybrid_scan_reader_impl::setup_chunking_for_filter_columns( auto data_page_mask = thrust::host_vector{}; if (mask_data_pages == use_data_page_mask::YES) { + _row_mask = set_nulls_to_true(row_mask, stream); data_page_mask = _extended_metadata->compute_data_page_mask( - row_mask, row_group_indices, _input_columns, _row_mask_offset, stream); + _row_mask, row_group_indices, _input_columns, stream); } prepare_data(read_mode::CHUNKED_READ, row_group_indices, column_chunk_data, data_page_mask); @@ -865,8 +869,9 @@ void hybrid_scan_reader_impl::setup_chunking_for_payload_columns( auto data_page_mask = thrust::host_vector{}; if (not row_mask.is_empty() and mask_data_pages == use_data_page_mask::YES) { + _row_mask = row_mask; data_page_mask = _extended_metadata->compute_data_page_mask( - row_mask, row_group_indices, _input_columns, _row_mask_offset, stream); + _row_mask, row_group_indices, _input_columns, stream); } prepare_data(read_mode::CHUNKED_READ, row_group_indices, column_chunk_data, data_page_mask); @@ -1074,7 +1079,6 @@ bool hybrid_scan_reader_impl::has_next_table_chunk() void hybrid_scan_reader_impl::reset_internal_state() { - _row_mask_offset = 0; _file_itm_data = file_intermediate_data{}; _file_preprocessed = false; _has_offset_index = false; @@ -1098,6 +1102,10 @@ void hybrid_scan_reader_impl::reset_internal_state() _output_chunk_read_limit = 0; _strings_to_categorical = false; _reader_column_schema.reset(); + + _row_mask = column_view{}; + _row_mask_offset = 0; + _expr_conv = parquet_filter_normalizer{}; _mr = cudf::get_current_device_resource_ref(); } @@ -1441,6 +1449,72 @@ void hybrid_scan_reader_impl::set_pass_page_mask(std::span data_page "Encountered mismatch in number of pass pages and page mask size"); } +thrust::host_vector hybrid_scan_reader_impl::compute_data_page_mask_with_page_headers() +{ + auto& pass = *_pass_itm_data; + pass.pages.device_to_host_async(_stream); + _stream.sync(); + + std::vector page_row_offsets; + page_row_offsets.reserve(pass.pages.size() * 2); + + // Maps each data page to its flat-page range; -1 keeps nested pages enabled. + std::vector row_range_map; + row_range_map.reserve(pass.pages.size()); + + cudf::size_type previous_chunk_idx = -1; + auto max_page_size = cudf::size_type{0}; + + for (auto const& page : pass.pages) { + // Ignore dictionary pages altogether + if (page.flags & parquet::detail::PAGEINFO_FLAGS_DICTIONARY) { continue; } + + auto const& chunk = pass.chunks[page.chunk_idx]; + + // Don't prune list column pages as rows may span page boundaries when offset index isn't + // present. + if (chunk.max_level[parquet::detail::level_type::REPETITION] > 0) { + row_range_map.push_back(-1); + continue; + } + + auto const page_start = chunk.start_row + page.chunk_row; + auto const page_end = page_start + page.num_rows; + max_page_size = std::max(max_page_size, page_end - page_start); + + // Starting a new column chunk. Push page start row + if (previous_chunk_idx == -1 or page.chunk_idx != previous_chunk_idx) { + page_row_offsets.push_back(page_start); + previous_chunk_idx = page.chunk_idx; + } + + // Push row range index and page end row + row_range_map.push_back(page_row_offsets.size() - 1); + page_row_offsets.push_back(page_end); + } + + auto data_page_mask = thrust::host_vector{}; + + // Compute the row range mask + CUDF_EXPECTS(std::cmp_equal(_row_mask.size(), pass.num_rows), + "Row mask must span across all rows in the pass"); + auto const row_range_mask = + compute_row_range_selection_mask(_row_mask, page_row_offsets, max_page_size, _stream); + + if (row_range_mask.empty()) { return data_page_mask; } + + CUDF_EXPECTS(row_range_mask.size() == page_row_offsets.size() - 1, + "Encountered invalid row range mask size"); + + data_page_mask.reserve(row_range_map.size()); + + // Scatter row range results while retaining list column pages. + for (auto const range_idx : row_range_map) { + data_page_mask.push_back(range_idx < 0 ? true : row_range_mask[range_idx]); + } + return data_page_mask; +} + void hybrid_scan_reader_impl::set_sparse_pass_page_mask( std::span const> page_data) { diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_impl.hpp b/cpp/src/io/parquet/experimental/hybrid_scan_impl.hpp index b7c38ac71a69..a7380cff3b68 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_impl.hpp +++ b/cpp/src/io/parquet/experimental/hybrid_scan_impl.hpp @@ -253,7 +253,7 @@ class hybrid_scan_reader_impl : public parquet::detail::reader_impl { std::size_t chunk_read_limit, std::size_t pass_read_limit, std::span const> row_group_indices, - cudf::column_view const& row_mask, + cudf::mutable_column_view const& row_mask, use_data_page_mask mask_data_pages, std::span const> column_chunk_data, parquet_reader_options const& options, @@ -394,6 +394,11 @@ class hybrid_scan_reader_impl : public parquet::detail::reader_impl { */ void set_sparse_pass_page_mask(std::span const> page_data); + /** + * @brief Compute a data page mask from the decoded page headers. + */ + [[nodiscard]] thrust::host_vector compute_data_page_mask_with_page_headers(); + /** * @brief Select the columns to be read based on the read mode * @@ -616,6 +621,7 @@ class hybrid_scan_reader_impl : public parquet::detail::reader_impl { std::optional> _filter_columns_names; + cudf::column_view _row_mask{}; cudf::size_type _row_mask_offset{0}; bool _output_chunk_produced{false}; diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_multifile.cpp b/cpp/src/io/parquet/experimental/hybrid_scan_multifile.cpp index 259435401aab..16360d0bc7d9 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_multifile.cpp +++ b/cpp/src/io/parquet/experimental/hybrid_scan_multifile.cpp @@ -197,7 +197,7 @@ void hybrid_scan_multifile::setup_chunking_for_filter_columns( std::size_t chunk_read_limit, std::size_t pass_read_limit, cudf::host_span const> row_group_indices, - cudf::column_view const& row_mask, + cudf::mutable_column_view const& row_mask, use_data_page_mask mask_data_pages, cudf::host_span const> column_chunk_data, parquet_reader_options const& options, diff --git a/cpp/src/io/parquet/experimental/page_index_filter.cu b/cpp/src/io/parquet/experimental/page_index_filter.cu index 7ec4aa859f0c..ad57c3417d61 100644 --- a/cpp/src/io/parquet/experimental/page_index_filter.cu +++ b/cpp/src/io/parquet/experimental/page_index_filter.cu @@ -10,19 +10,15 @@ #include #include -#include #include #include #include #include -#include #include #include -#include #include #include #include -#include #include #include #include @@ -41,7 +37,6 @@ #include #include -#include #include namespace cudf::io::parquet::experimental::detail { @@ -577,14 +572,13 @@ struct page_stats_to_row_mask_converter : public page_stats_caster { auto const page_mask_nullmask = page_mask->null_count() - ? cudf::detail::make_host_vector_async( + ? cudf::detail::make_host_vector( cudf::device_span{ page_mask->view().null_mask(), static_cast(num_bitmask_words(page_mask->size()))}, stream) : cudf::detail::make_empty_host_vector(0, stream); - stream.sync(); auto [row_mask_data, row_mask_bitmask] = build_data_and_nullmask(page_mask->mutable_view(), page_mask_nullmask.data(), @@ -606,234 +600,6 @@ struct page_stats_to_row_mask_converter : public page_stats_caster { } }; -/* - * @brief Functor to build a Fenwick tree level from the previous level data - * - * @param tree_level_ptrs Pointers to the start of Fenwick tree level data - * @param prev_level Previous tree level - * @param prev_level_size Size of the previous tree level - * @param current_level_size Size of the current tree level - */ -struct build_fenwick_tree_level_functor { - bool** tree_level_ptrs; - cudf::size_type prev_level; - cudf::size_type prev_level_size; - cudf::size_type current_level_size; - - /** - * @brief Builds the next Fenwick tree level from the current level data - * by ORing two elements at the current level. - * - * elem_current_level[idx] = elem_prev_level[idx * 2] OR elem_prev_level[idx * 2 + 1]; - * - * @param current_level_idx Current tree level element index - */ - __device__ void operator()(cudf::size_type current_level_idx) const noexcept - { - auto const prev_level_ptr = tree_level_ptrs[prev_level]; - auto current_level_ptr = tree_level_ptrs[prev_level + 1]; - - // Handle the odd-sized remaining element if prev_level_size is odd - if (prev_level_size % 2 and current_level_idx == current_level_size - 1) { - current_level_ptr[current_level_idx] = prev_level_ptr[prev_level_size - 1]; - } else { - current_level_ptr[current_level_idx] = - prev_level_ptr[(current_level_idx * 2)] or prev_level_ptr[(current_level_idx * 2) + 1]; - } - } -}; - -/** - * @brief Functor to binary search a `true` value in the Fenwick tree in range [start, end) - * - * @param tree_level_ptrs Pointers to the start of Fenwick tree level data - * @param page_offsets Pointer to page offsets describing each search range i as [page_offsets[i], - * page_offsets[i+1)) - * @param num_ranges Number of search ranges - */ -struct search_fenwick_tree_functor { - bool** tree_level_ptrs; - cudf::size_type const* page_offsets; - cudf::size_type num_ranges; - - /** - * @brief Enum class to represent which range boundary we are currently processing - */ - enum class boundary : uint8_t { - START = 0, - END = 1, - }; - - /** - * @brief Checks if a value is a power of two - * - * @param value Value to check - * @return Boolean indicating if the value is a power of two - */ - __device__ bool inline constexpr is_power_of_two(cudf::size_type value) const noexcept - { - return (value & (value - 1)) == 0; - } - - /** - * @brief Finds the smallest power of two in the range [start, end). If no power of two is - * found, returns a zero. - * - * @param start Range start - * @param end Range end - * @return Largest power of two in the range [start, end) or a zero if no power of two is found - */ - __device__ cudf::size_type inline constexpr smallest_power_of_two_in_range( - cudf::size_type start, cudf::size_type end) const noexcept - { - start--; - start |= start >> 1; - start |= start >> 2; - start |= start >> 4; - start |= start >> 8; - start |= start >> 16; - auto const result = start + 1; - return result < end ? result : 0; - } - - /** - * @brief Finds the largest power of two in the range (start, end]. If no power of two is found, - * returns a zero. - * - * @param start Range start - * @param end Range end - * @return Largest power of two in the range (start, end] or a zero if no power of two is found - */ - __device__ size_type inline constexpr largest_power_of_two_in_range(size_type start, - size_type end) const noexcept - { - auto constexpr nbits = cudf::detail::size_in_bits() - 1; - auto const result = size_type{1} << (nbits - cuda::std::countl_zero(end)); - return result > start ? result : 0; - } - - /** - * @brief Aligns a range boundary to the next power-of-two block - * - * @tparam Boundary Current boundary type (START or END) - * @param start Range start - * @param end Range end - * @return A pair of the tree level and block size - */ - template - __device__ auto inline constexpr align_range_boundary(cudf::size_type start, - cudf::size_type end) const noexcept - { - if constexpr (Boundary == boundary::START) { - if (start == 0 or is_power_of_two(start)) { - auto const block_size = - cuda::std::max(start & -start, largest_power_of_two_in_range(start, end)); - auto const tree_level = cuda::std::countr_zero(block_size); - return cuda::std::pair{tree_level, block_size}; - } else { - auto const tree_level = cuda::std::countr_zero(start); - return cuda::std::pair{tree_level, size_type{1} << tree_level}; - } - } else { - auto block_size = end & -end; - if (start > 0 and is_power_of_two(end)) { - auto const next_alignment = cuda::std::max(smallest_power_of_two_in_range(start, end), - largest_power_of_two_in_range(0, end - start)); - block_size = end - next_alignment; - } - return cuda::std::pair{cuda::std::countr_zero(block_size), block_size}; - } - } - - /** - * @brief Queries the Fenwick tree for the given boundary position, tree level and block size - * - * @tparam Boundary Current boundary type (START or END) - * @param boundary_pos Current boundary position - * @param tree_level Corresponding tree level to query - * @param block_size Alignment block size of the current boundary - * @return Boolean indicating if a `true` value is found in the fenwick tree - */ - template - __device__ bool inline constexpr query_fenwick_tree(cudf::size_type boundary_pos, - cudf::size_type tree_level, - cudf::size_type block_size) const noexcept - { - if constexpr (Boundary == boundary::START) { - auto const mask_index = boundary_pos >> tree_level; - return tree_level_ptrs[tree_level][mask_index]; - } else { - auto const mask_index = (boundary_pos - block_size) >> tree_level; - return tree_level_ptrs[tree_level][mask_index]; - } - } - - /** - * @brief Searches the Fenwick tree to find a `true` value in range [start, end) - * - * Algorithm: While `start` < `end`, align `start` UP and `end` DOWN to the next power-of-two - * searchable tree block. For the two aligned blocks, query the fenwick tree at corresponding - * levels for a `true` value (larger block first). If found, return. Else, move the boundaries - * to their alignments. - * - * @param range_idx Index of the range to search - * @return Boolean indicating if a `true` value is found in the range - */ - __device__ bool operator()(cudf::size_type range_idx) const noexcept - { - // Retrieve start and end for the current range [start, end) - size_type start = page_offsets[range_idx]; - size_type end = page_offsets[range_idx + 1]; - - // Return early if the range is empty or invalid - if (start >= end or range_idx >= num_ranges) { return false; } - - // Binary search decomposition loop - while (start < end) { - // Find the largest power-of-two block that aligns `start` up - auto const [start_tree_level, start_block_size] = - align_range_boundary(start, end); - - // Find the largest power-of-two block that aligns `end` down - auto const [end_tree_level, end_block_size] = align_range_boundary(start, end); - - // Check the larger block first to minimize the number of queries - if (start_block_size >= end_block_size) { - // Check the `start` side alignment block first - if (start + start_block_size <= end) { - if (query_fenwick_tree(start, start_tree_level, start_block_size)) { - return true; - } - start += start_block_size; - } - // Check the `end` side alignment block if it's still in range - if (end - end_block_size >= start) { - if (query_fenwick_tree(end, end_tree_level, end_block_size)) { - return true; - } - end -= end_block_size; - } - } else { - // Check the `end` side alignment block first - if (end - end_block_size >= start) { - if (query_fenwick_tree(end, end_tree_level, end_block_size)) { - return true; - } - end -= end_block_size; - } - // Check the `start` side alignment block if it's still in range - if (start + start_block_size <= end) { - if (query_fenwick_tree(start, start_tree_level, start_block_size)) { - return true; - } - start += start_block_size; - } - } - } - return false; - } -}; - } // namespace std::unique_ptr aggregate_reader_metadata::build_row_mask_with_page_index_stats( @@ -978,12 +744,10 @@ std::unique_ptr aggregate_reader_metadata::build_row_mask_with_pag page_stats_table, stats_expr.get_stats_expr().get(), stream, mr); } -template thrust::host_vector aggregate_reader_metadata::compute_data_page_mask( - ColumnView const& row_mask, + cudf::column_view const& row_mask, std::span const> row_group_indices, std::span input_columns, - cudf::size_type row_mask_offset, cuda::stream_ref stream) const { CUDF_FUNC_RANGE(); @@ -998,15 +762,15 @@ thrust::host_vector aggregate_reader_metadata::compute_data_page_mask( std::invalid_argument); CUDF_EXPECTS( - std::cmp_less_equal(row_mask_offset + total_rows, row_mask.size()), + std::cmp_equal(total_rows, row_mask.size()), "Encountered a mismatch in number of rows in the row group pass and the row mask size", std::overflow_error); + CUDF_EXPECTS( + row_mask.null_count() == 0, "Row mask must not contain nulls", std::invalid_argument); - // 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() + row_mask_offset, - row_mask.template begin() + row_mask_offset + total_rows, + // Return an empty vector if all rows are required + if (cudf::detail::all_of(row_mask.begin(), + row_mask.begin() + total_rows, cuda::std::identity{}, stream)) { return thrust::host_vector(0); @@ -1099,112 +863,26 @@ thrust::host_vector aggregate_reader_metadata::compute_data_page_mask( }); } - // Make sure all row_mask elements contain valid values even if they are nulls - if constexpr (cuda::std::is_same_v) { - if (row_mask.nullable() and row_mask.null_count() > 0) { - thrust::for_each(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), - cuda::counting_iterator(row_mask_offset), - cuda::counting_iterator(row_mask_offset + total_rows), - [row_mask = row_mask.template begin(), - null_mask = row_mask.null_mask()] __device__(auto const row_idx) { - if (not bit_is_set(null_mask, row_idx)) { row_mask[row_idx] = true; } - }); - } - } else { - CUDF_EXPECTS(not row_mask.nullable() or row_mask.null_count() == 0, - "Row mask must not contain nulls for payload columns"); - } - - auto const mr = cudf::get_current_device_resource_ref(); - - // Compute fenwick tree level offsets and total size (level 1 and higher) - auto const tree_level_offsets = compute_fenwick_tree_level_offsets(total_rows, max_page_size); - auto const num_levels = static_cast(tree_level_offsets.size()); - // Buffer to store Fenwick tree levels (level 1 and higher) data - auto tree_levels_data = rmm::device_uvector(tree_level_offsets.back(), stream, mr); - - // Pointers to each Fenwick tree level data - auto host_tree_level_ptrs = cudf::detail::make_pinned_vector_async(num_levels, stream); - // Zeroth level is just the row mask itself - host_tree_level_ptrs[0] = const_cast(row_mask.template begin()) + row_mask_offset; - std::for_each(cuda::counting_iterator{1}, - cuda::counting_iterator{num_levels}, - [&](auto const level_idx) { - host_tree_level_ptrs[level_idx] = - tree_levels_data.data() + tree_level_offsets[level_idx - 1]; - }); + auto data_page_mask = thrust::host_vector{}; - auto fenwick_tree_level_ptrs = - cudf::detail::make_device_uvector_async(host_tree_level_ptrs, stream, mr); - - // Build Fenwick tree levels (zeroth level is just the row mask itself) - auto prev_level_size = static_cast(total_rows); - std::for_each( - cuda::counting_iterator{0}, - cuda::counting_iterator{num_levels - 1}, - [&](auto const prev_level) { - auto const current_level_size = cudf::util::div_rounding_up_safe(prev_level_size, 2); - thrust::for_each( - rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), - cuda::counting_iterator{0}, - cuda::counting_iterator{current_level_size}, - build_fenwick_tree_level_functor{ - fenwick_tree_level_ptrs.data(), prev_level, prev_level_size, current_level_size}); - prev_level_size = current_level_size; - }); - - // Search the Fenwick tree to see if there's a surviving row in each page's row range - auto const num_ranges = static_cast(page_row_offsets.size() - 1); - rmm::device_uvector device_data_page_mask(num_ranges, stream, mr); - // Use a pinned bounce buffer to avoid pageable h2d copy - auto pinned_page_offsets = cudf::detail::make_pinned_vector( - cudf::host_span{page_row_offsets}, stream); - auto page_offsets = cudf::detail::make_device_uvector_async(pinned_page_offsets, stream, mr); - thrust::transform( - rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), - cuda::counting_iterator{0}, - cuda::counting_iterator{num_ranges}, - device_data_page_mask.begin(), - search_fenwick_tree_functor{fenwick_tree_level_ptrs.data(), page_offsets.data(), num_ranges}); - - // Copy over search results to host - auto host_results = cudf::detail::make_pinned_vector_async(device_data_page_mask, stream); - auto const total_pages = pinned_page_offsets.size() - num_columns; - auto data_page_mask = thrust::host_vector{}; - data_page_mask.reserve(total_pages); - auto host_results_iter = host_results.begin(); - stream.sync(); + auto const row_range_mask = + compute_row_range_selection_mask(row_mask, page_row_offsets, max_page_size, stream); + if (row_range_mask.empty()) { return data_page_mask; } + data_page_mask.reserve(page_row_offsets.size() - num_columns); // Discard results for invalid ranges. i.e. ranges starting at the last page of a column and // ending at the first page of the next column - auto num_pages_inserted = 0; std::for_each(cuda::counting_iterator{0}, cuda::counting_iterator{num_columns}, [&](auto col_idx) { auto const col_num_pages = col_page_offsets[col_idx + 1] - col_page_offsets[col_idx] - 1; - data_page_mask.insert(data_page_mask.begin() + num_pages_inserted, - host_results_iter, - host_results_iter + col_num_pages); - host_results_iter += col_num_pages + 1; - num_pages_inserted += col_num_pages; + auto const first_page_range = col_page_offsets[col_idx]; + data_page_mask.insert(data_page_mask.end(), + row_range_mask.begin() + first_page_range, + row_range_mask.begin() + first_page_range + col_num_pages); }); return data_page_mask; } -// Instantiate the templates with ColumnView as cudf::column_view and cudf::mutable_column_view -template thrust::host_vector aggregate_reader_metadata::compute_data_page_mask< - cudf::column_view>(cudf::column_view const& row_mask, - std::span const> row_group_indices, - std::span input_columns, - cudf::size_type row_mask_offset, - cuda::stream_ref stream) const; - -template thrust::host_vector aggregate_reader_metadata::compute_data_page_mask< - cudf::mutable_column_view>(cudf::mutable_column_view const& row_mask, - std::span const> row_group_indices, - std::span input_columns, - cudf::size_type row_mask_offset, - cuda::stream_ref stream) const; - } // namespace cudf::io::parquet::experimental::detail diff --git a/cpp/src/io/parquet/experimental/page_index_filter_utils.cu b/cpp/src/io/parquet/experimental/page_index_filter_utils.cu index 07c81c787afe..89332b889359 100644 --- a/cpp/src/io/parquet/experimental/page_index_filter_utils.cu +++ b/cpp/src/io/parquet/experimental/page_index_filter_utils.cu @@ -5,16 +5,24 @@ #include "page_index_filter_utils.hpp" +#include +#include +#include #include #include #include #include #include +#include #include +#include +#include #include -#include +#include +#include +#include #include #include @@ -22,6 +30,266 @@ namespace cudf::io::parquet::experimental::detail { +namespace { + +/* + * @brief Functor to build a Fenwick tree level from the previous level data + * + * @param tree_level_ptrs Pointers to the start of Fenwick tree level data + * @param prev_level Previous tree level + * @param prev_level_size Size of the previous tree level + * @param current_level_size Size of the current tree level + */ +struct build_fenwick_tree_level_functor { + bool** tree_level_ptrs; + cudf::size_type prev_level; + cudf::size_type prev_level_size; + cudf::size_type current_level_size; + + /** + * @brief Builds the next Fenwick tree level from the current level data + * by ORing two elements at the current level. + * + * elem_current_level[idx] = elem_prev_level[idx * 2] OR elem_prev_level[idx * 2 + 1]; + * + * @param current_level_idx Current tree level element index + */ + __device__ void operator()(cudf::size_type current_level_idx) const noexcept + { + auto const prev_level_ptr = tree_level_ptrs[prev_level]; + auto current_level_ptr = tree_level_ptrs[prev_level + 1]; + + // Handle the odd-sized remaining element if prev_level_size is odd + if (prev_level_size % 2 and current_level_idx == current_level_size - 1) { + current_level_ptr[current_level_idx] = prev_level_ptr[prev_level_size - 1]; + } else { + current_level_ptr[current_level_idx] = + prev_level_ptr[(current_level_idx * 2)] or prev_level_ptr[(current_level_idx * 2) + 1]; + } + } +}; + +/** + * @brief Functor to binary search a `true` value in the Fenwick tree in range [start, end) + * + * @param tree_level_ptrs Pointers to the start of Fenwick tree level data + * @param page_offsets Pointer to page offsets describing each search range i as [page_offsets[i], + * page_offsets[i+1)) + * @param num_ranges Number of search ranges + */ +struct search_fenwick_tree_functor { + bool** tree_level_ptrs; + cudf::size_type const* page_offsets; + cudf::size_type num_ranges; + + /** + * @brief Enum class to represent which range boundary we are currently processing + */ + enum class boundary : uint8_t { + START = 0, + END = 1, + }; + + /** + * @brief Checks if a value is a power of two + * + * @param value Value to check + * @return Boolean indicating if the value is a power of two + */ + __device__ bool inline constexpr is_power_of_two(cudf::size_type value) const noexcept + { + return (value & (value - 1)) == 0; + } + + /** + * @brief Finds the smallest power of two in the range [start, end). If no power of two is + * found, returns a zero. + * + * @param start Range start + * @param end Range end + * @return Largest power of two in the range [start, end) or a zero if no power of two is found + */ + __device__ cudf::size_type inline constexpr smallest_power_of_two_in_range( + cudf::size_type start, cudf::size_type end) const noexcept + { + start--; + start |= start >> 1; + start |= start >> 2; + start |= start >> 4; + start |= start >> 8; + start |= start >> 16; + auto const result = start + 1; + return result < end ? result : 0; + } + + /** + * @brief Finds the largest power of two in the range (start, end]. If no power of two is found, + * returns a zero. + * + * @param start Range start + * @param end Range end + * @return Largest power of two in the range (start, end] or a zero if no power of two is found + */ + __device__ size_type inline constexpr largest_power_of_two_in_range(size_type start, + size_type end) const noexcept + { + auto constexpr nbits = cudf::detail::size_in_bits() - 1; + auto const result = size_type{1} << (nbits - cuda::std::countl_zero(end)); + return result > start ? result : 0; + } + + /** + * @brief Aligns a range boundary to the next power-of-two block + * + * @tparam Boundary Current boundary type (START or END) + * @param start Range start + * @param end Range end + * @return A pair of the tree level and block size + */ + template + __device__ auto inline constexpr align_range_boundary(cudf::size_type start, + cudf::size_type end) const noexcept + { + if constexpr (Boundary == boundary::START) { + if (start == 0 or is_power_of_two(start)) { + auto const block_size = + cuda::std::max(start & -start, largest_power_of_two_in_range(start, end)); + auto const tree_level = cuda::std::countr_zero(block_size); + return cuda::std::pair{tree_level, block_size}; + } else { + auto const tree_level = cuda::std::countr_zero(start); + return cuda::std::pair{tree_level, size_type{1} << tree_level}; + } + } else { + auto block_size = end & -end; + if (start > 0 and is_power_of_two(end)) { + auto const next_alignment = cuda::std::max(smallest_power_of_two_in_range(start, end), + largest_power_of_two_in_range(0, end - start)); + block_size = end - next_alignment; + } + return cuda::std::pair{cuda::std::countr_zero(block_size), block_size}; + } + } + + /** + * @brief Queries the Fenwick tree for the given boundary position, tree level and block size + * + * @tparam Boundary Current boundary type (START or END) + * @param boundary_pos Current boundary position + * @param tree_level Corresponding tree level to query + * @param block_size Alignment block size of the current boundary + * @return Boolean indicating if a `true` value is found in the fenwick tree + */ + template + __device__ bool inline constexpr query_fenwick_tree(cudf::size_type boundary_pos, + cudf::size_type tree_level, + cudf::size_type block_size) const noexcept + { + if constexpr (Boundary == boundary::START) { + auto const mask_index = boundary_pos >> tree_level; + return tree_level_ptrs[tree_level][mask_index]; + } else { + auto const mask_index = (boundary_pos - block_size) >> tree_level; + return tree_level_ptrs[tree_level][mask_index]; + } + } + + /** + * @brief Searches the Fenwick tree to find a `true` value in range [start, end) + * + * Algorithm: While `start` < `end`, align `start` UP and `end` DOWN to the next power-of-two + * searchable tree block. For the two aligned blocks, query the fenwick tree at corresponding + * levels for a `true` value (larger block first). If found, return. Else, move the boundaries + * to their alignments. + * + * @param range_idx Index of the range to search + * @return Boolean indicating if a `true` value is found in the range + */ + __device__ bool operator()(cudf::size_type range_idx) const noexcept + { + // Retrieve start and end for the current range [start, end) + size_type start = page_offsets[range_idx]; + size_type end = page_offsets[range_idx + 1]; + + // Return early if the range is empty or invalid + if (start >= end or range_idx >= num_ranges) { return false; } + + // Binary search decomposition loop + while (start < end) { + // Find the largest power-of-two block that aligns `start` up + auto const [start_tree_level, start_block_size] = + align_range_boundary(start, end); + + // Find the largest power-of-two block that aligns `end` down + auto const [end_tree_level, end_block_size] = align_range_boundary(start, end); + + // Check the larger block first to minimize the number of queries + if (start_block_size >= end_block_size) { + // Check the `start` side alignment block first + if (start + start_block_size <= end) { + if (query_fenwick_tree(start, start_tree_level, start_block_size)) { + return true; + } + start += start_block_size; + } + // Check the `end` side alignment block if it's still in range + if (end - end_block_size >= start) { + if (query_fenwick_tree(end, end_tree_level, end_block_size)) { + return true; + } + end -= end_block_size; + } + } else { + // Check the `end` side alignment block first + if (end - end_block_size >= start) { + if (query_fenwick_tree(end, end_tree_level, end_block_size)) { + return true; + } + end -= end_block_size; + } + // Check the `start` side alignment block if it's still in range + if (start + start_block_size <= end) { + if (query_fenwick_tree(start, start_tree_level, start_block_size)) { + return true; + } + start += start_block_size; + } + } + } + return false; + } +}; + +/** + * @brief Computes the offsets of the Fenwick tree levels (level 1 and higher) until the tree level + * block size becomes larger than the maximum page (search range) size + * + * @param level0_size Size of the zeroth tree level (the row mask) + * @param max_page_size Maximum page (search range) size + * @return Fenwick tree level offsets + */ +std::vector compute_fenwick_tree_level_offsets(cudf::size_type level0_size, + cudf::size_type max_page_size) +{ + std::vector tree_level_offsets; + tree_level_offsets.push_back(0); + + cudf::size_type current_level_size = cudf::util::div_rounding_up_safe(level0_size, 2); + cudf::size_type current_level = 1; + + while (current_level_size > 0) { + auto const block_size = 1 << current_level; + if (std::cmp_greater(block_size, max_page_size)) { break; } + tree_level_offsets.push_back(tree_level_offsets.back() + current_level_size); + current_level_size = + current_level_size == 1 ? 0 : cudf::util::div_rounding_up_safe(current_level_size, 2); + current_level++; + } + return tree_level_offsets; +} + +} // namespace + std::pair, cudf::detail::host_vector> compute_page_row_offsets_and_colchunk_page_offsets( std::span per_file_metadata, @@ -158,24 +426,83 @@ rmm::device_uvector compute_page_indices_async( return page_indices; } -std::vector compute_fenwick_tree_level_offsets(cudf::size_type level0_size, - cudf::size_type max_page_size) +cudf::column_view set_nulls_to_true(cudf::mutable_column_view const& row_mask, + cuda::stream_ref stream) { - std::vector tree_level_offsets; - tree_level_offsets.push_back(0); + if (row_mask.has_nulls()) { + auto const d_row_mask = cudf::column_device_view::create(row_mask, stream); + auto const iter = cudf::detail::make_null_replacement_iterator(*d_row_mask, true, true); + thrust::copy(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + iter, + iter + row_mask.size(), + row_mask.begin()); + } - cudf::size_type current_level_size = cudf::util::div_rounding_up_safe(level0_size, 2); - cudf::size_type current_level = 1; + return cudf::column_view{ + row_mask.type(), row_mask.size(), row_mask.head(), nullptr, 0, row_mask.offset()}; +} - while (current_level_size > 0) { - auto const block_size = 1 << current_level; - if (std::cmp_greater(block_size, max_page_size)) { break; } - tree_level_offsets.push_back(tree_level_offsets.back() + current_level_size); - current_level_size = - current_level_size == 1 ? 0 : cudf::util::div_rounding_up_safe(current_level_size, 2); - current_level++; +thrust::host_vector compute_row_range_selection_mask( + cudf::column_view const& row_mask, + std::span page_row_offsets, + cudf::size_type max_page_size, + cuda::stream_ref stream) +{ + // Need at least two offsets (or one range) to search the Fenwick tree + if (page_row_offsets.size() < 2) return thrust::host_vector{}; + + auto const total_rows = row_mask.size(); + // Return early if all rows are needed. + if (cudf::detail::all_of(row_mask.begin(), + row_mask.begin() + total_rows, + cuda::std::identity{}, + stream)) { + return thrust::host_vector{}; } - return tree_level_offsets; + + auto const mr = cudf::get_current_device_resource_ref(); + auto const tree_level_offsets = compute_fenwick_tree_level_offsets(total_rows, max_page_size); + auto const num_levels = static_cast(tree_level_offsets.size()); + auto tree_levels_data = rmm::device_uvector(tree_level_offsets.back(), stream, mr); + auto host_tree_level_ptrs = cudf::detail::make_pinned_vector_async(num_levels, stream); + host_tree_level_ptrs[0] = const_cast(row_mask.begin()); + std::for_each(cuda::counting_iterator{1}, + cuda::counting_iterator{num_levels}, + [&](auto const level_idx) { + host_tree_level_ptrs[level_idx] = + tree_levels_data.data() + tree_level_offsets[level_idx - 1]; + }); + auto tree_level_ptrs = cudf::detail::make_device_uvector_async(host_tree_level_ptrs, stream, mr); + + auto prev_level_size = total_rows; + std::for_each( + cuda::counting_iterator{0}, + cuda::counting_iterator{num_levels - 1}, + [&](auto const prev_level) { + auto const current_level_size = cudf::util::div_rounding_up_safe(prev_level_size, 2); + thrust::for_each(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + cuda::counting_iterator{0}, + cuda::counting_iterator{current_level_size}, + build_fenwick_tree_level_functor{ + tree_level_ptrs.data(), prev_level, prev_level_size, current_level_size}); + prev_level_size = current_level_size; + }); + + auto const num_ranges = static_cast(page_row_offsets.size() - 1); + auto device_results = rmm::device_uvector(num_ranges, stream, mr); + auto pinned_page_offsets = cudf::detail::make_pinned_vector(page_row_offsets, stream); + auto page_offsets = cudf::detail::make_device_uvector_async(pinned_page_offsets, stream, mr); + thrust::transform( + rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + cuda::counting_iterator{0}, + cuda::counting_iterator{num_ranges}, + device_results.begin(), + search_fenwick_tree_functor{tree_level_ptrs.data(), page_offsets.data(), num_ranges}); + + auto results = cudf::detail::make_pinned_vector_async(device_results, stream); + stream.sync(); + + return thrust::host_vector(results.begin(), results.end()); } } // namespace cudf::io::parquet::experimental::detail diff --git a/cpp/src/io/parquet/experimental/page_index_filter_utils.hpp b/cpp/src/io/parquet/experimental/page_index_filter_utils.hpp index 264510ff48fd..94d65264b48f 100644 --- a/cpp/src/io/parquet/experimental/page_index_filter_utils.hpp +++ b/cpp/src/io/parquet/experimental/page_index_filter_utils.hpp @@ -7,6 +7,7 @@ #include "io/parquet/reader_impl_helpers.hpp" +#include #include #include #include @@ -47,8 +48,7 @@ compute_page_row_offsets_and_colchunk_page_offsets( * @param per_file_metadata Span of parquet footer metadata * @param row_group_indices Span of input row group indices * @param schema_idx Column's schema index - * @return A pair of page row offsets and the size of the largest page in this - * column + * @return A pair of page row offsets and the size of the largest page in this column */ [[nodiscard]] std::pair, size_type> compute_page_row_offsets( cudf::host_span per_file_metadata, @@ -71,14 +71,29 @@ compute_page_row_offsets_and_colchunk_page_offsets( rmm::device_async_resource_ref mr); /** - * @brief Computes the offsets of the Fenwick tree levels (level 1 and higher) until the tree level - * block size becomes larger than the maximum page (search range) size + * @brief Sets nulls in the row mask to true and returns a non-nullable row mask view * - * @param level0_size Size of the zeroth tree level (the row mask) - * @param max_page_size Maximum page (search range) size - * @return Fenwick tree level offsets + * @param row_mask Mutable row mask column view + * @param stream CUDA stream used for device memory + * operations and kernel launches + * @return Non-nullable column view of the resolved row mask */ -[[nodiscard]] std::vector compute_fenwick_tree_level_offsets( - cudf::size_type level0_size, cudf::size_type max_page_size); +[[nodiscard]] cudf::column_view set_nulls_to_true(cudf::mutable_column_view const& row_mask, + cuda::stream_ref stream); + +/** + * @brief Computes a mask indicating which row ranges contain at least one selected row + * + * @param row_mask Boolean column indicating selected rows + * @param page_row_offsets Page row offsets defining the row ranges + * @param max_page_size Size of the largest page row range + * @param stream CUDA stream used for device memory operations and kernel launches + * @return Boolean vector with one entry for each consecutive row range + */ +[[nodiscard]] thrust::host_vector compute_row_range_selection_mask( + cudf::column_view const& row_mask, + std::span page_row_offsets, + cudf::size_type max_page_size, + cuda::stream_ref stream); } // namespace cudf::io::parquet::experimental::detail diff --git a/cpp/tests/io/experimental/hybrid_scan_filters_test.cpp b/cpp/tests/io/experimental/hybrid_scan_filters_test.cpp index aa57288646c5..8f19f8492c62 100644 --- a/cpp/tests/io/experimental/hybrid_scan_filters_test.cpp +++ b/cpp/tests/io/experimental/hybrid_scan_filters_test.cpp @@ -509,7 +509,8 @@ TEST_F(HybridScanFiltersTest, FilterRowGroupsWithComplexExpressions) auto input_row_group_indices = reader->all_row_groups(options); auto stats_filtered = reader->filter_row_groups_with_stats( input_row_group_indices, options, cudf::get_default_stream()); - EXPECT_EQ(stats_filtered.size(), 2); + auto const expected = std::vector{1, 2}; + EXPECT_EQ(stats_filtered, expected); } // Filter: NOT(NOT(col0 < 100) OR col0 > 150) @@ -907,10 +908,11 @@ TEST_F(HybridScanFiltersTest, OffsetIndexOnlyDataPageMask) auto const expected = cudf::apply_boolean_mask(written_table->view(), row_mask_view, stream, mr); CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected->view(), result.tbl->view()); - // Without offset index, data-page pruning falls back to decoding all pages. + // Without an offset index, data-page pruning derives page ranges from decoded page headers. for (auto& row_group : metadata.row_groups) { for (auto& column : row_group.columns) { column.offset_index.reset(); + ASSERT_FALSE(column.offset_index.has_value()); } } auto no_index_reader = cudf::io::parquet::experimental::hybrid_scan_reader(metadata, options); diff --git a/java/src/main/java/ai/rapids/cudf/HybridScanReader.java b/java/src/main/java/ai/rapids/cudf/HybridScanReader.java index b783aa86cff4..2165abd9a085 100644 --- a/java/src/main/java/ai/rapids/cudf/HybridScanReader.java +++ b/java/src/main/java/ai/rapids/cudf/HybridScanReader.java @@ -39,10 +39,14 @@ * chunked reader pipeline. * *

The filter and payload materialization paths accept a boolean that toggles - * page-level pruning: skips decode of pages the filter (or row mask) proves empty, in - * exchange for a per-page stats scan and a carried row-mask column. Enable when the - * workload prunes many pages; requires prior {@link #setupPageIndex(HostMemoryBuffer)} to - * prune filter column pages using page-level statistics. + * page-level pruning: skips decode of pages the filter (or row mask) proves empty. Enable when the + * workload prunes many pages. Pruning requirements for the two column materializations differ: + *

    + *
  • Filter columns seed the row mask from page-index statistics, so + * {@link #setupPageIndex(HostMemoryBuffer)} must have been called first.
    • + *
  • Payload columns only need page row boundaries. These come from the + * {@code OffsetIndex} when page index has been set up, otherwise from the decoded page headers as fallback.
    • + *
* *

The reader is created with no filter expression installed. Filter-related APIs * behave as though nothing has been filtered out unless a filter is first supplied via @@ -382,8 +386,7 @@ public FilterMaterializationResult materializeFilterColumns(int[] rowGroupIndice * returned by {@link #payloadColumnChunksByteRanges(int[])} * @param rowMask row mask (read-only) * @param usePageLevelPruning enable the data page mask to skip decode of pages the row - * mask proves empty; requires prior - * {@link #setupPageIndex(HostMemoryBuffer)} to avoid fall back path + * mask proves empty * @return the materialized payload column table */ public Table materializePayloadColumns(int[] rowGroupIndices, @@ -529,8 +532,7 @@ public ColumnVector takeFilterRowMask() { * @param rowGroupIndices row groups to read * @param rowMask row mask (read-only) * @param usePageLevelPruning enable the data page mask to skip decode of pages the row - * mask proves empty; requires prior - * {@link #setupPageIndex(HostMemoryBuffer)} to avoid fall back path + * mask proves empty * @param columnChunkData device buffers holding the payload column chunks, in the order * returned by {@link #payloadColumnChunksByteRanges(int[])} */ diff --git a/java/src/main/native/src/HybridScanReaderJniMaterialize.cpp b/java/src/main/native/src/HybridScanReaderJniMaterialize.cpp index f6f671ae4e7d..172eb5acb947 100644 --- a/java/src/main/native/src/HybridScanReaderJniMaterialize.cpp +++ b/java/src/main/native/src/HybridScanReaderJniMaterialize.cpp @@ -183,7 +183,7 @@ Java_ai_rapids_cudf_HybridScanReader_setupChunkingForFilterColumns(JNIEnv* env, wrapper->reader->setup_chunking_for_filter_columns(chunk_limit, pass_limit, holder.span(), - row_mask_col->view(), + row_mask_col->mutable_view(), mode, spans, wrapper->options, diff --git a/java/src/test/java/ai/rapids/cudf/HybridScanReaderTest.java b/java/src/test/java/ai/rapids/cudf/HybridScanReaderTest.java index 5eed87e1e0e0..463323e380fe 100644 --- a/java/src/test/java/ai/rapids/cudf/HybridScanReaderTest.java +++ b/java/src/test/java/ai/rapids/cudf/HybridScanReaderTest.java @@ -593,6 +593,42 @@ void testMaterializePayloadColumnsExactRowCount(@TempDir Path tmp) throws IOExce } } + /** + * Verifies materializePayloadColumns() prunes pages from page-header row counts when the + * file has no page index: zip_code > 100,000 keeps only the last of row group 1's three + * pages, so the payload must be exactly the 19,999 rows with ids 100,001–119,999. + */ + @Test + void testMaterializePayloadColumnsPagePruningWithoutPageIndex(@TempDir Path tmp) + throws IOException { + try (OpenReader open = + OpenReader.multiPage(tmp).withFilter("zip_code", BinaryOperator.GREATER, 100000)) { + HybridScanReader reader = open.reader; + assertEquals(0L, reader.pageIndexByteRange().size(), + "Fixture must have no page index so the header-derived fallback is exercised"); + int[] survived = reader.filterRowGroupsWithStats(reader.allRowGroups()); + assertArrayEquals(new int[]{1}, survived, + "Group 0 (zip_code 0–59,999) cannot satisfy zip_code > 100,000"); + DeviceMemoryBuffer[] filterCols = copyRangesToDevice( + open.file, reader.filterColumnChunksByteRanges(survived)); + DeviceMemoryBuffer[] payloadCols = copyRangesToDevice( + open.file, reader.payloadColumnChunksByteRanges(survived)); + try (HybridScanReader.FilterMaterializationResult fr = + reader.materializeFilterColumns(survived, filterCols, false); + Table payload = reader.materializePayloadColumns(survived, payloadCols, + fr.rowMask(), true); + ColumnVector expectedIds = ColumnVector.fromInts( + IntStream.rangeClosed(100001, 119999).toArray())) { + assertEquals(2, payload.getNumberOfColumns(), "payload table contains id + num_units"); + assertEquals(19999L, payload.getRowCount()); + AssertUtils.assertColumnsAreEqual(expectedIds, payload.getColumn(0), "id"); + } finally { + closeAll(filterCols); + closeAll(payloadCols); + } + } + } + // -------------------------------------------------------------------- // Tests: materializeAllColumns() // -------------------------------------------------------------------- @@ -825,6 +861,49 @@ void testMaterializePayloadColumnsChunkExactTotal(@TempDir Path tmp) throws IOEx } } + /** + * Verifies the chunked payload pipeline drains the same 19,999 rows when page pruning + * falls back to page-header row counts on a file with no page index. + */ + @Test + void testMaterializePayloadColumnsChunkPagePruningWithoutPageIndex(@TempDir Path tmp) + throws IOException { + try (OpenReader open = + OpenReader.multiPage(tmp).withFilter("zip_code", BinaryOperator.GREATER, 100000)) { + HybridScanReader reader = open.reader; + assertEquals(0L, reader.pageIndexByteRange().size(), + "Fixture must have no page index so the header-derived fallback is exercised"); + int[] survived = reader.filterRowGroupsWithStats(reader.allRowGroups()); + DeviceMemoryBuffer[] filterCols = copyRangesToDevice( + open.file, reader.filterColumnChunksByteRanges(survived)); + DeviceMemoryBuffer[] payloadCols = copyRangesToDevice( + open.file, reader.payloadColumnChunksByteRanges(survived)); + try { + reader.setupChunkingForFilterColumns(0L, 0L, survived, false, filterCols); + while (reader.hasNextTableChunk()) { + reader.materializeFilterColumnsChunk().close(); + } + try (ColumnVector rowMask = reader.takeFilterRowMask()) { + assertEquals(60000L, rowMask.getRowCount(), "Mask spans row group 1"); + assertEquals(19999L, countTrue(rowMask), "zip_code 100,001–119,999 survive"); + reader.setupChunkingForPayloadColumns(0L, 0L, survived, rowMask, true, payloadCols); + long total = 0; + while (reader.hasNextTableChunk()) { + try (Table chunk = reader.materializePayloadColumnsChunk(rowMask)) { + assertEquals(2, chunk.getNumberOfColumns()); + total += chunk.getRowCount(); + } + } + assertEquals(19999L, total, + "Header-derived page pruning must not drop or duplicate selected rows"); + } + } finally { + closeAll(filterCols); + closeAll(payloadCols); + } + } + } + // -------------------------------------------------------------------- // Tests: setupChunkingForAllColumns() / materializeAllColumnsChunk() // @@ -1212,9 +1291,19 @@ static OpenReader pageIndex(Path tmp) throws IOException { return openFromFile(pq, DEFAULT_COLS); } + /** A single 100-row group, small enough that every column chunk holds one data page. */ static OpenReader rowGroupStats(Path tmp) throws IOException { File pq = tmp.resolve("fixture.parquet").toFile(); - writeRowGroupStatsParquet(pq); + writeNoPageIndexParquet(pq, 100, 1, + ParquetWriterOptions.StatisticsFrequency.ROWGROUP); + return openFromFile(pq, DEFAULT_COLS); + } + + /** Two 60,000-row groups, so each column chunk spans 3 data pages (20,000 rows each). */ + static OpenReader multiPage(Path tmp) throws IOException { + File pq = tmp.resolve("fixture.parquet").toFile(); + writeNoPageIndexParquet(pq, 60_000, 2, + ParquetWriterOptions.StatisticsFrequency.PAGE); return openFromFile(pq, DEFAULT_COLS); } @@ -1323,28 +1412,34 @@ private static int writeFixtureParquet(File path) { } /** - * Writes a small Parquet file with {@code ROWGROUP}-level statistics: row-group min/max - * are recorded but no page index (no {@code ColumnIndex}/{@code OffsetIndex}) is emitted. - * Includes a low-cardinality {@code num_units} column ({1, 2, 3} cycle) so the writer's - * ADAPTIVE dictionary policy emits a dictionary; this lets tests exercise the - * "no page index, dict exists" path (see - * {@link #testSecondaryFiltersByteRangesEmptyForRowGroupStats}). + * Writes a Parquet file with no page index; only {@code COLUMN} statistics emit one. Both + * {@code id} and {@code zip_code} hold the globally sequential row index, and + * {@code num_units} cycles over {1, 2, 3}, low-cardinality enough that the ADAPTIVE + * dictionary policy emits a dictionary (see + * {@link #testSecondaryFiltersByteRangesEmptyForRowGroupStats}). The writer caps a page at + * 20,000 rows, so exceed that per group for chunks spanning several pages. */ - private static void writeRowGroupStatsParquet(File path) { - int rows = 100; + private static void writeNoPageIndexParquet(File path, int rowsPerGroup, int numGroups, + ParquetWriterOptions.StatisticsFrequency stats) { ParquetWriterOptions opts = ParquetWriterOptions.builder() .withNonNullableColumns("id", "zip_code", "num_units") - .withRowGroupSizeRows(rows) - .withStatisticsFrequency(ParquetWriterOptions.StatisticsFrequency.ROWGROUP) + .withRowGroupSizeRows(rowsPerGroup) + .withStatisticsFrequency(stats) .build(); - try (TableWriter writer = Table.writeParquetChunked(opts, path); - ColumnVector id = ColumnVector.fromInts(IntStream.range(0, rows).toArray()); - ColumnVector zipCode = ColumnVector.fromInts( - IntStream.range(0, rows).map(i -> 10000 + i).toArray()); - ColumnVector numUnits = ColumnVector.fromInts( - IntStream.range(0, rows).map(i -> 1 + (i % 3)).toArray()); - Table t = new Table(id, zipCode, numUnits)) { - writer.write(t); + try (TableWriter writer = Table.writeParquetChunked(opts, path)) { + for (int g = 0; g < numGroups; g++) { + int start = g * rowsPerGroup; + try (ColumnVector id = ColumnVector.fromInts( + IntStream.range(start, start + rowsPerGroup).toArray()); + ColumnVector zipCode = ColumnVector.fromInts( + IntStream.range(start, start + rowsPerGroup).toArray()); + ColumnVector numUnits = ColumnVector.fromInts( + IntStream.range(start, start + rowsPerGroup) + .map(i -> 1 + (i % 3)).toArray()); + Table t = new Table(id, zipCode, numUnits)) { + writer.write(t); + } + } } } diff --git a/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pyx b/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pyx index 4cfc0214f4d8..70507ca66a4f 100644 --- a/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pyx +++ b/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pyx @@ -830,7 +830,7 @@ cdef class HybridScanReader: row_group_indices : list[int] Input row group indices row_mask : Column - Boolean column indicating surviving rows + Mutable boolean column indicating surviving rows mask_data_pages : UseDataPageMask Whether to use a data page mask column_chunk_data : Sequence @@ -853,7 +853,7 @@ cdef class HybridScanReader: # keep reference to avoid use-after-free of device spans self._filter_chunk_data = column_chunk_data - cdef column_view mask_view = row_mask.view() + cdef mutable_column_view mask_view = row_mask.mutable_view() with nogil: self.c_obj.get()[0].setup_chunking_for_filter_columns( chunk_read_limit, diff --git a/python/pylibcudf/pylibcudf/libcudf/io/hybrid_scan.pxd b/python/pylibcudf/pylibcudf/libcudf/io/hybrid_scan.pxd index 7a5aec269a56..04d84fb68a33 100644 --- a/python/pylibcudf/pylibcudf/libcudf/io/hybrid_scan.pxd +++ b/python/pylibcudf/pylibcudf/libcudf/io/hybrid_scan.pxd @@ -161,7 +161,7 @@ cdef extern from "cudf/io/experimental/hybrid_scan.hpp" \ size_t chunk_read_limit, size_t pass_read_limit, std_span[const_size_type] row_group_indices, - const column_view& row_mask, + const mutable_column_view& row_mask, use_data_page_mask mask_data_pages, std_span[const_device_span_const_uint8_t] column_chunk_data, const parquet_reader_options& options, diff --git a/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py b/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py index d1e4c1996985..b7a835da1bad 100644 --- a/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py +++ b/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py @@ -430,6 +430,73 @@ def test_hybrid_scan_materialize_columns( assert expected_arrow.equals(hybrid_arrow) +def test_hybrid_scan_payload_page_mask_without_page_index( + simple_parquet_bytes: bytes, + simple_hybrid_scan_reader: HybridScanReader, + simple_parquet_options: plc.io.parquet.ParquetReaderOptions, + simple_parquet_table: pa.Table, + num_rows: int, +) -> None: + """Test payload page pruning without a page index set up on the reader.""" + reader = simple_hybrid_scan_reader + row_groups = reader.all_row_groups(simple_parquet_options) + + # Keep the first half of the rows so the trailing data pages get pruned. + num_selected = num_rows // 2 + row_mask = plc.Column.from_arrow( + pa.array([i < num_selected for i in range(num_rows)], type=pa.bool_()) + ) + + payload_data = [ + plc.gpumemoryview( + rmm.DeviceBuffer.to_device( + simple_parquet_bytes[r.offset : r.offset + r.size], + plc.utils._get_stream(), + ) + ) + for r in reader.payload_column_chunks_byte_ranges( + row_groups, simple_parquet_options + ) + ] + synchronize_stream() + + # Chunks can disagree on field nullability, so compare row values only. + def to_rows(tbl: plc.Table) -> list: + return ( + tbl.to_arrow() + .rename_columns(simple_parquet_table.column_names) + .to_pylist() + ) + + expected_rows = simple_parquet_table.slice(0, num_selected).to_pylist() + + payload_result = reader.materialize_payload_columns( + row_groups, + payload_data, + row_mask, + UseDataPageMask.YES, + simple_parquet_options, + ) + synchronize_stream() + assert to_rows(payload_result.tbl) == expected_rows + + reader.setup_chunking_for_payload_columns( + 256, + 0, + row_groups, + row_mask, + UseDataPageMask.YES, + payload_data, + simple_parquet_options, + ) + chunked_rows = [] + while reader.has_next_table_chunk(): + chunk = reader.materialize_payload_columns_chunk(row_mask) + chunked_rows.extend(to_rows(chunk.tbl)) + synchronize_stream() + assert chunked_rows == expected_rows + + @pytest.mark.parametrize("stream", [None, Stream()]) def test_hybrid_scan_single_step_materialize( simple_parquet_bytes: bytes, From 1fd4b66c33cab555155d0972b2e974312c429359 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Tue, 25 Aug 2026 22:28:27 +0000 Subject: [PATCH 02/21] Address comments from vuule --- .../io/parquet/experimental/hybrid_scan_impl.cpp | 9 ++++++--- .../io/parquet/experimental/page_index_filter.cu | 11 +++-------- .../experimental/page_index_filter_utils.cu | 16 +++++++--------- .../experimental/page_index_filter_utils.hpp | 10 ++++++++++ 4 files changed, 26 insertions(+), 20 deletions(-) diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp b/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp index 6a67213551f0..f73a37efd82e 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp +++ b/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp @@ -1514,9 +1514,12 @@ void hybrid_scan_reader_impl::set_pass_page_mask(std::span data_page thrust::host_vector hybrid_scan_reader_impl::compute_data_page_mask_with_page_headers() { - auto& pass = *_pass_itm_data; - pass.pages.device_to_host_async(_stream); - _stream.sync(); + auto const& pass = *_pass_itm_data; + + // Return an empty vector if all rows are required + if (parquet::detail::are_all_rows_retained(_row_mask, _stream)) { + return thrust::host_vector(0); + } std::vector page_row_offsets; page_row_offsets.reserve(pass.pages.size() * 2); diff --git a/cpp/src/io/parquet/experimental/page_index_filter.cu b/cpp/src/io/parquet/experimental/page_index_filter.cu index ad57c3417d61..18fd9cebbf29 100644 --- a/cpp/src/io/parquet/experimental/page_index_filter.cu +++ b/cpp/src/io/parquet/experimental/page_index_filter.cu @@ -769,12 +769,7 @@ thrust::host_vector aggregate_reader_metadata::compute_data_page_mask( row_mask.null_count() == 0, "Row mask must not contain nulls", std::invalid_argument); // Return an empty vector if all rows are required - if (cudf::detail::all_of(row_mask.begin(), - row_mask.begin() + total_rows, - cuda::std::identity{}, - stream)) { - return thrust::host_vector(0); - } + if (are_all_rows_selected(row_mask, stream)) { return thrust::host_vector{}; } // Collect column schema indices from the input columns. auto column_schema_indices = std::vector(input_columns.size()); @@ -788,8 +783,8 @@ thrust::host_vector aggregate_reader_metadata::compute_data_page_mask( page_index_presence(row_group_indices, column_schema_indices).second; if (not has_offset_index) { CUDF_LOG_WARN( - "Encountered missing Parquet offset index for one or more output columns. Skipping page " - "pruning."); + "Encountered missing Parquet offset index for one or more output columns. Skipping " + "page-index based pruning."); return thrust::host_vector(0); } diff --git a/cpp/src/io/parquet/experimental/page_index_filter_utils.cu b/cpp/src/io/parquet/experimental/page_index_filter_utils.cu index 89332b889359..9df82e836c10 100644 --- a/cpp/src/io/parquet/experimental/page_index_filter_utils.cu +++ b/cpp/src/io/parquet/experimental/page_index_filter_utils.cu @@ -442,6 +442,12 @@ cudf::column_view set_nulls_to_true(cudf::mutable_column_view const& row_mask, row_mask.type(), row_mask.size(), row_mask.head(), nullptr, 0, row_mask.offset()}; } +bool are_all_rows_retained(cudf::column_view const& row_mask, cuda::stream_ref stream) +{ + return cudf::detail::all_of( + row_mask.begin(), row_mask.end(), cuda::std::identity{}, stream); +} + thrust::host_vector compute_row_range_selection_mask( cudf::column_view const& row_mask, std::span page_row_offsets, @@ -451,15 +457,7 @@ thrust::host_vector compute_row_range_selection_mask( // Need at least two offsets (or one range) to search the Fenwick tree if (page_row_offsets.size() < 2) return thrust::host_vector{}; - auto const total_rows = row_mask.size(); - // Return early if all rows are needed. - if (cudf::detail::all_of(row_mask.begin(), - row_mask.begin() + total_rows, - cuda::std::identity{}, - stream)) { - return thrust::host_vector{}; - } - + auto const total_rows = row_mask.size(); auto const mr = cudf::get_current_device_resource_ref(); auto const tree_level_offsets = compute_fenwick_tree_level_offsets(total_rows, max_page_size); auto const num_levels = static_cast(tree_level_offsets.size()); diff --git a/cpp/src/io/parquet/experimental/page_index_filter_utils.hpp b/cpp/src/io/parquet/experimental/page_index_filter_utils.hpp index 94d65264b48f..bc7b5eecfa5e 100644 --- a/cpp/src/io/parquet/experimental/page_index_filter_utils.hpp +++ b/cpp/src/io/parquet/experimental/page_index_filter_utils.hpp @@ -81,6 +81,16 @@ compute_page_row_offsets_and_colchunk_page_offsets( [[nodiscard]] cudf::column_view set_nulls_to_true(cudf::mutable_column_view const& row_mask, cuda::stream_ref stream); +/** + * @brief Checks whether every row is reatained by the boolean row mask + * + * @param retention_mask Boolean column indicating retained rows + * @param stream CUDA stream used for device memory operations and kernel launches + * @return Boolean indicating whether every row is retained + */ +[[nodiscard]] bool are_all_rows_retained(cudf::column_view const& retention_mask, + cuda::stream_ref stream); + /** * @brief Computes a mask indicating which row ranges contain at least one selected row * From ccc96b6dcf0c7132e6db6ad86b0d9feafb59d1f7 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Tue, 25 Aug 2026 22:35:10 +0000 Subject: [PATCH 03/21] Simplify java docs --- .../main/java/ai/rapids/cudf/HybridScanReader.java | 13 +++++++------ 1 file changed, 7 insertions(+), 6 deletions(-) diff --git a/java/src/main/java/ai/rapids/cudf/HybridScanReader.java b/java/src/main/java/ai/rapids/cudf/HybridScanReader.java index 2165abd9a085..ef8a31b6633c 100644 --- a/java/src/main/java/ai/rapids/cudf/HybridScanReader.java +++ b/java/src/main/java/ai/rapids/cudf/HybridScanReader.java @@ -42,10 +42,11 @@ * page-level pruning: skips decode of pages the filter (or row mask) proves empty. Enable when the * workload prunes many pages. Pruning requirements for the two column materializations differ: *

    - *
  • Filter columns seed the row mask from page-index statistics, so - * {@link #setupPageIndex(HostMemoryBuffer)} must have been called first.
    • - *
  • Payload columns only need page row boundaries. These come from the - * {@code OffsetIndex} when page index has been set up, otherwise from the decoded page headers as fallback.
    • + *
  • With page-level pruning enabled, filter columns require + * {@link #setupPageIndex(HostMemoryBuffer)} to seed the row mask from page-index + * statistics.
    • + *
  • Once a row mask exists, filter and payload columns use page row boundaries from the + * {@code OffsetIndex} when available, otherwise from decoded page headers.
    • *
* *

The reader is created with no filter expression installed. Filter-related APIs @@ -201,8 +202,8 @@ public ByteRange pageIndexByteRange() { /** * Materialize the {@code ColumnIndex} / {@code OffsetIndex} structs (collectively, the page - * index) from the supplied bytes. Required before any filter or payload column materialization - * call with {@code usePageLevelPruning == true}. + * index) from the supplied bytes. Required before any filter column materialization with + * {@code usePageLevelPruning == true}. * * @param pageIndexBuffer host-resident page index bytes */ From 2da7634495ea705f6768f1d013e7bbbacd871d9f Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Fri, 28 Aug 2026 00:11:20 +0000 Subject: [PATCH 04/21] Address comments from @pmattione-nvidia --- .../parquet/experimental/hybrid_scan_impl.cpp | 14 ++-- .../parquet/experimental/page_index_filter.cu | 2 +- .../experimental/hybrid_scan_filters_test.cpp | 49 ++++++++++++ .../ai/rapids/cudf/HybridScanReaderTest.java | 75 +++++++++++++++++++ 4 files changed, 131 insertions(+), 9 deletions(-) diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp b/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp index 3ccbee22ba62..0dacccfc0e56 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp +++ b/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp @@ -8,8 +8,8 @@ #include "cudf/io/text/byte_range_info.hpp" #include "hybrid_scan_helpers.hpp" #include "io/parquet/reader_impl_chunking_utils.cuh" -#include "page_index_filter_utils.hpp" #include "io/parquet/synthetic_column_helpers.hpp" +#include "page_index_filter_utils.hpp" #include #include @@ -1388,9 +1388,9 @@ table_with_metadata hybrid_scan_reader_impl::finalize_output( // Prepend the source and row index columns to filter columns only if (read_columns_mode == read_columns_mode::FILTER_COLUMNS) { if (_options.prepend_row_index_column) { - out_columns.emplace( - out_columns.begin(), - synthesize_row_index_column(_file_itm_data.row_groups, read_info, _stream, _mr)); + out_columns.emplace(out_columns.begin(), + parquet::detail::synthesize_row_index_column( + _file_itm_data.row_groups, read_info, _stream, _mr)); out_metadata.schema_info.emplace(out_metadata.schema_info.begin(), column_name_info{.name = "row_index", .is_nullable = false}); } @@ -1519,9 +1519,7 @@ thrust::host_vector hybrid_scan_reader_impl::compute_data_page_mask_with_p auto const& pass = *_pass_itm_data; // Return an empty vector if all rows are required - if (parquet::detail::are_all_rows_retained(_row_mask, _stream)) { - return thrust::host_vector(0); - } + if (are_all_rows_retained(_row_mask, _stream)) { return thrust::host_vector(0); } std::vector page_row_offsets; page_row_offsets.reserve(pass.pages.size() * 2); @@ -1551,7 +1549,7 @@ thrust::host_vector hybrid_scan_reader_impl::compute_data_page_mask_with_p max_page_size = std::max(max_page_size, page_end - page_start); // Starting a new column chunk. Push page start row - if (previous_chunk_idx == -1 or page.chunk_idx != previous_chunk_idx) { + if (page.chunk_idx != previous_chunk_idx) { page_row_offsets.push_back(page_start); previous_chunk_idx = page.chunk_idx; } diff --git a/cpp/src/io/parquet/experimental/page_index_filter.cu b/cpp/src/io/parquet/experimental/page_index_filter.cu index 18fd9cebbf29..6ec51293aa50 100644 --- a/cpp/src/io/parquet/experimental/page_index_filter.cu +++ b/cpp/src/io/parquet/experimental/page_index_filter.cu @@ -769,7 +769,7 @@ thrust::host_vector aggregate_reader_metadata::compute_data_page_mask( row_mask.null_count() == 0, "Row mask must not contain nulls", std::invalid_argument); // Return an empty vector if all rows are required - if (are_all_rows_selected(row_mask, stream)) { return thrust::host_vector{}; } + if (are_all_rows_retained(row_mask, stream)) { return thrust::host_vector{}; } // Collect column schema indices from the input columns. auto column_schema_indices = std::vector(input_columns.size()); diff --git a/cpp/tests/io/experimental/hybrid_scan_filters_test.cpp b/cpp/tests/io/experimental/hybrid_scan_filters_test.cpp index 8f19f8492c62..5b7f1cf9b1e2 100644 --- a/cpp/tests/io/experimental/hybrid_scan_filters_test.cpp +++ b/cpp/tests/io/experimental/hybrid_scan_filters_test.cpp @@ -933,6 +933,55 @@ TEST_F(HybridScanFiltersTest, OffsetIndexOnlyDataPageMask) CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected->view(), no_index_result.tbl->view()); } +TEST_F(HybridScanFiltersTest, NoOffsetIndexListColumns) +{ + // List pages cannot safely be mapped to rows from page headers alone, because a leaf page may + // begin in the middle of a list row. Verify the header-derived fallback retains correct output. + using T = uint32_t; + std::mt19937 gen(0xc0c0a); + auto list_col = make_list_str_column(gen, false, false); + auto scalar_col = testdata::ascending(); + auto list_table = cudf::table_view{{scalar_col, *list_col}}; + auto list_buffer = std::vector{}; + auto list_writer = + cudf::io::parquet_writer_options::builder(cudf::io::sink_info{&list_buffer}, list_table) + .row_group_size_rows(num_ordered_rows) + .max_page_size_rows(page_size_for_ordered_tests) + .stats_level(cudf::io::statistics_freq::STATISTICS_COLUMN) + .build(); + cudf::io::write_parquet(list_writer); + + auto const list_datasource = cudf::io::datasource::create(cudf::host_span( + reinterpret_cast(list_buffer.data()), list_buffer.size())); + auto const list_footer = cudf::io::parquet::fetch_footer_to_host(*list_datasource); + auto options = cudf::io::parquet_reader_options::builder().build(); + auto list_reader = cudf::io::parquet::experimental::hybrid_scan_reader(*list_footer, options); + auto const list_row_groups = list_reader.all_row_groups(options); + auto const list_row_mask_values = cudf::detail::make_counting_transform_iterator( + 0, [](auto const row) { return std::cmp_less(row, num_ordered_rows / 2); }); + auto list_row_mask = cudf::test::fixed_width_column_wrapper( + list_row_mask_values, list_row_mask_values + num_ordered_rows); + auto const list_row_mask_view = static_cast(list_row_mask); + auto const stream = cudf::get_default_stream(); + auto const mr = cudf::get_current_device_resource_ref(); + auto const list_byte_ranges = + list_reader.payload_column_chunks_byte_ranges(list_row_groups, options); + auto [list_buffers, list_data, list_tasks] = cudf::io::parquet::fetch_byte_ranges_to_device_async( + *list_datasource, list_byte_ranges, stream, mr); + list_tasks.get(); + + auto const list_result = list_reader.materialize_payload_columns( + list_row_groups, + list_data, + list_row_mask_view, + cudf::io::parquet::experimental::use_data_page_mask::YES, + options, + stream, + mr); + auto const list_expected = cudf::apply_boolean_mask(list_table, list_row_mask_view, stream, mr); + CUDF_TEST_EXPECT_TABLES_EQUIVALENT(list_expected->view(), list_result.tbl->view()); +} + template struct TimestampPageFiltering : public HybridScanFiltersTest {}; diff --git a/java/src/test/java/ai/rapids/cudf/HybridScanReaderTest.java b/java/src/test/java/ai/rapids/cudf/HybridScanReaderTest.java index 463323e380fe..8db6c6925522 100644 --- a/java/src/test/java/ai/rapids/cudf/HybridScanReaderTest.java +++ b/java/src/test/java/ai/rapids/cudf/HybridScanReaderTest.java @@ -38,6 +38,8 @@ public class HybridScanReaderTest extends CudfTestBase { private static final String[] DEFAULT_COLS = {"id", "zip_code", "num_units"}; + private static final HostColumnVector.ListType LIST_OF_INTS = + new HostColumnVector.ListType(true, new HostColumnVector.BasicType(true, DType.INT32)); private static final int[] ALL_ROW_GROUPS = {0, 1, 2}; // -------------------------------------------------------------------- @@ -904,6 +906,49 @@ void testMaterializePayloadColumnsChunkPagePruningWithoutPageIndex(@TempDir Path } } + /** + * Verifies that list payload columns remain readable when header-derived page pruning is used. + * List leaf pages may start mid-row, so they must not be pruned without an offset index. + */ + @Test + void testMaterializePayloadColumnsChunkListPagePruningWithoutPageIndex(@TempDir Path tmp) + throws IOException { + try (OpenReader open = OpenReader.multiPageWithList(tmp) + .withFilter("zip_code", BinaryOperator.GREATER, 30000)) { + HybridScanReader reader = open.reader; + assertEquals(0L, reader.pageIndexByteRange().size(), + "Fixture must have no page index so the header-derived fallback is exercised"); + int[] survived = reader.filterRowGroupsWithStats(reader.allRowGroups()); + DeviceMemoryBuffer[] filterCols = copyRangesToDevice( + open.file, reader.filterColumnChunksByteRanges(survived)); + DeviceMemoryBuffer[] payloadCols = copyRangesToDevice( + open.file, reader.payloadColumnChunksByteRanges(survived)); + try { + reader.setupChunkingForFilterColumns(0L, 0L, survived, false, filterCols); + while (reader.hasNextTableChunk()) { + reader.materializeFilterColumnsChunk().close(); + } + try (ColumnVector rowMask = reader.takeFilterRowMask()) { + assertEquals(60000L, rowMask.getRowCount(), "Mask spans the row group"); + assertEquals(29999L, countTrue(rowMask), "zip_code 30,001–59,999 survive"); + reader.setupChunkingForPayloadColumns(0L, 0L, survived, rowMask, true, payloadCols); + long total = 0; + while (reader.hasNextTableChunk()) { + try (Table chunk = reader.materializePayloadColumnsChunk(rowMask)) { + assertEquals(2, chunk.getNumberOfColumns()); + total += chunk.getRowCount(); + } + } + assertEquals(29999L, total, + "Header-derived pruning must retain all selected rows with list payloads"); + } + } finally { + closeAll(filterCols); + closeAll(payloadCols); + } + } + } + // -------------------------------------------------------------------- // Tests: setupChunkingForAllColumns() / materializeAllColumnsChunk() // @@ -1307,6 +1352,13 @@ static OpenReader multiPage(Path tmp) throws IOException { return openFromFile(pq, DEFAULT_COLS); } + /** A 60,000-row group with list payloads spanning multiple leaf pages. */ + static OpenReader multiPageWithList(Path tmp) throws IOException { + File pq = tmp.resolve("fixture.parquet").toFile(); + writeNoPageIndexListParquet(pq); + return openFromFile(pq, "id", "zip_code", "list_values"); + } + private static OpenReader openFromFile(File pq, String[] cols) throws IOException { HostMemoryBuffer file = readFileToHostBuffer(pq); HostMemoryBuffer footer = null; @@ -1443,6 +1495,29 @@ private static void writeNoPageIndexParquet(File path, int rowsPerGroup, int num } } + /** + * Writes one 60,000-row group with a list payload column. Each row holds three values, so + * leaf-page boundaries can fall within a logical row. + */ + private static void writeNoPageIndexListParquet(File path) { + int rows = 60_000; + ParquetWriterOptions opts = ParquetWriterOptions.builder() + .withNonNullableColumns("id", "zip_code", "list_values") + .withRowGroupSizeRows(rows) + .withStatisticsFrequency(ParquetWriterOptions.StatisticsFrequency.PAGE) + .build(); + Object[] listRows = IntStream.range(0, rows) + .mapToObj(i -> java.util.List.of(i * 3, i * 3 + 1, i * 3 + 2)) + .toArray(); + try (ColumnVector id = ColumnVector.fromInts(IntStream.range(0, rows).toArray()); + ColumnVector zipCode = ColumnVector.fromInts(IntStream.range(0, rows).toArray()); + ColumnVector listValues = ColumnVector.fromLists(LIST_OF_INTS, listRows); + Table table = new Table(id, zipCode, listValues); + TableWriter writer = Table.writeParquetChunked(opts, path)) { + writer.write(table); + } + } + /** * Writes a 3-row-group Parquet file with {@code COLUMN}-level statistics, guaranteeing * a non-empty page index. Each group gets a non-overlapping {@code zip_code} range From aa7be76b749a174d4459b2cb58a6df943968fe86 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Fri, 28 Aug 2026 02:56:47 +0000 Subject: [PATCH 05/21] minor fix --- .../hybrid_scan_io/hybrid_scan_composer.cpp | 2 +- .../ai/rapids/cudf/HybridScanReaderTest.java | 18 +++++++++++------- 2 files changed, 12 insertions(+), 8 deletions(-) diff --git a/cpp/examples/hybrid_scan_io/hybrid_scan_composer.cpp b/cpp/examples/hybrid_scan_io/hybrid_scan_composer.cpp index bd5e4765f5dc..47a2f36af349 100644 --- a/cpp/examples/hybrid_scan_io/hybrid_scan_composer.cpp +++ b/cpp/examples/hybrid_scan_io/hybrid_scan_composer.cpp @@ -465,7 +465,7 @@ template std::unique_ptr hybrid_scan( std::optional, std::unordered_set const&, bool, - rmm::cuda_stream_view, + cuda::stream_ref, rmm::device_async_resource_ref); template std::unique_ptr hybrid_scan( diff --git a/java/src/test/java/ai/rapids/cudf/HybridScanReaderTest.java b/java/src/test/java/ai/rapids/cudf/HybridScanReaderTest.java index 8db6c6925522..49a11c11ff41 100644 --- a/java/src/test/java/ai/rapids/cudf/HybridScanReaderTest.java +++ b/java/src/test/java/ai/rapids/cudf/HybridScanReaderTest.java @@ -28,6 +28,8 @@ import java.io.IOException; import java.nio.file.Files; import java.nio.file.Path; +import java.util.Arrays; +import java.util.List; import java.util.function.Consumer; import java.util.stream.IntStream; import java.util.stream.Stream; @@ -610,7 +612,7 @@ void testMaterializePayloadColumnsPagePruningWithoutPageIndex(@TempDir Path tmp) "Fixture must have no page index so the header-derived fallback is exercised"); int[] survived = reader.filterRowGroupsWithStats(reader.allRowGroups()); assertArrayEquals(new int[]{1}, survived, - "Group 0 (zip_code 0–59,999) cannot satisfy zip_code > 100,000"); + "Group 0 (zip_code 0-59,999) cannot satisfy zip_code > 100,000"); DeviceMemoryBuffer[] filterCols = copyRangesToDevice( open.file, reader.filterColumnChunksByteRanges(survived)); DeviceMemoryBuffer[] payloadCols = copyRangesToDevice( @@ -887,7 +889,7 @@ void testMaterializePayloadColumnsChunkPagePruningWithoutPageIndex(@TempDir Path } try (ColumnVector rowMask = reader.takeFilterRowMask()) { assertEquals(60000L, rowMask.getRowCount(), "Mask spans row group 1"); - assertEquals(19999L, countTrue(rowMask), "zip_code 100,001–119,999 survive"); + assertEquals(19999L, countTrue(rowMask), "zip_code 100,001-119,999 survive"); reader.setupChunkingForPayloadColumns(0L, 0L, survived, rowMask, true, payloadCols); long total = 0; while (reader.hasNextTableChunk()) { @@ -930,7 +932,7 @@ void testMaterializePayloadColumnsChunkListPagePruningWithoutPageIndex(@TempDir } try (ColumnVector rowMask = reader.takeFilterRowMask()) { assertEquals(60000L, rowMask.getRowCount(), "Mask spans the row group"); - assertEquals(29999L, countTrue(rowMask), "zip_code 30,001–59,999 survive"); + assertEquals(29999L, countTrue(rowMask), "zip_code 30,001-59,999 survive"); reader.setupChunkingForPayloadColumns(0L, 0L, survived, rowMask, true, payloadCols); long total = 0; while (reader.hasNextTableChunk()) { @@ -1359,7 +1361,7 @@ static OpenReader multiPageWithList(Path tmp) throws IOException { return openFromFile(pq, "id", "zip_code", "list_values"); } - private static OpenReader openFromFile(File pq, String[] cols) throws IOException { + private static OpenReader openFromFile(File pq, String... cols) throws IOException { HostMemoryBuffer file = readFileToHostBuffer(pq); HostMemoryBuffer footer = null; HybridScanReader reader = null; @@ -1499,6 +1501,7 @@ private static void writeNoPageIndexParquet(File path, int rowsPerGroup, int num * Writes one 60,000-row group with a list payload column. Each row holds three values, so * leaf-page boundaries can fall within a logical row. */ + @SuppressWarnings("unchecked") private static void writeNoPageIndexListParquet(File path) { int rows = 60_000; ParquetWriterOptions opts = ParquetWriterOptions.builder() @@ -1506,9 +1509,10 @@ private static void writeNoPageIndexListParquet(File path) { .withRowGroupSizeRows(rows) .withStatisticsFrequency(ParquetWriterOptions.StatisticsFrequency.PAGE) .build(); - Object[] listRows = IntStream.range(0, rows) - .mapToObj(i -> java.util.List.of(i * 3, i * 3 + 1, i * 3 + 2)) - .toArray(); + List[] listRows = (List[]) new List[rows]; + for (int i = 0; i < rows; i++) { + listRows[i] = Arrays.asList(i * 3, i * 3 + 1, i * 3 + 2); + } try (ColumnVector id = ColumnVector.fromInts(IntStream.range(0, rows).toArray()); ColumnVector zipCode = ColumnVector.fromInts(IntStream.range(0, rows).toArray()); ColumnVector listValues = ColumnVector.fromLists(LIST_OF_INTS, listRows); From ff88b4936613a7668c81540b25527ef70c3fedac Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Fri, 28 Aug 2026 17:12:40 +0000 Subject: [PATCH 06/21] Remove disabled overload --- .../hybrid_scan_io/hybrid_scan_composer.cpp | 15 --------------- .../java/ai/rapids/cudf/HybridScanReaderTest.java | 5 ++++- 2 files changed, 4 insertions(+), 16 deletions(-) diff --git a/cpp/examples/hybrid_scan_io/hybrid_scan_composer.cpp b/cpp/examples/hybrid_scan_io/hybrid_scan_composer.cpp index 47a2f36af349..bac0fed23ef4 100644 --- a/cpp/examples/hybrid_scan_io/hybrid_scan_composer.cpp +++ b/cpp/examples/hybrid_scan_io/hybrid_scan_composer.cpp @@ -426,21 +426,6 @@ std::unique_ptr hybrid_scan( } } -// Specialization for two-step read without page index -template - requires(not single_step_read and not use_page_index) -std::unique_ptr inline hybrid_scan( - io_source const& io_source, - std::optional filter_expression, - std::unordered_set const& filters, - bool verbose, - cuda::stream_ref stream, - rmm::device_async_resource_ref mr) -{ - static_assert(single_step_read or use_page_index, - "Hybrid scan requires parquet page index for two-step parquet read"); - return nullptr; -} // Instantiations for hybrid_scan template diff --git a/java/src/test/java/ai/rapids/cudf/HybridScanReaderTest.java b/java/src/test/java/ai/rapids/cudf/HybridScanReaderTest.java index 49a11c11ff41..a36282e6b2ee 100644 --- a/java/src/test/java/ai/rapids/cudf/HybridScanReaderTest.java +++ b/java/src/test/java/ai/rapids/cudf/HybridScanReaderTest.java @@ -1505,7 +1505,10 @@ private static void writeNoPageIndexParquet(File path, int rowsPerGroup, int num private static void writeNoPageIndexListParquet(File path) { int rows = 60_000; ParquetWriterOptions opts = ParquetWriterOptions.builder() - .withNonNullableColumns("id", "zip_code", "list_values") + .withNonNullableColumns("id", "zip_code") + .withListColumn(ColumnWriterOptions.listBuilder("list_values", false) + .withNonNullableColumns("element") + .build()) .withRowGroupSizeRows(rows) .withStatisticsFrequency(ParquetWriterOptions.StatisticsFrequency.PAGE) .build(); From b267d0b4d1f642f26174e5b3dfde4f8f1c0c1cdc Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Fri, 28 Aug 2026 20:25:00 +0000 Subject: [PATCH 07/21] formatting for the millionth time --- cpp/examples/hybrid_scan_io/hybrid_scan_composer.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/cpp/examples/hybrid_scan_io/hybrid_scan_composer.cpp b/cpp/examples/hybrid_scan_io/hybrid_scan_composer.cpp index bac0fed23ef4..0fedbac5e6f5 100644 --- a/cpp/examples/hybrid_scan_io/hybrid_scan_composer.cpp +++ b/cpp/examples/hybrid_scan_io/hybrid_scan_composer.cpp @@ -426,7 +426,6 @@ std::unique_ptr hybrid_scan( } } - // Instantiations for hybrid_scan template template std::unique_ptr hybrid_scan( From 4f786ef7db48b06232c197699e5cbbd6ccdc6126 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Fri, 28 Aug 2026 21:52:33 +0000 Subject: [PATCH 08/21] Fix JNI read option builder --- .../include/hybrid_scan_jni_internal.hpp | 1 - .../main/native/src/HybridScanReaderJni.cpp | 3 ++- .../src/HybridScanReaderJniInternal.cpp | 23 ------------------- 3 files changed, 2 insertions(+), 25 deletions(-) diff --git a/java/src/main/native/include/hybrid_scan_jni_internal.hpp b/java/src/main/native/include/hybrid_scan_jni_internal.hpp index b9c4ed2850ef..b4e986c027f0 100644 --- a/java/src/main/native/include/hybrid_scan_jni_internal.hpp +++ b/java/src/main/native/include/hybrid_scan_jni_internal.hpp @@ -74,7 +74,6 @@ struct row_group_span_holder { */ cudf::io::parquet_reader_options build_options(JNIEnv* env, jobjectArray j_column_names, - jbooleanArray j_read_binary_as_string, jint time_unit_type_id); row_group_span_holder make_row_group_span(JNIEnv* env, jintArray j_row_groups); diff --git a/java/src/main/native/src/HybridScanReaderJni.cpp b/java/src/main/native/src/HybridScanReaderJni.cpp index d73ba3493143..06830509cc50 100644 --- a/java/src/main/native/src/HybridScanReaderJni.cpp +++ b/java/src/main/native/src/HybridScanReaderJni.cpp @@ -41,7 +41,8 @@ Java_ai_rapids_cudf_HybridScanReader_createFromFooter(JNIEnv* env, { cudf::jni::auto_set_device(env); auto const len = checked_size_t(env, footer_length, "footerLength"); - auto opts = build_options(env, j_column_names, j_binary_as_str, time_unit_type_id); + (void)j_binary_as_str; + auto opts = build_options(env, j_column_names, time_unit_type_id); auto const* footer_ptr = reinterpret_cast(footer_address); cudf::host_span footer_bytes{footer_ptr, len}; auto wrapper = std::make_unique(footer_bytes, std::move(opts)); diff --git a/java/src/main/native/src/HybridScanReaderJniInternal.cpp b/java/src/main/native/src/HybridScanReaderJniInternal.cpp index e7974d2a650c..09fa5e9e4ac0 100644 --- a/java/src/main/native/src/HybridScanReaderJniInternal.cpp +++ b/java/src/main/native/src/HybridScanReaderJniInternal.cpp @@ -16,13 +16,8 @@ namespace hybrid_scan { cudf::io::parquet_reader_options build_options(JNIEnv* env, jobjectArray j_column_names, - jbooleanArray j_read_binary_as_string, jint time_unit_type_id) { - // The hybrid_scan_reader's options builder is constructed without a source_info because - // the reader works on already-parsed footer bytes (and on byte ranges fetched separately). - // Filter is not installed here; use HybridScanReader.setFilter (JNI setFilter) after - // construction. cudf::io::parquet_reader_options_builder builder; cudf::jni::native_jstringArray names(env, j_column_names); @@ -30,24 +25,6 @@ cudf::io::parquet_reader_options build_options(JNIEnv* env, builder = builder.column_names(names.as_cpp_vector()); } - // Translate Java's per-column "read binary as string" flags into the C++ schema override - // hooks. The reader_column_schema mechanism lets callers force binary→string conversion - // for the i-th projected column. - cudf::jni::native_jbooleanArray binary_as_str(env, j_read_binary_as_string); - if (!binary_as_str.is_null() && binary_as_str.size() > 0) { - std::vector schemas; - schemas.reserve(binary_as_str.size()); - for (int i = 0; i < binary_as_str.size(); ++i) { - cudf::io::reader_column_schema s; - s.set_convert_binary_to_strings(static_cast(binary_as_str[i])); - schemas.emplace_back(std::move(s)); - } - builder = builder.set_column_schema(std::move(schemas)); - binary_as_str.cancel(); - } - - // convert_strings_to_categories and ignore_missing_columns are fixed to match the - // standard cudf-java Parquet reader (see readParquet in TableJni.cpp). return builder.convert_strings_to_categories(false) .timestamp_type(cudf::data_type(static_cast(time_unit_type_id))) .ignore_missing_columns(true) From eaa693127b4b8a1352b55c2205c14f3b5eeb082d Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Mon, 31 Aug 2026 16:51:53 +0000 Subject: [PATCH 09/21] style --- java/src/main/native/src/HybridScanReaderJni.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/java/src/main/native/src/HybridScanReaderJni.cpp b/java/src/main/native/src/HybridScanReaderJni.cpp index 06830509cc50..6dad346d020d 100644 --- a/java/src/main/native/src/HybridScanReaderJni.cpp +++ b/java/src/main/native/src/HybridScanReaderJni.cpp @@ -40,9 +40,9 @@ Java_ai_rapids_cudf_HybridScanReader_createFromFooter(JNIEnv* env, JNI_TRY { cudf::jni::auto_set_device(env); - auto const len = checked_size_t(env, footer_length, "footerLength"); + auto const len = checked_size_t(env, footer_length, "footerLength"); (void)j_binary_as_str; - auto opts = build_options(env, j_column_names, time_unit_type_id); + auto opts = build_options(env, j_column_names, time_unit_type_id); auto const* footer_ptr = reinterpret_cast(footer_address); cudf::host_span footer_bytes{footer_ptr, len}; auto wrapper = std::make_unique(footer_bytes, std::move(opts)); From 5a57a4d6a800b9d9ac97e920488f72d9c8012ef8 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Mon, 31 Aug 2026 17:35:32 +0000 Subject: [PATCH 10/21] Apply suggestions from @pmattione-nvidia --- .../cudf/io/experimental/hybrid_scan.hpp | 2 +- .../io/experimental/hybrid_scan_multifile.hpp | 4 +- .../io/parquet/experimental/hybrid_scan.cpp | 2 +- .../experimental/hybrid_scan_helpers.hpp | 2 +- .../parquet/experimental/hybrid_scan_impl.cpp | 9 +- .../parquet/experimental/hybrid_scan_impl.hpp | 5 +- .../experimental/hybrid_scan_multifile.cpp | 2 +- .../experimental/page_index_filter_utils.cu | 115 +++++++++++------- .../experimental/page_index_filter_utils.hpp | 15 +-- .../src/HybridScanReaderJniMaterialize.cpp | 2 +- .../pylibcudf/io/experimental/hybrid_scan.pyx | 4 +- .../pylibcudf/libcudf/io/hybrid_scan.pxd | 2 +- 12 files changed, 98 insertions(+), 66 deletions(-) diff --git a/cpp/include/cudf/io/experimental/hybrid_scan.hpp b/cpp/include/cudf/io/experimental/hybrid_scan.hpp index 333fbfcf2c05..f0715f228ecf 100644 --- a/cpp/include/cudf/io/experimental/hybrid_scan.hpp +++ b/cpp/include/cudf/io/experimental/hybrid_scan.hpp @@ -656,7 +656,7 @@ class hybrid_scan_reader { std::size_t chunk_read_limit, std::size_t pass_read_limit, std::span row_group_indices, - cudf::mutable_column_view const& row_mask, + cudf::column_view const& row_mask, use_data_page_mask mask_data_pages, std::span const> column_chunk_data, parquet_reader_options const& options, diff --git a/cpp/include/cudf/io/experimental/hybrid_scan_multifile.hpp b/cpp/include/cudf/io/experimental/hybrid_scan_multifile.hpp index 0cc290270ff5..f770c5e77a07 100644 --- a/cpp/include/cudf/io/experimental/hybrid_scan_multifile.hpp +++ b/cpp/include/cudf/io/experimental/hybrid_scan_multifile.hpp @@ -389,7 +389,7 @@ class hybrid_scan_multifile { * @param pass_read_limit Limit on the memory used for reading and decompressing data. `0` if * there is no limit * @param row_group_indices Span of vectors of input row group indices, one per source - * @param[in,out] row_mask Mutable boolean column spanning all selected rows across all sources + * @param row_mask Mutable boolean column spanning all selected rows across all sources * indicating surviving rows * @param mask_data_pages Whether to build and use a data page mask using the row mask * @param column_chunk_data Flattened device spans of filter column chunk data returned in the @@ -402,7 +402,7 @@ class hybrid_scan_multifile { std::size_t chunk_read_limit, std::size_t pass_read_limit, cudf::host_span const> row_group_indices, - cudf::mutable_column_view const& row_mask, + cudf::column_view const& row_mask, use_data_page_mask mask_data_pages, cudf::host_span const> column_chunk_data, parquet_reader_options const& options, diff --git a/cpp/src/io/parquet/experimental/hybrid_scan.cpp b/cpp/src/io/parquet/experimental/hybrid_scan.cpp index 4f98d14ee92c..09faac8c261a 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan.cpp +++ b/cpp/src/io/parquet/experimental/hybrid_scan.cpp @@ -293,7 +293,7 @@ void hybrid_scan_reader::setup_chunking_for_filter_columns( std::size_t chunk_read_limit, std::size_t pass_read_limit, std::span row_group_indices, - cudf::mutable_column_view const& row_mask, + cudf::column_view const& row_mask, use_data_page_mask mask_data_pages, std::span const> column_chunk_data, parquet_reader_options const& options, diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_helpers.hpp b/cpp/src/io/parquet/experimental/hybrid_scan_helpers.hpp index 5134c3ee63d0..e1021eb8febb 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_helpers.hpp +++ b/cpp/src/io/parquet/experimental/hybrid_scan_helpers.hpp @@ -338,7 +338,7 @@ class aggregate_reader_metadata : public aggregate_reader_metadata_base { * Compute a vector of boolean vectors indicating which data pages need to be decoded to * construct each input column based on the row mask, one vector per column * - * @param row_mask Non-nullable boolean column view indicating surviving rows + * @param row_mask Boolean column view indicating surviving rows * @param row_group_indices Input row groups indices * @param input_columns Input column information * @param stream CUDA stream used for device memory operations and kernel launches diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp b/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp index 0dacccfc0e56..77239dd0c77e 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp +++ b/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp @@ -758,7 +758,7 @@ table_with_metadata hybrid_scan_reader_impl::materialize_filter_columns( auto data_page_mask = thrust::host_vector{}; if (mask_data_pages == use_data_page_mask::YES) { - _row_mask = set_nulls_to_true(row_mask, stream); + _row_mask = row_mask; data_page_mask = _extended_metadata->compute_data_page_mask( _row_mask, row_group_indices, _input_columns, stream); } @@ -837,7 +837,7 @@ void hybrid_scan_reader_impl::setup_chunking_for_filter_columns( std::size_t chunk_read_limit, std::size_t pass_read_limit, std::span const> row_group_indices, - cudf::mutable_column_view const& row_mask, + cudf::column_view const& row_mask, use_data_page_mask mask_data_pages, std::span const> column_chunk_data, parquet_reader_options const& options, @@ -869,7 +869,7 @@ void hybrid_scan_reader_impl::setup_chunking_for_filter_columns( auto data_page_mask = thrust::host_vector{}; if (mask_data_pages == use_data_page_mask::YES) { - _row_mask = set_nulls_to_true(row_mask, stream); + _row_mask = row_mask; data_page_mask = _extended_metadata->compute_data_page_mask( _row_mask, row_group_indices, _input_columns, stream); } @@ -1238,6 +1238,9 @@ void hybrid_scan_reader_impl::prepare_data( if (_file_itm_data._current_input_pass < _file_itm_data.num_passes()) { handle_chunking(mode, column_chunk_data, data_page_mask); } + + // Clear the cached row mask column view + _row_mask = cudf::column_view{}; } template diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_impl.hpp b/cpp/src/io/parquet/experimental/hybrid_scan_impl.hpp index 5e30e2b578f0..ce32d96b25e4 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_impl.hpp +++ b/cpp/src/io/parquet/experimental/hybrid_scan_impl.hpp @@ -253,7 +253,7 @@ class hybrid_scan_reader_impl : public parquet::detail::reader_impl { std::size_t chunk_read_limit, std::size_t pass_read_limit, std::span const> row_group_indices, - cudf::mutable_column_view const& row_mask, + cudf::column_view const& row_mask, use_data_page_mask mask_data_pages, std::span const> column_chunk_data, parquet_reader_options const& options, @@ -631,6 +631,9 @@ class hybrid_scan_reader_impl : public parquet::detail::reader_impl { std::optional> _filter_columns_names; + // Non-owning view of the caller's row mask, only valid for the duration of a single + // materialization or chunking setup call, during which the pass page mask is computed. Null + // entries mean the row could not be pruned and is therefore treated as a surviving row. cudf::column_view _row_mask{}; std::vector _original_output_buffers_template; diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_multifile.cpp b/cpp/src/io/parquet/experimental/hybrid_scan_multifile.cpp index 16360d0bc7d9..259435401aab 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_multifile.cpp +++ b/cpp/src/io/parquet/experimental/hybrid_scan_multifile.cpp @@ -197,7 +197,7 @@ void hybrid_scan_multifile::setup_chunking_for_filter_columns( std::size_t chunk_read_limit, std::size_t pass_read_limit, cudf::host_span const> row_group_indices, - cudf::mutable_column_view const& row_mask, + cudf::column_view const& row_mask, use_data_page_mask mask_data_pages, cudf::host_span const> column_chunk_data, parquet_reader_options const& options, diff --git a/cpp/src/io/parquet/experimental/page_index_filter_utils.cu b/cpp/src/io/parquet/experimental/page_index_filter_utils.cu index 9df82e836c10..de80e4e637a2 100644 --- a/cpp/src/io/parquet/experimental/page_index_filter_utils.cu +++ b/cpp/src/io/parquet/experimental/page_index_filter_utils.cu @@ -5,14 +5,13 @@ #include "page_index_filter_utils.hpp" -#include #include -#include #include #include #include #include #include +#include #include #include @@ -20,7 +19,6 @@ #include #include -#include #include #include @@ -32,20 +30,65 @@ namespace cudf::io::parquet::experimental::detail { namespace { +/** + * @brief Functor to read the row mask, i.e. the zeroth Fenwick tree level + * + * Nulls are read as retained rows here so that pages aren't accidentally pruned due to + * unavailable page-level statistics (represented as nulls) + */ +struct row_mask_accessor { + bool const* data; ///< Row mask data, adjusted for the column view offset + bitmask_type const* null_mask; ///< Null mask, or nullptr if there are no nulls + cudf::size_type offset; ///< Column view offset, needed to index into the null mask + + /** + * @brief Constructs a row mask accessor from a nullable boolean row mask column + * + * @param row_mask Boolean row mask column + */ + explicit row_mask_accessor(cudf::column_view const& row_mask) + : data{row_mask.begin()}, + // Skip the validity check altogether if there are no nulls to resolve + null_mask{row_mask.null_count() ? row_mask.null_mask() : nullptr}, + offset{row_mask.offset()} + { + } + + __device__ bool inline operator()(cudf::size_type row_idx) const noexcept + { + auto const is_null = null_mask != nullptr and not bit_is_set(null_mask, offset + row_idx); + return is_null or data[row_idx]; + } +}; + /* * @brief Functor to build a Fenwick tree level from the previous level data * * @param tree_level_ptrs Pointers to the start of Fenwick tree level data + * @param row_mask Accessor for the zeroth tree level (the row mask) * @param prev_level Previous tree level * @param prev_level_size Size of the previous tree level * @param current_level_size Size of the current tree level */ struct build_fenwick_tree_level_functor { bool** tree_level_ptrs; + row_mask_accessor row_mask; cudf::size_type prev_level; cudf::size_type prev_level_size; cudf::size_type current_level_size; + /** + * @brief Reads an element from the previous tree level + * + * @param prev_level_idx Previous tree level element index + * @return Value of the element at the previous tree level + */ + __device__ bool inline read_prev_level(cudf::size_type prev_level_idx) const noexcept + { + // `prev_level` is uniform across all threads so this branch is free + return prev_level == 0 ? row_mask(prev_level_idx) : tree_level_ptrs[prev_level][prev_level_idx]; + } + /** * @brief Builds the next Fenwick tree level from the current level data * by ORing two elements at the current level. @@ -56,15 +99,14 @@ struct build_fenwick_tree_level_functor { */ __device__ void operator()(cudf::size_type current_level_idx) const noexcept { - auto const prev_level_ptr = tree_level_ptrs[prev_level]; - auto current_level_ptr = tree_level_ptrs[prev_level + 1]; + auto current_level_ptr = tree_level_ptrs[prev_level + 1]; // Handle the odd-sized remaining element if prev_level_size is odd if (prev_level_size % 2 and current_level_idx == current_level_size - 1) { - current_level_ptr[current_level_idx] = prev_level_ptr[prev_level_size - 1]; + current_level_ptr[current_level_idx] = read_prev_level(prev_level_size - 1); } else { current_level_ptr[current_level_idx] = - prev_level_ptr[(current_level_idx * 2)] or prev_level_ptr[(current_level_idx * 2) + 1]; + read_prev_level(current_level_idx * 2) or read_prev_level((current_level_idx * 2) + 1); } } }; @@ -73,12 +115,14 @@ struct build_fenwick_tree_level_functor { * @brief Functor to binary search a `true` value in the Fenwick tree in range [start, end) * * @param tree_level_ptrs Pointers to the start of Fenwick tree level data + * @param row_mask Accessor for the zeroth tree level (the row mask) * @param page_offsets Pointer to page offsets describing each search range i as [page_offsets[i], * page_offsets[i+1)) * @param num_ranges Number of search ranges */ struct search_fenwick_tree_functor { bool** tree_level_ptrs; + row_mask_accessor row_mask; cudf::size_type const* page_offsets; cudf::size_type num_ranges; @@ -185,13 +229,9 @@ struct search_fenwick_tree_functor { cudf::size_type tree_level, cudf::size_type block_size) const noexcept { - if constexpr (Boundary == boundary::START) { - auto const mask_index = boundary_pos >> tree_level; - return tree_level_ptrs[tree_level][mask_index]; - } else { - auto const mask_index = (boundary_pos - block_size) >> tree_level; - return tree_level_ptrs[tree_level][mask_index]; - } + auto const position = (Boundary == boundary::START) ? boundary_pos : boundary_pos - block_size; + auto const mask_index = position >> tree_level; + return tree_level == 0 ? row_mask(mask_index) : tree_level_ptrs[tree_level][mask_index]; } /** @@ -426,26 +466,12 @@ rmm::device_uvector compute_page_indices_async( return page_indices; } -cudf::column_view set_nulls_to_true(cudf::mutable_column_view const& row_mask, - cuda::stream_ref stream) -{ - if (row_mask.has_nulls()) { - auto const d_row_mask = cudf::column_device_view::create(row_mask, stream); - auto const iter = cudf::detail::make_null_replacement_iterator(*d_row_mask, true, true); - thrust::copy(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), - iter, - iter + row_mask.size(), - row_mask.begin()); - } - - return cudf::column_view{ - row_mask.type(), row_mask.size(), row_mask.head(), nullptr, 0, row_mask.offset()}; -} - bool are_all_rows_retained(cudf::column_view const& row_mask, cuda::stream_ref stream) { - return cudf::detail::all_of( - row_mask.begin(), row_mask.end(), cuda::std::identity{}, stream); + return cudf::detail::all_of(cuda::counting_iterator{0}, + cuda::counting_iterator{row_mask.size()}, + row_mask_accessor{row_mask}, + stream); } thrust::host_vector compute_row_range_selection_mask( @@ -463,7 +489,9 @@ thrust::host_vector compute_row_range_selection_mask( auto const num_levels = static_cast(tree_level_offsets.size()); auto tree_levels_data = rmm::device_uvector(tree_level_offsets.back(), stream, mr); auto host_tree_level_ptrs = cudf::detail::make_pinned_vector_async(num_levels, stream); - host_tree_level_ptrs[0] = const_cast(row_mask.begin()); + // The zeroth level is the row mask itself, read through its accessor + auto const d_row_mask = row_mask_accessor{row_mask}; + host_tree_level_ptrs[0] = nullptr; std::for_each(cuda::counting_iterator{1}, cuda::counting_iterator{num_levels}, [&](auto const level_idx) { @@ -481,8 +509,11 @@ thrust::host_vector compute_row_range_selection_mask( thrust::for_each(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), cuda::counting_iterator{0}, cuda::counting_iterator{current_level_size}, - build_fenwick_tree_level_functor{ - tree_level_ptrs.data(), prev_level, prev_level_size, current_level_size}); + build_fenwick_tree_level_functor{.tree_level_ptrs = tree_level_ptrs.data(), + .row_mask = d_row_mask, + .prev_level = prev_level, + .prev_level_size = prev_level_size, + .current_level_size = current_level_size}); prev_level_size = current_level_size; }); @@ -490,12 +521,14 @@ thrust::host_vector compute_row_range_selection_mask( auto device_results = rmm::device_uvector(num_ranges, stream, mr); auto pinned_page_offsets = cudf::detail::make_pinned_vector(page_row_offsets, stream); auto page_offsets = cudf::detail::make_device_uvector_async(pinned_page_offsets, stream, mr); - thrust::transform( - rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), - cuda::counting_iterator{0}, - cuda::counting_iterator{num_ranges}, - device_results.begin(), - search_fenwick_tree_functor{tree_level_ptrs.data(), page_offsets.data(), num_ranges}); + thrust::transform(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + cuda::counting_iterator{0}, + cuda::counting_iterator{num_ranges}, + device_results.begin(), + search_fenwick_tree_functor{.tree_level_ptrs = tree_level_ptrs.data(), + .row_mask = d_row_mask, + .page_offsets = page_offsets.data(), + .num_ranges = num_ranges}); auto results = cudf::detail::make_pinned_vector_async(device_results, stream); stream.sync(); diff --git a/cpp/src/io/parquet/experimental/page_index_filter_utils.hpp b/cpp/src/io/parquet/experimental/page_index_filter_utils.hpp index bc7b5eecfa5e..407030380883 100644 --- a/cpp/src/io/parquet/experimental/page_index_filter_utils.hpp +++ b/cpp/src/io/parquet/experimental/page_index_filter_utils.hpp @@ -70,20 +70,11 @@ compute_page_row_offsets_and_colchunk_page_offsets( cuda::stream_ref stream, rmm::device_async_resource_ref mr); -/** - * @brief Sets nulls in the row mask to true and returns a non-nullable row mask view - * - * @param row_mask Mutable row mask column view - * @param stream CUDA stream used for device memory - * operations and kernel launches - * @return Non-nullable column view of the resolved row mask - */ -[[nodiscard]] cudf::column_view set_nulls_to_true(cudf::mutable_column_view const& row_mask, - cuda::stream_ref stream); - /** * @brief Checks whether every row is reatained by the boolean row mask * + * Null entries in the row mask are treated as surviving rows + * * @param retention_mask Boolean column indicating retained rows * @param stream CUDA stream used for device memory operations and kernel launches * @return Boolean indicating whether every row is retained @@ -94,6 +85,8 @@ compute_page_row_offsets_and_colchunk_page_offsets( /** * @brief Computes a mask indicating which row ranges contain at least one selected row * + * Null entries in the row mask are treated as surviving rows + * * @param row_mask Boolean column indicating selected rows * @param page_row_offsets Page row offsets defining the row ranges * @param max_page_size Size of the largest page row range diff --git a/java/src/main/native/src/HybridScanReaderJniMaterialize.cpp b/java/src/main/native/src/HybridScanReaderJniMaterialize.cpp index 172eb5acb947..f6f671ae4e7d 100644 --- a/java/src/main/native/src/HybridScanReaderJniMaterialize.cpp +++ b/java/src/main/native/src/HybridScanReaderJniMaterialize.cpp @@ -183,7 +183,7 @@ Java_ai_rapids_cudf_HybridScanReader_setupChunkingForFilterColumns(JNIEnv* env, wrapper->reader->setup_chunking_for_filter_columns(chunk_limit, pass_limit, holder.span(), - row_mask_col->mutable_view(), + row_mask_col->view(), mode, spans, wrapper->options, diff --git a/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pyx b/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pyx index 336528b06ffa..0d092541d84b 100644 --- a/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pyx +++ b/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pyx @@ -831,7 +831,7 @@ cdef class HybridScanReader: row_group_indices : list[int] Input row group indices row_mask : Column - Mutable boolean column indicating surviving rows + Boolean column indicating surviving rows mask_data_pages : UseDataPageMask Whether to use a data page mask column_chunk_data : Sequence @@ -854,7 +854,7 @@ cdef class HybridScanReader: # keep reference to avoid use-after-free of device spans self._filter_chunk_data = column_chunk_data - cdef mutable_column_view mask_view = row_mask.mutable_view() + cdef column_view mask_view = row_mask.view() with nogil: self.c_obj.get()[0].setup_chunking_for_filter_columns( chunk_read_limit, diff --git a/python/pylibcudf/pylibcudf/libcudf/io/hybrid_scan.pxd b/python/pylibcudf/pylibcudf/libcudf/io/hybrid_scan.pxd index 04d84fb68a33..7a5aec269a56 100644 --- a/python/pylibcudf/pylibcudf/libcudf/io/hybrid_scan.pxd +++ b/python/pylibcudf/pylibcudf/libcudf/io/hybrid_scan.pxd @@ -161,7 +161,7 @@ cdef extern from "cudf/io/experimental/hybrid_scan.hpp" \ size_t chunk_read_limit, size_t pass_read_limit, std_span[const_size_type] row_group_indices, - const mutable_column_view& row_mask, + const column_view& row_mask, use_data_page_mask mask_data_pages, std_span[const_device_span_const_uint8_t] column_chunk_data, const parquet_reader_options& options, From e2f376c1df94b47b6f0d5ae631653f03c2f69a02 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Tue, 1 Sep 2026 01:01:50 +0000 Subject: [PATCH 11/21] minor --- cpp/src/io/parquet/experimental/page_index_filter_utils.cu | 7 +++---- 1 file changed, 3 insertions(+), 4 deletions(-) diff --git a/cpp/src/io/parquet/experimental/page_index_filter_utils.cu b/cpp/src/io/parquet/experimental/page_index_filter_utils.cu index de80e4e637a2..1430e3103fe1 100644 --- a/cpp/src/io/parquet/experimental/page_index_filter_utils.cu +++ b/cpp/src/io/parquet/experimental/page_index_filter_utils.cu @@ -48,16 +48,15 @@ struct row_mask_accessor { */ explicit row_mask_accessor(cudf::column_view const& row_mask) : data{row_mask.begin()}, - // Skip the validity check altogether if there are no nulls to resolve - null_mask{row_mask.null_count() ? row_mask.null_mask() : nullptr}, + null_mask{row_mask.has_nulls() ? row_mask.null_mask() : nullptr}, offset{row_mask.offset()} { } __device__ bool inline operator()(cudf::size_type row_idx) const noexcept { - auto const is_null = null_mask != nullptr and not bit_is_set(null_mask, offset + row_idx); - return is_null or data[row_idx]; + if (data[row_idx]) { return true; } + return null_mask != nullptr and not bit_is_set(null_mask, offset + row_idx); } }; From ef23068f3bca6975f204ee1036c537c537965ee6 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Tue, 1 Sep 2026 01:03:46 +0000 Subject: [PATCH 12/21] minor --- cpp/src/io/parquet/experimental/page_index_filter_utils.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/src/io/parquet/experimental/page_index_filter_utils.cu b/cpp/src/io/parquet/experimental/page_index_filter_utils.cu index 1430e3103fe1..f143a8acabaf 100644 --- a/cpp/src/io/parquet/experimental/page_index_filter_utils.cu +++ b/cpp/src/io/parquet/experimental/page_index_filter_utils.cu @@ -84,7 +84,7 @@ struct build_fenwick_tree_level_functor { */ __device__ bool inline read_prev_level(cudf::size_type prev_level_idx) const noexcept { - // `prev_level` is uniform across all threads so this branch is free + // Use the row mask accessor for the zeroth tree level return prev_level == 0 ? row_mask(prev_level_idx) : tree_level_ptrs[prev_level][prev_level_idx]; } From 3f1f806244c343cd13a4a46aaa2c88b957b9020b Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Wed, 2 Sep 2026 00:11:53 +0000 Subject: [PATCH 13/21] Fix --- cpp/src/io/parquet/experimental/page_index_filter.cu | 2 -- 1 file changed, 2 deletions(-) diff --git a/cpp/src/io/parquet/experimental/page_index_filter.cu b/cpp/src/io/parquet/experimental/page_index_filter.cu index 6ec51293aa50..25d3878b6fde 100644 --- a/cpp/src/io/parquet/experimental/page_index_filter.cu +++ b/cpp/src/io/parquet/experimental/page_index_filter.cu @@ -765,8 +765,6 @@ thrust::host_vector aggregate_reader_metadata::compute_data_page_mask( std::cmp_equal(total_rows, row_mask.size()), "Encountered a mismatch in number of rows in the row group pass and the row mask size", std::overflow_error); - CUDF_EXPECTS( - row_mask.null_count() == 0, "Row mask must not contain nulls", std::invalid_argument); // Return an empty vector if all rows are required if (are_all_rows_retained(row_mask, stream)) { return thrust::host_vector{}; } From f2ce44d8402ec46cda08cdff9ae2d8ac62b03924 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Thu, 3 Sep 2026 07:34:47 +0000 Subject: [PATCH 14/21] Add suggested test --- .../experimental/hybrid_scan_filters_test.cpp | 59 +++++++++++++++++++ 1 file changed, 59 insertions(+) diff --git a/cpp/tests/io/experimental/hybrid_scan_filters_test.cpp b/cpp/tests/io/experimental/hybrid_scan_filters_test.cpp index 472b41d3f889..6579db00f550 100644 --- a/cpp/tests/io/experimental/hybrid_scan_filters_test.cpp +++ b/cpp/tests/io/experimental/hybrid_scan_filters_test.cpp @@ -29,6 +29,7 @@ #include #include #include +#include #include #include #include @@ -856,6 +857,64 @@ TYPED_TEST(PageFilteringWithPageIndexStats, FilterPages) } } +TEST_F(HybridScanFiltersTest, RowMaskNullsAreRetained) +{ + auto constexpr nan = std::numeric_limits::quiet_NaN(); + auto const input = cudf::test::fixed_width_column_wrapper{{3.0, nan, 1.0, 4.0}}; + auto const table = cudf::table_view{{input}}; + auto buffer = std::vector{}; + + auto metadata = cudf::io::table_input_metadata{table}; + metadata.column_metadata[0].set_name("col"); + auto writer_options = + cudf::io::parquet_writer_options::builder(cudf::io::sink_info{&buffer}, table) + .metadata(std::move(metadata)) + .stats_level(cudf::io::statistics_freq::STATISTICS_COLUMN) + .build(); + cudf::io::write_parquet(writer_options); + + auto scalar = cudf::numeric_scalar{2.0}; + auto literal = cudf::ast::literal{scalar}; + auto column_ref = cudf::ast::column_name_reference{"col"}; + auto filter = cudf::ast::operation{cudf::ast::ast_operator::GREATER, column_ref, literal}; + auto options = cudf::io::parquet_reader_options::builder().filter(filter).build(); + + auto datasource = cudf::io::datasource::create(cudf::host_span{ + reinterpret_cast(buffer.data()), buffer.size()}); + auto footer = cudf::io::parquet::fetch_footer_to_host(*datasource); + auto reader = cudf::io::parquet::experimental::hybrid_scan_reader{*footer, options}; + + auto page_index = + cudf::io::parquet::fetch_page_index_to_host(*datasource, reader.page_index_byte_range()); + reader.setup_page_index(*page_index); + + auto const stream = cudf::get_default_stream(); + auto const mr = cudf::get_current_device_resource_ref(); + auto const row_groups = reader.all_row_groups(options); + auto row_mask = reader.build_row_mask_with_page_index_stats(row_groups, options, stream, mr); + + ASSERT_EQ(row_mask->size(), table.num_rows()); + ASSERT_EQ(row_mask->null_count(), table.num_rows()); + + auto const byte_ranges = reader.filter_column_chunks_byte_ranges(row_groups, options); + auto [column_buffers, column_data, read_tasks] = + cudf::io::parquet::fetch_byte_ranges_to_device_async(*datasource, byte_ranges, stream, mr); + read_tasks.get(); + + auto row_mask_view = row_mask->mutable_view(); + auto result = + reader.materialize_filter_columns(row_groups, + column_data, + row_mask_view, + cudf::io::parquet::experimental::use_data_page_mask::YES, + options, + stream, + mr); + + auto const expected = cudf::test::fixed_width_column_wrapper{{3.0, 4.0}}; + CUDF_TEST_EXPECT_TABLES_EQUAL(cudf::table_view{{expected}}, result.tbl->view()); +} + TEST_F(HybridScanFiltersTest, OffsetIndexOnlyDataPageMask) { using T = uint32_t; From f36b1f9dc480b2e2efbbfbd5851d85f26a2b98a2 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Thu, 3 Sep 2026 17:43:19 +0000 Subject: [PATCH 15/21] fix gtest --- .../experimental/hybrid_scan_filters_test.cpp | 22 +++++++++++-------- 1 file changed, 13 insertions(+), 9 deletions(-) diff --git a/cpp/tests/io/experimental/hybrid_scan_filters_test.cpp b/cpp/tests/io/experimental/hybrid_scan_filters_test.cpp index 9716ce06ab63..529ee4a7bd38 100644 --- a/cpp/tests/io/experimental/hybrid_scan_filters_test.cpp +++ b/cpp/tests/io/experimental/hybrid_scan_filters_test.cpp @@ -859,8 +859,8 @@ TYPED_TEST(PageFilteringWithPageIndexStats, FilterPages) TEST_F(HybridScanFiltersTest, RowMaskNullsAreRetained) { - auto constexpr nan = std::numeric_limits::quiet_NaN(); - auto const input = cudf::test::fixed_width_column_wrapper{{3.0, nan, 1.0, 4.0}}; + auto const input = cudf::test::fixed_width_column_wrapper{ + {3.0, 0.0, 1.0, 0.0}, {true, false, true, false}}; auto const table = cudf::table_view{{input}}; auto buffer = std::vector{}; @@ -873,10 +873,8 @@ TEST_F(HybridScanFiltersTest, RowMaskNullsAreRetained) .build(); cudf::io::write_parquet(writer_options); - auto scalar = cudf::numeric_scalar{2.0}; - auto literal = cudf::ast::literal{scalar}; auto column_ref = cudf::ast::column_name_reference{"col"}; - auto filter = cudf::ast::operation{cudf::ast::ast_operator::GREATER, column_ref, literal}; + auto filter = cudf::ast::operation{cudf::ast::ast_operator::IS_NULL, column_ref}; auto options = cudf::io::parquet_reader_options::builder().filter(filter).build(); auto datasource = cudf::io::datasource::create(cudf::host_span{ @@ -898,7 +896,8 @@ TEST_F(HybridScanFiltersTest, RowMaskNullsAreRetained) auto const byte_ranges = reader.filter_column_chunks_byte_ranges(row_groups, options); auto [column_buffers, column_data, read_tasks] = - cudf::io::parquet::fetch_byte_ranges_to_device_async(*datasource, byte_ranges, stream, mr); + cudf::io::parquet::fetch_byte_ranges_to_device_async( + *datasource, byte_ranges, cudf::io::parquet::io_submission_policy::SERIALIZE, stream, mr); read_tasks.get(); auto row_mask_view = row_mask->mutable_view(); @@ -911,7 +910,8 @@ TEST_F(HybridScanFiltersTest, RowMaskNullsAreRetained) stream, mr); - auto const expected = cudf::test::fixed_width_column_wrapper{{3.0, 4.0}}; + auto const expected = + cudf::test::fixed_width_column_wrapper{{0.0, 0.0}, {false, false}}; CUDF_TEST_EXPECT_TABLES_EQUAL(cudf::table_view{{expected}}, result.tbl->view()); } @@ -1029,7 +1029,11 @@ TEST_F(HybridScanFiltersTest, NoOffsetIndexListColumns) auto const list_byte_ranges = list_reader.payload_column_chunks_byte_ranges(list_row_groups, options); auto [list_buffers, list_data, list_tasks] = cudf::io::parquet::fetch_byte_ranges_to_device_async( - *list_datasource, list_byte_ranges, stream, mr); + *list_datasource, + list_byte_ranges, + cudf::io::parquet::io_submission_policy::SERIALIZE, + stream, + mr); list_tasks.get(); auto const list_result = list_reader.materialize_payload_columns( @@ -1040,7 +1044,7 @@ TEST_F(HybridScanFiltersTest, NoOffsetIndexListColumns) options, stream, mr); - auto const list_expected = cudf::apply_boolean_mask(list_table, list_row_mask_view, stream, mr); + auto const list_expected = cudf::apply_retention_mask(list_table, list_row_mask_view, stream, mr); CUDF_TEST_EXPECT_TABLES_EQUIVALENT(list_expected->view(), list_result.tbl->view()); } From 3e576ad5f9abb8c230b468da6d850907290ee1d0 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Fri, 4 Sep 2026 23:46:09 +0000 Subject: [PATCH 16/21] Address comments --- .../cudf/io/experimental/hybrid_scan.hpp | 3 +- .../experimental/hybrid_scan_chunking.cu | 12 ++-- .../parquet/experimental/hybrid_scan_impl.cpp | 70 ++++++++----------- .../parquet/experimental/hybrid_scan_impl.hpp | 27 ++++--- .../experimental/page_index_filter_utils.cu | 2 +- .../experimental/page_index_filter_utils.hpp | 2 +- 6 files changed, 55 insertions(+), 61 deletions(-) diff --git a/cpp/include/cudf/io/experimental/hybrid_scan.hpp b/cpp/include/cudf/io/experimental/hybrid_scan.hpp index 336ac6d3b311..c9a2fff32d86 100644 --- a/cpp/include/cudf/io/experimental/hybrid_scan.hpp +++ b/cpp/include/cudf/io/experimental/hybrid_scan.hpp @@ -668,6 +668,7 @@ class hybrid_scan_reader { parquet_reader_options const& options, cuda::stream_ref stream, rmm::device_async_resource_ref mr) const; + /** * @brief Setup chunking information for filter columns and preprocess the input data pages * @@ -676,7 +677,7 @@ class hybrid_scan_reader { * @param pass_read_limit Limit on the memory used for reading and decompressing data. `0` if * there is no limit * @param row_group_indices Input row groups indices - * @param[in,out] row_mask Mutable boolean column indicating surviving rows + * @param row_mask Boolean column indicating surviving rows * @param mask_data_pages Whether to build and use a data page mask using the row mask * @param column_chunk_data Device spans of column chunk data of filter columns * @param options Parquet reader options diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu b/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu index 4918b8b1c532..eecf066478fd 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu +++ b/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu @@ -33,12 +33,13 @@ using parquet::detail::pass_intermediate_data; void hybrid_scan_reader_impl::handle_chunking( read_mode mode, std::span const> column_chunk_data, - host_span data_page_mask) + host_span data_page_mask, + std::optional row_mask) { // if this is our first time in here, setup the first pass. if (!_pass_itm_data) { // setup the next pass - setup_next_pass(column_chunk_data, data_page_mask); + setup_next_pass(column_chunk_data, data_page_mask, row_mask); } auto& pass = *_pass_itm_data; @@ -76,7 +77,8 @@ void hybrid_scan_reader_impl::handle_chunking( void hybrid_scan_reader_impl::setup_next_pass( std::span const> column_chunk_data, - std::span data_page_mask) + std::span data_page_mask, + std::optional row_mask) { auto const num_passes = _file_itm_data.num_passes(); CUDF_EXPECTS(num_passes == 1, @@ -129,8 +131,8 @@ void hybrid_scan_reader_impl::setup_next_pass( // When offset index is absent, compute and use the data page mask using the decoded page // headers from `setup_compressed_data`. auto const data_page_mask_pghdr = [&]() { - if (not _has_offset_index and not _row_mask.is_empty()) { - return compute_data_page_mask_with_page_headers(); + if (not _has_offset_index and row_mask.has_value()) { + return compute_data_page_mask_with_page_headers(row_mask.value()); } return thrust::host_vector{}; }(); diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp b/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp index dc2a687f8d0b..e1a97970c443 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp +++ b/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp @@ -730,14 +730,10 @@ table_with_metadata hybrid_scan_reader_impl::materialize_filter_columns( return read_chunk_internal(read_mode::READ_ALL, read_columns_mode::FILTER_COLUMNS, row_mask); } - auto data_page_mask = thrust::host_vector{}; - if (mask_data_pages == use_data_page_mask::YES) { - _row_mask = row_mask; - data_page_mask = _extended_metadata->compute_data_page_mask( - _row_mask, row_group_indices, _input_columns, stream); - } - - prepare_data(read_mode::READ_ALL, row_group_indices, column_chunk_data, data_page_mask); + auto const retention_mask = mask_data_pages == use_data_page_mask::YES + ? std::optional{cudf::column_view{row_mask}} + : std::nullopt; + prepare_data(read_mode::READ_ALL, row_group_indices, column_chunk_data, retention_mask); return read_chunk_internal(read_mode::READ_ALL, read_columns_mode::FILTER_COLUMNS, row_mask); } @@ -769,14 +765,10 @@ table_with_metadata hybrid_scan_reader_impl::materialize_payload_columns( return read_chunk_internal(read_mode::READ_ALL, read_columns_mode::PAYLOAD_COLUMNS, row_mask); } - auto data_page_mask = thrust::host_vector{}; - if (not row_mask.is_empty() and mask_data_pages == use_data_page_mask::YES) { - _row_mask = row_mask; - data_page_mask = _extended_metadata->compute_data_page_mask( - _row_mask, row_group_indices, _input_columns, stream); - } - - prepare_data(read_mode::READ_ALL, row_group_indices, column_chunk_data, data_page_mask); + auto const retention_mask = mask_data_pages == use_data_page_mask::YES + ? std::optional{row_mask} + : std::nullopt; + prepare_data(read_mode::READ_ALL, row_group_indices, column_chunk_data, retention_mask); return read_chunk_internal(read_mode::READ_ALL, read_columns_mode::PAYLOAD_COLUMNS, row_mask); } @@ -841,14 +833,9 @@ void hybrid_scan_reader_impl::setup_chunking_for_filter_columns( return; } - auto data_page_mask = thrust::host_vector{}; - if (mask_data_pages == use_data_page_mask::YES) { - _row_mask = row_mask; - data_page_mask = _extended_metadata->compute_data_page_mask( - _row_mask, row_group_indices, _input_columns, stream); - } - - prepare_data(read_mode::CHUNKED_READ, row_group_indices, column_chunk_data, data_page_mask); + auto const retention_mask = + mask_data_pages == use_data_page_mask::YES ? std::optional{row_mask} : std::nullopt; + prepare_data(read_mode::CHUNKED_READ, row_group_indices, column_chunk_data, retention_mask); } table_with_metadata hybrid_scan_reader_impl::materialize_filter_columns_chunk( @@ -902,14 +889,9 @@ void hybrid_scan_reader_impl::setup_chunking_for_payload_columns( return; } - auto data_page_mask = thrust::host_vector{}; - if (not row_mask.is_empty() and mask_data_pages == use_data_page_mask::YES) { - _row_mask = row_mask; - data_page_mask = _extended_metadata->compute_data_page_mask( - _row_mask, row_group_indices, _input_columns, stream); - } - - prepare_data(read_mode::CHUNKED_READ, row_group_indices, column_chunk_data, data_page_mask); + auto const retention_mask = + mask_data_pages == use_data_page_mask::YES ? std::optional{row_mask} : std::nullopt; + prepare_data(read_mode::CHUNKED_READ, row_group_indices, column_chunk_data, retention_mask); } void hybrid_scan_reader_impl::setup_chunking_for_payload_columns( @@ -1138,7 +1120,6 @@ void hybrid_scan_reader_impl::reset_internal_state() _strings_to_categorical = false; _reader_column_schema.reset(); - _row_mask = column_view{}; _row_mask_offset = 0; _expr_conv = parquet_filter_normalizer{}; @@ -1196,7 +1177,7 @@ void hybrid_scan_reader_impl::prepare_data( read_mode mode, std::span const> row_group_indices, std::span const> column_chunk_data, - host_span data_page_mask) + std::optional row_mask) { // if we have not preprocessed at the whole-file level, do that now if (not _file_preprocessed) { @@ -1207,14 +1188,18 @@ void hybrid_scan_reader_impl::prepare_data( prepare_row_groups(read_mode::READ_ALL, row_group_indices); } + // Compute data page mask from the row (retention) mask + auto data_page_mask = thrust::host_vector{}; + if (_has_offset_index and row_mask.has_value() and not _sparse_page_io) { + data_page_mask = _extended_metadata->compute_data_page_mask( + row_mask.value(), row_group_indices, _input_columns, _stream); + } + // handle any chunking work (ratcheting through the subpasses and chunks within // our current pass) if in bounds if (_file_itm_data._current_input_pass < _file_itm_data.num_passes()) { - handle_chunking(mode, column_chunk_data, data_page_mask); + handle_chunking(mode, column_chunk_data, data_page_mask, row_mask); } - - // Clear the cached row mask column view - _row_mask = cudf::column_view{}; } template @@ -1491,12 +1476,13 @@ void hybrid_scan_reader_impl::set_pass_page_mask(std::span data_page mark_buffers_nullable_for_pruned_pages(); } -thrust::host_vector hybrid_scan_reader_impl::compute_data_page_mask_with_page_headers() +thrust::host_vector hybrid_scan_reader_impl::compute_data_page_mask_with_page_headers( + cudf::column_view const& row_mask) { auto const& pass = *_pass_itm_data; // Return an empty vector if all rows are required - if (are_all_rows_retained(_row_mask, _stream)) { return thrust::host_vector(0); } + if (are_all_rows_retained(row_mask, _stream)) { return thrust::host_vector{}; } std::vector page_row_offsets; page_row_offsets.reserve(pass.pages.size() * 2); @@ -1539,10 +1525,10 @@ thrust::host_vector hybrid_scan_reader_impl::compute_data_page_mask_with_p auto data_page_mask = thrust::host_vector{}; // Compute the row range mask - CUDF_EXPECTS(std::cmp_equal(_row_mask.size(), pass.num_rows), + CUDF_EXPECTS(std::cmp_equal(row_mask.size(), pass.num_rows), "Row mask must span across all rows in the pass"); auto const row_range_mask = - compute_row_range_selection_mask(_row_mask, page_row_offsets, max_page_size, _stream); + compute_row_range_selection_mask(row_mask, page_row_offsets, max_page_size, _stream); if (row_range_mask.empty()) { return data_page_mask; } diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_impl.hpp b/cpp/src/io/parquet/experimental/hybrid_scan_impl.hpp index d61d5ca130f2..56fb387a9858 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_impl.hpp +++ b/cpp/src/io/parquet/experimental/hybrid_scan_impl.hpp @@ -389,8 +389,11 @@ class hybrid_scan_reader_impl : public parquet::detail::reader_impl { /** * @brief Compute a data page mask from the decoded page headers. + * + * @param row_mask Boolean column indicating which rows need to be read */ - [[nodiscard]] thrust::host_vector compute_data_page_mask_with_page_headers(); + [[nodiscard]] thrust::host_vector compute_data_page_mask_with_page_headers( + cudf::column_view const& row_mask); /** * @brief Mark output buffers nullable when page pruning synthesizes null rows @@ -455,14 +458,15 @@ class hybrid_scan_reader_impl : public parquet::detail::reader_impl { * @param row_group_indices Row group indices to read * @param column_chunk_data Device spans containing column chunk data, or page data when sparse * page I/O is enabled - * @param data_page_mask Input data page mask from page-pruning step + * @param row_mask Optional boolean column indicating surviving rows. `std::nullopt` indicates all + * rows are surviving */ void prepare_data(read_mode mode, std::span const> row_group_indices, std::span const> column_chunk_data, - host_span data_page_mask); + std::optional row_mask); - /** + /**row_mask * @brief Create descriptors for filter column chunks and decode dictionary page headers * * @param row_group_indices The row groups to read @@ -500,10 +504,13 @@ class hybrid_scan_reader_impl : public parquet::detail::reader_impl { * @param column_chunk_data Device spans containing column chunk data, or page data when sparse * page I/O is enabled * @param data_page_mask Input data page mask for the current pass + * @param row_mask Optional boolean column indicating surviving rows. `std::nullopt` indicates all + * rows are surviving */ void handle_chunking(read_mode mode, std::span const> column_chunk_data, - host_span data_page_mask); + host_span data_page_mask, + std::optional row_mask); /** * @brief Setup step for the next input read pass. @@ -514,9 +521,12 @@ class hybrid_scan_reader_impl : public parquet::detail::reader_impl { * @param column_chunk_data Device spans containing column chunk data, or page data when sparse * page I/O is enabled * @param data_page_mask Input data page mask for the current pass + * @param row_mask Optional boolean column indicating surviving rows. `std::nullopt` indicates all + * rows are surviving */ void setup_next_pass(std::span const> column_chunk_data, - std::span data_page_mask); + std::span data_page_mask, + std::optional row_mask); /** * @brief Setup pointers to columns chunks to be processed for this pass. @@ -624,11 +634,6 @@ class hybrid_scan_reader_impl : public parquet::detail::reader_impl { std::optional> _filter_columns_names; - // Non-owning view of the caller's row mask, only valid for the duration of a single - // materialization or chunking setup call, during which the pass page mask is computed. Null - // entries mean the row could not be pruned and is therefore treated as a surviving row. - cudf::column_view _row_mask{}; - std::vector _original_output_buffers_template; cudf::size_type _row_mask_offset{0}; diff --git a/cpp/src/io/parquet/experimental/page_index_filter_utils.cu b/cpp/src/io/parquet/experimental/page_index_filter_utils.cu index f143a8acabaf..d3a2f7f63282 100644 --- a/cpp/src/io/parquet/experimental/page_index_filter_utils.cu +++ b/cpp/src/io/parquet/experimental/page_index_filter_utils.cu @@ -480,7 +480,7 @@ thrust::host_vector compute_row_range_selection_mask( cuda::stream_ref stream) { // Need at least two offsets (or one range) to search the Fenwick tree - if (page_row_offsets.size() < 2) return thrust::host_vector{}; + if (page_row_offsets.size() < 2) { return thrust::host_vector{}; } auto const total_rows = row_mask.size(); auto const mr = cudf::get_current_device_resource_ref(); diff --git a/cpp/src/io/parquet/experimental/page_index_filter_utils.hpp b/cpp/src/io/parquet/experimental/page_index_filter_utils.hpp index 407030380883..12985a0f8a24 100644 --- a/cpp/src/io/parquet/experimental/page_index_filter_utils.hpp +++ b/cpp/src/io/parquet/experimental/page_index_filter_utils.hpp @@ -71,7 +71,7 @@ compute_page_row_offsets_and_colchunk_page_offsets( rmm::device_async_resource_ref mr); /** - * @brief Checks whether every row is reatained by the boolean row mask + * @brief Checks whether every row is retained by the boolean row mask * * Null entries in the row mask are treated as surviving rows * From 966e5ea136de82c535901d3f0d11dff04d2e5ef1 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Fri, 4 Sep 2026 23:52:49 +0000 Subject: [PATCH 17/21] minor docstring fix --- cpp/include/cudf/io/experimental/hybrid_scan_multifile.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/include/cudf/io/experimental/hybrid_scan_multifile.hpp b/cpp/include/cudf/io/experimental/hybrid_scan_multifile.hpp index f770c5e77a07..0ddf840c5a95 100644 --- a/cpp/include/cudf/io/experimental/hybrid_scan_multifile.hpp +++ b/cpp/include/cudf/io/experimental/hybrid_scan_multifile.hpp @@ -389,7 +389,7 @@ class hybrid_scan_multifile { * @param pass_read_limit Limit on the memory used for reading and decompressing data. `0` if * there is no limit * @param row_group_indices Span of vectors of input row group indices, one per source - * @param row_mask Mutable boolean column spanning all selected rows across all sources + * @param row_mask Boolean column spanning all selected rows across all sources * indicating surviving rows * @param mask_data_pages Whether to build and use a data page mask using the row mask * @param column_chunk_data Flattened device spans of filter column chunk data returned in the From 12ad95bbef781c5a20f447e4da927a158f117507 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Tue, 8 Sep 2026 18:52:17 +0000 Subject: [PATCH 18/21] style fix --- .../io/experimental/hybrid_scan_filters_test.cpp | 11 +++++------ 1 file changed, 5 insertions(+), 6 deletions(-) diff --git a/cpp/tests/io/experimental/hybrid_scan_filters_test.cpp b/cpp/tests/io/experimental/hybrid_scan_filters_test.cpp index 529ee4a7bd38..eb60d7a8df1e 100644 --- a/cpp/tests/io/experimental/hybrid_scan_filters_test.cpp +++ b/cpp/tests/io/experimental/hybrid_scan_filters_test.cpp @@ -859,10 +859,10 @@ TYPED_TEST(PageFilteringWithPageIndexStats, FilterPages) TEST_F(HybridScanFiltersTest, RowMaskNullsAreRetained) { - auto const input = cudf::test::fixed_width_column_wrapper{ - {3.0, 0.0, 1.0, 0.0}, {true, false, true, false}}; - auto const table = cudf::table_view{{input}}; - auto buffer = std::vector{}; + auto const input = cudf::test::fixed_width_column_wrapper{{3.0, 0.0, 1.0, 0.0}, + {true, false, true, false}}; + auto const table = cudf::table_view{{input}}; + auto buffer = std::vector{}; auto metadata = cudf::io::table_input_metadata{table}; metadata.column_metadata[0].set_name("col"); @@ -910,8 +910,7 @@ TEST_F(HybridScanFiltersTest, RowMaskNullsAreRetained) stream, mr); - auto const expected = - cudf::test::fixed_width_column_wrapper{{0.0, 0.0}, {false, false}}; + auto const expected = cudf::test::fixed_width_column_wrapper{{0.0, 0.0}, {false, false}}; CUDF_TEST_EXPECT_TABLES_EQUAL(cudf::table_view{{expected}}, result.tbl->view()); } From 3c8d576e22ce30675781e45e409866e05728ae1e Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Tue, 8 Sep 2026 12:03:06 -0700 Subject: [PATCH 19/21] Update python/pylibcudf/tests/io/test_experimental_hybrid_scan.py Co-authored-by: Matthew Murray <41342305+Matt711@users.noreply.github.com> --- python/pylibcudf/tests/io/test_experimental_hybrid_scan.py | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py b/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py index ca9239588d3d..daa764f0b5c4 100644 --- a/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py +++ b/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py @@ -448,10 +448,14 @@ def test_hybrid_scan_payload_page_mask_without_page_index( pa.array([i < num_selected for i in range(num_rows)], type=pa.bool_()) ) + # the caller is responsible for keeping the source bytes alive until + # synchronize_stream() below runs. + # See https://github.com/rapidsai/rmm/issues/2521 + src_bytes = simple_parquet_bytes[r.offset : r.offset + r.size] payload_data = [ plc.gpumemoryview( rmm.DeviceBuffer.to_device( - simple_parquet_bytes[r.offset : r.offset + r.size], + src_bytes, plc.utils._get_stream(), ) ) From 5330aad44b4373847db5b035b02ee47f731e9cc6 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Tue, 8 Sep 2026 19:09:53 +0000 Subject: [PATCH 20/21] ruff format --- .../tests/io/test_experimental_hybrid_scan.py | 35 +++++++++++-------- 1 file changed, 20 insertions(+), 15 deletions(-) diff --git a/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py b/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py index daa764f0b5c4..20202117001d 100644 --- a/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py +++ b/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py @@ -80,7 +80,8 @@ def simple_parquet_options( its own independent copy. """ # SourceInfo doesn't accept BytesIO, but that's fine for this test. - source = plc.io.SourceInfo([io.BytesIO(simple_parquet_bytes)]) # type: ignore[arg-type] + # type: ignore[arg-type] + source = plc.io.SourceInfo([io.BytesIO(simple_parquet_bytes)]) return plc.io.parquet.ParquetReaderOptions.builder(source).build() @@ -349,7 +350,7 @@ def test_hybrid_scan_materialize_columns( filter_data = [ plc.gpumemoryview( rmm.DeviceBuffer.to_device( - memoryview(simple_parquet_bytes)[r.offset : r.offset + r.size], + memoryview(simple_parquet_bytes)[r.offset: r.offset + r.size], plc.utils._get_stream(stream), ) ) @@ -384,7 +385,7 @@ def test_hybrid_scan_materialize_columns( payload_data = [ plc.gpumemoryview( rmm.DeviceBuffer.to_device( - memoryview(simple_parquet_bytes)[r.offset : r.offset + r.size], + memoryview(simple_parquet_bytes)[r.offset: r.offset + r.size], plc.utils._get_stream(stream), ) ) @@ -448,20 +449,23 @@ def test_hybrid_scan_payload_page_mask_without_page_index( pa.array([i < num_selected for i in range(num_rows)], type=pa.bool_()) ) - # the caller is responsible for keeping the source bytes alive until - # synchronize_stream() below runs. + # Caller is responsible for keeping the source bytes alive until + # synchronize_stream() is called below. # See https://github.com/rapidsai/rmm/issues/2521 - src_bytes = simple_parquet_bytes[r.offset : r.offset + r.size] + payload_ranges = [ + simple_parquet_bytes[r.offset: r.offset + r.size] + for r in reader.payload_column_chunks_byte_ranges( + row_groups, simple_parquet_options + ) + ] payload_data = [ plc.gpumemoryview( rmm.DeviceBuffer.to_device( - src_bytes, + src, plc.utils._get_stream(), ) ) - for r in reader.payload_column_chunks_byte_ranges( - row_groups, simple_parquet_options - ) + for src in payload_ranges ] synchronize_stream() @@ -483,6 +487,7 @@ def to_rows(tbl: plc.Table) -> list: simple_parquet_options, ) synchronize_stream() + assert to_rows(payload_result.tbl) == expected_rows reader.setup_chunking_for_payload_columns( @@ -546,7 +551,7 @@ def test_hybrid_scan_single_step_materialize( all_columns_data = [ plc.gpumemoryview( rmm.DeviceBuffer.to_device( - memoryview(simple_parquet_bytes)[r.offset : r.offset + r.size], + memoryview(simple_parquet_bytes)[r.offset: r.offset + r.size], plc.utils._get_stream(stream), ) ) @@ -628,7 +633,7 @@ def test_hybrid_scan_has_next_table_chunk( filter_data = [ plc.gpumemoryview( rmm.DeviceBuffer.to_device( - memoryview(simple_parquet_bytes)[r.offset : r.offset + r.size], + memoryview(simple_parquet_bytes)[r.offset: r.offset + r.size], plc.utils._get_stream(), ) ) @@ -698,7 +703,7 @@ def test_hybrid_scan_chunked_reading( filter_data = [ plc.gpumemoryview( rmm.DeviceBuffer.to_device( - memoryview(simple_parquet_bytes)[r.offset : r.offset + r.size], + memoryview(simple_parquet_bytes)[r.offset: r.offset + r.size], plc.utils._get_stream(stream), ) ) @@ -943,7 +948,7 @@ def prune(filter_expression: Operation) -> list[int]: # synchronize_stream() below runs. # See https://github.com/rapidsai/rmm/issues/2521 dict_page_bytes = [ - simple_parquet_bytes[r.offset : r.offset + r.size] + simple_parquet_bytes[r.offset: r.offset + r.size] for r in dictionary_ranges ] dictionary_data = [ @@ -1018,7 +1023,7 @@ def test_hybrid_scan_metadata_with_page_index( # Fetch page index bytes from the parquet file simple_parquet_mv = memoryview(simple_parquet_bytes) page_index_mv = simple_parquet_mv[ - page_index_byte_range.offset : page_index_byte_range.offset + page_index_byte_range.offset: page_index_byte_range.offset + page_index_byte_range.size ] From 7d85bfaabc7a23afc679abf152aaa40dbbb74a6c Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Tue, 8 Sep 2026 21:07:53 +0000 Subject: [PATCH 21/21] ruff-format --- .../tests/io/test_experimental_hybrid_scan.py | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) diff --git a/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py b/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py index 20202117001d..3f3a92ebf441 100644 --- a/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py +++ b/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py @@ -350,7 +350,7 @@ def test_hybrid_scan_materialize_columns( filter_data = [ plc.gpumemoryview( rmm.DeviceBuffer.to_device( - memoryview(simple_parquet_bytes)[r.offset: r.offset + r.size], + memoryview(simple_parquet_bytes)[r.offset : r.offset + r.size], plc.utils._get_stream(stream), ) ) @@ -385,7 +385,7 @@ def test_hybrid_scan_materialize_columns( payload_data = [ plc.gpumemoryview( rmm.DeviceBuffer.to_device( - memoryview(simple_parquet_bytes)[r.offset: r.offset + r.size], + memoryview(simple_parquet_bytes)[r.offset : r.offset + r.size], plc.utils._get_stream(stream), ) ) @@ -453,7 +453,7 @@ def test_hybrid_scan_payload_page_mask_without_page_index( # synchronize_stream() is called below. # See https://github.com/rapidsai/rmm/issues/2521 payload_ranges = [ - simple_parquet_bytes[r.offset: r.offset + r.size] + simple_parquet_bytes[r.offset : r.offset + r.size] for r in reader.payload_column_chunks_byte_ranges( row_groups, simple_parquet_options ) @@ -551,7 +551,7 @@ def test_hybrid_scan_single_step_materialize( all_columns_data = [ plc.gpumemoryview( rmm.DeviceBuffer.to_device( - memoryview(simple_parquet_bytes)[r.offset: r.offset + r.size], + memoryview(simple_parquet_bytes)[r.offset : r.offset + r.size], plc.utils._get_stream(stream), ) ) @@ -633,7 +633,7 @@ def test_hybrid_scan_has_next_table_chunk( filter_data = [ plc.gpumemoryview( rmm.DeviceBuffer.to_device( - memoryview(simple_parquet_bytes)[r.offset: r.offset + r.size], + memoryview(simple_parquet_bytes)[r.offset : r.offset + r.size], plc.utils._get_stream(), ) ) @@ -703,7 +703,7 @@ def test_hybrid_scan_chunked_reading( filter_data = [ plc.gpumemoryview( rmm.DeviceBuffer.to_device( - memoryview(simple_parquet_bytes)[r.offset: r.offset + r.size], + memoryview(simple_parquet_bytes)[r.offset : r.offset + r.size], plc.utils._get_stream(stream), ) ) @@ -948,7 +948,7 @@ def prune(filter_expression: Operation) -> list[int]: # synchronize_stream() below runs. # See https://github.com/rapidsai/rmm/issues/2521 dict_page_bytes = [ - simple_parquet_bytes[r.offset: r.offset + r.size] + simple_parquet_bytes[r.offset : r.offset + r.size] for r in dictionary_ranges ] dictionary_data = [ @@ -1023,7 +1023,7 @@ def test_hybrid_scan_metadata_with_page_index( # Fetch page index bytes from the parquet file simple_parquet_mv = memoryview(simple_parquet_bytes) page_index_mv = simple_parquet_mv[ - page_index_byte_range.offset: page_index_byte_range.offset + page_index_byte_range.offset : page_index_byte_range.offset + page_index_byte_range.size ]