From cf9226a49f96a635e6ea6e5574199a6a89182685 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb Date: Tue, 21 Jul 2026 17:39:20 +0000 Subject: [PATCH 01/10] Add page topology support for sparse reads Preserve page locations and variable-width offset state needed to safely reconstruct columns from a sparse subset of Parquet data pages. --- cpp/src/io/parquet/page_hdr.cu | 100 +++++++++++++++ cpp/src/io/parquet/parquet_gpu.hpp | 13 ++ cpp/src/io/parquet/reader_impl.cpp | 3 + cpp/src/io/parquet/reader_impl.hpp | 5 + cpp/src/io/parquet/reader_impl_helpers.cpp | 42 +++---- cpp/src/io/parquet/reader_impl_preprocess.cu | 114 ++++++++++++++++-- .../parquet/reader_impl_preprocess_utils.cu | 50 +++++++- .../parquet/reader_impl_preprocess_utils.cuh | 35 ++++-- 8 files changed, 312 insertions(+), 50 deletions(-) diff --git a/cpp/src/io/parquet/page_hdr.cu b/cpp/src/io/parquet/page_hdr.cu index 2565dab3ae31..147b5d3568b5 100644 --- a/cpp/src/io/parquet/page_hdr.cu +++ b/cpp/src/io/parquet/page_hdr.cu @@ -823,6 +823,86 @@ struct decode_page_headers_with_pgidx_fn { } }; +/** + * @brief Functor to decode indexed page headers from exact page spans + */ +struct decode_page_headers_with_pgidx_spans_fn { + cudf::device_span colchunks; + cudf::device_span pages; + cudf::device_span const> page_spans; + size_type* chunk_page_offsets; + kernel_error::pointer error_code; + + __device__ void operator()(size_type page_idx) const noexcept + { + auto const num_chunks = static_cast(colchunks.size()); + auto const chunk_idx = static_cast( + cuda::std::distance( + chunk_page_offsets, + thrust::upper_bound( + thrust::seq, chunk_page_offsets, chunk_page_offsets + num_chunks + 1, page_idx)) - + 1); + + if (chunk_idx < 0 or chunk_idx >= num_chunks) { + set_error(static_cast(decode_error::DATA_STREAM_OVERRUN), + error_code); + return; + } + + byte_stream_s bs{}; + bs.ck = colchunks[chunk_idx]; + zero_out_page_header_info(&bs); + bs.page.chunk_idx = chunk_idx; + bs.page.src_col_schema = bs.ck.src_col_schema; + + auto const span = page_spans[page_idx]; + if (span.empty()) { + // Preserve the logical page entry. Page-index metadata is filled in by fill_in_page_info(). + pages[page_idx] = bs.page; + return; + } + + bs.base = bs.cur = span.data(); + bs.end = span.data() + span.size(); + + if (not parse_valid_page_header(&bs)) { + set_error(static_cast(decode_error::INVALID_PAGE_HEADER), + error_code); + return; + } + if (not is_supported_encoding(bs.page.encoding)) { + set_error(static_cast(decode_error::UNSUPPORTED_ENCODING), + error_code); + return; + } + + switch (bs.page_type) { + case PageType::DATA_PAGE: bs.page.num_rows = bs.page.num_input_values; break; + case PageType::DATA_PAGE_V2: + bs.page.flags |= PAGEINFO_FLAGS_V2; + bs.page.definition_level_encoding = Encoding::RLE; + bs.page.repetition_level_encoding = Encoding::RLE; + break; + case PageType::DICTIONARY_PAGE: bs.page.flags |= PAGEINFO_FLAGS_DICTIONARY; break; + default: + set_error(static_cast(decode_error::INVALID_PAGE_TYPE), + error_code); + return; + } + + if (bs.page.compressed_page_size < 0 or + static_cast(bs.end - bs.cur) != static_cast(bs.page.compressed_page_size)) { + set_error(static_cast(decode_error::DATA_STREAM_OVERRUN), + error_code); + return; + } + + bs.page.page_data = const_cast(bs.cur); + bs.page.kernel_mask = kernel_mask_for_page(bs.page, bs.ck); + pages[page_idx] = bs.page; + } +}; + /** * @brief Kernel for building dictionary index for the specified column chunks * @@ -954,6 +1034,26 @@ void decode_page_headers_with_pgidx(cudf::device_span chu .error_code = error_code}); } +void decode_page_headers_with_pgidx_spans(cudf::device_span chunks, + cudf::device_span pages, + cudf::device_span const> + page_spans, + size_type* chunk_page_offsets, + kernel_error::pointer error_code, + rmm::cuda_stream_view stream) +{ + CUDF_EXPECTS(page_spans.size() == pages.size(), + "Page span count must match the number of logical pages"); + thrust::for_each(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + cuda::counting_iterator{0}, + cuda::counting_iterator{static_cast(pages.size())}, + decode_page_headers_with_pgidx_spans_fn{.colchunks = chunks, + .pages = pages, + .page_spans = page_spans, + .chunk_page_offsets = chunk_page_offsets, + .error_code = error_code}); +} + void build_string_dictionary_index(ColumnChunkDesc* chunks, int32_t num_chunks, kernel_error::pointer error_code, diff --git a/cpp/src/io/parquet/parquet_gpu.hpp b/cpp/src/io/parquet/parquet_gpu.hpp index 090431058c47..c821f14e2b9e 100644 --- a/cpp/src/io/parquet/parquet_gpu.hpp +++ b/cpp/src/io/parquet/parquet_gpu.hpp @@ -728,6 +728,19 @@ void decode_page_headers_with_pgidx(cudf::device_span chu kernel_error::pointer error_code, rmm::cuda_stream_view stream); +/** + * @brief Decode indexed page headers from exact, potentially discontiguous page spans + * + * Empty spans initialize the corresponding logical page descriptor but are not parsed. + */ +void decode_page_headers_with_pgidx_spans(cudf::device_span chunks, + cudf::device_span pages, + cudf::device_span const> + page_spans, + size_type* chunk_page_offsets, + kernel_error::pointer error_code, + rmm::cuda_stream_view stream); + /** * @brief Launches kernel for building the dictionary index for the column * chunks diff --git a/cpp/src/io/parquet/reader_impl.cpp b/cpp/src/io/parquet/reader_impl.cpp index 38ee9301d03f..bcad8cdccc40 100644 --- a/cpp/src/io/parquet/reader_impl.cpp +++ b/cpp/src/io/parquet/reader_impl.cpp @@ -453,6 +453,9 @@ void reader_impl::decode_page_data(read_mode mode, size_t skip_rows, size_t num_ cudf::detail::make_pinned_vector(cudf::host_span{out_buffers}, _stream); write_final_offsets(pinned_final_offsets, pinned_out_buffers, _stream); + // For page-level I/O, fill output string and list offsets for pruned pages + fill_pruned_offsets(skip_rows, num_rows); + // update null counts in the final column buffers for (size_t idx = 0; idx < subpass.pages.size(); idx++) { PageInfo* pi = &subpass.pages[idx]; diff --git a/cpp/src/io/parquet/reader_impl.hpp b/cpp/src/io/parquet/reader_impl.hpp index 61f85e047809..d3033518389e 100644 --- a/cpp/src/io/parquet/reader_impl.hpp +++ b/cpp/src/io/parquet/reader_impl.hpp @@ -350,6 +350,11 @@ class reader_impl { size_t skip_rows, size_t num_rows); + /** + * @brief Fill in string and list offsets for rows covered by pruned data pages. + */ + void fill_pruned_offsets(size_t skip_rows, size_t num_rows); + /** * @brief Creates file-wide parquet chunk information. * diff --git a/cpp/src/io/parquet/reader_impl_helpers.cpp b/cpp/src/io/parquet/reader_impl_helpers.cpp index ce06969cf0da..d25d48b2b876 100644 --- a/cpp/src/io/parquet/reader_impl_helpers.cpp +++ b/cpp/src/io/parquet/reader_impl_helpers.cpp @@ -719,14 +719,13 @@ void aggregate_reader_metadata::column_info_for_row_group(row_group_info& rg_inf auto const max_def_level = schema.max_definition_level; auto const max_rep_level = schema.max_repetition_level; - // If any columns lack the page indexes then just return without modifying the - // row_group_info. - if (not col_chunk.offset_index.has_value() or not col_chunk.column_index.has_value()) { - return; - } + // Page locations and row boundaries only require the offset index. Additional value-count + // metadata is populated below when a column index is also available. + if (not col_chunk.offset_index.has_value()) { return; } auto const& offset_index = col_chunk.offset_index.value(); - auto const& column_index = col_chunk.column_index.value(); + auto const* column_index = + col_chunk.column_index.has_value() ? &col_chunk.column_index.value() : nullptr; auto& chunk_info = chunks[col_idx]; auto const num_pages = offset_index.page_locations.size(); @@ -757,12 +756,14 @@ void aggregate_reader_metadata::column_info_for_row_group(row_group_info& rg_inf // definition_level_histogram is absent. // // In the future we might want the full histograms saved in the `column_info` struct. - int64_t const* const def_hist = column_index.definition_level_histogram.has_value() - ? column_index.definition_level_histogram.value().data() - : nullptr; - int64_t const* const rep_hist = column_index.repetition_level_histogram.has_value() - ? column_index.repetition_level_histogram.value().data() - : nullptr; + int64_t const* const def_hist = + column_index != nullptr and column_index->definition_level_histogram.has_value() + ? column_index->definition_level_histogram.value().data() + : nullptr; + int64_t const* const rep_hist = + column_index != nullptr and column_index->repetition_level_histogram.has_value() + ? column_index->repetition_level_histogram.value().data() + : nullptr; for (size_t pg_idx = 0; pg_idx < num_pages; pg_idx++) { auto const& page_loc = offset_index.page_locations[pg_idx]; @@ -777,8 +778,8 @@ void aggregate_reader_metadata::column_info_for_row_group(row_group_info& rg_inf page_info pg_info{.location = page_loc, .num_rows = num_rows}; // check to see if we already have null counts for each page - if (column_index.null_counts.has_value()) { - pg_info.num_nulls = column_index.null_counts.value()[pg_idx]; + if (column_index != nullptr and column_index->null_counts.has_value()) { + pg_info.num_nulls = column_index->null_counts.value()[pg_idx]; } // save variable length byte info if present @@ -820,19 +821,6 @@ void aggregate_reader_metadata::column_info_for_row_group(row_group_info& rg_inf } } - // If none of the ifs above triggered, then we have neither histogram (likely the writer - // doesn't produce them, the r:0 d:1 case should have been handled above). The column index - // doesn't give us value counts, so we'll have to rely on the page headers. If the histogram - // info is missing or insufficient, then just return without modifying the row_group_info. - if (not pg_info.num_nulls.has_value() or not pg_info.num_valid.has_value()) { return; } - - // Like above, if using older page indexes that lack size info, then return without modifying - // the row_group_info. - // TODO: cudf will still set the per-page var_bytes to '0' even for all null pages. Need to - // check the behavior of other implementations (once there are some). Some may not set the - // var bytes for all null pages, so check the `null_pages` field on the column index. - if (schema.type == Type::BYTE_ARRAY and not pg_info.var_bytes_size.has_value()) { return; } - chunk_info.pages.push_back(std::move(pg_info)); } } diff --git a/cpp/src/io/parquet/reader_impl_preprocess.cu b/cpp/src/io/parquet/reader_impl_preprocess.cu index 3eddddf8d0c0..000c18554d05 100644 --- a/cpp/src/io/parquet/reader_impl_preprocess.cu +++ b/cpp/src/io/parquet/reader_impl_preprocess.cu @@ -17,6 +17,7 @@ #include #include #include +#include #include #include #include @@ -35,6 +36,7 @@ #include #include #include +#include #include namespace cudf::io::parquet::detail { @@ -51,7 +53,44 @@ inline bool is_treat_fixed_length_as_string(cuda::std::optional con } struct set_str_bytes_all { - __device__ void operator()(PageInfo& p) { p.str_bytes_all = p.str_bytes; } + device_span pages; + device_span page_mask; + + __device__ void operator()(size_type index) const + { + pages[index].str_bytes_all = + page_mask.empty() or page_mask[index] ? pages[index].str_bytes : int32_t{0}; + } +}; + +struct set_pruned_string_offsets { + device_span pages; + device_span chunks; + device_span page_mask; + size_t skip_rows; + size_t num_rows; + + __device__ void operator()(size_type index) const + { + if (page_mask[index]) { return; } + auto const& page = pages[index]; + auto const& chunk = chunks[page.chunk_idx]; + if (chunk.max_level[level_type::REPETITION] != 0 or not is_string_col(chunk) or + chunk.is_large_string_col or chunk.column_data_base == nullptr) { + return; + } + auto offsets = static_cast(chunk.column_data_base[chunk.max_nesting_depth - 1]); + if (offsets == nullptr) { return; } + + auto const page_begin = chunk.start_row + page.chunk_row; + auto const page_end = page_begin + page.num_rows; + auto const read_end = skip_rows + num_rows; + auto const begin = cuda::std::max(page_begin, skip_rows); + auto const end = cuda::std::min(page_end, read_end); + for (auto row = begin; row < end; ++row) { + offsets[row - skip_rows] = static_cast(page.str_offset); + } + } }; } // namespace @@ -847,10 +886,11 @@ void reader_impl::preprocess_subpass_pages(read_mode mode, size_t chunk_read_lim _stream); } // set str_bytes_all - thrust::for_each(rmm::exec_policy_nosync(_stream, cudf::get_current_device_resource_ref()), - subpass.pages.device_begin(), - subpass.pages.device_end(), - set_str_bytes_all{}); + thrust::for_each( + rmm::exec_policy_nosync(_stream, cudf::get_current_device_resource_ref()), + cuda::counting_iterator{0}, + cuda::counting_iterator{static_cast(subpass.pages.size())}, + set_str_bytes_all{subpass.pages, subpass_page_mask_span()}); } // retrieve pages back @@ -932,6 +972,10 @@ void reader_impl::allocate_columns(read_mode mode, size_t skip_rows, size_t num_ bool has_lists = false; // Validity Buffer is a uint32_t pointer std::vector> nullmask_bufs; + auto const page_mask = subpass_page_mask_span(); + auto const has_pruned_page = + not page_mask.is_empty() and + std::any_of(page_mask.host_begin(), page_mask.host_end(), [](bool keep) { return not keep; }); for (auto const& input_col : _input_columns) { size_t const max_depth = input_col.nesting_depth(); @@ -955,8 +999,10 @@ void reader_impl::allocate_columns(read_mode mode, size_t skip_rows, size_t num_ CUDF_EXPECTS(out_buf_size <= std::numeric_limits::max(), "Number of rows exceeds cudf's column size limit", std::overflow_error); + auto const initialize_offsets = has_pruned_page and (out_buf.type.id() == type_id::STRING or + out_buf.type.id() == type_id::LIST); out_buf.create_with_mask( - out_buf_size, cudf::mask_state::UNINITIALIZED, false, _stream, _mr); + out_buf_size, cudf::mask_state::UNINITIALIZED, initialize_offsets, _stream, _mr); nullmask_bufs.emplace_back( out_buf.null_mask(), cudf::util::round_up_safe(out_buf.null_mask_size(), sizeof(cudf::bitmask_type)) / @@ -1074,8 +1120,11 @@ void reader_impl::allocate_columns(read_mode mode, size_t skip_rows, size_t num_ std::overflow_error); // allocate // we're going to start null mask as all valid and then turn bits off if necessary + auto const initialize_offsets = + has_pruned_page and + (out_buf.type.id() == type_id::STRING or out_buf.type.id() == type_id::LIST); out_buf.create_with_mask( - buffer_size, cudf::mask_state::UNINITIALIZED, false, _stream, _mr); + buffer_size, cudf::mask_state::UNINITIALIZED, initialize_offsets, _stream, _mr); nullmask_bufs.emplace_back( out_buf.null_mask(), cudf::util::round_up_safe(out_buf.null_mask_size(), sizeof(cudf::bitmask_type)) / @@ -1092,6 +1141,57 @@ void reader_impl::allocate_columns(read_mode mode, size_t skip_rows, size_t num_ pinned_nullmask_bufs, std::numeric_limits::max(), _stream); } +void reader_impl::fill_pruned_offsets(size_t skip_rows, size_t num_rows) +{ + auto& pass = *_pass_itm_data; + auto& subpass = *pass.subpass; + auto const page_mask = subpass_page_mask_span(); + if (page_mask.is_empty() or + std::all_of(page_mask.host_begin(), page_mask.host_end(), cuda::std::identity{})) { + return; + } + + auto const pages = device_span{subpass.pages.device_ptr(), subpass.pages.size()}; + auto const chunks = + device_span{pass.chunks.device_ptr(), pass.chunks.size()}; + auto const device_page_mask = static_cast>(page_mask); + thrust::for_each_n( + rmm::exec_policy_nosync(_stream, cudf::get_current_device_resource_ref()), + cuda::counting_iterator{0}, + static_cast(pages.size()), + set_pruned_string_offsets{pages, chunks, device_page_mask, skip_rows, num_rows}); + + auto offset_buffers = std::vector>{}; + auto collect_offsets = [&](auto&& self, auto& buffer) -> void { + auto const is_small_string = + buffer.type.id() == type_id::STRING and not buffer.is_large_strings_column(); + if (is_small_string or buffer.type.id() == type_id::LIST) { + offset_buffers.emplace_back(static_cast(buffer.data()), + buffer.size + (is_small_string ? 1 : 0)); + } + for (auto& child : buffer.children) { + self(self, child); + } + }; + for (auto& buffer : _output_buffers) { + collect_offsets(collect_offsets, buffer); + } + if (offset_buffers.empty()) { return; } + + auto const num_streams = std::min(offset_buffers.size(), 4); + auto const streams = cudf::detail::fork_streams(_stream, num_streams); + for (auto index = std::size_t{0}; index < offset_buffers.size(); ++index) { + auto const [offsets, num_items] = offset_buffers[index]; + auto const stream = streams[index % streams.size()]; + thrust::inclusive_scan(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + offsets, + offsets + num_items, + offsets, + cuda::maximum{}); + } + cudf::detail::join_streams(streams, _stream); +} + cudf::detail::host_vector reader_impl::calculate_page_string_offsets() { auto& pass = *_pass_itm_data; diff --git a/cpp/src/io/parquet/reader_impl_preprocess_utils.cu b/cpp/src/io/parquet/reader_impl_preprocess_utils.cu index a36f62f7af79..25af24a9487f 100644 --- a/cpp/src/io/parquet/reader_impl_preprocess_utils.cu +++ b/cpp/src/io/parquet/reader_impl_preprocess_utils.cu @@ -286,6 +286,11 @@ void fill_in_page_info(host_span chunks, page.num_nulls = chunk_info.pages[p].num_nulls.value_or(0); page.num_valids = chunk_info.pages[p].num_valid.value_or(0); page.str_bytes = chunk_info.pages[p].var_bytes_size.value_or(0); + page.has_value_info = + chunk_info.pages[p].num_nulls.has_value() and + chunk_info.pages[p].num_valid.has_value() and + (chunk.physical_type != Type::BYTE_ARRAY or + chunk_info.pages[p].var_bytes_size.has_value()); start_row += page.num_rows; } @@ -406,10 +411,13 @@ cudf::detail::hostdevice_vector sort_pages(device_span return pass_pages; } -void decode_page_headers(pass_intermediate_data& pass, - device_span unsorted_pages, - bool has_page_index, - rmm::cuda_stream_view stream) +namespace { + +void decode_page_headers_impl(pass_intermediate_data& pass, + device_span unsorted_pages, + bool has_page_index, + host_span const> page_spans, + rmm::cuda_stream_view stream) { CUDF_FUNC_RANGE(); @@ -439,9 +447,23 @@ void decode_page_headers(pass_intermediate_data& pass, kernel_error error_code(stream); + if (not page_spans.empty()) { + CUDF_EXPECTS(has_page_index, "Sparse page spans require Parquet page indexes"); + CUDF_EXPECTS(page_spans.size() == unsorted_pages.size(), + "Page span count must match the number of logical pages"); + auto device_page_spans = cudf::detail::make_device_uvector_async( + page_spans, stream, cudf::get_current_device_resource_ref()); + decode_page_headers_with_pgidx_spans( + device_span(pass.chunks.device_ptr(), pass.chunks.size()), + unsorted_pages, + device_page_spans, + chunk_page_offsets.begin(), + error_code.data(), + stream); + } // If page index is present, collect data ptrs for all pages and launch the accelerated decode // page headers kernel - if (has_page_index) { + else if (has_page_index) { auto host_page_locations = cudf::detail::make_pinned_vector_async(unsorted_pages.size(), stream); auto curr_page_idx = 0; @@ -576,4 +598,22 @@ void decode_page_headers(pass_intermediate_data& pass, stream.synchronize(); } +} // namespace + +void decode_page_headers(pass_intermediate_data& pass, + device_span unsorted_pages, + bool has_page_index, + rmm::cuda_stream_view stream) +{ + decode_page_headers_impl(pass, unsorted_pages, has_page_index, {}, stream); +} + +void decode_page_headers(pass_intermediate_data& pass, + device_span unsorted_pages, + host_span const> page_spans, + rmm::cuda_stream_view stream) +{ + decode_page_headers_impl(pass, unsorted_pages, true, page_spans, stream); +} + } // namespace cudf::io::parquet::detail diff --git a/cpp/src/io/parquet/reader_impl_preprocess_utils.cuh b/cpp/src/io/parquet/reader_impl_preprocess_utils.cuh index 3a0f54cd5cd6..b9183ad5351d 100644 --- a/cpp/src/io/parquet/reader_impl_preprocess_utils.cuh +++ b/cpp/src/io/parquet/reader_impl_preprocess_utils.cuh @@ -136,6 +136,16 @@ void decode_page_headers(pass_intermediate_data& pass, bool has_page_index, rmm::cuda_stream_view stream); +/** + * @brief Decode page information using one exact span per logical indexed page + * + * Empty spans represent masked pages and retain their logical page-index metadata. + */ +void decode_page_headers(pass_intermediate_data& pass, + device_span unsorted_pages, + host_span const> page_spans, + rmm::cuda_stream_view stream); + /** * @brief Check if the column chunk has a string (byte array or FLBA) type */ @@ -157,6 +167,7 @@ struct page_index_info { int32_t num_nulls; int32_t num_valids; int32_t str_bytes; + bool has_value_info; }; /** @@ -168,17 +179,19 @@ struct copy_page_info { __device__ constexpr void operator()(size_type idx) { - auto& pg = pages[idx]; - auto const& pi = page_indexes[idx]; - pg.num_rows = pi.num_rows; - pg.chunk_row = pi.chunk_row; - pg.has_page_index = true; - pg.num_nulls = pi.num_nulls; - pg.num_valids = pi.num_valids; - pg.str_bytes_from_index = pi.str_bytes; - pg.str_bytes = pi.str_bytes; - pg.start_val = 0; - pg.end_val = pg.num_valids; + auto& pg = pages[idx]; + auto const& pi = page_indexes[idx]; + pg.num_rows = pi.num_rows; + pg.chunk_row = pi.chunk_row; + pg.has_page_index = pi.has_value_info; + pg.start_val = 0; + if (pi.has_value_info) { + pg.num_nulls = pi.num_nulls; + pg.num_valids = pi.num_valids; + pg.str_bytes_from_index = pi.str_bytes; + pg.str_bytes = pi.str_bytes; + pg.end_val = pg.num_valids; + } } }; From f356b0d094c6f260c4d437463b4bc714e2189802 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb Date: Tue, 21 Jul 2026 17:39:59 +0000 Subject: [PATCH 02/10] Add sparse page I/O to hybrid scan Expose multifile page-range planning and consume the selected payload pages so hybrid scan avoids fetching pruned Parquet payload data. --- .../io/experimental/hybrid_scan_multifile.hpp | 48 ++ .../experimental/hybrid_scan_chunking.cu | 11 +- .../parquet/experimental/hybrid_scan_impl.cpp | 401 ++++++++++++- .../parquet/experimental/hybrid_scan_impl.hpp | 47 +- .../experimental/hybrid_scan_multifile.cpp | 36 ++ .../experimental/hybrid_scan_preprocess.cu | 37 ++ .../io/experimental/hybrid_scan_common.cpp | 9 + .../io/experimental/hybrid_scan_common.hpp | 9 + .../hybrid_scan_multifile_composer.cpp | 67 +++ .../hybrid_scan_multifile_composer.hpp | 24 + .../hybrid_scan_multifile_test.cpp | 525 +++++++++++++++++- 11 files changed, 1200 insertions(+), 14 deletions(-) diff --git a/cpp/include/cudf/io/experimental/hybrid_scan_multifile.hpp b/cpp/include/cudf/io/experimental/hybrid_scan_multifile.hpp index 1be7365c53f4..691e7ee6f516 100644 --- a/cpp/include/cudf/io/experimental/hybrid_scan_multifile.hpp +++ b/cpp/include/cudf/io/experimental/hybrid_scan_multifile.hpp @@ -260,6 +260,27 @@ class hybrid_scan_multifile { payload_column_chunks_byte_ranges(cudf::host_span const> row_group_indices, parquet_reader_options const& options) const; + /** + * @brief Plan page-level payload byte ranges grouped by source + * + * When page masking cannot be used, returns the legacy full-column-chunk ranges regrouped by + * source. The resulting plan must be consumed exactly once by the matching page-data setup + * overload. + * + * @param row_group_indices Input row group indices, one vector per source + * @param row_mask Boolean mask spanning the selected row groups + * @param mask_data_pages Whether to use the row mask to prune data pages + * @param options Parquet reader options + * @param stream CUDA stream used to compute the page mask + * @return Byte ranges to fetch, grouped by source + */ + [[nodiscard]] std::vector> payload_column_chunks_byte_ranges( + cudf::host_span const> row_group_indices, + cudf::column_view const& row_mask, + use_data_page_mask mask_data_pages, + parquet_reader_options const& options, + rmm::cuda_stream_view stream) const; + /** * @brief Materialize payload columns and applies the row mask to the output table * @@ -383,6 +404,33 @@ class hybrid_scan_multifile { rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) const; + /** + * @brief Setup payload chunking from source-grouped page-level fetch results + * + * Consumes the pending plan created by the page-level payload byte-range overload. Each input + * span must correspond to the byte range at the same source and range index. + * + * @param chunk_read_limit Maximum bytes returned per output table chunk, or zero + * @param pass_read_limit Maximum read/decompression memory, or zero + * @param row_group_indices Input row group indices, one vector per source + * @param row_mask Boolean mask spanning the selected row groups + * @param mask_data_pages Whether page masking was requested + * @param page_data_per_source Fetched device spans grouped by source + * @param options Parquet reader options + * @param stream CUDA stream used for preprocessing + * @param mr Device memory resource used for output table chunks + */ + void setup_chunking_for_payload_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, + use_data_page_mask mask_data_pages, + cudf::host_span> const> page_data_per_source, + parquet_reader_options const& options, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) const; + /** * @brief Materializes a chunk of payload columns and applies the corresponding range of input row * mask to the output table chunk diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu b/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu index 0f846261c35b..295b06d7d402 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu +++ b/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu @@ -38,10 +38,7 @@ void hybrid_scan_reader_impl::handle_chunking( // 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); - - // Must be called as soon as we create the pass - set_pass_page_mask(data_page_mask); + setup_next_pass(column_chunk_data, data_page_mask); } auto& pass = *_pass_itm_data; @@ -78,7 +75,8 @@ void hybrid_scan_reader_impl::handle_chunking( } void hybrid_scan_reader_impl::setup_next_pass( - std::span const> column_chunk_data) + std::span const> column_chunk_data, + host_span data_page_mask) { auto const num_passes = _file_itm_data.num_passes(); CUDF_EXPECTS(num_passes == 1, @@ -122,6 +120,9 @@ void hybrid_scan_reader_impl::setup_next_pass( // Setup page information for the chunk (which we can access without decompressing) setup_compressed_data(column_chunk_data); + // Establish the logical mask before malformed-page checks, size estimation, or subpass setup. + set_pass_page_mask(data_page_mask); + // detect malformed columns. // - we have seen some cases in the wild where we have a row group containing N // rows, but the total number of rows in the pages for column X is != N. while it diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp b/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp index 98353b88f432..631aa4197bfd 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp +++ b/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp @@ -24,8 +24,11 @@ #include #include +#include #include +#include #include +#include #include namespace cudf::io::parquet::experimental::detail { @@ -216,6 +219,8 @@ std::size_t hybrid_scan_reader_impl::total_rows_in_row_groups( void hybrid_scan_reader_impl::reset_column_selection() { + CUDF_EXPECTS(not _pending_payload_page_io_plan.has_value(), + "Cannot reset column selection while a payload page I/O plan is pending"); _is_all_columns_selected = false; _is_filter_columns_selected = false; _is_payload_columns_selected = false; @@ -241,6 +246,8 @@ void hybrid_scan_reader_impl::prepare_materialization(read_columns_mode read_col rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) { + CUDF_EXPECTS(not _pending_payload_page_io_plan.has_value(), + "Pending payload page I/O plan must be consumed by its setup overload"); reset_internal_state(); initialize_options(options, num_sources, stream, mr); select_columns(read_columns_mode, options); @@ -495,6 +502,264 @@ hybrid_scan_reader_impl::payload_column_chunks_byte_ranges( return get_input_column_chunk_byte_ranges(row_group_indices); } +std::vector> +hybrid_scan_reader_impl::payload_column_chunks_byte_ranges( + std::span const> row_group_indices, + cudf::column_view const& row_mask, + use_data_page_mask mask_data_pages, + parquet_reader_options const& options, + rmm::cuda_stream_view stream) +{ + CUDF_EXPECTS(row_group_indices.size() == _extended_metadata->get_num_sources(), + "Row group source count must match the number of input sources"); + CUDF_EXPECTS(std::cmp_equal(row_mask.size(), total_rows_in_row_groups(row_group_indices)), + "Row mask must span across all input row groups"); + CUDF_EXPECTS(row_mask.null_count() == 0, + "Row mask must not have any nulls when planning payload pages"); + CUDF_EXPECTS(not _pending_payload_page_io_plan.has_value(), + "The previous payload page I/O plan has not been consumed"); + + select_columns(read_columns_mode::PAYLOAD_COLUMNS, options); + + auto column_schemas = std::vector{}; + column_schemas.reserve(_input_columns.size()); + std::transform(_input_columns.begin(), + _input_columns.end(), + std::back_inserter(column_schemas), + [](auto const& col) { return col.schema_idx; }); + + auto make_full_chunk_plan = [&]() { + auto [flat_ranges, source_map] = get_input_column_chunk_byte_ranges(row_group_indices); + auto source_ranges = std::vector>(row_group_indices.size()); + CUDF_EXPECTS(flat_ranges.size() == source_map.size(), + "Column chunk range source map is invalid"); + for (std::size_t i = 0; i < flat_ranges.size(); ++i) { + CUDF_EXPECTS(std::cmp_less(source_map[i], source_ranges.size()), + "Column chunk range has an invalid source index"); + source_ranges[source_map[i]].push_back(flat_ranges[i]); + } + + _pending_payload_page_io_plan = payload_page_io_plan{ + .sparse = false, + .mask_data_pages = mask_data_pages, + .row_group_indices = {row_group_indices.begin(), row_group_indices.end()}, + .column_schema_indices = column_schemas, + .source_ranges = source_ranges, + .page_mappings = {}, + .resident_bytes_per_chunk = {}, + .dictionary_present_per_chunk = {}, + .data_page_mask = {}}; + return source_ranges; + }; + + if (mask_data_pages == use_data_page_mask::NO or row_mask.is_empty()) { + return make_full_chunk_plan(); + } + + // Sparse page planning only requires offset-index topology. Value counts and variable-width + // sizes can be derived from each retained page after it is fetched. + auto indexes_complete = true; + for (std::size_t source_idx = 0; source_idx < row_group_indices.size(); ++source_idx) { + for (auto const row_group_idx : row_group_indices[source_idx]) { + auto const& row_group = _extended_metadata->get_row_group(row_group_idx, source_idx); + for (auto const schema_idx : column_schemas) { + auto const candidate_it = std::find_if( + row_group.columns.begin(), row_group.columns.end(), [schema_idx](auto const& candidate) { + return candidate.schema_idx == schema_idx; + }); + if (candidate_it == row_group.columns.end() or + not candidate_it->offset_index.has_value()) { + indexes_complete = false; + break; + } + auto const& candidate = *candidate_it; + auto const& oi = candidate.offset_index.value(); + auto const num_pages = oi.page_locations.size(); + auto const index_vector_sizes_valid = + not oi.unencoded_byte_array_data_bytes.has_value() or + oi.unencoded_byte_array_data_bytes->size() == num_pages; + auto const dictionary_offsets_valid = + candidate.meta_data.dictionary_page_offset <= 0 or + candidate.meta_data.data_page_offset > candidate.meta_data.dictionary_page_offset; + auto const page_rows_valid = + num_pages > 0 and oi.page_locations.front().first_row_index == 0 and + std::is_sorted(oi.page_locations.begin(), + oi.page_locations.end(), + [](auto const& lhs, auto const& rhs) { + return lhs.first_row_index < rhs.first_row_index; + }) and + std::all_of( + oi.page_locations.begin(), oi.page_locations.end(), [&](auto const& location) { + return location.first_row_index >= 0 and + std::cmp_less_equal(location.first_row_index, row_group.num_rows); + }); + if (num_pages == 0 or not index_vector_sizes_valid or not page_rows_valid or + candidate.meta_data.data_page_offset <= 0 or not dictionary_offsets_valid or + std::any_of( + oi.page_locations.begin(), oi.page_locations.end(), [](auto const& location) { + return location.offset < 0 or location.compressed_page_size <= 0; + })) { + indexes_complete = false; + break; + } + } + if (not indexes_complete) { break; } + } + if (not indexes_complete) { break; } + } + if (not indexes_complete) { return make_full_chunk_plan(); } + + auto data_page_mask = _extended_metadata->compute_data_page_mask( + row_mask, row_group_indices, _input_columns, 0, stream); + // An empty mask is the established representation for "all pages retained". + if (data_page_mask.empty()) { return make_full_chunk_plan(); } + + auto const num_columns = _input_columns.size(); + auto const num_row_groups = + std::accumulate(row_group_indices.begin(), + row_group_indices.end(), + std::size_t{0}, + [](auto sum, auto const& groups) { return sum + groups.size(); }); + auto const num_chunks = num_row_groups * num_columns; + auto chunk_masks = std::vector>(num_chunks); + + // Translate the column-major mask into source-major/row-group-major chunk slots once. + std::size_t mask_idx = 0; + for (std::size_t col_idx = 0; col_idx < num_columns; ++col_idx) { + std::size_t row_group_ordinal = 0; + for (std::size_t source_idx = 0; source_idx < row_group_indices.size(); ++source_idx) { + for (auto const row_group_idx : row_group_indices[source_idx]) { + auto const& row_group = _extended_metadata->get_row_group(row_group_idx, source_idx); + auto const schema_idx = column_schemas[col_idx]; + auto const col = std::find_if( + row_group.columns.begin(), row_group.columns.end(), [schema_idx](auto const& candidate) { + return candidate.schema_idx == schema_idx; + }); + CUDF_EXPECTS(col != row_group.columns.end(), "Selected payload column is missing"); + auto const page_count = col->offset_index->page_locations.size(); + CUDF_EXPECTS(mask_idx + page_count <= data_page_mask.size(), + "Computed data page mask is incomplete"); + auto& mask = chunk_masks[row_group_ordinal * num_columns + col_idx]; + mask.reserve(page_count); + std::transform(data_page_mask.begin() + mask_idx, + data_page_mask.begin() + mask_idx + page_count, + std::back_inserter(mask), + [](bool retained) { return static_cast(retained); }); + mask_idx += page_count; + ++row_group_ordinal; + } + } + } + // compute_data_page_mask currently leaves unused trailing entries after the logical + // column-major page mask. Preserve the established consumer behavior by discarding them here. + data_page_mask.resize(mask_idx); + + struct exact_request { + int64_t offset; + int64_t size; + std::size_t mapping_idx; + }; + auto exact_requests = std::vector>(row_group_indices.size()); + auto page_mappings = std::vector{}; + auto resident_bytes = std::vector(num_chunks, 0); + auto dictionary_present = std::vector(num_chunks, 0); + + std::size_t row_group_ordinal = 0; + for (std::size_t source_idx = 0; source_idx < row_group_indices.size(); ++source_idx) { + for (auto const row_group_idx : row_group_indices[source_idx]) { + auto const& row_group = _extended_metadata->get_row_group(row_group_idx, source_idx); + for (std::size_t col_idx = 0; col_idx < num_columns; ++col_idx) { + auto const chunk_idx = row_group_ordinal * num_columns + col_idx; + auto const schema_idx = column_schemas[col_idx]; + auto const col = std::find_if( + row_group.columns.begin(), row_group.columns.end(), [schema_idx](auto const& candidate) { + return candidate.schema_idx == schema_idx; + }); + CUDF_EXPECTS(col != row_group.columns.end(), "Selected payload column is missing"); + auto const& page_locations = col->offset_index->page_locations; + auto const& retained = chunk_masks[chunk_idx]; + CUDF_EXPECTS(retained.size() == page_locations.size(), + "Data page mask does not match the offset index"); + auto const any_retained = + std::any_of(retained.begin(), retained.end(), [](auto value) { return value != 0; }); + + std::optional> dictionary_range; + if (col->meta_data.dictionary_page_offset > 0) { + auto const offset = col->meta_data.dictionary_page_offset; + auto const size = col->meta_data.data_page_offset - offset; + if (size > 0) { dictionary_range = std::pair{offset, size}; } + } else if (col->meta_data.data_page_offset < page_locations.front().offset) { + auto const offset = col->meta_data.data_page_offset; + dictionary_range = + std::pair{offset, page_locations.front().offset - col->meta_data.data_page_offset}; + } + + auto add_mapping = [&](bool fetched, int64_t offset, int64_t size) { + CUDF_EXPECTS( + offset >= 0 and size > 0 and offset <= std::numeric_limits::max() - size, + "Indexed page byte range is invalid"); + auto const mapping_idx = page_mappings.size(); + page_mappings.push_back( + page_range_mapping{.source_idx = static_cast(source_idx), + .range_idx = 0, + .range_offset = 0, + .size = fetched ? static_cast(size) : 0, + .fetched = fetched}); + if (fetched) { + exact_requests[source_idx].push_back(exact_request{offset, size, mapping_idx}); + resident_bytes[chunk_idx] += static_cast(size); + } + }; + + if (dictionary_range.has_value() and any_retained) { + add_mapping(true, dictionary_range->first, dictionary_range->second); + dictionary_present[chunk_idx] = 1; + } + for (std::size_t page_idx = 0; page_idx < page_locations.size(); ++page_idx) { + auto const& location = page_locations[page_idx]; + add_mapping(retained[page_idx] != 0, + location.offset, + static_cast(location.compressed_page_size)); + } + } + ++row_group_ordinal; + } + } + + auto source_ranges = std::vector>(row_group_indices.size()); + for (std::size_t source_idx = 0; source_idx < exact_requests.size(); ++source_idx) { + auto& requests = exact_requests[source_idx]; + std::stable_sort(requests.begin(), requests.end(), [](auto const& lhs, auto const& rhs) { + return std::tie(lhs.offset, lhs.size) < std::tie(rhs.offset, rhs.size); + }); + for (auto const& request : requests) { + auto& ranges = source_ranges[source_idx]; + if (ranges.empty() or request.offset > ranges.back().offset() + ranges.back().size()) { + ranges.emplace_back(request.offset, request.size); + } else { + auto const end = + std::max(ranges.back().offset() + ranges.back().size(), request.offset + request.size); + ranges.back() = byte_range_info{ranges.back().offset(), end - ranges.back().offset()}; + } + auto& mapping = page_mappings[request.mapping_idx]; + mapping.range_idx = ranges.size() - 1; + mapping.range_offset = static_cast(request.offset - ranges.back().offset()); + } + } + + _pending_payload_page_io_plan = + payload_page_io_plan{.sparse = true, + .mask_data_pages = mask_data_pages, + .row_group_indices = {row_group_indices.begin(), row_group_indices.end()}, + .column_schema_indices = std::move(column_schemas), + .source_ranges = source_ranges, + .page_mappings = std::move(page_mappings), + .resident_bytes_per_chunk = std::move(resident_bytes), + .dictionary_present_per_chunk = std::move(dictionary_present), + .data_page_mask = std::move(data_page_mask)}; + return source_ranges; +} + std::pair, std::vector> hybrid_scan_reader_impl::all_column_chunks_byte_ranges( std::span const> row_group_indices, parquet_reader_options const& options) @@ -712,6 +977,125 @@ void hybrid_scan_reader_impl::setup_chunking_for_payload_columns( prepare_data(read_mode::CHUNKED_READ, row_group_indices, column_chunk_data, data_page_mask); } +void hybrid_scan_reader_impl::setup_chunking_for_payload_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, + use_data_page_mask mask_data_pages, + std::span> const> page_data_per_source, + parquet_reader_options const& options, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + CUDF_EXPECTS(_pending_payload_page_io_plan.has_value(), + "No pending payload page I/O plan to consume"); + // Consume first so a failed setup cannot accidentally reuse stale pointer/range mappings. + auto plan = std::move(_pending_payload_page_io_plan.value()); + _pending_payload_page_io_plan.reset(); + + CUDF_EXPECTS(plan.mask_data_pages == mask_data_pages, + "Payload setup page-mask option does not match its pending plan"); + auto const setup_row_groups = + std::vector>{row_group_indices.begin(), row_group_indices.end()}; + CUDF_EXPECTS(plan.row_group_indices == setup_row_groups, + "Payload setup row groups do not match the pending page I/O plan"); + + reset_column_selection(); + select_columns(read_columns_mode::PAYLOAD_COLUMNS, options); + auto selected_schemas = std::vector{}; + selected_schemas.reserve(_input_columns.size()); + std::transform(_input_columns.begin(), + _input_columns.end(), + std::back_inserter(selected_schemas), + [](auto const& col) { return col.schema_idx; }); + CUDF_EXPECTS(selected_schemas == plan.column_schema_indices, + "Payload column selection does not match the pending page I/O plan"); + + CUDF_EXPECTS(page_data_per_source.size() == plan.source_ranges.size(), + "Fetched payload source count does not match the pending plan"); + for (std::size_t source_idx = 0; source_idx < page_data_per_source.size(); ++source_idx) { + CUDF_EXPECTS(page_data_per_source[source_idx].size() == plan.source_ranges[source_idx].size(), + "Fetched payload range count does not match the pending plan"); + for (std::size_t range_idx = 0; range_idx < page_data_per_source[source_idx].size(); + ++range_idx) { + auto const& data = page_data_per_source[source_idx][range_idx]; + auto const& range = plan.source_ranges[source_idx][range_idx]; + CUDF_EXPECTS(std::cmp_equal(data.size(), range.size()), + "Fetched payload span size does not match its planned byte range"); + CUDF_EXPECTS(data.size() == 0 or data.data() != nullptr, + "Fetched payload span has a null data pointer"); + } + } + + if (not plan.sparse) { + auto flat_chunk_data = std::vector>{}; + auto const span_count = + std::accumulate(page_data_per_source.begin(), + page_data_per_source.end(), + std::size_t{0}, + [](auto sum, auto const& spans) { return sum + spans.size(); }); + flat_chunk_data.reserve(span_count); + for (auto const& source_data : page_data_per_source) { + flat_chunk_data.insert(flat_chunk_data.end(), source_data.begin(), source_data.end()); + } + setup_chunking_for_payload_columns(chunk_read_limit, + pass_read_limit, + row_group_indices, + row_mask, + mask_data_pages, + flat_chunk_data, + options, + stream, + mr); + return; + } + + CUDF_EXPECTS(std::cmp_equal(row_mask.size(), total_rows_in_row_groups(row_group_indices)), + "Row mask must span across all input row groups"); + CUDF_EXPECTS(row_mask.null_count() == 0, + "Row mask must not have any nulls when materializing payload column"); + + prepare_materialization( + read_columns_mode::PAYLOAD_COLUMNS, row_group_indices.size(), options, stream, mr); + + _input_pass_read_limit = pass_read_limit; + _output_chunk_read_limit = chunk_read_limit; + + // Preserve the existing all-rows-pruned setup path. An all-false page plan has no byte ranges. + if (are_all_rows_pruned(row_mask, stream)) { + auto const empty_row_groups = + std::vector>(row_group_indices.size(), std::vector{}); + prepare_data(read_mode::CHUNKED_READ, empty_row_groups, {}, {}); + _file_itm_data.num_input_row_groups = count_row_groups(row_group_indices); + return; + } + + _sparse_page_spans.clear(); + _sparse_page_spans.reserve(plan.page_mappings.size()); + for (auto const& mapping : plan.page_mappings) { + if (not mapping.fetched) { + _sparse_page_spans.emplace_back(); + continue; + } + CUDF_EXPECTS(std::cmp_less(mapping.source_idx, page_data_per_source.size()), + "Sparse page mapping has an invalid source index"); + auto const& source_data = page_data_per_source[mapping.source_idx]; + CUDF_EXPECTS(mapping.range_idx < source_data.size(), + "Sparse page mapping has an invalid range index"); + auto const& range_data = source_data[mapping.range_idx]; + CUDF_EXPECTS(mapping.range_offset <= range_data.size() and + mapping.size <= range_data.size() - mapping.range_offset, + "Sparse page mapping exceeds its fetched range"); + _sparse_page_spans.emplace_back(range_data.data() + mapping.range_offset, mapping.size); + } + _sparse_resident_bytes_per_chunk = std::move(plan.resident_bytes_per_chunk); + _sparse_dictionary_present_per_chunk = std::move(plan.dictionary_present_per_chunk); + _sparse_page_io = true; + + prepare_data(read_mode::CHUNKED_READ, row_group_indices, {}, plan.data_page_mask); +} + table_with_metadata hybrid_scan_reader_impl::materialize_payload_columns_chunk( cudf::column_view const& row_mask) { @@ -882,6 +1266,10 @@ void hybrid_scan_reader_impl::reset_internal_state() _pass_page_mask.clear(); _subpass_page_mask.reset(); _output_metadata.reset(); + _sparse_page_spans.clear(); + _sparse_resident_bytes_per_chunk.clear(); + _sparse_dictionary_present_per_chunk.clear(); + _sparse_page_io = false; _options.timestamp_type = cudf::data_type{}; _options.decimal_width = type_id::EMPTY; @@ -1214,9 +1602,6 @@ void hybrid_scan_reader_impl::set_pass_page_mask(std::span data_page cuda::counting_iterator{_input_columns.size()}, [&](auto col_idx) { for (std::size_t chunk_idx = col_idx; chunk_idx < chunks.size(); chunk_idx += num_columns) { - // Insert a true value for each dictionary page - if (chunks[chunk_idx].num_dict_pages > 0) { _pass_page_mask.push_back(true); } - // Number of data pages in this column chunk auto const num_data_pages_this_col_chunk = chunks[chunk_idx].num_data_pages; @@ -1225,6 +1610,16 @@ void hybrid_scan_reader_impl::set_pass_page_mask(std::span data_page data_page_mask.size() >= num_inserted_data_pages + num_data_pages_this_col_chunk, "Encountered invalid data page mask size"); + // Sparse chunks omit dictionaries when every data page is pruned. The contiguous path + // retains its existing conservative dictionary behavior. + if (chunks[chunk_idx].num_dict_pages > 0) { + auto const chunk_has_retained_page = std::any_of( + data_page_mask.begin() + num_inserted_data_pages, + data_page_mask.begin() + num_inserted_data_pages + num_data_pages_this_col_chunk, + [](bool retained) { return retained; }); + _pass_page_mask.push_back(_sparse_page_io ? chunk_has_retained_page : true); + } + // Insert page mask for this column chunk _pass_page_mask.insert( _pass_page_mask.end(), diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_impl.hpp b/cpp/src/io/parquet/experimental/hybrid_scan_impl.hpp index 3dd241af35ff..5e21411d9fc9 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_impl.hpp +++ b/cpp/src/io/parquet/experimental/hybrid_scan_impl.hpp @@ -194,6 +194,13 @@ class hybrid_scan_reader_impl : public parquet::detail::reader_impl { payload_column_chunks_byte_ranges(std::span const> row_group_indices, parquet_reader_options const& options); + [[nodiscard]] std::vector> payload_column_chunks_byte_ranges( + std::span const> row_group_indices, + cudf::column_view const& row_mask, + use_data_page_mask mask_data_pages, + parquet_reader_options const& options, + rmm::cuda_stream_view stream); + /** * @copydoc cudf::io::parquet::experimental::hybrid_scan_multifile::materialize_payload_columns */ @@ -260,6 +267,17 @@ class hybrid_scan_reader_impl : public parquet::detail::reader_impl { rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr); + void setup_chunking_for_payload_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, + use_data_page_mask mask_data_pages, + std::span> const> page_data_per_source, + parquet_reader_options const& options, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr); + /** * @copydoc * cudf::io::parquet::experimental::hybrid_scan_multifile::materialize_payload_columns_chunk @@ -314,6 +332,26 @@ class hybrid_scan_reader_impl : public parquet::detail::reader_impl { */ enum class read_columns_mode { FILTER_COLUMNS, PAYLOAD_COLUMNS, ALL_COLUMNS }; + struct page_range_mapping { + cudf::size_type source_idx{}; + std::size_t range_idx{}; + std::size_t range_offset{}; + std::size_t size{}; + bool fetched{}; + }; + + struct payload_page_io_plan { + bool sparse{}; + use_data_page_mask mask_data_pages{}; + std::vector> row_group_indices; + std::vector column_schema_indices; + std::vector> source_ranges; + std::vector page_mappings; + std::vector resident_bytes_per_chunk; + std::vector dictionary_present_per_chunk; + thrust::host_vector data_page_mask; + }; + /** * @brief Populate the reader's `_options` config (and related members) from the user options. * @@ -460,7 +498,8 @@ class hybrid_scan_reader_impl : public parquet::detail::reader_impl { * * @param column_chunk_data Device spans of buffers containing column chunk data */ - void setup_next_pass(std::span const> column_chunk_data); + void setup_next_pass(std::span const> column_chunk_data, + host_span data_page_mask); /** * @brief Setup pointers to columns chunks to be processed for this pass. @@ -567,6 +606,12 @@ class hybrid_scan_reader_impl : public parquet::detail::reader_impl { bool _is_filter_columns_selected{false}; bool _is_payload_columns_selected{false}; bool _is_all_columns_selected{false}; + + std::optional _pending_payload_page_io_plan; + std::vector> _sparse_page_spans; + std::vector _sparse_resident_bytes_per_chunk; + std::vector _sparse_dictionary_present_per_chunk; + bool _sparse_page_io{false}; }; } // namespace cudf::io::parquet::experimental::detail diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_multifile.cpp b/cpp/src/io/parquet/experimental/hybrid_scan_multifile.cpp index 38c355a651d7..9f2d20cb70cf 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_multifile.cpp +++ b/cpp/src/io/parquet/experimental/hybrid_scan_multifile.cpp @@ -137,6 +137,19 @@ hybrid_scan_multifile::payload_column_chunks_byte_ranges( return _impl->payload_column_chunks_byte_ranges(row_group_indices, options); } +std::vector> +hybrid_scan_multifile::payload_column_chunks_byte_ranges( + cudf::host_span const> row_group_indices, + cudf::column_view const& row_mask, + use_data_page_mask mask_data_pages, + parquet_reader_options const& options, + rmm::cuda_stream_view stream) const +{ + CUDF_FUNC_RANGE(); + return _impl->payload_column_chunks_byte_ranges( + row_group_indices, row_mask, mask_data_pages, options, stream); +} + table_with_metadata hybrid_scan_multifile::materialize_payload_columns( cudf::host_span const> row_group_indices, cudf::host_span const> column_chunk_data, @@ -224,6 +237,29 @@ void hybrid_scan_multifile::setup_chunking_for_payload_columns( mr); } +void hybrid_scan_multifile::setup_chunking_for_payload_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, + use_data_page_mask mask_data_pages, + cudf::host_span> const> page_data_per_source, + parquet_reader_options const& options, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) const +{ + CUDF_FUNC_RANGE(); + _impl->setup_chunking_for_payload_columns(chunk_read_limit, + pass_read_limit, + row_group_indices, + row_mask, + mask_data_pages, + page_data_per_source, + options, + stream, + mr); +} + table_with_metadata hybrid_scan_multifile::materialize_payload_columns_chunk( cudf::column_view const& row_mask) const { diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu b/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu index 64037e0c87df..877ff78c6a94 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu +++ b/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu @@ -180,6 +180,43 @@ void hybrid_scan_reader_impl::setup_compressed_data( auto& chunks = pass.chunks; + if (_sparse_page_io) { + CUDF_EXPECTS(_has_page_index, "Sparse page I/O requires complete page indexes"); + CUDF_EXPECTS(_sparse_resident_bytes_per_chunk.size() == chunks.size(), + "Sparse resident-byte accounting does not match the logical chunks"); + CUDF_EXPECTS(_sparse_dictionary_present_per_chunk.size() == chunks.size(), + "Sparse dictionary mapping does not match the logical chunks"); + pass.has_compressed_data = false; + for (std::size_t chunk_idx = 0; chunk_idx < chunks.size(); ++chunk_idx) { + auto& chunk = chunks[chunk_idx]; + chunk.compressed_data = nullptr; + chunk.compressed_size = _sparse_resident_bytes_per_chunk[chunk_idx]; + pass.has_compressed_data |= + chunk.codec != Compression::UNCOMPRESSED and chunk.compressed_size > 0; + } + + auto const indexed_total_pages = count_page_headers_with_pgidx(chunks, _stream); + auto total_pages = std::size_t{0}; + for (std::size_t chunk_idx = 0; chunk_idx < chunks.size(); ++chunk_idx) { + chunks[chunk_idx].num_dict_pages = _sparse_dictionary_present_per_chunk[chunk_idx] ? 1 : 0; + total_pages += chunks[chunk_idx].num_data_pages + chunks[chunk_idx].num_dict_pages; + } + CUDF_EXPECTS(total_pages <= indexed_total_pages, + "Sparse dictionary mapping exceeds page-index metadata"); + chunks.host_to_device_async(_stream); + CUDF_EXPECTS(total_pages == _sparse_page_spans.size(), + "Sparse page span count does not match page-index metadata"); + if (total_pages <= 0) { return; } + // `decode_page_headers` may not write every byte of each PageInfo, and `sort_pages` copies + // PageInfo as whole objects. + auto unsorted_pages = cudf::detail::make_zeroed_device_uvector_async( + total_pages, _stream, cudf::get_current_device_resource_ref()); + parquet::detail::decode_page_headers(pass, unsorted_pages, _sparse_page_spans, _stream); + CUDF_EXPECTS(pass.page_offsets.size() - 1 == static_cast(_input_columns.size()), + "Encountered page_offsets / num_columns mismatch"); + return; + } + pass.has_compressed_data = setup_column_chunks(column_chunk_data); // Process dataset chunk pages into output columns diff --git a/cpp/tests/io/experimental/hybrid_scan_common.cpp b/cpp/tests/io/experimental/hybrid_scan_common.cpp index 339009ecf7e5..7f61477f6a8d 100644 --- a/cpp/tests/io/experimental/hybrid_scan_common.cpp +++ b/cpp/tests/io/experimental/hybrid_scan_common.cpp @@ -199,6 +199,15 @@ multisource_device_data fetch_multisource_device_data( { auto const byte_ranges_per_source = group_byte_ranges_by_source(byte_ranges_and_source_map, inputs.datasources.size()); + return fetch_multisource_device_data(inputs, byte_ranges_per_source, stream, mr); +} + +multisource_device_data fetch_multisource_device_data( + multifile_inputs const& inputs, + std::vector> const& byte_ranges_per_source, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ auto [buffers, per_source_spans, tasks] = cudf::io::parquet::fetch_byte_ranges_to_device_async( inputs.datasource_refs, cudf::host_span const>{byte_ranges_per_source}, diff --git a/cpp/tests/io/experimental/hybrid_scan_common.hpp b/cpp/tests/io/experimental/hybrid_scan_common.hpp index 8a8493977a4b..bba44361c8ae 100644 --- a/cpp/tests/io/experimental/hybrid_scan_common.hpp +++ b/cpp/tests/io/experimental/hybrid_scan_common.hpp @@ -98,6 +98,15 @@ void setup_page_indexes(cudf::io::parquet::experimental::hybrid_scan_multifile c rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr); +/** + * @brief Fetches per-source byte ranges and returns per-source and flattened spans + */ +[[nodiscard]] multisource_device_data fetch_multisource_device_data( + multifile_inputs const& inputs, + std::vector> const& byte_ranges_per_source, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr); + /** * @brief Concatenate a vector of tables and return the resultant table * diff --git a/cpp/tests/io/experimental/hybrid_scan_multifile_composer.cpp b/cpp/tests/io/experimental/hybrid_scan_multifile_composer.cpp index d3d52c2c5129..076b66c87bb6 100644 --- a/cpp/tests/io/experimental/hybrid_scan_multifile_composer.cpp +++ b/cpp/tests/io/experimental/hybrid_scan_multifile_composer.cpp @@ -158,6 +158,73 @@ chunked_hybrid_scan_multifile(cudf::io::source_info const& source_info, concatenate_tables(std::move(payload_tables), stream, mr)}; } +std::tuple, std::unique_ptr> +page_level_chunked_hybrid_scan_multifile( + cudf::io::source_info const& source_info, + cudf::ast::operation const& filter_expression, + std::optional> const& payload_column_names, + bool case_sensitive_names, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + auto options = cudf::io::parquet_reader_options::builder() + .filter(filter_expression) + .case_sensitive_names(case_sensitive_names) + .build(); + if (payload_column_names.has_value()) { options.set_column_names(payload_column_names.value()); } + + auto inputs = multifile_inputs(source_info); + auto reader = + cudf::io::parquet::experimental::hybrid_scan_multifile{inputs.footer_byte_spans, options}; + setup_page_indexes(reader, inputs); + + auto const input_row_groups = reader.all_row_groups(options); + auto const row_groups = reader.filter_row_groups_with_stats(input_row_groups, options, stream); + auto row_mask = reader.build_row_mask_with_page_index_stats(row_groups, options, stream, mr); + + auto constexpr chunk_read_limit = std::size_t{256 * 1024}; + auto constexpr pass_read_limit = std::size_t{1024 * 1024}; + + auto filter_tables = std::vector>{}; + auto payload_tables = std::vector>{}; + + auto filter_column_chunks = fetch_multisource_device_data( + inputs, reader.filter_column_chunks_byte_ranges(row_groups, options), stream, mr); + auto row_mask_view = row_mask->mutable_view(); + reader.setup_chunking_for_filter_columns(chunk_read_limit, + pass_read_limit, + row_groups, + row_mask_view, + use_data_page_mask::YES, + filter_column_chunks.flat_spans, + options, + stream, + mr); + while (reader.has_next_table_chunk()) { + filter_tables.push_back(reader.materialize_filter_columns_chunk(row_mask_view).tbl); + } + + auto const payload_page_ranges = reader.payload_column_chunks_byte_ranges( + row_groups, row_mask->view(), use_data_page_mask::YES, options, stream); + auto payload_page_data = fetch_multisource_device_data(inputs, payload_page_ranges, stream, mr); + + reader.setup_chunking_for_payload_columns(chunk_read_limit, + pass_read_limit, + row_groups, + row_mask->view(), + use_data_page_mask::YES, + payload_page_data.per_source_spans, + options, + stream, + mr); + while (reader.has_next_table_chunk()) { + payload_tables.push_back(reader.materialize_payload_columns_chunk(row_mask->view()).tbl); + } + + return std::tuple{concatenate_tables(std::move(filter_tables), stream, mr), + concatenate_tables(std::move(payload_tables), stream, mr)}; +} + std::unique_ptr chunked_hybrid_scan_multifile_single_step( cudf::io::source_info const& source_info, cudf::ast::operation const& filter_expression, diff --git a/cpp/tests/io/experimental/hybrid_scan_multifile_composer.hpp b/cpp/tests/io/experimental/hybrid_scan_multifile_composer.hpp index 50421f6f5ac2..fc87ed381575 100644 --- a/cpp/tests/io/experimental/hybrid_scan_multifile_composer.hpp +++ b/cpp/tests/io/experimental/hybrid_scan_multifile_composer.hpp @@ -79,6 +79,30 @@ chunked_hybrid_scan_multifile(cudf::io::source_info const& source_info, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr); +/** + * @brief Read parquet sources using chunked materialization and page-level payload I/O + * + * Filter columns continue to use the full-column-chunk path. Payload ranges are planned after + * filter materialization updates the row mask, then only requested pages are fetched. + * + * @param source_info Input source info containing one or more Parquet sources + * @param filter_expression Filter expression + * @param payload_column_names List of paths of select payload column names, if any + * @param case_sensitive_names Whether column names are case sensitive + * @param stream CUDA stream for hybrid scan reader + * @param mr Device memory resource + * + * @return Tuple of filter and payload tables + */ +std::tuple, std::unique_ptr> +page_level_chunked_hybrid_scan_multifile( + cudf::io::source_info const& source_info, + cudf::ast::operation const& filter_expression, + std::optional> const& payload_column_names, + bool case_sensitive_names, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr); + /** * @brief Read parquet sources with the hybrid scan multifile reader in a single step using chunked * materialization diff --git a/cpp/tests/io/experimental/hybrid_scan_multifile_test.cpp b/cpp/tests/io/experimental/hybrid_scan_multifile_test.cpp index 9a4d9daa2c97..84285e197350 100644 --- a/cpp/tests/io/experimental/hybrid_scan_multifile_test.cpp +++ b/cpp/tests/io/experimental/hybrid_scan_multifile_test.cpp @@ -18,6 +18,7 @@ #include #include #include +#include #include #include #include @@ -35,20 +36,126 @@ namespace { +using cudf::io::parquet::experimental::use_data_page_mask; + +std::pair payload_byte_range_sizes( + cudf::io::source_info const& source_info, + cudf::ast::operation const& filter_expression, + bool case_sensitive_names, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + auto options = cudf::io::parquet_reader_options::builder() + .filter(filter_expression) + .case_sensitive_names(case_sensitive_names) + .build(); + auto inputs = multifile_inputs(source_info); + auto reader = + cudf::io::parquet::experimental::hybrid_scan_multifile{inputs.footer_byte_spans, options}; + setup_page_indexes(reader, inputs); + + auto const input_row_groups = reader.all_row_groups(options); + auto const row_groups = reader.filter_row_groups_with_stats(input_row_groups, options, stream); + auto row_mask = reader.build_row_mask_with_page_index_stats(row_groups, options, stream, mr); + + auto filter_data = fetch_multisource_device_data( + inputs, reader.filter_column_chunks_byte_ranges(row_groups, options), stream, mr); + auto row_mask_view = row_mask->mutable_view(); + reader.setup_chunking_for_filter_columns(256 * 1024, + 1024 * 1024, + row_groups, + row_mask_view, + use_data_page_mask::YES, + filter_data.flat_spans, + options, + stream, + mr); + while (reader.has_next_table_chunk()) { + static_cast(reader.materialize_filter_columns_chunk(row_mask_view)); + } + + auto const full_ranges = reader.payload_column_chunks_byte_ranges(row_groups, options).first; + auto const page_ranges = reader.payload_column_chunks_byte_ranges( + row_groups, row_mask->view(), use_data_page_mask::YES, options, stream); + auto const full_bytes = std::accumulate( + full_ranges.begin(), full_ranges.end(), std::size_t{0}, [](auto sum, auto range) { + return sum + range.size(); + }); + auto const requested_bytes = std::accumulate( + page_ranges.begin(), + page_ranges.end(), + std::size_t{0}, + [](auto source_sum, auto const& source_ranges) { + return source_sum + std::accumulate(source_ranges.begin(), + source_ranges.end(), + std::size_t{0}, + [](auto sum, auto range) { return sum + range.size(); }); + }); + return {requested_bytes, full_bytes}; +} + +std::vector> make_plain_payload_parquet_buffers() +{ + auto constexpr num_sources = 2; + auto parquet_buffers = std::vector>(num_sources); + for (auto source_idx = 0; source_idx < num_sources; ++source_idx) { + auto filter_values = cuda::counting_iterator{0}; + auto payload_values = + cudf::detail::make_counting_transform_iterator(cudf::size_type{0}, [source_idx](auto i) { + return static_cast(i) + source_idx * int64_t{num_ordered_rows}; + }); + auto filter = cudf::test::fixed_width_column_wrapper( + filter_values, filter_values + num_ordered_rows); + auto payload = cudf::test::fixed_width_column_wrapper( + payload_values, payload_values + num_ordered_rows); + auto const table = cudf::table_view{{filter, payload}}; + + cudf::io::table_input_metadata metadata(table); + metadata.column_metadata[0].set_name("col0"); + metadata.column_metadata[1].set_name("col1").set_encoding(cudf::io::column_encoding::PLAIN); + auto options = cudf::io::parquet_writer_options::builder( + cudf::io::sink_info{&parquet_buffers[source_idx]}, table) + .metadata(metadata) + .row_group_size_rows(num_ordered_rows) + .max_page_size_rows(page_size_for_ordered_tests) + .max_page_size_bytes(64 * 1024 * 1024) + .compression(cudf::io::compression_type::NONE) + .dictionary_policy(cudf::io::dictionary_policy::NEVER) + .stats_level(cudf::io::statistics_freq::STATISTICS_COLUMN); + cudf::io::write_parquet(options); + } + return parquet_buffers; +} + +void expect_byte_ranges_equal( + std::vector> const& expected, + std::vector> const& actual) +{ + ASSERT_EQ(expected.size(), actual.size()); + for (std::size_t source_idx = 0; source_idx < expected.size(); ++source_idx) { + ASSERT_EQ(expected[source_idx].size(), actual[source_idx].size()); + for (std::size_t range_idx = 0; range_idx < expected[source_idx].size(); ++range_idx) { + EXPECT_EQ(expected[source_idx][range_idx].offset(), actual[source_idx][range_idx].offset()); + EXPECT_EQ(expected[source_idx][range_idx].size(), actual[source_idx][range_idx].size()); + } + } +} + /** * @brief Helper to test multifile hybrid scan single-shot materialization * * Writes the input table to multiple parquet sources and compares filter, payload, and all-column - * materialization output with the regular multi-source parquet reader. The filter expression used - * is `col0 >= 100`. + * materialization output with the regular multi-source parquet reader. The filter expression is + * `col0 >= literal_value`. * * @note The first column in the input table must be constructed with * `cudf::test::ascending()` */ template void test_hybrid_scan_multifile(std::vector const& columns, - bool case_sensitive_names = true, - uint32_t literal_value = 100) + bool case_sensitive_names = true, + uint32_t literal_value = 100, + bool expect_payload_byte_reduction = false) { auto const table = cudf::table_view{columns}; cudf::io::table_input_metadata expected_metadata(table); @@ -92,11 +199,17 @@ void test_hybrid_scan_multifile(std::vector const& columns, auto const [chunked_filter_table, chunked_payload_table] = chunked_hybrid_scan_multifile( source_info, filter_expression, {}, case_sensitive_names, stream, mr); + auto const [page_level_filter_table, page_level_payload_table] = + page_level_chunked_hybrid_scan_multifile( + source_info, filter_expression, {}, case_sensitive_names, stream, mr); + auto const chunked_all_table = chunked_hybrid_scan_multifile_single_step( source_info, filter_expression, {}, case_sensitive_names, stream, mr); CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected.tbl->select({0}), filter_table->view()); CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected.tbl->select({0}), chunked_filter_table->view()); + CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected.tbl->select({0}), page_level_filter_table->view()); + CUDF_TEST_EXPECT_TABLES_EQUIVALENT(chunked_filter_table->view(), page_level_filter_table->view()); auto payload_column_indices = std::vector(columns.size() - 1); std::iota(payload_column_indices.begin(), payload_column_indices.end(), 1); @@ -104,8 +217,18 @@ void test_hybrid_scan_multifile(std::vector const& columns, payload_table->view()); CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected.tbl->select(payload_column_indices), chunked_payload_table->view()); + CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected.tbl->select(payload_column_indices), + page_level_payload_table->view()); + CUDF_TEST_EXPECT_TABLES_EQUIVALENT(chunked_payload_table->view(), + page_level_payload_table->view()); CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected.tbl->view(), all_table->view()); CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected.tbl->view(), chunked_all_table->view()); + if (expect_payload_byte_reduction) { + auto const [requested_payload_bytes, full_payload_bytes] = + payload_byte_range_sizes(source_info, filter_expression, case_sensitive_names, stream, mr); + EXPECT_GT(requested_payload_bytes, 0); + EXPECT_LT(requested_payload_bytes, full_payload_bytes); + } } } // namespace @@ -122,7 +245,7 @@ TEST_F(HybridScanMultifileTest, EmptyResult) auto col3 = make_list_str_column(gen, true, false); auto col4 = make_list_str_column(gen, true, true); - auto constexpr literal_value = uint32_t(num_ordered_rows); + auto constexpr literal_value = static_cast(num_ordered_rows); test_hybrid_scan_multifile({col0, *col1, *col2, *col3, *col4}, false, literal_value); } @@ -159,6 +282,398 @@ TEST_F(HybridScanMultifileTest, MaterializeListsOfStrings) test_hybrid_scan_multifile({col0, *col1, *col2, *col3, *col4}, false); } +TEST_F(HybridScanMultifileTest, PageLevelDictionaryPayloadByteReduction) +{ + auto col0 = testdata::ascending(); + + auto payload_values = std::vector(num_ordered_rows); + for (auto i = std::size_t{0}; i < payload_values.size(); ++i) { + payload_values[i] = "dictionary value " + std::to_string(i % 8); + } + auto col1 = cudf::test::strings_column_wrapper(payload_values.begin(), payload_values.end()); + + // A page-aligned threshold retains two of four data pages. The writer's ALWAYS dictionary policy + // requires the page-I/O path to retain the dictionary while requesting fewer bytes than the + // legacy full-column-chunk path. + auto constexpr threshold = uint32_t{2 * page_size_for_ordered_tests / 100}; + test_hybrid_scan_multifile({col0, col1}, true, threshold, true); +} + +TEST_F(HybridScanMultifileTest, PageLevelStringsSeparatedByPrunedPages) +{ + auto filter_values = cudf::detail::make_counting_transform_iterator( + cudf::size_type{0}, [](auto i) { return (i / page_size_for_ordered_tests) % 2 == 0; }); + auto filter = + cudf::test::fixed_width_column_wrapper(filter_values, filter_values + num_ordered_rows); + + auto payload_values = std::vector(num_ordered_rows); + for (auto i = std::size_t{0}; i < payload_values.size(); ++i) { + payload_values[i] = "payload value " + std::to_string(i); + } + auto payload = cudf::test::strings_column_wrapper(payload_values.begin(), payload_values.end()); + auto table = cudf::table_view{{filter, payload}}; + + auto metadata = cudf::io::table_input_metadata(table); + metadata.column_metadata[0].set_name("filter"); + metadata.column_metadata[1].set_name("payload"); + + auto parquet_buffers = std::vector>(2); + for (auto& parquet_buffer : parquet_buffers) { + auto options = + cudf::io::parquet_writer_options::builder(cudf::io::sink_info{&parquet_buffer}, table) + .metadata(metadata) + .row_group_size_rows(num_ordered_rows) + .max_page_size_rows(page_size_for_ordered_tests) + .max_page_size_bytes(64 * 1024 * 1024) + .compression(cudf::io::compression_type::NONE) + .dictionary_policy(cudf::io::dictionary_policy::ALWAYS) + .stats_level(cudf::io::statistics_freq::STATISTICS_COLUMN); + cudf::io::write_parquet(options); + } + + auto const filter_ref = cudf::ast::column_name_reference("filter"); + auto filter_expression = cudf::ast::operation(cudf::ast::ast_operator::IDENTITY, filter_ref); + auto source_info = build_source_info(parquet_buffers); + auto const stream = cudf::get_default_stream(); + auto const mr = cudf::get_current_device_resource_ref(); + auto expected_options = + cudf::io::parquet_reader_options::builder(source_info).filter(filter_expression).build(); + auto expected = cudf::io::read_parquet(expected_options, stream, mr); + + auto const [filter_result, payload_result] = + page_level_chunked_hybrid_scan_multifile(source_info, filter_expression, {}, true, stream, mr); + + CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected.tbl->select({0}), filter_result->view()); + CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected.tbl->select({1}), payload_result->view()); +} + +TEST_F(HybridScanMultifileTest, PageLevelAsymmetricSourceRowGroupOrdering) +{ + auto constexpr rows_per_group = 2 * page_size_for_ordered_tests; + auto constexpr rows_source_0 = num_ordered_rows; + auto constexpr rows_source_1 = num_ordered_rows; + auto constexpr rows_per_page = rows_per_group / 4; + + auto source_0_filter_values = cudf::detail::make_counting_transform_iterator( + 0, [](auto i) { return static_cast(i / 100); }); + auto source_1_filter_values = cudf::detail::make_counting_transform_iterator( + 0, [](auto i) { return static_cast((num_ordered_rows - i) / 100); }); + auto source_0_filter = cudf::test::fixed_width_column_wrapper( + source_0_filter_values, source_0_filter_values + rows_source_0); + auto source_1_filter = cudf::test::fixed_width_column_wrapper( + source_1_filter_values, source_1_filter_values + rows_source_1); + + auto source_0_payload_values = std::vector(rows_source_0); + auto source_1_payload_values = std::vector(rows_source_1); + for (auto i = std::size_t{0}; i < source_0_payload_values.size(); ++i) { + source_0_payload_values[i] = "source 0 dictionary value " + std::to_string(i % 8); + } + for (auto i = std::size_t{0}; i < source_1_payload_values.size(); ++i) { + source_1_payload_values[i] = "source 1 dictionary value " + std::to_string(i % 8); + } + auto source_0_payload = cudf::test::strings_column_wrapper(source_0_payload_values.begin(), + source_0_payload_values.end()); + auto source_1_payload = cudf::test::strings_column_wrapper(source_1_payload_values.begin(), + source_1_payload_values.end()); + auto const source_0_table = cudf::table_view{{source_0_filter, source_0_payload}}; + auto const source_1_table = cudf::table_view{{source_1_filter, source_1_payload}}; + + auto parquet_buffers = std::vector>(2); + auto const write_source = [&](auto const& table, auto& buffer) { + cudf::io::table_input_metadata metadata(table); + metadata.column_metadata[0].set_name("col0"); + auto options = cudf::io::parquet_writer_options::builder(cudf::io::sink_info{&buffer}, table) + .metadata(metadata) + .row_group_size_rows(rows_per_group) + .max_page_size_rows(rows_per_page) + .dictionary_policy(cudf::io::dictionary_policy::ALWAYS) + .stats_level(cudf::io::statistics_freq::STATISTICS_COLUMN); + cudf::io::write_parquet(options); + }; + write_source(source_0_table, parquet_buffers[0]); + write_source(source_1_table, parquet_buffers[1]); + + auto constexpr threshold = uint32_t{75}; + auto scalar = cudf::numeric_scalar(threshold); + auto literal = cudf::ast::literal(scalar); + auto col_ref = cudf::ast::column_name_reference("col0"); + auto filter_expression = + cudf::ast::operation(cudf::ast::ast_operator::GREATER_EQUAL, col_ref, literal); + + auto const stream = cudf::get_default_stream(); + auto const mr = cudf::get_current_device_resource_ref(); + auto const source_info = build_source_info(parquet_buffers); + auto const expected = cudf::io::read_parquet( + cudf::io::parquet_reader_options::builder(source_info).filter(filter_expression), stream, mr); + + auto const [filter_table, payload_table] = + page_level_chunked_hybrid_scan_multifile(source_info, filter_expression, {}, true, stream, mr); + + CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected.tbl->select({0}), filter_table->view()); + CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected.tbl->select({1}), payload_table->view()); + auto const [requested_payload_bytes, full_payload_bytes] = + payload_byte_range_sizes(source_info, filter_expression, true, stream, mr); + EXPECT_GT(requested_payload_bytes, 0); + EXPECT_LT(requested_payload_bytes, full_payload_bytes); +} + +TEST_F(HybridScanMultifileTest, PageLevelPlainEncodingExactCoalescedRanges) +{ + auto parquet_buffers = make_plain_payload_parquet_buffers(); + auto const source_info = build_source_info(parquet_buffers); + auto inputs = multifile_inputs(source_info); + + auto scalar = cudf::numeric_scalar(0); + auto literal = cudf::ast::literal(scalar); + auto col_ref_0 = cudf::ast::column_name_reference("col0"); + auto filter_expression = + cudf::ast::operation(cudf::ast::ast_operator::GREATER_EQUAL, col_ref_0, literal); + auto options = cudf::io::parquet_reader_options::builder() + .column_names({"col1"}) + .filter(filter_expression) + .build(); + auto reader = + cudf::io::parquet::experimental::hybrid_scan_multifile{inputs.footer_byte_spans, options}; + setup_page_indexes(reader, inputs); + + auto const row_groups = reader.all_row_groups(options); + auto const metadatas = reader.parquet_metadatas(); + auto selected_rows = + std::vector(reader.total_rows_in_row_groups(row_groups), uint8_t{0}); + auto expected_ranges = + std::vector>(metadatas.size()); + + std::size_t source_row_offset = 0; + for (std::size_t source_idx = 0; source_idx < metadatas.size(); ++source_idx) { + auto const& metadata = metadatas[source_idx]; + ASSERT_EQ(metadata.row_groups.size(), 1); + auto const& payload_chunk = metadata.row_groups.front().columns[1]; + EXPECT_NE(std::find(payload_chunk.meta_data.encodings.begin(), + payload_chunk.meta_data.encodings.end(), + cudf::io::parquet::Encoding::PLAIN), + payload_chunk.meta_data.encodings.end()); + EXPECT_EQ(std::find(payload_chunk.meta_data.encodings.begin(), + payload_chunk.meta_data.encodings.end(), + cudf::io::parquet::Encoding::RLE_DICTIONARY), + payload_chunk.meta_data.encodings.end()); + EXPECT_EQ(payload_chunk.meta_data.dictionary_page_offset, 0); + ASSERT_TRUE(payload_chunk.offset_index.has_value()); + + auto const& pages = payload_chunk.offset_index->page_locations; + ASSERT_GE(pages.size(), 4); + ASSERT_EQ(pages[1].offset + pages[1].compressed_page_size, pages[2].offset); + auto const selected_begin = pages[1].first_row_index; + auto const selected_end = pages[3].first_row_index; + ASSERT_GE(selected_begin, 0); + ASSERT_LE(selected_end, metadata.row_groups.front().num_rows); + std::fill(selected_rows.begin() + source_row_offset + selected_begin, + selected_rows.begin() + source_row_offset + selected_end, + uint8_t{1}); + + expected_ranges[source_idx].emplace_back( + pages[1].offset, pages[2].offset + pages[2].compressed_page_size - pages[1].offset); + source_row_offset += metadata.row_groups.front().num_rows; + } + ASSERT_EQ(source_row_offset, selected_rows.size()); + auto row_mask = + cudf::test::fixed_width_column_wrapper(selected_rows.begin(), selected_rows.end()) + .release(); + + auto const stream = cudf::get_default_stream(); + auto const mr = cudf::get_current_device_resource_ref(); + auto const page_ranges = reader.payload_column_chunks_byte_ranges( + row_groups, row_mask->view(), use_data_page_mask::YES, options, stream); + expect_byte_ranges_equal(expected_ranges, page_ranges); + + auto page_data = fetch_multisource_device_data(inputs, page_ranges, stream, mr); + reader.setup_chunking_for_payload_columns(256 * 1024, + 1024 * 1024, + row_groups, + row_mask->view(), + use_data_page_mask::YES, + page_data.per_source_spans, + options, + stream, + mr); + auto payload_chunks = std::vector>{}; + while (reader.has_next_table_chunk()) { + payload_chunks.push_back( + std::move(reader.materialize_payload_columns_chunk(row_mask->view()).tbl)); + } + auto actual = concatenate_tables(std::move(payload_chunks), stream, mr); + + auto const full = + cudf::io::read_parquet(cudf::io::parquet_reader_options::builder(source_info), stream, mr); + auto const expected = cudf::apply_boolean_mask(full.tbl->select({1}), row_mask->view()); + CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected->view(), actual->view()); +} + +TEST_F(HybridScanMultifileTest, PageLevelAllFalseMaskHasNoRanges) +{ + auto parquet_buffers = make_plain_payload_parquet_buffers(); + auto const source_info = build_source_info(parquet_buffers); + auto inputs = multifile_inputs(source_info); + + auto scalar = cudf::numeric_scalar(0); + auto literal = cudf::ast::literal(scalar); + auto col_ref_0 = cudf::ast::column_name_reference("col0"); + auto filter_expression = + cudf::ast::operation(cudf::ast::ast_operator::GREATER_EQUAL, col_ref_0, literal); + auto options = cudf::io::parquet_reader_options::builder() + .column_names({"col1"}) + .filter(filter_expression) + .build(); + auto reader = + cudf::io::parquet::experimental::hybrid_scan_multifile{inputs.footer_byte_spans, options}; + setup_page_indexes(reader, inputs); + + auto const row_groups = reader.all_row_groups(options); + auto false_values = cuda::make_constant_iterator(false); + auto row_mask = cudf::test::fixed_width_column_wrapper( + false_values, false_values + reader.total_rows_in_row_groups(row_groups)) + .release(); + auto const stream = cudf::get_default_stream(); + auto const mr = cudf::get_current_device_resource_ref(); + + auto const page_ranges = reader.payload_column_chunks_byte_ranges( + row_groups, row_mask->view(), use_data_page_mask::YES, options, stream); + ASSERT_EQ(page_ranges.size(), parquet_buffers.size()); + EXPECT_TRUE(std::all_of( + page_ranges.begin(), page_ranges.end(), [](auto const& ranges) { return ranges.empty(); })); + + auto const empty_page_data = + std::vector>>(parquet_buffers.size()); + reader.setup_chunking_for_payload_columns(0, + 0, + row_groups, + row_mask->view(), + use_data_page_mask::YES, + empty_page_data, + options, + stream, + mr); + ASSERT_TRUE(reader.has_next_table_chunk()); + auto const result = reader.materialize_payload_columns_chunk(row_mask->view()); + EXPECT_EQ(result.tbl->num_rows(), 0); + EXPECT_EQ(result.tbl->num_columns(), 1); + EXPECT_EQ(result.metadata.num_input_row_groups, 2); + EXPECT_FALSE(reader.has_next_table_chunk()); +} + +TEST_F(HybridScanMultifileTest, PageLevelNoMaskFallbackAndPlanLifecycle) +{ + auto parquet_buffers = make_plain_payload_parquet_buffers(); + auto const source_info = build_source_info(parquet_buffers); + auto inputs = multifile_inputs(source_info); + + auto scalar = cudf::numeric_scalar(0); + auto literal = cudf::ast::literal(scalar); + auto col_ref_0 = cudf::ast::column_name_reference("col0"); + auto filter_expression = + cudf::ast::operation(cudf::ast::ast_operator::GREATER_EQUAL, col_ref_0, literal); + auto options = cudf::io::parquet_reader_options::builder() + .column_names({"col1"}) + .filter(filter_expression) + .build(); + auto reader = + cudf::io::parquet::experimental::hybrid_scan_multifile{inputs.footer_byte_spans, options}; + setup_page_indexes(reader, inputs); + + auto const row_groups = reader.all_row_groups(options); + auto const full_ranges = group_byte_ranges_by_source( + reader.payload_column_chunks_byte_ranges(row_groups, options), parquet_buffers.size()); + auto true_values = cuda::make_constant_iterator(true); + auto row_mask = cudf::test::fixed_width_column_wrapper( + true_values, true_values + reader.total_rows_in_row_groups(row_groups)) + .release(); + auto const stream = cudf::get_default_stream(); + auto const mr = cudf::get_current_device_resource_ref(); + + auto const planned_ranges = reader.payload_column_chunks_byte_ranges( + row_groups, row_mask->view(), use_data_page_mask::NO, options, stream); + expect_byte_ranges_equal(full_ranges, planned_ranges); + EXPECT_THROW(static_cast(reader.payload_column_chunks_byte_ranges( + row_groups, row_mask->view(), use_data_page_mask::NO, options, stream)), + cudf::logic_error); + + auto page_data = fetch_multisource_device_data(inputs, planned_ranges, stream, mr); + reader.setup_chunking_for_payload_columns(0, + 0, + row_groups, + row_mask->view(), + use_data_page_mask::NO, + page_data.per_source_spans, + options, + stream, + mr); + EXPECT_THROW(reader.setup_chunking_for_payload_columns(0, + 0, + row_groups, + row_mask->view(), + use_data_page_mask::NO, + page_data.per_source_spans, + options, + stream, + mr), + cudf::logic_error); +} + +TEST_F(HybridScanMultifileTest, PageLevelRejectsInvalidFetchedSpans) +{ + auto parquet_buffers = make_plain_payload_parquet_buffers(); + auto const source_info = build_source_info(parquet_buffers); + auto inputs = multifile_inputs(source_info); + + auto scalar = cudf::numeric_scalar(0); + auto literal = cudf::ast::literal(scalar); + auto col_ref_0 = cudf::ast::column_name_reference("col0"); + auto filter_expression = + cudf::ast::operation(cudf::ast::ast_operator::GREATER_EQUAL, col_ref_0, literal); + auto options = cudf::io::parquet_reader_options::builder() + .column_names({"col1"}) + .filter(filter_expression) + .build(); + auto reader = + cudf::io::parquet::experimental::hybrid_scan_multifile{inputs.footer_byte_spans, options}; + setup_page_indexes(reader, inputs); + + auto const row_groups = reader.all_row_groups(options); + auto selected_values = cudf::detail::make_counting_transform_iterator( + cudf::size_type{0}, [](auto i) { return (i % num_ordered_rows) < num_ordered_rows / 2; }); + auto row_mask = cudf::test::fixed_width_column_wrapper( + selected_values, selected_values + reader.total_rows_in_row_groups(row_groups)) + .release(); + auto const stream = cudf::get_default_stream(); + auto const mr = cudf::get_current_device_resource_ref(); + + auto page_ranges = reader.payload_column_chunks_byte_ranges( + row_groups, row_mask->view(), use_data_page_mask::YES, options, stream); + auto page_data = fetch_multisource_device_data(inputs, page_ranges, stream, mr); + auto bad_count = page_data.per_source_spans; + auto count_source = std::find_if( + bad_count.begin(), bad_count.end(), [](auto const& spans) { return not spans.empty(); }); + ASSERT_NE(count_source, bad_count.end()); + count_source->pop_back(); + EXPECT_THROW( + reader.setup_chunking_for_payload_columns( + 0, 0, row_groups, row_mask->view(), use_data_page_mask::YES, bad_count, options, stream, mr), + cudf::logic_error); + + page_ranges = reader.payload_column_chunks_byte_ranges( + row_groups, row_mask->view(), use_data_page_mask::YES, options, stream); + auto bad_size = page_data.per_source_spans; + auto size_source = std::find_if( + bad_size.begin(), bad_size.end(), [](auto const& spans) { return not spans.empty(); }); + ASSERT_NE(size_source, bad_size.end()); + ASSERT_GT(size_source->front().size(), 1); + size_source->front() = + cudf::device_span{size_source->front().data(), size_source->front().size() - 1}; + EXPECT_THROW( + reader.setup_chunking_for_payload_columns( + 0, 0, row_groups, row_mask->view(), use_data_page_mask::YES, bad_size, options, stream, mr), + cudf::logic_error); +} + TEST_F(HybridScanMultifileTest, PrependIndexColumns) { using T = int32_t; From af6c46664b67e8b9cce2946157df50314127792b Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Fri, 31 Jul 2026 23:26:14 +0000 Subject: [PATCH 03/10] Simplify and cleanup --- .../io/experimental/hybrid_scan_multifile.hpp | 40 +- .../experimental/hybrid_scan_chunking.cu | 38 +- .../experimental/hybrid_scan_helpers.hpp | 26 +- .../parquet/experimental/hybrid_scan_impl.cpp | 463 +++++---------- .../parquet/experimental/hybrid_scan_impl.hpp | 62 +- .../experimental/hybrid_scan_multifile.cpp | 23 +- .../experimental/hybrid_scan_preprocess.cu | 71 ++- .../parquet/experimental/page_index_filter.cu | 3 +- .../io/parquet/io_utils/parquet_io_utils.cpp | 3 +- cpp/src/io/parquet/page_hdr.cu | 100 ---- cpp/src/io/parquet/parquet_gpu.hpp | 13 - cpp/src/io/parquet/reader_impl.cpp | 3 - cpp/src/io/parquet/reader_impl.hpp | 3 + cpp/src/io/parquet/reader_impl_preprocess.cu | 17 +- .../parquet/reader_impl_preprocess_utils.cu | 6 +- .../parquet/reader_impl_preprocess_utils.cuh | 18 +- .../io/experimental/hybrid_scan_common.cpp | 9 - .../io/experimental/hybrid_scan_common.hpp | 9 - .../hybrid_scan_multifile_composer.cpp | 10 +- .../hybrid_scan_multifile_composer.hpp | 2 +- .../hybrid_scan_multifile_filters_test.cpp | 22 + .../hybrid_scan_multifile_test.cpp | 546 ++++-------------- 22 files changed, 416 insertions(+), 1071 deletions(-) diff --git a/cpp/include/cudf/io/experimental/hybrid_scan_multifile.hpp b/cpp/include/cudf/io/experimental/hybrid_scan_multifile.hpp index 691e7ee6f516..ad814b732466 100644 --- a/cpp/include/cudf/io/experimental/hybrid_scan_multifile.hpp +++ b/cpp/include/cudf/io/experimental/hybrid_scan_multifile.hpp @@ -261,25 +261,26 @@ class hybrid_scan_multifile { parquet_reader_options const& options) const; /** - * @brief Plan page-level payload byte ranges grouped by source + * @brief Get byte ranges of pages of payload columns * - * When page masking cannot be used, returns the legacy full-column-chunk ranges regrouped by - * source. The resulting plan must be consumed exactly once by the matching page-data setup - * overload. + * Byte ranges are flattened in source, row group, column chunk, and page order. The returned + * source map has one source index per byte range. Dictionary pages precede data pages within each + * column chunk. Pruned pages are represented by empty byte ranges. + * + * @throws std::invalid_argument if any selected column chunk does not have a valid offset + * index * * @param row_group_indices Input row group indices, one vector per source * @param row_mask Boolean mask spanning the selected row groups - * @param mask_data_pages Whether to use the row mask to prune data pages * @param options Parquet reader options * @param stream CUDA stream used to compute the page mask - * @return Byte ranges to fetch, grouped by source + * @return Pair of flattened payload page byte ranges and their corresponding source indices */ - [[nodiscard]] std::vector> payload_column_chunks_byte_ranges( - cudf::host_span const> row_group_indices, - cudf::column_view const& row_mask, - use_data_page_mask mask_data_pages, - parquet_reader_options const& options, - rmm::cuda_stream_view stream) const; + [[nodiscard]] std::pair, std::vector> + payload_pages_byte_ranges(cudf::host_span const> row_group_indices, + cudf::column_view const& row_mask, + parquet_reader_options const& options, + rmm::cuda_stream_view stream) const; /** * @brief Materialize payload columns and applies the row mask to the output table @@ -405,17 +406,16 @@ class hybrid_scan_multifile { rmm::device_async_resource_ref mr) const; /** - * @brief Setup payload chunking from source-grouped page-level fetch results + * @brief Setup chunking information for payload columns and preprocess the input data pages * - * Consumes the pending plan created by the page-level payload byte-range overload. Each input - * span must correspond to the byte range at the same source and range index. + * Input page data spans (including empty ones) must have the same shape as the byte ranges + * returned by `payload_pages_byte_ranges`. The data page mask is inferred from the data spans. * * @param chunk_read_limit Maximum bytes returned per output table chunk, or zero * @param pass_read_limit Maximum read/decompression memory, or zero * @param row_group_indices Input row group indices, one vector per source - * @param row_mask Boolean mask spanning the selected row groups - * @param mask_data_pages Whether page masking was requested - * @param page_data_per_source Fetched device spans grouped by source + * @param page_data Flattened device spans of payload page data in the same order as the byte + * ranges from `payload_pages_byte_ranges` * @param options Parquet reader options * @param stream CUDA stream used for preprocessing * @param mr Device memory resource used for output table chunks @@ -424,9 +424,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, - use_data_page_mask mask_data_pages, - cudf::host_span> const> page_data_per_source, + cudf::host_span const> page_data, parquet_reader_options const& options, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) const; diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu b/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu index 4d5174718b10..7c439c984e09 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu +++ b/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu @@ -76,7 +76,7 @@ void hybrid_scan_reader_impl::handle_chunking( void hybrid_scan_reader_impl::setup_next_pass( std::span const> column_chunk_data, - host_span data_page_mask) + std::span data_page_mask) { auto const num_passes = _file_itm_data.num_passes(); CUDF_EXPECTS(num_passes == 1, @@ -118,10 +118,16 @@ void hybrid_scan_reader_impl::setup_next_pass( pass.num_rows = _file_itm_data.global_num_rows; // Setup page information for the chunk (which we can access without decompressing) - setup_compressed_data(column_chunk_data); - - // Establish the logical mask before malformed-page checks, size estimation, or subpass setup. - set_pass_page_mask(data_page_mask); + if (_sparse_page_io) { + CUDF_EXPECTS(data_page_mask.empty(), + "Encountered a non-empty input data page mask in sparse I/O path.", + std::invalid_argument); + setup_sparse_compressed_data(column_chunk_data); + set_sparse_pass_page_mask(column_chunk_data); + } else { + setup_compressed_data(column_chunk_data); + set_pass_page_mask(data_page_mask); + } // detect malformed columns. // - we have seen some cases in the wild where we have a row group containing N @@ -149,12 +155,22 @@ void hybrid_scan_reader_impl::setup_next_pass( // store off how much memory we've used so far. This includes the compressed page data and the // decompressed dictionary data. we will subtract this from the available total memory for the // subpasses - auto chunk_iter = thrust::make_transform_iterator(pass.chunks.d_begin(), - parquet::detail::get_chunk_compressed_size{}); - pass.base_mem_size = - decomp_dict_data_size + - cudf::detail::reduce( - chunk_iter, chunk_iter + pass.chunks.size(), size_t{0}, cuda::std::plus{}, _stream); + auto const compressed_data_size = + _sparse_page_io + ? std::accumulate(column_chunk_data.begin(), + column_chunk_data.end(), + std::size_t{0}, + [](auto size, auto const& page) { return size + page.size(); }) + : [&] { + auto chunk_iter = thrust::make_transform_iterator( + pass.chunks.d_begin(), parquet::detail::get_chunk_compressed_size{}); + return cudf::detail::reduce(chunk_iter, + chunk_iter + pass.chunks.size(), + size_t{0}, + cuda::std::plus{}, + _stream); + }(); + pass.base_mem_size = decomp_dict_data_size + compressed_data_size; // if we are doing subpass reading, generate more accurate num_row estimates for list columns. // this helps us to generate more accurate subpass splits. diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_helpers.hpp b/cpp/src/io/parquet/experimental/hybrid_scan_helpers.hpp index 92287b55bb7a..89c0a42d7fc3 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_helpers.hpp +++ b/cpp/src/io/parquet/experimental/hybrid_scan_helpers.hpp @@ -47,19 +47,6 @@ struct metadata : public metadata_base { class aggregate_reader_metadata : public aggregate_reader_metadata_base { private: - /** - * @brief Check whether selected columns have column and offset indexes - * - * Schema indices are mapped to each source before locating the column chunks. - * - * @param row_group_indices Row group indices, one vector per source - * @param schema_indices Schema indices from the first source - * @return A pair indicating column-index and offset-index presence, respectively - */ - [[nodiscard]] std::pair page_index_presence( - std::span const> row_group_indices, - std::span schema_indices) const; - /** * @brief Filters the row groups using dictionary pages * @@ -91,6 +78,19 @@ class aggregate_reader_metadata : public aggregate_reader_metadata_base { rmm::cuda_stream_view stream) const; public: + /** + * @brief Check whether selected columns have column and offset indexes + * + * Schema indices are mapped to each source before locating the column chunks. + * + * @param row_group_indices Row group indices, one vector per source + * @param schema_indices Schema indices from the first source + * @return A pair indicating column-index and offset-index presence, respectively + */ + [[nodiscard]] std::pair page_index_presence( + std::span const> row_group_indices, + std::span schema_indices) const; + /** * @brief Constructor for aggregate_reader_metadata * diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp b/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp index f2fa1277b1e7..4d8393518a3e 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp +++ b/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp @@ -26,7 +26,6 @@ #include #include -#include #include #include #include @@ -95,6 +94,21 @@ namespace { [](auto sum, auto const& rgs) { return sum + static_cast(rgs.size()); }); } +[[nodiscard]] std::optional> dictionary_page_range( + ColumnChunk const& column) +{ + auto const& page_locations = column.offset_index->page_locations; + if (column.meta_data.dictionary_page_offset > 0) { + auto const offset = column.meta_data.dictionary_page_offset; + return std::pair{offset, column.meta_data.data_page_offset - offset}; + } + if (column.meta_data.data_page_offset < page_locations.front().offset) { + auto const offset = column.meta_data.data_page_offset; + return std::pair{offset, page_locations.front().offset - offset}; + } + return std::nullopt; +} + } // namespace hybrid_scan_reader_impl::hybrid_scan_reader_impl( @@ -219,8 +233,6 @@ std::size_t hybrid_scan_reader_impl::total_rows_in_row_groups( void hybrid_scan_reader_impl::reset_column_selection() { - CUDF_EXPECTS(not _pending_payload_page_io_plan.has_value(), - "Cannot reset column selection while a payload page I/O plan is pending"); _is_all_columns_selected = false; _is_filter_columns_selected = false; _is_payload_columns_selected = false; @@ -246,8 +258,6 @@ void hybrid_scan_reader_impl::prepare_materialization(read_columns_mode read_col rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) { - CUDF_EXPECTS(not _pending_payload_page_io_plan.has_value(), - "Pending payload page I/O plan must be consumed by its setup overload"); reset_internal_state(); initialize_options(options, num_sources, stream, mr); select_columns(read_columns_mode, options); @@ -495,11 +505,10 @@ hybrid_scan_reader_impl::payload_column_chunks_byte_ranges( return get_input_column_chunk_byte_ranges(row_group_indices); } -std::vector> -hybrid_scan_reader_impl::payload_column_chunks_byte_ranges( +std::pair, std::vector> +hybrid_scan_reader_impl::payload_pages_byte_ranges( std::span const> row_group_indices, cudf::column_view const& row_mask, - use_data_page_mask mask_data_pages, parquet_reader_options const& options, rmm::cuda_stream_view stream) { @@ -509,248 +518,103 @@ hybrid_scan_reader_impl::payload_column_chunks_byte_ranges( "Row mask must span across all input row groups"); CUDF_EXPECTS(row_mask.null_count() == 0, "Row mask must not have any nulls when planning payload pages"); - CUDF_EXPECTS(not _pending_payload_page_io_plan.has_value(), - "The previous payload page I/O plan has not been consumed"); select_columns(read_columns_mode::PAYLOAD_COLUMNS, options); - auto column_schemas = std::vector{}; + auto column_schemas = std::vector{}; column_schemas.reserve(_input_columns.size()); std::transform(_input_columns.begin(), _input_columns.end(), std::back_inserter(column_schemas), [](auto const& col) { return col.schema_idx; }); - auto make_full_chunk_plan = [&]() { - auto [flat_ranges, source_map] = get_input_column_chunk_byte_ranges(row_group_indices); - auto source_ranges = std::vector>(row_group_indices.size()); - CUDF_EXPECTS(flat_ranges.size() == source_map.size(), - "Column chunk range source map is invalid"); - for (std::size_t i = 0; i < flat_ranges.size(); ++i) { - CUDF_EXPECTS(std::cmp_less(source_map[i], source_ranges.size()), - "Column chunk range has an invalid source index"); - source_ranges[source_map[i]].push_back(flat_ranges[i]); - } + CUDF_EXPECTS(_extended_metadata->page_index_presence(row_group_indices, column_schemas).second, + "Page-level I/O for payload columns requires offset indexes to be present"); - _pending_payload_page_io_plan = payload_page_io_plan{ - .sparse = false, - .mask_data_pages = mask_data_pages, - .row_group_indices = {row_group_indices.begin(), row_group_indices.end()}, - .column_schema_indices = column_schemas, - .source_ranges = source_ranges, - .page_mappings = {}, - .resident_bytes_per_chunk = {}, - .dictionary_present_per_chunk = {}, - .data_page_mask = {}}; - return source_ranges; - }; - - if (mask_data_pages == use_data_page_mask::NO or row_mask.is_empty()) { - return make_full_chunk_plan(); - } + auto data_page_mask = _extended_metadata->compute_data_page_mask( + row_mask, row_group_indices, _input_columns, 0, stream); - // Sparse page planning only requires offset-index topology. Value counts and variable-width - // sizes can be derived from each retained page after it is fetched. - auto indexes_complete = true; + auto const num_columns = _input_columns.size(); + auto const num_chunks = + static_cast(count_row_groups(row_group_indices)) * num_columns; + auto chunk_page_counts = std::vector(num_chunks); + auto row_group_ordinal = std::size_t{0}; for (std::size_t source_idx = 0; source_idx < row_group_indices.size(); ++source_idx) { + auto colchunk_offsets = std::vector>(num_columns); for (auto const row_group_idx : row_group_indices[source_idx]) { auto const& row_group = _extended_metadata->get_row_group(row_group_idx, source_idx); - for (auto const schema_idx : column_schemas) { - auto const candidate_it = std::find_if( - row_group.columns.begin(), row_group.columns.end(), [schema_idx](auto const& candidate) { - return candidate.schema_idx == schema_idx; - }); - if (candidate_it == row_group.columns.end() or - not candidate_it->offset_index.has_value()) { - indexes_complete = false; - break; - } - auto const& candidate = *candidate_it; - auto const& oi = candidate.offset_index.value(); - auto const num_pages = oi.page_locations.size(); - auto const index_vector_sizes_valid = - not oi.unencoded_byte_array_data_bytes.has_value() or - oi.unencoded_byte_array_data_bytes->size() == num_pages; - auto const dictionary_offsets_valid = - candidate.meta_data.dictionary_page_offset <= 0 or - candidate.meta_data.data_page_offset > candidate.meta_data.dictionary_page_offset; - auto const page_rows_valid = - num_pages > 0 and oi.page_locations.front().first_row_index == 0 and - std::is_sorted(oi.page_locations.begin(), - oi.page_locations.end(), - [](auto const& lhs, auto const& rhs) { - return lhs.first_row_index < rhs.first_row_index; - }) and - std::all_of( - oi.page_locations.begin(), oi.page_locations.end(), [&](auto const& location) { - return location.first_row_index >= 0 and - std::cmp_less_equal(location.first_row_index, row_group.num_rows); - }); - if (num_pages == 0 or not index_vector_sizes_valid or not page_rows_valid or - candidate.meta_data.data_page_offset <= 0 or not dictionary_offsets_valid or - std::any_of( - oi.page_locations.begin(), oi.page_locations.end(), [](auto const& location) { - return location.offset < 0 or location.compressed_page_size <= 0; - })) { - indexes_complete = false; - break; - } + for (std::size_t col_idx = 0; col_idx < num_columns; ++col_idx) { + auto const schema_idx = + _extended_metadata->map_schema_index(column_schemas[col_idx], source_idx); + auto& colchunk_offset = colchunk_offsets[col_idx]; + colchunk_offset = + parquet::detail::find_colchunk_iter_offset(row_group, schema_idx, colchunk_offset); + chunk_page_counts[row_group_ordinal * num_columns + col_idx] = + row_group.columns[colchunk_offset.value()].offset_index->page_locations.size(); } - if (not indexes_complete) { break; } + ++row_group_ordinal; } - if (not indexes_complete) { break; } } - if (not indexes_complete) { return make_full_chunk_plan(); } - - auto data_page_mask = _extended_metadata->compute_data_page_mask( - row_mask, row_group_indices, _input_columns, 0, stream); - // An empty mask is the established representation for "all pages retained". - if (data_page_mask.empty()) { return make_full_chunk_plan(); } - auto const num_columns = _input_columns.size(); - auto const num_row_groups = - std::accumulate(row_group_indices.begin(), - row_group_indices.end(), - std::size_t{0}, - [](auto sum, auto const& groups) { return sum + groups.size(); }); - auto const num_chunks = num_row_groups * num_columns; - auto chunk_masks = std::vector>(num_chunks); - - // Translate the column-major mask into source-major/row-group-major chunk slots once. - std::size_t mask_idx = 0; + auto mask_offsets = std::vector(num_chunks); + std::size_t mask_size{0}; for (std::size_t col_idx = 0; col_idx < num_columns; ++col_idx) { - std::size_t row_group_ordinal = 0; - for (std::size_t source_idx = 0; source_idx < row_group_indices.size(); ++source_idx) { - for (auto const row_group_idx : row_group_indices[source_idx]) { - auto const& row_group = _extended_metadata->get_row_group(row_group_idx, source_idx); - auto const schema_idx = column_schemas[col_idx]; - auto const col = std::find_if( - row_group.columns.begin(), row_group.columns.end(), [schema_idx](auto const& candidate) { - return candidate.schema_idx == schema_idx; - }); - CUDF_EXPECTS(col != row_group.columns.end(), "Selected payload column is missing"); - auto const page_count = col->offset_index->page_locations.size(); - CUDF_EXPECTS(mask_idx + page_count <= data_page_mask.size(), - "Computed data page mask is incomplete"); - auto& mask = chunk_masks[row_group_ordinal * num_columns + col_idx]; - mask.reserve(page_count); - std::transform(data_page_mask.begin() + mask_idx, - data_page_mask.begin() + mask_idx + page_count, - std::back_inserter(mask), - [](bool retained) { return static_cast(retained); }); - mask_idx += page_count; - ++row_group_ordinal; - } + for (std::size_t chunk_idx = col_idx; chunk_idx < num_chunks; chunk_idx += num_columns) { + mask_offsets[chunk_idx] = mask_size; + mask_size += chunk_page_counts[chunk_idx]; } } - // compute_data_page_mask currently leaves unused trailing entries after the logical - // column-major page mask. Preserve the established consumer behavior by discarding them here. - data_page_mask.resize(mask_idx); - - struct exact_request { - int64_t offset; - int64_t size; - std::size_t mapping_idx; - }; - auto exact_requests = std::vector>(row_group_indices.size()); - auto page_mappings = std::vector{}; - auto resident_bytes = std::vector(num_chunks, 0); - auto dictionary_present = std::vector(num_chunks, 0); - - std::size_t row_group_ordinal = 0; + CUDF_EXPECTS(data_page_mask.empty() or data_page_mask.size() == mask_size, + "Computed data page mask does not match offset indexes"); + + auto page_ranges = std::vector{}; + auto source_map = std::vector{}; + page_ranges.reserve(mask_size + num_chunks); + source_map.reserve(mask_size + num_chunks); + + row_group_ordinal = 0; for (std::size_t source_idx = 0; source_idx < row_group_indices.size(); ++source_idx) { + auto colchunk_offsets = std::vector>(num_columns); for (auto const row_group_idx : row_group_indices[source_idx]) { auto const& row_group = _extended_metadata->get_row_group(row_group_idx, source_idx); for (std::size_t col_idx = 0; col_idx < num_columns; ++col_idx) { - auto const chunk_idx = row_group_ordinal * num_columns + col_idx; - auto const schema_idx = column_schemas[col_idx]; - auto const col = std::find_if( - row_group.columns.begin(), row_group.columns.end(), [schema_idx](auto const& candidate) { - return candidate.schema_idx == schema_idx; - }); - CUDF_EXPECTS(col != row_group.columns.end(), "Selected payload column is missing"); - auto const& page_locations = col->offset_index->page_locations; - auto const& retained = chunk_masks[chunk_idx]; - CUDF_EXPECTS(retained.size() == page_locations.size(), - "Data page mask does not match the offset index"); + auto const chunk_idx = row_group_ordinal * num_columns + col_idx; + auto const schema_idx = + _extended_metadata->map_schema_index(column_schemas[col_idx], source_idx); + auto& colchunk_offset = colchunk_offsets[col_idx]; + colchunk_offset = + parquet::detail::find_colchunk_iter_offset(row_group, schema_idx, colchunk_offset); + auto const& column_chunk = row_group.columns[colchunk_offset.value()]; + auto const& page_locations = column_chunk.offset_index->page_locations; + auto const mask_offset = mask_offsets[chunk_idx]; auto const any_retained = - std::any_of(retained.begin(), retained.end(), [](auto value) { return value != 0; }); - - std::optional> dictionary_range; - if (col->meta_data.dictionary_page_offset > 0) { - auto const offset = col->meta_data.dictionary_page_offset; - auto const size = col->meta_data.data_page_offset - offset; - if (size > 0) { dictionary_range = std::pair{offset, size}; } - } else if (col->meta_data.data_page_offset < page_locations.front().offset) { - auto const offset = col->meta_data.data_page_offset; - dictionary_range = - std::pair{offset, page_locations.front().offset - col->meta_data.data_page_offset}; - } + data_page_mask.empty() or + std::any_of(data_page_mask.begin() + mask_offset, + data_page_mask.begin() + mask_offset + page_locations.size(), + [](bool retained) { return retained; }); - auto add_mapping = [&](bool fetched, int64_t offset, int64_t size) { - CUDF_EXPECTS( - offset >= 0 and size > 0 and offset <= std::numeric_limits::max() - size, - "Indexed page byte range is invalid"); - auto const mapping_idx = page_mappings.size(); - page_mappings.push_back( - page_range_mapping{.source_idx = static_cast(source_idx), - .range_idx = 0, - .range_offset = 0, - .size = fetched ? static_cast(size) : 0, - .fetched = fetched}); - if (fetched) { - exact_requests[source_idx].push_back(exact_request{offset, size, mapping_idx}); - resident_bytes[chunk_idx] += static_cast(size); - } + auto const dictionary_range = dictionary_page_range(column_chunk); + + auto add_range = [&](bool retained_page, int64_t offset, int64_t size) { + page_ranges.emplace_back(offset, retained_page ? size : 0); + source_map.push_back(static_cast(source_idx)); }; - if (dictionary_range.has_value() and any_retained) { - add_mapping(true, dictionary_range->first, dictionary_range->second); - dictionary_present[chunk_idx] = 1; + if (dictionary_range.has_value()) { + add_range(any_retained, dictionary_range->first, dictionary_range->second); } for (std::size_t page_idx = 0; page_idx < page_locations.size(); ++page_idx) { auto const& location = page_locations[page_idx]; - add_mapping(retained[page_idx] != 0, - location.offset, - static_cast(location.compressed_page_size)); + add_range(data_page_mask.empty() or data_page_mask[mask_offset + page_idx], + location.offset, + static_cast(location.compressed_page_size)); } } ++row_group_ordinal; } } - auto source_ranges = std::vector>(row_group_indices.size()); - for (std::size_t source_idx = 0; source_idx < exact_requests.size(); ++source_idx) { - auto& requests = exact_requests[source_idx]; - std::stable_sort(requests.begin(), requests.end(), [](auto const& lhs, auto const& rhs) { - return std::tie(lhs.offset, lhs.size) < std::tie(rhs.offset, rhs.size); - }); - for (auto const& request : requests) { - auto& ranges = source_ranges[source_idx]; - if (ranges.empty() or request.offset > ranges.back().offset() + ranges.back().size()) { - ranges.emplace_back(request.offset, request.size); - } else { - auto const end = - std::max(ranges.back().offset() + ranges.back().size(), request.offset + request.size); - ranges.back() = byte_range_info{ranges.back().offset(), end - ranges.back().offset()}; - } - auto& mapping = page_mappings[request.mapping_idx]; - mapping.range_idx = ranges.size() - 1; - mapping.range_offset = static_cast(request.offset - ranges.back().offset()); - } - } - - _pending_payload_page_io_plan = - payload_page_io_plan{.sparse = true, - .mask_data_pages = mask_data_pages, - .row_group_indices = {row_group_indices.begin(), row_group_indices.end()}, - .column_schema_indices = std::move(column_schemas), - .source_ranges = source_ranges, - .page_mappings = std::move(page_mappings), - .resident_bytes_per_chunk = std::move(resident_bytes), - .dictionary_present_per_chunk = std::move(dictionary_present), - .data_page_mask = std::move(data_page_mask)}; - return source_ranges; + return {std::move(page_ranges), std::move(source_map)}; } std::pair, std::vector> @@ -974,89 +838,21 @@ void hybrid_scan_reader_impl::setup_chunking_for_payload_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, - use_data_page_mask mask_data_pages, - std::span> const> page_data_per_source, + std::span const> page_data, parquet_reader_options const& options, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) { - CUDF_EXPECTS(_pending_payload_page_io_plan.has_value(), - "No pending payload page I/O plan to consume"); - // Consume first so a failed setup cannot accidentally reuse stale pointer/range mappings. - auto plan = std::move(_pending_payload_page_io_plan.value()); - _pending_payload_page_io_plan.reset(); - - CUDF_EXPECTS(plan.mask_data_pages == mask_data_pages, - "Payload setup page-mask option does not match its pending plan"); - auto const setup_row_groups = - std::vector>{row_group_indices.begin(), row_group_indices.end()}; - CUDF_EXPECTS(plan.row_group_indices == setup_row_groups, - "Payload setup row groups do not match the pending page I/O plan"); - reset_column_selection(); - select_columns(read_columns_mode::PAYLOAD_COLUMNS, options); - auto selected_schemas = std::vector{}; - selected_schemas.reserve(_input_columns.size()); - std::transform(_input_columns.begin(), - _input_columns.end(), - std::back_inserter(selected_schemas), - [](auto const& col) { return col.schema_idx; }); - CUDF_EXPECTS(selected_schemas == plan.column_schema_indices, - "Payload column selection does not match the pending page I/O plan"); - - CUDF_EXPECTS(page_data_per_source.size() == plan.source_ranges.size(), - "Fetched payload source count does not match the pending plan"); - for (std::size_t source_idx = 0; source_idx < page_data_per_source.size(); ++source_idx) { - CUDF_EXPECTS(page_data_per_source[source_idx].size() == plan.source_ranges[source_idx].size(), - "Fetched payload range count does not match the pending plan"); - for (std::size_t range_idx = 0; range_idx < page_data_per_source[source_idx].size(); - ++range_idx) { - auto const& data = page_data_per_source[source_idx][range_idx]; - auto const& range = plan.source_ranges[source_idx][range_idx]; - CUDF_EXPECTS(std::cmp_equal(data.size(), range.size()), - "Fetched payload span size does not match its planned byte range"); - CUDF_EXPECTS(data.size() == 0 or data.data() != nullptr, - "Fetched payload span has a null data pointer"); - } - } - - if (not plan.sparse) { - auto flat_chunk_data = std::vector>{}; - auto const span_count = - std::accumulate(page_data_per_source.begin(), - page_data_per_source.end(), - std::size_t{0}, - [](auto sum, auto const& spans) { return sum + spans.size(); }); - flat_chunk_data.reserve(span_count); - for (auto const& source_data : page_data_per_source) { - flat_chunk_data.insert(flat_chunk_data.end(), source_data.begin(), source_data.end()); - } - setup_chunking_for_payload_columns(chunk_read_limit, - pass_read_limit, - row_group_indices, - row_mask, - mask_data_pages, - flat_chunk_data, - options, - stream, - mr); - return; - } - - CUDF_EXPECTS(std::cmp_equal(row_mask.size(), total_rows_in_row_groups(row_group_indices)), - "Row mask must span across all input row groups"); - CUDF_EXPECTS(row_mask.null_count() == 0, - "Row mask must not have any nulls when materializing payload column"); - prepare_materialization( read_columns_mode::PAYLOAD_COLUMNS, row_group_indices.size(), options, stream, mr); _input_pass_read_limit = pass_read_limit; _output_chunk_read_limit = chunk_read_limit; - // Preserve the existing all-rows-pruned setup path. An all-false page plan has no byte ranges. - if (are_all_rows_pruned(row_mask, stream)) { + // Return early if all payload pages are pruned. + if (std::all_of( + page_data.begin(), page_data.end(), [](auto const& page) { return page.empty(); })) { auto const empty_row_groups = std::vector>(row_group_indices.size(), std::vector{}); prepare_data(read_mode::CHUNKED_READ, empty_row_groups, {}, {}); @@ -1064,29 +860,23 @@ void hybrid_scan_reader_impl::setup_chunking_for_payload_columns( return; } - _sparse_page_spans.clear(); - _sparse_page_spans.reserve(plan.page_mappings.size()); - for (auto const& mapping : plan.page_mappings) { - if (not mapping.fetched) { - _sparse_page_spans.emplace_back(); - continue; - } - CUDF_EXPECTS(std::cmp_less(mapping.source_idx, page_data_per_source.size()), - "Sparse page mapping has an invalid source index"); - auto const& source_data = page_data_per_source[mapping.source_idx]; - CUDF_EXPECTS(mapping.range_idx < source_data.size(), - "Sparse page mapping has an invalid range index"); - auto const& range_data = source_data[mapping.range_idx]; - CUDF_EXPECTS(mapping.range_offset <= range_data.size() and - mapping.size <= range_data.size() - mapping.range_offset, - "Sparse page mapping exceeds its fetched range"); - _sparse_page_spans.emplace_back(range_data.data() + mapping.range_offset, mapping.size); - } - _sparse_resident_bytes_per_chunk = std::move(plan.resident_bytes_per_chunk); - _sparse_dictionary_present_per_chunk = std::move(plan.dictionary_present_per_chunk); - _sparse_page_io = true; + // Check if offset indexes are present + auto const num_columns = _input_columns.size(); + auto column_schemas = std::vector{}; + column_schemas.reserve(num_columns); + std::transform(_input_columns.begin(), + _input_columns.end(), + std::back_inserter(column_schemas), + [](auto const& col) { return col.schema_idx; }); + CUDF_EXPECTS(_extended_metadata->page_index_presence(row_group_indices, column_schemas).second, + "Page-level I/O for payload columns requires offset indexes to be present"); - prepare_data(read_mode::CHUNKED_READ, row_group_indices, {}, plan.data_page_mask); + // Mark that we are using page-level I/O for payload columns + _sparse_page_io = true; + + // Data page mask in sparse mode will be computed directly from the page data span inside + // `prepare_data() -> setup_sparse_compressed_data()` + prepare_data(read_mode::CHUNKED_READ, row_group_indices, page_data, {}); } table_with_metadata hybrid_scan_reader_impl::materialize_payload_columns_chunk( @@ -1259,9 +1049,6 @@ void hybrid_scan_reader_impl::reset_internal_state() _pass_page_mask.clear(); _subpass_page_mask.reset(); _output_metadata.reset(); - _sparse_page_spans.clear(); - _sparse_resident_bytes_per_chunk.clear(); - _sparse_dictionary_present_per_chunk.clear(); _sparse_page_io = false; _options.timestamp_type = cudf::data_type{}; @@ -1603,15 +1390,7 @@ void hybrid_scan_reader_impl::set_pass_page_mask(std::span data_page data_page_mask.size() >= num_inserted_data_pages + num_data_pages_this_col_chunk, "Encountered invalid data page mask size"); - // Sparse chunks omit dictionaries when every data page is pruned. The contiguous path - // retains its existing conservative dictionary behavior. - if (chunks[chunk_idx].num_dict_pages > 0) { - auto const chunk_has_retained_page = std::any_of( - data_page_mask.begin() + num_inserted_data_pages, - data_page_mask.begin() + num_inserted_data_pages + num_data_pages_this_col_chunk, - [](bool retained) { return retained; }); - _pass_page_mask.push_back(_sparse_page_io ? chunk_has_retained_page : true); - } + if (chunks[chunk_idx].num_dict_pages > 0) { _pass_page_mask.push_back(true); } // Insert page mask for this column chunk _pass_page_mask.insert( @@ -1628,4 +1407,52 @@ void hybrid_scan_reader_impl::set_pass_page_mask(std::span data_page "Encountered mismatch in number of pass pages and page mask size"); } +void hybrid_scan_reader_impl::set_sparse_pass_page_mask( + std::span const> page_data) +{ + auto const& pass = _pass_itm_data; + auto const& chunks = pass->chunks; + + _pass_page_mask = cudf::detail::make_empty_host_vector(pass->pages.size(), _stream); + + // Find the first logical page-data span for every column chunk. + auto page_offsets = std::vector(chunks.size()); + auto chunk_idx = std::size_t{0}; + auto const num_logical_pages = std::accumulate( + chunks.begin(), chunks.end(), std::size_t{0}, [&](auto offset, auto const& chunk) { + page_offsets[chunk_idx++] = offset; + return offset + chunk.num_dict_pages + chunk.num_data_pages; + }); + CUDF_EXPECTS(page_data.size() == num_logical_pages, + "Sparse page span count does not match the number of logical pages"); + + auto const num_columns = _input_columns.size(); + // Build the internal mask in column/chunk order. + std::for_each( + cuda::counting_iterator{0}, + cuda::counting_iterator{num_columns}, + [&](auto col_idx) { + for (std::size_t chunk_idx = col_idx; chunk_idx < chunks.size(); chunk_idx += num_columns) { + auto const& chunk = chunks[chunk_idx]; + auto const data_page_idx = page_offsets[chunk_idx] + chunk.num_dict_pages; + auto const data_page_end = data_page_idx + chunk.num_data_pages; + + // Retain a dictionary whenever the column chunk has a retained data page. + if (chunk.num_dict_pages > 0) { + _pass_page_mask.push_back(std::any_of(page_data.begin() + data_page_idx, + page_data.begin() + data_page_end, + [](auto const& page) { return not page.empty(); })); + } + // Insert page-mask values directly from the corresponding data-page spans. + std::transform(page_data.begin() + data_page_idx, + page_data.begin() + data_page_end, + std::back_inserter(_pass_page_mask), + [](auto const& page) { return not page.empty(); }); + } + }); + // Make sure we inserted exactly the number of pages for this pass. + CUDF_EXPECTS(_pass_page_mask.size() == pass->pages.size(), + "Encountered mismatch in number of pass pages and page mask size"); +} + } // namespace cudf::io::parquet::experimental::detail diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_impl.hpp b/cpp/src/io/parquet/experimental/hybrid_scan_impl.hpp index 5e21411d9fc9..3f3307e49417 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_impl.hpp +++ b/cpp/src/io/parquet/experimental/hybrid_scan_impl.hpp @@ -194,12 +194,11 @@ class hybrid_scan_reader_impl : public parquet::detail::reader_impl { payload_column_chunks_byte_ranges(std::span const> row_group_indices, parquet_reader_options const& options); - [[nodiscard]] std::vector> payload_column_chunks_byte_ranges( - std::span const> row_group_indices, - cudf::column_view const& row_mask, - use_data_page_mask mask_data_pages, - parquet_reader_options const& options, - rmm::cuda_stream_view stream); + [[nodiscard]] std::pair, std::vector> + payload_pages_byte_ranges(std::span const> row_group_indices, + cudf::column_view const& row_mask, + parquet_reader_options const& options, + rmm::cuda_stream_view stream); /** * @copydoc cudf::io::parquet::experimental::hybrid_scan_multifile::materialize_payload_columns @@ -271,9 +270,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, - use_data_page_mask mask_data_pages, - std::span> const> page_data_per_source, + std::span const> page_data, parquet_reader_options const& options, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr); @@ -332,26 +329,6 @@ class hybrid_scan_reader_impl : public parquet::detail::reader_impl { */ enum class read_columns_mode { FILTER_COLUMNS, PAYLOAD_COLUMNS, ALL_COLUMNS }; - struct page_range_mapping { - cudf::size_type source_idx{}; - std::size_t range_idx{}; - std::size_t range_offset{}; - std::size_t size{}; - bool fetched{}; - }; - - struct payload_page_io_plan { - bool sparse{}; - use_data_page_mask mask_data_pages{}; - std::vector> row_group_indices; - std::vector column_schema_indices; - std::vector> source_ranges; - std::vector page_mappings; - std::vector resident_bytes_per_chunk; - std::vector dictionary_present_per_chunk; - thrust::host_vector data_page_mask; - }; - /** * @brief Populate the reader's `_options` config (and related members) from the user options. * @@ -385,10 +362,17 @@ class hybrid_scan_reader_impl : public parquet::detail::reader_impl { /** * @brief Set the page mask for the pass pages * - * @param data_page_mask Input data page mask from page-pruning step + * @param data_page_mask Input data page mask for the current pass */ void set_pass_page_mask(std::span data_page_mask); + /** + * @brief Set the page mask using sparse (page-level) data spans for the current pass + * + * @param page_data Span of device spans of sparse page data + */ + void set_sparse_pass_page_mask(std::span const> page_data); + /** * @brief Select the columns to be read based on the read mode * @@ -484,7 +468,7 @@ class hybrid_scan_reader_impl : public parquet::detail::reader_impl { * * @param mode Value indicating if the data sources are read all at once or chunk by chunk * @param column_chunk_data Device spans of buffers containing column chunk data - * @param data_page_mask Input data page mask from page-pruning step for the current pass + * @param data_page_mask Input data page mask for the current pass */ void handle_chunking(read_mode mode, std::span const> column_chunk_data, @@ -497,9 +481,10 @@ class hybrid_scan_reader_impl : public parquet::detail::reader_impl { * requested set of all row groups. * * @param column_chunk_data Device spans of buffers containing column chunk data + * @param data_page_mask Input data page mask for the current pass */ void setup_next_pass(std::span const> column_chunk_data, - host_span data_page_mask); + std::span data_page_mask); /** * @brief Setup pointers to columns chunks to be processed for this pass. @@ -518,6 +503,13 @@ class hybrid_scan_reader_impl : public parquet::detail::reader_impl { */ void setup_compressed_data(std::span const> column_chunk_data); + /** + * @brief Setup sparse (page-level) data and decode page headers for the current pass. + * + * @param page_data Span of device spans of sparse page data + */ + void setup_sparse_compressed_data(std::span const> page_data); + /** * @brief Reset the internal state of the reader. */ @@ -606,12 +598,6 @@ class hybrid_scan_reader_impl : public parquet::detail::reader_impl { bool _is_filter_columns_selected{false}; bool _is_payload_columns_selected{false}; bool _is_all_columns_selected{false}; - - std::optional _pending_payload_page_io_plan; - std::vector> _sparse_page_spans; - std::vector _sparse_resident_bytes_per_chunk; - std::vector _sparse_dictionary_present_per_chunk; - bool _sparse_page_io{false}; }; } // namespace cudf::io::parquet::experimental::detail diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_multifile.cpp b/cpp/src/io/parquet/experimental/hybrid_scan_multifile.cpp index 9f2d20cb70cf..ea6f4bf8204b 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_multifile.cpp +++ b/cpp/src/io/parquet/experimental/hybrid_scan_multifile.cpp @@ -137,17 +137,15 @@ hybrid_scan_multifile::payload_column_chunks_byte_ranges( return _impl->payload_column_chunks_byte_ranges(row_group_indices, options); } -std::vector> -hybrid_scan_multifile::payload_column_chunks_byte_ranges( +std::pair, std::vector> +hybrid_scan_multifile::payload_pages_byte_ranges( cudf::host_span const> row_group_indices, cudf::column_view const& row_mask, - use_data_page_mask mask_data_pages, parquet_reader_options const& options, rmm::cuda_stream_view stream) const { CUDF_FUNC_RANGE(); - return _impl->payload_column_chunks_byte_ranges( - row_group_indices, row_mask, mask_data_pages, options, stream); + return _impl->payload_pages_byte_ranges(row_group_indices, row_mask, options, stream); } table_with_metadata hybrid_scan_multifile::materialize_payload_columns( @@ -241,23 +239,14 @@ void hybrid_scan_multifile::setup_chunking_for_payload_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, - use_data_page_mask mask_data_pages, - cudf::host_span> const> page_data_per_source, + cudf::host_span const> page_data, parquet_reader_options const& options, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) const { CUDF_FUNC_RANGE(); - _impl->setup_chunking_for_payload_columns(chunk_read_limit, - pass_read_limit, - row_group_indices, - row_mask, - mask_data_pages, - page_data_per_source, - options, - stream, - mr); + _impl->setup_chunking_for_payload_columns( + chunk_read_limit, pass_read_limit, row_group_indices, page_data, options, stream, mr); } table_with_metadata hybrid_scan_multifile::materialize_payload_columns_chunk( diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu b/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu index 6cc8cc236e62..300fc69f1926 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu +++ b/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu @@ -179,43 +179,6 @@ void hybrid_scan_reader_impl::setup_compressed_data( auto& chunks = pass.chunks; - if (_sparse_page_io) { - CUDF_EXPECTS(_has_page_index, "Sparse page I/O requires complete page indexes"); - CUDF_EXPECTS(_sparse_resident_bytes_per_chunk.size() == chunks.size(), - "Sparse resident-byte accounting does not match the logical chunks"); - CUDF_EXPECTS(_sparse_dictionary_present_per_chunk.size() == chunks.size(), - "Sparse dictionary mapping does not match the logical chunks"); - pass.has_compressed_data = false; - for (std::size_t chunk_idx = 0; chunk_idx < chunks.size(); ++chunk_idx) { - auto& chunk = chunks[chunk_idx]; - chunk.compressed_data = nullptr; - chunk.compressed_size = _sparse_resident_bytes_per_chunk[chunk_idx]; - pass.has_compressed_data |= - chunk.codec != Compression::UNCOMPRESSED and chunk.compressed_size > 0; - } - - auto const indexed_total_pages = count_page_headers_with_pgidx(chunks, _stream); - auto total_pages = std::size_t{0}; - for (std::size_t chunk_idx = 0; chunk_idx < chunks.size(); ++chunk_idx) { - chunks[chunk_idx].num_dict_pages = _sparse_dictionary_present_per_chunk[chunk_idx] ? 1 : 0; - total_pages += chunks[chunk_idx].num_data_pages + chunks[chunk_idx].num_dict_pages; - } - CUDF_EXPECTS(total_pages <= indexed_total_pages, - "Sparse dictionary mapping exceeds page-index metadata"); - chunks.host_to_device_async(_stream); - CUDF_EXPECTS(total_pages == _sparse_page_spans.size(), - "Sparse page span count does not match page-index metadata"); - if (total_pages <= 0) { return; } - // `decode_page_headers` may not write every byte of each PageInfo, and `sort_pages` copies - // PageInfo as whole objects. - auto unsorted_pages = cudf::detail::make_zeroed_device_uvector_async( - total_pages, _stream, cudf::get_current_device_resource_ref()); - parquet::detail::decode_page_headers(pass, unsorted_pages, _sparse_page_spans, _stream); - CUDF_EXPECTS(pass.page_offsets.size() - 1 == static_cast(_input_columns.size()), - "Encountered page_offsets / num_columns mismatch"); - return; - } - pass.has_compressed_data = setup_column_chunks(column_chunk_data); // Process dataset chunk pages into output columns @@ -230,6 +193,40 @@ void hybrid_scan_reader_impl::setup_compressed_data( "Encountered page_offsets / num_columns mismatch"); } +void hybrid_scan_reader_impl::setup_sparse_compressed_data( + std::span const> page_data) +{ + auto& pass = *_pass_itm_data; + + CUDF_EXPECTS(_has_offset_index, "Sparse page I/O requires complete offset indexes"); + + auto& chunks = pass.chunks; + auto const total_pages = count_page_headers_with_pgidx(chunks, _stream); + CUDF_EXPECTS(total_pages == page_data.size(), + "Sparse page span count does not match page-index metadata"); + if (total_pages == 0) { return; } + + pass.has_compressed_data = false; + std::size_t page_idx = 0; + + for (auto const& chunk : chunks) { + auto const num_pages = chunk.num_data_pages + chunk.num_dict_pages; + auto const has_resident_page = std::any_of(page_data.begin() + page_idx, + page_data.begin() + page_idx + num_pages, + [](auto const& page) { return not page.empty(); }); + pass.has_compressed_data |= chunk.codec != Compression::UNCOMPRESSED and has_resident_page; + page_idx += num_pages; + } + + // `decode_page_headers` may not write every byte of each PageInfo, and `sort_pages` copies + // PageInfo as whole objects. + auto unsorted_pages = cudf::detail::make_zeroed_device_uvector_async( + total_pages, _stream, cudf::get_current_device_resource_ref()); + parquet::detail::decode_page_headers(pass, unsorted_pages, page_data, _stream); + CUDF_EXPECTS(pass.page_offsets.size() - 1 == static_cast(_input_columns.size()), + "Encountered page_offsets / num_columns mismatch"); +} + std::tuple, cudf::detail::hostdevice_vector> diff --git a/cpp/src/io/parquet/experimental/page_index_filter.cu b/cpp/src/io/parquet/experimental/page_index_filter.cu index c8374da7b2bd..337359c15f67 100644 --- a/cpp/src/io/parquet/experimental/page_index_filter.cu +++ b/cpp/src/io/parquet/experimental/page_index_filter.cu @@ -1176,7 +1176,8 @@ thrust::host_vector aggregate_reader_metadata::compute_data_page_mask( // 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(total_pages, stream); + auto data_page_mask = thrust::host_vector(0, stream); + data_page_mask.reserve(total_pages); auto host_results_iter = host_results.begin(); stream.synchronize(); diff --git a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp index b6f0413c57c1..e6d12f0bba98 100644 --- a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp +++ b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp @@ -341,7 +341,8 @@ fetch_byte_ranges_to_device_async_impl( column_chunk_data.reserve(byte_ranges.size()); std::ignore = std::accumulate( byte_ranges.begin(), byte_ranges.end(), std::size_t{0}, [&](auto acc, auto const& range) { - column_chunk_data.emplace_back(buffer_data + acc, static_cast(range.size())); + auto const data = buffer_data == nullptr ? nullptr : buffer_data + acc; + column_chunk_data.emplace_back(data, static_cast(range.size())); return acc + range.size(); }); diff --git a/cpp/src/io/parquet/page_hdr.cu b/cpp/src/io/parquet/page_hdr.cu index ecad3d02f6a4..b69c26a4d5eb 100644 --- a/cpp/src/io/parquet/page_hdr.cu +++ b/cpp/src/io/parquet/page_hdr.cu @@ -833,86 +833,6 @@ struct decode_from_page_data_fn { } }; -/** - * @brief Functor to decode indexed page headers from exact page spans - */ -struct decode_page_headers_with_pgidx_spans_fn { - cudf::device_span colchunks; - cudf::device_span pages; - cudf::device_span const> page_spans; - size_type* chunk_page_offsets; - kernel_error::pointer error_code; - - __device__ void operator()(size_type page_idx) const noexcept - { - auto const num_chunks = static_cast(colchunks.size()); - auto const chunk_idx = static_cast( - cuda::std::distance( - chunk_page_offsets, - thrust::upper_bound( - thrust::seq, chunk_page_offsets, chunk_page_offsets + num_chunks + 1, page_idx)) - - 1); - - if (chunk_idx < 0 or chunk_idx >= num_chunks) { - set_error(static_cast(decode_error::DATA_STREAM_OVERRUN), - error_code); - return; - } - - byte_stream_s bs{}; - bs.ck = colchunks[chunk_idx]; - zero_out_page_header_info(&bs); - bs.page.chunk_idx = chunk_idx; - bs.page.src_col_schema = bs.ck.src_col_schema; - - auto const span = page_spans[page_idx]; - if (span.empty()) { - // Preserve the logical page entry. Page-index metadata is filled in by fill_in_page_info(). - pages[page_idx] = bs.page; - return; - } - - bs.base = bs.cur = span.data(); - bs.end = span.data() + span.size(); - - if (not parse_valid_page_header(&bs)) { - set_error(static_cast(decode_error::INVALID_PAGE_HEADER), - error_code); - return; - } - if (not is_supported_encoding(bs.page.encoding)) { - set_error(static_cast(decode_error::UNSUPPORTED_ENCODING), - error_code); - return; - } - - switch (bs.page_type) { - case PageType::DATA_PAGE: bs.page.num_rows = bs.page.num_input_values; break; - case PageType::DATA_PAGE_V2: - bs.page.flags |= PAGEINFO_FLAGS_V2; - bs.page.definition_level_encoding = Encoding::RLE; - bs.page.repetition_level_encoding = Encoding::RLE; - break; - case PageType::DICTIONARY_PAGE: bs.page.flags |= PAGEINFO_FLAGS_DICTIONARY; break; - default: - set_error(static_cast(decode_error::INVALID_PAGE_TYPE), - error_code); - return; - } - - if (bs.page.compressed_page_size < 0 or - static_cast(bs.end - bs.cur) != static_cast(bs.page.compressed_page_size)) { - set_error(static_cast(decode_error::DATA_STREAM_OVERRUN), - error_code); - return; - } - - bs.page.page_data = const_cast(bs.cur); - bs.page.kernel_mask = kernel_mask_for_page(bs.page, bs.ck); - pages[page_idx] = bs.page; - } -}; - /** * @brief Kernel for building dictionary index for the specified column chunks * @@ -1049,26 +969,6 @@ void decode_page_headers_from_page_data( .error_code = error_code}); } -void decode_page_headers_with_pgidx_spans(cudf::device_span chunks, - cudf::device_span pages, - cudf::device_span const> - page_spans, - size_type* chunk_page_offsets, - kernel_error::pointer error_code, - rmm::cuda_stream_view stream) -{ - CUDF_EXPECTS(page_spans.size() == pages.size(), - "Page span count must match the number of logical pages"); - thrust::for_each(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), - cuda::counting_iterator{0}, - cuda::counting_iterator{static_cast(pages.size())}, - decode_page_headers_with_pgidx_spans_fn{.colchunks = chunks, - .pages = pages, - .page_spans = page_spans, - .chunk_page_offsets = chunk_page_offsets, - .error_code = error_code}); -} - void build_string_dictionary_index(ColumnChunkDesc* chunks, int32_t num_chunks, kernel_error::pointer error_code, diff --git a/cpp/src/io/parquet/parquet_gpu.hpp b/cpp/src/io/parquet/parquet_gpu.hpp index e448fc2a1b55..b8324b3f570a 100644 --- a/cpp/src/io/parquet/parquet_gpu.hpp +++ b/cpp/src/io/parquet/parquet_gpu.hpp @@ -731,19 +731,6 @@ void decode_page_headers_from_page_data( kernel_error::pointer error_code, rmm::cuda_stream_view stream); -/** - * @brief Decode indexed page headers from exact, potentially discontiguous page spans - * - * Empty spans initialize the corresponding logical page descriptor but are not parsed. - */ -void decode_page_headers_with_pgidx_spans(cudf::device_span chunks, - cudf::device_span pages, - cudf::device_span const> - page_spans, - size_type* chunk_page_offsets, - kernel_error::pointer error_code, - rmm::cuda_stream_view stream); - /** * @brief Launches kernel for building the dictionary index for the column * chunks diff --git a/cpp/src/io/parquet/reader_impl.cpp b/cpp/src/io/parquet/reader_impl.cpp index 347fd9185a6c..02485a65637c 100644 --- a/cpp/src/io/parquet/reader_impl.cpp +++ b/cpp/src/io/parquet/reader_impl.cpp @@ -462,9 +462,6 @@ void reader_impl::decode_page_data(read_mode mode, size_t skip_rows, size_t num_ cudf::detail::make_pinned_vector(cudf::host_span{out_buffers}, _stream); write_final_offsets(pinned_final_offsets, pinned_out_buffers, _stream); - // For page-level I/O, fill output string and list offsets for pruned pages - fill_pruned_offsets(skip_rows, num_rows); - // update null counts in the final column buffers for (size_t idx = 0; idx < subpass.pages.size(); idx++) { PageInfo* pi = &subpass.pages[idx]; diff --git a/cpp/src/io/parquet/reader_impl.hpp b/cpp/src/io/parquet/reader_impl.hpp index 76d88f52e310..e62cc5f59a30 100644 --- a/cpp/src/io/parquet/reader_impl.hpp +++ b/cpp/src/io/parquet/reader_impl.hpp @@ -582,6 +582,9 @@ class reader_impl { // are offset indexes available for selected row groups bool _has_offset_index = false; + // whether sparse page I/O is enabled + bool _sparse_page_io = false; + std::optional> _reader_column_schema; // chunked reading happens in 2 parts: diff --git a/cpp/src/io/parquet/reader_impl_preprocess.cu b/cpp/src/io/parquet/reader_impl_preprocess.cu index ef5a30f8c55e..0dcfc4f5d447 100644 --- a/cpp/src/io/parquet/reader_impl_preprocess.cu +++ b/cpp/src/io/parquet/reader_impl_preprocess.cu @@ -17,7 +17,6 @@ #include #include #include -#include #include #include #include @@ -107,7 +106,8 @@ void reader_impl::build_string_dict_indices() rmm::exec_policy_nosync(_stream, cudf::get_current_device_resource_ref()), iter, iter + pass.chunks.size(), - set_str_dict_index_ptr{pass.str_dict_index.data(), str_dict_index_offsets, pass.chunks}); + set_str_dict_index_ptr{ + pass.str_dict_index.data(), str_dict_index_offsets, pass.chunks, _sparse_page_io}); // compute the indices kernel_error error_code(_stream); @@ -947,10 +947,6 @@ void reader_impl::allocate_columns(read_mode mode, size_t skip_rows, size_t num_ bool has_lists = false; // Validity Buffer is a uint32_t pointer std::vector> nullmask_bufs; - auto const page_mask = subpass_page_mask_span(); - auto const has_pruned_page = - not page_mask.is_empty() and - std::any_of(page_mask.host_begin(), page_mask.host_end(), [](bool keep) { return not keep; }); for (auto const& input_col : _input_columns) { size_t const max_depth = input_col.nesting_depth(); @@ -974,10 +970,8 @@ void reader_impl::allocate_columns(read_mode mode, size_t skip_rows, size_t num_ CUDF_EXPECTS(out_buf_size <= std::numeric_limits::max(), "Number of rows exceeds cudf's column size limit", std::overflow_error); - auto const initialize_offsets = has_pruned_page and (out_buf.type.id() == type_id::STRING or - out_buf.type.id() == type_id::LIST); out_buf.create_with_mask( - out_buf_size, cudf::mask_state::UNINITIALIZED, initialize_offsets, _stream, _mr); + out_buf_size, cudf::mask_state::UNINITIALIZED, false, _stream, _mr); nullmask_bufs.emplace_back( out_buf.null_mask(), cudf::util::round_up_safe(out_buf.null_mask_size(), sizeof(cudf::bitmask_type)) / @@ -1095,11 +1089,8 @@ void reader_impl::allocate_columns(read_mode mode, size_t skip_rows, size_t num_ std::overflow_error); // allocate // we're going to start null mask as all valid and then turn bits off if necessary - auto const initialize_offsets = - has_pruned_page and - (out_buf.type.id() == type_id::STRING or out_buf.type.id() == type_id::LIST); out_buf.create_with_mask( - buffer_size, cudf::mask_state::UNINITIALIZED, initialize_offsets, _stream, _mr); + buffer_size, cudf::mask_state::UNINITIALIZED, false, _stream, _mr); nullmask_bufs.emplace_back( out_buf.null_mask(), cudf::util::round_up_safe(out_buf.null_mask_size(), sizeof(cudf::bitmask_type)) / diff --git a/cpp/src/io/parquet/reader_impl_preprocess_utils.cu b/cpp/src/io/parquet/reader_impl_preprocess_utils.cu index 9ee4ceecf87d..8236cb0ef66e 100644 --- a/cpp/src/io/parquet/reader_impl_preprocess_utils.cu +++ b/cpp/src/io/parquet/reader_impl_preprocess_utils.cu @@ -429,13 +429,13 @@ enum class page_data_source_type : uint8_t { * * @param pass Struct containing pass information * @param unsorted_pages Device span of page information to decode - * @param page_data Host span of page data spans (only used for PAGE_SPANS source) + * @param page_data Span of page data spans (only used for PAGE_SPANS source) * @param stream Stream to use */ template void decode_page_headers_impl(pass_intermediate_data& pass, device_span unsorted_pages, - host_span const> page_data, + std::span const> page_data, rmm::cuda_stream_view stream) { CUDF_FUNC_RANGE(); @@ -641,7 +641,7 @@ void decode_page_headers(pass_intermediate_data& pass, void decode_page_headers(pass_intermediate_data& pass, device_span unsorted_pages, - host_span const> page_data, + std::span const> page_data, rmm::cuda_stream_view stream) { decode_page_headers_impl( diff --git a/cpp/src/io/parquet/reader_impl_preprocess_utils.cuh b/cpp/src/io/parquet/reader_impl_preprocess_utils.cuh index 518890b5db67..16ca0946898a 100644 --- a/cpp/src/io/parquet/reader_impl_preprocess_utils.cuh +++ b/cpp/src/io/parquet/reader_impl_preprocess_utils.cuh @@ -15,6 +15,7 @@ #include #include +#include #include namespace cudf::io::parquet::detail { @@ -148,17 +149,7 @@ void decode_page_headers(pass_intermediate_data& pass, */ void decode_page_headers(pass_intermediate_data& pass, device_span unsorted_pages, - host_span const> page_data, - rmm::cuda_stream_view stream); - -/** - * @brief Decode page information using one exact span per logical indexed page - * - * Empty spans represent masked pages and retain their logical page-index metadata. - */ -void decode_page_headers(pass_intermediate_data& pass, - device_span unsorted_pages, - host_span const> page_spans, + std::span const> page_data, rmm::cuda_stream_view stream); /** @@ -235,10 +226,15 @@ struct set_str_dict_index_ptr { string_index_pair* const base; device_span str_dict_index_offsets; device_span chunks; + bool const sparse_page_io; __device__ constexpr inline void operator()(size_t i) { auto& chunk = chunks[i]; + // In Sparse I/O case, the dictionary page may be null if all data pages were pruned. + if (sparse_page_io and (chunk.dict_page == nullptr or chunk.dict_page->page_data == nullptr)) { + return; + } if (chunk.num_dict_pages > 0 and is_string_chunk(chunk)) { chunk.str_dict_index = base + str_dict_index_offsets[i]; } diff --git a/cpp/tests/io/experimental/hybrid_scan_common.cpp b/cpp/tests/io/experimental/hybrid_scan_common.cpp index 7f61477f6a8d..339009ecf7e5 100644 --- a/cpp/tests/io/experimental/hybrid_scan_common.cpp +++ b/cpp/tests/io/experimental/hybrid_scan_common.cpp @@ -199,15 +199,6 @@ multisource_device_data fetch_multisource_device_data( { auto const byte_ranges_per_source = group_byte_ranges_by_source(byte_ranges_and_source_map, inputs.datasources.size()); - return fetch_multisource_device_data(inputs, byte_ranges_per_source, stream, mr); -} - -multisource_device_data fetch_multisource_device_data( - multifile_inputs const& inputs, - std::vector> const& byte_ranges_per_source, - rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr) -{ auto [buffers, per_source_spans, tasks] = cudf::io::parquet::fetch_byte_ranges_to_device_async( inputs.datasource_refs, cudf::host_span const>{byte_ranges_per_source}, diff --git a/cpp/tests/io/experimental/hybrid_scan_common.hpp b/cpp/tests/io/experimental/hybrid_scan_common.hpp index bba44361c8ae..8a8493977a4b 100644 --- a/cpp/tests/io/experimental/hybrid_scan_common.hpp +++ b/cpp/tests/io/experimental/hybrid_scan_common.hpp @@ -98,15 +98,6 @@ void setup_page_indexes(cudf::io::parquet::experimental::hybrid_scan_multifile c rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr); -/** - * @brief Fetches per-source byte ranges and returns per-source and flattened spans - */ -[[nodiscard]] multisource_device_data fetch_multisource_device_data( - multifile_inputs const& inputs, - std::vector> const& byte_ranges_per_source, - rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr); - /** * @brief Concatenate a vector of tables and return the resultant table * diff --git a/cpp/tests/io/experimental/hybrid_scan_multifile_composer.cpp b/cpp/tests/io/experimental/hybrid_scan_multifile_composer.cpp index 076b66c87bb6..9832ec782f2e 100644 --- a/cpp/tests/io/experimental/hybrid_scan_multifile_composer.cpp +++ b/cpp/tests/io/experimental/hybrid_scan_multifile_composer.cpp @@ -159,7 +159,7 @@ chunked_hybrid_scan_multifile(cudf::io::source_info const& source_info, } std::tuple, std::unique_ptr> -page_level_chunked_hybrid_scan_multifile( +chunked_sparse_hybrid_scan_multifile( cudf::io::source_info const& source_info, cudf::ast::operation const& filter_expression, std::optional> const& payload_column_names, @@ -204,16 +204,14 @@ page_level_chunked_hybrid_scan_multifile( filter_tables.push_back(reader.materialize_filter_columns_chunk(row_mask_view).tbl); } - auto const payload_page_ranges = reader.payload_column_chunks_byte_ranges( - row_groups, row_mask->view(), use_data_page_mask::YES, options, stream); + auto const payload_page_ranges = + reader.payload_pages_byte_ranges(row_groups, row_mask->view(), options, stream); auto payload_page_data = fetch_multisource_device_data(inputs, payload_page_ranges, stream, mr); reader.setup_chunking_for_payload_columns(chunk_read_limit, pass_read_limit, row_groups, - row_mask->view(), - use_data_page_mask::YES, - payload_page_data.per_source_spans, + payload_page_data.flat_spans, options, stream, mr); diff --git a/cpp/tests/io/experimental/hybrid_scan_multifile_composer.hpp b/cpp/tests/io/experimental/hybrid_scan_multifile_composer.hpp index fc87ed381575..5287cd684f7b 100644 --- a/cpp/tests/io/experimental/hybrid_scan_multifile_composer.hpp +++ b/cpp/tests/io/experimental/hybrid_scan_multifile_composer.hpp @@ -95,7 +95,7 @@ chunked_hybrid_scan_multifile(cudf::io::source_info const& source_info, * @return Tuple of filter and payload tables */ std::tuple, std::unique_ptr> -page_level_chunked_hybrid_scan_multifile( +chunked_sparse_hybrid_scan_multifile( cudf::io::source_info const& source_info, cudf::ast::operation const& filter_expression, std::optional> const& payload_column_names, diff --git a/cpp/tests/io/experimental/hybrid_scan_multifile_filters_test.cpp b/cpp/tests/io/experimental/hybrid_scan_multifile_filters_test.cpp index baf05e6188b9..82ee5c972bd2 100644 --- a/cpp/tests/io/experimental/hybrid_scan_multifile_filters_test.cpp +++ b/cpp/tests/io/experimental/hybrid_scan_multifile_filters_test.cpp @@ -528,6 +528,28 @@ TEST_F(HybridScanMultifileFiltersTest, BuildAllTrueRowMask) test_all_true_row_mask(row_group_indices); } +TEST_F(HybridScanMultifileFiltersTest, SparsePayloadPagesWithoutOffsetIndexes) +{ + using T = uint32_t; + + auto file_buffers = std::vector>{}; + file_buffers.emplace_back(std::get<1>(create_parquet_with_stats())); + auto inputs = multifile_inputs(build_source_info(file_buffers)); + + auto const options = cudf::io::parquet_reader_options::builder().column_names({"col1"}).build(); + auto reader = + cudf::io::parquet::experimental::hybrid_scan_multifile{inputs.footer_byte_spans, options}; + + 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 const row_mask = reader.build_all_true_row_mask(row_groups, stream, mr); + + EXPECT_THROW( + std::ignore = reader.payload_pages_byte_ranges(row_groups, row_mask->view(), options, stream), + cudf::logic_error); +} + template struct HybridScanMultifilePageIndexRowMaskTest : public HybridScanMultifileFiltersTest {}; diff --git a/cpp/tests/io/experimental/hybrid_scan_multifile_test.cpp b/cpp/tests/io/experimental/hybrid_scan_multifile_test.cpp index 84285e197350..8652f71878f5 100644 --- a/cpp/tests/io/experimental/hybrid_scan_multifile_test.cpp +++ b/cpp/tests/io/experimental/hybrid_scan_multifile_test.cpp @@ -16,7 +16,6 @@ #include #include #include -#include #include #include #include @@ -38,62 +37,6 @@ namespace { using cudf::io::parquet::experimental::use_data_page_mask; -std::pair payload_byte_range_sizes( - cudf::io::source_info const& source_info, - cudf::ast::operation const& filter_expression, - bool case_sensitive_names, - rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr) -{ - auto options = cudf::io::parquet_reader_options::builder() - .filter(filter_expression) - .case_sensitive_names(case_sensitive_names) - .build(); - auto inputs = multifile_inputs(source_info); - auto reader = - cudf::io::parquet::experimental::hybrid_scan_multifile{inputs.footer_byte_spans, options}; - setup_page_indexes(reader, inputs); - - auto const input_row_groups = reader.all_row_groups(options); - auto const row_groups = reader.filter_row_groups_with_stats(input_row_groups, options, stream); - auto row_mask = reader.build_row_mask_with_page_index_stats(row_groups, options, stream, mr); - - auto filter_data = fetch_multisource_device_data( - inputs, reader.filter_column_chunks_byte_ranges(row_groups, options), stream, mr); - auto row_mask_view = row_mask->mutable_view(); - reader.setup_chunking_for_filter_columns(256 * 1024, - 1024 * 1024, - row_groups, - row_mask_view, - use_data_page_mask::YES, - filter_data.flat_spans, - options, - stream, - mr); - while (reader.has_next_table_chunk()) { - static_cast(reader.materialize_filter_columns_chunk(row_mask_view)); - } - - auto const full_ranges = reader.payload_column_chunks_byte_ranges(row_groups, options).first; - auto const page_ranges = reader.payload_column_chunks_byte_ranges( - row_groups, row_mask->view(), use_data_page_mask::YES, options, stream); - auto const full_bytes = std::accumulate( - full_ranges.begin(), full_ranges.end(), std::size_t{0}, [](auto sum, auto range) { - return sum + range.size(); - }); - auto const requested_bytes = std::accumulate( - page_ranges.begin(), - page_ranges.end(), - std::size_t{0}, - [](auto source_sum, auto const& source_ranges) { - return source_sum + std::accumulate(source_ranges.begin(), - source_ranges.end(), - std::size_t{0}, - [](auto sum, auto range) { return sum + range.size(); }); - }); - return {requested_bytes, full_bytes}; -} - std::vector> make_plain_payload_parquet_buffers() { auto constexpr num_sources = 2; @@ -127,20 +70,6 @@ std::vector> make_plain_payload_parquet_buffers() return parquet_buffers; } -void expect_byte_ranges_equal( - std::vector> const& expected, - std::vector> const& actual) -{ - ASSERT_EQ(expected.size(), actual.size()); - for (std::size_t source_idx = 0; source_idx < expected.size(); ++source_idx) { - ASSERT_EQ(expected[source_idx].size(), actual[source_idx].size()); - for (std::size_t range_idx = 0; range_idx < expected[source_idx].size(); ++range_idx) { - EXPECT_EQ(expected[source_idx][range_idx].offset(), actual[source_idx][range_idx].offset()); - EXPECT_EQ(expected[source_idx][range_idx].size(), actual[source_idx][range_idx].size()); - } - } -} - /** * @brief Helper to test multifile hybrid scan single-shot materialization * @@ -153,9 +82,8 @@ void expect_byte_ranges_equal( */ template void test_hybrid_scan_multifile(std::vector const& columns, - bool case_sensitive_names = true, - uint32_t literal_value = 100, - bool expect_payload_byte_reduction = false) + bool case_sensitive_names = true, + uint32_t literal_value = 100) { auto const table = cudf::table_view{columns}; cudf::io::table_input_metadata expected_metadata(table); @@ -199,17 +127,16 @@ void test_hybrid_scan_multifile(std::vector const& columns, auto const [chunked_filter_table, chunked_payload_table] = chunked_hybrid_scan_multifile( source_info, filter_expression, {}, case_sensitive_names, stream, mr); - auto const [page_level_filter_table, page_level_payload_table] = - page_level_chunked_hybrid_scan_multifile( - source_info, filter_expression, {}, case_sensitive_names, stream, mr); + auto const [sparse_filter_table, sparse_payload_table] = chunked_sparse_hybrid_scan_multifile( + source_info, filter_expression, {}, case_sensitive_names, stream, mr); auto const chunked_all_table = chunked_hybrid_scan_multifile_single_step( source_info, filter_expression, {}, case_sensitive_names, stream, mr); CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected.tbl->select({0}), filter_table->view()); CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected.tbl->select({0}), chunked_filter_table->view()); - CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected.tbl->select({0}), page_level_filter_table->view()); - CUDF_TEST_EXPECT_TABLES_EQUIVALENT(chunked_filter_table->view(), page_level_filter_table->view()); + CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected.tbl->select({0}), sparse_filter_table->view()); + CUDF_TEST_EXPECT_TABLES_EQUIVALENT(chunked_filter_table->view(), sparse_filter_table->view()); auto payload_column_indices = std::vector(columns.size() - 1); std::iota(payload_column_indices.begin(), payload_column_indices.end(), 1); @@ -218,17 +145,10 @@ void test_hybrid_scan_multifile(std::vector const& columns, CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected.tbl->select(payload_column_indices), chunked_payload_table->view()); CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected.tbl->select(payload_column_indices), - page_level_payload_table->view()); - CUDF_TEST_EXPECT_TABLES_EQUIVALENT(chunked_payload_table->view(), - page_level_payload_table->view()); + sparse_payload_table->view()); + CUDF_TEST_EXPECT_TABLES_EQUIVALENT(chunked_payload_table->view(), sparse_payload_table->view()); CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected.tbl->view(), all_table->view()); CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected.tbl->view(), chunked_all_table->view()); - if (expect_payload_byte_reduction) { - auto const [requested_payload_bytes, full_payload_bytes] = - payload_byte_range_sizes(source_info, filter_expression, case_sensitive_names, stream, mr); - EXPECT_GT(requested_payload_bytes, 0); - EXPECT_LT(requested_payload_bytes, full_payload_bytes); - } } } // namespace @@ -282,233 +202,7 @@ TEST_F(HybridScanMultifileTest, MaterializeListsOfStrings) test_hybrid_scan_multifile({col0, *col1, *col2, *col3, *col4}, false); } -TEST_F(HybridScanMultifileTest, PageLevelDictionaryPayloadByteReduction) -{ - auto col0 = testdata::ascending(); - - auto payload_values = std::vector(num_ordered_rows); - for (auto i = std::size_t{0}; i < payload_values.size(); ++i) { - payload_values[i] = "dictionary value " + std::to_string(i % 8); - } - auto col1 = cudf::test::strings_column_wrapper(payload_values.begin(), payload_values.end()); - - // A page-aligned threshold retains two of four data pages. The writer's ALWAYS dictionary policy - // requires the page-I/O path to retain the dictionary while requesting fewer bytes than the - // legacy full-column-chunk path. - auto constexpr threshold = uint32_t{2 * page_size_for_ordered_tests / 100}; - test_hybrid_scan_multifile({col0, col1}, true, threshold, true); -} - -TEST_F(HybridScanMultifileTest, PageLevelStringsSeparatedByPrunedPages) -{ - auto filter_values = cudf::detail::make_counting_transform_iterator( - cudf::size_type{0}, [](auto i) { return (i / page_size_for_ordered_tests) % 2 == 0; }); - auto filter = - cudf::test::fixed_width_column_wrapper(filter_values, filter_values + num_ordered_rows); - - auto payload_values = std::vector(num_ordered_rows); - for (auto i = std::size_t{0}; i < payload_values.size(); ++i) { - payload_values[i] = "payload value " + std::to_string(i); - } - auto payload = cudf::test::strings_column_wrapper(payload_values.begin(), payload_values.end()); - auto table = cudf::table_view{{filter, payload}}; - - auto metadata = cudf::io::table_input_metadata(table); - metadata.column_metadata[0].set_name("filter"); - metadata.column_metadata[1].set_name("payload"); - - auto parquet_buffers = std::vector>(2); - for (auto& parquet_buffer : parquet_buffers) { - auto options = - cudf::io::parquet_writer_options::builder(cudf::io::sink_info{&parquet_buffer}, table) - .metadata(metadata) - .row_group_size_rows(num_ordered_rows) - .max_page_size_rows(page_size_for_ordered_tests) - .max_page_size_bytes(64 * 1024 * 1024) - .compression(cudf::io::compression_type::NONE) - .dictionary_policy(cudf::io::dictionary_policy::ALWAYS) - .stats_level(cudf::io::statistics_freq::STATISTICS_COLUMN); - cudf::io::write_parquet(options); - } - - auto const filter_ref = cudf::ast::column_name_reference("filter"); - auto filter_expression = cudf::ast::operation(cudf::ast::ast_operator::IDENTITY, filter_ref); - auto source_info = build_source_info(parquet_buffers); - auto const stream = cudf::get_default_stream(); - auto const mr = cudf::get_current_device_resource_ref(); - auto expected_options = - cudf::io::parquet_reader_options::builder(source_info).filter(filter_expression).build(); - auto expected = cudf::io::read_parquet(expected_options, stream, mr); - - auto const [filter_result, payload_result] = - page_level_chunked_hybrid_scan_multifile(source_info, filter_expression, {}, true, stream, mr); - - CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected.tbl->select({0}), filter_result->view()); - CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected.tbl->select({1}), payload_result->view()); -} - -TEST_F(HybridScanMultifileTest, PageLevelAsymmetricSourceRowGroupOrdering) -{ - auto constexpr rows_per_group = 2 * page_size_for_ordered_tests; - auto constexpr rows_source_0 = num_ordered_rows; - auto constexpr rows_source_1 = num_ordered_rows; - auto constexpr rows_per_page = rows_per_group / 4; - - auto source_0_filter_values = cudf::detail::make_counting_transform_iterator( - 0, [](auto i) { return static_cast(i / 100); }); - auto source_1_filter_values = cudf::detail::make_counting_transform_iterator( - 0, [](auto i) { return static_cast((num_ordered_rows - i) / 100); }); - auto source_0_filter = cudf::test::fixed_width_column_wrapper( - source_0_filter_values, source_0_filter_values + rows_source_0); - auto source_1_filter = cudf::test::fixed_width_column_wrapper( - source_1_filter_values, source_1_filter_values + rows_source_1); - - auto source_0_payload_values = std::vector(rows_source_0); - auto source_1_payload_values = std::vector(rows_source_1); - for (auto i = std::size_t{0}; i < source_0_payload_values.size(); ++i) { - source_0_payload_values[i] = "source 0 dictionary value " + std::to_string(i % 8); - } - for (auto i = std::size_t{0}; i < source_1_payload_values.size(); ++i) { - source_1_payload_values[i] = "source 1 dictionary value " + std::to_string(i % 8); - } - auto source_0_payload = cudf::test::strings_column_wrapper(source_0_payload_values.begin(), - source_0_payload_values.end()); - auto source_1_payload = cudf::test::strings_column_wrapper(source_1_payload_values.begin(), - source_1_payload_values.end()); - auto const source_0_table = cudf::table_view{{source_0_filter, source_0_payload}}; - auto const source_1_table = cudf::table_view{{source_1_filter, source_1_payload}}; - - auto parquet_buffers = std::vector>(2); - auto const write_source = [&](auto const& table, auto& buffer) { - cudf::io::table_input_metadata metadata(table); - metadata.column_metadata[0].set_name("col0"); - auto options = cudf::io::parquet_writer_options::builder(cudf::io::sink_info{&buffer}, table) - .metadata(metadata) - .row_group_size_rows(rows_per_group) - .max_page_size_rows(rows_per_page) - .dictionary_policy(cudf::io::dictionary_policy::ALWAYS) - .stats_level(cudf::io::statistics_freq::STATISTICS_COLUMN); - cudf::io::write_parquet(options); - }; - write_source(source_0_table, parquet_buffers[0]); - write_source(source_1_table, parquet_buffers[1]); - - auto constexpr threshold = uint32_t{75}; - auto scalar = cudf::numeric_scalar(threshold); - auto literal = cudf::ast::literal(scalar); - auto col_ref = cudf::ast::column_name_reference("col0"); - auto filter_expression = - cudf::ast::operation(cudf::ast::ast_operator::GREATER_EQUAL, col_ref, literal); - - auto const stream = cudf::get_default_stream(); - auto const mr = cudf::get_current_device_resource_ref(); - auto const source_info = build_source_info(parquet_buffers); - auto const expected = cudf::io::read_parquet( - cudf::io::parquet_reader_options::builder(source_info).filter(filter_expression), stream, mr); - - auto const [filter_table, payload_table] = - page_level_chunked_hybrid_scan_multifile(source_info, filter_expression, {}, true, stream, mr); - - CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected.tbl->select({0}), filter_table->view()); - CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected.tbl->select({1}), payload_table->view()); - auto const [requested_payload_bytes, full_payload_bytes] = - payload_byte_range_sizes(source_info, filter_expression, true, stream, mr); - EXPECT_GT(requested_payload_bytes, 0); - EXPECT_LT(requested_payload_bytes, full_payload_bytes); -} - -TEST_F(HybridScanMultifileTest, PageLevelPlainEncodingExactCoalescedRanges) -{ - auto parquet_buffers = make_plain_payload_parquet_buffers(); - auto const source_info = build_source_info(parquet_buffers); - auto inputs = multifile_inputs(source_info); - - auto scalar = cudf::numeric_scalar(0); - auto literal = cudf::ast::literal(scalar); - auto col_ref_0 = cudf::ast::column_name_reference("col0"); - auto filter_expression = - cudf::ast::operation(cudf::ast::ast_operator::GREATER_EQUAL, col_ref_0, literal); - auto options = cudf::io::parquet_reader_options::builder() - .column_names({"col1"}) - .filter(filter_expression) - .build(); - auto reader = - cudf::io::parquet::experimental::hybrid_scan_multifile{inputs.footer_byte_spans, options}; - setup_page_indexes(reader, inputs); - - auto const row_groups = reader.all_row_groups(options); - auto const metadatas = reader.parquet_metadatas(); - auto selected_rows = - std::vector(reader.total_rows_in_row_groups(row_groups), uint8_t{0}); - auto expected_ranges = - std::vector>(metadatas.size()); - - std::size_t source_row_offset = 0; - for (std::size_t source_idx = 0; source_idx < metadatas.size(); ++source_idx) { - auto const& metadata = metadatas[source_idx]; - ASSERT_EQ(metadata.row_groups.size(), 1); - auto const& payload_chunk = metadata.row_groups.front().columns[1]; - EXPECT_NE(std::find(payload_chunk.meta_data.encodings.begin(), - payload_chunk.meta_data.encodings.end(), - cudf::io::parquet::Encoding::PLAIN), - payload_chunk.meta_data.encodings.end()); - EXPECT_EQ(std::find(payload_chunk.meta_data.encodings.begin(), - payload_chunk.meta_data.encodings.end(), - cudf::io::parquet::Encoding::RLE_DICTIONARY), - payload_chunk.meta_data.encodings.end()); - EXPECT_EQ(payload_chunk.meta_data.dictionary_page_offset, 0); - ASSERT_TRUE(payload_chunk.offset_index.has_value()); - - auto const& pages = payload_chunk.offset_index->page_locations; - ASSERT_GE(pages.size(), 4); - ASSERT_EQ(pages[1].offset + pages[1].compressed_page_size, pages[2].offset); - auto const selected_begin = pages[1].first_row_index; - auto const selected_end = pages[3].first_row_index; - ASSERT_GE(selected_begin, 0); - ASSERT_LE(selected_end, metadata.row_groups.front().num_rows); - std::fill(selected_rows.begin() + source_row_offset + selected_begin, - selected_rows.begin() + source_row_offset + selected_end, - uint8_t{1}); - - expected_ranges[source_idx].emplace_back( - pages[1].offset, pages[2].offset + pages[2].compressed_page_size - pages[1].offset); - source_row_offset += metadata.row_groups.front().num_rows; - } - ASSERT_EQ(source_row_offset, selected_rows.size()); - auto row_mask = - cudf::test::fixed_width_column_wrapper(selected_rows.begin(), selected_rows.end()) - .release(); - - auto const stream = cudf::get_default_stream(); - auto const mr = cudf::get_current_device_resource_ref(); - auto const page_ranges = reader.payload_column_chunks_byte_ranges( - row_groups, row_mask->view(), use_data_page_mask::YES, options, stream); - expect_byte_ranges_equal(expected_ranges, page_ranges); - - auto page_data = fetch_multisource_device_data(inputs, page_ranges, stream, mr); - reader.setup_chunking_for_payload_columns(256 * 1024, - 1024 * 1024, - row_groups, - row_mask->view(), - use_data_page_mask::YES, - page_data.per_source_spans, - options, - stream, - mr); - auto payload_chunks = std::vector>{}; - while (reader.has_next_table_chunk()) { - payload_chunks.push_back( - std::move(reader.materialize_payload_columns_chunk(row_mask->view()).tbl)); - } - auto actual = concatenate_tables(std::move(payload_chunks), stream, mr); - - auto const full = - cudf::io::read_parquet(cudf::io::parquet_reader_options::builder(source_info), stream, mr); - auto const expected = cudf::apply_boolean_mask(full.tbl->select({1}), row_mask->view()); - CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected->view(), actual->view()); -} - -TEST_F(HybridScanMultifileTest, PageLevelAllFalseMaskHasNoRanges) +TEST_F(HybridScanMultifileTest, PageLevelAllFalseMaskHasEmptyRangeSlots) { auto parquet_buffers = make_plain_payload_parquet_buffers(); auto const source_info = build_source_info(parquet_buffers); @@ -535,23 +229,21 @@ TEST_F(HybridScanMultifileTest, PageLevelAllFalseMaskHasNoRanges) auto const stream = cudf::get_default_stream(); auto const mr = cudf::get_current_device_resource_ref(); - auto const page_ranges = reader.payload_column_chunks_byte_ranges( - row_groups, row_mask->view(), use_data_page_mask::YES, options, stream); - ASSERT_EQ(page_ranges.size(), parquet_buffers.size()); - EXPECT_TRUE(std::all_of( - page_ranges.begin(), page_ranges.end(), [](auto const& ranges) { return ranges.empty(); })); - - auto const empty_page_data = - std::vector>>(parquet_buffers.size()); - reader.setup_chunking_for_payload_columns(0, - 0, - row_groups, - row_mask->view(), - use_data_page_mask::YES, - empty_page_data, - options, - stream, - mr); + auto const page_ranges = + reader.payload_pages_byte_ranges(row_groups, row_mask->view(), options, stream); + ASSERT_FALSE(page_ranges.first.empty()); + ASSERT_EQ(page_ranges.first.size(), page_ranges.second.size()); + EXPECT_TRUE(std::all_of(page_ranges.first.begin(), + page_ranges.first.end(), + [](auto const& range) { return range.is_empty(); })); + + auto page_data = fetch_multisource_device_data(inputs, page_ranges, stream, mr); + ASSERT_EQ(page_data.flat_spans.size(), page_ranges.first.size()); + EXPECT_TRUE(std::all_of(page_data.flat_spans.begin(), + page_data.flat_spans.end(), + [](auto const& span) { return span.empty(); })); + reader.setup_chunking_for_payload_columns( + 0, 0, row_groups, page_data.flat_spans, options, stream, mr); ASSERT_TRUE(reader.has_next_table_chunk()); auto const result = reader.materialize_payload_columns_chunk(row_mask->view()); EXPECT_EQ(result.tbl->num_rows(), 0); @@ -560,120 +252,6 @@ TEST_F(HybridScanMultifileTest, PageLevelAllFalseMaskHasNoRanges) EXPECT_FALSE(reader.has_next_table_chunk()); } -TEST_F(HybridScanMultifileTest, PageLevelNoMaskFallbackAndPlanLifecycle) -{ - auto parquet_buffers = make_plain_payload_parquet_buffers(); - auto const source_info = build_source_info(parquet_buffers); - auto inputs = multifile_inputs(source_info); - - auto scalar = cudf::numeric_scalar(0); - auto literal = cudf::ast::literal(scalar); - auto col_ref_0 = cudf::ast::column_name_reference("col0"); - auto filter_expression = - cudf::ast::operation(cudf::ast::ast_operator::GREATER_EQUAL, col_ref_0, literal); - auto options = cudf::io::parquet_reader_options::builder() - .column_names({"col1"}) - .filter(filter_expression) - .build(); - auto reader = - cudf::io::parquet::experimental::hybrid_scan_multifile{inputs.footer_byte_spans, options}; - setup_page_indexes(reader, inputs); - - auto const row_groups = reader.all_row_groups(options); - auto const full_ranges = group_byte_ranges_by_source( - reader.payload_column_chunks_byte_ranges(row_groups, options), parquet_buffers.size()); - auto true_values = cuda::make_constant_iterator(true); - auto row_mask = cudf::test::fixed_width_column_wrapper( - true_values, true_values + reader.total_rows_in_row_groups(row_groups)) - .release(); - auto const stream = cudf::get_default_stream(); - auto const mr = cudf::get_current_device_resource_ref(); - - auto const planned_ranges = reader.payload_column_chunks_byte_ranges( - row_groups, row_mask->view(), use_data_page_mask::NO, options, stream); - expect_byte_ranges_equal(full_ranges, planned_ranges); - EXPECT_THROW(static_cast(reader.payload_column_chunks_byte_ranges( - row_groups, row_mask->view(), use_data_page_mask::NO, options, stream)), - cudf::logic_error); - - auto page_data = fetch_multisource_device_data(inputs, planned_ranges, stream, mr); - reader.setup_chunking_for_payload_columns(0, - 0, - row_groups, - row_mask->view(), - use_data_page_mask::NO, - page_data.per_source_spans, - options, - stream, - mr); - EXPECT_THROW(reader.setup_chunking_for_payload_columns(0, - 0, - row_groups, - row_mask->view(), - use_data_page_mask::NO, - page_data.per_source_spans, - options, - stream, - mr), - cudf::logic_error); -} - -TEST_F(HybridScanMultifileTest, PageLevelRejectsInvalidFetchedSpans) -{ - auto parquet_buffers = make_plain_payload_parquet_buffers(); - auto const source_info = build_source_info(parquet_buffers); - auto inputs = multifile_inputs(source_info); - - auto scalar = cudf::numeric_scalar(0); - auto literal = cudf::ast::literal(scalar); - auto col_ref_0 = cudf::ast::column_name_reference("col0"); - auto filter_expression = - cudf::ast::operation(cudf::ast::ast_operator::GREATER_EQUAL, col_ref_0, literal); - auto options = cudf::io::parquet_reader_options::builder() - .column_names({"col1"}) - .filter(filter_expression) - .build(); - auto reader = - cudf::io::parquet::experimental::hybrid_scan_multifile{inputs.footer_byte_spans, options}; - setup_page_indexes(reader, inputs); - - auto const row_groups = reader.all_row_groups(options); - auto selected_values = cudf::detail::make_counting_transform_iterator( - cudf::size_type{0}, [](auto i) { return (i % num_ordered_rows) < num_ordered_rows / 2; }); - auto row_mask = cudf::test::fixed_width_column_wrapper( - selected_values, selected_values + reader.total_rows_in_row_groups(row_groups)) - .release(); - auto const stream = cudf::get_default_stream(); - auto const mr = cudf::get_current_device_resource_ref(); - - auto page_ranges = reader.payload_column_chunks_byte_ranges( - row_groups, row_mask->view(), use_data_page_mask::YES, options, stream); - auto page_data = fetch_multisource_device_data(inputs, page_ranges, stream, mr); - auto bad_count = page_data.per_source_spans; - auto count_source = std::find_if( - bad_count.begin(), bad_count.end(), [](auto const& spans) { return not spans.empty(); }); - ASSERT_NE(count_source, bad_count.end()); - count_source->pop_back(); - EXPECT_THROW( - reader.setup_chunking_for_payload_columns( - 0, 0, row_groups, row_mask->view(), use_data_page_mask::YES, bad_count, options, stream, mr), - cudf::logic_error); - - page_ranges = reader.payload_column_chunks_byte_ranges( - row_groups, row_mask->view(), use_data_page_mask::YES, options, stream); - auto bad_size = page_data.per_source_spans; - auto size_source = std::find_if( - bad_size.begin(), bad_size.end(), [](auto const& spans) { return not spans.empty(); }); - ASSERT_NE(size_source, bad_size.end()); - ASSERT_GT(size_source->front().size(), 1); - size_source->front() = - cudf::device_span{size_source->front().data(), size_source->front().size() - 1}; - EXPECT_THROW( - reader.setup_chunking_for_payload_columns( - 0, 0, row_groups, row_mask->view(), use_data_page_mask::YES, bad_size, options, stream, mr), - cudf::logic_error); -} - TEST_F(HybridScanMultifileTest, PrependIndexColumns) { using T = int32_t; @@ -823,3 +401,79 @@ TEST_F(HybridScanMultifileTest, MaterializeStructs) test_hybrid_scan_multifile({col0, *col1, *col2}); } + +TEST_F(HybridScanMultifileTest, SparseDictionaryEncodedPages) +{ + auto constexpr num_sources = 2; + auto const [table, parquet_buffer] = create_parquet_with_stats(); + auto const payload_table = table->view().select({2}); + auto parquet_buffers = std::vector>(num_sources, parquet_buffer); + + auto const source_info = build_source_info(parquet_buffers); + auto const reader_options = + cudf::io::parquet_reader_options::builder().column_names({"col2"}).build(); + auto inputs = multifile_inputs(source_info); + auto reader = cudf::io::parquet::experimental::hybrid_scan_multifile{inputs.footer_byte_spans, + reader_options}; + setup_page_indexes(reader, inputs); + + auto const row_groups = reader.all_row_groups(reader_options); + auto row_mask_values = cudf::detail::make_counting_transform_iterator( + cudf::size_type{0}, [](auto i) { return (i / page_size_for_ordered_tests) % 2 == 0; }); + auto row_mask = cudf::test::fixed_width_column_wrapper( + row_mask_values, row_mask_values + reader.total_rows_in_row_groups(row_groups)); + auto const stream = cudf::get_default_stream(); + auto const mr = cudf::get_current_device_resource_ref(); + + auto const page_ranges = + reader.payload_pages_byte_ranges(row_groups, row_mask, reader_options, stream); + auto page_data = fetch_multisource_device_data(inputs, page_ranges, stream, mr); + reader.setup_chunking_for_payload_columns( + 0, 0, row_groups, page_data.flat_spans, reader_options, stream, mr); + + ASSERT_TRUE(reader.has_next_table_chunk()); + auto const result = reader.materialize_payload_columns_chunk(row_mask); + EXPECT_FALSE(reader.has_next_table_chunk()); + + auto const input = + cudf::concatenate(std::vector(num_sources, payload_table), stream, mr); + auto const expected = cudf::apply_boolean_mask(input->view(), row_mask, stream, mr); + CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected->view(), result.tbl->view()); +} + +TEST_F(HybridScanMultifileTest, SparsePayloadWithAsymmetricRowGroupOrdering) +{ + using T = uint32_t; + + auto parquet_buffers = std::vector>{}; + parquet_buffers.emplace_back(std::get<1>(create_parquet_with_stats())); + // Name the descending column `col0` so this source retains earlier row groups. + parquet_buffers.emplace_back(std::get<1>(create_parquet_with_stats( + 100, cudf::io::compression_type::AUTO, {"col0", "col1", "col2"}, {1, 0, 2}))); + + auto const source_info = build_source_info(parquet_buffers); + auto const stream = cudf::get_default_stream(); + auto const mr = cudf::get_current_device_resource_ref(); + auto scalar = cudf::numeric_scalar(75); + auto literal = cudf::ast::literal(scalar); + auto col_ref = cudf::ast::column_name_reference("col0"); + auto filter_expression = + cudf::ast::operation(cudf::ast::ast_operator::GREATER_EQUAL, col_ref, literal); + auto const options = + cudf::io::parquet_reader_options::builder().filter(filter_expression).build(); + + auto inputs = multifile_inputs(source_info); + auto reader = + cudf::io::parquet::experimental::hybrid_scan_multifile{inputs.footer_byte_spans, options}; + auto const row_groups = + reader.filter_row_groups_with_stats(reader.all_row_groups(options), options, stream); + EXPECT_EQ(row_groups, (std::vector>{{1, 2, 3}, {0, 1, 2}})); + + auto const expected = cudf::io::read_parquet( + cudf::io::parquet_reader_options::builder(source_info).filter(filter_expression), stream, mr); + auto const [filter_table, payload_table] = + chunked_sparse_hybrid_scan_multifile(source_info, filter_expression, {}, true, stream, mr); + + CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected.tbl->select({0}), filter_table->view()); + CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected.tbl->select({1, 2}), payload_table->view()); +} \ No newline at end of file From d277ca905d2b24a15a7ddeaf1ee2f0f23f81fcf1 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Mon, 3 Aug 2026 17:18:00 +0000 Subject: [PATCH 04/10] More simplifications --- .../io/experimental/hybrid_scan_multifile.hpp | 6 +- .../experimental/hybrid_scan_chunking.cu | 28 ++- .../parquet/experimental/hybrid_scan_impl.cpp | 27 ++- .../parquet/experimental/hybrid_scan_impl.hpp | 17 +- .../experimental/hybrid_scan_multifile.cpp | 3 +- .../experimental/hybrid_scan_preprocess.cu | 2 + .../hybrid_scan_multifile_composer.cpp | 1 + .../hybrid_scan_multifile_test.cpp | 169 +++++++++--------- 8 files changed, 138 insertions(+), 115 deletions(-) diff --git a/cpp/include/cudf/io/experimental/hybrid_scan_multifile.hpp b/cpp/include/cudf/io/experimental/hybrid_scan_multifile.hpp index ad814b732466..f66cbdf93fde 100644 --- a/cpp/include/cudf/io/experimental/hybrid_scan_multifile.hpp +++ b/cpp/include/cudf/io/experimental/hybrid_scan_multifile.hpp @@ -267,8 +267,7 @@ class hybrid_scan_multifile { * source map has one source index per byte range. Dictionary pages precede data pages within each * column chunk. Pruned pages are represented by empty byte ranges. * - * @throws std::invalid_argument if any selected column chunk does not have a valid offset - * index + * @throws cudf::logic_error if any selected column chunk does not have a valid offset index * * @param row_group_indices Input row group indices, one vector per source * @param row_mask Boolean mask spanning the selected row groups @@ -414,6 +413,8 @@ class hybrid_scan_multifile { * @param chunk_read_limit Maximum bytes returned per output table chunk, or zero * @param pass_read_limit Maximum read/decompression memory, or zero * @param row_group_indices Input row group indices, one vector per source + * @param row_mask Boolean column spanning all selected rows across all sources and indicating + * which rows need to be read * @param page_data Flattened device spans of payload page data in the same order as the byte * ranges from `payload_pages_byte_ranges` * @param options Parquet reader options @@ -424,6 +425,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::host_span const> page_data, parquet_reader_options const& options, rmm::cuda_stream_view stream, diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu b/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu index 7c439c984e09..b142e1aff241 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu +++ b/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu @@ -155,21 +155,19 @@ void hybrid_scan_reader_impl::setup_next_pass( // store off how much memory we've used so far. This includes the compressed page data and the // decompressed dictionary data. we will subtract this from the available total memory for the // subpasses - auto const compressed_data_size = - _sparse_page_io - ? std::accumulate(column_chunk_data.begin(), - column_chunk_data.end(), - std::size_t{0}, - [](auto size, auto const& page) { return size + page.size(); }) - : [&] { - auto chunk_iter = thrust::make_transform_iterator( - pass.chunks.d_begin(), parquet::detail::get_chunk_compressed_size{}); - return cudf::detail::reduce(chunk_iter, - chunk_iter + pass.chunks.size(), - size_t{0}, - cuda::std::plus{}, - _stream); - }(); + auto const compressed_data_size = [&] { + // In Sparse I/O case, compressed chunk size is the sum of its page data span sizes + if (_sparse_page_io) { + return std::accumulate(column_chunk_data.begin(), + column_chunk_data.end(), + std::size_t{0}, + [](auto size, auto const& page) { return size + page.size(); }); + } + auto chunk_iter = thrust::make_transform_iterator( + pass.chunks.d_begin(), parquet::detail::get_chunk_compressed_size{}); + return cudf::detail::reduce( + chunk_iter, chunk_iter + pass.chunks.size(), size_t{0}, cuda::std::plus{}, _stream); + }(); pass.base_mem_size = decomp_dict_data_size + compressed_data_size; // if we are doing subpass reading, generate more accurate num_row estimates for list columns. diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp b/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp index 4d8393518a3e..ed2752674274 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp +++ b/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp @@ -94,6 +94,12 @@ namespace { [](auto sum, auto const& rgs) { return sum + static_cast(rgs.size()); }); } +/** + * @brief Get the byte range of a column chunk's dictionary page, if present + * + * @param column Column chunk metadata with a valid offset index + * @return Dictionary page offset and size, or `std::nullopt` when no dictionary page is present + */ [[nodiscard]] std::optional> dictionary_page_range( ColumnChunk const& column) { @@ -102,7 +108,8 @@ namespace { auto const offset = column.meta_data.dictionary_page_offset; return std::pair{offset, column.meta_data.data_page_offset - offset}; } - if (column.meta_data.data_page_offset < page_locations.front().offset) { + if (not page_locations.empty() and + column.meta_data.data_page_offset < page_locations.front().offset) { auto const offset = column.meta_data.data_page_offset; return std::pair{offset, page_locations.front().offset - offset}; } @@ -838,11 +845,17 @@ void hybrid_scan_reader_impl::setup_chunking_for_payload_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, std::span const> page_data, parquet_reader_options const& options, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) { + CUDF_EXPECTS(std::cmp_equal(row_mask.size(), total_rows_in_row_groups(row_group_indices)), + "Row mask must span across all input row groups"); + CUDF_EXPECTS(row_mask.null_count() == 0, + "Row mask must not have any nulls when materializing payload column"); + reset_column_selection(); prepare_materialization( read_columns_mode::PAYLOAD_COLUMNS, row_group_indices.size(), options, stream, mr); @@ -850,12 +863,12 @@ void hybrid_scan_reader_impl::setup_chunking_for_payload_columns( _input_pass_read_limit = pass_read_limit; _output_chunk_read_limit = chunk_read_limit; - // Return early if all payload pages are pruned. - if (std::all_of( - page_data.begin(), page_data.end(), [](auto const& page) { return page.empty(); })) { + // Return early if all rows are pruned + if (are_all_rows_pruned(row_mask, stream)) { auto const empty_row_groups = std::vector>(row_group_indices.size(), std::vector{}); prepare_data(read_mode::CHUNKED_READ, empty_row_groups, {}, {}); + // Set correct number of input row groups to the output metadata _file_itm_data.num_input_row_groups = count_row_groups(row_group_indices); return; } @@ -1416,11 +1429,11 @@ void hybrid_scan_reader_impl::set_sparse_pass_page_mask( _pass_page_mask = cudf::detail::make_empty_host_vector(pass->pages.size(), _stream); // Find the first logical page-data span for every column chunk. - auto page_offsets = std::vector(chunks.size()); - auto chunk_idx = std::size_t{0}; + auto page_offsets = std::vector{}; + page_offsets.reserve(chunks.size()); auto const num_logical_pages = std::accumulate( chunks.begin(), chunks.end(), std::size_t{0}, [&](auto offset, auto const& chunk) { - page_offsets[chunk_idx++] = offset; + page_offsets.push_back(offset); return offset + chunk.num_dict_pages + chunk.num_data_pages; }); CUDF_EXPECTS(page_data.size() == num_logical_pages, diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_impl.hpp b/cpp/src/io/parquet/experimental/hybrid_scan_impl.hpp index 3f3307e49417..3e8604e71cce 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_impl.hpp +++ b/cpp/src/io/parquet/experimental/hybrid_scan_impl.hpp @@ -194,6 +194,9 @@ class hybrid_scan_reader_impl : public parquet::detail::reader_impl { payload_column_chunks_byte_ranges(std::span const> row_group_indices, parquet_reader_options const& options); + /** + * @copydoc cudf::io::parquet::experimental::hybrid_scan_multifile::payload_pages_byte_ranges + */ [[nodiscard]] std::pair, std::vector> payload_pages_byte_ranges(std::span const> row_group_indices, cudf::column_view const& row_mask, @@ -266,10 +269,15 @@ class hybrid_scan_reader_impl : public parquet::detail::reader_impl { rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr); + /** + * @copydoc + * cudf::io::parquet::experimental::hybrid_scan_multifile::setup_chunking_for_payload_columns + */ void setup_chunking_for_payload_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, std::span const> page_data, parquet_reader_options const& options, rmm::cuda_stream_view stream, @@ -424,7 +432,8 @@ class hybrid_scan_reader_impl : public parquet::detail::reader_impl { * * @param mode Value indicating if the data sources are read all at once or chunk by chunk * @param row_group_indices Row group indices to read - * @param column_chunk_data Device spans of buffers containing column chunk data + * @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 */ void prepare_data(read_mode mode, @@ -467,7 +476,8 @@ class hybrid_scan_reader_impl : public parquet::detail::reader_impl { * @brief Ratchet the pass/subpass/chunk process forward. * * @param mode Value indicating if the data sources are read all at once or chunk by chunk - * @param column_chunk_data Device spans of buffers containing column chunk data + * @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 */ void handle_chunking(read_mode mode, @@ -480,7 +490,8 @@ class hybrid_scan_reader_impl : public parquet::detail::reader_impl { * A 'pass' is defined as a subset of row groups read out of the globally * requested set of all row groups. * - * @param column_chunk_data Device spans of buffers containing column chunk data + * @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 */ void setup_next_pass(std::span const> column_chunk_data, diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_multifile.cpp b/cpp/src/io/parquet/experimental/hybrid_scan_multifile.cpp index ea6f4bf8204b..524721dd1cc4 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_multifile.cpp +++ b/cpp/src/io/parquet/experimental/hybrid_scan_multifile.cpp @@ -239,6 +239,7 @@ void hybrid_scan_multifile::setup_chunking_for_payload_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::host_span const> page_data, parquet_reader_options const& options, rmm::cuda_stream_view stream, @@ -246,7 +247,7 @@ void hybrid_scan_multifile::setup_chunking_for_payload_columns( { CUDF_FUNC_RANGE(); _impl->setup_chunking_for_payload_columns( - chunk_read_limit, pass_read_limit, row_group_indices, page_data, options, stream, mr); + chunk_read_limit, pass_read_limit, row_group_indices, row_mask, page_data, options, stream, mr); } table_with_metadata hybrid_scan_multifile::materialize_payload_columns_chunk( diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu b/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu index 300fc69f1926..820ca68f33ce 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu +++ b/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu @@ -198,6 +198,8 @@ void hybrid_scan_reader_impl::setup_sparse_compressed_data( { auto& pass = *_pass_itm_data; + // This function should never be called if `num_rows == 0`. + CUDF_EXPECTS(_pass_itm_data->num_rows > 0, "Number of reading rows must not be zero."); CUDF_EXPECTS(_has_offset_index, "Sparse page I/O requires complete offset indexes"); auto& chunks = pass.chunks; diff --git a/cpp/tests/io/experimental/hybrid_scan_multifile_composer.cpp b/cpp/tests/io/experimental/hybrid_scan_multifile_composer.cpp index 9832ec782f2e..003ebfbc3d80 100644 --- a/cpp/tests/io/experimental/hybrid_scan_multifile_composer.cpp +++ b/cpp/tests/io/experimental/hybrid_scan_multifile_composer.cpp @@ -211,6 +211,7 @@ chunked_sparse_hybrid_scan_multifile( reader.setup_chunking_for_payload_columns(chunk_read_limit, pass_read_limit, row_groups, + row_mask->view(), payload_page_data.flat_spans, options, stream, diff --git a/cpp/tests/io/experimental/hybrid_scan_multifile_test.cpp b/cpp/tests/io/experimental/hybrid_scan_multifile_test.cpp index 8652f71878f5..47c8164468b6 100644 --- a/cpp/tests/io/experimental/hybrid_scan_multifile_test.cpp +++ b/cpp/tests/io/experimental/hybrid_scan_multifile_test.cpp @@ -12,17 +12,16 @@ #include #include +#include #include #include #include -#include #include #include #include #include #include #include -#include #include @@ -37,39 +36,6 @@ namespace { using cudf::io::parquet::experimental::use_data_page_mask; -std::vector> make_plain_payload_parquet_buffers() -{ - auto constexpr num_sources = 2; - auto parquet_buffers = std::vector>(num_sources); - for (auto source_idx = 0; source_idx < num_sources; ++source_idx) { - auto filter_values = cuda::counting_iterator{0}; - auto payload_values = - cudf::detail::make_counting_transform_iterator(cudf::size_type{0}, [source_idx](auto i) { - return static_cast(i) + source_idx * int64_t{num_ordered_rows}; - }); - auto filter = cudf::test::fixed_width_column_wrapper( - filter_values, filter_values + num_ordered_rows); - auto payload = cudf::test::fixed_width_column_wrapper( - payload_values, payload_values + num_ordered_rows); - auto const table = cudf::table_view{{filter, payload}}; - - cudf::io::table_input_metadata metadata(table); - metadata.column_metadata[0].set_name("col0"); - metadata.column_metadata[1].set_name("col1").set_encoding(cudf::io::column_encoding::PLAIN); - auto options = cudf::io::parquet_writer_options::builder( - cudf::io::sink_info{&parquet_buffers[source_idx]}, table) - .metadata(metadata) - .row_group_size_rows(num_ordered_rows) - .max_page_size_rows(page_size_for_ordered_tests) - .max_page_size_bytes(64 * 1024 * 1024) - .compression(cudf::io::compression_type::NONE) - .dictionary_policy(cudf::io::dictionary_policy::NEVER) - .stats_level(cudf::io::statistics_freq::STATISTICS_COLUMN); - cudf::io::write_parquet(options); - } - return parquet_buffers; -} - /** * @brief Helper to test multifile hybrid scan single-shot materialization * @@ -202,56 +168,6 @@ TEST_F(HybridScanMultifileTest, MaterializeListsOfStrings) test_hybrid_scan_multifile({col0, *col1, *col2, *col3, *col4}, false); } -TEST_F(HybridScanMultifileTest, PageLevelAllFalseMaskHasEmptyRangeSlots) -{ - auto parquet_buffers = make_plain_payload_parquet_buffers(); - auto const source_info = build_source_info(parquet_buffers); - auto inputs = multifile_inputs(source_info); - - auto scalar = cudf::numeric_scalar(0); - auto literal = cudf::ast::literal(scalar); - auto col_ref_0 = cudf::ast::column_name_reference("col0"); - auto filter_expression = - cudf::ast::operation(cudf::ast::ast_operator::GREATER_EQUAL, col_ref_0, literal); - auto options = cudf::io::parquet_reader_options::builder() - .column_names({"col1"}) - .filter(filter_expression) - .build(); - auto reader = - cudf::io::parquet::experimental::hybrid_scan_multifile{inputs.footer_byte_spans, options}; - setup_page_indexes(reader, inputs); - - auto const row_groups = reader.all_row_groups(options); - auto false_values = cuda::make_constant_iterator(false); - auto row_mask = cudf::test::fixed_width_column_wrapper( - false_values, false_values + reader.total_rows_in_row_groups(row_groups)) - .release(); - auto const stream = cudf::get_default_stream(); - auto const mr = cudf::get_current_device_resource_ref(); - - auto const page_ranges = - reader.payload_pages_byte_ranges(row_groups, row_mask->view(), options, stream); - ASSERT_FALSE(page_ranges.first.empty()); - ASSERT_EQ(page_ranges.first.size(), page_ranges.second.size()); - EXPECT_TRUE(std::all_of(page_ranges.first.begin(), - page_ranges.first.end(), - [](auto const& range) { return range.is_empty(); })); - - auto page_data = fetch_multisource_device_data(inputs, page_ranges, stream, mr); - ASSERT_EQ(page_data.flat_spans.size(), page_ranges.first.size()); - EXPECT_TRUE(std::all_of(page_data.flat_spans.begin(), - page_data.flat_spans.end(), - [](auto const& span) { return span.empty(); })); - reader.setup_chunking_for_payload_columns( - 0, 0, row_groups, page_data.flat_spans, options, stream, mr); - ASSERT_TRUE(reader.has_next_table_chunk()); - auto const result = reader.materialize_payload_columns_chunk(row_mask->view()); - EXPECT_EQ(result.tbl->num_rows(), 0); - EXPECT_EQ(result.tbl->num_columns(), 1); - EXPECT_EQ(result.metadata.num_input_row_groups, 2); - EXPECT_FALSE(reader.has_next_table_chunk()); -} - TEST_F(HybridScanMultifileTest, PrependIndexColumns) { using T = int32_t; @@ -429,7 +345,7 @@ TEST_F(HybridScanMultifileTest, SparseDictionaryEncodedPages) reader.payload_pages_byte_ranges(row_groups, row_mask, reader_options, stream); auto page_data = fetch_multisource_device_data(inputs, page_ranges, stream, mr); reader.setup_chunking_for_payload_columns( - 0, 0, row_groups, page_data.flat_spans, reader_options, stream, mr); + 0, 0, row_groups, row_mask, page_data.flat_spans, reader_options, stream, mr); ASSERT_TRUE(reader.has_next_table_chunk()); auto const result = reader.materialize_payload_columns_chunk(row_mask); @@ -476,4 +392,83 @@ TEST_F(HybridScanMultifileTest, SparsePayloadWithAsymmetricRowGroupOrdering) CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected.tbl->select({0}), filter_table->view()); CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected.tbl->select({1, 2}), payload_table->view()); -} \ No newline at end of file +} + +TEST_F(HybridScanMultifileTest, SparsePayloadEmptyAndAllPrunedPageData) +{ + using T = uint32_t; + + // Create two sources with page indexes for sparse payload materialization. + auto file_buffers = std::vector>{}; + file_buffers.emplace_back(std::get<1>(create_parquet_with_stats())); + file_buffers.emplace_back(std::get<1>(create_parquet_with_stats())); + auto inputs = multifile_inputs(build_source_info(file_buffers)); + + auto const options = cudf::io::parquet_reader_options::builder().column_names({"col1"}).build(); + auto const stream = cudf::get_default_stream(); + auto const mr = cudf::get_current_device_resource_ref(); + + // Empty row-group selection accepts an empty outer page-data span. + { + auto reader = std::make_unique( + inputs.footer_byte_spans, options); + auto const row_groups = + std::vector>(inputs.footer_byte_spans.size()); + auto const empty_page_data = std::vector>{}; + auto false_scalar = cudf::numeric_scalar{false}; + auto row_mask = cudf::make_column_from_scalar(false_scalar, 0); + + EXPECT_NO_THROW(reader->setup_chunking_for_payload_columns( + 0, 0, row_groups, row_mask->view(), empty_page_data, options, stream, mr)); + ASSERT_TRUE(reader->has_next_table_chunk()); + auto const result = reader->materialize_payload_columns_chunk(row_mask->view()); + EXPECT_EQ(result.tbl->num_rows(), 0); + EXPECT_EQ(result.metadata.num_input_row_groups, 0); + EXPECT_FALSE(reader->has_next_table_chunk()); + } + + // An empty outer span is invalid when the selected row groups contain pages. + { + auto reader = std::make_unique( + inputs.footer_byte_spans, options); + setup_page_indexes(*reader, inputs); + + auto const row_groups = reader->all_row_groups(options); + auto const row_mask = reader->build_all_true_row_mask(row_groups, stream, mr); + auto const empty_page_data = std::vector>{}; + + EXPECT_THROW(reader->setup_chunking_for_payload_columns( + 0, 0, row_groups, row_mask->view(), empty_page_data, options, stream, mr), + cudf::logic_error); + } + + // An all-false row mask prunes every page and yields an empty payload table. + { + auto reader = std::make_unique( + inputs.footer_byte_spans, options); + setup_page_indexes(*reader, inputs); + + auto const row_groups = reader->all_row_groups(options); + auto false_scalar = cudf::numeric_scalar{false}; + auto row_mask = + cudf::make_column_from_scalar(false_scalar, reader->total_rows_in_row_groups(row_groups)); + auto const page_ranges = + reader->payload_pages_byte_ranges(row_groups, row_mask->view(), options, stream); + + ASSERT_FALSE(page_ranges.first.empty()); + EXPECT_TRUE(std::all_of(page_ranges.first.begin(), + page_ranges.first.end(), + [](auto const& range) { return range.is_empty(); })); + auto const all_pruned_page_data = + std::vector>(page_ranges.first.size()); + reader->setup_chunking_for_payload_columns( + 0, 0, row_groups, row_mask->view(), all_pruned_page_data, options, stream, mr); + ASSERT_TRUE(reader->has_next_table_chunk()); + auto const result = reader->materialize_payload_columns_chunk(row_mask->view()); + EXPECT_EQ(result.tbl->num_rows(), 0); + EXPECT_EQ(result.tbl->num_columns(), 1); + // Two sources with four row groups each + EXPECT_EQ(result.metadata.num_input_row_groups, 8); + EXPECT_FALSE(reader->has_next_table_chunk()); + } +} From 7db024fadf20e3610d15c10f9fba5df2d863b145 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Mon, 3 Aug 2026 19:18:23 +0000 Subject: [PATCH 05/10] Simplify further and add comments --- .../parquet/experimental/hybrid_scan_impl.cpp | 76 ++++++++++--------- 1 file changed, 42 insertions(+), 34 deletions(-) diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp b/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp index ed2752674274..3c5460488d71 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp +++ b/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp @@ -534,21 +534,23 @@ hybrid_scan_reader_impl::payload_pages_byte_ranges( _input_columns.end(), std::back_inserter(column_schemas), [](auto const& col) { return col.schema_idx; }); - CUDF_EXPECTS(_extended_metadata->page_index_presence(row_group_indices, column_schemas).second, "Page-level I/O for payload columns requires offset indexes to be present"); - auto data_page_mask = _extended_metadata->compute_data_page_mask( - row_mask, row_group_indices, _input_columns, 0, stream); - auto const num_columns = _input_columns.size(); auto const num_chunks = static_cast(count_row_groups(row_group_indices)) * num_columns; - auto chunk_page_counts = std::vector(num_chunks); - auto row_group_ordinal = std::size_t{0}; + + // The data page mask is ordered by column (all pages of a column, then the next column) so + // accumulate page counts per column to locate each column's portion of the mask. + auto mask_offsets = std::vector(num_columns + 1, 0); + + // For each source for (std::size_t source_idx = 0; source_idx < row_group_indices.size(); ++source_idx) { + // For each row group in the source auto colchunk_offsets = std::vector>(num_columns); for (auto const row_group_idx : row_group_indices[source_idx]) { + // For each selected column chunk in the row group auto const& row_group = _extended_metadata->get_row_group(row_group_idx, source_idx); for (std::size_t col_idx = 0; col_idx < num_columns; ++col_idx) { auto const schema_idx = @@ -556,36 +558,37 @@ hybrid_scan_reader_impl::payload_pages_byte_ranges( auto& colchunk_offset = colchunk_offsets[col_idx]; colchunk_offset = parquet::detail::find_colchunk_iter_offset(row_group, schema_idx, colchunk_offset); - chunk_page_counts[row_group_ordinal * num_columns + col_idx] = + // Accumulate page counts per column + mask_offsets[col_idx + 1] += row_group.columns[colchunk_offset.value()].offset_index->page_locations.size(); } - ++row_group_ordinal; } } - auto mask_offsets = std::vector(num_chunks); - std::size_t mask_size{0}; - for (std::size_t col_idx = 0; col_idx < num_columns; ++col_idx) { - for (std::size_t chunk_idx = col_idx; chunk_idx < num_chunks; chunk_idx += num_columns) { - mask_offsets[chunk_idx] = mask_size; - mask_size += chunk_page_counts[chunk_idx]; - } - } + // Accumulate page counts per column + std::partial_sum(mask_offsets.begin(), mask_offsets.end(), mask_offsets.begin()); + + // 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); CUDF_EXPECTS(data_page_mask.empty() or data_page_mask.size() == mask_size, "Computed data page mask does not match offset indexes"); + // Generate page ranges (row group wise) and the corresponding source map auto page_ranges = std::vector{}; auto source_map = std::vector{}; page_ranges.reserve(mask_size + num_chunks); source_map.reserve(mask_size + num_chunks); - row_group_ordinal = 0; + // For each source for (std::size_t source_idx = 0; source_idx < row_group_indices.size(); ++source_idx) { auto colchunk_offsets = std::vector>(num_columns); + // For each row group in the source for (auto const row_group_idx : row_group_indices[source_idx]) { auto const& row_group = _extended_metadata->get_row_group(row_group_idx, source_idx); + // For each selected column chunk in the row group for (std::size_t col_idx = 0; col_idx < num_columns; ++col_idx) { - auto const chunk_idx = row_group_ordinal * num_columns + col_idx; auto const schema_idx = _extended_metadata->map_schema_index(column_schemas[col_idx], source_idx); auto& colchunk_offset = colchunk_offsets[col_idx]; @@ -593,31 +596,35 @@ hybrid_scan_reader_impl::payload_pages_byte_ranges( parquet::detail::find_colchunk_iter_offset(row_group, schema_idx, colchunk_offset); auto const& column_chunk = row_group.columns[colchunk_offset.value()]; auto const& page_locations = column_chunk.offset_index->page_locations; - auto const mask_offset = mask_offsets[chunk_idx]; - auto const any_retained = + auto const mask_offset = mask_offsets[col_idx]; + auto const any_data_page_retained = data_page_mask.empty() or std::any_of(data_page_mask.begin() + mask_offset, data_page_mask.begin() + mask_offset + page_locations.size(), - [](bool retained) { return retained; }); + cuda::std::identity{}); - auto const dictionary_range = dictionary_page_range(column_chunk); - - auto add_range = [&](bool retained_page, int64_t offset, int64_t size) { - page_ranges.emplace_back(offset, retained_page ? size : 0); + // Helper lambda to add a page's byte range and source index to the output vectors + auto add_page_range = [&](bool is_page_retained, int64_t offset, int64_t size) { + page_ranges.emplace_back(offset, is_page_retained ? size : 0); source_map.push_back(static_cast(source_idx)); }; - if (dictionary_range.has_value()) { - add_range(any_retained, dictionary_range->first, dictionary_range->second); + // Add dictionary page range if any of the data pages are also retained + if (auto const dict_page_range = dictionary_page_range(column_chunk); + dict_page_range.has_value()) { + add_page_range(any_data_page_retained, dict_page_range->first, dict_page_range->second); } - for (std::size_t page_idx = 0; page_idx < page_locations.size(); ++page_idx) { - auto const& location = page_locations[page_idx]; - add_range(data_page_mask.empty() or data_page_mask[mask_offset + page_idx], - location.offset, - static_cast(location.compressed_page_size)); + + // Add data page ranges + auto mask_iter = data_page_mask.cbegin() + mask_offset; + for (auto const& location : page_locations) { + add_page_range(data_page_mask.empty() or *mask_iter++, + location.offset, + static_cast(location.compressed_page_size)); } + // Update the mask offset for the next column + mask_offsets[col_idx] += page_locations.size(); } - ++row_group_ordinal; } } @@ -1403,7 +1410,8 @@ void hybrid_scan_reader_impl::set_pass_page_mask(std::span data_page data_page_mask.size() >= num_inserted_data_pages + num_data_pages_this_col_chunk, "Encountered invalid data page mask size"); - if (chunks[chunk_idx].num_dict_pages > 0) { _pass_page_mask.push_back(true); } + // Insert a true value for each dictionary page + _pass_page_mask.insert(_pass_page_mask.end(), chunks[chunk_idx].num_dict_pages, true); // Insert page mask for this column chunk _pass_page_mask.insert( From d6a6f4c58939c81a40c95f3449f3af0140ddcffe Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Fri, 7 Aug 2026 14:59:40 -0700 Subject: [PATCH 06/10] Update cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu --- cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu b/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu index b142e1aff241..73207b76cb93 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu +++ b/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu @@ -163,7 +163,7 @@ void hybrid_scan_reader_impl::setup_next_pass( std::size_t{0}, [](auto size, auto const& page) { return size + page.size(); }); } - auto chunk_iter = thrust::make_transform_iterator( + auto chunk_iter = cuda::make_transform_iterator( pass.chunks.d_begin(), parquet::detail::get_chunk_compressed_size{}); return cudf::detail::reduce( chunk_iter, chunk_iter + pass.chunks.size(), size_t{0}, cuda::std::plus{}, _stream); From b563b7402f0505c9c382db26f6c3bdb1a860a3e7 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Mon, 17 Aug 2026 23:23:42 +0000 Subject: [PATCH 07/10] Use `cuda::stream_ref` --- cpp/include/cudf/io/experimental/hybrid_scan_multifile.hpp | 4 ++-- cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu | 4 ++-- cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp | 4 ++-- cpp/src/io/parquet/experimental/hybrid_scan_impl.hpp | 4 ++-- cpp/src/io/parquet/experimental/hybrid_scan_multifile.cpp | 4 ++-- cpp/tests/io/experimental/hybrid_scan_multifile_composer.cpp | 2 +- cpp/tests/io/experimental/hybrid_scan_multifile_composer.hpp | 2 +- 7 files changed, 12 insertions(+), 12 deletions(-) diff --git a/cpp/include/cudf/io/experimental/hybrid_scan_multifile.hpp b/cpp/include/cudf/io/experimental/hybrid_scan_multifile.hpp index eaa55120fbdb..1343ebb3e789 100644 --- a/cpp/include/cudf/io/experimental/hybrid_scan_multifile.hpp +++ b/cpp/include/cudf/io/experimental/hybrid_scan_multifile.hpp @@ -280,7 +280,7 @@ class hybrid_scan_multifile { payload_pages_byte_ranges(cudf::host_span const> row_group_indices, cudf::column_view const& row_mask, parquet_reader_options const& options, - rmm::cuda_stream_view stream) const; + cuda::stream_ref stream) const; /** * @brief Materialize payload columns and applies the row mask to the output table @@ -429,7 +429,7 @@ class hybrid_scan_multifile { cudf::column_view const& row_mask, cudf::host_span const> page_data, parquet_reader_options const& options, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const; /** diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu b/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu index 039e13480f9b..3a07dc14806b 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu +++ b/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu @@ -163,8 +163,8 @@ void hybrid_scan_reader_impl::setup_next_pass( std::size_t{0}, [](auto size, auto const& page) { return size + page.size(); }); } - auto chunk_iter = cuda::make_transform_iterator( - pass.chunks.d_begin(), parquet::detail::get_chunk_compressed_size{}); + auto chunk_iter = cuda::make_transform_iterator(pass.chunks.d_begin(), + parquet::detail::get_chunk_compressed_size{}); return cudf::detail::reduce( chunk_iter, chunk_iter + pass.chunks.size(), size_t{0}, cuda::std::plus{}, _stream); }(); diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp b/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp index 46eb20f4254f..06a7c479984f 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp +++ b/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp @@ -517,7 +517,7 @@ hybrid_scan_reader_impl::payload_pages_byte_ranges( std::span const> row_group_indices, cudf::column_view const& row_mask, parquet_reader_options const& options, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { CUDF_EXPECTS(row_group_indices.size() == _extended_metadata->get_num_sources(), "Row group source count must match the number of input sources"); @@ -855,7 +855,7 @@ void hybrid_scan_reader_impl::setup_chunking_for_payload_columns( cudf::column_view const& row_mask, std::span const> page_data, parquet_reader_options const& options, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(std::cmp_equal(row_mask.size(), total_rows_in_row_groups(row_group_indices)), diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_impl.hpp b/cpp/src/io/parquet/experimental/hybrid_scan_impl.hpp index 39b8e7360eaa..d6b14b2aaa15 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_impl.hpp +++ b/cpp/src/io/parquet/experimental/hybrid_scan_impl.hpp @@ -200,7 +200,7 @@ class hybrid_scan_reader_impl : public parquet::detail::reader_impl { payload_pages_byte_ranges(std::span const> row_group_indices, cudf::column_view const& row_mask, parquet_reader_options const& options, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); /** * @copydoc cudf::io::parquet::experimental::hybrid_scan_multifile::materialize_payload_columns @@ -279,7 +279,7 @@ class hybrid_scan_reader_impl : public parquet::detail::reader_impl { cudf::column_view const& row_mask, std::span const> page_data, parquet_reader_options const& options, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_multifile.cpp b/cpp/src/io/parquet/experimental/hybrid_scan_multifile.cpp index e4262cec32d1..10fad9df9439 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_multifile.cpp +++ b/cpp/src/io/parquet/experimental/hybrid_scan_multifile.cpp @@ -142,7 +142,7 @@ hybrid_scan_multifile::payload_pages_byte_ranges( cudf::host_span const> row_group_indices, cudf::column_view const& row_mask, parquet_reader_options const& options, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { CUDF_FUNC_RANGE(); return _impl->payload_pages_byte_ranges(row_group_indices, row_mask, options, stream); @@ -242,7 +242,7 @@ void hybrid_scan_multifile::setup_chunking_for_payload_columns( cudf::column_view const& row_mask, cudf::host_span const> page_data, parquet_reader_options const& options, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { CUDF_FUNC_RANGE(); diff --git a/cpp/tests/io/experimental/hybrid_scan_multifile_composer.cpp b/cpp/tests/io/experimental/hybrid_scan_multifile_composer.cpp index e17cd530963f..de22feb0da8f 100644 --- a/cpp/tests/io/experimental/hybrid_scan_multifile_composer.cpp +++ b/cpp/tests/io/experimental/hybrid_scan_multifile_composer.cpp @@ -164,7 +164,7 @@ chunked_sparse_hybrid_scan_multifile( cudf::ast::operation const& filter_expression, std::optional> const& payload_column_names, bool case_sensitive_names, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto options = cudf::io::parquet_reader_options::builder() diff --git a/cpp/tests/io/experimental/hybrid_scan_multifile_composer.hpp b/cpp/tests/io/experimental/hybrid_scan_multifile_composer.hpp index 19a93f7b9a0e..90bf42405a1e 100644 --- a/cpp/tests/io/experimental/hybrid_scan_multifile_composer.hpp +++ b/cpp/tests/io/experimental/hybrid_scan_multifile_composer.hpp @@ -100,7 +100,7 @@ chunked_sparse_hybrid_scan_multifile( cudf::ast::operation const& filter_expression, std::optional> const& payload_column_names, bool case_sensitive_names, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** From d202c7197be15c6f107b81540add89cc22065fad Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Mon, 17 Aug 2026 23:39:13 +0000 Subject: [PATCH 08/10] Minor bug fix --- cpp/src/io/parquet/experimental/page_index_filter.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/src/io/parquet/experimental/page_index_filter.cu b/cpp/src/io/parquet/experimental/page_index_filter.cu index 0f0473eb9341..0ddbc5d252c2 100644 --- a/cpp/src/io/parquet/experimental/page_index_filter.cu +++ b/cpp/src/io/parquet/experimental/page_index_filter.cu @@ -1176,7 +1176,7 @@ thrust::host_vector aggregate_reader_metadata::compute_data_page_mask( // 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(0, stream); + auto data_page_mask = thrust::host_vector{}; data_page_mask.reserve(total_pages); auto host_results_iter = host_results.begin(); stream.sync(); From 9c1eeb3e5cda2baeb3c58fa736e09dae9b944ae5 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Tue, 18 Aug 2026 01:14:41 +0000 Subject: [PATCH 09/10] style fix --- cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu b/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu index b75e2df80acb..970f790da3e3 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu +++ b/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu @@ -164,7 +164,7 @@ void hybrid_scan_reader_impl::setup_next_pass( [](auto size, auto const& page) { return size + page.size(); }); } auto chunk_iter = cuda::transform_iterator(pass.chunks.d_begin(), - parquet::detail::get_chunk_compressed_size{}); + parquet::detail::get_chunk_compressed_size{}); return cudf::detail::reduce( chunk_iter, chunk_iter + pass.chunks.size(), size_t{0}, cuda::std::plus{}, _stream); }(); From 37ee40499ce7994c3622e46d47804c337d25330a Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Tue, 18 Aug 2026 20:36:22 +0000 Subject: [PATCH 10/10] Remove unnecessary column selection --- cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp b/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp index 06a7c479984f..006d11450fc9 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp +++ b/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp @@ -863,7 +863,6 @@ void hybrid_scan_reader_impl::setup_chunking_for_payload_columns( CUDF_EXPECTS(row_mask.null_count() == 0, "Row mask must not have any nulls when materializing payload column"); - reset_column_selection(); prepare_materialization( read_columns_mode::PAYLOAD_COLUMNS, row_group_indices.size(), options, stream, mr);