From 6e4040ab4b6565d3cfa87039d9bc8c2e17447da4 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb Date: Tue, 21 Jul 2026 17:39:20 +0000 Subject: [PATCH 01/16] 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 cebbcd5b7ef4bda8a18ab9056ad08c6c3ffc76e4 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb Date: Wed, 22 Jul 2026 00:39:44 +0000 Subject: [PATCH 02/16] Improve sparse Parquet page decoding Refine page topology preprocessing and add pruned-page decoding with hybrid scan coverage. --- cpp/CMakeLists.txt | 1 + cpp/src/io/parquet/decode_pruned_pages.cu | 103 +++++++++ .../experimental/hybrid_scan_preprocess.cu | 2 +- cpp/src/io/parquet/page_hdr.cu | 206 ++++++++---------- cpp/src/io/parquet/parquet_gpu.hpp | 56 +++-- cpp/src/io/parquet/reader_impl_helpers.cpp | 9 +- cpp/src/io/parquet/reader_impl_preprocess.cu | 99 ++------- .../parquet/reader_impl_preprocess_utils.cu | 93 +++++--- .../parquet/reader_impl_preprocess_utils.cuh | 21 +- .../io/experimental/hybrid_scan_test.cpp | 107 +++++++++ 10 files changed, 439 insertions(+), 258 deletions(-) create mode 100644 cpp/src/io/parquet/decode_pruned_pages.cu diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index f542bc4b6dc8..3790478a2481 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -818,6 +818,7 @@ add_library( src/io/parquet/writer_impl.cu src/io/parquet/writer_impl_helpers.cpp src/io/parquet/decode_fixed.cu + src/io/parquet/decode_pruned_pages.cu src/io/statistics/orc_column_statistics.cu src/io/statistics/parquet_column_statistics.cu src/io/text/byte_range_info.cpp diff --git a/cpp/src/io/parquet/decode_pruned_pages.cu b/cpp/src/io/parquet/decode_pruned_pages.cu new file mode 100644 index 000000000000..127952de12a2 --- /dev/null +++ b/cpp/src/io/parquet/decode_pruned_pages.cu @@ -0,0 +1,103 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include "parquet_gpu.hpp" + +#include + +#include +#include + +namespace cudf::io::parquet::detail { + +namespace { + +auto constexpr block_size = 4 * cudf::detail::warp_size; + +/** + * @brief Initialize output entries for pruned string and list pages. + * + * String entries receive either the page's initial offset for small strings or a zero size for + * large strings. List offsets receive the start value of their child nesting level. These targeted + * writes avoid zero-initializing entire output buffers. + */ +CUDF_KERNEL void __launch_bounds__(block_size) + fill_pruned_offsets_kernel(device_span pages, + device_span chunks, + device_span page_mask, + size_t skip_rows, + size_t num_rows) +{ + namespace cg = cooperative_groups; + + auto const block = cg::this_thread_block(); + auto const page_idx = cg::this_grid().block_rank(); + auto const t = static_cast(block.thread_rank()); + if (page_mask[page_idx]) { return; } + + auto const& page = pages[page_idx]; + auto const& chunk = chunks[page.chunk_idx]; + + if (chunk.column_data_base == nullptr) { return; } + + auto const is_list_col = chunk.max_level[level_type::REPETITION] != 0; + + // Write offsets for pruned non-list (flat) string columns. + // Mirrors `update_string_offsets_for_pruned_pages` in page_string_utils.cuh. + if (not is_list_col and is_string_col(chunk)) { + auto data = static_cast(chunk.column_data_base[chunk.max_nesting_depth - 1]); + if (data == 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); + // Write zeros for large strings and the page's initial offset otherwise. + auto const value = + chunk.is_large_string_col ? size_type{0} : static_cast(page.str_offset); + for (auto row = begin + t; row < end; row += block.size()) { + data[row - skip_rows] = value; + } + return; + } + + // Write offsets to list locations at each depth. + // Mirrors `update_list_offsets_for_pruned_pages` in page_decode.cuh. + if (is_list_col and page.nesting != nullptr and page.nesting_decode != nullptr) { + for (auto depth = 0; depth < chunk.max_nesting_depth - 1; depth++) { + auto offsets = static_cast(chunk.column_data_base[depth]); + auto& nesting_info = page.nesting[depth]; + // Pruned list pages retain rows through the first list level but contribute no child values. + // The preprocessing pass computes the output range and child start value at every list depth. + if (nesting_info.type != type_id::LIST or offsets == nullptr) { continue; } + // Emit an offset for the current nesting level equal to current length of the next nesting + // level + auto const output_begin = page.nesting_decode[depth].page_start_value; + auto const offset = page.nesting_decode[depth + 1].page_start_value; + for (auto offset_idx = t; offset_idx < nesting_info.batch_size; + offset_idx += static_cast(block.size())) { + offsets[output_begin + offset_idx] = offset; + } + } + } +} + +} // namespace + +void fill_pruned_offsets(cudf::device_span pages, + cudf::device_span chunks, + cudf::device_span page_mask, + size_t skip_rows, + size_t num_rows, + rmm::cuda_stream_view stream) +{ + CUDF_EXPECTS(pages.size() == page_mask.size(), "Page mask size does not match page count"); + fill_pruned_offsets_kernel<<>>( + pages, chunks, page_mask, skip_rows, num_rows); + CUDF_CUDA_TRY(cudaGetLastError()); +} + +} // namespace cudf::io::parquet::detail diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu b/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu index 64037e0c87df..060ff53bf7e1 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu +++ b/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu @@ -57,7 +57,7 @@ void decode_dictionary_page_headers(cudf::detail::hostdevice_spanpage.kernel_mask = decode_kernel_mask::NONE; } +/** + * @brief Decode a page header from an initialized byte stream. + * + * @param bs Byte stream + * @param chunk_idx Index of the chunk containing the page + * @param page Pointer to the page info to decode + * @param error_code Pointer to the error code for kernel failures + */ +__device__ void decode_page_header(byte_stream_s* bs, + cudf::size_type chunk_idx, + PageInfo* page, + kernel_error::pointer error_code) +{ + bs->page.chunk_idx = chunk_idx; + bs->page.src_col_schema = bs->ck.src_col_schema; + + 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; + } +} + /** * @brief Kernel for outputting page headers from the specified column chunks * @@ -530,7 +570,7 @@ void __forceinline__ __device__ zero_out_page_header_info(byte_stream_s* bs) CUDF_KERNEL void __launch_bounds__(decode_page_headers_block_size) decode_page_headers_kernel(device_span chunks, - chunk_page_info* chunk_pages, + device_span chunk_pages, kernel_error::pointer error_code) { auto constexpr num_warps_per_block = decode_page_headers_block_size / cudf::detail::warp_size; @@ -739,26 +779,22 @@ CUDF_KERNEL void __launch_bounds__(count_page_headers_block_size) /** * @brief Functor to decode page headers from specified page locations */ -struct decode_page_headers_with_pgidx_fn { +struct decode_using_page_index_fn { cudf::device_span colchunks; cudf::device_span pages; + cudf::device_span chunk_page_offsets; uint8_t** page_locations; - 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()); - - // Binary search the the column chunk index for this page - auto const chunk_idx = static_cast( + auto const chunk_idx = static_cast( cuda::std::distance( - chunk_page_offsets, + chunk_page_offsets.begin(), thrust::upper_bound( - thrust::seq, chunk_page_offsets, chunk_page_offsets + num_chunks + 1, page_idx)) - + thrust::seq, chunk_page_offsets.begin(), chunk_page_offsets.end(), page_idx)) - 1); - - // Check if the chunk index is valid if (chunk_idx < 0 or chunk_idx >= num_chunks) { set_error(static_cast(decode_error::DATA_STREAM_OVERRUN), error_code); @@ -778,59 +814,24 @@ struct decode_page_headers_with_pgidx_fn { // Clear page header info before writing known fields zero_out_page_header_info(&bs); - bs.page.chunk_idx = chunk_idx; - bs.page.src_col_schema = bs.ck.src_col_schema; - - // bs.page.chunk_row not computed here and will be filled in later by - // `fill_in_page_info()`. - - // Parsed page must be valid and not empty - 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: - // this computation is only valid for flat schemas. for nested schemas, - // they will be recomputed in the preprocess step by examining repetition and - // definition levels - bs.page.num_rows = bs.page.num_input_values; - break; - case PageType::DATA_PAGE_V2: - bs.page.flags |= PAGEINFO_FLAGS_V2; - // V2 only uses RLE, so it was removed from the header - 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; - } + decode_page_header(&bs, chunk_idx, &pages[page_idx], error_code); bs.page.page_data = const_cast(bs.cur); bs.page.kernel_mask = kernel_mask_for_page(bs.page, bs.ck); - // Copy over the page info from byte stream + // Copy the page info to the output span pages[page_idx] = bs.page; } }; /** - * @brief Functor to decode indexed page headers from exact page spans + * @brief Functor to decode specified page headers from corresponding page data spans */ -struct decode_page_headers_with_pgidx_spans_fn { +struct decode_from_page_data_fn { cudf::device_span colchunks; cudf::device_span pages; - cudf::device_span const> page_spans; - size_type* chunk_page_offsets; + cudf::device_span const> page_data; + cudf::device_span chunk_page_offsets; kernel_error::pointer error_code; __device__ void operator()(size_type page_idx) const noexcept @@ -838,11 +839,10 @@ struct decode_page_headers_with_pgidx_spans_fn { auto const num_chunks = static_cast(colchunks.size()); auto const chunk_idx = static_cast( cuda::std::distance( - chunk_page_offsets, + chunk_page_offsets.begin(), thrust::upper_bound( - thrust::seq, chunk_page_offsets, chunk_page_offsets + num_chunks + 1, page_idx)) - + thrust::seq, chunk_page_offsets.begin(), chunk_page_offsets.end(), page_idx)) - 1); - if (chunk_idx < 0 or chunk_idx >= num_chunks) { set_error(static_cast(decode_error::DATA_STREAM_OVERRUN), error_code); @@ -852,43 +852,21 @@ struct decode_page_headers_with_pgidx_spans_fn { 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; + auto const page_span = page_data[page_idx]; + if (page_span.empty()) { + // Initialize the logical page descriptor. Page-index fields are populated by + // fill_in_page_info(). + bs.page.chunk_idx = chunk_idx; + bs.page.src_col_schema = bs.ck.src_col_schema; + 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; - } + bs.base = bs.cur = page_span.data(); + bs.end = page_span.data() + page_span.size(); - 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; - } + decode_page_header(&bs, chunk_idx, &pages[page_idx], error_code); if (bs.page.compressed_page_size < 0 or static_cast(bs.end - bs.cur) != static_cast(bs.page.compressed_page_size)) { @@ -997,12 +975,14 @@ void count_page_headers(cudf::detail::hostdevice_span chunks, } void decode_page_headers(cudf::device_span chunks, - chunk_page_info* chunk_pages, + cudf::device_span chunk_pages, kernel_error::pointer error_code, rmm::cuda_stream_view stream) { static_assert(decode_page_headers_block_size % cudf::detail::warp_size == 0, "Block size for decode page headers kernel must be a multiple of warp size"); + CUDF_EXPECTS(chunk_pages.size() == chunks.size(), + "Chunk page info must contain one entry per chunk"); auto const num_chunks = static_cast(chunks.size()); auto constexpr num_warps_per_block = decode_page_headers_block_size / cudf::detail::warp_size; @@ -1017,41 +997,45 @@ void decode_page_headers(cudf::device_span chunks, CUDF_CUDA_TRY(cudaGetLastError()); } -void decode_page_headers_with_pgidx(cudf::device_span chunks, - cudf::device_span pages, - uint8_t** page_locations, - size_type* chunk_page_offsets, - kernel_error::pointer error_code, - rmm::cuda_stream_view stream) +void decode_page_headers_using_page_index(cudf::device_span chunks, + cudf::device_span pages, + uint8_t** page_locations, + cudf::device_span chunk_page_offsets, + kernel_error::pointer error_code, + rmm::cuda_stream_view stream) { + CUDF_EXPECTS(chunk_page_offsets.size() == chunks.size() + 1, + "Chunk page offsets must cover all chunks"); 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_fn{.colchunks = chunks, - .pages = pages, - .page_locations = page_locations, - .chunk_page_offsets = chunk_page_offsets, - .error_code = error_code}); + decode_using_page_index_fn{.colchunks = chunks, + .pages = pages, + .chunk_page_offsets = chunk_page_offsets, + .page_locations = page_locations, + .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) +void decode_page_headers_from_page_data( + cudf::device_span chunks, + cudf::device_span pages, + cudf::device_span const> page_data, + cudf::device_span chunk_page_offsets, + kernel_error::pointer error_code, + rmm::cuda_stream_view stream) { - CUDF_EXPECTS(page_spans.size() == pages.size(), + CUDF_EXPECTS(page_data.size() == pages.size(), "Page span count must match the number of logical pages"); + CUDF_EXPECTS(chunk_page_offsets.size() == chunks.size() + 1, + "Chunk page offsets must cover all chunks"); 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}); + decode_from_page_data_fn{.colchunks = chunks, + .pages = pages, + .page_data = page_data, + .chunk_page_offsets = chunk_page_offsets, + .error_code = error_code}); } void build_string_dictionary_index(ColumnChunkDesc* chunks, diff --git a/cpp/src/io/parquet/parquet_gpu.hpp b/cpp/src/io/parquet/parquet_gpu.hpp index c821f14e2b9e..4a305b6ed9c1 100644 --- a/cpp/src/io/parquet/parquet_gpu.hpp +++ b/cpp/src/io/parquet/parquet_gpu.hpp @@ -707,7 +707,7 @@ void count_page_headers(cudf::detail::hostdevice_span chunks, * @param[in] stream CUDA stream to use */ void decode_page_headers(cudf::device_span chunks, - chunk_page_info* chunk_pages, + cudf::device_span chunk_pages, kernel_error::pointer error_code, rmm::cuda_stream_view stream); @@ -721,25 +721,32 @@ void decode_page_headers(cudf::device_span chunks, * @param[out] error_code Error code for kernel failures * @param[in] stream CUDA stream to use */ -void decode_page_headers_with_pgidx(cudf::device_span chunks, - cudf::device_span pages, - uint8_t** page_locations, - size_type* chunk_page_offsets, - kernel_error::pointer error_code, - rmm::cuda_stream_view stream); +void decode_page_headers_using_page_index(cudf::device_span chunks, + cudf::device_span pages, + uint8_t** page_locations, + cudf::device_span chunk_page_offsets, + kernel_error::pointer error_code, + rmm::cuda_stream_view stream); /** - * @brief Decode indexed page headers from exact, potentially discontiguous page spans + * @brief Decode specified page headers from corresponding page data spans. + * + * Empty spans initialize the corresponding logical page descriptor but are not decoded. * - * Empty spans initialize the corresponding logical page descriptor but are not parsed. + * @param[in] chunks Device span of column chunks + * @param[out] pages Device span of pages + * @param[in] page_data Device span of page data + * @param[in] chunk_page_offsets List of running count of page locations per column chunk + * @param[out] error_code Error code for kernel failures + * @param[in] stream CUDA stream to use */ -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); +void decode_page_headers_from_page_data( + cudf::device_span chunks, + cudf::device_span pages, + cudf::device_span const> page_data, + cudf::device_span chunk_page_offsets, + kernel_error::pointer error_code, + rmm::cuda_stream_view stream); /** * @brief Launches kernel for building the dictionary index for the column @@ -1030,6 +1037,23 @@ void preprocess_levels(cudf::detail::hostdevice_span pages, int level_type_size, rmm::cuda_stream_view stream); +/** + * @brief Fills output offset entries for pruned string and list pages + * + * @param[in] pages All pages to be processed + * @param[in] chunks All chunks to be processed + * @param[in] page_mask Boolean vector indicating which pages are decoded + * @param[in] skip_rows Number of rows to skip + * @param[in] num_rows Number of rows to read + * @param[in] stream CUDA stream to use + */ +void fill_pruned_offsets(cudf::device_span pages, + cudf::device_span chunks, + cudf::device_span page_mask, + size_t skip_rows, + size_t num_rows, + rmm::cuda_stream_view stream); + /** * @brief Launches kernel for reading non-dictionary fixed width column data stored in the pages * diff --git a/cpp/src/io/parquet/reader_impl_helpers.cpp b/cpp/src/io/parquet/reader_impl_helpers.cpp index d25d48b2b876..ed60a4db0e2d 100644 --- a/cpp/src/io/parquet/reader_impl_helpers.cpp +++ b/cpp/src/io/parquet/reader_impl_helpers.cpp @@ -719,13 +719,10 @@ 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; - // Page locations and row boundaries only require the offset index. Additional value-count - // metadata is populated below when a column index is also available. + // Return early if any columns lack the offset index. 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.has_value() ? &col_chunk.column_index.value() : nullptr; auto& chunk_info = chunks[col_idx]; auto const num_pages = offset_index.page_locations.size(); @@ -750,6 +747,10 @@ void aggregate_reader_metadata::column_info_for_row_group(row_group_info& rg_inf } } + // Populate additional value-count metadata if column index is also available. + auto const* column_index = + col_chunk.column_index.has_value() ? &col_chunk.column_index.value() : nullptr; + // Use the definition_level_histogram to get num_valid and num_null. For now, these are // only ever used for byte array columns. The repetition_level_histogram might be // necessary to determine the total number of values in the page if the diff --git a/cpp/src/io/parquet/reader_impl_preprocess.cu b/cpp/src/io/parquet/reader_impl_preprocess.cu index 000c18554d05..91add706b279 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 @@ -63,36 +62,6 @@ struct set_str_bytes_all { } }; -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 void reader_impl::build_string_dict_indices() @@ -886,11 +855,10 @@ 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()), - cuda::counting_iterator{0}, - cuda::counting_iterator{static_cast(subpass.pages.size())}, - set_str_bytes_all{subpass.pages, subpass_page_mask_span()}); + thrust::for_each(rmm::exec_policy_nosync(_stream, cudf::get_current_device_resource_ref()), + cuda::counting_iterator(0), + cuda::counting_iterator(subpass.pages.size()), + set_str_bytes_all{subpass.pages, subpass_page_mask_span()}); } // retrieve pages back @@ -972,10 +940,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(); @@ -999,10 +963,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)) / @@ -1120,11 +1082,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)) / @@ -1143,53 +1102,23 @@ void reader_impl::allocate_columns(read_mode mode, size_t skip_rows, size_t num_ void reader_impl::fill_pruned_offsets(size_t skip_rows, size_t num_rows) { - auto& pass = *_pass_itm_data; - auto& subpass = *pass.subpass; + // Return early if there are no pruned pages 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& pass = *_pass_itm_data; + auto const& subpass = *pass.subpass; + 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); + + // Set offsets for pruned string and list pages. + parquet::detail::fill_pruned_offsets( + pages, chunks, device_page_mask, skip_rows, num_rows, _stream); } cudf::detail::host_vector reader_impl::calculate_page_string_offsets() diff --git a/cpp/src/io/parquet/reader_impl_preprocess_utils.cu b/cpp/src/io/parquet/reader_impl_preprocess_utils.cu index 25af24a9487f..ae88f300ee81 100644 --- a/cpp/src/io/parquet/reader_impl_preprocess_utils.cu +++ b/cpp/src/io/parquet/reader_impl_preprocess_utils.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -272,6 +272,7 @@ void fill_in_page_info(host_span chunks, { auto const num_pages = pages.size(); auto page_indexes = cudf::detail::make_pinned_vector_async(num_pages, stream); + std::fill(page_indexes.begin(), page_indexes.end(), page_index_info{}); for (size_t c = 0, page_count = 0; c < chunks.size(); c++) { auto const& chunk = chunks[c]; @@ -280,17 +281,16 @@ void fill_in_page_info(host_span chunks, size_t start_row = 0; page_count += chunk.num_dict_pages; for (size_t p = 0; p < chunk_info.pages.size(); p++, page_count++) { - auto& page = page_indexes[page_count]; - page.num_rows = chunk_info.pages[p].num_rows; - page.chunk_row = start_row; - 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()); + auto& page = page_indexes[page_count]; + page.num_rows = chunk_info.pages[p].num_rows; + page.chunk_row = start_row; + 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 = static_cast( + (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; } @@ -413,10 +413,29 @@ cudf::detail::hostdevice_vector sort_pages(device_span namespace { +/** + * @brief Page data source type + */ +enum class page_data_source_type : uint8_t { + COLUMN_CHUNKS = 0, + OFFSET_INDEX = 1, + PAGE_SPANS = 2, +}; + +/** + * @brief Dispatch decode page headers base on the type of page data source + * + * @tparam data_source_type Type of page data source + * + * @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 stream Stream to use + */ +template void decode_page_headers_impl(pass_intermediate_data& pass, device_span unsorted_pages, - bool has_page_index, - host_span const> page_spans, + host_span const> page_data, rmm::cuda_stream_view stream) { CUDF_FUNC_RANGE(); @@ -447,23 +466,22 @@ void decode_page_headers_impl(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(), + if constexpr (data_source_type == page_data_source_type::PAGE_SPANS) { + CUDF_EXPECTS(page_data.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( + auto device_page_data = cudf::detail::make_device_uvector_async( + page_data, stream, cudf::get_current_device_resource_ref()); + decode_page_headers_from_page_data( device_span(pass.chunks.device_ptr(), pass.chunks.size()), unsorted_pages, - device_page_spans, - chunk_page_offsets.begin(), + device_page_data, + device_span(chunk_page_offsets.data(), chunk_page_offsets.size()), error_code.data(), stream); } - // If page index is present, collect data ptrs for all pages and launch the accelerated decode + // If offset index is present, collect data ptrs for all pages and launch the accelerated decode // page headers kernel - else if (has_page_index) { + else if (data_source_type == page_data_source_type::OFFSET_INDEX) { auto host_page_locations = cudf::detail::make_pinned_vector_async(unsorted_pages.size(), stream); auto curr_page_idx = 0; @@ -514,18 +532,18 @@ void decode_page_headers_impl(pass_intermediate_data& pass, host_page_locations, stream, cudf::get_current_device_resource_ref()); // Accelerated decode page headers, one thread per page - decode_page_headers_with_pgidx( + decode_page_headers_using_page_index( device_span(pass.chunks.device_ptr(), pass.chunks.size()), unsorted_pages, page_locations.begin(), - chunk_page_offsets.begin(), + device_span(chunk_page_offsets.data(), chunk_page_offsets.size()), error_code.data(), stream); } else { - // (Slow) decode page headers, one warp (lane) per pages of a chunk + // (Slow) decode page headers, one warp (lane) per pages of a column chunk decode_page_headers( device_span(pass.chunks.device_ptr(), pass.chunks.size()), - d_chunk_page_info.begin(), + device_span(d_chunk_page_info.data(), d_chunk_page_info.size()), error_code.data(), stream); } @@ -541,7 +559,10 @@ void decode_page_headers_impl(pass_intermediate_data& pass, } } - if (has_page_index) { fill_in_page_info(pass.chunks, unsorted_pages, stream); } + if (data_source_type == page_data_source_type::OFFSET_INDEX or + data_source_type == page_data_source_type::PAGE_SPANS) { + fill_in_page_info(pass.chunks, unsorted_pages, stream); + } // compute max bytes needed for level data auto level_bit_size = cudf::detail::make_counting_transform_iterator( @@ -602,18 +623,24 @@ void decode_page_headers_impl(pass_intermediate_data& pass, void decode_page_headers(pass_intermediate_data& pass, device_span unsorted_pages, - bool has_page_index, + bool has_offset_index, rmm::cuda_stream_view stream) { - decode_page_headers_impl(pass, unsorted_pages, has_page_index, {}, stream); + if (has_offset_index) { + decode_page_headers_impl(pass, unsorted_pages, {}, stream); + } else { + decode_page_headers_impl( + pass, unsorted_pages, {}, stream); + } } void decode_page_headers(pass_intermediate_data& pass, device_span unsorted_pages, - host_span const> page_spans, + host_span const> page_data, rmm::cuda_stream_view stream) { - decode_page_headers_impl(pass, unsorted_pages, true, page_spans, stream); + decode_page_headers_impl( + pass, unsorted_pages, page_data, 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 b9183ad5351d..a7f0e7187197 100644 --- a/cpp/src/io/parquet/reader_impl_preprocess_utils.cuh +++ b/cpp/src/io/parquet/reader_impl_preprocess_utils.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -128,22 +128,27 @@ std::string encoding_to_string(Encoding encoding); * * @param pass The struct containing pass information * @param unsorted_pages Device span of page information to decode - * @param has_page_index Boolean indicating if the page index is available + * @param has_offset_index Boolean indicating if the page index is available * @param stream CUDA stream used for device memory operations and kernel launches */ void decode_page_headers(pass_intermediate_data& pass, device_span unsorted_pages, - bool has_page_index, + bool has_offset_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. + * Empty data spans represent masked pages and retain their logical page-index metadata. + * + * @param pass Struct containing pass information + * @param unsorted_pages Device span of page information to decode + * @param page_data Host span of page data device spans, one per logical indexed page + * @param stream CUDA stream used for device memory operations and kernel launches */ void decode_page_headers(pass_intermediate_data& pass, device_span unsorted_pages, - host_span const> page_spans, + host_span const> page_data, rmm::cuda_stream_view stream); /** @@ -167,7 +172,7 @@ struct page_index_info { int32_t num_nulls; int32_t num_valids; int32_t str_bytes; - bool has_value_info; + int32_t has_value_info; }; /** @@ -183,9 +188,9 @@ struct copy_page_info { 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.has_page_index = pi.has_value_info != 0; pg.start_val = 0; - if (pi.has_value_info) { + if (pg.has_page_index) { pg.num_nulls = pi.num_nulls; pg.num_valids = pi.num_valids; pg.str_bytes_from_index = pi.str_bytes; diff --git a/cpp/tests/io/experimental/hybrid_scan_test.cpp b/cpp/tests/io/experimental/hybrid_scan_test.cpp index e580994b6d34..edf43fc6bdff 100644 --- a/cpp/tests/io/experimental/hybrid_scan_test.cpp +++ b/cpp/tests/io/experimental/hybrid_scan_test.cpp @@ -13,11 +13,14 @@ #include #include +#include #include #include #include #include +#include #include +#include #include #include #include @@ -438,6 +441,110 @@ TEST_F(HybridScanTest, MaterializeListsOfStrings) test_hybrid_scan({col0, *col1, *col2, *col3, *col4}, false); } +TEST_F(HybridScanTest, ConsecutivePrunedPageOffsets) +{ + std::mt19937 gen(0x5ca1e); + auto constexpr num_rows = num_ordered_rows; + + auto col0 = testdata::ascending(); + auto col1 = testdata::ascending(); + auto col2 = make_parquet_list_col(gen, num_rows, 3, true); + col2 = cudf::purge_nonempty_nulls(col2->view()); + auto col3 = make_parquet_list_list_col(0, num_rows, 2, 3, true); + col3 = cudf::purge_nonempty_nulls(col3->view()); + auto col4 = make_list_str_column(gen, true, true); + col4 = cudf::purge_nonempty_nulls(col4->view()); + + auto const input = cudf::table_view{{col0, col1, *col2, *col3, *col4}}; + + std::string filepath = "ConsecutivePrunedPageOffsets.parquet"; + { + auto metadata = cudf::io::table_input_metadata(input); + metadata.column_metadata[0].set_name("col0"); + + auto const write_options = + cudf::io::parquet_writer_options::builder(cudf::io::sink_info{filepath}, input) + .metadata(std::move(metadata)) + .row_group_size_rows(num_rows) + .max_page_size_rows(page_size_for_ordered_tests) + .stats_level(cudf::io::statistics_freq::STATISTICS_COLUMN) + .build(); + cudf::io::write_parquet(write_options); + } + + auto const stream = cudf::get_default_stream(); + auto const mr = cudf::get_current_device_resource_ref(); + auto aligned_mr = rmm::mr::aligned_resource_adaptor(mr, bloom_filter_alignment); + + // Helper to validate monotonic string and list offsets + auto const expect_monotonic_offsets = [](auto& self, cudf::column_view const& column) -> void { + EXPECT_FALSE(cudf::has_nonempty_nulls(column)); + + if (column.type().id() == cudf::type_id::STRING) { + auto const offsets = cudf::strings_column_view{column}.offsets(); + auto const host_offsets = cudf::test::to_host(offsets).first; + EXPECT_TRUE(std::is_sorted(host_offsets.begin(), host_offsets.end())); + return; + } + + if (column.type().id() == cudf::type_id::LIST) { + auto const lists = cudf::lists_column_view{column}; + auto const host_offsets = cudf::test::to_host(lists.offsets()).first; + EXPECT_TRUE(std::is_sorted(host_offsets.begin(), host_offsets.end())); + self(self, lists.child()); + return; + } + + for (auto index = cudf::size_type{0}; index < column.num_children(); ++index) { + self(self, column.child(index)); + } + }; + + // Helper to validate the `offsets` children of string and list column + auto const validate = [&](cudf::ast::operation const& filter, + std::vector const& expected_slices) { + auto datasource = cudf::io::datasource::create({filepath}); + + auto const expected = cudf::concatenate(cudf::slice(input, expected_slices)); + std::unique_ptr filter_table; + std::unique_ptr payload_table; + ASSERT_NO_THROW(std::tie(filter_table, payload_table) = + hybrid_scan(*datasource, filter, {}, true, stream, mr, aligned_mr)); + + CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected->view().select({0}), filter_table->view()); + CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected->view().select({1, 2, 3, 4}), + payload_table->view()); + for (auto const& column : payload_table->view()) { + expect_monotonic_offsets(expect_monotonic_offsets, column); + } + }; + + auto const col_ref = cudf::ast::column_name_reference{"col0"}; + + // Prune two leading pages. + auto middle_value = cudf::numeric_scalar{2 * page_size_for_ordered_tests / 100}; + auto middle = cudf::ast::literal{middle_value}; + auto const keep_trailing = + cudf::ast::operation{cudf::ast::ast_operator::GREATER_EQUAL, col_ref, middle}; + validate(keep_trailing, {2 * page_size_for_ordered_tests, num_rows}); + + // Prune two trailing pages. + auto const keep_leading = cudf::ast::operation{cudf::ast::ast_operator::LESS, col_ref, middle}; + validate(keep_leading, {0, 2 * page_size_for_ordered_tests}); + + // Prune two middle pages. + auto lower_value = cudf::numeric_scalar{page_size_for_ordered_tests / 100}; + auto upper_value = cudf::numeric_scalar{3 * page_size_for_ordered_tests / 100}; + auto lower = cudf::ast::literal{lower_value}; + auto upper = cudf::ast::literal{upper_value}; + auto const keep_first = cudf::ast::operation{cudf::ast::ast_operator::LESS, col_ref, lower}; + auto const keep_last = + cudf::ast::operation{cudf::ast::ast_operator::GREATER_EQUAL, col_ref, upper}; + auto const keep_outer = + cudf::ast::operation{cudf::ast::ast_operator::LOGICAL_OR, keep_first, keep_last}; + validate(keep_outer, {0, page_size_for_ordered_tests, 3 * page_size_for_ordered_tests, num_rows}); +} + TEST_F(HybridScanTest, MaterializeStructs) { std::mt19937 gen(0xbaLL); From 0612e616aa6a9f137f6a86a6ce0d8a8a2a4c560d Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb Date: Sat, 18 Jul 2026 05:15:47 +0000 Subject: [PATCH 03/16] Sparse page IO --- .../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 | 34 ++ cpp/src/io/parquet/reader_impl_helpers.cpp | 3 +- .../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 | 477 +++++++++++++++++- 12 files changed, 1151 insertions(+), 15 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 060ff53bf7e1..794de5edee97 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu +++ b/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu @@ -180,6 +180,40 @@ 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; } + rmm::device_uvector unsorted_pages(total_pages, _stream); + 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/src/io/parquet/reader_impl_helpers.cpp b/cpp/src/io/parquet/reader_impl_helpers.cpp index ed60a4db0e2d..4bd87050abc0 100644 --- a/cpp/src/io/parquet/reader_impl_helpers.cpp +++ b/cpp/src/io/parquet/reader_impl_helpers.cpp @@ -719,7 +719,8 @@ 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; - // Return early if any columns lack the offset index. + // 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(); 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..8c11621cc4ec 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,350 @@ 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, 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 df762d816605dd6f5e0af47e58ffd940d7ece96e Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb Date: Wed, 22 Jul 2026 03:55:23 +0000 Subject: [PATCH 04/16] Fix sparse page I/O string offsets Initialize sparse page state and carry variable-width offsets across pruned pages so page-level reads produce valid columns without overestimating string storage. --- .../experimental/hybrid_scan_preprocess.cu | 5 +- .../hybrid_scan_multifile_test.cpp | 48 +++++++++++++++++++ 2 files changed, 52 insertions(+), 1 deletion(-) diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu b/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu index 794de5edee97..bf072bcee4f3 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu +++ b/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu @@ -207,7 +207,10 @@ void hybrid_scan_reader_impl::setup_compressed_data( CUDF_EXPECTS(total_pages == _sparse_page_spans.size(), "Sparse page span count does not match page-index metadata"); if (total_pages <= 0) { return; } - rmm::device_uvector unsorted_pages(total_pages, _stream); + // `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"); diff --git a/cpp/tests/io/experimental/hybrid_scan_multifile_test.cpp b/cpp/tests/io/experimental/hybrid_scan_multifile_test.cpp index 8c11621cc4ec..84285e197350 100644 --- a/cpp/tests/io/experimental/hybrid_scan_multifile_test.cpp +++ b/cpp/tests/io/experimental/hybrid_scan_multifile_test.cpp @@ -299,6 +299,54 @@ TEST_F(HybridScanMultifileTest, PageLevelDictionaryPayloadByteReduction) 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; From 00e187ef1e79e35371dc3841acc0cce8a33c7875 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb Date: Wed, 22 Jul 2026 04:27:32 +0000 Subject: [PATCH 05/16] Use sparse page I/O in MINT-1T example Exercise source-grouped sparse payload reads and make dataset selection configurable for validated A/B benchmarks. --- cpp/examples/hybrid_scan_io/CMakeLists.txt | 1 + cpp/examples/hybrid_scan_io/io_utils.cpp | 72 ++ cpp/examples/hybrid_scan_io/io_utils.hpp | 53 ++ .../hybrid_scan_io/mint1t_hybrid_scan.cpp | 832 ++++++++++++++++++ 4 files changed, 958 insertions(+) create mode 100644 cpp/examples/hybrid_scan_io/mint1t_hybrid_scan.cpp diff --git a/cpp/examples/hybrid_scan_io/CMakeLists.txt b/cpp/examples/hybrid_scan_io/CMakeLists.txt index b0d12692fa5e..90fae1005891 100644 --- a/cpp/examples/hybrid_scan_io/CMakeLists.txt +++ b/cpp/examples/hybrid_scan_io/CMakeLists.txt @@ -50,6 +50,7 @@ add_hybrid_scan_example(hybrid_scan_io hybrid_scan_io.cpp) add_hybrid_scan_example(hybrid_scan_pipeline hybrid_scan_pipeline.cpp) add_hybrid_scan_example(hybrid_scan_multifile_single_step hybrid_scan_multifile_single_step.cpp) add_hybrid_scan_example(hybrid_scan_multifile_two_step hybrid_scan_multifile_two_step.cpp) +add_hybrid_scan_example(mint1t_hybrid_scan mint1t_hybrid_scan.cpp) # Install the example.parquet file install(FILES ${CMAKE_CURRENT_LIST_DIR}/example.parquet diff --git a/cpp/examples/hybrid_scan_io/io_utils.cpp b/cpp/examples/hybrid_scan_io/io_utils.cpp index f21bb5a47645..fec6f43ba930 100644 --- a/cpp/examples/hybrid_scan_io/io_utils.cpp +++ b/cpp/examples/hybrid_scan_io/io_utils.cpp @@ -3,6 +3,8 @@ * SPDX-License-Identifier: Apache-2.0 */ +#include "io_utils.hpp" + #include #include #include @@ -39,3 +41,73 @@ fetch_byte_ranges_async(cudf::io::datasource& datasource, // Using libcudf utility but may have custom implementation in the future return cudf::io::parquet::fetch_byte_ranges_to_device_async(datasource, byte_ranges, stream, mr); } + +multifile_inputs::multifile_inputs(cudf::io::source_info const& source_info) + : datasources{cudf::io::make_datasources(source_info)} +{ + datasource_refs.reserve(datasources.size()); + std::transform(datasources.begin(), + datasources.end(), + std::back_inserter(datasource_refs), + [](auto const& datasource) { return std::ref(*datasource); }); +} + +void multifile_inputs::fetch_footers() +{ + footer_buffers = cudf::io::parquet::fetch_footers_to_host(datasource_refs); + footer_byte_spans.clear(); + footer_byte_spans.reserve(footer_buffers.size()); + std::transform(footer_buffers.begin(), + footer_buffers.end(), + std::back_inserter(footer_byte_spans), + [](auto const& buffer) { return cudf::host_span{*buffer}; }); +} + +std::vector> group_byte_ranges_by_source( + std::pair, std::vector> const& + byte_ranges_and_source_map, + std::size_t num_sources) +{ + auto const& [byte_ranges, source_map] = byte_ranges_and_source_map; + CUDF_EXPECTS(byte_ranges.size() == source_map.size(), "Invalid source map size"); + + auto byte_ranges_per_source = + std::vector>(num_sources); + for (auto range_index = std::size_t{0}; range_index < byte_ranges.size(); ++range_index) { + auto const source_index = source_map[range_index]; + CUDF_EXPECTS( + source_index >= 0 and static_cast(source_index) < byte_ranges_per_source.size(), + "Invalid source index"); + byte_ranges_per_source[source_index].push_back(byte_ranges[range_index]); + } + return byte_ranges_per_source; +} + +multisource_device_data fetch_multisource_device_data( + multifile_inputs const& inputs, + std::pair, std::vector> const& + byte_ranges_and_source_map, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + 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, byte_ranges_per_source, stream, mr); + tasks.get(); + + auto flat_spans = std::vector>{}; + for (auto const& source_spans : per_source_spans) { + flat_spans.insert(flat_spans.end(), source_spans.begin(), source_spans.end()); + } + return {std::move(buffers), std::move(per_source_spans), std::move(flat_spans)}; +} diff --git a/cpp/examples/hybrid_scan_io/io_utils.hpp b/cpp/examples/hybrid_scan_io/io_utils.hpp index 3a977b16616b..5814a416ee87 100644 --- a/cpp/examples/hybrid_scan_io/io_utils.hpp +++ b/cpp/examples/hybrid_scan_io/io_utils.hpp @@ -57,3 +57,56 @@ fetch_byte_ranges_async(cudf::io::datasource& datasource, cudf::host_span byte_ranges, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr); + +/** + * @brief Owns the datasources, footer buffers, and byte spans for a multifile read. + * + * Call `fetch_footers()` after construction. Keeping datasource and footer setup separate allows + * examples to time those operations independently. + */ +struct multifile_inputs { + explicit multifile_inputs(cudf::io::source_info const& source_info); + + void fetch_footers(); + + std::vector> datasources; + std::vector> datasource_refs; + std::vector> footer_buffers; + std::vector> footer_byte_spans; +}; + +/** + * @brief Owns multifile device buffers and the corresponding per-source and flattened spans. + */ +struct multisource_device_data { + std::vector buffers; + std::vector>> per_source_spans; + std::vector> flat_spans; +}; + +/** + * @brief Regroups flattened byte ranges using the source map returned by Hybrid Scan. + */ +[[nodiscard]] std::vector> group_byte_ranges_by_source( + std::pair, std::vector> const& + byte_ranges_and_source_map, + std::size_t num_sources); + +/** + * @brief Fetches source-mapped multifile byte ranges and flattens the resulting device spans. + */ +[[nodiscard]] multisource_device_data fetch_multisource_device_data( + multifile_inputs const& inputs, + std::pair, std::vector> const& + byte_ranges_and_source_map, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr); + +/** + * @brief Fetches source-grouped multifile byte ranges. + */ +[[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); diff --git a/cpp/examples/hybrid_scan_io/mint1t_hybrid_scan.cpp b/cpp/examples/hybrid_scan_io/mint1t_hybrid_scan.cpp new file mode 100644 index 000000000000..5ea368ffb5da --- /dev/null +++ b/cpp/examples/hybrid_scan_io/mint1t_hybrid_scan.cpp @@ -0,0 +1,832 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include "common_utils.hpp" +#include "io_utils.hpp" + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +#include + +#include +#include +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace { + +using clock_type = std::chrono::steady_clock; +using cudf::io::parquet::experimental::hybrid_scan_multifile; +using cudf::io::parquet::experimental::use_data_page_mask; + +auto constexpr original_dataset_root = "/MINT1T_benchmarking_subset"; +auto constexpr rewritten_dataset_root = "/MINT1T_rewritten"; +auto constexpr default_refs = + "/MINT1T_benchmarking_subset/parquet/interleaved_image_url_refs/image_url_refs.parquet"; +auto constexpr default_image_dir = + "/MINT1T_rewritten/parquet/mint_1t_html_images_stable_row_ids/data"; +auto constexpr default_expected = "/MINT1T_rewritten/output/image_url_payloads.parquet"; +auto constexpr writable_root = "/MINT1T_rewritten"; + +std::vector const payload_columns{ + "image", "image_format", "mime_type", "image_size_bytes", "md5", "sha256", "width", "height"}; + +struct arguments { + std::filesystem::path dataset_root{rewritten_dataset_root}; + std::filesystem::path refs{default_refs}; + std::filesystem::path image_dir{default_image_dir}; + std::optional limit; + std::size_t pass_read_limit{}; + bool pass_read_limit_set{false}; + bool use_page_mask{true}; + bool use_sparse_page_io{true}; + bool drop_cache{false}; + bool validate{false}; + std::filesystem::path expected_output{default_expected}; + std::optional output; +}; + +struct reference { + std::string image_parquet; + uint32_t row_offset; + + bool operator<(reference const& other) const + { + return std::tie(image_parquet, row_offset) < std::tie(other.image_parquet, other.row_offset); + } + + bool operator==(reference const& other) const + { + return image_parquet == other.image_parquet and row_offset == other.row_offset; + } +}; + +struct timing_data { + std::vector> stages; + + void add_seconds(std::string name, double seconds) + { + std::cout << "TIMING " << std::left << std::setw(38) << name << std::right << std::fixed + << std::setprecision(6) << seconds << " s\n" + << std::flush; + stages.emplace_back(std::move(name), seconds); + } + + void add(std::string name, clock_type::time_point start) + { + auto const seconds = std::chrono::duration(clock_type::now() - start).count(); + auto const existing = std::find_if( + stages.begin(), stages.end(), [&](auto const& item) { return item.first == name; }); + if (existing == stages.end()) { + add_seconds(std::move(name), seconds); + } else { + existing->second += seconds; + } + } + + [[nodiscard]] double get(std::string_view name) const + { + auto const it = std::find_if( + stages.begin(), stages.end(), [&](auto const& item) { return item.first == name; }); + return it == stages.end() ? 0.0 : it->second; + } + + [[nodiscard]] double sum(std::span names) const + { + return std::accumulate( + names.begin(), names.end(), 0.0, [&](double total, auto name) { return total + get(name); }); + } + + void print() const + { + std::cout << "\nStage timings:\n"; + for (auto const& [name, seconds] : stages) { + std::cout << " " << std::left << std::setw(38) << name << std::right << std::fixed + << std::setprecision(6) << seconds << " s\n"; + } + } +}; + +struct cache_drop_result { + bool attempted{}; + std::size_t files{}; + std::size_t errors{}; +}; + +void print_usage() +{ + std::cout + << "Usage: mint1t_hybrid_scan [options]\n\n" + << " --dataset-root PATH Dataset root; sets the default image-shard directory\n" + << " --refs PATH Reference Parquet (default: original manifest)\n" + << " --image-parquet-dir PATH Directory containing image Parquet shards\n" + << " --limit N Use the first N manifest rows\n" + << " --pass-read-limit-mib N Per-pass limit; default is 105% of GPU memory, 0 unbounded\n" + << " --use-data-page-mask YES|NO Toggle payload data-page pruning (default YES)\n" + << " --use-sparse-page-io YES|NO Toggle sparse physical payload I/O (default YES)\n" + << " --best-effort-drop-cache Apply POSIX_FADV_DONTNEED before timed source setup\n" + << " --validate Compare with the projected Python output\n" + << " --expected-output PATH Validation Parquet used by --validate\n" + << " --output PATH Write the payload-only C++ result under /MINT1T_rewritten\n" + << " -h, --help Show this message\n"; +} + +arguments parse_args(int argc, char const** argv) +{ + arguments args; + auto image_dir_set = false; + auto dataset_root_set = false; + auto expected_output_set = false; + auto require_value = [&](int& index, std::string_view option) -> std::string { + if (++index >= argc) { + throw std::invalid_argument("Missing value for " + std::string{option}); + } + return argv[index]; + }; + + for (int index = 1; index < argc; ++index) { + auto const option = std::string_view{argv[index]}; + if (option == "-h" or option == "--help") { + print_usage(); + std::exit(0); + } else if (option == "--dataset-root") { + args.dataset_root = require_value(index, option); + dataset_root_set = true; + } else if (option == "--refs") { + args.refs = require_value(index, option); + } else if (option == "--image-parquet-dir") { + args.image_dir = require_value(index, option); + image_dir_set = true; + } else if (option == "--limit") { + auto const value = std::stoll(require_value(index, option)); + CUDF_EXPECTS(value >= 0 and value <= std::numeric_limits::max(), + "Invalid --limit"); + args.limit = static_cast(value); + } else if (option == "--pass-read-limit-mib") { + auto const value = std::stoull(require_value(index, option)); + CUDF_EXPECTS(value <= std::numeric_limits::max() / (1024 * 1024), + "Invalid --pass-read-limit-mib"); + args.pass_read_limit = value * 1024 * 1024; + args.pass_read_limit_set = true; + } else if (option == "--use-data-page-mask") { + args.use_page_mask = get_boolean(require_value(index, option)); + } else if (option == "--use-sparse-page-io") { + args.use_sparse_page_io = get_boolean(require_value(index, option)); + } else if (option == "--best-effort-drop-cache") { + args.drop_cache = true; + } else if (option == "--validate") { + args.validate = true; + } else if (option == "--expected-output") { + args.expected_output = require_value(index, option); + args.validate = true; + expected_output_set = true; + } else if (option == "--output") { + args.output = require_value(index, option); + } else { + throw std::invalid_argument("Unknown option: " + std::string{option}); + } + } + if (not image_dir_set) { + args.image_dir = + args.dataset_root / "parquet/mint_1t_html_images_stable_row_ids/data"; + } + if (dataset_root_set and not expected_output_set) { + auto const output_dir = + args.dataset_root == std::filesystem::path{original_dataset_root} + ? std::filesystem::path{writable_root} / "original" + : std::filesystem::path{writable_root} / "output"; + args.expected_output = output_dir / "image_url_payloads.parquet"; + } + return args; +} + +void require_writable_output_path(std::filesystem::path const& path) +{ + auto const output = std::filesystem::absolute(path).lexically_normal(); + auto const root = std::filesystem::path{writable_root}.lexically_normal(); + auto const rel = output.lexically_relative(root); + CUDF_EXPECTS(not rel.empty() and *rel.begin() != "..", + "Output paths must be under " + root.string()); +} + +std::vector refs_to_host(cudf::table_view refs, rmm::cuda_stream_view stream) +{ + CUDF_EXPECTS(refs.num_columns() == 2, "Unexpected reference table column count"); + CUDF_EXPECTS(refs.column(0).null_count() == 0 and refs.column(1).null_count() == 0, + "Null reference fields are not supported"); + + auto const strings = cudf::strings_column_view{refs.column(0)}; + auto host_offsets = std::vector(strings.size() + 1); + CUDF_CUDA_TRY(cudaMemcpyAsync(host_offsets.data(), + strings.offsets().data() + strings.offset(), + host_offsets.size() * sizeof(cudf::size_type), + cudaMemcpyDeviceToHost, + stream.value())); + stream.synchronize(); + + auto const first_char = host_offsets.front(); + auto const chars_size = host_offsets.back() - first_char; + auto host_chars = std::vector(chars_size); + CUDF_CUDA_TRY(cudaMemcpyAsync(host_chars.data(), + strings.chars_begin(stream) + first_char, + host_chars.size(), + cudaMemcpyDeviceToHost, + stream.value())); + + auto host_row_offsets = std::vector(refs.num_rows()); + CUDF_CUDA_TRY(cudaMemcpyAsync(host_row_offsets.data(), + refs.column(1).data() + refs.column(1).offset(), + host_row_offsets.size() * sizeof(uint32_t), + cudaMemcpyDeviceToHost, + stream.value())); + stream.synchronize(); + + auto result = std::vector{}; + result.reserve(refs.num_rows()); + for (cudf::size_type row = 0; row < refs.num_rows(); ++row) { + auto const begin = host_offsets[row] - first_char; + auto const end = host_offsets[row + 1] - first_char; + result.push_back({std::string{host_chars.data() + begin, static_cast(end - begin)}, + host_row_offsets[row]}); + } + return result; +} + +cache_drop_result drop_file_cache(std::vector const& paths) +{ + cache_drop_result result{.attempted = true}; + for (auto const& path : paths) { + auto const fd = ::open(path.c_str(), O_RDONLY); + if (fd < 0) { + ++result.errors; + continue; + } + if (::posix_fadvise(fd, 0, 0, POSIX_FADV_DONTNEED) == 0) { + ++result.files; + } else { + ++result.errors; + } + ::close(fd); + } + return result; +} + +std::unique_ptr make_device_column(std::span host_data, + cudf::data_type type, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + auto result = + cudf::make_numeric_column(type, host_data.size(), cudf::mask_state::UNALLOCATED, stream, mr); + CUDF_CUDA_TRY(cudaMemcpyAsync(result->mutable_view().data(), + host_data.data(), + host_data.size(), + cudaMemcpyHostToDevice, + stream.value())); + stream.synchronize(); + return result; +} + +std::unique_ptr make_gather_map(std::span host_data, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + auto result = cudf::make_numeric_column(cudf::data_type{cudf::type_id::INT32}, + host_data.size(), + cudf::mask_state::UNALLOCATED, + stream, + mr); + CUDF_CUDA_TRY(cudaMemcpyAsync(result->mutable_view().data(), + host_data.data(), + host_data.size_bytes(), + cudaMemcpyHostToDevice, + stream.value())); + stream.synchronize(); + return result; +} + +uint64_t payload_bytes(cudf::table_view table, rmm::cuda_stream_view stream) +{ + CUDF_EXPECTS(table.num_columns() == static_cast(payload_columns.size()), + "Unexpected payload table schema"); + auto const& sizes = table.column(3); + auto host_sizes = std::vector(table.num_rows()); + CUDF_CUDA_TRY(cudaMemcpyAsync(host_sizes.data(), + sizes.data() + sizes.offset(), + host_sizes.size() * sizeof(int64_t), + cudaMemcpyDeviceToHost, + stream.value())); + stream.synchronize(); + return std::accumulate(host_sizes.begin(), host_sizes.end(), uint64_t{0}); +} + +void write_output(std::filesystem::path const& path, cudf::table_view table) +{ + require_writable_output_path(path); + std::filesystem::create_directories(path.parent_path()); + auto metadata = cudf::io::table_input_metadata{table}; + for (std::size_t index = 0; index < payload_columns.size(); ++index) { + metadata.column_metadata[index].set_name(payload_columns[index]); + } + metadata.column_metadata.front().set_output_as_binary(true); + auto options = + cudf::io::parquet_writer_options::builder(cudf::io::sink_info{path.string()}, table) + .metadata(std::move(metadata)) + .compression(cudf::io::compression_type::ZSTD) + .stats_level(cudf::io::statistics_freq::STATISTICS_COLUMN) + .build(); + cudf::io::write_parquet(options); +} + +std::unique_ptr read_expected(std::filesystem::path const& path) +{ + auto options = cudf::io::parquet_reader_options::builder(cudf::io::source_info{path.string()}) + .column_names(payload_columns) + .build(); + return std::move(cudf::io::read_parquet(options).tbl); +} + +void print_json(timing_data const& timings, + cache_drop_result const& cache, + bool use_page_mask, + bool use_sparse_page_io, + std::size_t files, + std::size_t row_groups, + cudf::size_type rows, + uint64_t selected_payload_bytes, + uint64_t requested_bytes, + std::size_t output_allocated_bytes, + double benchmark_elapsed, + double end_to_end_elapsed) +{ + auto const payload_mib = selected_payload_bytes / double{1024 * 1024}; + auto const fetch_time = timings.get("payload byte-range fetch"); + auto const decode_time = timings.get("payload materialization/decode"); + + std::cout << "\n{" + << R"("method":"hybrid_scan_multifile_payload_mask",)" + << "\"use_data_page_mask\":" << std::boolalpha << use_page_mask << "," + << "\"use_sparse_page_io\":" << use_sparse_page_io << "," + << "\"elapsed_s\":" << benchmark_elapsed << "," + << "\"end_to_end_elapsed_s\":" << end_to_end_elapsed << "," + << "\"rows\":" << rows << "," + << "\"payload_bytes\":" << selected_payload_bytes << "," + << "\"payload_mib\":" << payload_mib << "," + << "\"rows_per_s\":" << (benchmark_elapsed > 0 ? rows / benchmark_elapsed : 0) << "," + << "\"payload_mib_per_s\":" + << (benchmark_elapsed > 0 ? payload_mib / benchmark_elapsed : 0) << "," + << "\"referenced_files\":" << files << "," + << "\"row_groups_read\":" << row_groups << "," + << "\"compressed_payload_bytes_requested\":" << requested_bytes << "," + << "\"output_allocated_bytes\":" << output_allocated_bytes << "," + << "\"payload_fetch_mib_per_s\":" + << (fetch_time > 0 ? requested_bytes / double{1024 * 1024} / fetch_time : 0) << "," + << "\"materialization_mib_per_s\":" + << (decode_time > 0 ? output_allocated_bytes / double{1024 * 1024} / decode_time : 0) + << "," + << R"("best_effort_cache_drop":{"attempted":)" << cache.attempted + << ",\"files\":" << cache.files << ",\"errors\":" << cache.errors << "}," + << "\"stage_seconds\":{"; + for (std::size_t index = 0; index < timings.stages.size(); ++index) { + if (index != 0) { std::cout << ","; } + std::cout << "\"" << timings.stages[index].first << "\":" << timings.stages[index].second; + } + std::cout << "}}\n"; +} + +} // namespace + +int main(int argc, char const** argv) +{ + auto const program_start = clock_type::now(); + auto timings = timing_data{}; + auto args = parse_args(argc, argv); + if (not args.pass_read_limit_set) { + auto free_bytes = std::size_t{}; + auto total_bytes = std::size_t{}; + CUDF_CUDA_TRY(cudaMemGetInfo(&free_bytes, &total_bytes)); + args.pass_read_limit = total_bytes + total_bytes / 20; + } + auto const stream = cudf::get_default_stream(); + auto resource = create_memory_resource(false); + auto stats_mr = rmm::mr::statistics_resource_adaptor{resource}; + rmm::mr::set_current_device_resource(stats_mr); + + auto start = clock_type::now(); + auto refs_options = + cudf::io::parquet_reader_options::builder(cudf::io::source_info{args.refs.string()}) + .column_names({"image_parquet", "_row_offset"}) + .build(); + if (args.limit.has_value()) { refs_options.set_num_rows(args.limit.value()); } + auto refs_table = std::move(cudf::io::read_parquet(refs_options, stream).tbl); + stream.synchronize(); + timings.add("reference parquet read", start); + + start = clock_type::now(); + auto refs = refs_to_host(refs_table->view(), stream); + std::stable_sort(refs.begin(), refs.end()); + CUDF_EXPECTS(not refs.empty(), "No references to read"); + + auto unique_refs = std::vector{}; + auto gather_map = std::vector{}; + unique_refs.reserve(refs.size()); + gather_map.reserve(refs.size()); + for (auto const& ref : refs) { + if (unique_refs.empty() or not(unique_refs.back() == ref)) { unique_refs.push_back(ref); } + gather_map.push_back(static_cast(unique_refs.size() - 1)); + } + + auto source_names = std::vector{}; + auto source_offsets = std::vector>{}; + for (auto const& ref : unique_refs) { + if (source_names.empty() or source_names.back() != ref.image_parquet) { + source_names.push_back(ref.image_parquet); + source_offsets.emplace_back(); + } + source_offsets.back().push_back(ref.row_offset); + } + auto source_paths = std::vector{}; + auto source_files = std::vector{}; + source_paths.reserve(source_names.size()); + source_files.reserve(source_names.size()); + for (auto const& name : source_names) { + auto const path = args.image_dir / name; + CUDF_EXPECTS(std::filesystem::is_regular_file(path), "Missing source file: " + path.string()); + source_paths.push_back(path); + source_files.push_back(path.string()); + } + timings.add("reference host extraction/sort/group", start); + + auto cache_result = cache_drop_result{}; + start = clock_type::now(); + if (args.drop_cache) { + auto cache_paths = source_paths; + cache_paths.insert(cache_paths.begin(), args.refs); + cache_result = drop_file_cache(cache_paths); + } + timings.add("best-effort cache drop", start); + + auto const benchmark_start = clock_type::now(); + + start = clock_type::now(); + auto inputs = multifile_inputs{cudf::io::source_info{source_files}}; + timings.add("datasource construction", start); + + start = clock_type::now(); + inputs.fetch_footers(); + timings.add("footer fetch", start); + + auto options = cudf::io::parquet_reader_options::builder().column_names(payload_columns).build(); + start = clock_type::now(); + auto reader = hybrid_scan_multifile{inputs.footer_byte_spans, options}; + timings.add("hybrid reader construction", start); + + start = clock_type::now(); + auto const metadatas = reader.parquet_metadatas(); + auto row_groups = std::vector>(source_names.size()); + auto host_row_mask = std::vector{}; + std::size_t selected_row_groups{}; + for (std::size_t source = 0; source < source_names.size(); ++source) { + auto const& metadata = metadatas[source]; + auto const& offsets = source_offsets[source]; + auto offset_index = std::size_t{0}; + auto file_row_start = uint64_t{0}; + for (std::size_t rg = 0; rg < metadata.row_groups.size() and offset_index < offsets.size(); + ++rg) { + auto const rows = metadata.row_groups[rg].num_rows; + CUDF_EXPECTS(rows >= 0, "Negative row count in Parquet metadata"); + auto const rows_unsigned = static_cast(rows); + auto const file_row_stop = file_row_start + rows_unsigned; + if (offsets[offset_index] < file_row_start) { + throw std::logic_error("References are not ordered"); + } + if (offsets[offset_index] < file_row_stop) { + row_groups[source].push_back(static_cast(rg)); + ++selected_row_groups; + auto const mask_start = host_row_mask.size(); + host_row_mask.resize(mask_start + static_cast(rows_unsigned), uint8_t{0}); + while (offset_index < offsets.size() and offsets[offset_index] < file_row_stop) { + auto const local_offset = + static_cast(offsets[offset_index] - file_row_start); + host_row_mask[mask_start + local_offset] = uint8_t{1}; + ++offset_index; + } + } + file_row_start = file_row_stop; + } + CUDF_EXPECTS(offset_index == offsets.size(), + "Out-of-range row offset in " + source_names[source]); + } + CUDF_EXPECTS( + host_row_mask.size() <= static_cast(std::numeric_limits::max()), + "Selected row groups exceed cudf column size limit"); + timings.add("row-group planning/host mask", start); + + start = clock_type::now(); + auto page_ranges = reader.page_index_byte_ranges(); + auto missing_page = std::find_if( + page_ranges.begin(), page_ranges.end(), [](auto const& range) { return range.is_empty(); }); + CUDF_EXPECTS(missing_page == page_ranges.end(), "A referenced source has no Parquet page index"); + timings.add("page-index range planning", start); + + start = clock_type::now(); + auto page_buffers = + cudf::io::parquet::fetch_page_indexes_to_host(inputs.datasource_refs, page_ranges); + auto page_spans = std::vector>{}; + page_spans.reserve(page_buffers.size()); + std::transform(page_buffers.begin(), + page_buffers.end(), + std::back_inserter(page_spans), + [](auto const& buffer) { return cudf::host_span{*buffer}; }); + timings.add("page-index fetch", start); + + start = clock_type::now(); + reader.setup_page_indexes(page_spans); + timings.add("page-index setup", start); + + start = clock_type::now(); + auto const indexed_metadatas = reader.parquet_metadatas(); + for (std::size_t source = 0; source < row_groups.size(); ++source) { + for (auto const rg : row_groups[source]) { + auto const& chunks = indexed_metadatas[source].row_groups[rg].columns; + for (auto const& name : payload_columns) { + auto const chunk = std::find_if(chunks.begin(), chunks.end(), [&](auto const& candidate) { + auto const& path = candidate.meta_data.path_in_schema; + return not path.empty() and path.back() == name; + }); + CUDF_EXPECTS(chunk != chunks.end(), + "Missing payload column " + name + " in " + source_names[source]); + CUDF_EXPECTS(chunk->offset_index.has_value(), + "Missing offset index for " + source_names[source] + " row group " + + std::to_string(rg) + " column " + name); + } + } + } + timings.add("selected index validation", start); + + start = clock_type::now(); + auto row_mask = + make_device_column(host_row_mask, cudf::data_type{cudf::type_id::BOOL8}, stream, stats_mr); + timings.add("row-mask host-to-device", start); + + start = clock_type::now(); + auto const passes = reader.construct_row_group_passes(row_groups, args.pass_read_limit); + timings.add("row-group pass construction", start); + + auto requested_bytes = uint64_t{0}; + auto selected_payload_bytes = uint64_t{0}; + auto output_allocated_bytes = std::size_t{0}; + auto global_mask_start = cudf::size_type{0}; + auto pass_tables = std::vector>{}; + auto const retain_output = + args.validate or args.output.has_value() or refs.size() != unique_refs.size(); + + auto range_planning_seconds = double{0}; + auto range_fetch_seconds = double{0}; + auto chunking_setup_seconds = double{0}; + auto materialize_seconds = double{0}; + auto stat_seconds = double{0}; + + std::cout << "PROGRESS constructed " << passes.size() << " multifile passes (limit " + << args.pass_read_limit / double{1024 * 1024} << " MiB)\n" + << std::flush; + for (std::size_t pass_index = 0; pass_index < passes.size(); ++pass_index) { + auto const& pass = passes[pass_index]; + auto const pass_start = clock_type::now(); + auto const pass_row_groups = + std::accumulate(pass.begin(), + pass.end(), + std::size_t{0}, + [](auto count, auto const& source_row_groups) { + return count + source_row_groups.size(); + }); + auto const pass_rows = reader.total_rows_in_row_groups(pass); + auto const pass_mask = cudf::column_view( + row_mask->type(), pass_rows, row_mask->view().data(), nullptr, 0, global_mask_start); + std::cout << "PROGRESS pass " << pass_index + 1 << "/" << passes.size() << " started (" + << pass_row_groups << " row groups, " << pass_rows << " input rows)\n" + << std::flush; + + auto payload_data = multisource_device_data{}; + if (args.use_page_mask and args.use_sparse_page_io) { + start = clock_type::now(); + auto const payload_ranges = reader.payload_column_chunks_byte_ranges( + pass, pass_mask, use_data_page_mask::YES, options, stream); + for (auto const& source_ranges : payload_ranges) { + requested_bytes = + std::accumulate(source_ranges.begin(), + source_ranges.end(), + requested_bytes, + [](uint64_t sum, auto const& range) { return sum + range.size(); }); + } + range_planning_seconds += std::chrono::duration(clock_type::now() - start).count(); + + start = clock_type::now(); + payload_data = fetch_multisource_device_data(inputs, payload_ranges, stream, stats_mr); + range_fetch_seconds += std::chrono::duration(clock_type::now() - start).count(); + + start = clock_type::now(); + reader.setup_chunking_for_payload_columns(0, + 0, + pass, + pass_mask, + use_data_page_mask::YES, + payload_data.per_source_spans, + options, + stream, + stats_mr); + } else { + start = clock_type::now(); + auto payload_ranges = reader.payload_column_chunks_byte_ranges(pass, options); + requested_bytes += + std::accumulate(payload_ranges.first.begin(), + payload_ranges.first.end(), + uint64_t{0}, + [](uint64_t sum, auto const& range) { return sum + range.size(); }); + range_planning_seconds += std::chrono::duration(clock_type::now() - start).count(); + + start = clock_type::now(); + payload_data = fetch_multisource_device_data(inputs, payload_ranges, stream, stats_mr); + range_fetch_seconds += std::chrono::duration(clock_type::now() - start).count(); + + start = clock_type::now(); + reader.setup_chunking_for_payload_columns(0, + 0, + pass, + pass_mask, + args.use_page_mask ? use_data_page_mask::YES + : use_data_page_mask::NO, + payload_data.flat_spans, + options, + stream, + stats_mr); + } + stream.synchronize(); + chunking_setup_seconds += std::chrono::duration(clock_type::now() - start).count(); + + auto current_tables = std::vector>{}; + start = clock_type::now(); + while (reader.has_next_table_chunk()) { + current_tables.push_back(reader.materialize_payload_columns_chunk(pass_mask).tbl); + } + stream.synchronize(); + materialize_seconds += std::chrono::duration(clock_type::now() - start).count(); + + if (retain_output) { + std::move(current_tables.begin(), current_tables.end(), std::back_inserter(pass_tables)); + } else { + start = clock_type::now(); + for (auto const& table : current_tables) { + selected_payload_bytes += payload_bytes(table->view(), stream); + output_allocated_bytes += table->alloc_size(); + } + stream.synchronize(); + stat_seconds += std::chrono::duration(clock_type::now() - start).count(); + } + global_mask_start += pass_rows; + auto const pass_seconds = + std::chrono::duration(clock_type::now() - pass_start).count(); + std::cout << "PROGRESS pass " << pass_index + 1 << "/" << passes.size() << " completed in " + << std::fixed << std::setprecision(3) << pass_seconds << " s\n" + << std::flush; + } + CUDF_EXPECTS(std::cmp_equal(global_mask_start, host_row_mask.size()), + "Row-group passes do not span the global row mask"); + + timings.add_seconds("payload byte-range planning", range_planning_seconds); + timings.add_seconds("payload byte-range fetch", range_fetch_seconds); + timings.add_seconds("payload chunking setup", chunking_setup_seconds); + timings.add_seconds("payload materialization/decode", materialize_seconds); + + start = clock_type::now(); + auto output = [&]() -> std::unique_ptr { + if (not retain_output) { return nullptr; } + CUDF_EXPECTS(not pass_tables.empty(), "Payload materialization produced no table chunks"); + if (pass_tables.size() == 1) { return std::move(pass_tables.front()); } + auto views = std::vector{}; + views.reserve(pass_tables.size()); + std::transform( + pass_tables.begin(), pass_tables.end(), std::back_inserter(views), [](auto const& table) { + return table->view(); + }); + return cudf::concatenate(views, stream, stats_mr); + }(); + stream.synchronize(); + timings.add("payload pass concatenation", start); + + start = clock_type::now(); + if (refs.size() != unique_refs.size()) { + CUDF_EXPECTS(output != nullptr, "Duplicate references require a retained output table"); + auto device_gather_map = make_gather_map(gather_map, stream, stats_mr); + output = cudf::gather(output->view(), + device_gather_map->view(), + cudf::out_of_bounds_policy::DONT_CHECK, + cudf::negative_index_policy::NOT_ALLOWED, + stream, + stats_mr); + stream.synchronize(); + } + timings.add("duplicate-order gather", start); + + if (retain_output) { + start = clock_type::now(); + selected_payload_bytes = payload_bytes(output->view(), stream); + output_allocated_bytes = output->alloc_size(); + stream.synchronize(); + stat_seconds = std::chrono::duration(clock_type::now() - start).count(); + } + timings.add_seconds("payload byte/stat extraction", stat_seconds); + + start = clock_type::now(); + std::unique_ptr expected; + if (args.validate) { + CUDF_EXPECTS(std::filesystem::is_regular_file(args.expected_output), + "Expected output does not exist: " + args.expected_output.string()); + expected = read_expected(args.expected_output); + stream.synchronize(); + } + timings.add("expected-output read", start); + + start = clock_type::now(); + if (args.validate) { + auto const equal = + cudf::tables_equal(output->view(), expected->view(), cudf::null_equality::EQUAL, stream); + stream.synchronize(); + CUDF_EXPECTS(equal, "Hybrid Scan output differs from Python output"); + } + timings.add("table comparison", start); + + auto const benchmark_elapsed = + std::chrono::duration(clock_type::now() - benchmark_start).count(); + + start = clock_type::now(); + if (args.output.has_value()) { write_output(args.output.value(), output->view()); } + stream.synchronize(); + timings.add("output parquet write", start); + + start = clock_type::now(); + expected.reset(); + output.reset(); + pass_tables.clear(); + row_mask.reset(); + stream.synchronize(); + timings.add("result cleanup", start); + + auto const end_to_end_elapsed = + std::chrono::duration(clock_type::now() - program_start).count(); + timings.print(); + std::cout << "\nSummary:\n" + << " setup + payload benchmark: " << benchmark_elapsed << " s\n" + << " end-to-end: " << end_to_end_elapsed << " s\n" + << " referenced files: " << source_names.size() << "\n" + << " selected row groups: " << selected_row_groups << "\n" + << " selected rows: " << refs.size() << "\n" + << " selected image bytes: " << selected_payload_bytes << "\n" + << " requested compressed bytes: " << requested_bytes << "\n" + << " use data page mask: " << std::boolalpha << args.use_page_mask << "\n" + << " use sparse page I/O: " + << (args.use_page_mask and args.use_sparse_page_io) << "\n"; + print_json(timings, + cache_result, + args.use_page_mask, + args.use_page_mask and args.use_sparse_page_io, + source_names.size(), + selected_row_groups, + static_cast(refs.size()), + selected_payload_bytes, + requested_bytes, + output_allocated_bytes, + benchmark_elapsed, + end_to_end_elapsed); + std::cout << "Peak device memory: " << stats_mr.get_bytes_counter().peak / double{1024 * 1024} + << " MiB\n"; + return 0; +} From 512c354eaa58ca2d0d5ad667025105fa3060a635 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb Date: Wed, 22 Jul 2026 04:40:14 +0000 Subject: [PATCH 06/16] Update tests and add a new composer --- .../cudf/io/experimental/hybrid_scan.hpp | 48 ++++++++ .../io/parquet/experimental/hybrid_scan.cpp | 45 ++++++++ .../io/experimental/hybrid_scan_composer.cpp | 104 ++++++++++++++++++ .../io/experimental/hybrid_scan_composer.hpp | 25 +++++ .../io/experimental/hybrid_scan_test.cpp | 2 +- 5 files changed, 223 insertions(+), 1 deletion(-) diff --git a/cpp/include/cudf/io/experimental/hybrid_scan.hpp b/cpp/include/cudf/io/experimental/hybrid_scan.hpp index 6a3b2059b55d..6462150841b3 100644 --- a/cpp/include/cudf/io/experimental/hybrid_scan.hpp +++ b/cpp/include/cudf/io/experimental/hybrid_scan.hpp @@ -510,6 +510,26 @@ class hybrid_scan_reader { [[nodiscard]] std::vector payload_column_chunks_byte_ranges( std::span row_group_indices, parquet_reader_options const& options) const; + /** + * @brief Plan page-level payload byte ranges using the row mask + * + * When page masking cannot be used, returns the full payload column-chunk ranges. The resulting + * plan must be consumed exactly once by the matching page-data setup overload. + * + * @param row_group_indices Input row group indices + * @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 + */ + [[nodiscard]] std::vector payload_column_chunks_byte_ranges( + std::span 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 * @@ -620,6 +640,34 @@ class hybrid_scan_reader { rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) const; + /** + * @brief Setup chunking for payload columns using page-level payload data + * + * The `page_data` must correspond to the byte ranges returned by the matching page-level + * `payload_column_chunks_byte_ranges` overload. + * + * @param chunk_read_limit Limit on total number of bytes to be returned per table chunk + * @param pass_read_limit Limit on the memory used for reading and decompressing data + * @param row_group_indices Input row group indices + * @param row_mask Boolean column indicating which rows need to be read + * @param mask_data_pages Whether to use the row mask to prune data pages + * @param page_data_per_source Device spans of page data returned by the page-level range plan, + * grouped by source + * @param options Parquet reader options + * @param stream CUDA stream used for device memory operations and kernel launches + * @param mr Device memory resource used to allocate the output table chunks + */ + void setup_chunking_for_payload_columns( + std::size_t chunk_read_limit, + std::size_t pass_read_limit, + std::span 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.cpp b/cpp/src/io/parquet/experimental/hybrid_scan.cpp index 36bc5de06dc4..dfa1a5fd29a1 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan.cpp +++ b/cpp/src/io/parquet/experimental/hybrid_scan.cpp @@ -217,6 +217,24 @@ hybrid_scan_reader::payload_column_chunks_byte_ranges(std::span return _impl->payload_column_chunks_byte_ranges(input_row_group_indices, options).first; } +std::vector hybrid_scan_reader::payload_column_chunks_byte_ranges( + std::span 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(); + + auto const input_row_group_indices = + std::vector>{{row_group_indices.begin(), row_group_indices.end()}}; + + return _impl + ->payload_column_chunks_byte_ranges( + input_row_group_indices, row_mask, mask_data_pages, options, stream) + .front(); +} + table_with_metadata hybrid_scan_reader::materialize_payload_columns( std::span row_group_indices, std::span const> column_chunk_data, @@ -328,6 +346,33 @@ void hybrid_scan_reader::setup_chunking_for_payload_columns( mr); } +void hybrid_scan_reader::setup_chunking_for_payload_columns( + std::size_t chunk_read_limit, + std::size_t pass_read_limit, + std::span 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(); + + auto const input_row_group_indices = + std::vector>{{row_group_indices.begin(), row_group_indices.end()}}; + + _impl->setup_chunking_for_payload_columns(chunk_read_limit, + pass_read_limit, + input_row_group_indices, + row_mask, + mask_data_pages, + page_data_per_source, + options, + stream, + mr); +} + table_with_metadata hybrid_scan_reader::materialize_payload_columns_chunk( cudf::column_view const& row_mask) const { diff --git a/cpp/tests/io/experimental/hybrid_scan_composer.cpp b/cpp/tests/io/experimental/hybrid_scan_composer.cpp index 57a6f6cb7e3f..752f47132080 100644 --- a/cpp/tests/io/experimental/hybrid_scan_composer.cpp +++ b/cpp/tests/io/experimental/hybrid_scan_composer.cpp @@ -319,6 +319,110 @@ std::tuple, std::unique_ptr> chunked_h return std::tuple{std::move(filter_table), std::move(payload_table)}; } +std::tuple, std::unique_ptr> sparse_chunked_hybrid_scan( + cudf::io::datasource& datasource, + 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, + rmm::mr::aligned_resource_adaptor& aligned_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 const reader = setup_reader(datasource, options); + auto reader_ref = std::ref(*reader); + auto const filtered_row_group_indices = + apply_hybrid_scan_filters(datasource, reader_ref, options, stream, mr); + auto const current_row_group_indices = + cudf::host_span(filtered_row_group_indices); + auto row_mask = + options.get_filter().has_value() + ? reader->build_row_mask_with_page_index_stats(current_row_group_indices, options, stream, mr) + : reader->build_all_true_row_mask(current_row_group_indices, stream, mr); + + auto filter_tables = std::vector>{}; + auto payload_tables = std::vector>{}; + std::size_t rows_materialized = 0; + auto const materialize_pass = [&](cudf::host_span row_group_indices) { + auto const rows_in_pass = reader->total_rows_in_row_groups(row_group_indices); + auto* null_mask = row_mask->nullable() ? row_mask->mutable_view().null_mask() : nullptr; + auto const slice_null_count = + cudf::null_count(null_mask, rows_materialized, rows_materialized + rows_in_pass, stream); + auto row_mask_view = cudf::mutable_column_view(row_mask->type(), + rows_in_pass, + row_mask->mutable_view().data(), + null_mask, + slice_null_count, + rows_materialized); + + auto const filter_byte_ranges = + reader->filter_column_chunks_byte_ranges(row_group_indices, options); + auto [filter_buffers, filter_data, filter_tasks] = + cudf::io::parquet::fetch_byte_ranges_to_device_async( + datasource, filter_byte_ranges, stream, mr); + filter_tasks.get(); + reader->setup_chunking_for_filter_columns( + 1024, + 10240, + row_group_indices, + row_mask_view, + cudf::io::parquet::experimental::use_data_page_mask::YES, + filter_data, + 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_group_indices, + row_mask_view, + cudf::io::parquet::experimental::use_data_page_mask::YES, + options, + stream); + auto [payload_buffers, payload_data, payload_tasks] = + cudf::io::parquet::fetch_byte_ranges_to_device_async( + datasource, payload_page_ranges, stream, mr); + payload_tasks.get(); + auto const page_data_per_source = + std::vector>>{{payload_data.begin(), + payload_data.end()}}; + reader->setup_chunking_for_payload_columns( + 1024, + 10240, + row_group_indices, + row_mask_view, + cudf::io::parquet::experimental::use_data_page_mask::YES, + page_data_per_source, + options, + stream, + mr); + while (reader->has_next_table_chunk()) { + payload_tables.push_back(reader->materialize_payload_columns_chunk(row_mask_view).tbl); + } + + rows_materialized += rows_in_pass; + }; + + if (current_row_group_indices.size() > 1) { + auto const row_group_split = current_row_group_indices.size() / 2; + materialize_pass(current_row_group_indices.subspan(0, row_group_split)); + materialize_pass(current_row_group_indices.subspan( + row_group_split, current_row_group_indices.size() - row_group_split)); + } else { + materialize_pass(current_row_group_indices); + } + + return std::tuple{concatenate_tables(std::move(filter_tables), stream, mr), + concatenate_tables(std::move(payload_tables), stream, mr)}; +} + std::unique_ptr hybrid_scan_single_step( cudf::io::datasource& datasource, cudf::ast::operation const& filter_expression, diff --git a/cpp/tests/io/experimental/hybrid_scan_composer.hpp b/cpp/tests/io/experimental/hybrid_scan_composer.hpp index 80519cdd3130..31eece7d1e29 100644 --- a/cpp/tests/io/experimental/hybrid_scan_composer.hpp +++ b/cpp/tests/io/experimental/hybrid_scan_composer.hpp @@ -58,6 +58,31 @@ std::tuple, std::unique_ptr> chunked_h rmm::device_async_resource_ref mr, rmm::mr::aligned_resource_adaptor& aligned_mr); +/** + * @brief Read parquet file with chunked hybrid scan and sparse page-level payload I/O + * + * Filter columns use full-column-chunk I/O. Payload ranges are planned only after filter + * materialization updates the row mask, then only retained data pages are fetched. + * + * @param datasource Input datasource + * @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 + * @param aligned_mr Device memory resource to allocate aligned memory for bloom filters + * + * @return Tuple of filter and payload tables + */ +std::tuple, std::unique_ptr> sparse_chunked_hybrid_scan( + cudf::io::datasource& datasource, + 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, + rmm::mr::aligned_resource_adaptor& aligned_mr); + /** * @brief Read parquet file with the hybrid scan reader in a single step * diff --git a/cpp/tests/io/experimental/hybrid_scan_test.cpp b/cpp/tests/io/experimental/hybrid_scan_test.cpp index edf43fc6bdff..5c316f2d3d25 100644 --- a/cpp/tests/io/experimental/hybrid_scan_test.cpp +++ b/cpp/tests/io/experimental/hybrid_scan_test.cpp @@ -509,7 +509,7 @@ TEST_F(HybridScanTest, ConsecutivePrunedPageOffsets) std::unique_ptr filter_table; std::unique_ptr payload_table; ASSERT_NO_THROW(std::tie(filter_table, payload_table) = - hybrid_scan(*datasource, filter, {}, true, stream, mr, aligned_mr)); + sparse_chunked_hybrid_scan(*datasource, filter, {}, true, stream, mr, aligned_mr)); CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected->view().select({0}), filter_table->view()); CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected->view().select({1, 2, 3, 4}), From 0265b76a0a063a61005e0b262fb7dbfdda033e3a Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb Date: Wed, 22 Jul 2026 19:16:21 +0000 Subject: [PATCH 07/16] Simplify pruned page handling --- cpp/src/io/parquet/decode_fixed.cu | 31 ++--------- cpp/src/io/parquet/decode_pruned_pages.cu | 2 - cpp/src/io/parquet/page_data.cu | 63 +++-------------------- cpp/src/io/parquet/page_decode.cuh | 32 ------------ cpp/src/io/parquet/page_delta_decode.cu | 61 +++------------------- cpp/src/io/parquet/page_string_utils.cuh | 50 ------------------ 6 files changed, 16 insertions(+), 223 deletions(-) diff --git a/cpp/src/io/parquet/decode_fixed.cu b/cpp/src/io/parquet/decode_fixed.cu index 45a7aeb4ec6d..5e3122d9fa37 100644 --- a/cpp/src/io/parquet/decode_fixed.cu +++ b/cpp/src/io/parquet/decode_fixed.cu @@ -1015,20 +1015,12 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size_t, 8) if (!(BitAnd(pages[page_idx].kernel_mask, kernel_mask_t))) { return; } + // Exit early if the page is pruned + if (page_mask.size() > 0 and not page_mask[page_idx]) { return; } + // must come after the kernel mask check [[maybe_unused]] null_count_back_copier _{s, t}; - // Exit super early for simple types if the page does not need to be decoded - if constexpr (not has_lists_t and not has_strings_t and not has_nesting_t) { - if (not page_mask[page_idx]) { - pp->num_nulls = pp->nesting[0].batch_size; - pp->num_valids = 0; - // Set s->nesting info = nullptr to bypass `null_count_back_copier` at return - s->nesting_info = nullptr; - return; - } - } - // Setup local page info if (!setup_local_page_info(s, pp, @@ -1040,23 +1032,6 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size_t, 8) return; } - // Write list and/or string offsets and exit if the page does not need to be decoded - if (not page_mask[page_idx]) { - // Update offsets for all list depth levels - if constexpr (has_lists_t) { update_list_offsets_for_pruned_pages(s); } - // Update string offsets or write string sizes for small and large strings respectively - if constexpr (has_strings_t) { - update_string_offsets_for_pruned_pages( - s, initial_str_offsets, pages[page_idx]); - } - // Must be set after computing above list and string offsets - pp->num_nulls = pp->nesting[s->col.max_nesting_depth - 1].batch_size; - if constexpr (not has_lists_t) { pp->num_nulls -= s->first_row; } - pp->num_valids = 0; - - return; - } - bool const process_nulls = should_process_nulls(s); // shared buffer. all shared memory is suballocated out of here diff --git a/cpp/src/io/parquet/decode_pruned_pages.cu b/cpp/src/io/parquet/decode_pruned_pages.cu index 127952de12a2..657a8c366957 100644 --- a/cpp/src/io/parquet/decode_pruned_pages.cu +++ b/cpp/src/io/parquet/decode_pruned_pages.cu @@ -45,7 +45,6 @@ CUDF_KERNEL void __launch_bounds__(block_size) auto const is_list_col = chunk.max_level[level_type::REPETITION] != 0; // Write offsets for pruned non-list (flat) string columns. - // Mirrors `update_string_offsets_for_pruned_pages` in page_string_utils.cuh. if (not is_list_col and is_string_col(chunk)) { auto data = static_cast(chunk.column_data_base[chunk.max_nesting_depth - 1]); if (data == nullptr) { return; } @@ -65,7 +64,6 @@ CUDF_KERNEL void __launch_bounds__(block_size) } // Write offsets to list locations at each depth. - // Mirrors `update_list_offsets_for_pruned_pages` in page_decode.cuh. if (is_list_col and page.nesting != nullptr and page.nesting_decode != nullptr) { for (auto depth = 0; depth < chunk.max_nesting_depth - 1; depth++) { auto offsets = static_cast(chunk.column_data_base[depth]); diff --git a/cpp/src/io/parquet/page_data.cu b/cpp/src/io/parquet/page_data.cu index 5b2f854fce08..b6740d0dee71 100644 --- a/cpp/src/io/parquet/page_data.cu +++ b/cpp/src/io/parquet/page_data.cu @@ -61,6 +61,9 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size) auto const block = cg::this_thread_block(); auto const warp = cg::tiled_partition(block); + // Exit early if the page is pruned + if (not page_mask.empty() and not page_mask[page_idx]) { return; } + [[maybe_unused]] null_count_back_copier _{s, static_cast(block.thread_rank())}; // Setup local page info @@ -78,21 +81,6 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size) bool const has_repetition = s->col.max_level[level_type::REPETITION] > 0; bool const process_nulls = should_process_nulls(s); - // Write list offsets and exit if the page does not need to be decoded - if (not page_mask[page_idx]) { - auto& page = pages[page_idx]; - // Update offsets for all list depth levels - if (has_repetition) { update_list_offsets_for_pruned_pages(s); } - - // Must be set after computing above list offsets - cg::invoke_one(block, [&]() { - page.num_nulls = page.nesting[s->col.max_nesting_depth - 1].batch_size; - page.num_nulls -= has_repetition ? 0 : s->first_row; - page.num_valids = 0; - }); - return; - } - auto const data_len = cuda::std::distance(s->data_start, s->data_end); auto const num_values = data_len / s->dtype_len_in; @@ -280,6 +268,10 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size) auto const block = cg::this_thread_block(); auto const warp = cg::tiled_partition(block); int out_warp_id; + + // Exit early if the page is pruned + if (not page_mask.empty() and not page_mask[page_idx]) { return; } + [[maybe_unused]] null_count_back_copier _{s, static_cast(block.thread_rank())}; // Setup local page info @@ -297,47 +289,6 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size) bool const has_repetition = s->col.max_level[level_type::REPETITION] > 0; bool const process_nulls = should_process_nulls(s); - // Write list offsets and exit if the page does not need to be decoded - if (not page_mask[page_idx]) { - auto& page = pages[page_idx]; - - // Update offsets for all list depth levels - if (has_repetition) { update_list_offsets_for_pruned_pages(s); } - - // Fill offsets with the initial `str_offset` to indicate empty strings for BYTE_ARRAY and - // FIXED_LEN_BYTE_ARRAY types. These types are now decoded by `decode_page_data_generic()` - // anyway so the following code should never be reached. Also note that this decoder does not - // handle large strings either and should eventually be removed. - Type const dtype = s->col.physical_type; - auto const is_decimal = - s->col.logical_type.has_value() and s->col.logical_type->type == LogicalType::DECIMAL; - if (dtype == Type::FIXED_LEN_BYTE_ARRAY or (dtype == Type::BYTE_ARRAY and not is_decimal)) { - // Initial string offset - auto const initial_value = page.str_offset; - - // We must use the batch size from the nesting info (the size of the page for this batch) - auto value_count = page.nesting[s->col.max_nesting_depth - 1].batch_size; - - // If no repetition we haven't calculated start/end bounds and instead just skipped - // values until we reach first_row. account for that here. - if (not has_repetition) { value_count -= s->first_row; } - - auto& ni = s->nesting_info[s->col.max_nesting_depth - 1]; - auto offptr = reinterpret_cast(ni.data_out); - - // Write the initial string offset at all positions to indicate empty strings - for (int idx = block.thread_rank(); idx < value_count; idx += block.size()) { - offptr[idx] = initial_value; - } - } - - page.num_nulls = page.nesting[s->col.max_nesting_depth - 1].batch_size; - page.num_nulls -= has_repetition ? 0 : s->first_row; - page.num_valids = 0; - - return; - } - PageNestingDecodeInfo* nesting_info_base = s->nesting_info; // Capture initial valid_map_offset before any processing that might modify it diff --git a/cpp/src/io/parquet/page_decode.cuh b/cpp/src/io/parquet/page_decode.cuh index ef0a8f9e502d..31a9f4a05d05 100644 --- a/cpp/src/io/parquet/page_decode.cuh +++ b/cpp/src/io/parquet/page_decode.cuh @@ -717,38 +717,6 @@ inline __device__ void get_nesting_bounds(int& start_depth, } } -/** - * @brief Updates nesting level offsets for pruned pages of a list column - * - * This function iterates through the nesting levels of a column and updates the offsets for a list - * column. The offset for the current nesting level equals the length of the next nesting level - * - * @tparam block_size The size of the block used for decoding. - * @param[in,out] state Pointer to page state containing column and nesting information. - */ -template -static __device__ void update_list_offsets_for_pruned_pages(page_state_s* state) -{ - int const max_depth = state->col.max_nesting_depth - 1; - bool const in_nesting_bounds = max_depth >= 0; - auto const tid = cg::this_thread_block().thread_rank(); - - // Iterate by depth and store offset(s) to the list location(s) - for (int depth = 0; depth < max_depth; depth++) { - auto& nesting_info = state->nesting_info[depth]; - // If we're -not- at a leaf column and we're within nesting/row bounds and we have a valid - // data_out pointer, it implies this is a list column, so emit an offset for the current nesting - // level equal to current length of the next nesting level - if (in_nesting_bounds and nesting_info.data_out != nullptr) { - auto const& next_nesting_info = state->nesting_info[depth + 1]; - auto const offset = next_nesting_info.page_start_value; - for (int idx = tid; idx < state->page.nesting[depth].batch_size; idx += block_size) { - (reinterpret_cast(nesting_info.data_out))[idx] = offset; - } - } - } -} - /** * @brief Process a batch of incoming repetition/definition level values and generate * validity, nested column offsets (where appropriate) and decoding indices. diff --git a/cpp/src/io/parquet/page_delta_decode.cu b/cpp/src/io/parquet/page_delta_decode.cu index 85bf383bb43f..7cb5fbd59d26 100644 --- a/cpp/src/io/parquet/page_delta_decode.cu +++ b/cpp/src/io/parquet/page_delta_decode.cu @@ -317,6 +317,10 @@ CUDF_KERNEL void __launch_bounds__(decode_delta_binary_block_size) auto const block = cg::this_thread_block(); auto const warp = cg::tiled_partition(block); auto* const db = &db_state; + + // Exit early if the page is pruned + if (page_mask.size() > 0 and not page_mask[page_idx]) { return; } + [[maybe_unused]] null_count_back_copier _{s, static_cast(block.thread_rank())}; // Setup local page info @@ -337,17 +341,6 @@ CUDF_KERNEL void __launch_bounds__(decode_delta_binary_block_size) // Capture initial valid_map_offset before any processing that might modify it int const init_valid_map_offset = s->nesting_info[s->col.max_nesting_depth - 1].valid_map_offset; - // Write list offsets and exit if the page does not need to be decoded - if (not page_mask[page_idx]) { - auto& page = pages[page_idx]; - // Update offsets for all list depth levels - if (has_repetition) { update_list_offsets_for_pruned_pages(s); } - page.num_nulls = page.nesting[s->col.max_nesting_depth - 1].batch_size; - page.num_nulls -= has_repetition ? 0 : s->first_row; - page.num_valids = 0; - return; - } - // copying logic from gpuDecodePageData. PageNestingDecodeInfo const* nesting_info_base = s->nesting_info; @@ -478,6 +471,7 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size) auto* const prefix_db = &db_state.prefixes; auto* const suffix_db = &db_state.suffixes; auto* const dba = &db_state; + if (page_mask.size() > 0 and not page_mask[page_idx]) { return; } [[maybe_unused]] null_count_back_copier _{s, static_cast(block.thread_rank())}; if (!setup_local_page_info(s, @@ -504,28 +498,6 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size) // Capture initial valid_map_offset before any processing that might modify it int const init_valid_map_offset = s->nesting_info[s->col.max_nesting_depth - 1].valid_map_offset; - // Write list/string offsets and exit if the page does not need to be decoded - if (not page_mask[page_idx]) { - auto page = &pages[page_idx]; - // Update list offsets and string offsets or sizes depending on the large-string property - if (has_repetition) { - // Update list offsets - update_list_offsets_for_pruned_pages(s); - // Update string offsets or sizes - update_string_offsets_for_pruned_pages( - s, initial_str_offsets, pages[page_idx]); - } else { - // Update string offsets or sizes - update_string_offsets_for_pruned_pages( - s, initial_str_offsets, pages[page_idx]); - } - page->num_nulls = page->nesting[s->col.max_nesting_depth - 1].batch_size; - page->num_nulls -= has_repetition ? 0 : s->first_row; - page->num_valids = 0; - - return; - } - // choose a character parallel string copy when the average string is longer than a warp auto const use_char_ll = (s->page.str_bytes / s->page.num_valids) > cudf::detail::warp_size; @@ -704,6 +676,7 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size) auto const block = cg::this_thread_block(); auto const warp = cg::tiled_partition(block); auto* const db = &db_state; + if (page_mask.size() > 0 and not page_mask[page_idx]) { return; } [[maybe_unused]] null_count_back_copier _{s, static_cast(block.thread_rank())}; auto const mask = decode_kernel_mask::DELTA_LENGTH_BA; @@ -731,28 +704,6 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size) // Capture initial valid_map_offset before any processing that might modify it int const init_valid_map_offset = s->nesting_info[s->col.max_nesting_depth - 1].valid_map_offset; - // Write list/string offsets and exit if the page does not need to be decoded - if (not page_mask[page_idx]) { - auto page = &pages[page_idx]; - // Update list offsets and string offsets or sizes depending on the large-string property - if (has_repetition) { - // Update list offsets - update_list_offsets_for_pruned_pages(s); - // Update string offsets or sizes - update_string_offsets_for_pruned_pages( - s, initial_str_offsets, pages[page_idx]); - } else { - // Update string offsets or sizes - update_string_offsets_for_pruned_pages( - s, initial_str_offsets, pages[page_idx]); - } - page->num_nulls = page->nesting[s->col.max_nesting_depth - 1].batch_size; - page->num_nulls -= has_repetition ? 0 : s->first_row; - page->num_valids = 0; - - return; - } - // copying logic from gpuDecodePageData. PageNestingDecodeInfo const* nesting_info_base = s->nesting_info; diff --git a/cpp/src/io/parquet/page_string_utils.cuh b/cpp/src/io/parquet/page_string_utils.cuh index 084d9334602f..c846db4c65d2 100644 --- a/cpp/src/io/parquet/page_string_utils.cuh +++ b/cpp/src/io/parquet/page_string_utils.cuh @@ -138,56 +138,6 @@ inline __device__ void compute_initial_large_strings_offset(page_state_s const* } } -/** - * @brief Update offsets with either zeros if this is a large string column, `initial_value` - * otherwise. - * - * For large string columns, fill zeros (sizes) at all offsets and atomically update the initial - * string offset. Otherwise, fill `initial_value` at all offsets. - * - * @tparam block_size Thread block size - * @tparam has_lists Whether the column is a list column - * @param[in,out] state page state - * @param[out] initial_str_offsets Initial string offsets - * @param[in] page Page information - */ -template -__device__ void update_string_offsets_for_pruned_pages( - page_state_s* state, cudf::device_span initial_str_offsets, PageInfo const& page) -{ - namespace cg = cooperative_groups; - - // Initial string offset - auto const initial_value = page.str_offset; - // The value count is either the leaf-level batch size in case of lists or the number of - // effective rows being read by this page - auto const value_count = - has_lists ? page.nesting[state->col.max_nesting_depth - 1].batch_size : state->num_rows; - auto const tid = cg::this_thread_block().thread_rank(); - - // Offsets pointer contains string sizes in case of large strings and actual offsets - // otherwise - auto& ni = state->nesting_info[state->col.max_nesting_depth - 1]; - auto offptr = reinterpret_cast(ni.data_out); - // For large strings, update the initial string buffer offset to be used during large string - // column construction. Otherwise, convert string sizes to final offsets - if (state->col.is_large_string_col) { - // Write zero string sizes - for (int idx = tid; idx < value_count; idx += block_size) { - offptr[idx] = 0; - } - // page.chunk_idx are ordered by input_col_idx and row_group_idx respectively - auto const chunks_per_rowgroup = initial_str_offsets.size(); - auto const input_col_idx = page.chunk_idx % chunks_per_rowgroup; - compute_initial_large_strings_offset(state, initial_str_offsets[input_col_idx]); - } else { - // Write the initial offset at all positions to indicate zero sized strings - for (int idx = tid; idx < value_count; idx += block_size) { - offptr[idx] = initial_value; - } - } -} - template CUDF_HOST_DEVICE constexpr int log2_int() { From 26e24fb06cb1c48954bd10034200b3c04de804a8 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb Date: Wed, 22 Jul 2026 20:26:52 +0000 Subject: [PATCH 08/16] Handle pruned large strings --- cpp/src/io/parquet/decode_pruned_pages.cu | 17 ++++++++++++++++- cpp/src/io/parquet/page_data.cu | 2 +- cpp/src/io/parquet/page_string_utils.cuh | 2 +- cpp/src/io/parquet/parquet_gpu.hpp | 2 ++ cpp/src/io/parquet/reader_impl.cpp | 18 ++++++++++++------ cpp/src/io/parquet/reader_impl.hpp | 8 +++++++- cpp/src/io/parquet/reader_impl_preprocess.cu | 6 ++++-- 7 files changed, 43 insertions(+), 12 deletions(-) diff --git a/cpp/src/io/parquet/decode_pruned_pages.cu b/cpp/src/io/parquet/decode_pruned_pages.cu index 657a8c366957..c668ebdb8530 100644 --- a/cpp/src/io/parquet/decode_pruned_pages.cu +++ b/cpp/src/io/parquet/decode_pruned_pages.cu @@ -8,6 +8,7 @@ #include #include +#include #include namespace cudf::io::parquet::detail { @@ -27,6 +28,7 @@ CUDF_KERNEL void __launch_bounds__(block_size) fill_pruned_offsets_kernel(device_span pages, device_span chunks, device_span page_mask, + device_span initial_str_offsets, size_t skip_rows, size_t num_rows) { @@ -54,6 +56,18 @@ CUDF_KERNEL void __launch_bounds__(block_size) 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); + if (begin >= end) { return; } + + // Large strings needs the first page string offset, including if the page was + // pruned. Record it here. + if (chunk.is_large_string_col and t == 0) { + auto const chunks_per_rowgroup = initial_str_offsets.size(); + auto const input_col_idx = page.chunk_idx % chunks_per_rowgroup; + cuda::atomic_ref initial_str_offset{ + initial_str_offsets[input_col_idx]}; + initial_str_offset.fetch_min(page.str_offset, cuda::std::memory_order_relaxed); + } + // Write zeros for large strings and the page's initial offset otherwise. auto const value = chunk.is_large_string_col ? size_type{0} : static_cast(page.str_offset); @@ -88,13 +102,14 @@ CUDF_KERNEL void __launch_bounds__(block_size) void fill_pruned_offsets(cudf::device_span pages, cudf::device_span chunks, cudf::device_span page_mask, + cudf::device_span initial_str_offsets, size_t skip_rows, size_t num_rows, rmm::cuda_stream_view stream) { CUDF_EXPECTS(pages.size() == page_mask.size(), "Page mask size does not match page count"); fill_pruned_offsets_kernel<<>>( - pages, chunks, page_mask, skip_rows, num_rows); + pages, chunks, page_mask, initial_str_offsets, skip_rows, num_rows); CUDF_CUDA_TRY(cudaGetLastError()); } diff --git a/cpp/src/io/parquet/page_data.cu b/cpp/src/io/parquet/page_data.cu index b6740d0dee71..e5737b790b3c 100644 --- a/cpp/src/io/parquet/page_data.cu +++ b/cpp/src/io/parquet/page_data.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2018-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2018-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ diff --git a/cpp/src/io/parquet/page_string_utils.cuh b/cpp/src/io/parquet/page_string_utils.cuh index c846db4c65d2..85f8bf06200e 100644 --- a/cpp/src/io/parquet/page_string_utils.cuh +++ b/cpp/src/io/parquet/page_string_utils.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ diff --git a/cpp/src/io/parquet/parquet_gpu.hpp b/cpp/src/io/parquet/parquet_gpu.hpp index 4a305b6ed9c1..5c9c0f62dfcb 100644 --- a/cpp/src/io/parquet/parquet_gpu.hpp +++ b/cpp/src/io/parquet/parquet_gpu.hpp @@ -1043,6 +1043,7 @@ void preprocess_levels(cudf::detail::hostdevice_span pages, * @param[in] pages All pages to be processed * @param[in] chunks All chunks to be processed * @param[in] page_mask Boolean vector indicating which pages are decoded + * @param[in,out] initial_str_offsets Initial offsets used to construct large nested strings * @param[in] skip_rows Number of rows to skip * @param[in] num_rows Number of rows to read * @param[in] stream CUDA stream to use @@ -1050,6 +1051,7 @@ void preprocess_levels(cudf::detail::hostdevice_span pages, void fill_pruned_offsets(cudf::device_span pages, cudf::device_span chunks, cudf::device_span page_mask, + cudf::device_span initial_str_offsets, size_t skip_rows, size_t num_rows, rmm::cuda_stream_view stream); diff --git a/cpp/src/io/parquet/reader_impl.cpp b/cpp/src/io/parquet/reader_impl.cpp index bcad8cdccc40..02485a65637c 100644 --- a/cpp/src/io/parquet/reader_impl.cpp +++ b/cpp/src/io/parquet/reader_impl.cpp @@ -397,6 +397,9 @@ void reader_impl::decode_page_data(read_mode mode, size_t skip_rows, size_t num_ // Invalidate output buffer nullmasks at row indices spanned by pruned pages update_output_nullmasks_for_pruned_pages(subpass_page_mask_span(), skip_rows, num_rows); + // Fill output offsets for pruned pages before retrieving large string initial offsets. + fill_pruned_offsets(skip_rows, num_rows, initial_str_offsets); + // Copy over initial string offsets from device auto h_initial_str_offsets = cudf::detail::make_pinned_vector_async(initial_str_offsets, _stream); @@ -439,9 +442,15 @@ void reader_impl::decode_page_data(read_mode mode, size_t skip_rows, size_t num_ } // Nested large strings column else if (input_col.nesting_depth() > 0) { - CUDF_EXPECTS(h_initial_str_offsets[idx] != std::numeric_limits::max(), - "Encountered invalid initial offset for large string column"); - out_buf.set_initial_string_offset(h_initial_str_offsets[idx]); + // A fully pruned list may have no string child values and therefore no page from which + // to record an initial offset. + if (out_buf.size == 0) { + out_buf.set_initial_string_offset(0); + } else { + CUDF_EXPECTS(h_initial_str_offsets[idx] != std::numeric_limits::max(), + "Encountered invalid initial offset for large string column"); + out_buf.set_initial_string_offset(h_initial_str_offsets[idx]); + } } } } @@ -453,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 d3033518389e..a49ca49b2710 100644 --- a/cpp/src/io/parquet/reader_impl.hpp +++ b/cpp/src/io/parquet/reader_impl.hpp @@ -352,8 +352,14 @@ class reader_impl { /** * @brief Fill in string and list offsets for rows covered by pruned data pages. + * + * @param skip_rows Offset of the first row in the table chunk + * @param num_rows Number of rows in the table chunk + * @param initial_str_offsets Initial offsets used to construct large nested strings */ - void fill_pruned_offsets(size_t skip_rows, size_t num_rows); + void fill_pruned_offsets(size_t skip_rows, + size_t num_rows, + cudf::device_span initial_str_offsets); /** * @brief Creates file-wide parquet chunk information. diff --git a/cpp/src/io/parquet/reader_impl_preprocess.cu b/cpp/src/io/parquet/reader_impl_preprocess.cu index 91add706b279..3ac3415c2dc4 100644 --- a/cpp/src/io/parquet/reader_impl_preprocess.cu +++ b/cpp/src/io/parquet/reader_impl_preprocess.cu @@ -1100,7 +1100,9 @@ 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) +void reader_impl::fill_pruned_offsets(size_t skip_rows, + size_t num_rows, + cudf::device_span initial_str_offsets) { // Return early if there are no pruned pages auto const page_mask = subpass_page_mask_span(); @@ -1118,7 +1120,7 @@ void reader_impl::fill_pruned_offsets(size_t skip_rows, size_t num_rows) // Set offsets for pruned string and list pages. parquet::detail::fill_pruned_offsets( - pages, chunks, device_page_mask, skip_rows, num_rows, _stream); + pages, chunks, device_page_mask, initial_str_offsets, skip_rows, num_rows, _stream); } cudf::detail::host_vector reader_impl::calculate_page_string_offsets() From 3ce085078f6fda5c6c6e9e19306e1d8b4511751f Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb Date: Wed, 22 Jul 2026 23:36:48 +0000 Subject: [PATCH 09/16] Fix string pages with missing value info --- .../experimental/hybrid_scan_chunking.cu | 2 +- .../parquet/experimental/hybrid_scan_impl.cpp | 2 +- .../experimental/hybrid_scan_preprocess.cu | 15 ++++++----- cpp/src/io/parquet/page_string_decode.cu | 10 +++---- cpp/src/io/parquet/parquet_gpu.hpp | 2 +- cpp/src/io/parquet/reader_impl.hpp | 4 +-- cpp/src/io/parquet/reader_impl_chunking.cu | 10 +++---- .../io/parquet/reader_impl_chunking_utils.cu | 20 ++++++++------ .../io/parquet/reader_impl_chunking_utils.cuh | 16 +++++------ cpp/src/io/parquet/reader_impl_helpers.hpp | 4 +-- cpp/src/io/parquet/reader_impl_preprocess.cu | 27 ++++++++++++------- .../parquet/reader_impl_preprocess_utils.cuh | 4 +-- 12 files changed, 65 insertions(+), 51 deletions(-) diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu b/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu index 295b06d7d402..4d5174718b10 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu +++ b/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu @@ -159,7 +159,7 @@ void hybrid_scan_reader_impl::setup_next_pass( // if we are doing subpass reading, generate more accurate num_row estimates for list columns. // this helps us to generate more accurate subpass splits. if (pass.has_compressed_data && _input_pass_read_limit != 0) { - if (_has_page_index) { + if (_has_offset_index) { generate_list_column_row_counts(is_estimate_row_counts::NO); } else { generate_list_column_row_counts(is_estimate_row_counts::YES); diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp b/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp index 631aa4197bfd..95b00837be2c 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp +++ b/cpp/src/io/parquet/experimental/hybrid_scan_impl.cpp @@ -1261,7 +1261,7 @@ void hybrid_scan_reader_impl::reset_internal_state() _row_mask_offset = 0; _file_itm_data = file_intermediate_data{}; _file_preprocessed = false; - _has_page_index = false; + _has_offset_index = false; _pass_itm_data.reset(); _pass_page_mask.clear(); _subpass_page_mask.reset(); diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu b/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu index bf072bcee4f3..65aeb0881aa9 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu +++ b/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu @@ -115,10 +115,11 @@ void hybrid_scan_reader_impl::prepare_row_groups( _file_itm_data.num_rows_per_source.cend(), _file_itm_data.exclusive_sum_num_rows_per_source.begin()); - // check for page indexes - _has_page_index = std::all_of(_file_itm_data.row_groups.cbegin(), - _file_itm_data.row_groups.cend(), - [](auto const& row_group) { return row_group.has_page_index(); }); + // Check for offset indexes. + _has_offset_index = + std::all_of(_file_itm_data.row_groups.cbegin(), + _file_itm_data.row_groups.cend(), + [](auto const& row_group) { return row_group.has_offset_index(); }); if (_file_itm_data.global_num_rows > 0 && not _file_itm_data.row_groups.empty() && not _input_columns.empty()) { @@ -220,13 +221,13 @@ void hybrid_scan_reader_impl::setup_compressed_data( pass.has_compressed_data = setup_column_chunks(column_chunk_data); // Process dataset chunk pages into output columns - auto const total_pages = _has_page_index ? count_page_headers_with_pgidx(chunks, _stream) - : count_page_headers(chunks, _stream); + auto const total_pages = _has_offset_index ? count_page_headers_with_pgidx(chunks, _stream) + : count_page_headers(chunks, _stream); if (total_pages <= 0) { return; } rmm::device_uvector unsorted_pages(total_pages, _stream); // decoding of column/page information - parquet::detail::decode_page_headers(pass, unsorted_pages, _has_page_index, _stream); + parquet::detail::decode_page_headers(pass, unsorted_pages, _has_offset_index, _stream); CUDF_EXPECTS(pass.page_offsets.size() - 1 == static_cast(_input_columns.size()), "Encountered page_offsets / num_columns mismatch"); } diff --git a/cpp/src/io/parquet/page_string_decode.cu b/cpp/src/io/parquet/page_string_decode.cu index 4a53c717cbcd..adc5c375b8f7 100644 --- a/cpp/src/io/parquet/page_string_decode.cu +++ b/cpp/src/io/parquet/page_string_decode.cu @@ -520,7 +520,7 @@ CUDF_KERNEL void __launch_bounds__(preprocess_block_size) if (t == 0) { // don't clobber these if they're already computed from the index - if (!pp->has_page_index) { + if (!pp->has_value_info) { pp->num_nulls = 0; pp->num_valids = 0; } @@ -553,7 +553,7 @@ CUDF_KERNEL void __launch_bounds__(preprocess_block_size) is_bounds_page(s->page, s->col.start_row, min_row, num_rows, has_repetition); // if we have size info, then we only need to do this for bounds pages - if (pp->has_page_index && !is_bounds_pg) { return; } + if (pp->has_value_info && !is_bounds_pg) { return; } // Zero out everything and return early if the page is pruned if (not page_mask.empty() and not page_mask[page_idx]) { @@ -646,7 +646,7 @@ CUDF_KERNEL void __launch_bounds__(delta_preproc_block_size) is_bounds_page(s->page, s->col.start_row, min_row, num_rows, has_repetition); // if we have size info, then we only need to do this for bounds pages - if (pp->has_page_index && !is_bounds_pg) { + if (pp->has_value_info && !is_bounds_pg) { // check if we need to store values from the index if (t == 0 && is_page_contained(s->page, s->col.start_row, min_row, num_rows)) { pp->str_bytes = pp->str_bytes_from_index; @@ -728,7 +728,7 @@ CUDF_KERNEL void __launch_bounds__(delta_length_block_size) is_bounds_page(s->page, s->col.start_row, min_row, num_rows, has_repetition); // if we have size info, then we only need to do this for bounds pages - if (pp->has_page_index && !is_bounds_pg) { + if (pp->has_value_info && !is_bounds_pg) { // check if we need to store values from the index if (t == 0 && is_page_contained(s->page, s->col.start_row, min_row, num_rows)) { pp->str_bytes = pp->str_bytes_from_index; @@ -842,7 +842,7 @@ CUDF_KERNEL void __launch_bounds__(preprocess_block_size) is_bounds_page(s->page, s->col.start_row, min_row, num_rows, has_repetition); // if we have size info, then we only need to do this for bounds pages - if (pp->has_page_index && !is_bounds_pg) { + if (pp->has_value_info && !is_bounds_pg) { // check if we need to store values from the index if (t == 0 && is_page_contained(s->page, s->col.start_row, min_row, num_rows)) { pp->str_bytes = pp->str_bytes_from_index; diff --git a/cpp/src/io/parquet/parquet_gpu.hpp b/cpp/src/io/parquet/parquet_gpu.hpp index 5c9c0f62dfcb..e69415fa8003 100644 --- a/cpp/src/io/parquet/parquet_gpu.hpp +++ b/cpp/src/io/parquet/parquet_gpu.hpp @@ -374,7 +374,7 @@ struct PageInfo { Encoding definition_level_encoding; // Encoding used for definition levels (data page) Encoding repetition_level_encoding; // Encoding used for repetition levels (data page) bool is_compressed; // Whether the page is compressed (V2 header) - bool has_page_index; // true if str_bytes, num_valids, etc are derivable from page indexes + bool has_value_info; // true if str_bytes, num_valids, etc are derivable from page indexes }; // forward declaration diff --git a/cpp/src/io/parquet/reader_impl.hpp b/cpp/src/io/parquet/reader_impl.hpp index a49ca49b2710..76d88f52e310 100644 --- a/cpp/src/io/parquet/reader_impl.hpp +++ b/cpp/src/io/parquet/reader_impl.hpp @@ -579,8 +579,8 @@ class reader_impl { bool _strings_to_categorical = false; - // are there usable page indexes available - bool _has_page_index = false; + // are offset indexes available for selected row groups + bool _has_offset_index = false; std::optional> _reader_column_schema; diff --git a/cpp/src/io/parquet/reader_impl_chunking.cu b/cpp/src/io/parquet/reader_impl_chunking.cu index 9e4957dbea05..c2a53826bb05 100644 --- a/cpp/src/io/parquet/reader_impl_chunking.cu +++ b/cpp/src/io/parquet/reader_impl_chunking.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -177,7 +177,7 @@ void reader_impl::setup_next_pass(read_mode mode) // if we are doing subpass reading, generate more accurate num_row estimates for list columns. // this helps us to generate more accurate subpass splits. if (pass.has_compressed_data && _input_pass_read_limit != 0) { - if (not _has_page_index) { + if (not _has_offset_index) { generate_list_column_row_counts(is_estimate_row_counts::YES); } else { generate_list_column_row_counts(is_estimate_row_counts::NO); @@ -295,7 +295,7 @@ void reader_impl::setup_next_subpass(read_mode mode) remaining_read_limit, num_columns, is_first_subpass, - _has_page_index, + _has_offset_index, _stream); }(); @@ -422,7 +422,7 @@ void reader_impl::create_global_chunk_info() // Mapping of input column to page index column std::vector column_mapping; - if (_has_page_index and not row_groups_info.empty()) { + if (_has_offset_index and not row_groups_info.empty()) { // use first row group to define mappings (assumes same schema for each file) auto const& rg = row_groups_info[0]; auto const& columns = _metadata->get_row_group(rg.index, rg.source_index).columns; @@ -480,7 +480,7 @@ void reader_impl::create_global_chunk_info() // grab the column_chunk_info for each chunk (if it exists) column_chunk_info const* const chunk_info = - _has_page_index ? &rg.column_chunks.value()[column_mapping[i]] : nullptr; + _has_offset_index ? &rg.column_chunks.value()[column_mapping[i]] : nullptr; chunks.emplace_back(col_meta.total_compressed_size, nullptr, diff --git a/cpp/src/io/parquet/reader_impl_chunking_utils.cu b/cpp/src/io/parquet/reader_impl_chunking_utils.cu index 96d3f1a67cc2..ebca8918171c 100644 --- a/cpp/src/io/parquet/reader_impl_chunking_utils.cu +++ b/cpp/src/io/parquet/reader_impl_chunking_utils.cu @@ -356,7 +356,7 @@ std::tuple, size_t, size_t> compute_next_subpass( size_t size_limit, size_t num_columns, bool is_first_subpass, - bool has_page_index, + bool has_offset_index, rmm::cuda_stream_view stream) { auto [aggregated_info, page_keys_by_split] = adjust_cumulative_sizes(c_info, pages, stream); @@ -387,13 +387,17 @@ std::tuple, size_t, size_t> compute_next_subpass( auto iter = cuda::counting_iterator{size_t{0}}; auto page_row_index = cudf::detail::make_counting_transform_iterator(0, get_page_end_row_index{c_info}); - thrust::transform( - rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), - iter, - iter + num_columns, - page_bounds.begin(), - get_page_span{ - page_offsets, chunks, page_row_index, start_row, end_row, is_first_subpass, has_page_index}); + thrust::transform(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + iter, + iter + num_columns, + page_bounds.begin(), + get_page_span{page_offsets, + chunks, + page_row_index, + start_row, + end_row, + is_first_subpass, + has_offset_index}); // total page count over all columns auto page_count_iter = cuda::make_transform_iterator(page_bounds.begin(), get_span_size{}); diff --git a/cpp/src/io/parquet/reader_impl_chunking_utils.cuh b/cpp/src/io/parquet/reader_impl_chunking_utils.cuh index 10a172c75c9b..347f874e11a5 100644 --- a/cpp/src/io/parquet/reader_impl_chunking_utils.cuh +++ b/cpp/src/io/parquet/reader_impl_chunking_utils.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -128,7 +128,7 @@ adjust_cumulative_sizes(device_span c_info, * @param size_limit The size limit in bytes of the subpass * @param num_columns The number of columns * @param is_first_subpass Boolean indicating if this is the first subpass - * @param has_page_index Boolean indicating if we have a page index + * @param has_offset_index Boolean indicating if offset indexes are available * @param stream The stream to execute cuda operations on * @returns A tuple containing a vector of page_span structs indicating the page indices to include * for each column to be processed, the total number of pages over all columns, and the total @@ -144,7 +144,7 @@ std::tuple, size_t, size_t> compute_next_subpass( size_t size_limit, size_t num_columns, bool is_first_subpass, - bool has_page_index, + bool has_offset_index, rmm::cuda_stream_view stream); /** @@ -626,7 +626,7 @@ struct get_page_span { size_t const start_row; size_t const end_row; bool const is_first_subpass; - bool const has_page_index; + bool const has_offset_index; get_page_span(device_span _page_offsets, device_span _chunks, @@ -634,14 +634,14 @@ struct get_page_span { size_t _start_row, size_t _end_row, bool _is_first_subpass, - bool _has_page_index) + bool _has_offset_index) : page_offsets(_page_offsets), chunks(_chunks), page_row_index(_page_row_index), start_row(_start_row), end_row(_end_row), is_first_subpass(_is_first_subpass), - has_page_index(_has_page_index) + has_offset_index(_has_offset_index) { } @@ -656,13 +656,13 @@ struct get_page_span { // For list columns, the row counts are estimates so we need all prefix pages to correctly // compute page bounds. For non-list columns, we can get an exact span of pages. auto start_page = first_page_index; - auto const update_start_page = has_page_index or (not is_list) or (not is_first_subpass); + auto const update_start_page = has_offset_index or (not is_list) or (not is_first_subpass); if (update_start_page) { start_page += cuda::std::distance( column_page_start, thrust::lower_bound(thrust::seq, column_page_start, column_page_end, start_row)); } - if (page_row_index[start_page] == start_row and (has_page_index or not is_list)) { + if (page_row_index[start_page] == start_row and (has_offset_index or not is_list)) { start_page++; } diff --git a/cpp/src/io/parquet/reader_impl_helpers.hpp b/cpp/src/io/parquet/reader_impl_helpers.hpp index dacaf2bfa22f..40c70babf2b8 100644 --- a/cpp/src/io/parquet/reader_impl_helpers.hpp +++ b/cpp/src/io/parquet/reader_impl_helpers.hpp @@ -77,9 +77,9 @@ struct row_group_info { std::optional> column_chunks; /** - * @brief Indicates the presence of page-level indexes. + * @brief Indicates the presence of page-level offset indexes. */ - [[nodiscard]] bool has_page_index() const { return column_chunks.has_value(); } + [[nodiscard]] bool has_offset_index() const { return column_chunks.has_value(); } }; /** diff --git a/cpp/src/io/parquet/reader_impl_preprocess.cu b/cpp/src/io/parquet/reader_impl_preprocess.cu index 3ac3415c2dc4..5ade86bfdcf2 100644 --- a/cpp/src/io/parquet/reader_impl_preprocess.cu +++ b/cpp/src/io/parquet/reader_impl_preprocess.cu @@ -570,8 +570,8 @@ void reader_impl::read_compressed_data() read_chunks_tasks.get(); // Process dataset chunk pages into output columns - auto const total_pages = _has_page_index ? count_page_headers_with_pgidx(chunks, _stream) - : count_page_headers(chunks, _stream); + auto const total_pages = _has_offset_index ? count_page_headers_with_pgidx(chunks, _stream) + : count_page_headers(chunks, _stream); if (total_pages <= 0) { return; } // Zero out the vector before `decode_page_headers` as it may not write every byte of the buffer, @@ -580,7 +580,7 @@ void reader_impl::read_compressed_data() total_pages, _stream, cudf::get_current_device_resource_ref()); // decoding of column/page information - decode_page_headers(pass, unsorted_pages, _has_page_index, _stream); + decode_page_headers(pass, unsorted_pages, _has_offset_index, _stream); CUDF_EXPECTS(pass.page_offsets.size() - 1 == static_cast(_input_columns.size()), "Encountered page_offsets / num_columns mismatch"); } @@ -632,10 +632,11 @@ void reader_impl::preprocess_file(read_mode mode) _file_itm_data.exclusive_sum_num_rows_per_source.begin()); } - // check for page indexes - _has_page_index = std::all_of(_file_itm_data.row_groups.cbegin(), - _file_itm_data.row_groups.cend(), - [](auto const& row_group) { return row_group.has_page_index(); }); + // Check for offset indexes. + _has_offset_index = + std::all_of(_file_itm_data.row_groups.cbegin(), + _file_itm_data.row_groups.cend(), + [](auto const& row_group) { return row_group.has_offset_index(); }); if (_file_itm_data.global_num_rows > 0 && not _file_itm_data.row_groups.empty() && not _input_columns.empty()) { @@ -841,7 +842,15 @@ void reader_impl::preprocess_subpass_pages(read_mode mode, size_t chunk_read_lim is_treat_fixed_length_as_string(chunk.logical_type); }); - if (!_has_page_index || has_flba) { + // String pages with missing value info do not have the value counts or string byte sizes needed + // for chunking. Scan them so chunk boundaries and output allocations use their actual sizes. + auto const has_string_page_without_info = + std::any_of(subpass.pages.host_begin(), subpass.pages.host_end(), [](auto const& page) { + return (static_cast(page.kernel_mask) & STRINGS_MASK) != 0 and + (page.flags & PAGEINFO_FLAGS_DICTIONARY) == 0 and not page.has_value_info; + }); + + if (has_string_page_without_info || has_flba) { constexpr bool compute_all_string_sizes = true; compute_page_string_sizes_pass1(subpass.pages, pass.chunks, @@ -891,7 +900,7 @@ void reader_impl::preprocess_subpass_pages(read_mode mode, size_t chunk_read_lim // corner case: only decode up to the second-to-last row, except if this is the last page in the // entire pass or if we have the page index. this handles the case where we only have 1 chunk, 1 // page, and potentially even just 1 row. - if (is_list and std::cmp_less(max_col_row, last_pass_row) and not _has_page_index) { + if (is_list and std::cmp_less(max_col_row, last_pass_row) and not _has_offset_index) { // compute min row for this column in the subpass auto const& first_page = subpass.pages[first_page_index]; auto const& first_chunk = pass.chunks[first_page.chunk_idx]; diff --git a/cpp/src/io/parquet/reader_impl_preprocess_utils.cuh b/cpp/src/io/parquet/reader_impl_preprocess_utils.cuh index a7f0e7187197..c7244a029bb8 100644 --- a/cpp/src/io/parquet/reader_impl_preprocess_utils.cuh +++ b/cpp/src/io/parquet/reader_impl_preprocess_utils.cuh @@ -188,9 +188,9 @@ struct copy_page_info { 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 != 0; + pg.has_value_info = pi.has_value_info != 0; pg.start_val = 0; - if (pg.has_page_index) { + if (pg.has_value_info) { pg.num_nulls = pi.num_nulls; pg.num_valids = pi.num_valids; pg.str_bytes_from_index = pi.str_bytes; From 8d7d5ed9ee8e5b1e03c6b86298e15d33a44ab28a Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb Date: Thu, 23 Jul 2026 18:09:36 +0000 Subject: [PATCH 10/16] Simplify --- cpp/src/io/parquet/page_hdr.cu | 172 +++++------------- cpp/src/io/parquet/parquet_gpu.hpp | 19 +- .../parquet/reader_impl_preprocess_utils.cu | 49 ++--- 3 files changed, 78 insertions(+), 162 deletions(-) diff --git a/cpp/src/io/parquet/page_hdr.cu b/cpp/src/io/parquet/page_hdr.cu index a7071b8979f5..84a429c2a88e 100644 --- a/cpp/src/io/parquet/page_hdr.cu +++ b/cpp/src/io/parquet/page_hdr.cu @@ -3,6 +3,7 @@ * SPDX-License-Identifier: Apache-2.0 */ +#include "cuda/std/__utility/cmp.h" #include "error.hpp" #include "io/utilities/block_utils.cuh" #include "parquet_gpu.hpp" @@ -520,46 +521,6 @@ void __forceinline__ __device__ zero_out_page_header_info(byte_stream_s* bs) bs->page.kernel_mask = decode_kernel_mask::NONE; } -/** - * @brief Decode a page header from an initialized byte stream. - * - * @param bs Byte stream - * @param chunk_idx Index of the chunk containing the page - * @param page Pointer to the page info to decode - * @param error_code Pointer to the error code for kernel failures - */ -__device__ void decode_page_header(byte_stream_s* bs, - cudf::size_type chunk_idx, - PageInfo* page, - kernel_error::pointer error_code) -{ - bs->page.chunk_idx = chunk_idx; - bs->page.src_col_schema = bs->ck.src_col_schema; - - 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; - } -} - /** * @brief Kernel for outputting page headers from the specified column chunks * @@ -776,54 +737,6 @@ CUDF_KERNEL void __launch_bounds__(count_page_headers_block_size) }); } -/** - * @brief Functor to decode page headers from specified page locations - */ -struct decode_using_page_index_fn { - cudf::device_span colchunks; - cudf::device_span pages; - cudf::device_span chunk_page_offsets; - uint8_t** page_locations; - 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.begin(), - thrust::upper_bound( - thrust::seq, chunk_page_offsets.begin(), chunk_page_offsets.end(), 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]; - bs.base = bs.cur = page_locations[page_idx]; - bs.end = bs.ck.compressed_data + bs.ck.compressed_size; - // Check if byte stream pointers are valid. - if (bs.end < bs.cur) { - set_error(static_cast(decode_error::DATA_STREAM_OVERRUN), - error_code); - return; - } - // Clear page header info before writing known fields - zero_out_page_header_info(&bs); - - decode_page_header(&bs, chunk_idx, &pages[page_idx], error_code); - - bs.page.page_data = const_cast(bs.cur); - bs.page.kernel_mask = kernel_mask_for_page(bs.page, bs.ck); - - // Copy the page info to the output span - pages[page_idx] = bs.page; - } -}; - /** * @brief Functor to decode specified page headers from corresponding page data spans */ @@ -837,12 +750,16 @@ struct decode_from_page_data_fn { __device__ void operator()(size_type page_idx) const noexcept { auto const num_chunks = static_cast(colchunks.size()); - auto const chunk_idx = static_cast( + + // Binary search the column chunk index for this page + auto const chunk_idx = static_cast( cuda::std::distance( chunk_page_offsets.begin(), thrust::upper_bound( thrust::seq, chunk_page_offsets.begin(), chunk_page_offsets.end(), page_idx)) - 1); + + // Check if the chunk index is valid. if (chunk_idx < 0 or chunk_idx >= num_chunks) { set_error(static_cast(decode_error::DATA_STREAM_OVERRUN), error_code); @@ -851,25 +768,55 @@ struct decode_from_page_data_fn { byte_stream_s bs{}; bs.ck = colchunks[chunk_idx]; + // Clear the logical page descriptor. zero_out_page_header_info(&bs); + bs.page.chunk_idx = chunk_idx; + bs.page.src_col_schema = bs.ck.src_col_schema; + + // bs.page.chunk_row not computed here and will be filled in later by + // `fill_in_page_info()`. + + // Return if empty page span (pruned page) auto const page_span = page_data[page_idx]; if (page_span.empty()) { - // Initialize the logical page descriptor. Page-index fields are populated by - // fill_in_page_info(). - bs.page.chunk_idx = chunk_idx; - bs.page.src_col_schema = bs.ck.src_col_schema; - pages[page_idx] = bs.page; + pages[page_idx] = bs.page; return; } - bs.base = bs.cur = page_span.data(); - bs.end = page_span.data() + page_span.size(); - - decode_page_header(&bs, chunk_idx, &pages[page_idx], error_code); + // Parsed page must be valid and not empty + 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: + // this computation is only valid for flat schemas. for nested schemas, + // they will be recomputed in the preprocess step by examining repetition and + // definition levels + bs.page.num_rows = bs.page.num_input_values; + break; + case PageType::DATA_PAGE_V2: + bs.page.flags |= PAGEINFO_FLAGS_V2; + // V2 only uses RLE, so it was removed from the header + 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)) { + // Ensure we read the entire page. + if (cuda::std::cmp_not_equal(bs.end - bs.cur, bs.page.compressed_page_size)) { set_error(static_cast(decode_error::DATA_STREAM_OVERRUN), error_code); return; @@ -877,7 +824,9 @@ struct decode_from_page_data_fn { 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; + + // Copy over the page info from byte stream + pages[page_idx] = bs.page; } }; @@ -997,25 +946,6 @@ void decode_page_headers(cudf::device_span chunks, CUDF_CUDA_TRY(cudaGetLastError()); } -void decode_page_headers_using_page_index(cudf::device_span chunks, - cudf::device_span pages, - uint8_t** page_locations, - cudf::device_span chunk_page_offsets, - kernel_error::pointer error_code, - rmm::cuda_stream_view stream) -{ - CUDF_EXPECTS(chunk_page_offsets.size() == chunks.size() + 1, - "Chunk page offsets must cover all chunks"); - 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_using_page_index_fn{.colchunks = chunks, - .pages = pages, - .chunk_page_offsets = chunk_page_offsets, - .page_locations = page_locations, - .error_code = error_code}); -} - void decode_page_headers_from_page_data( cudf::device_span chunks, cudf::device_span pages, @@ -1024,8 +954,6 @@ void decode_page_headers_from_page_data( kernel_error::pointer error_code, rmm::cuda_stream_view stream) { - CUDF_EXPECTS(page_data.size() == pages.size(), - "Page span count must match the number of logical pages"); CUDF_EXPECTS(chunk_page_offsets.size() == chunks.size() + 1, "Chunk page offsets must cover all chunks"); thrust::for_each(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), diff --git a/cpp/src/io/parquet/parquet_gpu.hpp b/cpp/src/io/parquet/parquet_gpu.hpp index e69415fa8003..b8324b3f570a 100644 --- a/cpp/src/io/parquet/parquet_gpu.hpp +++ b/cpp/src/io/parquet/parquet_gpu.hpp @@ -712,24 +712,7 @@ void decode_page_headers(cudf::device_span chunks, rmm::cuda_stream_view stream); /** - * @brief Decode page headers from specified page locations from the page index - * - * @param[in] chunks Device span of column chunks - * @param[out] pages Device span of pages - * @param[in] page_locations List of page locations - * @param[in] chunk_page_offsets List of running count of page locations per column chunk - * @param[out] error_code Error code for kernel failures - * @param[in] stream CUDA stream to use - */ -void decode_page_headers_using_page_index(cudf::device_span chunks, - cudf::device_span pages, - uint8_t** page_locations, - cudf::device_span chunk_page_offsets, - kernel_error::pointer error_code, - rmm::cuda_stream_view stream); - -/** - * @brief Decode specified page headers from corresponding page data spans. + * @brief Decode page headers from corresponding specified page data spans. * * Empty spans initialize the corresponding logical page descriptor but are not decoded. * diff --git a/cpp/src/io/parquet/reader_impl_preprocess_utils.cu b/cpp/src/io/parquet/reader_impl_preprocess_utils.cu index ae88f300ee81..3c3767fa6eae 100644 --- a/cpp/src/io/parquet/reader_impl_preprocess_utils.cu +++ b/cpp/src/io/parquet/reader_impl_preprocess_utils.cu @@ -479,11 +479,11 @@ void decode_page_headers_impl(pass_intermediate_data& pass, error_code.data(), stream); } - // If offset index is present, collect data ptrs for all pages and launch the accelerated decode + // If offset index is present, collect data spans for all pages and launch the accelerated decode // page headers kernel - else if (data_source_type == page_data_source_type::OFFSET_INDEX) { - auto host_page_locations = - cudf::detail::make_pinned_vector_async(unsorted_pages.size(), stream); + else if constexpr (data_source_type == page_data_source_type::OFFSET_INDEX) { + auto host_page_data = cudf::detail::make_pinned_vector_async>( + unsorted_pages.size(), stream); auto curr_page_idx = 0; std::for_each(pass.chunks.begin(), pass.chunks.end(), [&](auto const& chunk) { @@ -501,9 +501,11 @@ void decode_page_headers_impl(pass_intermediate_data& pass, CUDF_EXPECTS(std::cmp_less(chunk.h_chunk_info->dictionary_offset.value(), chunk.h_chunk_info->pages.front().location.offset), "Encountered dictionary page located beyond the first data page"); - host_page_locations[curr_page_idx] = data_ptr; + auto const dictionary_size = chunk.h_chunk_info->dictionary_size.value(); + CUDF_EXPECTS(dictionary_size >= 0, "Encountered invalid dictionary page size"); + host_page_data[curr_page_idx] = {data_ptr, static_cast(dictionary_size)}; ++curr_page_idx; - data_ptr += chunk.h_chunk_info->dictionary_size.value(); + data_ptr += dictionary_size; } // Data pages @@ -511,31 +513,34 @@ void decode_page_headers_impl(pass_intermediate_data& pass, std::cmp_equal(chunk.h_chunk_info->pages.size(), chunk.num_data_pages), "Encountered invalid sized data page information in the page index"); auto const num_data_pages = chunk.num_data_pages; - std::for_each(cuda::counting_iterator{0}, - cuda::counting_iterator{num_data_pages}, - [&](auto const page_idx) { - host_page_locations[curr_page_idx] = data_ptr; - ++curr_page_idx; - if (page_idx < num_data_pages - 1) { - data_ptr += chunk.h_chunk_info->pages[page_idx + 1].location.offset - - chunk.h_chunk_info->pages[page_idx].location.offset; - } - }); + std::for_each( + cuda::counting_iterator{0}, + cuda::counting_iterator{num_data_pages}, + [&](auto const page_idx) { + auto const page_size = chunk.h_chunk_info->pages[page_idx].location.compressed_page_size; + CUDF_EXPECTS(page_size >= 0, "Encountered invalid data page size"); + host_page_data[curr_page_idx] = {data_ptr, static_cast(page_size)}; + ++curr_page_idx; + if (page_idx < num_data_pages - 1) { + data_ptr += chunk.h_chunk_info->pages[page_idx + 1].location.offset - + chunk.h_chunk_info->pages[page_idx].location.offset; + } + }); }); - // Check if we have data ptrs for all input pages + // Check if we have data spans for all input pages CUDF_EXPECTS(std::cmp_equal(curr_page_idx, unsorted_pages.size()), "Expected page offsets to match total pages"); - // Copy page data ptrs to device - auto page_locations = cudf::detail::make_device_uvector_async( - host_page_locations, stream, cudf::get_current_device_resource_ref()); + // Copy page data spans to device + auto page_data = cudf::detail::make_device_uvector_async( + host_page_data, stream, cudf::get_current_device_resource_ref()); // Accelerated decode page headers, one thread per page - decode_page_headers_using_page_index( + decode_page_headers_from_page_data( device_span(pass.chunks.device_ptr(), pass.chunks.size()), unsorted_pages, - page_locations.begin(), + page_data, device_span(chunk_page_offsets.data(), chunk_page_offsets.size()), error_code.data(), stream); From 503e917d7ee6efdb4583dd5746f4f88add8cad61 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb Date: Thu, 23 Jul 2026 18:12:52 +0000 Subject: [PATCH 11/16] Minor --- cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu b/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu index 65aeb0881aa9..8b9bbd70a08b 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu +++ b/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu @@ -182,7 +182,7 @@ 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(_has_offset_index, "Sparse page I/O requires offset 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(), From 465e5030846c6ee445098a28b5f3963b30c1e62a Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb Date: Thu, 23 Jul 2026 18:16:59 +0000 Subject: [PATCH 12/16] Minor --- cpp/src/io/parquet/page_hdr.cu | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/cpp/src/io/parquet/page_hdr.cu b/cpp/src/io/parquet/page_hdr.cu index 84a429c2a88e..b69c26a4d5eb 100644 --- a/cpp/src/io/parquet/page_hdr.cu +++ b/cpp/src/io/parquet/page_hdr.cu @@ -766,19 +766,22 @@ struct decode_from_page_data_fn { return; } + auto const page_span = page_data[page_idx]; + byte_stream_s bs{}; - bs.ck = colchunks[chunk_idx]; + bs.ck = colchunks[chunk_idx]; + bs.base = bs.cur = page_span.data(); + bs.end = page_span.data() + page_span.size(); + // Clear the logical page descriptor. zero_out_page_header_info(&bs); bs.page.chunk_idx = chunk_idx; bs.page.src_col_schema = bs.ck.src_col_schema; - // bs.page.chunk_row not computed here and will be filled in later by // `fill_in_page_info()`. // Return if empty page span (pruned page) - auto const page_span = page_data[page_idx]; if (page_span.empty()) { pages[page_idx] = bs.page; return; From 148a007360a8adae0b4a71eeadb87fec45946cd4 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb Date: Thu, 30 Jul 2026 22:30:19 +0000 Subject: [PATCH 13/16] Add python bindings --- .../pylibcudf/io/experimental/__init__.pxd | 1 + .../pylibcudf/io/experimental/__init__.py | 2 + .../pylibcudf/io/experimental/hybrid_scan.pxd | 7 + .../pylibcudf/io/experimental/hybrid_scan.pyi | 51 ++++ .../pylibcudf/io/experimental/hybrid_scan.pyx | 253 +++++++++++++++++- .../pylibcudf/libcudf/io/hybrid_scan.pxd | 68 +++++ .../tests/io/test_experimental_hybrid_scan.py | 96 +++++++ 7 files changed, 477 insertions(+), 1 deletion(-) diff --git a/python/pylibcudf/pylibcudf/io/experimental/__init__.pxd b/python/pylibcudf/pylibcudf/io/experimental/__init__.pxd index 87cc217ebf94..b36338238833 100644 --- a/python/pylibcudf/pylibcudf/io/experimental/__init__.pxd +++ b/python/pylibcudf/pylibcudf/io/experimental/__init__.pxd @@ -3,5 +3,6 @@ from pylibcudf.io.experimental.hybrid_scan cimport ( FileMetaData, + HybridScanMultiFile, HybridScanReader, ) diff --git a/python/pylibcudf/pylibcudf/io/experimental/__init__.py b/python/pylibcudf/pylibcudf/io/experimental/__init__.py index 6c64231eb1e9..5f31ee5efae2 100644 --- a/python/pylibcudf/pylibcudf/io/experimental/__init__.py +++ b/python/pylibcudf/pylibcudf/io/experimental/__init__.py @@ -2,6 +2,7 @@ # SPDX-License-Identifier: Apache-2.0 from pylibcudf.io.experimental.hybrid_scan import ( + HybridScanMultiFile, HybridScanReader, UseDataPageMask, ) @@ -9,6 +10,7 @@ __all__ = [ "FileMetaData", # backwards compatibility + "HybridScanMultiFile", "HybridScanReader", "UseDataPageMask", ] diff --git a/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pxd b/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pxd index a19cd7db8bf8..91a26bb53d29 100644 --- a/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pxd +++ b/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pxd @@ -14,6 +14,7 @@ from pylibcudf.io.parquet_metadata cimport FileMetaData as c_FileMetaData from pylibcudf.io.types cimport TableWithMetadata from pylibcudf.libcudf.io.hybrid_scan cimport ( hybrid_scan_reader as cpp_hybrid_scan_reader, + hybrid_scan_multifile as cpp_hybrid_scan_multifile, use_data_page_mask, ) from pylibcudf.libcudf.io.hybrid_scan cimport const_uint8_t @@ -27,3 +28,9 @@ cdef class HybridScanReader: cdef unique_ptr[cpp_hybrid_scan_reader] c_obj cdef Stream _stream cdef DeviceMemoryResource mr + +cdef class HybridScanMultiFile: + cdef unique_ptr[cpp_hybrid_scan_multifile] c_obj + cdef Stream _stream + cdef DeviceMemoryResource mr + cdef object _page_data_keepalive diff --git a/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pyi b/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pyi index f95dc8b054d3..b1229306862e 100644 --- a/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pyi +++ b/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pyi @@ -142,3 +142,54 @@ class HybridScanReader: pass_read_limit: int, ) -> list[list[int]]: ... def has_next_table_chunk(self) -> bool: ... + +class HybridScanMultiFile: + @staticmethod + def from_parquet_metadatas( + parquet_metadatas: list[FileMetaData], + options: ParquetReaderOptions, + ) -> HybridScanMultiFile: ... + def parquet_metadatas(self) -> list[FileMetaData]: ... + def page_index_byte_ranges(self) -> list[ByteRangeInfo]: ... + def setup_page_indexes(self, page_index_bytes: list[Buffer]) -> None: ... + def all_row_groups( + self, options: ParquetReaderOptions + ) -> list[list[int]]: ... + def total_rows_in_row_groups( + self, row_group_indices: list[list[int]] + ) -> int: ... + def build_all_true_row_mask( + self, + row_group_indices: list[list[int]], + stream: CudaStreamLike | None = None, + mr: DeviceMemoryResource | None = None, + ) -> Column: ... + def payload_column_chunks_byte_ranges( + self, + row_group_indices: list[list[int]], + row_mask: Column, + mask_data_pages: UseDataPageMask, + options: ParquetReaderOptions, + stream: CudaStreamLike | None = None, + ) -> list[list[ByteRangeInfo]]: ... + def setup_chunking_for_payload_columns( + self, + chunk_read_limit: int, + pass_read_limit: int, + row_group_indices: list[list[int]], + row_mask: Column, + mask_data_pages: UseDataPageMask, + page_data_per_source: list[list], + options: ParquetReaderOptions, + stream: CudaStreamLike | None = None, + mr: DeviceMemoryResource | None = None, + ) -> None: ... + def materialize_payload_columns_chunk( + self, row_mask: Column + ) -> TableWithMetadata: ... + def construct_row_group_passes( + self, + row_group_indices: list[list[int]], + pass_read_limit: int, + ) -> list[list[list[int]]]: ... + def has_next_table_chunk(self) -> bool: ... diff --git a/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pyx b/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pyx index 664eb489428c..2f063bcde601 100644 --- a/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pyx +++ b/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pyx @@ -21,11 +21,19 @@ from pylibcudf.libcudf.column.column cimport column from pylibcudf.libcudf.column.column_view cimport column_view, mutable_column_view from pylibcudf.libcudf.io.hybrid_scan cimport ( const_device_span_const_uint8_t, + const_FileMetaData, + const_host_span_const_uint8_t, const_size_type, const_uint8_t, + const_vector_device_span_const_uint8_t, + const_vector_size_type, + hybrid_scan_multifile as cpp_hybrid_scan_multifile, hybrid_scan_reader as cpp_hybrid_scan_reader, use_data_page_mask as cpp_use_data_page_mask, ) +from pylibcudf.libcudf.io.parquet_schema cimport ( + FileMetaData as cpp_FileMetaData, +) from pylibcudf.libcudf.io.text cimport byte_range_info from pylibcudf.libcudf.io.types cimport table_with_metadata from pylibcudf.libcudf.types cimport size_type @@ -39,7 +47,12 @@ import pylibcudf.libcudf.io.hybrid_scan UseDataPageMask = pylibcudf.libcudf.io.hybrid_scan.use_data_page_mask -__all__ = ["FileMetaData", "HybridScanReader", "UseDataPageMask"] +__all__ = [ + "FileMetaData", + "HybridScanMultiFile", + "HybridScanReader", + "UseDataPageMask", +] cdef device_span[const_uint8_t] _get_device_span(object obj) except *: @@ -53,6 +66,18 @@ cdef device_span[const_uint8_t] _get_device_span(object obj) except *: obj.size) +cdef vector[vector[size_type]] _get_row_groups(object row_groups) except *: + """Convert Python per-source row-group indices to C++ vectors.""" + cdef vector[vector[size_type]] result + cdef vector[size_type] source_row_groups + for source in row_groups: + for row_group in source: + source_row_groups.push_back(row_group) + result.push_back(source_row_groups) + source_row_groups.clear() + return result + + cdef class HybridScanReader: """Experimental Parquet reader optimized for highly selective filters. @@ -835,4 +860,230 @@ cdef class HybridScanReader: return self.c_obj.get()[0].has_next_table_chunk() +cdef class HybridScanMultiFile: + """Experimental Hybrid Scan reader for multiple Parquet sources. + + The per-source inputs and outputs use source order. A row mask spans all + selected rows in source order, and then in row-group order within a source. + This API is experimental. + """ + + @staticmethod + def from_parquet_metadatas(object parquet_metadatas, ParquetReaderOptions options): + """Create a reader from one ``FileMetaData`` per Parquet source.""" + cdef HybridScanMultiFile reader = HybridScanMultiFile.__new__( + HybridScanMultiFile + ) + cdef vector[cpp_FileMetaData] metadatas + cdef object metadata + for metadata in parquet_metadatas: + if not isinstance(metadata, FileMetaData): + raise TypeError( + "parquet_metadatas must contain only FileMetaData objects" + ) + metadatas.push_back((metadata).c_obj) + if metadatas.empty(): + raise ValueError("parquet_metadatas must not be empty") + reader.c_obj = make_unique[cpp_hybrid_scan_multifile]( + host_span[const_FileMetaData]( + metadatas.data(), metadatas.size() + ), + options.c_obj, + ) + return reader + + def parquet_metadatas(self): + """Return one ``FileMetaData`` object per source.""" + cdef vector[cpp_FileMetaData] metadatas = ( + self.c_obj.get()[0].parquet_metadatas() + ) + cdef cpp_FileMetaData metadata + cdef list result = [] + for metadata in metadatas: + result.append(c_FileMetaData.from_cpp(metadata)) + return result + + def page_index_byte_ranges(self): + """Return the page-index byte range for each source.""" + cdef vector[byte_range_info] ranges = ( + self.c_obj.get()[0].page_index_byte_ranges() + ) + return [ByteRangeInfo(r.offset(), r.size()) for r in ranges] + + def setup_page_indexes(self, list page_index_bytes): + """Install one host page-index buffer for each source.""" + cdef vector[host_span[const_uint8_t]] spans + cdef const uint8_t[::1] page_index + for page_index in page_index_bytes: + if len(page_index) == 0: + spans.push_back( + host_span[const_uint8_t](0, 0) + ) + else: + spans.push_back( + host_span[const_uint8_t](&page_index[0], len(page_index)) + ) + self.c_obj.get()[0].setup_page_indexes( + host_span[const_host_span_const_uint8_t]( + spans.data(), spans.size() + ) + ) + + def all_row_groups(self, ParquetReaderOptions options): + """Return row-group indices for every source.""" + cdef vector[vector[size_type]] row_groups = ( + self.c_obj.get()[0].all_row_groups(options.c_obj) + ) + return [list(row_groups[source]) for source in range(row_groups.size())] + + def total_rows_in_row_groups(self, object row_group_indices): + """Return total selected rows across all sources.""" + cdef vector[vector[size_type]] indices = _get_row_groups( + row_group_indices + ) + return self.c_obj.get()[0].total_rows_in_row_groups( + host_span[const_vector_size_type]( + indices.data(), indices.size() + ) + ) + + def build_all_true_row_mask( + self, + object row_group_indices, + object stream=None, + DeviceMemoryResource mr=None, + ): + """Build an all-true BOOL8 mask spanning selected rows.""" + cdef vector[vector[size_type]] indices = _get_row_groups( + row_group_indices + ) + cdef Stream _stream = _get_stream(stream) + mr = _get_memory_resource(mr) + cdef unique_ptr[column] c_result = ( + self.c_obj.get()[0].build_all_true_row_mask( + host_span[const_vector_size_type]( + indices.data(), indices.size() + ), + _stream.view().value(), + mr.get_mr(), + ) + ) + return Column.from_libcudf(move(c_result), _stream, mr) + + def payload_column_chunks_byte_ranges( + self, + object row_group_indices, + Column row_mask, + cpp_use_data_page_mask mask_data_pages, + ParquetReaderOptions options, + object stream=None, + ): + """Plan payload page ranges, grouped by source.""" + cdef vector[vector[size_type]] indices = _get_row_groups( + row_group_indices + ) + cdef Stream _stream = _get_stream(stream) + cdef column_view mask_view = row_mask.view() + cdef vector[vector[byte_range_info]] ranges = ( + self.c_obj.get()[0].payload_column_chunks_byte_ranges( + host_span[const_vector_size_type]( + indices.data(), indices.size() + ), + mask_view, + mask_data_pages, + options.c_obj, + _stream.view().value(), + ) + ) + return [ + [ + ByteRangeInfo(byte_range.offset(), byte_range.size()) + for byte_range in ranges[source] + ] + for source in range(ranges.size()) + ] + + def setup_chunking_for_payload_columns( + self, + size_t chunk_read_limit, + size_t pass_read_limit, + object row_group_indices, + Column row_mask, + cpp_use_data_page_mask mask_data_pages, + list page_data_per_source, + ParquetReaderOptions options, + object stream=None, + DeviceMemoryResource mr=None, + ): + """Configure payload chunking from source-grouped device page data.""" + cdef vector[vector[size_type]] indices = _get_row_groups( + row_group_indices + ) + cdef vector[vector[device_span[const_uint8_t]]] source_spans + cdef vector[device_span[const_uint8_t]] spans + cdef object source_data + cdef object data + for source_data in page_data_per_source: + for data in source_data: + spans.push_back(_get_device_span(data)) + source_spans.push_back(spans) + spans.clear() + + self._stream = _get_stream(stream) + self.mr = _get_memory_resource(mr) + self._page_data_keepalive = page_data_per_source + cdef column_view mask_view = row_mask.view() + self.c_obj.get()[0].setup_chunking_for_payload_columns( + chunk_read_limit, + pass_read_limit, + host_span[const_vector_size_type]( + indices.data(), indices.size() + ), + mask_view, + mask_data_pages, + host_span[const_vector_device_span_const_uint8_t]( + source_spans.data(), + source_spans.size(), + ), + options.c_obj, + self._stream.view().value(), + self.mr.get_mr(), + ) + + def materialize_payload_columns_chunk(self, Column row_mask): + """Materialize the next configured payload output chunk.""" + cdef column_view mask_view = row_mask.view() + cdef table_with_metadata c_result = ( + self.c_obj.get()[0].materialize_payload_columns_chunk(mask_view) + ) + return TableWithMetadata.from_libcudf(c_result, self._stream, self.mr) + + def construct_row_group_passes( + self, object row_group_indices, size_t pass_read_limit + ): + """Partition per-source row groups into bounded-memory passes.""" + cdef vector[vector[size_type]] indices = _get_row_groups( + row_group_indices + ) + cdef vector[vector[vector[size_type]]] passes = ( + self.c_obj.get()[0].construct_row_group_passes( + host_span[const_vector_size_type]( + indices.data(), indices.size() + ), + pass_read_limit, + ) + ) + return [ + [ + list(row_groups) + for row_groups in passes[pass_index] + ] + for pass_index in range(passes.size()) + ] + + def has_next_table_chunk(self): + """Return whether a configured chunked read has output remaining.""" + return self.c_obj.get()[0].has_next_table_chunk() + + UseDataPageMask.__str__ = UseDataPageMask.__repr__ diff --git a/python/pylibcudf/pylibcudf/libcudf/io/hybrid_scan.pxd b/python/pylibcudf/pylibcudf/libcudf/io/hybrid_scan.pxd index 36201d545de6..2c62a19d007c 100644 --- a/python/pylibcudf/pylibcudf/libcudf/io/hybrid_scan.pxd +++ b/python/pylibcudf/pylibcudf/libcudf/io/hybrid_scan.pxd @@ -21,7 +21,12 @@ from rmm.librmm.memory_resource cimport device_async_resource_ref ctypedef const uint8_t const_uint8_t ctypedef const size_type const_size_type +ctypedef const FileMetaData const_FileMetaData ctypedef const device_span[const_uint8_t] const_device_span_const_uint8_t +ctypedef const vector[size_type] const_vector_size_type +ctypedef const vector[device_span[const_uint8_t]] const_vector_device_span_const_uint8_t +ctypedef host_span[const_uint8_t] host_span_const_uint8_t +ctypedef const host_span_const_uint8_t const_host_span_const_uint8_t cdef extern from "cudf/io/experimental/hybrid_scan.hpp" \ namespace "cudf::io::parquet::experimental" nogil: @@ -174,3 +179,66 @@ cdef extern from "cudf/io/experimental/hybrid_scan.hpp" \ ) except +libcudf_exception_handler bool has_next_table_chunk() except +libcudf_exception_handler + + +cdef extern from "cudf/io/experimental/hybrid_scan_multifile.hpp" \ + namespace "cudf::io::parquet::experimental" nogil: + + cdef cppclass hybrid_scan_multifile: + hybrid_scan_multifile( + host_span[const_FileMetaData] parquet_metadata, + const parquet_reader_options& options + ) except +libcudf_exception_handler + + vector[FileMetaData] parquet_metadatas() except +libcudf_exception_handler + + vector[byte_range_info] page_index_byte_ranges() except +libcudf_exception_handler + + void setup_page_indexes( + host_span[const_host_span_const_uint8_t] page_index_bytes + ) except +libcudf_exception_handler + + vector[vector[size_type]] all_row_groups( + const parquet_reader_options& options + ) except +libcudf_exception_handler + + size_type total_rows_in_row_groups( + host_span[const_vector_size_type] row_group_indices + ) except +libcudf_exception_handler + + unique_ptr[column] build_all_true_row_mask( + host_span[const_vector_size_type] row_group_indices, + cudaStream_t stream, + device_async_resource_ref mr + ) except +libcudf_exception_handler + + vector[vector[byte_range_info]] payload_column_chunks_byte_ranges( + host_span[const_vector_size_type] row_group_indices, + const column_view& row_mask, + use_data_page_mask mask_data_pages, + const parquet_reader_options& options, + cudaStream_t stream + ) except +libcudf_exception_handler + + void setup_chunking_for_payload_columns( + size_t chunk_read_limit, + size_t pass_read_limit, + host_span[const_vector_size_type] row_group_indices, + const column_view& row_mask, + use_data_page_mask mask_data_pages, + host_span[const_vector_device_span_const_uint8_t] page_data_per_source, + const parquet_reader_options& options, + cudaStream_t stream, + device_async_resource_ref mr + ) except +libcudf_exception_handler + + table_with_metadata materialize_payload_columns_chunk( + const column_view& row_mask + ) except +libcudf_exception_handler + + vector[vector[vector[size_type]]] construct_row_group_passes( + host_span[const_vector_size_type] row_group_indices, + size_t pass_read_limit + ) except +libcudf_exception_handler + + bool has_next_table_chunk() except +libcudf_exception_handler diff --git a/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py b/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py index 7bf3a19e1d13..2b631be985bb 100644 --- a/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py +++ b/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py @@ -781,3 +781,99 @@ def test_hybrid_scan_metadata_with_page_index( assert row_mask is not None assert row_mask.size() > 0 assert row_mask.type().id() == plc.types.TypeId.BOOL8 + + +@pytest.mark.parametrize("stream", [None, Stream()]) +def test_hybrid_scan_multifile_payload_page_scan( + tmp_path, + simple_parquet_table: pa.Table, + row_group_size: int, + stream: Stream | None, +) -> None: + """Read payload pages from multiple sources using the multifile reader.""" + paths = [tmp_path / f"source-{i}.parquet" for i in range(2)] + for index, path in enumerate(paths): + table = simple_parquet_table.slice( + index * 500, 500 + ) + pq.write_table( + table, + path, + row_group_size=row_group_size, + use_dictionary=True, + write_page_index=True, + ) + + source = plc.io.SourceInfo([str(path) for path in paths]) + options = plc.io.parquet.ParquetReaderOptions.builder(source).build() + options.set_column_names(["col1"]) + metadatas = plc.io.parquet_metadata.read_parquet_footers(source) + reader = plc.io.experimental.HybridScanMultiFile.from_parquet_metadatas( + metadatas, options + ) + + assert [metadata.num_rows for metadata in reader.parquet_metadatas()] == [ + 500, + 500, + ] + row_groups = reader.all_row_groups(options) + assert row_groups == [[0, 1], [0, 1]] + assert reader.total_rows_in_row_groups(row_groups) == 1000 + + page_ranges = reader.page_index_byte_ranges() + page_indexes = [] + for path, byte_range in zip(paths, page_ranges, strict=True): + with path.open("rb") as file: + file.seek(byte_range.offset) + page_indexes.append(file.read(byte_range.size)) + reader.setup_page_indexes(page_indexes) + + row_mask = reader.build_all_true_row_mask(row_groups, stream) + assert row_mask.size() == 1000 + + passes = reader.construct_row_group_passes(row_groups, 0) + assert passes == [row_groups] + page_ranges_per_source = reader.payload_column_chunks_byte_ranges( + passes[0], + row_mask, + UseDataPageMask.YES, + options, + stream, + ) + assert len(page_ranges_per_source) == len(paths) + assert all(ranges for ranges in page_ranges_per_source) + + page_data_per_source = [] + for path, byte_ranges in zip(paths, page_ranges_per_source, strict=True): + source_data = [] + with path.open("rb") as file: + for byte_range in byte_ranges: + file.seek(byte_range.offset) + source_data.append( + plc.gpumemoryview( + rmm.DeviceBuffer.to_device( + file.read(byte_range.size), + plc.utils._get_stream(stream), + ) + ) + ) + page_data_per_source.append(source_data) + + synchronize_stream(stream) + reader.setup_chunking_for_payload_columns( + 0, + 0, + passes[0], + row_mask, + UseDataPageMask.YES, + page_data_per_source, + options, + stream, + ) + output_rows = 0 + while reader.has_next_table_chunk(): + output = reader.materialize_payload_columns_chunk(row_mask) + assert output.tbl.num_columns() == 1 + output_rows += output.tbl.num_rows() + synchronize_stream(stream) + assert output_rows == 1000 From 9947bb2a63ffa73e567ee0b9446fe84f259b9dd4 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Thu, 20 Aug 2026 03:42:16 +0000 Subject: [PATCH 14/16] Revert unnecessary stuff --- .../cudf/io/experimental/hybrid_scan.hpp | 48 ---- .../io/parquet/experimental/hybrid_scan.cpp | 45 --- .../experimental/hybrid_scan_chunking.cu | 3 - .../parquet/experimental/hybrid_scan_impl.cpp | 263 ------------------ .../parquet/experimental/hybrid_scan_impl.hpp | 44 --- .../experimental/hybrid_scan_preprocess.cu | 37 --- .../io/experimental/hybrid_scan_common.cpp | 9 - .../io/experimental/hybrid_scan_common.hpp | 9 - .../io/experimental/hybrid_scan_composer.cpp | 9 +- 9 files changed, 4 insertions(+), 463 deletions(-) diff --git a/cpp/include/cudf/io/experimental/hybrid_scan.hpp b/cpp/include/cudf/io/experimental/hybrid_scan.hpp index e60b5baa24a1..bf8991b8557b 100644 --- a/cpp/include/cudf/io/experimental/hybrid_scan.hpp +++ b/cpp/include/cudf/io/experimental/hybrid_scan.hpp @@ -512,26 +512,6 @@ class hybrid_scan_reader { [[nodiscard]] std::vector payload_column_chunks_byte_ranges( std::span row_group_indices, parquet_reader_options const& options) const; - /** - * @brief Plan page-level payload byte ranges using the row mask - * - * When page masking cannot be used, returns the full payload column-chunk ranges. The resulting - * plan must be consumed exactly once by the matching page-data setup overload. - * - * @param row_group_indices Input row group indices - * @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 - */ - [[nodiscard]] std::vector payload_column_chunks_byte_ranges( - std::span 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 * @@ -642,34 +622,6 @@ class hybrid_scan_reader { cuda::stream_ref stream, rmm::device_async_resource_ref mr) const; - /** - * @brief Setup chunking for payload columns using page-level payload data - * - * The `page_data` must correspond to the byte ranges returned by the matching page-level - * `payload_column_chunks_byte_ranges` overload. - * - * @param chunk_read_limit Limit on total number of bytes to be returned per table chunk - * @param pass_read_limit Limit on the memory used for reading and decompressing data - * @param row_group_indices Input row group indices - * @param row_mask Boolean column indicating which rows need to be read - * @param mask_data_pages Whether to use the row mask to prune data pages - * @param page_data_per_source Device spans of page data returned by the page-level range plan, - * grouped by source - * @param options Parquet reader options - * @param stream CUDA stream used for device memory operations and kernel launches - * @param mr Device memory resource used to allocate the output table chunks - */ - void setup_chunking_for_payload_columns( - std::size_t chunk_read_limit, - std::size_t pass_read_limit, - std::span 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.cpp b/cpp/src/io/parquet/experimental/hybrid_scan.cpp index fa3e8bd08b69..a97135ca8480 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan.cpp +++ b/cpp/src/io/parquet/experimental/hybrid_scan.cpp @@ -217,24 +217,6 @@ hybrid_scan_reader::payload_column_chunks_byte_ranges(std::span return _impl->payload_column_chunks_byte_ranges(input_row_group_indices, options).first; } -std::vector hybrid_scan_reader::payload_column_chunks_byte_ranges( - std::span 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(); - - auto const input_row_group_indices = - std::vector>{{row_group_indices.begin(), row_group_indices.end()}}; - - return _impl - ->payload_column_chunks_byte_ranges( - input_row_group_indices, row_mask, mask_data_pages, options, stream) - .front(); -} - table_with_metadata hybrid_scan_reader::materialize_payload_columns( std::span row_group_indices, std::span const> column_chunk_data, @@ -346,33 +328,6 @@ void hybrid_scan_reader::setup_chunking_for_payload_columns( mr); } -void hybrid_scan_reader::setup_chunking_for_payload_columns( - std::size_t chunk_read_limit, - std::size_t pass_read_limit, - std::span 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(); - - auto const input_row_group_indices = - std::vector>{{row_group_indices.begin(), row_group_indices.end()}}; - - _impl->setup_chunking_for_payload_columns(chunk_read_limit, - pass_read_limit, - input_row_group_indices, - row_mask, - mask_data_pages, - page_data_per_source, - options, - stream, - mr); -} - table_with_metadata hybrid_scan_reader::materialize_payload_columns_chunk( cudf::column_view const& row_mask) const { diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu b/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu index 2568ef2abd20..970f790da3e3 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu +++ b/cpp/src/io/parquet/experimental/hybrid_scan_chunking.cu @@ -129,9 +129,6 @@ void hybrid_scan_reader_impl::setup_next_pass( set_pass_page_mask(data_page_mask); } - // 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 86c018120b7c..96ae63b61409 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 @@ -241,8 +240,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; @@ -268,8 +265,6 @@ void hybrid_scan_reader_impl::prepare_materialization(read_columns_mode read_col cuda::stream_ref 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); @@ -533,264 +528,6 @@ 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::payload_pages_byte_ranges( std::span const> 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 fc67a70584fc..a6588f5c2ee8 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_impl.hpp +++ b/cpp/src/io/parquet/experimental/hybrid_scan_impl.hpp @@ -200,13 +200,6 @@ 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::payload_pages_byte_ranges */ @@ -296,17 +289,6 @@ class hybrid_scan_reader_impl : public parquet::detail::reader_impl { cuda::stream_ref 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 @@ -361,26 +343,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. * @@ -653,12 +615,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_preprocess.cu b/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu index e880b5b2f2a7..f2c56007b49b 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu +++ b/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu @@ -196,43 +196,6 @@ void hybrid_scan_reader_impl::setup_compressed_data( auto& chunks = pass.chunks; - if (_sparse_page_io) { - CUDF_EXPECTS(_has_offset_index, "Sparse page I/O requires offset 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 7af6c81c16d5..f478c09e2eac 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 074e51fda19a..31eec8b23631 100644 --- a/cpp/tests/io/experimental/hybrid_scan_common.hpp +++ b/cpp/tests/io/experimental/hybrid_scan_common.hpp @@ -99,15 +99,6 @@ void setup_page_indexes(cudf::io::parquet::experimental::hybrid_scan_multifile c cuda::stream_ref 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_composer.cpp b/cpp/tests/io/experimental/hybrid_scan_composer.cpp index 36fccc1d30e9..3f3d7e0c3522 100644 --- a/cpp/tests/io/experimental/hybrid_scan_composer.cpp +++ b/cpp/tests/io/experimental/hybrid_scan_composer.cpp @@ -345,8 +345,8 @@ std::tuple, std::unique_ptr> sparse_ch ? reader->build_row_mask_with_page_index_stats(current_row_group_indices, options, stream, mr) : reader->build_all_true_row_mask(current_row_group_indices, stream, mr); - auto filter_tables = std::vector>{}; - auto payload_tables = std::vector>{}; + auto filter_tables = std::vector>{}; + auto payload_tables = std::vector>{}; std::size_t rows_materialized = 0; auto const materialize_pass = [&](cudf::host_span row_group_indices) { auto const rows_in_pass = reader->total_rows_in_row_groups(row_group_indices); @@ -390,9 +390,8 @@ std::tuple, std::unique_ptr> sparse_ch cudf::io::parquet::fetch_byte_ranges_to_device_async( datasource, payload_page_ranges, stream, mr); payload_tasks.get(); - auto const page_data_per_source = - std::vector>>{{payload_data.begin(), - payload_data.end()}}; + auto const page_data_per_source = std::vector>>{ + {payload_data.begin(), payload_data.end()}}; reader->setup_chunking_for_payload_columns( 1024, 10240, From 91e56265c1c9546b11dd2f3c72368f3f5c75d7f8 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Thu, 20 Aug 2026 03:43:02 +0000 Subject: [PATCH 15/16] More revert --- .../io/experimental/hybrid_scan_multifile.hpp | 27 ------------------- 1 file changed, 27 deletions(-) diff --git a/cpp/include/cudf/io/experimental/hybrid_scan_multifile.hpp b/cpp/include/cudf/io/experimental/hybrid_scan_multifile.hpp index 91e4bfd01120..c75fa3d186d3 100644 --- a/cpp/include/cudf/io/experimental/hybrid_scan_multifile.hpp +++ b/cpp/include/cudf/io/experimental/hybrid_scan_multifile.hpp @@ -478,33 +478,6 @@ class hybrid_scan_multifile { cuda::stream_ref 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 From 9033aa9ce15ee5202582c593140fab32b2f2aa80 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Thu, 20 Aug 2026 03:45:21 +0000 Subject: [PATCH 16/16] Revert --- .../io/experimental/hybrid_scan_multifile_test.cpp | 11 ++--------- 1 file changed, 2 insertions(+), 9 deletions(-) diff --git a/cpp/tests/io/experimental/hybrid_scan_multifile_test.cpp b/cpp/tests/io/experimental/hybrid_scan_multifile_test.cpp index 7b28e090aded..157546ee9518 100644 --- a/cpp/tests/io/experimental/hybrid_scan_multifile_test.cpp +++ b/cpp/tests/io/experimental/hybrid_scan_multifile_test.cpp @@ -48,9 +48,8 @@ using cudf::io::parquet::experimental::use_data_page_mask; */ 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); @@ -116,12 +115,6 @@ void test_hybrid_scan_multifile(std::vector const& columns, 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