From a18c546391ce045d099dbe9860ec27d343095e44 Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Sun, 14 Jun 2026 14:37:28 +0200 Subject: [PATCH 01/39] Add test for filtering row groups with Bloom filters on real data --- .../experimental/hybrid_scan_filters_test.cpp | 52 +++++++++++++++++++ 1 file changed, 52 insertions(+) diff --git a/cpp/tests/io/experimental/hybrid_scan_filters_test.cpp b/cpp/tests/io/experimental/hybrid_scan_filters_test.cpp index 10960c185ef9..ce0c0a20ec7e 100644 --- a/cpp/tests/io/experimental/hybrid_scan_filters_test.cpp +++ b/cpp/tests/io/experimental/hybrid_scan_filters_test.cpp @@ -19,11 +19,15 @@ #include +#include +#include + #include #include #include #include #include +#include #include namespace { @@ -1331,6 +1335,54 @@ TEST_F(HybridScanFiltersTest, FilterRowGroupsWithDictionary) } } +TEST_F(HybridScanFiltersTest, FilterRowGroupsWithBloomFiltersRealData) +{ + auto const stream = cudf::get_default_stream(); + auto aligned_mr = rmm::mr::aligned_resource_adaptor(cudf::get_current_device_resource_ref(), + rmm::CUDA_ALLOCATION_ALIGNMENT); + + auto literal_value = cudf::string_scalar("Did not like the color", true, stream); + auto literal = cudf::ast::literal(literal_value); + auto column_ref = cudf::ast::column_name_reference("r_reason_desc"); + auto filter = cudf::ast::operation(cudf::ast::ast_operator::EQUAL, column_ref, literal); + auto options = cudf::io::parquet_reader_options::builder().filter(filter).build(); + + auto const parquet_filepath = + std::filesystem::path{__FILE__} + .parent_path() + .parent_path() + .parent_path() + .parent_path() + .parent_path() / + "python/cudf/cudf/tests/data/parquet/bloom_filter_alignment.parquet"; + auto const datasource_ptr = cudf::io::datasource::create(parquet_filepath.string()); + auto datasource = std::ref(*datasource_ptr); + auto const footer = cudf::io::parquet::fetch_footer_to_host(datasource); + auto const reader = + std::make_unique(*footer, options); + + auto row_group_indices = reader->all_row_groups(options); + auto current_row_group_indices = cudf::host_span{row_group_indices}; + auto const expected_row_groups = std::vector{0}; + ASSERT_EQ(row_group_indices, expected_row_groups); + + auto const bloom_filter_byte_ranges = + std::get<0>(reader->secondary_filters_byte_ranges(current_row_group_indices, options)); + ASSERT_FALSE(bloom_filter_byte_ranges.empty()); + + auto [bloom_filter_buffers, bloom_filter_data, bloom_filter_tasks] = + cudf::io::parquet::fetch_byte_ranges_to_device_async( + datasource, bloom_filter_byte_ranges, stream, aligned_mr); + bloom_filter_tasks.get(); + + auto const surviving_row_groups = + reader->filter_row_groups_with_bloom_filters( + bloom_filter_data, current_row_group_indices, options, stream); + + EXPECT_EQ(surviving_row_groups, expected_row_groups) + << "BUG: hybrid scan bloom filtering pruned a row group containing the queried value."; +} + template struct RowGroupFilteringWithDictTest : public HybridScanFiltersTest {}; From b6097558c2d6e20efda34cc84a1a7797c6fa91e1 Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Tue, 16 Jun 2026 11:03:06 +0200 Subject: [PATCH 02/39] [WIP] Implement debug logging for Bloom filter data in parquet reader; remove outdated test for Bloom filters with real data Need to rebuild all the code for A/B test --- cpp/src/io/parquet/bloom_filter_reader.cu | 19 ++++++ .../experimental/hybrid_scan_filters_test.cpp | 48 -------------- .../tests/io/test_experimental_hybrid_scan.py | 66 +++++++++++++++++++ 3 files changed, 85 insertions(+), 48 deletions(-) diff --git a/cpp/src/io/parquet/bloom_filter_reader.cu b/cpp/src/io/parquet/bloom_filter_reader.cu index 2160b7179cc7..02791c2d7c6b 100644 --- a/cpp/src/io/parquet/bloom_filter_reader.cu +++ b/cpp/src/io/parquet/bloom_filter_reader.cu @@ -29,6 +29,7 @@ #include #include +#include #include #include #include @@ -516,6 +517,24 @@ std::optional>> aggregate_reader_metadata::ap static_cast(total_row_groups), bloom_filter_col_schemas.size()}; + // [bloom-dbg] dev-only: remove before merge. Hexdump the exact bytes each path feeds bloom filter + // construction. read_parquet strips the BloomFilterHeader; hybrid scan currently does not, so its + // spans are 16 bytes longer (header + bitset) and the caster misreads the header as a filter block. + for (std::size_t dbg_i = 0; dbg_i < bloom_filter_data.size(); ++dbg_i) { + auto const& dbg_span = bloom_filter_data[dbg_i]; + std::vector dbg_host(dbg_span.size()); + if (not dbg_span.empty()) { + cudaMemcpyAsync( + dbg_host.data(), dbg_span.data(), dbg_span.size(), cudaMemcpyDeviceToHost, stream.value()); + stream.synchronize(); + } + std::fprintf(stderr, "[bloom-dbg] bloom span[%zu] size=%zu raw=", dbg_i, dbg_span.size()); + for (auto const byte : dbg_host) { + std::fprintf(stderr, "%02x", static_cast(byte)); + } + std::fprintf(stderr, "\n"); + } + // Converts bloom filter membership for equality predicate columns to a table // containing a column for each `col[i] == literal` predicate to be evaluated. // The table contains #sources * #column_chunks_per_src rows. diff --git a/cpp/tests/io/experimental/hybrid_scan_filters_test.cpp b/cpp/tests/io/experimental/hybrid_scan_filters_test.cpp index ce0c0a20ec7e..be7f60736430 100644 --- a/cpp/tests/io/experimental/hybrid_scan_filters_test.cpp +++ b/cpp/tests/io/experimental/hybrid_scan_filters_test.cpp @@ -1335,54 +1335,6 @@ TEST_F(HybridScanFiltersTest, FilterRowGroupsWithDictionary) } } -TEST_F(HybridScanFiltersTest, FilterRowGroupsWithBloomFiltersRealData) -{ - auto const stream = cudf::get_default_stream(); - auto aligned_mr = rmm::mr::aligned_resource_adaptor(cudf::get_current_device_resource_ref(), - rmm::CUDA_ALLOCATION_ALIGNMENT); - - auto literal_value = cudf::string_scalar("Did not like the color", true, stream); - auto literal = cudf::ast::literal(literal_value); - auto column_ref = cudf::ast::column_name_reference("r_reason_desc"); - auto filter = cudf::ast::operation(cudf::ast::ast_operator::EQUAL, column_ref, literal); - auto options = cudf::io::parquet_reader_options::builder().filter(filter).build(); - - auto const parquet_filepath = - std::filesystem::path{__FILE__} - .parent_path() - .parent_path() - .parent_path() - .parent_path() - .parent_path() / - "python/cudf/cudf/tests/data/parquet/bloom_filter_alignment.parquet"; - auto const datasource_ptr = cudf::io::datasource::create(parquet_filepath.string()); - auto datasource = std::ref(*datasource_ptr); - auto const footer = cudf::io::parquet::fetch_footer_to_host(datasource); - auto const reader = - std::make_unique(*footer, options); - - auto row_group_indices = reader->all_row_groups(options); - auto current_row_group_indices = cudf::host_span{row_group_indices}; - auto const expected_row_groups = std::vector{0}; - ASSERT_EQ(row_group_indices, expected_row_groups); - - auto const bloom_filter_byte_ranges = - std::get<0>(reader->secondary_filters_byte_ranges(current_row_group_indices, options)); - ASSERT_FALSE(bloom_filter_byte_ranges.empty()); - - auto [bloom_filter_buffers, bloom_filter_data, bloom_filter_tasks] = - cudf::io::parquet::fetch_byte_ranges_to_device_async( - datasource, bloom_filter_byte_ranges, stream, aligned_mr); - bloom_filter_tasks.get(); - - auto const surviving_row_groups = - reader->filter_row_groups_with_bloom_filters( - bloom_filter_data, current_row_group_indices, options, stream); - - EXPECT_EQ(surviving_row_groups, expected_row_groups) - << "BUG: hybrid scan bloom filtering pruned a row group containing the queried value."; -} - template struct RowGroupFilteringWithDictTest : public HybridScanFiltersTest {}; diff --git a/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py b/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py index 74f467f16193..db21bc0343db 100644 --- a/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py +++ b/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py @@ -1,6 +1,7 @@ # SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: Apache-2.0 import io +from pathlib import Path import pyarrow as pa import pyarrow.parquet as pq @@ -251,6 +252,71 @@ def test_hybrid_scan_secondary_filters_byte_ranges( assert isinstance(dict_ranges, list) +@pytest.mark.xfail( + reason="hybrid scan does not strip the Parquet BloomFilterHeader before probing", + strict=True, +) +def test_hybrid_scan_bloom_filter_matches_read_parquet(): + """Hybrid-scan bloom filtering must keep the same row group as ``read_parquet``. + + pyarrow cannot write bloom filters, so this reads the committed DuckDB-written fixture + (bloom filter on ``r_reason_desc``; the value "Did not like the color" is present). The + hybrid-scan path forwards the fetched bytes (``BloomFilterHeader`` + bitset) to the filter + without stripping the header, so it wrongly prunes the matching row group. + """ + fixture = ( + Path(__file__).parents[4] + / "python/cudf/cudf/tests/data/parquet/bloom_filter_alignment.parquet" + ) + if not fixture.exists(): + pytest.skip(f"bloom fixture not found: {fixture}") + data = fixture.read_bytes() + + bloom_filter = Operation( + ASTOperator.EQUAL, + ColumnNameReference("r_reason_desc"), + Literal(plc.Scalar.from_arrow(pa.scalar("Did not like the color"))), + ) + + def make_options(): + options = plc.io.parquet.ParquetReaderOptions.builder( + plc.io.SourceInfo([io.BytesIO(data)]) + ).build() + options.set_filter(bloom_filter) + return options + + # A: standard read_parquet keeps the only row group after bloom filtering. + table_w_meta = plc.io.parquet.read_parquet(make_options()) + assert table_w_meta.num_input_row_groups == 1 + assert table_w_meta.num_row_groups_after_bloom_filter == 1 + + # B: hybrid scan should keep the same row group. + options = make_options() + suffix = 8 # 4-byte footer length + "PAR1" + mv = memoryview(data) + footer_size = int.from_bytes(mv[-suffix:-4], byteorder="little") + reader = HybridScanReader(mv[-suffix - footer_size : -suffix], options) + + row_groups = reader.all_row_groups(options) + bloom_ranges, _ = reader.secondary_filters_byte_ranges(row_groups, options) + assert bloom_ranges # the equality predicate makes r_reason_desc bloom-eligible + + stream = plc.utils._get_stream(None) + bloom_data = [ + plc.gpumemoryview( + rmm.DeviceBuffer.to_device(data[r.offset : r.offset + r.size], stream) + ) + for r in bloom_ranges + ] + synchronize_stream(None) + surviving = reader.filter_row_groups_with_bloom_filters( + bloom_data, row_groups, options + ) + + # The queried value is present, so the hybrid path must match read_parquet (row group kept). + assert surviving == row_groups == [0] + + def test_hybrid_scan_column_chunk_byte_ranges( simple_hybrid_scan_reader: HybridScanReader, simple_parquet_options: plc.io.parquet.ParquetReaderOptions, From a32a95ffe76b96d011761b7066a93feef49847dc Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Tue, 16 Jun 2026 14:13:17 +0200 Subject: [PATCH 03/39] Refactor hybrid scan bloom filter test to clarify A/B comparison with read_parquet Updated the test for hybrid scan bloom filtering to explicitly compare its behavior against the read_parquet method. Enhanced debug logging to provide detailed output of the bloom filter data being processed, ensuring consistency in row group retention between the two methods. Adjusted comments and structure for better readability and understanding of the test's purpose. --- cpp/src/io/parquet/bloom_filter_reader.cu | 7 +- .../tests/io/test_experimental_hybrid_scan.py | 71 +++++++++++++++---- 2 files changed, 60 insertions(+), 18 deletions(-) diff --git a/cpp/src/io/parquet/bloom_filter_reader.cu b/cpp/src/io/parquet/bloom_filter_reader.cu index 02791c2d7c6b..748c9f424a28 100644 --- a/cpp/src/io/parquet/bloom_filter_reader.cu +++ b/cpp/src/io/parquet/bloom_filter_reader.cu @@ -517,9 +517,10 @@ std::optional>> aggregate_reader_metadata::ap static_cast(total_row_groups), bloom_filter_col_schemas.size()}; - // [bloom-dbg] dev-only: remove before merge. Hexdump the exact bytes each path feeds bloom filter - // construction. read_parquet strips the BloomFilterHeader; hybrid scan currently does not, so its - // spans are 16 bytes longer (header + bitset) and the caster misreads the header as a filter block. + // [bloom-dbg] dev-only: remove before merge. Hexdump the exact bytes each path feeds bloom + // filter construction. read_parquet strips the BloomFilterHeader; hybrid scan currently does not, + // so its spans are 16 bytes longer (header + bitset) and the caster misreads the header as a + // filter block. for (std::size_t dbg_i = 0; dbg_i < bloom_filter_data.size(); ++dbg_i) { auto const& dbg_span = bloom_filter_data[dbg_i]; std::vector dbg_host(dbg_span.size()); diff --git a/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py b/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py index db21bc0343db..ae418144566d 100644 --- a/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py +++ b/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py @@ -257,12 +257,12 @@ def test_hybrid_scan_secondary_filters_byte_ranges( strict=True, ) def test_hybrid_scan_bloom_filter_matches_read_parquet(): - """Hybrid-scan bloom filtering must keep the same row group as ``read_parquet``. + """A/B test hybrid scan bloom filtering vs read_parquet. - pyarrow cannot write bloom filters, so this reads the committed DuckDB-written fixture - (bloom filter on ``r_reason_desc``; the value "Did not like the color" is present). The - hybrid-scan path forwards the fetched bytes (``BloomFilterHeader`` + bitset) to the filter - without stripping the header, so it wrongly prunes the matching row group. + A: read_parquet path (libcudf strips the BloomFilterHeader before probing) + B: hybrid-scan path (fetches BloomFilterHeader + bitset from secondary_filters_byte_ranges) + + The hybrid path must keep the same row groups as read_parquet. """ fixture = ( Path(__file__).parents[4] @@ -271,12 +271,14 @@ def test_hybrid_scan_bloom_filter_matches_read_parquet(): if not fixture.exists(): pytest.skip(f"bloom fixture not found: {fixture}") data = fixture.read_bytes() + column_name = "r_reason_desc" + query_value = "Did not like the color" + expected_row_groups = list(range(pq.ParquetFile(fixture).metadata.num_row_groups)) - bloom_filter = Operation( - ASTOperator.EQUAL, - ColumnNameReference("r_reason_desc"), - Literal(plc.Scalar.from_arrow(pa.scalar("Did not like the color"))), - ) + literal_value = plc.Scalar.from_arrow(pa.scalar(query_value)) + literal = Literal(literal_value) + column_ref = ColumnNameReference(column_name) + bloom_filter = Operation(ASTOperator.EQUAL, column_ref, literal) def make_options(): options = plc.io.parquet.ParquetReaderOptions.builder( @@ -285,12 +287,43 @@ def make_options(): options.set_filter(bloom_filter) return options + def make_path_options(): + options = plc.io.parquet.ParquetReaderOptions.builder( + plc.io.SourceInfo([fixture]) + ).build() + options.set_filter(bloom_filter) + return options + # A: standard read_parquet keeps the only row group after bloom filtering. - table_w_meta = plc.io.parquet.read_parquet(make_options()) - assert table_w_meta.num_input_row_groups == 1 - assert table_w_meta.num_row_groups_after_bloom_filter == 1 + print( + f"\n[bloom-demo] fixture={fixture.name} predicate={column_name} == {query_value!r}\n", + flush=True, + ) + print( + "[bloom-demo] A: read_parquet path (libcudf strips the " + "BloomFilterHeader before probing)", + flush=True, + ) + table_w_meta = plc.io.parquet.read_parquet(make_path_options()) + print( + "[bloom-demo] A result: " + f"rows={table_w_meta.tbl.num_rows()} " + f"num_input_row_groups={table_w_meta.num_input_row_groups} " + f"num_row_groups_after_bloom_filter=" + f"{table_w_meta.num_row_groups_after_bloom_filter}\n", + flush=True, + ) + assert table_w_meta.num_input_row_groups == len(expected_row_groups) + assert table_w_meta.num_row_groups_after_bloom_filter == len( + expected_row_groups + ) # B: hybrid scan should keep the same row group. + print( + "[bloom-demo] B: hybrid-scan path (fetches " + "BloomFilterHeader + bitset from secondary_filters_byte_ranges)", + flush=True, + ) options = make_options() suffix = 8 # 4-byte footer length + "PAR1" mv = memoryview(data) @@ -300,6 +333,13 @@ def make_options(): row_groups = reader.all_row_groups(options) bloom_ranges, _ = reader.secondary_filters_byte_ranges(row_groups, options) assert bloom_ranges # the equality predicate makes r_reason_desc bloom-eligible + for idx, r in enumerate(bloom_ranges): + raw = data[r.offset : r.offset + r.size] + print( + f"[bloom-demo] B fetched range[{idx}]: " + f"offset={r.offset} size={r.size} raw={raw.hex()}", + flush=True, + ) stream = plc.utils._get_stream(None) bloom_data = [ @@ -312,9 +352,10 @@ def make_options(): surviving = reader.filter_row_groups_with_bloom_filters( bloom_data, row_groups, options ) + print(f"[bloom-demo] B result: surviving_row_groups={surviving}", flush=True) - # The queried value is present, so the hybrid path must match read_parquet (row group kept). - assert surviving == row_groups == [0] + # The queried value is present, so the hybrid path must match read_parquet. + assert surviving == row_groups == expected_row_groups def test_hybrid_scan_column_chunk_byte_ranges( From 03c18343bcd29f581203f1473cd287c6dad9bdb8 Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Tue, 16 Jun 2026 17:00:37 +0200 Subject: [PATCH 04/39] Enhance hybrid scan bloom filter test with parameterized fixtures Refactored the hybrid scan bloom filter test to utilize parameterized fixtures for improved flexibility and clarity. This change allows for multiple test cases to be run with different input values, enhancing the robustness of the test against various scenarios. Updated the test structure and comments for better readability and understanding of the A/B comparison with the read_parquet method. --- .../tests/io/test_experimental_hybrid_scan.py | 24 +++++++++++++++---- 1 file changed, 20 insertions(+), 4 deletions(-) diff --git a/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py b/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py index ae418144566d..b58a76b198fb 100644 --- a/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py +++ b/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py @@ -256,7 +256,24 @@ def test_hybrid_scan_secondary_filters_byte_ranges( reason="hybrid scan does not strip the Parquet BloomFilterHeader before probing", strict=True, ) -def test_hybrid_scan_bloom_filter_matches_read_parquet(): +@pytest.mark.parametrize( + "fixture_name,column_name,query_value", + [ + ( + "bloom_filter_alignment.parquet", + "r_reason_desc", + "Did not like the color", + ), + ( + "mixed_card_ndv_100_bf_fpp0.1_nostats.snappy.parquet", + "str", + "tbRpIyVJZX", + ), + ], +) +def test_hybrid_scan_bloom_filter_matches_read_parquet( + fixture_name, column_name, query_value +): """A/B test hybrid scan bloom filtering vs read_parquet. A: read_parquet path (libcudf strips the BloomFilterHeader before probing) @@ -266,13 +283,12 @@ def test_hybrid_scan_bloom_filter_matches_read_parquet(): """ fixture = ( Path(__file__).parents[4] - / "python/cudf/cudf/tests/data/parquet/bloom_filter_alignment.parquet" + / "python/cudf/cudf/tests/data/parquet" + / fixture_name ) if not fixture.exists(): pytest.skip(f"bloom fixture not found: {fixture}") data = fixture.read_bytes() - column_name = "r_reason_desc" - query_value = "Did not like the color" expected_row_groups = list(range(pq.ParquetFile(fixture).metadata.num_row_groups)) literal_value = plc.Scalar.from_arrow(pa.scalar(query_value)) From df8fc3301850156a92fe009be4d94bd0f5b1f75e Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Wed, 24 Jun 2026 16:14:10 +0200 Subject: [PATCH 05/39] Refactor bloom filter handling in hybrid scan to improve performance and clarity Updated the hybrid scan implementation to replace the previous method of fetching bloom filter data with a new function, `fetch_bloom_filters_to_device_async`, which strips headers and ensures 32-byte alignment. This change enhances the efficiency of bloom filter processing and simplifies the codebase. Additionally, updated related tests to reflect these changes and ensure consistency in behavior across different methods of bloom filter handling. --- .../hybrid_scan/hybrid_scan_composer.cpp | 10 +- .../hybrid_scan_io/hybrid_scan_composer.cpp | 9 +- cpp/include/cudf/io/parquet_io_utils.hpp | 21 +++ cpp/src/io/parquet/bloom_filter_reader.cu | 54 +++--- .../io/parquet/io_utils/parquet_io_utils.cpp | 167 ++++++++++++++++++ cpp/src/io/parquet/reader_impl_helpers.hpp | 18 ++ .../io/experimental/hybrid_scan_composer.cpp | 9 +- .../pylibcudf/io/experimental/hybrid_scan.pyi | 10 +- .../pylibcudf/io/experimental/hybrid_scan.pyx | 137 +++++++++++++- .../pylibcudf/libcudf/io/hybrid_scan.pxd | 43 +++++ .../tests/io/test_experimental_hybrid_scan.py | 26 +-- 11 files changed, 449 insertions(+), 55 deletions(-) diff --git a/cpp/benchmarks/io/parquet/experimental/hybrid_scan/hybrid_scan_composer.cpp b/cpp/benchmarks/io/parquet/experimental/hybrid_scan/hybrid_scan_composer.cpp index 13b9c28914ae..98db682a709c 100644 --- a/cpp/benchmarks/io/parquet/experimental/hybrid_scan/hybrid_scan_composer.cpp +++ b/cpp/benchmarks/io/parquet/experimental/hybrid_scan/hybrid_scan_composer.cpp @@ -12,8 +12,6 @@ #include #include -#include - #include #include @@ -104,12 +102,10 @@ std::vector apply_row_group_filters( if (filters.contains(hybrid_scan_filter_type::ROW_GROUPS_WITH_BLOOM_FILTERS) and bloom_filter_byte_ranges.size()) { - // Fetch 32-byte aligned bloom filter data buffers from the input file buffer - auto constexpr bloom_filter_alignment = rmm::CUDA_ALLOCATION_ALIGNMENT; - auto aligned_mr = rmm::mr::aligned_resource_adaptor(mr, bloom_filter_alignment); + // Fetch the header-stripped, 32-byte-aligned bloom filter bitsets from the input file auto [bloom_filter_buffers, bloom_filter_data, bloom_read_tasks] = - cudf::io::parquet::fetch_byte_ranges_to_device_async( - datasource, bloom_filter_byte_ranges, stream, aligned_mr); + cudf::io::parquet::fetch_bloom_filters_to_device_async( + datasource, bloom_filter_byte_ranges, stream, mr); bloom_read_tasks.get(); bloom_filtered_row_groups = reader.filter_row_groups_with_bloom_filters( diff --git a/cpp/examples/hybrid_scan_io/hybrid_scan_composer.cpp b/cpp/examples/hybrid_scan_io/hybrid_scan_composer.cpp index a26839d9ffd6..fdc6f3735c86 100644 --- a/cpp/examples/hybrid_scan_io/hybrid_scan_composer.cpp +++ b/cpp/examples/hybrid_scan_io/hybrid_scan_composer.cpp @@ -16,8 +16,6 @@ #include #include -#include - #include #include @@ -175,14 +173,13 @@ std::vector apply_row_group_filters( bloom_filtered_row_group_indices.reserve(current_row_group_indices.size()); if (filters.contains(hybrid_scan_filter_type::ROW_GROUPS_WITH_BLOOM_FILTERS) and bloom_filter_byte_ranges.size()) { - // Fetch 32-byte aligned bloom filter data buffers from the input file buffer - auto constexpr bloom_filter_alignment = rmm::CUDA_ALLOCATION_ALIGNMENT; - auto aligned_mr = rmm::mr::aligned_resource_adaptor(temp_mr, bloom_filter_alignment); if (verbose) { std::cout << "READER: Filter row groups with bloom filters...\n"; } timer.reset(); nvtxRangePush("fetch_bloom_filter_byte_ranges"); + // Fetch the header-stripped, 32-byte-aligned bloom filter bitsets from the input file auto [bloom_filter_buffers, bloom_filter_data, bloom_read_tasks] = - fetch_byte_ranges_async(datasource, bloom_filter_byte_ranges, stream, aligned_mr); + cudf::io::parquet::fetch_bloom_filters_to_device_async( + datasource, bloom_filter_byte_ranges, stream, temp_mr); bloom_read_tasks.get(); nvtxRangePop(); diff --git a/cpp/include/cudf/io/parquet_io_utils.hpp b/cpp/include/cudf/io/parquet_io_utils.hpp index ba557f1f8ca6..9c0c759d95aa 100644 --- a/cpp/include/cudf/io/parquet_io_utils.hpp +++ b/cpp/include/cudf/io/parquet_io_utils.hpp @@ -151,6 +151,27 @@ fetch_byte_ranges_to_device_async( rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr); +/** + * @brief Fetches Parquet bloom filter bitsets from a datasource into device buffers + * + * @ingroup io_utils + * + * @param datasource Input datasource + * @param bloom_filter_byte_ranges Byte ranges of complete bloom filters to fetch, must span a complete bloom filter + * @param stream CUDA stream + * @param mr Device memory resource + * + * @return A tuple containing the device buffers, the device spans of the bitset data (one per input + * range; empty for ranges without a bloom filter), and a future to wait on the read tasks + */ +std::tuple, + std::vector>, + std::future> +fetch_bloom_filters_to_device_async(cudf::io::datasource& datasource, + cudf::host_span bloom_filter_byte_ranges, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr); + /** @} */ // end of group } // namespace io::parquet } // namespace CUDF_EXPORT cudf diff --git a/cpp/src/io/parquet/bloom_filter_reader.cu b/cpp/src/io/parquet/bloom_filter_reader.cu index 748c9f424a28..edc00f79a003 100644 --- a/cpp/src/io/parquet/bloom_filter_reader.cu +++ b/cpp/src/io/parquet/bloom_filter_reader.cu @@ -33,6 +33,7 @@ #include #include #include +#include namespace cudf::io::parquet::detail { namespace { @@ -305,7 +306,6 @@ void read_bloom_filter_data(host_span const> sources cuco::extent, cuco::thread_scope_thread, policy_type>::filter_block_type); - auto constexpr words_per_block = policy_type::words_per_block; // Read tasks for bloom filter data std::vector> read_tasks; @@ -323,39 +323,28 @@ void read_bloom_filter_data(host_span const> sources // Read bloom filter iff present auto const bloom_filter_offset = bloom_filter_offsets[chunk].value(); - // If Bloom filter size (header + bitset) is available, just read the entire thing. - // Else just read 256 bytes which will contain the entire header and may contain the - // entire bitset as well. - auto constexpr bloom_filter_size_guess = 256; + // If the bloom filter size (header + bitset) is available, read the entire thing. Else read + // the max header size, which contains the entire header and may contain the entire bitset. auto const initial_read_size = - static_cast(bloom_filter_sizes[chunk].value_or(bloom_filter_size_guess)); + static_cast(bloom_filter_sizes[chunk].value_or(bloom_filter_header_max_size)); // Read an initial buffer from source auto& source = sources[chunk_source_map[chunk]]; auto buffer = source->host_read(bloom_filter_offset, initial_read_size); - // Deserialize the Bloom filter header from the buffer. - BloomFilterHeader header; - CompactProtocolReader cp{buffer->data(), buffer->size()}; - cp.read(&header); - - // Check if the bloom filter header is valid. - auto const is_header_valid = - (header.num_bytes % words_per_block) == 0 and - header.compression.compression == BloomFilterCompression::UNCOMPRESSED and - header.algorithm.algorithm == BloomFilterAlgorithm::SPLIT_BLOCK and - header.hash.hash == BloomFilterHash::XXHASH; + // Deserialize and validate the bloom filter header from the buffer. + auto const header_info = parse_bloom_filter_header({buffer->data(), buffer->size()}); // Do not read if the bloom filter is invalid - if (not is_header_valid) { + if (not header_info.has_value()) { bloom_filter_data[chunk] = {}; CUDF_LOG_WARN("Encountered an invalid bloom filter header. Skipping"); return; } - // Bloom filter header size - auto const bloom_filter_header_size = static_cast(cp.bytecount()); - auto const bitset_size = static_cast(header.num_bytes); + // Bloom filter header and bitset sizes + auto const bloom_filter_header_size = header_info->first; + auto const bitset_size = header_info->second; // Check if we already read in the filter bitset in the initial read. if (initial_read_size >= bloom_filter_header_size + bitset_size) { @@ -406,6 +395,29 @@ void read_bloom_filter_data(host_span const> sources } // namespace +std::optional> parse_bloom_filter_header( + host_span bytes) +{ + using policy_type = arrow_filter_policy; + auto constexpr words_per_block = policy_type::words_per_block; + + // Deserialize the bloom filter header from the front of the buffer + BloomFilterHeader header; + CompactProtocolReader cp{bytes.data(), bytes.size()}; + cp.read(&header); + + // Check if the bloom filter header is valid + auto const is_header_valid = + (header.num_bytes % words_per_block) == 0 and + header.compression.compression == BloomFilterCompression::UNCOMPRESSED and + header.algorithm.algorithm == BloomFilterAlgorithm::SPLIT_BLOCK and + header.hash.hash == BloomFilterHash::XXHASH; + if (not is_header_valid) { return std::nullopt; } + + return std::pair{static_cast(cp.bytecount()), + static_cast(header.num_bytes)}; +} + std::size_t aggregate_reader_metadata::get_bloom_filter_alignment() const { // Required alignment: diff --git a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp index 61ca6c7f4038..ea2efff68ff6 100644 --- a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp +++ b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp @@ -5,6 +5,7 @@ #include "io/comp/common.hpp" #include "io/parquet/parquet_common.hpp" +#include "io/parquet/reader_impl_helpers.hpp" #include #include @@ -25,15 +26,19 @@ #include #include +#include #include #include #include +#include #include #include +#include #include #include #include #include +#include #include /** @@ -87,6 +92,71 @@ auto dispatch_fetch_tasks(std::size_t num_sources, Task fetch_task) return results; } +/** + * @brief Dispatches a grouped fetch task for each `(group_idx, item_idx)` and collects the results + * based on the total number of items across all groups. + * + * @tparam Task Callable invocable as `fetch_task(std::size_t group_idx, std::size_t item_idx)` + * @param items_per_group Number of items in each group + * @param fetch_task Task to run for each `(group_idx, item_idx)` + * @return Vector (one per group) of vectors of results, in input order + */ +template +auto dispatch_fetch_tasks(cudf::host_span items_per_group, Task fetch_task) +{ + using result_type = std::invoke_result_t; + + auto constexpr parallel_threshold = 32; + + auto const num_groups = items_per_group.size(); + auto const total_items = + std::accumulate(items_per_group.begin(), items_per_group.end(), std::size_t{0}); + + std::vector> results(num_groups); + + if (total_items < parallel_threshold) { + // Run sequentially to avoid task dispatch overhead + std::for_each(cuda::counting_iterator(0), + cuda::counting_iterator(num_groups), + [&](std::size_t group_idx) { + results[group_idx].reserve(items_per_group[group_idx]); + std::for_each( + cuda::counting_iterator(0), + cuda::counting_iterator(items_per_group[group_idx]), + [&](std::size_t item_idx) { + results[group_idx].emplace_back(fetch_task(group_idx, item_idx)); + }); + }); + } else { + // Dispatch every item to the host worker pool, keeping futures grouped to preserve input order + std::vector>> tasks(num_groups); + std::for_each( + cuda::counting_iterator(0), + cuda::counting_iterator(num_groups), + [&](std::size_t group_idx) { + tasks[group_idx].reserve(items_per_group[group_idx]); + std::for_each(cuda::counting_iterator(0), + cuda::counting_iterator(items_per_group[group_idx]), + [&](std::size_t item_idx) { + tasks[group_idx].emplace_back(cudf::detail::host_worker_pool().submit_task( + [&fetch_task, group_idx, item_idx]() { + return fetch_task(group_idx, item_idx); + })); + }); + }); + std::for_each(cuda::counting_iterator(0), + cuda::counting_iterator(num_groups), + [&](std::size_t group_idx) { + results[group_idx].reserve(items_per_group[group_idx]); + std::transform(tasks[group_idx].begin(), + tasks[group_idx].end(), + std::back_inserter(results[group_idx]), + [](auto& task) { return task.get(); }); + }); + } + return results; +} + /** * @copydoc cudf::io::parquet::fetch_footers_to_host */ @@ -375,6 +445,78 @@ fetch_byte_ranges_to_device_async_impl( std::async(std::launch::deferred, sync_function, std::move(device_read_tasks))}; } +std::tuple, + std::vector, + std::future> +fetch_bloom_filters_to_device_async_impl( + cudf::host_span const> datasources, + cudf::host_span const> + bloom_filter_byte_ranges_per_source, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + auto const num_sources = datasources.size(); + CUDF_EXPECTS( + num_sources == bloom_filter_byte_ranges_per_source.size(), + "Encountered mismatch in number of datasources and bloom filter byte range spans"); + + // Number of bloom filters per source (no input flattening needed). The grouped dispatch below + // decides sequential vs. host-worker-pool execution from the TOTAL count across sources, so reads + // parallelize even within a single source (the common case). + std::vector bloom_filters_per_source(num_sources); + std::transform(bloom_filter_byte_ranges_per_source.begin(), + bloom_filter_byte_ranges_per_source.end(), + bloom_filters_per_source.begin(), + [](auto const& ranges) { return ranges.size(); }); + + // Read + parse a single bloom filter header to host and return its bitset-only byte range + // `(offset + header_size, num_bytes)`. The reader emits an empty `{0, 0}` placeholder for a chunk + // whose column has no bloom filter written (to keep the per-source ranges aligned with the row + // group x column grid), so empty ranges pass through unchanged. Only the header prefix is read to + // host so a (potentially large) bitset is never staged on the host. + auto const fetch_bitset_range = + [](cudf::io::datasource& datasource, + cudf::io::text::byte_range_info const& bloom_range) -> cudf::io::text::byte_range_info { + if (bloom_range.is_empty()) { return {0, 0}; } + auto const header_read_size = + std::min(bloom_range.size(), detail::bloom_filter_header_max_size); + auto const header = datasource.host_read(static_cast(bloom_range.offset()), + static_cast(header_read_size)); + auto const header_info = detail::parse_bloom_filter_header({header->data(), header->size()}); + CUDF_EXPECTS(header_info.has_value(), "Encountered an invalid bloom filter header"); + auto const [header_size, bitset_size] = header_info.value(); + return {bloom_range.offset() + header_size, static_cast(bitset_size)}; + }; + + // Read + parse headers per `(source, bloom filter)`; results come back grouped per source. + auto const bitset_byte_ranges_per_source = dispatch_fetch_tasks( + bloom_filters_per_source, [&](std::size_t source_idx, std::size_t bloom_idx) { + return fetch_bitset_range(datasources[source_idx].get(), + bloom_filter_byte_ranges_per_source[source_idx][bloom_idx]); + }); + + // Fetch only the header-free bitsets to device. Delegate to the multi-source + // `fetch_byte_ranges_to_device_async`, which already performs the `vector -> host_span` + // conversion, so the grouped ranges are passed through directly. cuco's bloom filter probe + // requires a 32-byte-aligned bitset; the rmm buffer base is 256-byte aligned and bitset sizes are + // multiples of 32, so each bitset lands at a 32-byte-aligned address. + auto result = + fetch_byte_ranges_to_device_async(datasources, bitset_byte_ranges_per_source, stream, mr); + + auto const& bitset_spans_per_source = std::get<1>(result); + CUDF_EXPECTS(std::all_of(bitset_spans_per_source.begin(), + bitset_spans_per_source.end(), + [](auto const& spans) { + return std::all_of(spans.begin(), spans.end(), [](auto const& span) { + return span.empty() or + (reinterpret_cast(span.data()) % 32) == 0; + }); + }), + "Bloom filter bitset is not 32-byte aligned"); + + return result; +} + } // namespace [[nodiscard]] std::size_t metadata_size_hint() @@ -471,4 +613,29 @@ fetch_byte_ranges_to_device_async( mr); } +std::tuple, + std::vector>, + std::future> +fetch_bloom_filters_to_device_async( + cudf::io::datasource& datasource, + cudf::host_span bloom_filter_byte_ranges, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + CUDF_FUNC_RANGE(); + + // Wrap the inputs into arrays and delegate to the multi-source implementation + std::array, 1> datasources{std::ref(datasource)}; + std::array, 1> + bloom_filter_byte_ranges_per_source{bloom_filter_byte_ranges}; + + auto [buffers, fetched_byte_ranges, fut] = fetch_bloom_filters_to_device_async_impl( + {datasources.data(), datasources.size()}, + {bloom_filter_byte_ranges_per_source.data(), bloom_filter_byte_ranges_per_source.size()}, + stream, + mr); + + return {std::move(buffers), std::move(fetched_byte_ranges.front()), std::move(fut)}; +} + } // namespace cudf::io::parquet diff --git a/cpp/src/io/parquet/reader_impl_helpers.hpp b/cpp/src/io/parquet/reader_impl_helpers.hpp index 9417416c1e57..a4de3c0cc596 100644 --- a/cpp/src/io/parquet/reader_impl_helpers.hpp +++ b/cpp/src/io/parquet/reader_impl_helpers.hpp @@ -15,10 +15,12 @@ #include #include +#include #include #include #include #include +#include #include namespace cudf::io::parquet::detail { @@ -142,6 +144,22 @@ struct surviving_row_group_metrics { std::optional after_bloom_filter; // number of surviving row groups after bloom filter }; +/** + * @brief Upper bound on the size in bytes of a Parquet `BloomFilterHeader` + */ +inline constexpr int64_t bloom_filter_header_max_size = 256; + +/** + * @brief Parses and validates a Parquet `BloomFilterHeader` from the front of `bytes` + * + * @param bytes Host bytes starting at the beginning of a bloom filter (header followed by bitset) + * + * @return A pair of the bloom filter header size and the bitset size in bytes, or `std::nullopt` + * if the header is missing or unsupported + */ +[[nodiscard]] std::optional> parse_bloom_filter_header( + host_span bytes); + class aggregate_reader_metadata { protected: std::vector per_file_metadata; diff --git a/cpp/tests/io/experimental/hybrid_scan_composer.cpp b/cpp/tests/io/experimental/hybrid_scan_composer.cpp index 5365988220d9..98f41975dce5 100644 --- a/cpp/tests/io/experimental/hybrid_scan_composer.cpp +++ b/cpp/tests/io/experimental/hybrid_scan_composer.cpp @@ -104,13 +104,10 @@ auto apply_hybrid_scan_filters(cudf::io::datasource& datasource, std::vector bloom_filtered_row_group_indices; bloom_filtered_row_group_indices.reserve(current_row_group_indices.size()); if (bloom_filter_byte_ranges.size()) { - // Fetch 32 byte aligned bloom filter data buffers from the input file buffer - auto aligned_mr = rmm::mr::aligned_resource_adaptor(cudf::get_current_device_resource_ref(), - bloom_filter_alignment); - + // Fetch the header-stripped, 32-byte-aligned bloom filter bitsets from the input file auto [bloom_filter_buffers, bloom_filter_data, bloom_read_tasks] = - cudf::io::parquet::fetch_byte_ranges_to_device_async( - datasource, bloom_filter_byte_ranges, stream, aligned_mr); + cudf::io::parquet::fetch_bloom_filters_to_device_async( + datasource, bloom_filter_byte_ranges, stream, mr); bloom_read_tasks.get(); // Filter row groups with bloom filters diff --git a/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pyi b/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pyi index f95dc8b054d3..739e46994203 100644 --- a/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pyi +++ b/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pyi @@ -6,10 +6,11 @@ from enum import IntEnum from rmm.pylibrmm.memory_resource import DeviceMemoryResource from pylibcudf.column import Column +from pylibcudf.gpumemoryview import gpumemoryview from pylibcudf.io.parquet import ParquetReaderOptions from pylibcudf.io.parquet_metadata import FileMetaData from pylibcudf.io.text import ByteRangeInfo -from pylibcudf.io.types import TableWithMetadata +from pylibcudf.io.types import SourceInfo, TableWithMetadata from pylibcudf.utils import CudaStreamLike try: @@ -142,3 +143,10 @@ class HybridScanReader: pass_read_limit: int, ) -> list[list[int]]: ... def has_next_table_chunk(self) -> bool: ... + +def fetch_bloom_filters_to_device_async( + source_info: SourceInfo, + bloom_filter_byte_ranges: list[ByteRangeInfo], + stream: CudaStreamLike | None = None, + mr: DeviceMemoryResource | None = None, +) -> list[gpumemoryview]: ... diff --git a/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pyx b/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pyx index 30aa33c0bf2c..048482e58cd8 100644 --- a/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pyx +++ b/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pyx @@ -18,10 +18,20 @@ from pylibcudf.io.text cimport ByteRangeInfo from pylibcudf.io.types cimport TableWithMetadata from pylibcudf.libcudf.column.column cimport column from pylibcudf.libcudf.column.column_view cimport column_view, mutable_column_view +from pylibcudf.gpumemoryview cimport gpumemoryview +from pylibcudf.io.types cimport SourceInfo +from pylibcudf.libcudf.io.datasource cimport datasource, make_datasources from pylibcudf.libcudf.io.hybrid_scan cimport ( + bloom_fetch_get_buffers, + bloom_fetch_get_future, + bloom_fetch_get_spans, + bloom_filter_fetch_result, + const_byte_range_info, const_device_span_const_uint8_t, const_size_type, const_uint8_t, + device_span_u8, + fetch_bloom_filters_to_device_async as cpp_fetch_bloom_filters_to_device_async, hybrid_scan_reader as cpp_hybrid_scan_reader, use_data_page_mask as cpp_use_data_page_mask, ) @@ -30,6 +40,7 @@ from pylibcudf.libcudf.io.types cimport table_with_metadata from pylibcudf.libcudf.types cimport size_type from pylibcudf.libcudf.utilities.span cimport device_span, host_span from pylibcudf.utils cimport _get_memory_resource, _get_stream +from rmm.librmm.device_buffer cimport device_buffer from pylibcudf.span import is_span from pylibcudf.io.parquet_metadata import FileMetaData @@ -38,7 +49,12 @@ import pylibcudf.libcudf.io.hybrid_scan UseDataPageMask = pylibcudf.libcudf.io.hybrid_scan.use_data_page_mask -__all__ = ["FileMetaData", "HybridScanReader", "UseDataPageMask"] +__all__ = [ + "FileMetaData", + "HybridScanReader", + "UseDataPageMask", + "fetch_bloom_filters_to_device_async", +] cdef device_span[const_uint8_t] _get_device_span(object obj) except *: @@ -52,6 +68,125 @@ cdef device_span[const_uint8_t] _get_device_span(object obj) except *: obj.size) +cdef class _BloomFilterBitsetBuffers: + """Owns the packed device buffer(s) backing fetched bloom filter bitsets. + + The bitset views handed to Python are non-owning, so this holder keeps the + underlying device memory alive for as long as any view references it. + """ + cdef vector[device_buffer] buffers + + +cdef class _BloomFilterBitsetView: + """Non-owning device view of a single bloom filter bitset. + + Exposes the CUDA array interface (and hence the Span protocol via + :class:`gpumemoryview`) over one bitset while keeping its backing buffers + alive through ``_owner``. + """ + cdef readonly uintptr_t ptr + cdef readonly size_t nbytes + cdef object _owner + + @property + def size(self): + return self.nbytes + + @property + def __cuda_array_interface__(self): + return { + "shape": (int(self.nbytes),), + "typestr": "|u1", + "data": (int(self.ptr), False), + "version": 3, + } + + +def fetch_bloom_filters_to_device_async( + SourceInfo source_info, + list bloom_filter_byte_ranges, + object stream=None, + object mr=None, +): + """Fetch header-stripped, 32-byte-aligned bloom filter bitsets to device. + + Each input byte range must span a complete bloom filter (its Thrift + ``BloomFilterHeader`` followed by the bitset), as returned by + :meth:`HybridScanReader.secondary_filters_byte_ranges`. The headers are read + to host and parsed to locate each bitset; only the header-free bitsets are + fetched to device, at 32-byte-aligned addresses as required by the bloom + filter probe. + + For details, see + :cpp:func:`cudf::io::parquet::fetch_bloom_filters_to_device_async` + + Parameters + ---------- + source_info : SourceInfo + Source describing the Parquet file to read bloom filters from. Must + describe a single source. + bloom_filter_byte_ranges : list[ByteRangeInfo] + Byte ranges of complete bloom filters (header + bitset) to fetch. + stream : Stream, optional + CUDA stream. + mr : DeviceMemoryResource, optional + Device memory resource used to allocate the returned buffers. + + Returns + ------- + list[gpumemoryview] + One device view per input range, each holding a header-free bitset. + """ + cdef Stream _stream = _get_stream(stream) + cdef DeviceMemoryResource _mr = _get_memory_resource(mr) + + cdef vector[unique_ptr[datasource]] datasources = \ + make_datasources(source_info.c_obj) + if datasources.size() != 1: + raise ValueError( + "fetch_bloom_filters_to_device_async expects a single source, got " + f"{datasources.size()}" + ) + + cdef vector[byte_range_info] c_ranges + cdef ByteRangeInfo br + for obj in bloom_filter_byte_ranges: + br = obj + c_ranges.push_back(br.c_obj) + + cdef bloom_filter_fetch_result result = \ + cpp_fetch_bloom_filters_to_device_async( + datasources[0].get()[0], + host_span[const_byte_range_info]( + c_ranges.data(), + c_ranges.size(), + ), + _stream.view(), + _mr.get_mr(), + ) + # Block until the device reads launched by the helper complete. + bloom_fetch_get_future(result).get() + + # Take ownership of the packed device buffer(s); the views below are + # non-owning and must keep this alive. + cdef _BloomFilterBitsetBuffers owner = \ + _BloomFilterBitsetBuffers.__new__(_BloomFilterBitsetBuffers) + owner.buffers = move(bloom_fetch_get_buffers(result)) + + # The bitset spans are lightweight views into the buffers above. + cdef vector[device_span_u8] spans = bloom_fetch_get_spans(result) + cdef size_t i + cdef _BloomFilterBitsetView view + bloom_data = [] + for i in range(spans.size()): + view = _BloomFilterBitsetView.__new__(_BloomFilterBitsetView) + view.ptr = spans[i].data() + view.nbytes = spans[i].size() + view._owner = owner + bloom_data.append(gpumemoryview(view)) + return bloom_data + + cdef class HybridScanReader: """Experimental Parquet reader optimized for highly selective filters. diff --git a/python/pylibcudf/pylibcudf/libcudf/io/hybrid_scan.pxd b/python/pylibcudf/pylibcudf/libcudf/io/hybrid_scan.pxd index 8578908fc43c..82998d6da0e9 100644 --- a/python/pylibcudf/pylibcudf/libcudf/io/hybrid_scan.pxd +++ b/python/pylibcudf/pylibcudf/libcudf/io/hybrid_scan.pxd @@ -1,6 +1,7 @@ # SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION. # SPDX-License-Identifier: Apache-2.0 +from libc.stddef cimport size_t from libc.stdint cimport uint8_t from libcpp cimport bool from libcpp.memory cimport unique_ptr @@ -9,6 +10,7 @@ from libcpp.vector cimport vector from pylibcudf.exception_handler cimport libcudf_exception_handler from pylibcudf.libcudf.column.column cimport column from pylibcudf.libcudf.column.column_view cimport column_view, mutable_column_view +from pylibcudf.libcudf.io.datasource cimport datasource from pylibcudf.libcudf.io.parquet cimport parquet_reader_options from pylibcudf.libcudf.io.parquet_schema cimport FileMetaData from pylibcudf.libcudf.io.text cimport byte_range_info @@ -16,11 +18,14 @@ from pylibcudf.libcudf.io.types cimport table_with_metadata from pylibcudf.libcudf.types cimport size_type from pylibcudf.libcudf.utilities.span cimport device_span, host_span from cuda.bindings.cyruntime cimport cudaStream_t +from rmm.librmm.cuda_stream_view cimport cuda_stream_view +from rmm.librmm.device_buffer cimport device_buffer 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 device_span[const_uint8_t] const_device_span_const_uint8_t +ctypedef const byte_range_info const_byte_range_info cdef extern from "cudf/io/experimental/hybrid_scan.hpp" \ namespace "cudf::io::parquet::experimental" nogil: @@ -173,3 +178,41 @@ cdef extern from "cudf/io/experimental/hybrid_scan.hpp" \ ) except +libcudf_exception_handler bool has_next_table_chunk() except +libcudf_exception_handler + + +# Bloom filter fetch IO util (cudf/io/parquet_io_utils.hpp). Co-located here as +# the hybrid scan flow is its only consumer. Binding the tuple/future return +# requires a little std:: glue since there is no Cython precedent for either. +cdef extern from "cudf/utilities/span.hpp" namespace "cudf" nogil: + cdef cppclass device_span_u8 "cudf::device_span": + const_uint8_t* data() + size_t size() + + +cdef extern from "" namespace "std" nogil: + cdef cppclass future_void "std::future": + void get() except +libcudf_exception_handler + + +cdef extern from "cudf/io/parquet_io_utils.hpp" namespace "cudf::io::parquet" nogil: + cdef cppclass bloom_filter_fetch_result "std::tuple, std::vector >, std::future >": # noqa: E501 + bloom_filter_fetch_result() + + bloom_filter_fetch_result fetch_bloom_filters_to_device_async( + datasource& source, + host_span[const_byte_range_info] bloom_filter_byte_ranges, + cuda_stream_view stream, + device_async_resource_ref mr, + ) except +libcudf_exception_handler + + +cdef extern from "" namespace "std" nogil: + vector[device_buffer]& bloom_fetch_get_buffers "std::get<0>"( + bloom_filter_fetch_result& result + ) + vector[device_span_u8]& bloom_fetch_get_spans "std::get<1>"( + bloom_filter_fetch_result& result + ) + future_void& bloom_fetch_get_future "std::get<2>"( + bloom_filter_fetch_result& result + ) diff --git a/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py b/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py index b58a76b198fb..58eb44a66023 100644 --- a/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py +++ b/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py @@ -252,10 +252,6 @@ def test_hybrid_scan_secondary_filters_byte_ranges( assert isinstance(dict_ranges, list) -@pytest.mark.xfail( - reason="hybrid scan does not strip the Parquet BloomFilterHeader before probing", - strict=True, -) @pytest.mark.parametrize( "fixture_name,column_name,query_value", [ @@ -277,7 +273,8 @@ def test_hybrid_scan_bloom_filter_matches_read_parquet( """A/B test hybrid scan bloom filtering vs read_parquet. A: read_parquet path (libcudf strips the BloomFilterHeader before probing) - B: hybrid-scan path (fetches BloomFilterHeader + bitset from secondary_filters_byte_ranges) + B: hybrid-scan path (fetch_bloom_filters_to_device_async strips the header, + returning header-free, 32-byte-aligned bitsets) The hybrid path must keep the same row groups as read_parquet. """ @@ -336,8 +333,8 @@ def make_path_options(): # B: hybrid scan should keep the same row group. print( - "[bloom-demo] B: hybrid-scan path (fetches " - "BloomFilterHeader + bitset from secondary_filters_byte_ranges)", + "[bloom-demo] B: hybrid-scan path " + "(fetch_bloom_filters_to_device_async strips the header)", flush=True, ) options = make_options() @@ -358,13 +355,16 @@ def make_path_options(): ) stream = plc.utils._get_stream(None) - bloom_data = [ - plc.gpumemoryview( - rmm.DeviceBuffer.to_device(data[r.offset : r.offset + r.size], stream) - ) - for r in bloom_ranges - ] + # The fix: fetch header-stripped, 32-byte-aligned bitsets via the IO helper + # (this is what the C++ hybrid-scan callers now use), instead of copying the + # whole BloomFilterHeader+bitset range to device. + bloom_data = plc.io.experimental.hybrid_scan.fetch_bloom_filters_to_device_async( + plc.io.SourceInfo([io.BytesIO(data)]), bloom_ranges, stream + ) synchronize_stream(None) + # The bloom filter probe requires 32-byte-aligned bitsets. + for bitset in bloom_data: + assert bitset.ptr % 32 == 0 surviving = reader.filter_row_groups_with_bloom_filters( bloom_data, row_groups, options ) From be4fde68b835ab0fbb03f7260c38af88d3e6f9dc Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Fri, 26 Jun 2026 22:24:32 +0200 Subject: [PATCH 06/39] Refactor hybrid scan implementation to streamline bloom filter handling Removed unused imports and functions related to bloom filter processing in the hybrid scan module. This cleanup enhances code clarity and reduces complexity. Updated the test suite to reflect these changes, ensuring that the functionality remains intact and consistent with previous implementations. --- .../experimental/hybrid_scan_filters_test.cpp | 4 - .../pylibcudf/io/experimental/hybrid_scan.pyi | 10 +- .../pylibcudf/io/experimental/hybrid_scan.pyx | 137 +----------------- .../pylibcudf/libcudf/io/hybrid_scan.pxd | 43 ------ .../tests/io/test_experimental_hybrid_scan.py | 123 ---------------- 5 files changed, 2 insertions(+), 315 deletions(-) diff --git a/cpp/tests/io/experimental/hybrid_scan_filters_test.cpp b/cpp/tests/io/experimental/hybrid_scan_filters_test.cpp index be7f60736430..10960c185ef9 100644 --- a/cpp/tests/io/experimental/hybrid_scan_filters_test.cpp +++ b/cpp/tests/io/experimental/hybrid_scan_filters_test.cpp @@ -19,15 +19,11 @@ #include -#include -#include - #include #include #include #include #include -#include #include namespace { diff --git a/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pyi b/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pyi index 739e46994203..f95dc8b054d3 100644 --- a/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pyi +++ b/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pyi @@ -6,11 +6,10 @@ from enum import IntEnum from rmm.pylibrmm.memory_resource import DeviceMemoryResource from pylibcudf.column import Column -from pylibcudf.gpumemoryview import gpumemoryview from pylibcudf.io.parquet import ParquetReaderOptions from pylibcudf.io.parquet_metadata import FileMetaData from pylibcudf.io.text import ByteRangeInfo -from pylibcudf.io.types import SourceInfo, TableWithMetadata +from pylibcudf.io.types import TableWithMetadata from pylibcudf.utils import CudaStreamLike try: @@ -143,10 +142,3 @@ class HybridScanReader: pass_read_limit: int, ) -> list[list[int]]: ... def has_next_table_chunk(self) -> bool: ... - -def fetch_bloom_filters_to_device_async( - source_info: SourceInfo, - bloom_filter_byte_ranges: list[ByteRangeInfo], - stream: CudaStreamLike | None = None, - mr: DeviceMemoryResource | None = None, -) -> list[gpumemoryview]: ... diff --git a/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pyx b/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pyx index 048482e58cd8..30aa33c0bf2c 100644 --- a/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pyx +++ b/python/pylibcudf/pylibcudf/io/experimental/hybrid_scan.pyx @@ -18,20 +18,10 @@ from pylibcudf.io.text cimport ByteRangeInfo from pylibcudf.io.types cimport TableWithMetadata from pylibcudf.libcudf.column.column cimport column from pylibcudf.libcudf.column.column_view cimport column_view, mutable_column_view -from pylibcudf.gpumemoryview cimport gpumemoryview -from pylibcudf.io.types cimport SourceInfo -from pylibcudf.libcudf.io.datasource cimport datasource, make_datasources from pylibcudf.libcudf.io.hybrid_scan cimport ( - bloom_fetch_get_buffers, - bloom_fetch_get_future, - bloom_fetch_get_spans, - bloom_filter_fetch_result, - const_byte_range_info, const_device_span_const_uint8_t, const_size_type, const_uint8_t, - device_span_u8, - fetch_bloom_filters_to_device_async as cpp_fetch_bloom_filters_to_device_async, hybrid_scan_reader as cpp_hybrid_scan_reader, use_data_page_mask as cpp_use_data_page_mask, ) @@ -40,7 +30,6 @@ from pylibcudf.libcudf.io.types cimport table_with_metadata from pylibcudf.libcudf.types cimport size_type from pylibcudf.libcudf.utilities.span cimport device_span, host_span from pylibcudf.utils cimport _get_memory_resource, _get_stream -from rmm.librmm.device_buffer cimport device_buffer from pylibcudf.span import is_span from pylibcudf.io.parquet_metadata import FileMetaData @@ -49,12 +38,7 @@ import pylibcudf.libcudf.io.hybrid_scan UseDataPageMask = pylibcudf.libcudf.io.hybrid_scan.use_data_page_mask -__all__ = [ - "FileMetaData", - "HybridScanReader", - "UseDataPageMask", - "fetch_bloom_filters_to_device_async", -] +__all__ = ["FileMetaData", "HybridScanReader", "UseDataPageMask"] cdef device_span[const_uint8_t] _get_device_span(object obj) except *: @@ -68,125 +52,6 @@ cdef device_span[const_uint8_t] _get_device_span(object obj) except *: obj.size) -cdef class _BloomFilterBitsetBuffers: - """Owns the packed device buffer(s) backing fetched bloom filter bitsets. - - The bitset views handed to Python are non-owning, so this holder keeps the - underlying device memory alive for as long as any view references it. - """ - cdef vector[device_buffer] buffers - - -cdef class _BloomFilterBitsetView: - """Non-owning device view of a single bloom filter bitset. - - Exposes the CUDA array interface (and hence the Span protocol via - :class:`gpumemoryview`) over one bitset while keeping its backing buffers - alive through ``_owner``. - """ - cdef readonly uintptr_t ptr - cdef readonly size_t nbytes - cdef object _owner - - @property - def size(self): - return self.nbytes - - @property - def __cuda_array_interface__(self): - return { - "shape": (int(self.nbytes),), - "typestr": "|u1", - "data": (int(self.ptr), False), - "version": 3, - } - - -def fetch_bloom_filters_to_device_async( - SourceInfo source_info, - list bloom_filter_byte_ranges, - object stream=None, - object mr=None, -): - """Fetch header-stripped, 32-byte-aligned bloom filter bitsets to device. - - Each input byte range must span a complete bloom filter (its Thrift - ``BloomFilterHeader`` followed by the bitset), as returned by - :meth:`HybridScanReader.secondary_filters_byte_ranges`. The headers are read - to host and parsed to locate each bitset; only the header-free bitsets are - fetched to device, at 32-byte-aligned addresses as required by the bloom - filter probe. - - For details, see - :cpp:func:`cudf::io::parquet::fetch_bloom_filters_to_device_async` - - Parameters - ---------- - source_info : SourceInfo - Source describing the Parquet file to read bloom filters from. Must - describe a single source. - bloom_filter_byte_ranges : list[ByteRangeInfo] - Byte ranges of complete bloom filters (header + bitset) to fetch. - stream : Stream, optional - CUDA stream. - mr : DeviceMemoryResource, optional - Device memory resource used to allocate the returned buffers. - - Returns - ------- - list[gpumemoryview] - One device view per input range, each holding a header-free bitset. - """ - cdef Stream _stream = _get_stream(stream) - cdef DeviceMemoryResource _mr = _get_memory_resource(mr) - - cdef vector[unique_ptr[datasource]] datasources = \ - make_datasources(source_info.c_obj) - if datasources.size() != 1: - raise ValueError( - "fetch_bloom_filters_to_device_async expects a single source, got " - f"{datasources.size()}" - ) - - cdef vector[byte_range_info] c_ranges - cdef ByteRangeInfo br - for obj in bloom_filter_byte_ranges: - br = obj - c_ranges.push_back(br.c_obj) - - cdef bloom_filter_fetch_result result = \ - cpp_fetch_bloom_filters_to_device_async( - datasources[0].get()[0], - host_span[const_byte_range_info]( - c_ranges.data(), - c_ranges.size(), - ), - _stream.view(), - _mr.get_mr(), - ) - # Block until the device reads launched by the helper complete. - bloom_fetch_get_future(result).get() - - # Take ownership of the packed device buffer(s); the views below are - # non-owning and must keep this alive. - cdef _BloomFilterBitsetBuffers owner = \ - _BloomFilterBitsetBuffers.__new__(_BloomFilterBitsetBuffers) - owner.buffers = move(bloom_fetch_get_buffers(result)) - - # The bitset spans are lightweight views into the buffers above. - cdef vector[device_span_u8] spans = bloom_fetch_get_spans(result) - cdef size_t i - cdef _BloomFilterBitsetView view - bloom_data = [] - for i in range(spans.size()): - view = _BloomFilterBitsetView.__new__(_BloomFilterBitsetView) - view.ptr = spans[i].data() - view.nbytes = spans[i].size() - view._owner = owner - bloom_data.append(gpumemoryview(view)) - return bloom_data - - cdef class HybridScanReader: """Experimental Parquet reader optimized for highly selective filters. diff --git a/python/pylibcudf/pylibcudf/libcudf/io/hybrid_scan.pxd b/python/pylibcudf/pylibcudf/libcudf/io/hybrid_scan.pxd index 82998d6da0e9..8578908fc43c 100644 --- a/python/pylibcudf/pylibcudf/libcudf/io/hybrid_scan.pxd +++ b/python/pylibcudf/pylibcudf/libcudf/io/hybrid_scan.pxd @@ -1,7 +1,6 @@ # SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION. # SPDX-License-Identifier: Apache-2.0 -from libc.stddef cimport size_t from libc.stdint cimport uint8_t from libcpp cimport bool from libcpp.memory cimport unique_ptr @@ -10,7 +9,6 @@ from libcpp.vector cimport vector from pylibcudf.exception_handler cimport libcudf_exception_handler from pylibcudf.libcudf.column.column cimport column from pylibcudf.libcudf.column.column_view cimport column_view, mutable_column_view -from pylibcudf.libcudf.io.datasource cimport datasource from pylibcudf.libcudf.io.parquet cimport parquet_reader_options from pylibcudf.libcudf.io.parquet_schema cimport FileMetaData from pylibcudf.libcudf.io.text cimport byte_range_info @@ -18,14 +16,11 @@ from pylibcudf.libcudf.io.types cimport table_with_metadata from pylibcudf.libcudf.types cimport size_type from pylibcudf.libcudf.utilities.span cimport device_span, host_span from cuda.bindings.cyruntime cimport cudaStream_t -from rmm.librmm.cuda_stream_view cimport cuda_stream_view -from rmm.librmm.device_buffer cimport device_buffer 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 device_span[const_uint8_t] const_device_span_const_uint8_t -ctypedef const byte_range_info const_byte_range_info cdef extern from "cudf/io/experimental/hybrid_scan.hpp" \ namespace "cudf::io::parquet::experimental" nogil: @@ -178,41 +173,3 @@ cdef extern from "cudf/io/experimental/hybrid_scan.hpp" \ ) except +libcudf_exception_handler bool has_next_table_chunk() except +libcudf_exception_handler - - -# Bloom filter fetch IO util (cudf/io/parquet_io_utils.hpp). Co-located here as -# the hybrid scan flow is its only consumer. Binding the tuple/future return -# requires a little std:: glue since there is no Cython precedent for either. -cdef extern from "cudf/utilities/span.hpp" namespace "cudf" nogil: - cdef cppclass device_span_u8 "cudf::device_span": - const_uint8_t* data() - size_t size() - - -cdef extern from "" namespace "std" nogil: - cdef cppclass future_void "std::future": - void get() except +libcudf_exception_handler - - -cdef extern from "cudf/io/parquet_io_utils.hpp" namespace "cudf::io::parquet" nogil: - cdef cppclass bloom_filter_fetch_result "std::tuple, std::vector >, std::future >": # noqa: E501 - bloom_filter_fetch_result() - - bloom_filter_fetch_result fetch_bloom_filters_to_device_async( - datasource& source, - host_span[const_byte_range_info] bloom_filter_byte_ranges, - cuda_stream_view stream, - device_async_resource_ref mr, - ) except +libcudf_exception_handler - - -cdef extern from "" namespace "std" nogil: - vector[device_buffer]& bloom_fetch_get_buffers "std::get<0>"( - bloom_filter_fetch_result& result - ) - vector[device_span_u8]& bloom_fetch_get_spans "std::get<1>"( - bloom_filter_fetch_result& result - ) - future_void& bloom_fetch_get_future "std::get<2>"( - bloom_filter_fetch_result& result - ) diff --git a/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py b/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py index 58eb44a66023..74f467f16193 100644 --- a/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py +++ b/python/pylibcudf/tests/io/test_experimental_hybrid_scan.py @@ -1,7 +1,6 @@ # SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: Apache-2.0 import io -from pathlib import Path import pyarrow as pa import pyarrow.parquet as pq @@ -252,128 +251,6 @@ def test_hybrid_scan_secondary_filters_byte_ranges( assert isinstance(dict_ranges, list) -@pytest.mark.parametrize( - "fixture_name,column_name,query_value", - [ - ( - "bloom_filter_alignment.parquet", - "r_reason_desc", - "Did not like the color", - ), - ( - "mixed_card_ndv_100_bf_fpp0.1_nostats.snappy.parquet", - "str", - "tbRpIyVJZX", - ), - ], -) -def test_hybrid_scan_bloom_filter_matches_read_parquet( - fixture_name, column_name, query_value -): - """A/B test hybrid scan bloom filtering vs read_parquet. - - A: read_parquet path (libcudf strips the BloomFilterHeader before probing) - B: hybrid-scan path (fetch_bloom_filters_to_device_async strips the header, - returning header-free, 32-byte-aligned bitsets) - - The hybrid path must keep the same row groups as read_parquet. - """ - fixture = ( - Path(__file__).parents[4] - / "python/cudf/cudf/tests/data/parquet" - / fixture_name - ) - if not fixture.exists(): - pytest.skip(f"bloom fixture not found: {fixture}") - data = fixture.read_bytes() - expected_row_groups = list(range(pq.ParquetFile(fixture).metadata.num_row_groups)) - - literal_value = plc.Scalar.from_arrow(pa.scalar(query_value)) - literal = Literal(literal_value) - column_ref = ColumnNameReference(column_name) - bloom_filter = Operation(ASTOperator.EQUAL, column_ref, literal) - - def make_options(): - options = plc.io.parquet.ParquetReaderOptions.builder( - plc.io.SourceInfo([io.BytesIO(data)]) - ).build() - options.set_filter(bloom_filter) - return options - - def make_path_options(): - options = plc.io.parquet.ParquetReaderOptions.builder( - plc.io.SourceInfo([fixture]) - ).build() - options.set_filter(bloom_filter) - return options - - # A: standard read_parquet keeps the only row group after bloom filtering. - print( - f"\n[bloom-demo] fixture={fixture.name} predicate={column_name} == {query_value!r}\n", - flush=True, - ) - print( - "[bloom-demo] A: read_parquet path (libcudf strips the " - "BloomFilterHeader before probing)", - flush=True, - ) - table_w_meta = plc.io.parquet.read_parquet(make_path_options()) - print( - "[bloom-demo] A result: " - f"rows={table_w_meta.tbl.num_rows()} " - f"num_input_row_groups={table_w_meta.num_input_row_groups} " - f"num_row_groups_after_bloom_filter=" - f"{table_w_meta.num_row_groups_after_bloom_filter}\n", - flush=True, - ) - assert table_w_meta.num_input_row_groups == len(expected_row_groups) - assert table_w_meta.num_row_groups_after_bloom_filter == len( - expected_row_groups - ) - - # B: hybrid scan should keep the same row group. - print( - "[bloom-demo] B: hybrid-scan path " - "(fetch_bloom_filters_to_device_async strips the header)", - flush=True, - ) - options = make_options() - suffix = 8 # 4-byte footer length + "PAR1" - mv = memoryview(data) - footer_size = int.from_bytes(mv[-suffix:-4], byteorder="little") - reader = HybridScanReader(mv[-suffix - footer_size : -suffix], options) - - row_groups = reader.all_row_groups(options) - bloom_ranges, _ = reader.secondary_filters_byte_ranges(row_groups, options) - assert bloom_ranges # the equality predicate makes r_reason_desc bloom-eligible - for idx, r in enumerate(bloom_ranges): - raw = data[r.offset : r.offset + r.size] - print( - f"[bloom-demo] B fetched range[{idx}]: " - f"offset={r.offset} size={r.size} raw={raw.hex()}", - flush=True, - ) - - stream = plc.utils._get_stream(None) - # The fix: fetch header-stripped, 32-byte-aligned bitsets via the IO helper - # (this is what the C++ hybrid-scan callers now use), instead of copying the - # whole BloomFilterHeader+bitset range to device. - bloom_data = plc.io.experimental.hybrid_scan.fetch_bloom_filters_to_device_async( - plc.io.SourceInfo([io.BytesIO(data)]), bloom_ranges, stream - ) - synchronize_stream(None) - # The bloom filter probe requires 32-byte-aligned bitsets. - for bitset in bloom_data: - assert bitset.ptr % 32 == 0 - surviving = reader.filter_row_groups_with_bloom_filters( - bloom_data, row_groups, options - ) - print(f"[bloom-demo] B result: surviving_row_groups={surviving}", flush=True) - - # The queried value is present, so the hybrid path must match read_parquet. - assert surviving == row_groups == expected_row_groups - - def test_hybrid_scan_column_chunk_byte_ranges( simple_hybrid_scan_reader: HybridScanReader, simple_parquet_options: plc.io.parquet.ParquetReaderOptions, From 48517bc371ba15a779338bf179879b101b5c1d0a Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Fri, 26 Jun 2026 22:24:42 +0200 Subject: [PATCH 07/39] Update copyright notices and improve code clarity in hybrid scan files Modified copyright statements in multiple files to include "NVIDIA CORPORATION & AFFILIATES." Enhanced code readability by breaking long parameter descriptions into multiple lines in the `parquet_io_utils.hpp` file. Removed outdated debug logging and unnecessary comments in the `bloom_filter_reader.cu` and `parquet_io_utils.cpp` files to streamline the codebase. --- .../hybrid_scan/hybrid_scan_composer.cpp | 2 +- .../hybrid_scan_io/hybrid_scan_composer.cpp | 2 +- cpp/include/cudf/io/parquet_io_utils.hpp | 3 +- cpp/src/io/parquet/bloom_filter_reader.cu | 22 +--- .../io/parquet/io_utils/parquet_io_utils.cpp | 106 ++++-------------- .../io/experimental/hybrid_scan_composer.cpp | 2 +- 6 files changed, 29 insertions(+), 108 deletions(-) diff --git a/cpp/benchmarks/io/parquet/experimental/hybrid_scan/hybrid_scan_composer.cpp b/cpp/benchmarks/io/parquet/experimental/hybrid_scan/hybrid_scan_composer.cpp index 98db682a709c..cd5dfd54fe1e 100644 --- a/cpp/benchmarks/io/parquet/experimental/hybrid_scan/hybrid_scan_composer.cpp +++ b/cpp/benchmarks/io/parquet/experimental/hybrid_scan/hybrid_scan_composer.cpp @@ -1,6 +1,6 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ diff --git a/cpp/examples/hybrid_scan_io/hybrid_scan_composer.cpp b/cpp/examples/hybrid_scan_io/hybrid_scan_composer.cpp index fdc6f3735c86..8eac7ea80f25 100644 --- a/cpp/examples/hybrid_scan_io/hybrid_scan_composer.cpp +++ b/cpp/examples/hybrid_scan_io/hybrid_scan_composer.cpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ diff --git a/cpp/include/cudf/io/parquet_io_utils.hpp b/cpp/include/cudf/io/parquet_io_utils.hpp index 9c0c759d95aa..58ff6529264b 100644 --- a/cpp/include/cudf/io/parquet_io_utils.hpp +++ b/cpp/include/cudf/io/parquet_io_utils.hpp @@ -157,7 +157,8 @@ fetch_byte_ranges_to_device_async( * @ingroup io_utils * * @param datasource Input datasource - * @param bloom_filter_byte_ranges Byte ranges of complete bloom filters to fetch, must span a complete bloom filter + * @param bloom_filter_byte_ranges Byte ranges of complete bloom filters to fetch, must span a + * complete bloom filter * @param stream CUDA stream * @param mr Device memory resource * diff --git a/cpp/src/io/parquet/bloom_filter_reader.cu b/cpp/src/io/parquet/bloom_filter_reader.cu index edc00f79a003..7e91f95a35b9 100644 --- a/cpp/src/io/parquet/bloom_filter_reader.cu +++ b/cpp/src/io/parquet/bloom_filter_reader.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 */ @@ -29,7 +29,6 @@ #include #include -#include #include #include #include @@ -529,25 +528,6 @@ std::optional>> aggregate_reader_metadata::ap static_cast(total_row_groups), bloom_filter_col_schemas.size()}; - // [bloom-dbg] dev-only: remove before merge. Hexdump the exact bytes each path feeds bloom - // filter construction. read_parquet strips the BloomFilterHeader; hybrid scan currently does not, - // so its spans are 16 bytes longer (header + bitset) and the caster misreads the header as a - // filter block. - for (std::size_t dbg_i = 0; dbg_i < bloom_filter_data.size(); ++dbg_i) { - auto const& dbg_span = bloom_filter_data[dbg_i]; - std::vector dbg_host(dbg_span.size()); - if (not dbg_span.empty()) { - cudaMemcpyAsync( - dbg_host.data(), dbg_span.data(), dbg_span.size(), cudaMemcpyDeviceToHost, stream.value()); - stream.synchronize(); - } - std::fprintf(stderr, "[bloom-dbg] bloom span[%zu] size=%zu raw=", dbg_i, dbg_span.size()); - for (auto const byte : dbg_host) { - std::fprintf(stderr, "%02x", static_cast(byte)); - } - std::fprintf(stderr, "\n"); - } - // Converts bloom filter membership for equality predicate columns to a table // containing a column for each `col[i] == literal` predicate to be evaluated. // The table contains #sources * #column_chunks_per_src rows. diff --git a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp index ea2efff68ff6..ba9da844d27f 100644 --- a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp +++ b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp @@ -92,71 +92,6 @@ auto dispatch_fetch_tasks(std::size_t num_sources, Task fetch_task) return results; } -/** - * @brief Dispatches a grouped fetch task for each `(group_idx, item_idx)` and collects the results - * based on the total number of items across all groups. - * - * @tparam Task Callable invocable as `fetch_task(std::size_t group_idx, std::size_t item_idx)` - * @param items_per_group Number of items in each group - * @param fetch_task Task to run for each `(group_idx, item_idx)` - * @return Vector (one per group) of vectors of results, in input order - */ -template -auto dispatch_fetch_tasks(cudf::host_span items_per_group, Task fetch_task) -{ - using result_type = std::invoke_result_t; - - auto constexpr parallel_threshold = 32; - - auto const num_groups = items_per_group.size(); - auto const total_items = - std::accumulate(items_per_group.begin(), items_per_group.end(), std::size_t{0}); - - std::vector> results(num_groups); - - if (total_items < parallel_threshold) { - // Run sequentially to avoid task dispatch overhead - std::for_each(cuda::counting_iterator(0), - cuda::counting_iterator(num_groups), - [&](std::size_t group_idx) { - results[group_idx].reserve(items_per_group[group_idx]); - std::for_each( - cuda::counting_iterator(0), - cuda::counting_iterator(items_per_group[group_idx]), - [&](std::size_t item_idx) { - results[group_idx].emplace_back(fetch_task(group_idx, item_idx)); - }); - }); - } else { - // Dispatch every item to the host worker pool, keeping futures grouped to preserve input order - std::vector>> tasks(num_groups); - std::for_each( - cuda::counting_iterator(0), - cuda::counting_iterator(num_groups), - [&](std::size_t group_idx) { - tasks[group_idx].reserve(items_per_group[group_idx]); - std::for_each(cuda::counting_iterator(0), - cuda::counting_iterator(items_per_group[group_idx]), - [&](std::size_t item_idx) { - tasks[group_idx].emplace_back(cudf::detail::host_worker_pool().submit_task( - [&fetch_task, group_idx, item_idx]() { - return fetch_task(group_idx, item_idx); - })); - }); - }); - std::for_each(cuda::counting_iterator(0), - cuda::counting_iterator(num_groups), - [&](std::size_t group_idx) { - results[group_idx].reserve(items_per_group[group_idx]); - std::transform(tasks[group_idx].begin(), - tasks[group_idx].end(), - std::back_inserter(results[group_idx]), - [](auto& task) { return task.get(); }); - }); - } - return results; -} - /** * @copydoc cudf::io::parquet::fetch_footers_to_host */ @@ -456,18 +391,8 @@ fetch_bloom_filters_to_device_async_impl( rmm::device_async_resource_ref mr) { auto const num_sources = datasources.size(); - CUDF_EXPECTS( - num_sources == bloom_filter_byte_ranges_per_source.size(), - "Encountered mismatch in number of datasources and bloom filter byte range spans"); - - // Number of bloom filters per source (no input flattening needed). The grouped dispatch below - // decides sequential vs. host-worker-pool execution from the TOTAL count across sources, so reads - // parallelize even within a single source (the common case). - std::vector bloom_filters_per_source(num_sources); - std::transform(bloom_filter_byte_ranges_per_source.begin(), - bloom_filter_byte_ranges_per_source.end(), - bloom_filters_per_source.begin(), - [](auto const& ranges) { return ranges.size(); }); + CUDF_EXPECTS(num_sources == bloom_filter_byte_ranges_per_source.size(), + "Encountered mismatch in number of datasources and bloom filter byte range spans"); // Read + parse a single bloom filter header to host and return its bitset-only byte range // `(offset + header_size, num_bytes)`. The reader emits an empty `{0, 0}` placeholder for a chunk @@ -480,7 +405,7 @@ fetch_bloom_filters_to_device_async_impl( if (bloom_range.is_empty()) { return {0, 0}; } auto const header_read_size = std::min(bloom_range.size(), detail::bloom_filter_header_max_size); - auto const header = datasource.host_read(static_cast(bloom_range.offset()), + auto const header = datasource.host_read(static_cast(bloom_range.offset()), static_cast(header_read_size)); auto const header_info = detail::parse_bloom_filter_header({header->data(), header->size()}); CUDF_EXPECTS(header_info.has_value(), "Encountered an invalid bloom filter header"); @@ -488,11 +413,26 @@ fetch_bloom_filters_to_device_async_impl( return {bloom_range.offset() + header_size, static_cast(bitset_size)}; }; - // Read + parse headers per `(source, bloom filter)`; results come back grouped per source. - auto const bitset_byte_ranges_per_source = dispatch_fetch_tasks( - bloom_filters_per_source, [&](std::size_t source_idx, std::size_t bloom_idx) { - return fetch_bitset_range(datasources[source_idx].get(), - bloom_filter_byte_ranges_per_source[source_idx][bloom_idx]); + // Parse each source's bloom filter headers with one task per source, reusing the shared + // `dispatch_fetch_tasks` (sequential for few sources, host worker pool for many). + // + // Trade-off of reusing the per-source dispatcher: a source's bloom-filter header reads are issued + // one at a time within its task - we do NOT split them into separate per-(source, bloom filter) + // `host_read` tasks, so the individual header reads of a single source are not parallelized. Each + // task also reads only its own datasource, so there are no concurrent reads on a single + // datasource. + auto const bitset_byte_ranges_per_source = + dispatch_fetch_tasks(num_sources, [&](std::size_t source_idx) { + auto const& bloom_ranges = bloom_filter_byte_ranges_per_source[source_idx]; + auto& datasource = datasources[source_idx].get(); + std::vector bitset_ranges; + bitset_ranges.reserve(bloom_ranges.size()); + std::transform( + bloom_ranges.begin(), + bloom_ranges.end(), + std::back_inserter(bitset_ranges), + [&](auto const& bloom_range) { return fetch_bitset_range(datasource, bloom_range); }); + return bitset_ranges; }); // Fetch only the header-free bitsets to device. Delegate to the multi-source diff --git a/cpp/tests/io/experimental/hybrid_scan_composer.cpp b/cpp/tests/io/experimental/hybrid_scan_composer.cpp index 98f41975dce5..3f045ab99a77 100644 --- a/cpp/tests/io/experimental/hybrid_scan_composer.cpp +++ b/cpp/tests/io/experimental/hybrid_scan_composer.cpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ From f59589e2bd6872acc655aaea6a39c709b7ad8ec8 Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Fri, 26 Jun 2026 22:33:47 +0200 Subject: [PATCH 08/39] Refactor bloom filter header processing in `fetch_bloom_filters_to_device_async_impl` Simplified comments and improved code clarity by removing outdated explanations regarding bloom filter header handling. The changes focus on enhancing readability while maintaining the functionality of fetching bloom filter data to the device. This aligns with previous efforts to streamline the hybrid scan implementation and improve performance. --- .../io/parquet/io_utils/parquet_io_utils.cpp | 20 +++---------------- 1 file changed, 3 insertions(+), 17 deletions(-) diff --git a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp index ba9da844d27f..aa8f170539cd 100644 --- a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp +++ b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp @@ -395,13 +395,10 @@ fetch_bloom_filters_to_device_async_impl( "Encountered mismatch in number of datasources and bloom filter byte range spans"); // Read + parse a single bloom filter header to host and return its bitset-only byte range - // `(offset + header_size, num_bytes)`. The reader emits an empty `{0, 0}` placeholder for a chunk - // whose column has no bloom filter written (to keep the per-source ranges aligned with the row - // group x column grid), so empty ranges pass through unchanged. Only the header prefix is read to - // host so a (potentially large) bitset is never staged on the host. auto const fetch_bitset_range = [](cudf::io::datasource& datasource, cudf::io::text::byte_range_info const& bloom_range) -> cudf::io::text::byte_range_info { + // placeholder for a chunk whose column has no bloom filter written if (bloom_range.is_empty()) { return {0, 0}; } auto const header_read_size = std::min(bloom_range.size(), detail::bloom_filter_header_max_size); @@ -413,14 +410,7 @@ fetch_bloom_filters_to_device_async_impl( return {bloom_range.offset() + header_size, static_cast(bitset_size)}; }; - // Parse each source's bloom filter headers with one task per source, reusing the shared - // `dispatch_fetch_tasks` (sequential for few sources, host worker pool for many). - // - // Trade-off of reusing the per-source dispatcher: a source's bloom-filter header reads are issued - // one at a time within its task - we do NOT split them into separate per-(source, bloom filter) - // `host_read` tasks, so the individual header reads of a single source are not parallelized. Each - // task also reads only its own datasource, so there are no concurrent reads on a single - // datasource. + // Parse each source's bloom filter headers with one task per source auto const bitset_byte_ranges_per_source = dispatch_fetch_tasks(num_sources, [&](std::size_t source_idx) { auto const& bloom_ranges = bloom_filter_byte_ranges_per_source[source_idx]; @@ -435,11 +425,7 @@ fetch_bloom_filters_to_device_async_impl( return bitset_ranges; }); - // Fetch only the header-free bitsets to device. Delegate to the multi-source - // `fetch_byte_ranges_to_device_async`, which already performs the `vector -> host_span` - // conversion, so the grouped ranges are passed through directly. cuco's bloom filter probe - // requires a 32-byte-aligned bitset; the rmm buffer base is 256-byte aligned and bitset sizes are - // multiples of 32, so each bitset lands at a 32-byte-aligned address. + // Fetch only the header-free bitsets to device auto result = fetch_byte_ranges_to_device_async(datasources, bitset_byte_ranges_per_source, stream, mr); From 518d2782542013f75d0a794e67fa434db01e2681 Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Fri, 26 Jun 2026 23:42:20 +0200 Subject: [PATCH 09/39] Add new function to fetch bloom filters from multiple datasources Implemented a new overload of `fetch_bloom_filters_to_device_async` to retrieve bloom filter bitsets from multiple datasources into device buffers. This function enhances the existing functionality by allowing for the specification of byte ranges for bloom filters per datasource, improving flexibility in data handling. Updated the implementation to convert input vectors into host spans for better performance and clarity. --- cpp/include/cudf/io/parquet_io_utils.hpp | 24 +++++++++++++++++ .../io/parquet/io_utils/parquet_io_utils.cpp | 27 +++++++++++++++++++ 2 files changed, 51 insertions(+) diff --git a/cpp/include/cudf/io/parquet_io_utils.hpp b/cpp/include/cudf/io/parquet_io_utils.hpp index 58ff6529264b..e9bb65d2864c 100644 --- a/cpp/include/cudf/io/parquet_io_utils.hpp +++ b/cpp/include/cudf/io/parquet_io_utils.hpp @@ -173,6 +173,30 @@ fetch_bloom_filters_to_device_async(cudf::io::datasource& datasource, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr); +/** + * @brief Fetches Parquet bloom filter bitsets from multiple datasources into device buffers + * + * @ingroup io_utils + * + * @param datasources Input datasources + * @param bloom_filter_byte_ranges_per_source Byte ranges of complete bloom filters to fetch, one + * vector per datasource. Each byte range must span a complete bloom filter. + * @param stream CUDA stream + * @param mr Device memory resource + * + * @return A tuple containing a vector of device buffers, a vector of vectors of device spans (one + * bitset span per input range per datasource; empty for ranges without a bloom filter), and a + * future to wait on the read tasks + */ +std::tuple, + std::vector>>, + std::future> +fetch_bloom_filters_to_device_async( + cudf::host_span const> datasources, + cudf::host_span const> bloom_filter_byte_ranges_per_source, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr); + /** @} */ // end of group } // namespace io::parquet } // namespace CUDF_EXPORT cudf diff --git a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp index aa8f170539cd..aa0f5099605f 100644 --- a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp +++ b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp @@ -564,4 +564,31 @@ fetch_bloom_filters_to_device_async( return {std::move(buffers), std::move(fetched_byte_ranges.front()), std::move(fut)}; } +std::tuple, + std::vector>>, + std::future> +fetch_bloom_filters_to_device_async( + cudf::host_span const> datasources, + cudf::host_span const> + bloom_filter_byte_ranges_per_source, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + CUDF_FUNC_RANGE(); + + // Convert input vectors into host spans for the implementation + std::vector> + bloom_filter_byte_range_spans_per_source; + bloom_filter_byte_range_spans_per_source.reserve(bloom_filter_byte_ranges_per_source.size()); + for (auto const& ranges : bloom_filter_byte_ranges_per_source) { + bloom_filter_byte_range_spans_per_source.emplace_back(ranges); + } + return fetch_bloom_filters_to_device_async_impl( + datasources, + {bloom_filter_byte_range_spans_per_source.data(), + bloom_filter_byte_range_spans_per_source.size()}, + stream, + mr); +} + } // namespace cudf::io::parquet From d9f37e2dde9ddd59f8c4123b83277e819fc725f0 Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Fri, 26 Jun 2026 23:49:22 +0200 Subject: [PATCH 10/39] Refactor comments in parquet_io_utils to improve clarity Updated the documentation comments in `parquet_io_utils.hpp` and `parquet_io_utils.cpp` to enhance clarity by removing redundant phrases. This change simplifies the return descriptions for the `fetch_byte_ranges_to_device_async` and `fetch_bloom_filters_to_device_async` functions, improving overall readability without altering functionality. --- cpp/include/cudf/io/parquet_io_utils.hpp | 7 +++---- cpp/src/io/parquet/io_utils/parquet_io_utils.cpp | 11 +++++------ 2 files changed, 8 insertions(+), 10 deletions(-) diff --git a/cpp/include/cudf/io/parquet_io_utils.hpp b/cpp/include/cudf/io/parquet_io_utils.hpp index e9bb65d2864c..d6ce8749b22e 100644 --- a/cpp/include/cudf/io/parquet_io_utils.hpp +++ b/cpp/include/cudf/io/parquet_io_utils.hpp @@ -162,8 +162,8 @@ fetch_byte_ranges_to_device_async( * @param stream CUDA stream * @param mr Device memory resource * - * @return A tuple containing the device buffers, the device spans of the bitset data (one per input - * range; empty for ranges without a bloom filter), and a future to wait on the read tasks + * @return A tuple containing the device buffers, the device spans of the bitset data, and a future + * to wait on the read tasks */ std::tuple, std::vector>, @@ -184,8 +184,7 @@ fetch_bloom_filters_to_device_async(cudf::io::datasource& datasource, * @param stream CUDA stream * @param mr Device memory resource * - * @return A tuple containing a vector of device buffers, a vector of vectors of device spans (one - * bitset span per input range per datasource; empty for ranges without a bloom filter), and a + * @return A tuple containing a vector of device buffers, a vector of vectors of device spans, and a * future to wait on the read tasks */ std::tuple, diff --git a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp index aa0f5099605f..3c44bac1356a 100644 --- a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp +++ b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp @@ -583,12 +583,11 @@ fetch_bloom_filters_to_device_async( for (auto const& ranges : bloom_filter_byte_ranges_per_source) { bloom_filter_byte_range_spans_per_source.emplace_back(ranges); } - return fetch_bloom_filters_to_device_async_impl( - datasources, - {bloom_filter_byte_range_spans_per_source.data(), - bloom_filter_byte_range_spans_per_source.size()}, - stream, - mr); + return fetch_bloom_filters_to_device_async_impl(datasources, + {bloom_filter_byte_range_spans_per_source.data(), + bloom_filter_byte_range_spans_per_source.size()}, + stream, + mr); } } // namespace cudf::io::parquet From c568dc2e5b5a6a8a58316831b6c3fe6ec57cdc8e Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Sat, 27 Jun 2026 00:01:59 +0200 Subject: [PATCH 11/39] Updated the stale hybrid scan docs. --- .../cudf/io/experimental/hybrid_scan.hpp | 19 ++++++++++++------- 1 file changed, 12 insertions(+), 7 deletions(-) diff --git a/cpp/include/cudf/io/experimental/hybrid_scan.hpp b/cpp/include/cudf/io/experimental/hybrid_scan.hpp index bab4bf3fb93e..29f48e66f1e0 100644 --- a/cpp/include/cudf/io/experimental/hybrid_scan.hpp +++ b/cpp/include/cudf/io/experimental/hybrid_scan.hpp @@ -117,9 +117,11 @@ enum class use_data_page_mask : bool { * Row group pruning (OPTIONAL): Start with either a list of custom or all row group indices in the * parquet file and optionally filter it using a byte range and/or the filter expression using * column chunk statistics, dictionaries and bloom filters. Byte ranges for column chunk dictionary - * pages and bloom filters within parquet file may be obtained via `secondary_filters_byte_ranges()` - * function. The byte ranges may be read into device buffers and their device spans may be passed - * to the row group filtration functions. + * pages and complete bloom filters (header + bitset) within parquet file may be obtained via + * `secondary_filters_byte_ranges()` function. Dictionary page ranges may be read directly into + * device buffers. Bloom filter ranges should be fetched with + * `parquet::fetch_bloom_filters_to_device_async()`, which strips the serialized + * BloomFilterHeader and returns 32-byte-aligned bitset spans for row group filtration. * @code{.cpp} * // Start with a list of all parquet row group indices from the file footer * auto all_row_group_indices = reader->all_row_groups(options); @@ -166,9 +168,10 @@ enum class use_data_page_mask : bool { * auto bloom_filtered_row_group_indices = std::vector{}; * * if (bloom_filter_byte_ranges.size()) { - * // Fetch bloom filter byte ranges into device buffers and create spans + * // Fetch bloom filter bitsets into device buffers and create spans * auto [bloom_filter_buffers, bloom_filter_data, bloom_filter_tasks] = - * parquet::fetch_byte_ranges_to_device_async(datasource, bloom_filter_byte_ranges, stream, mr); + * parquet::fetch_bloom_filters_to_device_async( + * datasource, bloom_filter_byte_ranges, stream, mr); * bloom_filter_tasks.get(); * * // Prune row groups using bloom filters @@ -386,8 +389,10 @@ class hybrid_scan_reader { * @brief Get byte ranges of bloom filters and dictionary pages (secondary filters) for row group * pruning * - * @note Device buffers for bloom filter byte ranges must be allocated using a 32 byte - * aligned memory resource + * @note Bloom filter byte ranges include the serialized BloomFilterHeader. Before calling + * `filter_row_groups_with_bloom_filters()`, fetch these ranges with + * `parquet::fetch_bloom_filters_to_device_async()` to strip the header and produce + * 32-byte-aligned bitset spans. * * @param row_group_indices Input row groups indices * @param options Parquet reader options From 6d876527737d1a5ff7969fc3dfcd213caf421688 Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Sat, 27 Jun 2026 00:08:46 +0200 Subject: [PATCH 12/39] Formatting --- cpp/include/cudf/io/experimental/hybrid_scan.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/include/cudf/io/experimental/hybrid_scan.hpp b/cpp/include/cudf/io/experimental/hybrid_scan.hpp index 29f48e66f1e0..3057a665bc71 100644 --- a/cpp/include/cudf/io/experimental/hybrid_scan.hpp +++ b/cpp/include/cudf/io/experimental/hybrid_scan.hpp @@ -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 */ From 15ddaeec7a9806b979ee4f43d7f64ce0cd25a134 Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Tue, 30 Jun 2026 20:53:54 +0200 Subject: [PATCH 13/39] Refactor bloom filter fetching in hybrid scan to use aligned memory resources Updated the hybrid scan implementation to fetch 32-byte aligned bloom filter data buffers from input files. This change includes modifications to the fetch functions and related comments for clarity. The copyright notice has also been simplified across multiple files. Changes include: - Replaced `fetch_bloom_filters_to_device_async` with `fetch_byte_ranges_to_device_async` using aligned memory resources. - Updated comments to reflect the new fetching mechanism. - Simplified copyright statements in several files. This refactor enhances memory alignment handling for improved performance. --- .../hybrid_scan/hybrid_scan_composer.cpp | 12 ++- .../hybrid_scan_io/hybrid_scan_composer.cpp | 11 ++- .../cudf/io/experimental/hybrid_scan.hpp | 21 ++--- .../io/parquet/io_utils/parquet_io_utils.cpp | 77 ++++++++++++------- .../io/experimental/hybrid_scan_composer.cpp | 11 ++- 5 files changed, 79 insertions(+), 53 deletions(-) diff --git a/cpp/benchmarks/io/parquet/experimental/hybrid_scan/hybrid_scan_composer.cpp b/cpp/benchmarks/io/parquet/experimental/hybrid_scan/hybrid_scan_composer.cpp index cd5dfd54fe1e..13b9c28914ae 100644 --- a/cpp/benchmarks/io/parquet/experimental/hybrid_scan/hybrid_scan_composer.cpp +++ b/cpp/benchmarks/io/parquet/experimental/hybrid_scan/hybrid_scan_composer.cpp @@ -1,6 +1,6 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. * SPDX-License-Identifier: Apache-2.0 */ @@ -12,6 +12,8 @@ #include #include +#include + #include #include @@ -102,10 +104,12 @@ std::vector apply_row_group_filters( if (filters.contains(hybrid_scan_filter_type::ROW_GROUPS_WITH_BLOOM_FILTERS) and bloom_filter_byte_ranges.size()) { - // Fetch the header-stripped, 32-byte-aligned bloom filter bitsets from the input file + // Fetch 32-byte aligned bloom filter data buffers from the input file buffer + auto constexpr bloom_filter_alignment = rmm::CUDA_ALLOCATION_ALIGNMENT; + auto aligned_mr = rmm::mr::aligned_resource_adaptor(mr, bloom_filter_alignment); auto [bloom_filter_buffers, bloom_filter_data, bloom_read_tasks] = - cudf::io::parquet::fetch_bloom_filters_to_device_async( - datasource, bloom_filter_byte_ranges, stream, mr); + cudf::io::parquet::fetch_byte_ranges_to_device_async( + datasource, bloom_filter_byte_ranges, stream, aligned_mr); bloom_read_tasks.get(); bloom_filtered_row_groups = reader.filter_row_groups_with_bloom_filters( diff --git a/cpp/examples/hybrid_scan_io/hybrid_scan_composer.cpp b/cpp/examples/hybrid_scan_io/hybrid_scan_composer.cpp index 8eac7ea80f25..a26839d9ffd6 100644 --- a/cpp/examples/hybrid_scan_io/hybrid_scan_composer.cpp +++ b/cpp/examples/hybrid_scan_io/hybrid_scan_composer.cpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. * SPDX-License-Identifier: Apache-2.0 */ @@ -16,6 +16,8 @@ #include #include +#include + #include #include @@ -173,13 +175,14 @@ std::vector apply_row_group_filters( bloom_filtered_row_group_indices.reserve(current_row_group_indices.size()); if (filters.contains(hybrid_scan_filter_type::ROW_GROUPS_WITH_BLOOM_FILTERS) and bloom_filter_byte_ranges.size()) { + // Fetch 32-byte aligned bloom filter data buffers from the input file buffer + auto constexpr bloom_filter_alignment = rmm::CUDA_ALLOCATION_ALIGNMENT; + auto aligned_mr = rmm::mr::aligned_resource_adaptor(temp_mr, bloom_filter_alignment); if (verbose) { std::cout << "READER: Filter row groups with bloom filters...\n"; } timer.reset(); nvtxRangePush("fetch_bloom_filter_byte_ranges"); - // Fetch the header-stripped, 32-byte-aligned bloom filter bitsets from the input file auto [bloom_filter_buffers, bloom_filter_data, bloom_read_tasks] = - cudf::io::parquet::fetch_bloom_filters_to_device_async( - datasource, bloom_filter_byte_ranges, stream, temp_mr); + fetch_byte_ranges_async(datasource, bloom_filter_byte_ranges, stream, aligned_mr); bloom_read_tasks.get(); nvtxRangePop(); diff --git a/cpp/include/cudf/io/experimental/hybrid_scan.hpp b/cpp/include/cudf/io/experimental/hybrid_scan.hpp index 3057a665bc71..bab4bf3fb93e 100644 --- a/cpp/include/cudf/io/experimental/hybrid_scan.hpp +++ b/cpp/include/cudf/io/experimental/hybrid_scan.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION. * SPDX-License-Identifier: Apache-2.0 */ @@ -117,11 +117,9 @@ enum class use_data_page_mask : bool { * Row group pruning (OPTIONAL): Start with either a list of custom or all row group indices in the * parquet file and optionally filter it using a byte range and/or the filter expression using * column chunk statistics, dictionaries and bloom filters. Byte ranges for column chunk dictionary - * pages and complete bloom filters (header + bitset) within parquet file may be obtained via - * `secondary_filters_byte_ranges()` function. Dictionary page ranges may be read directly into - * device buffers. Bloom filter ranges should be fetched with - * `parquet::fetch_bloom_filters_to_device_async()`, which strips the serialized - * BloomFilterHeader and returns 32-byte-aligned bitset spans for row group filtration. + * pages and bloom filters within parquet file may be obtained via `secondary_filters_byte_ranges()` + * function. The byte ranges may be read into device buffers and their device spans may be passed + * to the row group filtration functions. * @code{.cpp} * // Start with a list of all parquet row group indices from the file footer * auto all_row_group_indices = reader->all_row_groups(options); @@ -168,10 +166,9 @@ enum class use_data_page_mask : bool { * auto bloom_filtered_row_group_indices = std::vector{}; * * if (bloom_filter_byte_ranges.size()) { - * // Fetch bloom filter bitsets into device buffers and create spans + * // Fetch bloom filter byte ranges into device buffers and create spans * auto [bloom_filter_buffers, bloom_filter_data, bloom_filter_tasks] = - * parquet::fetch_bloom_filters_to_device_async( - * datasource, bloom_filter_byte_ranges, stream, mr); + * parquet::fetch_byte_ranges_to_device_async(datasource, bloom_filter_byte_ranges, stream, mr); * bloom_filter_tasks.get(); * * // Prune row groups using bloom filters @@ -389,10 +386,8 @@ class hybrid_scan_reader { * @brief Get byte ranges of bloom filters and dictionary pages (secondary filters) for row group * pruning * - * @note Bloom filter byte ranges include the serialized BloomFilterHeader. Before calling - * `filter_row_groups_with_bloom_filters()`, fetch these ranges with - * `parquet::fetch_bloom_filters_to_device_async()` to strip the header and produce - * 32-byte-aligned bitset spans. + * @note Device buffers for bloom filter byte ranges must be allocated using a 32 byte + * aligned memory resource * * @param row_group_indices Input row groups indices * @param options Parquet reader options diff --git a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp index 3c44bac1356a..7f1658ea9ec1 100644 --- a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp +++ b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp @@ -51,39 +51,39 @@ namespace cudf::io::parquet { namespace { /** - * @brief Dispatches the fetch task for each source index and collects the results + * @brief Dispatches a fetch task for each task index and collects the results * - * Dispatches sequentially or using host worker pool depending on the number of sources. + * Dispatches sequentially or using host worker pool depending on the number of tasks. * - * @tparam Task Callable invocable as `fetch_task(std::size_t source_idx)` - * @param num_sources Number of sources to process - * @param fetch_task Task to run for each source index - * @return Vector of results, one per source, in source order + * @tparam Task Callable invocable as `fetch_task(std::size_t task_idx)` + * @param num_tasks Number of tasks to dispatch + * @param fetch_task Task to run for each task index + * @return Vector of results, one per task, in task-index order */ template -auto dispatch_fetch_tasks(std::size_t num_sources, Task fetch_task) +auto dispatch_fetch_tasks(std::size_t num_tasks, Task fetch_task) { using result_type = std::invoke_result_t; auto constexpr parallel_threshold = 32; std::vector results; - results.reserve(num_sources); + results.reserve(num_tasks); - if (num_sources < parallel_threshold) { + if (num_tasks < parallel_threshold) { // Run sequentially to avoid task dispatch overhead std::for_each(cuda::counting_iterator(0), - cuda::counting_iterator(num_sources), - [&](std::size_t source_idx) { results.emplace_back(fetch_task(source_idx)); }); + cuda::counting_iterator(num_tasks), + [&](std::size_t task_idx) { results.emplace_back(fetch_task(task_idx)); }); } else { // Dispatch the tasks to the host worker pool std::vector> tasks; - tasks.reserve(num_sources); + tasks.reserve(num_tasks); std::for_each(cuda::counting_iterator(0), - cuda::counting_iterator(num_sources), - [&](std::size_t source_idx) { + cuda::counting_iterator(num_tasks), + [&](std::size_t task_idx) { tasks.emplace_back(cudf::detail::host_worker_pool().submit_task( - [&fetch_task, source_idx]() { return fetch_task(source_idx); })); + [&fetch_task, task_idx]() { return fetch_task(task_idx); })); }); std::transform(tasks.begin(), tasks.end(), std::back_inserter(results), [](auto& task) { return task.get(); @@ -410,21 +410,42 @@ fetch_bloom_filters_to_device_async_impl( return {bloom_range.offset() + header_size, static_cast(bitset_size)}; }; - // Parse each source's bloom filter headers with one task per source - auto const bitset_byte_ranges_per_source = - dispatch_fetch_tasks(num_sources, [&](std::size_t source_idx) { - auto const& bloom_ranges = bloom_filter_byte_ranges_per_source[source_idx]; - auto& datasource = datasources[source_idx].get(); - std::vector bitset_ranges; - bitset_ranges.reserve(bloom_ranges.size()); - std::transform( - bloom_ranges.begin(), - bloom_ranges.end(), - std::back_inserter(bitset_ranges), - [&](auto const& bloom_range) { return fetch_bitset_range(datasource, bloom_range); }); - return bitset_ranges; + // Flatten to one (source index, bloom filter byte range) entry per bloom filter + std::vector> bloom_header_tasks; + bloom_header_tasks.reserve(std::accumulate( + bloom_filter_byte_ranges_per_source.begin(), + bloom_filter_byte_ranges_per_source.end(), + std::size_t{0}, + [](auto acc, auto const& bloom_ranges) { return acc + bloom_ranges.size(); })); + for (std::size_t source_idx = 0; source_idx < num_sources; ++source_idx) { + auto const& bloom_ranges = bloom_filter_byte_ranges_per_source[source_idx]; + std::transform(bloom_ranges.begin(), + bloom_ranges.end(), + std::back_inserter(bloom_header_tasks), + [source_idx](auto const& bloom_range) { + return std::tuple{source_idx, bloom_range}; + }); + } + + // Dispatch one header read+parse per bloom filter + auto const flat_bitset_ranges = + dispatch_fetch_tasks(bloom_header_tasks.size(), [&](std::size_t task_idx) { + auto const& [source_idx, bloom_range] = bloom_header_tasks[task_idx]; + return fetch_bitset_range(datasources[source_idx].get(), bloom_range); }); + // Regroup the flat results into per-source bitset ranges + std::vector> bitset_byte_ranges_per_source; + bitset_byte_ranges_per_source.reserve(num_sources); + std::transform(bloom_filter_byte_ranges_per_source.begin(), + bloom_filter_byte_ranges_per_source.end(), + std::back_inserter(bitset_byte_ranges_per_source), + [next = flat_bitset_ranges.begin()](auto const& bloom_ranges) mutable { + auto const first = next; + std::advance(next, bloom_ranges.size()); + return std::vector(first, next); + }); + // Fetch only the header-free bitsets to device auto result = fetch_byte_ranges_to_device_async(datasources, bitset_byte_ranges_per_source, stream, mr); diff --git a/cpp/tests/io/experimental/hybrid_scan_composer.cpp b/cpp/tests/io/experimental/hybrid_scan_composer.cpp index 3f045ab99a77..5365988220d9 100644 --- a/cpp/tests/io/experimental/hybrid_scan_composer.cpp +++ b/cpp/tests/io/experimental/hybrid_scan_composer.cpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. * SPDX-License-Identifier: Apache-2.0 */ @@ -104,10 +104,13 @@ auto apply_hybrid_scan_filters(cudf::io::datasource& datasource, std::vector bloom_filtered_row_group_indices; bloom_filtered_row_group_indices.reserve(current_row_group_indices.size()); if (bloom_filter_byte_ranges.size()) { - // Fetch the header-stripped, 32-byte-aligned bloom filter bitsets from the input file + // Fetch 32 byte aligned bloom filter data buffers from the input file buffer + auto aligned_mr = rmm::mr::aligned_resource_adaptor(cudf::get_current_device_resource_ref(), + bloom_filter_alignment); + auto [bloom_filter_buffers, bloom_filter_data, bloom_read_tasks] = - cudf::io::parquet::fetch_bloom_filters_to_device_async( - datasource, bloom_filter_byte_ranges, stream, mr); + cudf::io::parquet::fetch_byte_ranges_to_device_async( + datasource, bloom_filter_byte_ranges, stream, aligned_mr); bloom_read_tasks.get(); // Filter row groups with bloom filters From 4ee65aa3eaa47cb8e0e1968b15386d74a826c789 Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Wed, 1 Jul 2026 10:50:27 +0200 Subject: [PATCH 14/39] Refactor bloom filter handling in Parquet I/O - Updated function signatures in `parquet_io_utils.hpp` and `reader_impl_helpers.hpp` to use aligned memory resources for bloom filter buffers. - Modified `read_bloom_filters` to return both device buffers and spans for bloom filter bitsets. - Adjusted `fetch_bloom_filters_to_device_async_impl` to accommodate new bloom filter byte range handling. - Removed deprecated code related to bloom filter data reading in `bloom_filter_reader.cu`. This change enhances memory alignment and improves the efficiency of bloom filter operations in the Parquet I/O module. --- cpp/include/cudf/io/parquet_io_utils.hpp | 14 +- cpp/src/io/parquet/bloom_filter_reader.cu | 230 +++++------------- .../io/parquet/io_utils/parquet_io_utils.cpp | 126 +++++----- cpp/src/io/parquet/predicate_pushdown.cpp | 30 +-- cpp/src/io/parquet/reader_impl_helpers.hpp | 21 +- 5 files changed, 154 insertions(+), 267 deletions(-) diff --git a/cpp/include/cudf/io/parquet_io_utils.hpp b/cpp/include/cudf/io/parquet_io_utils.hpp index ea26873b5b83..df8288b56dae 100644 --- a/cpp/include/cudf/io/parquet_io_utils.hpp +++ b/cpp/include/cudf/io/parquet_io_utils.hpp @@ -157,11 +157,14 @@ fetch_byte_ranges_to_device_async( * * @ingroup io_utils * + * @note Device buffers for bloom filter byte ranges must be allocated using a 32 byte aligned + * memory resource + * * @param datasource Input datasource * @param bloom_filter_byte_ranges Byte ranges of complete bloom filters to fetch, must span a * complete bloom filter * @param stream CUDA stream - * @param mr Device memory resource + * @param aligned_mr Device memory resource to allocate aligned memory for bloom filters * * @return A tuple containing the device buffers, the device spans of the bitset data, and a future * to wait on the read tasks @@ -172,18 +175,21 @@ std::tuple, fetch_bloom_filters_to_device_async(cudf::io::datasource& datasource, cudf::host_span bloom_filter_byte_ranges, rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr); + rmm::device_async_resource_ref aligned_mr); /** * @brief Fetches Parquet bloom filter bitsets from multiple datasources into device buffers * * @ingroup io_utils * + * @note Device buffers for bloom filter byte ranges must be allocated using a 32 byte aligned + * memory resource + * * @param datasources Input datasources * @param bloom_filter_byte_ranges_per_source Byte ranges of complete bloom filters to fetch, one * vector per datasource. Each byte range must span a complete bloom filter. * @param stream CUDA stream - * @param mr Device memory resource + * @param aligned_mr Device memory resource to allocate aligned memory for bloom filters * * @return A tuple containing a vector of device buffers, a vector of vectors of device spans, and a * future to wait on the read tasks @@ -195,7 +201,7 @@ fetch_bloom_filters_to_device_async( cudf::host_span const> datasources, cudf::host_span const> bloom_filter_byte_ranges_per_source, rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr); + rmm::device_async_resource_ref aligned_mr); /** @} */ // end of group } // namespace io::parquet diff --git a/cpp/src/io/parquet/bloom_filter_reader.cu b/cpp/src/io/parquet/bloom_filter_reader.cu index 7e91f95a35b9..18c6f981799a 100644 --- a/cpp/src/io/parquet/bloom_filter_reader.cu +++ b/cpp/src/io/parquet/bloom_filter_reader.cu @@ -15,6 +15,7 @@ #include #include #include +#include #include #include #include @@ -29,6 +30,7 @@ #include #include +#include #include #include #include @@ -276,122 +278,6 @@ class bloom_filter_expression_converter : public equality_literals_collector { std::unique_ptr _always_true; }; -/** - * @brief Reads bloom filter data to device. - * - * @param sources Dataset sources - * @param num_chunks Number of total column chunks to read - * @param bloom_filter_data Device buffers to hold bloom filter bitsets for each chunk - * @param bloom_filter_offsets Bloom filter offsets for all chunks - * @param bloom_filter_sizes Bloom filter sizes for all chunks - * @param chunk_source_map Association between each column chunk and its source - * @param stream CUDA stream used for device memory operations and kernel launches - * @param aligned_mr Aligned device memory resource to allocate bloom filter buffers - */ -void read_bloom_filter_data(host_span const> sources, - std::size_t num_chunks, - cudf::host_span bloom_filter_data, - cudf::host_span> bloom_filter_offsets, - cudf::host_span> bloom_filter_sizes, - std::vector const& chunk_source_map, - rmm::cuda_stream_view stream, - rmm::device_async_resource_ref aligned_mr) -{ - // Using `arrow_filter_policy` with a temporary `cuda::std::byte` key type to extract bloom - // filter properties - using policy_type = arrow_filter_policy; - auto constexpr filter_block_alignment = - alignof(cuco::bloom_filter_ref, - cuco::thread_scope_thread, - policy_type>::filter_block_type); - - // Read tasks for bloom filter data - std::vector> read_tasks; - - // Read bloom filters for all column chunks - std::for_each( - cuda::counting_iterator{0}, - cuda::counting_iterator{num_chunks}, - [&](auto const chunk) { - // If bloom filter offset absent, fill in an empty buffer and skip ahead - if (not bloom_filter_offsets[chunk].has_value()) { - bloom_filter_data[chunk] = {}; - return; - } - // Read bloom filter iff present - auto const bloom_filter_offset = bloom_filter_offsets[chunk].value(); - - // If the bloom filter size (header + bitset) is available, read the entire thing. Else read - // the max header size, which contains the entire header and may contain the entire bitset. - auto const initial_read_size = - static_cast(bloom_filter_sizes[chunk].value_or(bloom_filter_header_max_size)); - - // Read an initial buffer from source - auto& source = sources[chunk_source_map[chunk]]; - auto buffer = source->host_read(bloom_filter_offset, initial_read_size); - - // Deserialize and validate the bloom filter header from the buffer. - auto const header_info = parse_bloom_filter_header({buffer->data(), buffer->size()}); - - // Do not read if the bloom filter is invalid - if (not header_info.has_value()) { - bloom_filter_data[chunk] = {}; - CUDF_LOG_WARN("Encountered an invalid bloom filter header. Skipping"); - return; - } - - // Bloom filter header and bitset sizes - auto const bloom_filter_header_size = header_info->first; - auto const bitset_size = header_info->second; - - // Check if we already read in the filter bitset in the initial read. - if (initial_read_size >= bloom_filter_header_size + bitset_size) { - bloom_filter_data[chunk] = rmm::device_buffer{ - buffer->data() + bloom_filter_header_size, bitset_size, stream, aligned_mr}; - // The allocated bloom filter buffer must be aligned - CUDF_EXPECTS(reinterpret_cast(bloom_filter_data[chunk].data()) % - filter_block_alignment == - 0, - "Encountered misaligned bloom filter block"); - } - // Read the bitset from datasource. - else { - auto const bitset_offset = bloom_filter_offset + bloom_filter_header_size; - // Directly read to device if preferred - if (source->is_device_read_preferred(bitset_size)) { - bloom_filter_data[chunk] = rmm::device_buffer{bitset_size, stream, aligned_mr}; - // The allocated bloom filter buffer must be aligned - CUDF_EXPECTS(reinterpret_cast(bloom_filter_data[chunk].data()) % - filter_block_alignment == - 0, - "Encountered misaligned bloom filter block"); - auto future_read_size = - source->device_read_async(bitset_offset, - bitset_size, - static_cast(bloom_filter_data[chunk].data()), - stream); - - read_tasks.emplace_back(std::move(future_read_size)); - } else { - buffer = source->host_read(bitset_offset, bitset_size); - bloom_filter_data[chunk] = - rmm::device_buffer{buffer->data(), buffer->size(), stream, aligned_mr}; - // The allocated bloom filter buffer must be aligned - CUDF_EXPECTS(reinterpret_cast(bloom_filter_data[chunk].data()) % - filter_block_alignment == - 0, - "Encountered misaligned bloom filter block"); - } - } - }); - - // Read task sync function - for (auto& task : read_tasks) { - task.get(); - } -} - } // namespace std::optional> parse_bloom_filter_header( @@ -430,7 +316,8 @@ std::size_t aggregate_reader_metadata::get_bloom_filter_alignment() const return std::max(alignment, rmm::CUDA_ALLOCATION_ALIGNMENT); } -std::vector aggregate_reader_metadata::read_bloom_filters( +std::pair, std::vector>> +aggregate_reader_metadata::read_bloom_filters( host_span const> sources, host_span const> row_group_indices, host_span column_schemas, @@ -442,64 +329,71 @@ std::vector aggregate_reader_metadata::read_bloom_filters( auto const num_input_columns = column_schemas.size(); auto const num_chunks = total_row_groups * num_input_columns; - // Association between each column chunk and its source - std::vector chunk_source_map(num_chunks); - - // Keep track of column chunk file offsets - std::vector> bloom_filter_offsets(num_chunks); - std::vector> bloom_filter_sizes(num_chunks); - - // Gather all bloom filter offsets and sizes. - size_type chunk_count = 0; - // Flag to check if we have at least one valid bloom filter offset auto have_bloom_filters = false; - + // Build complete bloom filter byte ranges (header + bitset) for every column chunk + std::vector> bloom_filter_byte_ranges_per_source( + row_group_indices.size()); // For all data sources - std::for_each(cuda::counting_iterator{0}, - cuda::counting_iterator{row_group_indices.size()}, - [&](auto const src_index) { - // Get all row group indices in the data source - auto const& rg_indices = row_group_indices[src_index]; - // For all row groups - std::for_each(rg_indices.cbegin(), rg_indices.cend(), [&](auto const rg_index) { - // For all column chunks - std::for_each( - column_schemas.begin(), column_schemas.end(), [&](auto const schema_idx) { - auto& col_meta = get_column_metadata(rg_index, src_index, schema_idx); - - // Get bloom filter offsets and sizes - bloom_filter_offsets[chunk_count] = col_meta.bloom_filter_offset; - bloom_filter_sizes[chunk_count] = col_meta.bloom_filter_length; - - // Set `have_bloom_filters` if `bloom_filter_offset` is valid - if (col_meta.bloom_filter_offset.has_value()) { have_bloom_filters = true; } - - // Map each column chunk to its source index - chunk_source_map[chunk_count] = src_index; - chunk_count++; - }); - }); - }); + std::for_each( + cuda::counting_iterator{0}, + cuda::counting_iterator{row_group_indices.size()}, + [&](auto const src_index) { + auto const& rg_indices = row_group_indices[src_index]; + auto& source_ranges = bloom_filter_byte_ranges_per_source[src_index]; + source_ranges.reserve(rg_indices.size() * num_input_columns); + // For all row groups in the source + std::for_each(rg_indices.cbegin(), rg_indices.cend(), [&](auto const rg_index) { + // For all column chunks in the row group + std::for_each(column_schemas.begin(), column_schemas.end(), [&](auto const schema_idx) { + auto const& col_meta = get_column_metadata(rg_index, src_index, schema_idx); + if (col_meta.bloom_filter_offset.has_value()) { + have_bloom_filters = true; + // When the length is absent, read up to the max header size to recover the bitset size. + auto const length = col_meta.bloom_filter_length.has_value() + ? static_cast(col_meta.bloom_filter_length.value()) + : bloom_filter_header_max_size; + source_ranges.push_back( + cudf::io::text::byte_range_info{col_meta.bloom_filter_offset.value(), length}); + } else { + source_ranges.push_back(cudf::io::text::byte_range_info{0, 0}); + } + }); + }); + }); // Exit early if we don't have any bloom filters if (not have_bloom_filters) { return {}; } - // Vector to hold bloom filter data - std::vector bloom_filter_data(num_chunks); - - // Read bloom filter data - read_bloom_filter_data(sources, - num_chunks, - bloom_filter_data, - bloom_filter_offsets, - bloom_filter_sizes, - chunk_source_map, - stream, - aligned_mr); - - // Return bloom filter data - return bloom_filter_data; + // Fetch the header-stripped, 32-byte-aligned bloom filter bitsets to device via the shared + // utility + std::vector> datasource_refs; + datasource_refs.reserve(sources.size()); + std::transform( + sources.begin(), sources.end(), std::back_inserter(datasource_refs), [](auto const& source) { + return std::ref(*source); + }); + + auto [bloom_filter_buffers, bitset_spans_per_source, fetch_task] = + fetch_bloom_filters_to_device_async( + datasource_refs, bloom_filter_byte_ranges_per_source, stream, aligned_mr); + fetch_task.get(); + + // Flatten the per-source bitset spans into per-chunk order + std::vector> bloom_filter_data; + bloom_filter_data.reserve(num_chunks); + std::for_each( + bitset_spans_per_source.begin(), bitset_spans_per_source.end(), [&](auto const& source_spans) { + std::transform(source_spans.begin(), + source_spans.end(), + std::back_inserter(bloom_filter_data), + [](auto const& span) { + return cudf::device_span{ + reinterpret_cast(span.data()), span.size()}; + }); + }); + + return {std::move(bloom_filter_buffers), std::move(bloom_filter_data)}; } std::optional>> aggregate_reader_metadata::apply_bloom_filters( diff --git a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp index 0ca86a354b71..82b7ef1e6242 100644 --- a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp +++ b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp @@ -51,17 +51,17 @@ namespace cudf::io::parquet { namespace { /** - * @brief Dispatches a fetch task for each task index and collects the results + * @brief Dispatches a task for each task index and collects the results * * Dispatches sequentially or using host worker pool depending on the number of tasks. * - * @tparam Task Callable invocable as `fetch_task(std::size_t task_idx)` + * @tparam Task Callable invocable as `task(std::size_t task_idx)` * @param num_tasks Number of tasks to dispatch - * @param fetch_task Task to run for each task index + * @param task Task to run for each task index * @return Vector of results, one per task, in task-index order */ template -auto dispatch_fetch_tasks(std::size_t num_tasks, Task fetch_task) +auto dispatch_tasks(std::size_t num_tasks, Task task) { using result_type = std::invoke_result_t; @@ -74,19 +74,19 @@ auto dispatch_fetch_tasks(std::size_t num_tasks, Task fetch_task) // Run sequentially to avoid task dispatch overhead std::for_each(cuda::counting_iterator(0), cuda::counting_iterator(num_tasks), - [&](std::size_t task_idx) { results.emplace_back(fetch_task(task_idx)); }); + [&](std::size_t task_idx) { results.emplace_back(task(task_idx)); }); } else { // Dispatch the tasks to the host worker pool - std::vector> tasks; - tasks.reserve(num_tasks); + std::vector> futures; + futures.reserve(num_tasks); std::for_each(cuda::counting_iterator(0), cuda::counting_iterator(num_tasks), [&](std::size_t task_idx) { - tasks.emplace_back(cudf::detail::host_worker_pool().submit_task( - [&fetch_task, task_idx]() { return fetch_task(task_idx); })); + futures.emplace_back(cudf::detail::host_worker_pool().submit_task( + [&task, task_idx]() { return task(task_idx); })); }); - std::transform(tasks.begin(), tasks.end(), std::back_inserter(results), [](auto& task) { - return task.get(); + std::transform(futures.begin(), futures.end(), std::back_inserter(results), [](auto& fut) { + return fut.get(); }); } return results; @@ -152,7 +152,7 @@ std::vector> fetch_footers_to_host return cudf::io::datasource::buffer::create(std::move(footer_bytes)); }; - return dispatch_fetch_tasks(datasources.size(), [&](std::size_t source_idx) { + return dispatch_tasks(datasources.size(), [&](std::size_t source_idx) { return fetch_footer(datasources[source_idx].get()); }); } @@ -181,7 +181,7 @@ std::vector> fetch_page_indexes_to return datasource.host_read(page_index_bytes.offset(), page_index_bytes.size()); }; - return dispatch_fetch_tasks(datasources.size(), [&](std::size_t source_idx) { + return dispatch_tasks(datasources.size(), [&](std::size_t source_idx) { return fetch_page_index(datasources[source_idx].get(), page_index_bytes_per_source[source_idx]); }); } @@ -388,7 +388,7 @@ fetch_bloom_filters_to_device_async_impl( cudf::host_span const> bloom_filter_byte_ranges_per_source, rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr) + rmm::device_async_resource_ref aligned_mr) { auto const num_sources = datasources.size(); CUDF_EXPECTS(num_sources == bloom_filter_byte_ranges_per_source.size(), @@ -410,58 +410,54 @@ fetch_bloom_filters_to_device_async_impl( return {bloom_range.offset() + header_size, static_cast(bitset_size)}; }; - // Flatten to one (source index, bloom filter byte range) entry per bloom filter - std::vector> bloom_header_tasks; - bloom_header_tasks.reserve(std::accumulate( - bloom_filter_byte_ranges_per_source.begin(), - bloom_filter_byte_ranges_per_source.end(), - std::size_t{0}, - [](auto acc, auto const& bloom_ranges) { return acc + bloom_ranges.size(); })); - for (std::size_t source_idx = 0; source_idx < num_sources; ++source_idx) { - auto const& bloom_ranges = bloom_filter_byte_ranges_per_source[source_idx]; - std::transform(bloom_ranges.begin(), - bloom_ranges.end(), - std::back_inserter(bloom_header_tasks), - [source_idx](auto const& bloom_range) { - return std::tuple{source_idx, bloom_range}; - }); - } - - // Dispatch one header read+parse per bloom filter - auto const flat_bitset_ranges = - dispatch_fetch_tasks(bloom_header_tasks.size(), [&](std::size_t task_idx) { - auto const& [source_idx, bloom_range] = bloom_header_tasks[task_idx]; - return fetch_bitset_range(datasources[source_idx].get(), bloom_range); - }); - - // Regroup the flat results into per-source bitset ranges - std::vector> bitset_byte_ranges_per_source; - bitset_byte_ranges_per_source.reserve(num_sources); + // Flatten to one (source index, index within source, bloom filter byte range) entry per bloom + // filter + std::vector> + bloom_header_tasks; + bloom_header_tasks.reserve( + std::accumulate(bloom_filter_byte_ranges_per_source.begin(), + bloom_filter_byte_ranges_per_source.end(), + std::size_t{0}, + [](auto acc, auto const& bloom_ranges) { return acc + bloom_ranges.size(); })); + std::for_each(cuda::counting_iterator(0), + cuda::counting_iterator(num_sources), + [&](std::size_t outer_idx) { + auto const& bloom_ranges = bloom_filter_byte_ranges_per_source[outer_idx]; + std::transform(cuda::counting_iterator(0), + cuda::counting_iterator(bloom_ranges.size()), + std::back_inserter(bloom_header_tasks), + [&](std::size_t inner_idx) { + return std::tuple{outer_idx, inner_idx, bloom_ranges[inner_idx]}; + }); + }); + + // Mirror the input so each fetched range can be gathered back to its slot + std::vector> bitset_byte_ranges_per_source( + num_sources); std::transform(bloom_filter_byte_ranges_per_source.begin(), bloom_filter_byte_ranges_per_source.end(), - std::back_inserter(bitset_byte_ranges_per_source), - [next = flat_bitset_ranges.begin()](auto const& bloom_ranges) mutable { - auto const first = next; - std::advance(next, bloom_ranges.size()); - return std::vector(first, next); + bitset_byte_ranges_per_source.begin(), + [](auto const& bloom_ranges) { + return std::vector(bloom_ranges.size()); }); + // Dispatch one header read+parse per bloom filter, then gather each bitset range into its slot + auto const flat_bitset_ranges = + dispatch_tasks(bloom_header_tasks.size(), [&](std::size_t task_idx) { + auto const& task = bloom_header_tasks[task_idx]; + return fetch_bitset_range(datasources[std::get<0>(task)].get(), std::get<2>(task)); + }); + std::for_each(cuda::counting_iterator(0), + cuda::counting_iterator(bloom_header_tasks.size()), + [&](std::size_t task_idx) { + auto const& task = bloom_header_tasks[task_idx]; + bitset_byte_ranges_per_source[std::get<0>(task)][std::get<1>(task)] = + flat_bitset_ranges[task_idx]; + }); + // Fetch only the header-free bitsets to device - auto result = - fetch_byte_ranges_to_device_async(datasources, bitset_byte_ranges_per_source, stream, mr); - - auto const& bitset_spans_per_source = std::get<1>(result); - CUDF_EXPECTS(std::all_of(bitset_spans_per_source.begin(), - bitset_spans_per_source.end(), - [](auto const& spans) { - return std::all_of(spans.begin(), spans.end(), [](auto const& span) { - return span.empty() or - (reinterpret_cast(span.data()) % 32) == 0; - }); - }), - "Bloom filter bitset is not 32-byte aligned"); - - return result; + return fetch_byte_ranges_to_device_async( + datasources, bitset_byte_ranges_per_source, stream, aligned_mr); } } // namespace @@ -566,7 +562,7 @@ fetch_bloom_filters_to_device_async( cudf::io::datasource& datasource, cudf::host_span bloom_filter_byte_ranges, rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr) + rmm::device_async_resource_ref aligned_mr) { CUDF_FUNC_RANGE(); @@ -579,7 +575,7 @@ fetch_bloom_filters_to_device_async( {datasources.data(), datasources.size()}, {bloom_filter_byte_ranges_per_source.data(), bloom_filter_byte_ranges_per_source.size()}, stream, - mr); + aligned_mr); return {std::move(buffers), std::move(fetched_byte_ranges.front()), std::move(fut)}; } @@ -592,7 +588,7 @@ fetch_bloom_filters_to_device_async( cudf::host_span const> bloom_filter_byte_ranges_per_source, rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr) + rmm::device_async_resource_ref aligned_mr) { CUDF_FUNC_RANGE(); @@ -607,7 +603,7 @@ fetch_bloom_filters_to_device_async( {bloom_filter_byte_range_spans_per_source.data(), bloom_filter_byte_range_spans_per_source.size()}, stream, - mr); + aligned_mr); } } // namespace cudf::io::parquet diff --git a/cpp/src/io/parquet/predicate_pushdown.cpp b/cpp/src/io/parquet/predicate_pushdown.cpp index 9aab5ae708cd..eecb9b8d7a65 100644 --- a/cpp/src/io/parquet/predicate_pushdown.cpp +++ b/cpp/src/io/parquet/predicate_pushdown.cpp @@ -336,30 +336,20 @@ aggregate_reader_metadata::filter_row_groups( // Read a vector of bloom filter bitset device buffers for all columns with equality // predicate(s) across all row groups - auto bloom_filter_buffers = read_bloom_filters(sources, - bloom_filter_input_row_groups, - equality_col_schemas, - num_stats_filtered_row_groups, - stream, - aligned_mr); - - // No bloom filter buffers, return early - if (bloom_filter_buffers.empty()) { + auto const [bloom_filter_buffers, bloom_filter_data] = + read_bloom_filters(sources, + bloom_filter_input_row_groups, + equality_col_schemas, + num_stats_filtered_row_groups, + stream, + aligned_mr); + + // No bloom filters, return early + if (bloom_filter_data.empty()) { return {stats_filtered_row_groups, {std::make_optional(num_stats_filtered_row_groups), std::nullopt}}; } - // Create spans from bloom filter buffers - std::vector> bloom_filter_data; - bloom_filter_data.reserve(bloom_filter_buffers.size()); - std::transform(bloom_filter_buffers.begin(), - bloom_filter_buffers.end(), - std::back_inserter(bloom_filter_data), - [](auto& buffer) { - return cudf::device_span( - static_cast(buffer.data()), buffer.size()); - }); - // Apply bloom filtering on the output row groups from stats filter auto const bloom_filtered_row_groups = apply_bloom_filters(bloom_filter_data, bloom_filter_input_row_groups, diff --git a/cpp/src/io/parquet/reader_impl_helpers.hpp b/cpp/src/io/parquet/reader_impl_helpers.hpp index a4de3c0cc596..78378d47f6a3 100644 --- a/cpp/src/io/parquet/reader_impl_helpers.hpp +++ b/cpp/src/io/parquet/reader_impl_helpers.hpp @@ -247,16 +247,17 @@ class aggregate_reader_metadata { * @param stream CUDA stream used for device memory operations and kernel launches * @param aligned_mr Aligned device memory resource to allocate bloom filter buffers * - * @return A flattened list of bloom filter bitset device buffers for each predicate column across - * row group - */ - [[nodiscard]] std::vector read_bloom_filters( - host_span const> sources, - host_span const> row_group_indices, - host_span column_schemas, - size_type num_row_groups, - rmm::cuda_stream_view stream, - rmm::device_async_resource_ref aligned_mr) const; + * @return A pair of the device buffers backing the bloom filter bitsets and a flattened, + * per-chunk list of bitset device spans (empty spans for chunks without a bloom filter) + */ + [[nodiscard]] std::pair, + std::vector>> + read_bloom_filters(host_span const> sources, + host_span const> row_group_indices, + host_span column_schemas, + size_type num_row_groups, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref aligned_mr) const; /** * @brief Collects Parquet types for the columns with the specified schema indices From cb9067f9f737610466adfc809ae20862061e1986 Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Wed, 1 Jul 2026 19:11:41 +0200 Subject: [PATCH 15/39] Use speculative single-read for Parquet bloom filter fetch Replace the two-phase bloom filter fetch (a header read followed by a separate bitset read) with a speculative single read: read the whole filter in one host read when its serialized size is within the reader's speculative read size (metadata_size_hint) and copy the header-stripped bitset to device; larger filters read only the header and stream the bitset to device. The two-phase path issues two reads per filter, which roughly doubles the fetch time on high-latency/remote storage where the extra round trip dominates. The public interface and all callers are unchanged. --- .../io/parquet/io_utils/parquet_io_utils.cpp | 132 ++++++++++++------ 1 file changed, 87 insertions(+), 45 deletions(-) diff --git a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp index 82b7ef1e6242..48d2427513c8 100644 --- a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp +++ b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp @@ -394,27 +394,16 @@ fetch_bloom_filters_to_device_async_impl( CUDF_EXPECTS(num_sources == bloom_filter_byte_ranges_per_source.size(), "Encountered mismatch in number of datasources and bloom filter byte range spans"); - // Read + parse a single bloom filter header to host and return its bitset-only byte range - auto const fetch_bitset_range = - [](cudf::io::datasource& datasource, - cudf::io::text::byte_range_info const& bloom_range) -> cudf::io::text::byte_range_info { - // placeholder for a chunk whose column has no bloom filter written - if (bloom_range.is_empty()) { return {0, 0}; } - auto const header_read_size = - std::min(bloom_range.size(), detail::bloom_filter_header_max_size); - auto const header = datasource.host_read(static_cast(bloom_range.offset()), - static_cast(header_read_size)); - auto const header_info = detail::parse_bloom_filter_header({header->data(), header->size()}); - CUDF_EXPECTS(header_info.has_value(), "Encountered an invalid bloom filter header"); - auto const [header_size, bitset_size] = header_info.value(); - return {bloom_range.offset() + header_size, static_cast(bitset_size)}; - }; + // Read a filter whole in one host read (up to the reader's speculative read size) and copy its + // bitset to device; larger filters read only the header and stream the bitset to device. + auto const speculative_read_size = + std::max(cudf::io::parquet::metadata_size_hint(), + static_cast(detail::bloom_filter_header_max_size)); - // Flatten to one (source index, index within source, bloom filter byte range) entry per bloom - // filter + // Flatten to one (source index, index within source, bloom filter byte range) task per filter std::vector> - bloom_header_tasks; - bloom_header_tasks.reserve( + bloom_filter_tasks; + bloom_filter_tasks.reserve( std::accumulate(bloom_filter_byte_ranges_per_source.begin(), bloom_filter_byte_ranges_per_source.end(), std::size_t{0}, @@ -425,39 +414,92 @@ fetch_bloom_filters_to_device_async_impl( auto const& bloom_ranges = bloom_filter_byte_ranges_per_source[outer_idx]; std::transform(cuda::counting_iterator(0), cuda::counting_iterator(bloom_ranges.size()), - std::back_inserter(bloom_header_tasks), + std::back_inserter(bloom_filter_tasks), [&](std::size_t inner_idx) { return std::tuple{outer_idx, inner_idx, bloom_ranges[inner_idx]}; }); }); - // Mirror the input so each fetched range can be gathered back to its slot - std::vector> bitset_byte_ranges_per_source( - num_sources); - std::transform(bloom_filter_byte_ranges_per_source.begin(), - bloom_filter_byte_ranges_per_source.end(), - bitset_byte_ranges_per_source.begin(), - [](auto const& bloom_ranges) { - return std::vector(bloom_ranges.size()); - }); - - // Dispatch one header read+parse per bloom filter, then gather each bitset range into its slot - auto const flat_bitset_ranges = - dispatch_tasks(bloom_header_tasks.size(), [&](std::size_t task_idx) { - auto const& task = bloom_header_tasks[task_idx]; - return fetch_bitset_range(datasources[std::get<0>(task)].get(), std::get<2>(task)); + // Phase 1: dispatch one speculative host read + header parse per bloom filter (source -> host) + struct speculative_read { + std::unique_ptr host_buffer; // speculative prefix, null if empty + int64_t header_size{}; + std::size_t bitset_size{}; + bool covered{}; // the whole bitset is already present in `host_buffer` + }; + auto speculative_reads = + dispatch_tasks(bloom_filter_tasks.size(), [&](std::size_t task_idx) -> speculative_read { + auto const& [source_idx, inner_idx, bloom_range] = bloom_filter_tasks[task_idx]; + // placeholder for a chunk whose column has no bloom filter written + if (bloom_range.is_empty()) { return {}; } + auto const total_size = static_cast(bloom_range.size()); + // Whole filter in one read when its known length is within the cap; else just the header. + auto const read_size = total_size <= speculative_read_size + ? total_size + : static_cast(detail::bloom_filter_header_max_size); + auto host_buffer = datasources[source_idx].get().host_read( + static_cast(bloom_range.offset()), read_size); + auto const header_info = + detail::parse_bloom_filter_header({host_buffer->data(), host_buffer->size()}); + CUDF_EXPECTS(header_info.has_value(), "Encountered an invalid bloom filter header"); + auto const [header_size, bitset_size] = header_info.value(); + auto const covered = read_size >= static_cast(header_size) + bitset_size; + return {std::move(host_buffer), header_size, bitset_size, covered}; }); - std::for_each(cuda::counting_iterator(0), - cuda::counting_iterator(bloom_header_tasks.size()), - [&](std::size_t task_idx) { - auto const& task = bloom_header_tasks[task_idx]; - bitset_byte_ranges_per_source[std::get<0>(task)][std::get<1>(task)] = - flat_bitset_ranges[task_idx]; - }); - // Fetch only the header-free bitsets to device - return fetch_byte_ranges_to_device_async( - datasources, bitset_byte_ranges_per_source, stream, aligned_mr); + // Phase 2: materialize aligned device bitsets — copy covered filters from the host read, stream + // the rest to device. + std::vector bloom_filter_buffers; + bloom_filter_buffers.reserve(bloom_filter_tasks.size()); + std::vector bitset_spans_per_source(num_sources); + std::transform( + bloom_filter_byte_ranges_per_source.begin(), + bloom_filter_byte_ranges_per_source.end(), + bitset_spans_per_source.begin(), + [](auto const& bloom_ranges) { return device_spans_per_source_type(bloom_ranges.size()); }); + std::vector> device_read_tasks; + + std::for_each( + cuda::counting_iterator(0), + cuda::counting_iterator(bloom_filter_tasks.size()), + [&](std::size_t task_idx) { + auto const& [source_idx, inner_idx, bloom_range] = bloom_filter_tasks[task_idx]; + auto& spec = speculative_reads[task_idx]; + // empty bloom filter -> leave an empty span in place + if (spec.host_buffer == nullptr) { return; } + auto const bitset_offset = + static_cast(bloom_range.offset()) + static_cast(spec.header_size); + if (spec.covered) { + // Whole bitset already read to host: copy the header-stripped bitset to device + bloom_filter_buffers.emplace_back( + spec.host_buffer->data() + spec.header_size, spec.bitset_size, stream, aligned_mr); + } else if (datasources[source_idx].get().is_device_read_preferred(spec.bitset_size)) { + // device-capable source: stream the bitset to device (kvikio picks GDS vs host bounce) + bloom_filter_buffers.emplace_back(spec.bitset_size, stream, aligned_mr); + device_read_tasks.emplace_back(datasources[source_idx].get().device_read_async( + bitset_offset, + spec.bitset_size, + static_cast(bloom_filter_buffers.back().data()), + stream)); + } else { + // host-only source (e.g. host buffer): read the bitset to host, then copy to device + auto const bitset_buffer = + datasources[source_idx].get().host_read(bitset_offset, spec.bitset_size); + bloom_filter_buffers.emplace_back( + bitset_buffer->data(), spec.bitset_size, stream, aligned_mr); + } + bitset_spans_per_source[source_idx][inner_idx] = cudf::device_span( + static_cast(bloom_filter_buffers.back().data()), spec.bitset_size); + }); + + auto sync_function = [](decltype(device_read_tasks) tasks) { + for (auto& task : tasks) { + task.get(); + } + }; + return {std::move(bloom_filter_buffers), + std::move(bitset_spans_per_source), + std::async(std::launch::deferred, sync_function, std::move(device_read_tasks))}; } } // namespace From 442cb67fd82852e1219527a13cc2db2f9e8fbacd Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Wed, 1 Jul 2026 20:00:54 +0200 Subject: [PATCH 16/39] Refine bloom filter fetch logic in Parquet I/O Updated the `fetch_bloom_filters_to_device_async_impl` function to enhance the speculative read approach for bloom filters. The logic now reads each filter in one host read based on its serialized length, improving efficiency by eliminating unnecessary header reads for known lengths. This change maintains the existing public interface and improves performance in high-latency storage scenarios. --- cpp/src/io/parquet/io_utils/parquet_io_utils.cpp | 14 ++++---------- 1 file changed, 4 insertions(+), 10 deletions(-) diff --git a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp index 48d2427513c8..ae40f317d427 100644 --- a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp +++ b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp @@ -394,11 +394,8 @@ fetch_bloom_filters_to_device_async_impl( CUDF_EXPECTS(num_sources == bloom_filter_byte_ranges_per_source.size(), "Encountered mismatch in number of datasources and bloom filter byte range spans"); - // Read a filter whole in one host read (up to the reader's speculative read size) and copy its - // bitset to device; larger filters read only the header and stream the bitset to device. - auto const speculative_read_size = - std::max(cudf::io::parquet::metadata_size_hint(), - static_cast(detail::bloom_filter_header_max_size)); + // Speculatively read each filter in one host read (its serialized length when known, else a + // header-sized guess), then copy the bitset to device if covered, otherwise stream the bitset. // Flatten to one (source index, index within source, bloom filter byte range) task per filter std::vector> @@ -432,11 +429,8 @@ fetch_bloom_filters_to_device_async_impl( auto const& [source_idx, inner_idx, bloom_range] = bloom_filter_tasks[task_idx]; // placeholder for a chunk whose column has no bloom filter written if (bloom_range.is_empty()) { return {}; } - auto const total_size = static_cast(bloom_range.size()); - // Whole filter in one read when its known length is within the cap; else just the header. - auto const read_size = total_size <= speculative_read_size - ? total_size - : static_cast(detail::bloom_filter_header_max_size); + // Read the whole filter when its length is known, else guess + auto const read_size = static_cast(bloom_range.size()); auto host_buffer = datasources[source_idx].get().host_read( static_cast(bloom_range.offset()), read_size); auto const header_info = From 8981d57841c4d48f302581824820f5b0ed9740d7 Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Wed, 1 Jul 2026 23:40:54 +0200 Subject: [PATCH 17/39] Refactor bloom filter reading in Parquet I/O Simplified the `fetch_bloom_filters_to_device_async_impl` function by consolidating the speculative read logic into a single lambda function. This change enhances code clarity and maintains the existing functionality while improving the efficiency of bloom filter fetching. Additionally, removed outdated comments to streamline the codebase. --- cpp/src/io/parquet/bloom_filter_reader.cu | 3 +- .../io/parquet/io_utils/parquet_io_utils.cpp | 45 ++++++++++--------- 2 files changed, 24 insertions(+), 24 deletions(-) diff --git a/cpp/src/io/parquet/bloom_filter_reader.cu b/cpp/src/io/parquet/bloom_filter_reader.cu index 18c6f981799a..5cb11710b963 100644 --- a/cpp/src/io/parquet/bloom_filter_reader.cu +++ b/cpp/src/io/parquet/bloom_filter_reader.cu @@ -365,8 +365,7 @@ aggregate_reader_metadata::read_bloom_filters( // Exit early if we don't have any bloom filters if (not have_bloom_filters) { return {}; } - // Fetch the header-stripped, 32-byte-aligned bloom filter bitsets to device via the shared - // utility + // Fetch the header-stripped, 32-byte-aligned bloom filter bitsets to device std::vector> datasource_refs; datasource_refs.reserve(sources.size()); std::transform( diff --git a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp index ae40f317d427..4740187d8455 100644 --- a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp +++ b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp @@ -394,9 +394,6 @@ fetch_bloom_filters_to_device_async_impl( CUDF_EXPECTS(num_sources == bloom_filter_byte_ranges_per_source.size(), "Encountered mismatch in number of datasources and bloom filter byte range spans"); - // Speculatively read each filter in one host read (its serialized length when known, else a - // header-sized guess), then copy the bitset to device if covered, otherwise stream the bitset. - // Flatten to one (source index, index within source, bloom filter byte range) task per filter std::vector> bloom_filter_tasks; @@ -424,25 +421,29 @@ fetch_bloom_filters_to_device_async_impl( std::size_t bitset_size{}; bool covered{}; // the whole bitset is already present in `host_buffer` }; - auto speculative_reads = - dispatch_tasks(bloom_filter_tasks.size(), [&](std::size_t task_idx) -> speculative_read { - auto const& [source_idx, inner_idx, bloom_range] = bloom_filter_tasks[task_idx]; - // placeholder for a chunk whose column has no bloom filter written - if (bloom_range.is_empty()) { return {}; } - // Read the whole filter when its length is known, else guess - auto const read_size = static_cast(bloom_range.size()); - auto host_buffer = datasources[source_idx].get().host_read( - static_cast(bloom_range.offset()), read_size); - auto const header_info = - detail::parse_bloom_filter_header({host_buffer->data(), host_buffer->size()}); - CUDF_EXPECTS(header_info.has_value(), "Encountered an invalid bloom filter header"); - auto const [header_size, bitset_size] = header_info.value(); - auto const covered = read_size >= static_cast(header_size) + bitset_size; - return {std::move(host_buffer), header_size, bitset_size, covered}; - }); + // Speculatively read one bloom filter from a source + auto const read_speculative = + [](cudf::io::datasource& datasource, + cudf::io::text::byte_range_info const& bloom_range) -> speculative_read { + // placeholder for a chunk whose column has no bloom filter written + if (bloom_range.is_empty()) { return {}; } + auto const read_size = static_cast(bloom_range.size()); + auto host_buffer = + datasource.host_read(static_cast(bloom_range.offset()), read_size); + auto const header_info = + detail::parse_bloom_filter_header({host_buffer->data(), host_buffer->size()}); + CUDF_EXPECTS(header_info.has_value(), "Encountered an invalid bloom filter header"); + auto const [header_size, bitset_size] = header_info.value(); + auto const covered = read_size >= static_cast(header_size) + bitset_size; + return {std::move(host_buffer), header_size, bitset_size, covered}; + }; + + auto speculative_reads = dispatch_tasks(bloom_filter_tasks.size(), [&](std::size_t task_idx) { + auto const& [source_idx, inner_idx, bloom_range] = bloom_filter_tasks[task_idx]; + return read_speculative(datasources[source_idx].get(), bloom_range); + }); - // Phase 2: materialize aligned device bitsets — copy covered filters from the host read, stream - // the rest to device. + // Phase 2: materialize aligned device bitsets std::vector bloom_filter_buffers; bloom_filter_buffers.reserve(bloom_filter_tasks.size()); std::vector bitset_spans_per_source(num_sources); @@ -468,7 +469,7 @@ fetch_bloom_filters_to_device_async_impl( bloom_filter_buffers.emplace_back( spec.host_buffer->data() + spec.header_size, spec.bitset_size, stream, aligned_mr); } else if (datasources[source_idx].get().is_device_read_preferred(spec.bitset_size)) { - // device-capable source: stream the bitset to device (kvikio picks GDS vs host bounce) + // device-capable source: stream the bitset to device bloom_filter_buffers.emplace_back(spec.bitset_size, stream, aligned_mr); device_read_tasks.emplace_back(datasources[source_idx].get().device_read_async( bitset_offset, From 6e6443addc1421f4db515354bceff242edb56823 Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Fri, 3 Jul 2026 11:37:02 +0200 Subject: [PATCH 18/39] Refactor bloom filter fetching in Parquet I/O Updated the `fetch_bloom_filters_to_device_async` function to `fetch_bloom_filters_to_device`, changing the return type from a tuple to a pair for improved clarity. This change simplifies the function signature and enhances code readability. Additionally, adjusted the corresponding implementation in `bloom_filter_reader.cu` to align with the new function signature, maintaining existing functionality while streamlining the codebase. --- cpp/include/cudf/io/parquet_io_utils.hpp | 26 +-- cpp/src/io/parquet/bloom_filter_reader.cu | 6 +- .../io/parquet/io_utils/parquet_io_utils.cpp | 220 ++++++++++-------- 3 files changed, 140 insertions(+), 112 deletions(-) diff --git a/cpp/include/cudf/io/parquet_io_utils.hpp b/cpp/include/cudf/io/parquet_io_utils.hpp index df8288b56dae..e1c8d97881d9 100644 --- a/cpp/include/cudf/io/parquet_io_utils.hpp +++ b/cpp/include/cudf/io/parquet_io_utils.hpp @@ -17,6 +17,7 @@ #include #include #include +#include #include /** @@ -166,16 +167,13 @@ fetch_byte_ranges_to_device_async( * @param stream CUDA stream * @param aligned_mr Device memory resource to allocate aligned memory for bloom filters * - * @return A tuple containing the device buffers, the device spans of the bitset data, and a future - * to wait on the read tasks + * @return A pair containing the device buffers and the device spans of the bitset data */ -std::tuple, - std::vector>, - std::future> -fetch_bloom_filters_to_device_async(cudf::io::datasource& datasource, - cudf::host_span bloom_filter_byte_ranges, - rmm::cuda_stream_view stream, - rmm::device_async_resource_ref aligned_mr); +std::pair, std::vector>> +fetch_bloom_filters_to_device(cudf::io::datasource& datasource, + cudf::host_span bloom_filter_byte_ranges, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref aligned_mr); /** * @brief Fetches Parquet bloom filter bitsets from multiple datasources into device buffers @@ -191,13 +189,11 @@ fetch_bloom_filters_to_device_async(cudf::io::datasource& datasource, * @param stream CUDA stream * @param aligned_mr Device memory resource to allocate aligned memory for bloom filters * - * @return A tuple containing a vector of device buffers, a vector of vectors of device spans, and a - * future to wait on the read tasks + * @return A pair containing a vector of device buffers and a vector of vectors of device spans */ -std::tuple, - std::vector>>, - std::future> -fetch_bloom_filters_to_device_async( +std::pair, + std::vector>>> +fetch_bloom_filters_to_device( cudf::host_span const> datasources, cudf::host_span const> bloom_filter_byte_ranges_per_source, rmm::cuda_stream_view stream, diff --git a/cpp/src/io/parquet/bloom_filter_reader.cu b/cpp/src/io/parquet/bloom_filter_reader.cu index 5cb11710b963..94310c26e534 100644 --- a/cpp/src/io/parquet/bloom_filter_reader.cu +++ b/cpp/src/io/parquet/bloom_filter_reader.cu @@ -373,10 +373,8 @@ aggregate_reader_metadata::read_bloom_filters( return std::ref(*source); }); - auto [bloom_filter_buffers, bitset_spans_per_source, fetch_task] = - fetch_bloom_filters_to_device_async( - datasource_refs, bloom_filter_byte_ranges_per_source, stream, aligned_mr); - fetch_task.get(); + auto [bloom_filter_buffers, bitset_spans_per_source] = fetch_bloom_filters_to_device( + datasource_refs, bloom_filter_byte_ranges_per_source, stream, aligned_mr); // Flatten the per-source bitset spans into per-chunk order std::vector> bloom_filter_data; diff --git a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp index 4740187d8455..68544b748f4d 100644 --- a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp +++ b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp @@ -18,6 +18,7 @@ #include #include +#include #include #include @@ -31,6 +32,7 @@ #include #include #include +#include #include #include #include @@ -380,10 +382,8 @@ fetch_byte_ranges_to_device_async_impl( std::async(std::launch::deferred, sync_function, std::move(device_read_tasks))}; } -std::tuple, - std::vector, - std::future> -fetch_bloom_filters_to_device_async_impl( +std::pair, std::vector> +fetch_bloom_filters_to_device_impl( cudf::host_span const> datasources, cudf::host_span const> bloom_filter_byte_ranges_per_source, @@ -394,107 +394,144 @@ fetch_bloom_filters_to_device_async_impl( CUDF_EXPECTS(num_sources == bloom_filter_byte_ranges_per_source.size(), "Encountered mismatch in number of datasources and bloom filter byte range spans"); - // Flatten to one (source index, index within source, bloom filter byte range) task per filter - std::vector> - bloom_filter_tasks; + // One task per bloom filter, flattened across sources + struct bloom_filter_task { + std::size_t source_idx; + std::size_t inner_idx; + cudf::io::text::byte_range_info bloom_range; + std::size_t host_offset{}; + std::size_t device_offset{}; + std::size_t header_size{}; + std::size_t bitset_size{}; + bool covered{}; // the whole bitset is already present in the source's host buffer + bool present{}; // false for a chunk whose column has no bloom filter + }; + + // Phase 1: read bloom filters to host and parse their headers + std::vector bloom_filter_tasks; bloom_filter_tasks.reserve( std::accumulate(bloom_filter_byte_ranges_per_source.begin(), bloom_filter_byte_ranges_per_source.end(), std::size_t{0}, [](auto acc, auto const& bloom_ranges) { return acc + bloom_ranges.size(); })); + + // Flatten per-source ranges while assigning each filter a task slot + std::vector per_source_host_size(num_sources, 0); std::for_each(cuda::counting_iterator(0), cuda::counting_iterator(num_sources), - [&](std::size_t outer_idx) { - auto const& bloom_ranges = bloom_filter_byte_ranges_per_source[outer_idx]; - std::transform(cuda::counting_iterator(0), - cuda::counting_iterator(bloom_ranges.size()), - std::back_inserter(bloom_filter_tasks), - [&](std::size_t inner_idx) { - return std::tuple{outer_idx, inner_idx, bloom_ranges[inner_idx]}; - }); + [&](std::size_t source_idx) { + auto const& bloom_ranges = bloom_filter_byte_ranges_per_source[source_idx]; + std::for_each(cuda::counting_iterator(0), + cuda::counting_iterator(bloom_ranges.size()), + [&](std::size_t inner_idx) { + auto const& bloom_range = bloom_ranges[inner_idx]; + auto const host_offset = per_source_host_size[source_idx]; + auto const read_size = + static_cast(bloom_range.size()); + bloom_filter_tasks.push_back( + {source_idx, inner_idx, bloom_range, host_offset}); + per_source_host_size[source_idx] = host_offset + read_size; + }); }); - // Phase 1: dispatch one speculative host read + header parse per bloom filter (source -> host) - struct speculative_read { - std::unique_ptr host_buffer; // speculative prefix, null if empty - int64_t header_size{}; - std::size_t bitset_size{}; - bool covered{}; // the whole bitset is already present in `host_buffer` - }; - // Speculatively read one bloom filter from a source - auto const read_speculative = - [](cudf::io::datasource& datasource, - cudf::io::text::byte_range_info const& bloom_range) -> speculative_read { - // placeholder for a chunk whose column has no bloom filter written - if (bloom_range.is_empty()) { return {}; } - auto const read_size = static_cast(bloom_range.size()); - auto host_buffer = - datasource.host_read(static_cast(bloom_range.offset()), read_size); - auto const header_info = - detail::parse_bloom_filter_header({host_buffer->data(), host_buffer->size()}); + // Allocate one host buffer per source + std::vector> host_buffers(num_sources); + std::transform(per_source_host_size.begin(), + per_source_host_size.end(), + host_buffers.begin(), + [](std::size_t size) { return std::vector(size); }); + + // Speculatively read to host buffer & parse header + auto const read_speculative = [&](std::size_t task_idx) -> bloom_filter_task { + auto task = bloom_filter_tasks[task_idx]; + if (task.bloom_range.is_empty()) { return task; } + auto const read_size = static_cast(task.bloom_range.size()); + auto const host_slot = + std::span{host_buffers[task.source_idx]}.subspan(task.host_offset, read_size); + auto const bytes_read = datasources[task.source_idx].get().host_read( + static_cast(task.bloom_range.offset()), read_size, host_slot.data()); + auto const header_info = detail::parse_bloom_filter_header({host_slot.data(), bytes_read}); CUDF_EXPECTS(header_info.has_value(), "Encountered an invalid bloom filter header"); - auto const [header_size, bitset_size] = header_info.value(); - auto const covered = read_size >= static_cast(header_size) + bitset_size; - return {std::move(host_buffer), header_size, bitset_size, covered}; + auto const [header_bytes, bitset_size] = header_info.value(); + task.header_size = static_cast(header_bytes); + task.bitset_size = bitset_size; + task.covered = bytes_read >= task.header_size + bitset_size; + task.present = true; + return task; }; + bloom_filter_tasks = dispatch_tasks(bloom_filter_tasks.size(), read_speculative); - auto speculative_reads = dispatch_tasks(bloom_filter_tasks.size(), [&](std::size_t task_idx) { - auto const& [source_idx, inner_idx, bloom_range] = bloom_filter_tasks[task_idx]; - return read_speculative(datasources[source_idx].get(), bloom_range); + // Phase 2: allocate one aligned device buffer per source + auto constexpr bitset_alignment = rmm::CUDA_ALLOCATION_ALIGNMENT; + + std::vector per_source_device_size(num_sources, 0); + std::for_each(bloom_filter_tasks.begin(), bloom_filter_tasks.end(), [&](bloom_filter_task& task) { + if (not task.present) { return; } + task.device_offset = per_source_device_size[task.source_idx]; + per_source_device_size[task.source_idx] += + cudf::util::round_up_safe(task.bitset_size, bitset_alignment); }); - // Phase 2: materialize aligned device bitsets std::vector bloom_filter_buffers; - bloom_filter_buffers.reserve(bloom_filter_tasks.size()); + bloom_filter_buffers.reserve(num_sources); + std::transform(per_source_device_size.begin(), + per_source_device_size.end(), + std::back_inserter(bloom_filter_buffers), + [&](std::size_t size) { return rmm::device_buffer(size, stream, aligned_mr); }); + + // Phase 3: make all bitsets host-resident, then batch-copy them to device + std::vector copy_dsts; + std::vector copy_srcs; + std::vector copy_sizes; + copy_dsts.reserve(bloom_filter_tasks.size()); + copy_srcs.reserve(bloom_filter_tasks.size()); + copy_sizes.reserve(bloom_filter_tasks.size()); + std::vector> bounce_buffers; + + // Allocate device span table std::vector bitset_spans_per_source(num_sources); std::transform( bloom_filter_byte_ranges_per_source.begin(), bloom_filter_byte_ranges_per_source.end(), bitset_spans_per_source.begin(), [](auto const& bloom_ranges) { return device_spans_per_source_type(bloom_ranges.size()); }); - std::vector> device_read_tasks; std::for_each( - cuda::counting_iterator(0), - cuda::counting_iterator(bloom_filter_tasks.size()), - [&](std::size_t task_idx) { - auto const& [source_idx, inner_idx, bloom_range] = bloom_filter_tasks[task_idx]; - auto& spec = speculative_reads[task_idx]; - // empty bloom filter -> leave an empty span in place - if (spec.host_buffer == nullptr) { return; } - auto const bitset_offset = - static_cast(bloom_range.offset()) + static_cast(spec.header_size); - if (spec.covered) { - // Whole bitset already read to host: copy the header-stripped bitset to device - bloom_filter_buffers.emplace_back( - spec.host_buffer->data() + spec.header_size, spec.bitset_size, stream, aligned_mr); - } else if (datasources[source_idx].get().is_device_read_preferred(spec.bitset_size)) { - // device-capable source: stream the bitset to device - bloom_filter_buffers.emplace_back(spec.bitset_size, stream, aligned_mr); - device_read_tasks.emplace_back(datasources[source_idx].get().device_read_async( - bitset_offset, - spec.bitset_size, - static_cast(bloom_filter_buffers.back().data()), - stream)); + bloom_filter_tasks.begin(), bloom_filter_tasks.end(), [&](bloom_filter_task const& task) { + // Empty bloom filter: leave an empty span in place + if (not task.present) { return; } + // Aligned destination slot within this source's device buffer. + auto& dst_buffer = bloom_filter_buffers[task.source_idx]; + auto const dst = + cudf::device_span{static_cast(dst_buffer.data()), dst_buffer.size()} + .subspan(task.device_offset, task.bitset_size); + CUDF_EXPECTS(reinterpret_cast(dst.data()) % bitset_alignment == 0, + "Encountered a misaligned bloom filter bitset"); + if (task.covered) { + // Bitset fully covered: direct copy to device + auto const src = std::span{host_buffers[task.source_idx]}.subspan( + task.host_offset + task.header_size, task.bitset_size); + copy_srcs.push_back(src.data()); } else { - // host-only source (e.g. host buffer): read the bitset to host, then copy to device - auto const bitset_buffer = - datasources[source_idx].get().host_read(bitset_offset, spec.bitset_size); - bloom_filter_buffers.emplace_back( - bitset_buffer->data(), spec.bitset_size, stream, aligned_mr); + // Bitset not fully covered: read from source, then copy to device + auto const bitset_offset = + static_cast(task.bloom_range.offset()) + task.header_size; + bounce_buffers.emplace_back( + datasources[task.source_idx].get().host_read(bitset_offset, task.bitset_size)); + copy_srcs.push_back(bounce_buffers.back()->data()); } - bitset_spans_per_source[source_idx][inner_idx] = cudf::device_span( - static_cast(bloom_filter_buffers.back().data()), spec.bitset_size); + copy_dsts.push_back(dst.data()); + copy_sizes.push_back(task.bitset_size); + bitset_spans_per_source[task.source_idx][task.inner_idx] = dst; }); - auto sync_function = [](decltype(device_read_tasks) tasks) { - for (auto& task : tasks) { - task.get(); - } - }; - return {std::move(bloom_filter_buffers), - std::move(bitset_spans_per_source), - std::async(std::launch::deferred, sync_function, std::move(device_read_tasks))}; + // One batched host-to-device copy for every bitset + if (not copy_dsts.empty()) { + CUDF_CUDA_TRY(cudf::detail::memcpy_batch_async( + copy_dsts.data(), copy_srcs.data(), copy_sizes.data(), copy_dsts.size(), stream)); + stream.synchronize(); + } + return {std::move(bloom_filter_buffers), std::move(bitset_spans_per_source)}; } } // namespace @@ -592,10 +629,8 @@ fetch_byte_ranges_to_device_async( mr); } -std::tuple, - std::vector>, - std::future> -fetch_bloom_filters_to_device_async( +std::pair, std::vector>> +fetch_bloom_filters_to_device( cudf::io::datasource& datasource, cudf::host_span bloom_filter_byte_ranges, rmm::cuda_stream_view stream, @@ -608,19 +643,18 @@ fetch_bloom_filters_to_device_async( std::array, 1> bloom_filter_byte_ranges_per_source{bloom_filter_byte_ranges}; - auto [buffers, fetched_byte_ranges, fut] = fetch_bloom_filters_to_device_async_impl( + auto [buffers, fetched_byte_ranges] = fetch_bloom_filters_to_device_impl( {datasources.data(), datasources.size()}, {bloom_filter_byte_ranges_per_source.data(), bloom_filter_byte_ranges_per_source.size()}, stream, aligned_mr); - return {std::move(buffers), std::move(fetched_byte_ranges.front()), std::move(fut)}; + return {std::move(buffers), std::move(fetched_byte_ranges.front())}; } -std::tuple, - std::vector>>, - std::future> -fetch_bloom_filters_to_device_async( +std::pair, + std::vector>>> +fetch_bloom_filters_to_device( cudf::host_span const> datasources, cudf::host_span const> bloom_filter_byte_ranges_per_source, @@ -636,11 +670,11 @@ fetch_bloom_filters_to_device_async( for (auto const& ranges : bloom_filter_byte_ranges_per_source) { bloom_filter_byte_range_spans_per_source.emplace_back(ranges); } - return fetch_bloom_filters_to_device_async_impl(datasources, - {bloom_filter_byte_range_spans_per_source.data(), - bloom_filter_byte_range_spans_per_source.size()}, - stream, - aligned_mr); + return fetch_bloom_filters_to_device_impl(datasources, + {bloom_filter_byte_range_spans_per_source.data(), + bloom_filter_byte_range_spans_per_source.size()}, + stream, + aligned_mr); } } // namespace cudf::io::parquet From b3cd38b3dc5d693e81cbdfd26a80505e0c6db302 Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Fri, 3 Jul 2026 16:12:11 +0200 Subject: [PATCH 19/39] Update task dispatch documentation in Parquet I/O Refined the comment for the task dispatch function to clarify its purpose. The updated description now emphasizes that it dispatches each indexed task and collects the results, improving code readability and understanding for future developers. --- cpp/src/io/parquet/io_utils/parquet_io_utils.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp index 68544b748f4d..3e74fed40c46 100644 --- a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp +++ b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp @@ -53,7 +53,7 @@ namespace cudf::io::parquet { namespace { /** - * @brief Dispatches a task for each task index and collects the results + * @brief Dispatches each indexed task and collects the results * * Dispatches sequentially or using host worker pool depending on the number of tasks. * From 3969eed2a1ed42b2d8826dce40e398a39d06fbd7 Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Mon, 6 Jul 2026 10:46:07 +0200 Subject: [PATCH 20/39] Enhance bloom filter reading logic in Parquet I/O Updated the speculative read size for bloom filters to improve efficiency. The maximum header size constant has been replaced with a new constant that allows for a more accurate speculative read, ensuring better performance when recovering bitset sizes. This change clarifies the intent of the code and aligns with recent refactoring efforts in bloom filter handling. --- cpp/src/io/parquet/bloom_filter_reader.cu | 4 ++-- cpp/src/io/parquet/reader_impl_helpers.hpp | 4 ++-- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/cpp/src/io/parquet/bloom_filter_reader.cu b/cpp/src/io/parquet/bloom_filter_reader.cu index 94310c26e534..3a7e56967f25 100644 --- a/cpp/src/io/parquet/bloom_filter_reader.cu +++ b/cpp/src/io/parquet/bloom_filter_reader.cu @@ -349,10 +349,10 @@ aggregate_reader_metadata::read_bloom_filters( auto const& col_meta = get_column_metadata(rg_index, src_index, schema_idx); if (col_meta.bloom_filter_offset.has_value()) { have_bloom_filters = true; - // When the length is absent, read up to the max header size to recover the bitset size. + // Length absent: read a speculative chunk to recover the bitset size. auto const length = col_meta.bloom_filter_length.has_value() ? static_cast(col_meta.bloom_filter_length.value()) - : bloom_filter_header_max_size; + : bloom_filter_speculative_read_size; source_ranges.push_back( cudf::io::text::byte_range_info{col_meta.bloom_filter_offset.value(), length}); } else { diff --git a/cpp/src/io/parquet/reader_impl_helpers.hpp b/cpp/src/io/parquet/reader_impl_helpers.hpp index 78378d47f6a3..ac48140e2eb9 100644 --- a/cpp/src/io/parquet/reader_impl_helpers.hpp +++ b/cpp/src/io/parquet/reader_impl_helpers.hpp @@ -145,9 +145,9 @@ struct surviving_row_group_metrics { }; /** - * @brief Upper bound on the size in bytes of a Parquet `BloomFilterHeader` + * @brief Bytes enough to recover the header (and often the whole bitset) in one speculative read. */ -inline constexpr int64_t bloom_filter_header_max_size = 256; +inline constexpr int64_t bloom_filter_speculative_read_size = 512; /** * @brief Parses and validates a Parquet `BloomFilterHeader` from the front of `bytes` From 3d438af0e6bb8a92b928ce3f3909e48ca4e2de19 Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Mon, 6 Jul 2026 12:37:19 +0200 Subject: [PATCH 21/39] Refactor bloom filter speculative read size in Parquet I/O Updated the speculative read size constant for bloom filters in `bloom_filter_reader.cu` to improve clarity and maintainability. The change replaces the previous constant with a more descriptive name, enhancing code readability and aligning with recent refactoring efforts in bloom filter handling. --- cpp/src/io/parquet/bloom_filter_reader.cu | 5 ++++- cpp/src/io/parquet/reader_impl_helpers.hpp | 5 ----- 2 files changed, 4 insertions(+), 6 deletions(-) diff --git a/cpp/src/io/parquet/bloom_filter_reader.cu b/cpp/src/io/parquet/bloom_filter_reader.cu index 3a7e56967f25..2982aef27aa3 100644 --- a/cpp/src/io/parquet/bloom_filter_reader.cu +++ b/cpp/src/io/parquet/bloom_filter_reader.cu @@ -331,6 +331,9 @@ aggregate_reader_metadata::read_bloom_filters( // Flag to check if we have at least one valid bloom filter offset auto have_bloom_filters = false; + // Speculatively read when a bloom filter's length is absent, enough to cover the header (and + // often the whole bitset). + auto constexpr speculative_read_size = int64_t{512}; // Build complete bloom filter byte ranges (header + bitset) for every column chunk std::vector> bloom_filter_byte_ranges_per_source( row_group_indices.size()); @@ -352,7 +355,7 @@ aggregate_reader_metadata::read_bloom_filters( // Length absent: read a speculative chunk to recover the bitset size. auto const length = col_meta.bloom_filter_length.has_value() ? static_cast(col_meta.bloom_filter_length.value()) - : bloom_filter_speculative_read_size; + : speculative_read_size; source_ranges.push_back( cudf::io::text::byte_range_info{col_meta.bloom_filter_offset.value(), length}); } else { diff --git a/cpp/src/io/parquet/reader_impl_helpers.hpp b/cpp/src/io/parquet/reader_impl_helpers.hpp index ac48140e2eb9..2595e933f9c8 100644 --- a/cpp/src/io/parquet/reader_impl_helpers.hpp +++ b/cpp/src/io/parquet/reader_impl_helpers.hpp @@ -144,11 +144,6 @@ struct surviving_row_group_metrics { std::optional after_bloom_filter; // number of surviving row groups after bloom filter }; -/** - * @brief Bytes enough to recover the header (and often the whole bitset) in one speculative read. - */ -inline constexpr int64_t bloom_filter_speculative_read_size = 512; - /** * @brief Parses and validates a Parquet `BloomFilterHeader` from the front of `bytes` * From b82a758255145ea32f9fd336677a626c8d586b3b Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Mon, 6 Jul 2026 17:21:40 +0200 Subject: [PATCH 22/39] Enhance bloom filter validation in Parquet I/O Added checks to ensure that the bloom filter bitset size is a multiple of 32 bytes and that the complete bloom filter bitset is read successfully. These enhancements improve error handling and maintain the integrity of bloom filter operations, aligning with recent refactoring efforts in the Parquet I/O module. --- cpp/src/io/parquet/io_utils/parquet_io_utils.cpp | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp index 3e74fed40c46..9700848dc8c5 100644 --- a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp +++ b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp @@ -457,6 +457,10 @@ fetch_bloom_filters_to_device_impl( task.bitset_size = bitset_size; task.covered = bytes_read >= task.header_size + bitset_size; task.present = true; + // Arrow bloom filter bitsets should be a whole number of 32-byte blocks + auto constexpr bloom_filter_block_bytes = std::size_t{32}; + CUDF_EXPECTS(bitset_size % bloom_filter_block_bytes == 0, + "Bloom filter bitset size must be a multiple of 32 bytes"); return task; }; bloom_filter_tasks = dispatch_tasks(bloom_filter_tasks.size(), read_speculative); @@ -518,6 +522,8 @@ fetch_bloom_filters_to_device_impl( static_cast(task.bloom_range.offset()) + task.header_size; bounce_buffers.emplace_back( datasources[task.source_idx].get().host_read(bitset_offset, task.bitset_size)); + CUDF_EXPECTS(bounce_buffers.back()->size() == task.bitset_size, + "Failed to read complete bloom filter bitset"); copy_srcs.push_back(bounce_buffers.back()->data()); } copy_dsts.push_back(dst.data()); From be4f00f9411cc36ca308e3cf40cc2ae670c3ba1e Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Tue, 7 Jul 2026 17:32:25 +0200 Subject: [PATCH 23/39] Enhance bloom filter reading and host data fetching in Parquet I/O Added a new helper function `read_ranges_to_host` to streamline the process of reading data ranges to host memory while ensuring thread safety. Updated the bloom filter reading logic to include checks for valid offsets and sizes, improving error handling. This change enhances the efficiency and clarity of the bloom filter fetching process, aligning with recent refactoring efforts in the Parquet I/O module. --- cpp/src/io/parquet/bloom_filter_reader.cu | 11 +- .../io/parquet/io_utils/parquet_io_utils.cpp | 301 +++++++++--------- 2 files changed, 162 insertions(+), 150 deletions(-) diff --git a/cpp/src/io/parquet/bloom_filter_reader.cu b/cpp/src/io/parquet/bloom_filter_reader.cu index 2982aef27aa3..da409a692075 100644 --- a/cpp/src/io/parquet/bloom_filter_reader.cu +++ b/cpp/src/io/parquet/bloom_filter_reader.cu @@ -344,6 +344,7 @@ aggregate_reader_metadata::read_bloom_filters( [&](auto const src_index) { auto const& rg_indices = row_group_indices[src_index]; auto& source_ranges = bloom_filter_byte_ranges_per_source[src_index]; + auto const source_size = static_cast(sources[src_index]->size()); source_ranges.reserve(rg_indices.size() * num_input_columns); // For all row groups in the source std::for_each(rg_indices.cbegin(), rg_indices.cend(), [&](auto const rg_index) { @@ -352,12 +353,14 @@ aggregate_reader_metadata::read_bloom_filters( auto const& col_meta = get_column_metadata(rg_index, src_index, schema_idx); if (col_meta.bloom_filter_offset.has_value()) { have_bloom_filters = true; - // Length absent: read a speculative chunk to recover the bitset size. + auto const offset = col_meta.bloom_filter_offset.value(); + CUDF_EXPECTS(offset >= 0 and offset < source_size, + "Bloom filter offset is out of datasource bounds"); + // Length absent: speculatively read enough to recover the header, clamped at EOF auto const length = col_meta.bloom_filter_length.has_value() ? static_cast(col_meta.bloom_filter_length.value()) - : speculative_read_size; - source_ranges.push_back( - cudf::io::text::byte_range_info{col_meta.bloom_filter_offset.value(), length}); + : std::min(speculative_read_size, source_size - offset); + source_ranges.push_back(cudf::io::text::byte_range_info{offset, length}); } else { source_ranges.push_back(cudf::io::text::byte_range_info{0, 0}); } diff --git a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp index 9700848dc8c5..f9e4df49df56 100644 --- a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp +++ b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp @@ -190,6 +190,44 @@ std::vector> fetch_page_indexes_to using device_spans_per_source_type = std::vector>; +// Reads the given (source, offset, size) ranges to host, holding a shared mutex while scheduling so +// that one caller thread's reads are dispatched contiguously without interleaving with reads from +// other threads (better pipelining, avoids cross-thread mis-sync). Returns one host buffer per +// range, in input order. Shared by the byte-range and bloom-filter device fetch paths. +std::vector> read_ranges_to_host( + cudf::host_span const> datasources, + cudf::host_span source_indices, + cudf::host_span offsets, + cudf::host_span sizes) +{ + static std::mutex host_read_mutex; + using host_read_buffer = std::unique_ptr; + + std::vector> host_read_tasks; + host_read_tasks.reserve(source_indices.size()); + { + std::scoped_lock lock(host_read_mutex); + auto iter = cuda::make_zip_iterator(source_indices.begin(), offsets.begin(), sizes.begin()); + std::for_each(iter, iter + source_indices.size(), [&](auto const& tuple) { + auto& datasource = datasources[cuda::std::get<0>(tuple)].get(); + auto const io_offset = cuda::std::get<1>(tuple); + auto const io_size = cuda::std::get<2>(tuple); + host_read_tasks.emplace_back(cudf::detail::host_worker_pool().submit_task( + [&datasource, io_offset, io_size]() -> host_read_buffer { + return datasource.host_read(io_offset, io_size); + })); + }); + } + + std::vector host_buffers; + host_buffers.reserve(host_read_tasks.size()); + std::transform(host_read_tasks.begin(), + host_read_tasks.end(), + std::back_inserter(host_buffers), + [](auto& task) { return task.get(); }); + return host_buffers; +} + std::tuple, std::vector, std::future> @@ -200,7 +238,6 @@ fetch_byte_ranges_to_device_async_impl( rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) { - static std::mutex host_read_mutex; static std::mutex device_read_mutex; auto const num_sources = datasources.size(); @@ -290,9 +327,7 @@ fetch_byte_ranges_to_device_async_impl( // Vectors to hold futures from datasource std::vector> device_read_tasks{}; - std::vector> host_read_tasks{}; device_read_tasks.reserve(io_offsets.size()); - host_read_tasks.reserve(io_offsets.size()); // Vectors to store intermediate host buffers and relevant pointers std::vector host_buffers{}; @@ -305,39 +340,30 @@ fetch_byte_ranges_to_device_async_impl( auto iter = cuda::make_zip_iterator( io_source_indices.begin(), io_offsets.begin(), io_sizes.begin(), destinations.begin()); - // Schedule host reads holding the `host_read_mutex` so that all reads for a caller thread - // are scheduled without interleaving with reads from other threads yielding better pipelining - { - std::scoped_lock lock(host_read_mutex); - - std::for_each(iter, iter + io_offsets.size(), [&](auto const& tuple) { - auto const src_idx = cuda::std::get<0>(tuple); - auto const io_offset = cuda::std::get<1>(tuple); - auto const io_size = cuda::std::get<2>(tuple); - auto const dest = cuda::std::get<3>(tuple); - - auto& datasource = datasources[src_idx].get(); - if (not datasource.is_device_read_preferred(io_size)) { - // Asynchronously read column chunk data to a host buffer - host_read_tasks.emplace_back(cudf::detail::host_worker_pool().submit_task( - [&datasource, io_offset, io_size]() -> host_read_buffer { - return datasource.host_read(io_offset, io_size); - })); - copy_dsts.push_back(static_cast(dest)); - copy_sizes.push_back(io_size); - } - }); - } - - // Complete host reads - if (not host_read_tasks.empty()) { - copy_srcs.reserve(host_read_tasks.size()); - host_buffers.reserve(host_read_tasks.size()); - - for (auto& task : host_read_tasks) { - host_buffers.emplace_back(task.get()); - copy_srcs.push_back(host_buffers.back().get()->data()); + // Collect the host-preferred ranges and read them to host through the shared helper + std::vector host_source_indices; + std::vector host_offsets; + std::vector host_sizes; + std::for_each(iter, iter + io_offsets.size(), [&](auto const& tuple) { + auto const src_idx = cuda::std::get<0>(tuple); + auto const io_offset = cuda::std::get<1>(tuple); + auto const io_size = cuda::std::get<2>(tuple); + auto const dest = cuda::std::get<3>(tuple); + + auto& datasource = datasources[src_idx].get(); + if (not datasource.is_device_read_preferred(io_size)) { + host_source_indices.push_back(src_idx); + host_offsets.push_back(io_offset); + host_sizes.push_back(io_size); + copy_dsts.push_back(static_cast(dest)); + copy_sizes.push_back(io_size); } + }); + + host_buffers = read_ranges_to_host(datasources, host_source_indices, host_offsets, host_sizes); + copy_srcs.reserve(host_buffers.size()); + for (auto const& buffer : host_buffers) { + copy_srcs.push_back(buffer->data()); } // `device_read_async` is not guaranteed to follow stream-ordering (see datasource API docs) @@ -394,88 +420,79 @@ fetch_bloom_filters_to_device_impl( CUDF_EXPECTS(num_sources == bloom_filter_byte_ranges_per_source.size(), "Encountered mismatch in number of datasources and bloom filter byte range spans"); - // One task per bloom filter, flattened across sources - struct bloom_filter_task { - std::size_t source_idx; - std::size_t inner_idx; - cudf::io::text::byte_range_info bloom_range; - std::size_t host_offset{}; - std::size_t device_offset{}; - std::size_t header_size{}; - std::size_t bitset_size{}; - bool covered{}; // the whole bitset is already present in the source's host buffer - bool present{}; // false for a chunk whose column has no bloom filter - }; + // Each bitset must align to 32-byte boundaries, as required by cuco's Arrow bloom filter bitsets. + auto constexpr bloom_filter_block_bytes = std::size_t{32}; - // Phase 1: read bloom filters to host and parse their headers - std::vector bloom_filter_tasks; - bloom_filter_tasks.reserve( + auto const total_filters = std::accumulate(bloom_filter_byte_ranges_per_source.begin(), bloom_filter_byte_ranges_per_source.end(), std::size_t{0}, - [](auto acc, auto const& bloom_ranges) { return acc + bloom_ranges.size(); })); - - // Flatten per-source ranges while assigning each filter a task slot - std::vector per_source_host_size(num_sources, 0); + [](auto acc, auto const& bloom_ranges) { return acc + bloom_ranges.size(); }); + + // Phase 1: speculatively read every filter to host. Absent filters keep a zero-length slot so the + // read buffers stay positionally aligned with the flattened filters + std::vector spec_source_indices; + std::vector spec_offsets; + std::vector spec_sizes; + spec_source_indices.reserve(total_filters); + spec_offsets.reserve(total_filters); + spec_sizes.reserve(total_filters); std::for_each(cuda::counting_iterator(0), cuda::counting_iterator(num_sources), [&](std::size_t source_idx) { auto const& bloom_ranges = bloom_filter_byte_ranges_per_source[source_idx]; - std::for_each(cuda::counting_iterator(0), - cuda::counting_iterator(bloom_ranges.size()), - [&](std::size_t inner_idx) { - auto const& bloom_range = bloom_ranges[inner_idx]; - auto const host_offset = per_source_host_size[source_idx]; - auto const read_size = - static_cast(bloom_range.size()); - bloom_filter_tasks.push_back( - {source_idx, inner_idx, bloom_range, host_offset}); - per_source_host_size[source_idx] = host_offset + read_size; - }); + std::for_each( + bloom_ranges.begin(), bloom_ranges.end(), [&](auto const& bloom_range) { + spec_source_indices.push_back(source_idx); + spec_offsets.push_back(static_cast(bloom_range.offset())); + spec_sizes.push_back(static_cast(bloom_range.size())); + }); }); + auto const spec_buffers = + read_ranges_to_host(datasources, spec_source_indices, spec_offsets, spec_sizes); - // Allocate one host buffer per source - std::vector> host_buffers(num_sources); - std::transform(per_source_host_size.begin(), - per_source_host_size.end(), - host_buffers.begin(), - [](std::size_t size) { return std::vector(size); }); - - // Speculatively read to host buffer & parse header - auto const read_speculative = [&](std::size_t task_idx) -> bloom_filter_task { - auto task = bloom_filter_tasks[task_idx]; - if (task.bloom_range.is_empty()) { return task; } - auto const read_size = static_cast(task.bloom_range.size()); - auto const host_slot = - std::span{host_buffers[task.source_idx]}.subspan(task.host_offset, read_size); - auto const bytes_read = datasources[task.source_idx].get().host_read( - static_cast(task.bloom_range.offset()), read_size, host_slot.data()); - auto const header_info = detail::parse_bloom_filter_header({host_slot.data(), bytes_read}); - CUDF_EXPECTS(header_info.has_value(), "Encountered an invalid bloom filter header"); - auto const [header_bytes, bitset_size] = header_info.value(); - task.header_size = static_cast(header_bytes); - task.bitset_size = bitset_size; - task.covered = bytes_read >= task.header_size + bitset_size; - task.present = true; - // Arrow bloom filter bitsets should be a whole number of 32-byte blocks - auto constexpr bloom_filter_block_bytes = std::size_t{32}; - CUDF_EXPECTS(bitset_size % bloom_filter_block_bytes == 0, - "Bloom filter bitset size must be a multiple of 32 bytes"); - return task; - }; - bloom_filter_tasks = dispatch_tasks(bloom_filter_tasks.size(), read_speculative); - - // Phase 2: allocate one aligned device buffer per source - auto constexpr bitset_alignment = rmm::CUDA_ALLOCATION_ALIGNMENT; + // Phase 2: parse each present filter's header for its sizes, total the aligned device bytes per + // source, and queue a follow-up read for any bitset the speculative read did not fully cover. + std::vector bloom_header_sizes(total_filters, 0); + std::vector bloom_bitset_sizes(total_filters, 0); + std::vector bloom_total_sizes(total_filters, 0); std::vector per_source_device_size(num_sources, 0); - std::for_each(bloom_filter_tasks.begin(), bloom_filter_tasks.end(), [&](bloom_filter_task& task) { - if (not task.present) { return; } - task.device_offset = per_source_device_size[task.source_idx]; - per_source_device_size[task.source_idx] += - cudf::util::round_up_safe(task.bitset_size, bitset_alignment); - }); + std::vector followup_source_indices; + std::vector followup_offsets; + std::vector followup_sizes; + std::for_each( + cuda::counting_iterator(0), + cuda::counting_iterator(total_filters), + [&](std::size_t filter_idx) { + // Absent chunk: no bloom filter to read + if (spec_sizes[filter_idx] == 0) { return; } + + auto const bytes_read = spec_buffers[filter_idx]->size(); + auto const header_info = + detail::parse_bloom_filter_header({spec_buffers[filter_idx]->data(), bytes_read}); + CUDF_EXPECTS(header_info.has_value(), "Encountered an invalid bloom filter header"); + auto const [header_bytes, bitset_bytes] = header_info.value(); + CUDF_EXPECTS(bitset_bytes % bloom_filter_block_bytes == 0, + "Bloom filter bitset size must be a multiple of 32 bytes"); + bloom_header_sizes[filter_idx] = static_cast(header_bytes); + bloom_bitset_sizes[filter_idx] = static_cast(bitset_bytes); + bloom_total_sizes[filter_idx] = + bloom_header_sizes[filter_idx] + bloom_bitset_sizes[filter_idx]; + per_source_device_size[spec_source_indices[filter_idx]] += + cudf::util::round_up_safe(bloom_bitset_sizes[filter_idx], bloom_filter_block_bytes); + + // Speculative read did not reach the whole bitset: queue a follow-up read for it + if (bytes_read < bloom_total_sizes[filter_idx]) { + followup_source_indices.push_back(spec_source_indices[filter_idx]); + followup_offsets.push_back(spec_offsets[filter_idx] + bloom_header_sizes[filter_idx]); + followup_sizes.push_back(bloom_bitset_sizes[filter_idx]); + } + }); + auto const followup_buffers = + read_ranges_to_host(datasources, followup_source_indices, followup_offsets, followup_sizes); + // Phase 3: one aligned device buffer per source std::vector bloom_filter_buffers; bloom_filter_buffers.reserve(num_sources); std::transform(per_source_device_size.begin(), @@ -483,55 +500,47 @@ fetch_bloom_filters_to_device_impl( std::back_inserter(bloom_filter_buffers), [&](std::size_t size) { return rmm::device_buffer(size, stream, aligned_mr); }); - // Phase 3: make all bitsets host-resident, then batch-copy them to device + // Phase 4: copy every bitset straight to its device slot in one batched transfer + std::vector bitset_spans_per_source(num_sources); + std::vector copy_dsts; std::vector copy_srcs; std::vector copy_sizes; - copy_dsts.reserve(bloom_filter_tasks.size()); - copy_srcs.reserve(bloom_filter_tasks.size()); - copy_sizes.reserve(bloom_filter_tasks.size()); - std::vector> bounce_buffers; - - // Allocate device span table - std::vector bitset_spans_per_source(num_sources); - std::transform( - bloom_filter_byte_ranges_per_source.begin(), - bloom_filter_byte_ranges_per_source.end(), - bitset_spans_per_source.begin(), - [](auto const& bloom_ranges) { return device_spans_per_source_type(bloom_ranges.size()); }); - - std::for_each( - bloom_filter_tasks.begin(), bloom_filter_tasks.end(), [&](bloom_filter_task const& task) { - // Empty bloom filter: leave an empty span in place - if (not task.present) { return; } - // Aligned destination slot within this source's device buffer. - auto& dst_buffer = bloom_filter_buffers[task.source_idx]; - auto const dst = - cudf::device_span{static_cast(dst_buffer.data()), dst_buffer.size()} - .subspan(task.device_offset, task.bitset_size); - CUDF_EXPECTS(reinterpret_cast(dst.data()) % bitset_alignment == 0, + copy_dsts.reserve(total_filters); + copy_srcs.reserve(total_filters); + copy_sizes.reserve(total_filters); + std::size_t filter_idx = 0; + std::size_t followup_index = 0; + for (std::size_t source_idx = 0; source_idx < num_sources; ++source_idx) { + auto const& bloom_ranges = bloom_filter_byte_ranges_per_source[source_idx]; + bitset_spans_per_source[source_idx] = device_spans_per_source_type(bloom_ranges.size()); + auto* const device_base = static_cast(bloom_filter_buffers[source_idx].data()); + std::size_t device_offset = 0; + for (std::size_t inner_idx = 0; inner_idx < bloom_ranges.size(); ++inner_idx, ++filter_idx) { + auto const bitset_size = bloom_bitset_sizes[filter_idx]; + // Absent filter: leave an empty span + if (bitset_size == 0) { continue; } + auto* const device_dst = device_base + device_offset; + CUDF_EXPECTS(reinterpret_cast(device_dst) % bloom_filter_block_bytes == 0, "Encountered a misaligned bloom filter bitset"); - if (task.covered) { - // Bitset fully covered: direct copy to device - auto const src = std::span{host_buffers[task.source_idx]}.subspan( - task.host_offset + task.header_size, task.bitset_size); - copy_srcs.push_back(src.data()); + bitset_spans_per_source[source_idx][inner_idx] = + cudf::device_span{device_dst, bitset_size}; + if (spec_buffers[filter_idx]->size() >= bloom_total_sizes[filter_idx]) { + // Whole bitset already read speculatively: copy it directly, stripping the header + copy_srcs.push_back(spec_buffers[filter_idx]->data() + bloom_header_sizes[filter_idx]); } else { - // Bitset not fully covered: read from source, then copy to device - auto const bitset_offset = - static_cast(task.bloom_range.offset()) + task.header_size; - bounce_buffers.emplace_back( - datasources[task.source_idx].get().host_read(bitset_offset, task.bitset_size)); - CUDF_EXPECTS(bounce_buffers.back()->size() == task.bitset_size, - "Failed to read complete bloom filter bitset"); - copy_srcs.push_back(bounce_buffers.back()->data()); + // Bitset came from the follow-up read (queued in the same order) + auto const& buffer = followup_buffers[followup_index++]; + CUDF_EXPECTS(buffer->size() == bitset_size, "Failed to read complete bloom filter bitset"); + copy_srcs.push_back(buffer->data()); } - copy_dsts.push_back(dst.data()); - copy_sizes.push_back(task.bitset_size); - bitset_spans_per_source[task.source_idx][task.inner_idx] = dst; - }); + copy_dsts.push_back(device_dst); + copy_sizes.push_back(bitset_size); + device_offset += cudf::util::round_up_safe(bitset_size, bloom_filter_block_bytes); + } + } - // One batched host-to-device copy for every bitset + // One batched host-to-device copy for every bitset (bloom filters never use direct device reads). if (not copy_dsts.empty()) { CUDF_CUDA_TRY(cudf::detail::memcpy_batch_async( copy_dsts.data(), copy_srcs.data(), copy_sizes.data(), copy_dsts.size(), stream)); From 177f978a1270fcc4395b5adaef13f404315054eb Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Wed, 8 Jul 2026 17:06:29 +0200 Subject: [PATCH 24/39] Share host/device fetch helpers between byte-range and bloom paths Extract read_ranges_to_host and copy_ranges_to_device so both fetch_byte_ranges_to_device_async_impl and fetch_bloom_filters_to_device_impl schedule their reads and copies through the same helpers. copy_ranges_to_device owns device_read_mutex and runs the direct device reads plus the batched H2D copy under it, so the bloom bitset copy is now guarded by the same mutex as the column-chunk fetch, preventing cross-thread read interleaving (mis-sync). Split the byte-range fetch into host- and device-preferred buckets up front (fuse_split), dropping the per-range branch in the copy path. No public API change and no device-memory delta. --- .../io/parquet/io_utils/parquet_io_utils.cpp | 297 +++++++++++------- 1 file changed, 187 insertions(+), 110 deletions(-) diff --git a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp index f9e4df49df56..8893de1919cf 100644 --- a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp +++ b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp @@ -17,6 +17,7 @@ #include #include #include +#include #include #include @@ -189,43 +190,117 @@ std::vector> fetch_page_indexes_to } using device_spans_per_source_type = std::vector>; +using host_read_buffer = std::unique_ptr; -// Reads the given (source, offset, size) ranges to host, holding a shared mutex while scheduling so -// that one caller thread's reads are dispatched contiguously without interleaving with reads from -// other threads (better pipelining, avoids cross-thread mis-sync). Returns one host buffer per -// range, in input order. Shared by the byte-range and bloom-filter device fetch paths. -std::vector> read_ranges_to_host( +/** + * @brief Schedules host reads for the given byte ranges. + * + * Holds a mutex while scheduling so each thread's host reads are submitted contiguously. + * + * @param datasources Input datasources + * @param source_indices Datasource index for each byte range + * @param offsets Offset for each byte range + * @param sizes Size for each byte range + * @return Pending host-read tasks, one per range in input order + */ +std::vector> read_ranges_to_host( cudf::host_span const> datasources, cudf::host_span source_indices, cudf::host_span offsets, cudf::host_span sizes) { static std::mutex host_read_mutex; - using host_read_buffer = std::unique_ptr; std::vector> host_read_tasks; host_read_tasks.reserve(source_indices.size()); + + auto iter = cuda::make_zip_iterator(source_indices.begin(), offsets.begin(), sizes.begin()); + + // Schedule host reads holding the `host_read_mutex` so that all reads for a caller thread + // are scheduled without interleaving with reads from other threads yielding better pipelining { std::scoped_lock lock(host_read_mutex); - auto iter = cuda::make_zip_iterator(source_indices.begin(), offsets.begin(), sizes.begin()); + std::for_each(iter, iter + source_indices.size(), [&](auto const& tuple) { - auto& datasource = datasources[cuda::std::get<0>(tuple)].get(); + auto const src_idx = cuda::std::get<0>(tuple); auto const io_offset = cuda::std::get<1>(tuple); auto const io_size = cuda::std::get<2>(tuple); + + auto& datasource = datasources[src_idx].get(); host_read_tasks.emplace_back(cudf::detail::host_worker_pool().submit_task( [&datasource, io_offset, io_size]() -> host_read_buffer { return datasource.host_read(io_offset, io_size); })); }); } + return host_read_tasks; +} + +/** + * @brief Reads the given ranges directly to device and copies host-resident ranges to device. + * + * Holds the mutex while scheduling so each thread's device reads and its batched host-to-device + * copy are submitted contiguously without interleaving with other threads. + * + * @param datasources Input datasources + * @param device_source_indices Datasource index for each device-read-preferred range + * @param device_offsets Offset for each device-read-preferred range + * @param device_sizes Size for each device-read-preferred range + * @param device_destinations Device destination for each device-read-preferred range + * @param copy_dsts Device destinations for the host-to-device copies + * @param copy_srcs Host sources for the host-to-device copies + * @param copy_sizes Sizes for the host-to-device copies + * @param stream CUDA stream used to schedule the device reads and copies + * @return Futures for the scheduled device reads + */ +std::vector> copy_ranges_to_device( + cudf::host_span const> datasources, + cudf::host_span device_source_indices, + cudf::host_span device_offsets, + cudf::host_span device_sizes, + cudf::host_span device_destinations, + cudf::host_span copy_dsts, + cudf::host_span copy_srcs, + cudf::host_span copy_sizes, + rmm::cuda_stream_view stream) +{ + static std::mutex device_read_mutex; + + std::vector> device_read_tasks{}; + device_read_tasks.reserve(device_source_indices.size()); + + // `device_read_async` is not guaranteed to follow stream-ordering (see datasource API docs) + if (not device_source_indices.empty()) { stream.synchronize(); } + + auto iter = cuda::make_zip_iterator(device_source_indices.begin(), + device_offsets.begin(), + device_sizes.begin(), + device_destinations.begin()); + + // Schedule device reads holding the `device_read_mutex` so that all reads for a caller thread + // are scheduled without interleaving with reads from other threads yielding better pipelining + { + std::scoped_lock lock(device_read_mutex); + + std::for_each(iter, iter + device_source_indices.size(), [&](auto const& tuple) { + auto const src_idx = cuda::std::get<0>(tuple); + auto const io_offset = cuda::std::get<1>(tuple); + auto const io_size = cuda::std::get<2>(tuple); + auto const dest = cuda::std::get<3>(tuple); + + auto& datasource = datasources[src_idx].get(); + device_read_tasks.emplace_back( + datasource.device_read_async(io_offset, io_size, dest, stream)); + }); + + // Schedule a batched memcpy from host buffers to device + if (not copy_dsts.empty()) { + CUDF_CUDA_TRY(cudf::detail::memcpy_batch_async( + copy_dsts.data(), copy_srcs.data(), copy_sizes.data(), copy_dsts.size(), stream)); + } + } - std::vector host_buffers; - host_buffers.reserve(host_read_tasks.size()); - std::transform(host_read_tasks.begin(), - host_read_tasks.end(), - std::back_inserter(host_buffers), - [](auto& task) { return task.get(); }); - return host_buffers; + return device_read_tasks; } std::tuple, @@ -238,8 +313,6 @@ fetch_byte_ranges_to_device_async_impl( rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) { - static std::mutex device_read_mutex; - auto const num_sources = datasources.size(); CUDF_EXPECTS(num_sources == byte_ranges_per_source.size(), @@ -252,15 +325,23 @@ fetch_byte_ranges_to_device_async_impl( std::size_t{0}, [](auto acc, auto const& ranges) { return acc + ranges.size(); }); - // IO descriptors - std::vector io_source_indices; - std::vector io_offsets; - std::vector io_sizes; - std::vector destinations; - io_source_indices.reserve(total_byte_ranges); - io_offsets.reserve(total_byte_ranges); - io_sizes.reserve(total_byte_ranges); - destinations.reserve(total_byte_ranges); + // IO descriptors, split up front into device-read-preferred and host-read buckets + std::vector device_source_indices; + std::vector device_offsets; + std::vector device_sizes; + std::vector device_destinations; + std::vector host_source_indices; + std::vector host_offsets; + std::vector host_sizes; + std::vector host_destinations; // H2D copy destinations (memcpy wants void*) + device_source_indices.reserve(total_byte_ranges); + device_offsets.reserve(total_byte_ranges); + device_sizes.reserve(total_byte_ranges); + device_destinations.reserve(total_byte_ranges); + host_source_indices.reserve(total_byte_ranges); + host_offsets.reserve(total_byte_ranges); + host_sizes.reserve(total_byte_ranges); + host_destinations.reserve(total_byte_ranges); // Allocate one device buffer per byte ranges of a datasource std::vector column_chunk_buffers{}; @@ -310,90 +391,61 @@ fetch_byte_ranges_to_device_async_impl( next_chunk++; } if (io_size != 0) { - io_source_indices.push_back(source_idx); - io_offsets.push_back(io_offset); - io_sizes.push_back(io_size); - destinations.push_back(const_cast(column_chunk_data[chunk].data())); + auto* const dest = const_cast(column_chunk_data[chunk].data()); + // Route each coalesced request to the device-read or host-read bucket up front + if (datasources[source_idx].get().is_device_read_preferred(io_size)) { + device_source_indices.push_back(source_idx); + device_offsets.push_back(io_offset); + device_sizes.push_back(io_size); + device_destinations.push_back(dest); + } else { + host_source_indices.push_back(source_idx); + host_offsets.push_back(io_offset); + host_sizes.push_back(io_size); + host_destinations.push_back(dest); + } } chunk = next_chunk; } }); - CUDF_EXPECTS(io_offsets.size() == io_sizes.size() and io_sizes.size() == destinations.size() and - io_source_indices.size() == io_offsets.size(), + CUDF_EXPECTS(device_source_indices.size() == device_offsets.size() and + device_offsets.size() == device_sizes.size() and + device_sizes.size() == device_destinations.size() and + host_source_indices.size() == host_offsets.size() and + host_offsets.size() == host_sizes.size() and + host_sizes.size() == host_destinations.size(), "Unexpected number of IO source indices, offsets, sizes, or destinations"); - using host_read_buffer = std::unique_ptr; - - // Vectors to hold futures from datasource - std::vector> device_read_tasks{}; - device_read_tasks.reserve(io_offsets.size()); + // Schedule host reads for the host-preferred ranges + auto host_read_tasks = + read_ranges_to_host(datasources, host_source_indices, host_offsets, host_sizes); // Vectors to store intermediate host buffers and relevant pointers std::vector host_buffers{}; std::vector copy_srcs{}; - std::vector copy_dsts{}; - std::vector copy_sizes{}; - copy_dsts.reserve(io_offsets.size()); - copy_sizes.reserve(io_offsets.size()); - auto iter = cuda::make_zip_iterator( - io_source_indices.begin(), io_offsets.begin(), io_sizes.begin(), destinations.begin()); + // Complete host reads + if (not host_read_tasks.empty()) { + host_buffers.reserve(host_read_tasks.size()); + copy_srcs.reserve(host_read_tasks.size()); - // Collect the host-preferred ranges and read them to host through the shared helper - std::vector host_source_indices; - std::vector host_offsets; - std::vector host_sizes; - std::for_each(iter, iter + io_offsets.size(), [&](auto const& tuple) { - auto const src_idx = cuda::std::get<0>(tuple); - auto const io_offset = cuda::std::get<1>(tuple); - auto const io_size = cuda::std::get<2>(tuple); - auto const dest = cuda::std::get<3>(tuple); - - auto& datasource = datasources[src_idx].get(); - if (not datasource.is_device_read_preferred(io_size)) { - host_source_indices.push_back(src_idx); - host_offsets.push_back(io_offset); - host_sizes.push_back(io_size); - copy_dsts.push_back(static_cast(dest)); - copy_sizes.push_back(io_size); + for (auto& task : host_read_tasks) { + host_buffers.emplace_back(task.get()); + copy_srcs.push_back(host_buffers.back().get()->data()); } - }); - - host_buffers = read_ranges_to_host(datasources, host_source_indices, host_offsets, host_sizes); - copy_srcs.reserve(host_buffers.size()); - for (auto const& buffer : host_buffers) { - copy_srcs.push_back(buffer->data()); } - // `device_read_async` is not guaranteed to follow stream-ordering (see datasource API docs) - stream.synchronize(); - - // Schedule device reads holding the `device_read_mutex` so that all reads for a caller thread - // are scheduled without interleaving with reads from other threads yielding better pipelining - { - std::scoped_lock lock(device_read_mutex); - - std::for_each(iter, iter + io_offsets.size(), [&](auto const& tuple) { - auto const src_idx = cuda::std::get<0>(tuple); - auto const io_offset = cuda::std::get<1>(tuple); - auto const io_size = cuda::std::get<2>(tuple); - auto const dest = cuda::std::get<3>(tuple); - - auto& datasource = datasources[src_idx].get(); - // Directly read the column chunk data to the device buffer if supported - if (datasource.is_device_read_preferred(io_size)) { - device_read_tasks.emplace_back( - datasource.device_read_async(io_offset, io_size, dest, stream)); - } - }); - - // Schedule a batched memcpy from host buffers to device - if (not host_buffers.empty()) { - CUDF_CUDA_TRY(cudf::detail::memcpy_batch_async( - copy_dsts.data(), copy_srcs.data(), copy_sizes.data(), copy_dsts.size(), stream)); - } - } + // Schedule the device reads and the batched host-to-device copy under the shared device mutex + auto device_read_tasks = copy_ranges_to_device(datasources, + device_source_indices, + device_offsets, + device_sizes, + device_destinations, + host_destinations, + copy_srcs, + host_sizes, + stream); // Synchronize stream if `memcpy_batch_async` was called to safely discard the host buffers if (not host_buffers.empty()) { stream.synchronize(); } @@ -429,14 +481,15 @@ fetch_bloom_filters_to_device_impl( std::size_t{0}, [](auto acc, auto const& bloom_ranges) { return acc + bloom_ranges.size(); }); - // Phase 1: speculatively read every filter to host. Absent filters keep a zero-length slot so the - // read buffers stay positionally aligned with the flattened filters + // Phase 1: Speculative reads; absent filters keep a zero-length slot std::vector spec_source_indices; std::vector spec_offsets; std::vector spec_sizes; spec_source_indices.reserve(total_filters); spec_offsets.reserve(total_filters); spec_sizes.reserve(total_filters); + + // Queue speculative reads std::for_each(cuda::counting_iterator(0), cuda::counting_iterator(num_sources), [&](std::size_t source_idx) { @@ -448,12 +501,17 @@ fetch_bloom_filters_to_device_impl( spec_sizes.push_back(static_cast(bloom_range.size())); }); }); - auto const spec_buffers = - read_ranges_to_host(datasources, spec_source_indices, spec_offsets, spec_sizes); - // Phase 2: parse each present filter's header for its sizes, total the aligned device bytes per - // source, and queue a follow-up read for any bitset the speculative read did not fully cover. + // Complete speculative reads + auto spec_read_tasks = + read_ranges_to_host(datasources, spec_source_indices, spec_offsets, spec_sizes); + std::vector spec_buffers{}; + spec_buffers.reserve(spec_read_tasks.size()); + for (auto& task : spec_read_tasks) { + spec_buffers.emplace_back(task.get()); + } + // Phase 2: Follow-up reads; absent filters keep a zero-length slot std::vector bloom_header_sizes(total_filters, 0); std::vector bloom_bitset_sizes(total_filters, 0); std::vector bloom_total_sizes(total_filters, 0); @@ -461,6 +519,8 @@ fetch_bloom_filters_to_device_impl( std::vector followup_source_indices; std::vector followup_offsets; std::vector followup_sizes; + + // Parse filters' headers and bitsets, and queue follow-up reads if needed std::for_each( cuda::counting_iterator(0), cuda::counting_iterator(total_filters), @@ -468,13 +528,20 @@ fetch_bloom_filters_to_device_impl( // Absent chunk: no bloom filter to read if (spec_sizes[filter_idx] == 0) { return; } + // Parse the filter's header auto const bytes_read = spec_buffers[filter_idx]->size(); auto const header_info = detail::parse_bloom_filter_header({spec_buffers[filter_idx]->data(), bytes_read}); - CUDF_EXPECTS(header_info.has_value(), "Encountered an invalid bloom filter header"); + if (not header_info.has_value()) { + CUDF_LOG_WARN("Encountered an invalid bloom filter header. Skipping"); + return; + } auto const [header_bytes, bitset_bytes] = header_info.value(); - CUDF_EXPECTS(bitset_bytes % bloom_filter_block_bytes == 0, - "Bloom filter bitset size must be a multiple of 32 bytes"); + if (bitset_bytes % bloom_filter_block_bytes != 0) { + CUDF_LOG_WARN( + "Encountered a bloom filter bitset size that is not a multiple of 32 bytes. Skipping"); + return; + } bloom_header_sizes[filter_idx] = static_cast(header_bytes); bloom_bitset_sizes[filter_idx] = static_cast(bitset_bytes); bloom_total_sizes[filter_idx] = @@ -489,10 +556,17 @@ fetch_bloom_filters_to_device_impl( followup_sizes.push_back(bloom_bitset_sizes[filter_idx]); } }); - auto const followup_buffers = + + // Complete follow-up reads + auto followup_read_tasks = read_ranges_to_host(datasources, followup_source_indices, followup_offsets, followup_sizes); + std::vector followup_buffers{}; + followup_buffers.reserve(followup_read_tasks.size()); + for (auto& task : followup_read_tasks) { + followup_buffers.emplace_back(task.get()); + } - // Phase 3: one aligned device buffer per source + // Phase 3: Create one aligned device buffer per source std::vector bloom_filter_buffers; bloom_filter_buffers.reserve(num_sources); std::transform(per_source_device_size.begin(), @@ -500,15 +574,16 @@ fetch_bloom_filters_to_device_impl( std::back_inserter(bloom_filter_buffers), [&](std::size_t size) { return rmm::device_buffer(size, stream, aligned_mr); }); - // Phase 4: copy every bitset straight to its device slot in one batched transfer + // Phase 4: Batch copy all bitsets to their device slots std::vector bitset_spans_per_source(num_sources); - std::vector copy_dsts; std::vector copy_srcs; std::vector copy_sizes; copy_dsts.reserve(total_filters); copy_srcs.reserve(total_filters); copy_sizes.reserve(total_filters); + + // Queue bitset copies std::size_t filter_idx = 0; std::size_t followup_index = 0; for (std::size_t source_idx = 0; source_idx < num_sources; ++source_idx) { @@ -518,13 +593,16 @@ fetch_bloom_filters_to_device_impl( std::size_t device_offset = 0; for (std::size_t inner_idx = 0; inner_idx < bloom_ranges.size(); ++inner_idx, ++filter_idx) { auto const bitset_size = bloom_bitset_sizes[filter_idx]; + // Absent filter: leave an empty span if (bitset_size == 0) { continue; } + auto* const device_dst = device_base + device_offset; CUDF_EXPECTS(reinterpret_cast(device_dst) % bloom_filter_block_bytes == 0, "Encountered a misaligned bloom filter bitset"); bitset_spans_per_source[source_idx][inner_idx] = cudf::device_span{device_dst, bitset_size}; + if (spec_buffers[filter_idx]->size() >= bloom_total_sizes[filter_idx]) { // Whole bitset already read speculatively: copy it directly, stripping the header copy_srcs.push_back(spec_buffers[filter_idx]->data() + bloom_header_sizes[filter_idx]); @@ -540,10 +618,9 @@ fetch_bloom_filters_to_device_impl( } } - // One batched host-to-device copy for every bitset (bloom filters never use direct device reads). + // Reuse the device-copy helper to serialize this batched H2D copy with device-read submissions. if (not copy_dsts.empty()) { - CUDF_CUDA_TRY(cudf::detail::memcpy_batch_async( - copy_dsts.data(), copy_srcs.data(), copy_sizes.data(), copy_dsts.size(), stream)); + copy_ranges_to_device(datasources, {}, {}, {}, {}, copy_dsts, copy_srcs, copy_sizes, stream); stream.synchronize(); } return {std::move(bloom_filter_buffers), std::move(bitset_spans_per_source)}; From 7606a90d0cd2347b96348a5b982b99bd18b839b0 Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Thu, 16 Jul 2026 11:29:59 +0200 Subject: [PATCH 25/39] Revert `fetch_byte_ranges_to_device_async_impl` while use mutex --- .../io/parquet/io_utils/parquet_io_utils.cpp | 225 ++++++++---------- 1 file changed, 103 insertions(+), 122 deletions(-) diff --git a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp index 8893de1919cf..3d73a7eef809 100644 --- a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp +++ b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp @@ -192,6 +192,20 @@ std::vector> fetch_page_indexes_to using device_spans_per_source_type = std::vector>; using host_read_buffer = std::unique_ptr; +// Shared across the byte-range and bloom-filter fetch paths to serialize their host reads and their +// device reads/copies against each other. +std::mutex& host_read_mutex() +{ + static std::mutex mutex; + return mutex; +} + +std::mutex& device_read_mutex() +{ + static std::mutex mutex; + return mutex; +} + /** * @brief Schedules host reads for the given byte ranges. * @@ -209,8 +223,6 @@ std::vector> read_ranges_to_host( cudf::host_span offsets, cudf::host_span sizes) { - static std::mutex host_read_mutex; - std::vector> host_read_tasks; host_read_tasks.reserve(source_indices.size()); @@ -219,7 +231,7 @@ std::vector> read_ranges_to_host( // Schedule host reads holding the `host_read_mutex` so that all reads for a caller thread // are scheduled without interleaving with reads from other threads yielding better pipelining { - std::scoped_lock lock(host_read_mutex); + std::scoped_lock lock(host_read_mutex()); std::for_each(iter, iter + source_indices.size(), [&](auto const& tuple) { auto const src_idx = cuda::std::get<0>(tuple); @@ -236,73 +248,6 @@ std::vector> read_ranges_to_host( return host_read_tasks; } -/** - * @brief Reads the given ranges directly to device and copies host-resident ranges to device. - * - * Holds the mutex while scheduling so each thread's device reads and its batched host-to-device - * copy are submitted contiguously without interleaving with other threads. - * - * @param datasources Input datasources - * @param device_source_indices Datasource index for each device-read-preferred range - * @param device_offsets Offset for each device-read-preferred range - * @param device_sizes Size for each device-read-preferred range - * @param device_destinations Device destination for each device-read-preferred range - * @param copy_dsts Device destinations for the host-to-device copies - * @param copy_srcs Host sources for the host-to-device copies - * @param copy_sizes Sizes for the host-to-device copies - * @param stream CUDA stream used to schedule the device reads and copies - * @return Futures for the scheduled device reads - */ -std::vector> copy_ranges_to_device( - cudf::host_span const> datasources, - cudf::host_span device_source_indices, - cudf::host_span device_offsets, - cudf::host_span device_sizes, - cudf::host_span device_destinations, - cudf::host_span copy_dsts, - cudf::host_span copy_srcs, - cudf::host_span copy_sizes, - rmm::cuda_stream_view stream) -{ - static std::mutex device_read_mutex; - - std::vector> device_read_tasks{}; - device_read_tasks.reserve(device_source_indices.size()); - - // `device_read_async` is not guaranteed to follow stream-ordering (see datasource API docs) - if (not device_source_indices.empty()) { stream.synchronize(); } - - auto iter = cuda::make_zip_iterator(device_source_indices.begin(), - device_offsets.begin(), - device_sizes.begin(), - device_destinations.begin()); - - // Schedule device reads holding the `device_read_mutex` so that all reads for a caller thread - // are scheduled without interleaving with reads from other threads yielding better pipelining - { - std::scoped_lock lock(device_read_mutex); - - std::for_each(iter, iter + device_source_indices.size(), [&](auto const& tuple) { - auto const src_idx = cuda::std::get<0>(tuple); - auto const io_offset = cuda::std::get<1>(tuple); - auto const io_size = cuda::std::get<2>(tuple); - auto const dest = cuda::std::get<3>(tuple); - - auto& datasource = datasources[src_idx].get(); - device_read_tasks.emplace_back( - datasource.device_read_async(io_offset, io_size, dest, stream)); - }); - - // Schedule a batched memcpy from host buffers to device - if (not copy_dsts.empty()) { - CUDF_CUDA_TRY(cudf::detail::memcpy_batch_async( - copy_dsts.data(), copy_srcs.data(), copy_sizes.data(), copy_dsts.size(), stream)); - } - } - - return device_read_tasks; -} - std::tuple, std::vector, std::future> @@ -325,23 +270,15 @@ fetch_byte_ranges_to_device_async_impl( std::size_t{0}, [](auto acc, auto const& ranges) { return acc + ranges.size(); }); - // IO descriptors, split up front into device-read-preferred and host-read buckets - std::vector device_source_indices; - std::vector device_offsets; - std::vector device_sizes; - std::vector device_destinations; - std::vector host_source_indices; - std::vector host_offsets; - std::vector host_sizes; - std::vector host_destinations; // H2D copy destinations (memcpy wants void*) - device_source_indices.reserve(total_byte_ranges); - device_offsets.reserve(total_byte_ranges); - device_sizes.reserve(total_byte_ranges); - device_destinations.reserve(total_byte_ranges); - host_source_indices.reserve(total_byte_ranges); - host_offsets.reserve(total_byte_ranges); - host_sizes.reserve(total_byte_ranges); - host_destinations.reserve(total_byte_ranges); + // IO descriptors + std::vector io_source_indices; + std::vector io_offsets; + std::vector io_sizes; + std::vector destinations; + io_source_indices.reserve(total_byte_ranges); + io_offsets.reserve(total_byte_ranges); + io_sizes.reserve(total_byte_ranges); + destinations.reserve(total_byte_ranges); // Allocate one device buffer per byte ranges of a datasource std::vector column_chunk_buffers{}; @@ -391,44 +328,66 @@ fetch_byte_ranges_to_device_async_impl( next_chunk++; } if (io_size != 0) { - auto* const dest = const_cast(column_chunk_data[chunk].data()); - // Route each coalesced request to the device-read or host-read bucket up front - if (datasources[source_idx].get().is_device_read_preferred(io_size)) { - device_source_indices.push_back(source_idx); - device_offsets.push_back(io_offset); - device_sizes.push_back(io_size); - device_destinations.push_back(dest); - } else { - host_source_indices.push_back(source_idx); - host_offsets.push_back(io_offset); - host_sizes.push_back(io_size); - host_destinations.push_back(dest); - } + io_source_indices.push_back(source_idx); + io_offsets.push_back(io_offset); + io_sizes.push_back(io_size); + destinations.push_back(const_cast(column_chunk_data[chunk].data())); } chunk = next_chunk; } }); - CUDF_EXPECTS(device_source_indices.size() == device_offsets.size() and - device_offsets.size() == device_sizes.size() and - device_sizes.size() == device_destinations.size() and - host_source_indices.size() == host_offsets.size() and - host_offsets.size() == host_sizes.size() and - host_sizes.size() == host_destinations.size(), + CUDF_EXPECTS(io_offsets.size() == io_sizes.size() and io_sizes.size() == destinations.size() and + io_source_indices.size() == io_offsets.size(), "Unexpected number of IO source indices, offsets, sizes, or destinations"); - // Schedule host reads for the host-preferred ranges - auto host_read_tasks = - read_ranges_to_host(datasources, host_source_indices, host_offsets, host_sizes); + using host_read_buffer = std::unique_ptr; + + // Vectors to hold futures from datasource + std::vector> device_read_tasks{}; + std::vector> host_read_tasks{}; + device_read_tasks.reserve(io_offsets.size()); + host_read_tasks.reserve(io_offsets.size()); // Vectors to store intermediate host buffers and relevant pointers std::vector host_buffers{}; std::vector copy_srcs{}; + std::vector copy_dsts{}; + std::vector copy_sizes{}; + copy_dsts.reserve(io_offsets.size()); + copy_sizes.reserve(io_offsets.size()); + + auto iter = cuda::make_zip_iterator( + io_source_indices.begin(), io_offsets.begin(), io_sizes.begin(), destinations.begin()); + + // Schedule host reads holding the `host_read_mutex` so that all reads for a caller thread + // are scheduled without interleaving with reads from other threads yielding better pipelining + { + std::scoped_lock lock(host_read_mutex()); + + std::for_each(iter, iter + io_offsets.size(), [&](auto const& tuple) { + auto const src_idx = cuda::std::get<0>(tuple); + auto const io_offset = cuda::std::get<1>(tuple); + auto const io_size = cuda::std::get<2>(tuple); + auto const dest = cuda::std::get<3>(tuple); + + auto& datasource = datasources[src_idx].get(); + if (not datasource.is_device_read_preferred(io_size)) { + // Asynchronously read column chunk data to a host buffer + host_read_tasks.emplace_back(cudf::detail::host_worker_pool().submit_task( + [&datasource, io_offset, io_size]() -> host_read_buffer { + return datasource.host_read(io_offset, io_size); + })); + copy_dsts.push_back(static_cast(dest)); + copy_sizes.push_back(io_size); + } + }); + } // Complete host reads if (not host_read_tasks.empty()) { - host_buffers.reserve(host_read_tasks.size()); copy_srcs.reserve(host_read_tasks.size()); + host_buffers.reserve(host_read_tasks.size()); for (auto& task : host_read_tasks) { host_buffers.emplace_back(task.get()); @@ -436,16 +395,34 @@ fetch_byte_ranges_to_device_async_impl( } } - // Schedule the device reads and the batched host-to-device copy under the shared device mutex - auto device_read_tasks = copy_ranges_to_device(datasources, - device_source_indices, - device_offsets, - device_sizes, - device_destinations, - host_destinations, - copy_srcs, - host_sizes, - stream); + // `device_read_async` is not guaranteed to follow stream-ordering (see datasource API docs) + stream.synchronize(); + + // Schedule device reads holding the `device_read_mutex` so that all reads for a caller thread + // are scheduled without interleaving with reads from other threads yielding better pipelining + { + std::scoped_lock lock(device_read_mutex()); + + std::for_each(iter, iter + io_offsets.size(), [&](auto const& tuple) { + auto const src_idx = cuda::std::get<0>(tuple); + auto const io_offset = cuda::std::get<1>(tuple); + auto const io_size = cuda::std::get<2>(tuple); + auto const dest = cuda::std::get<3>(tuple); + + auto& datasource = datasources[src_idx].get(); + // Directly read the column chunk data to the device buffer if supported + if (datasource.is_device_read_preferred(io_size)) { + device_read_tasks.emplace_back( + datasource.device_read_async(io_offset, io_size, dest, stream)); + } + }); + + // Schedule a batched memcpy from host buffers to device + if (not host_buffers.empty()) { + CUDF_CUDA_TRY(cudf::detail::memcpy_batch_async( + copy_dsts.data(), copy_srcs.data(), copy_sizes.data(), copy_dsts.size(), stream)); + } + } // Synchronize stream if `memcpy_batch_async` was called to safely discard the host buffers if (not host_buffers.empty()) { stream.synchronize(); } @@ -618,9 +595,13 @@ fetch_bloom_filters_to_device_impl( } } - // Reuse the device-copy helper to serialize this batched H2D copy with device-read submissions. + // One batched host-to-device copy for every bitset (bloom filters never use direct device reads). if (not copy_dsts.empty()) { - copy_ranges_to_device(datasources, {}, {}, {}, {}, copy_dsts, copy_srcs, copy_sizes, stream); + { + std::scoped_lock lock(device_read_mutex()); + CUDF_CUDA_TRY(cudf::detail::memcpy_batch_async( + copy_dsts.data(), copy_srcs.data(), copy_sizes.data(), copy_dsts.size(), stream)); + } stream.synchronize(); } return {std::move(bloom_filter_buffers), std::move(bitset_spans_per_source)}; From 44a57435203ae5e5a64d32f5d90115c2e7ee66f1 Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Thu, 16 Jul 2026 11:42:09 +0200 Subject: [PATCH 26/39] Refactor `read_ranges_to_host` to return host buffers directly instead of futures. Update comments for clarity and improve mutex usage documentation. --- .../io/parquet/io_utils/parquet_io_utils.cpp | 34 ++++++++----------- 1 file changed, 14 insertions(+), 20 deletions(-) diff --git a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp index 3d73a7eef809..a046c8f342d1 100644 --- a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp +++ b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp @@ -192,8 +192,7 @@ std::vector> fetch_page_indexes_to using device_spans_per_source_type = std::vector>; using host_read_buffer = std::unique_ptr; -// Shared across the byte-range and bloom-filter fetch paths to serialize their host reads and their -// device reads/copies against each other. +// Shared across the byte-range and bloom-filter fetch paths to serialize their readings std::mutex& host_read_mutex() { static std::mutex mutex; @@ -207,7 +206,7 @@ std::mutex& device_read_mutex() } /** - * @brief Schedules host reads for the given byte ranges. + * @brief Reads the given byte ranges to host. * * Holds a mutex while scheduling so each thread's host reads are submitted contiguously. * @@ -215,9 +214,9 @@ std::mutex& device_read_mutex() * @param source_indices Datasource index for each byte range * @param offsets Offset for each byte range * @param sizes Size for each byte range - * @return Pending host-read tasks, one per range in input order + * @return Host buffers for the read ranges, in input order */ -std::vector> read_ranges_to_host( +std::vector read_ranges_to_host( cudf::host_span const> datasources, cudf::host_span source_indices, cudf::host_span offsets, @@ -245,7 +244,14 @@ std::vector> read_ranges_to_host( })); }); } - return host_read_tasks; + + std::vector host_buffers; + host_buffers.reserve(host_read_tasks.size()); + std::transform(host_read_tasks.begin(), + host_read_tasks.end(), + std::back_inserter(host_buffers), + [](auto& task) { return task.get(); }); + return host_buffers; } std::tuple, @@ -479,14 +485,8 @@ fetch_bloom_filters_to_device_impl( }); }); - // Complete speculative reads - auto spec_read_tasks = + auto const spec_buffers = read_ranges_to_host(datasources, spec_source_indices, spec_offsets, spec_sizes); - std::vector spec_buffers{}; - spec_buffers.reserve(spec_read_tasks.size()); - for (auto& task : spec_read_tasks) { - spec_buffers.emplace_back(task.get()); - } // Phase 2: Follow-up reads; absent filters keep a zero-length slot std::vector bloom_header_sizes(total_filters, 0); @@ -534,14 +534,8 @@ fetch_bloom_filters_to_device_impl( } }); - // Complete follow-up reads - auto followup_read_tasks = + auto const followup_buffers = read_ranges_to_host(datasources, followup_source_indices, followup_offsets, followup_sizes); - std::vector followup_buffers{}; - followup_buffers.reserve(followup_read_tasks.size()); - for (auto& task : followup_read_tasks) { - followup_buffers.emplace_back(task.get()); - } // Phase 3: Create one aligned device buffer per source std::vector bloom_filter_buffers; From b3741f9966eac62a5c877c247f99b64a2d8907dd Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Thu, 16 Jul 2026 16:49:18 +0200 Subject: [PATCH 27/39] Refactor `fetch_bloom_filters_to_device_impl` to improve clarity and efficiency. Replace `per_source_device_size` with `bitset_buffer_size_per_source` for better alignment of device buffers, and update comments to reflect changes in the batching process for bitset copying. --- .../io/parquet/io_utils/parquet_io_utils.cpp | 23 ++++++++++--------- 1 file changed, 12 insertions(+), 11 deletions(-) diff --git a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp index a046c8f342d1..f1246bb6a22c 100644 --- a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp +++ b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp @@ -492,11 +492,13 @@ fetch_bloom_filters_to_device_impl( std::vector bloom_header_sizes(total_filters, 0); std::vector bloom_bitset_sizes(total_filters, 0); std::vector bloom_total_sizes(total_filters, 0); - std::vector per_source_device_size(num_sources, 0); std::vector followup_source_indices; std::vector followup_offsets; std::vector followup_sizes; + // Accumulated bitset sizes for each source, rounded up to 32-byte boundaries + std::vector bitset_buffer_size_per_source(num_sources, 0); + // Parse filters' headers and bitsets, and queue follow-up reads if needed std::for_each( cuda::counting_iterator(0), @@ -523,7 +525,7 @@ fetch_bloom_filters_to_device_impl( bloom_bitset_sizes[filter_idx] = static_cast(bitset_bytes); bloom_total_sizes[filter_idx] = bloom_header_sizes[filter_idx] + bloom_bitset_sizes[filter_idx]; - per_source_device_size[spec_source_indices[filter_idx]] += + bitset_buffer_size_per_source[spec_source_indices[filter_idx]] += cudf::util::round_up_safe(bloom_bitset_sizes[filter_idx], bloom_filter_block_bytes); // Speculative read did not reach the whole bitset: queue a follow-up read for it @@ -537,15 +539,14 @@ fetch_bloom_filters_to_device_impl( auto const followup_buffers = read_ranges_to_host(datasources, followup_source_indices, followup_offsets, followup_sizes); - // Phase 3: Create one aligned device buffer per source - std::vector bloom_filter_buffers; - bloom_filter_buffers.reserve(num_sources); - std::transform(per_source_device_size.begin(), - per_source_device_size.end(), - std::back_inserter(bloom_filter_buffers), + // Phase 3: Batch copy all bitsets to their device slots + std::vector bitset_buffers_per_source; + bitset_buffers_per_source.reserve(num_sources); + std::transform(bitset_buffer_size_per_source.begin(), + bitset_buffer_size_per_source.end(), + std::back_inserter(bitset_buffers_per_source), [&](std::size_t size) { return rmm::device_buffer(size, stream, aligned_mr); }); - // Phase 4: Batch copy all bitsets to their device slots std::vector bitset_spans_per_source(num_sources); std::vector copy_dsts; std::vector copy_srcs; @@ -560,7 +561,7 @@ fetch_bloom_filters_to_device_impl( for (std::size_t source_idx = 0; source_idx < num_sources; ++source_idx) { auto const& bloom_ranges = bloom_filter_byte_ranges_per_source[source_idx]; bitset_spans_per_source[source_idx] = device_spans_per_source_type(bloom_ranges.size()); - auto* const device_base = static_cast(bloom_filter_buffers[source_idx].data()); + auto* const device_base = static_cast(bitset_buffers_per_source[source_idx].data()); std::size_t device_offset = 0; for (std::size_t inner_idx = 0; inner_idx < bloom_ranges.size(); ++inner_idx, ++filter_idx) { auto const bitset_size = bloom_bitset_sizes[filter_idx]; @@ -598,7 +599,7 @@ fetch_bloom_filters_to_device_impl( } stream.synchronize(); } - return {std::move(bloom_filter_buffers), std::move(bitset_spans_per_source)}; + return {std::move(bitset_buffers_per_source), std::move(bitset_spans_per_source)}; } } // namespace From 344a3b351fd9217570fb4b09426fd39667db93cd Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Fri, 17 Jul 2026 11:28:45 +0200 Subject: [PATCH 28/39] Enhance `read_bloom_filters` and `read_ranges_to_host` functions for improved error handling and clarity. Added bounds checking for bloom filter lengths and updated the `read_ranges_to_host` function to accept destination offsets and a host buffer, ensuring complete reads. Refactored mutex usage documentation for better understanding. --- cpp/src/io/parquet/bloom_filter_reader.cu | 6 +- .../io/parquet/io_utils/parquet_io_utils.cpp | 311 ++++++++++-------- 2 files changed, 180 insertions(+), 137 deletions(-) diff --git a/cpp/src/io/parquet/bloom_filter_reader.cu b/cpp/src/io/parquet/bloom_filter_reader.cu index 39cdd67830b7..a8a3fe5a5eb8 100644 --- a/cpp/src/io/parquet/bloom_filter_reader.cu +++ b/cpp/src/io/parquet/bloom_filter_reader.cu @@ -384,9 +384,11 @@ aggregate_reader_metadata::read_bloom_filters( auto const length = col_meta.bloom_filter_length.has_value() ? static_cast(col_meta.bloom_filter_length.value()) : std::min(speculative_read_size, source_size - offset); - source_ranges.push_back(cudf::io::text::byte_range_info{offset, length}); + CUDF_EXPECTS(length >= 0 and offset + length <= source_size, + "Bloom filter length is out of datasource bounds"); + source_ranges.push_back({offset, length}); } else { - source_ranges.push_back(cudf::io::text::byte_range_info{0, 0}); + source_ranges.push_back({0, 0}); } }); }); diff --git a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp index f1246bb6a22c..6bdb29ec2cdd 100644 --- a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp +++ b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp @@ -12,6 +12,7 @@ #include #include #include +#include #include #include #include @@ -53,6 +54,30 @@ namespace cudf::io::parquet { namespace { +using device_spans_per_source_type = std::vector>; +using host_read_buffer = std::unique_ptr; + +/** + * @brief Mutex serializing host reads across the byte-range and bloom-filter fetch paths + * + * Function-local static (shared, not a global) so each caller thread's host reads are scheduled + * contiguously without interleaving with other threads. + */ +std::mutex& host_read_mutex() +{ + static std::mutex mutex; + return mutex; +} + +/** + * @brief Mutex serializing device reads and host-to-device copies across those same fetch paths + */ +std::mutex& device_read_mutex() +{ + static std::mutex mutex; + return mutex; +} + /** * @brief Dispatches each indexed task and collects the results * @@ -189,43 +214,34 @@ std::vector> fetch_page_indexes_to }); } -using device_spans_per_source_type = std::vector>; -using host_read_buffer = std::unique_ptr; - -// Shared across the byte-range and bloom-filter fetch paths to serialize their readings -std::mutex& host_read_mutex() -{ - static std::mutex mutex; - return mutex; -} - -std::mutex& device_read_mutex() -{ - static std::mutex mutex; - return mutex; -} - /** - * @brief Reads the given byte ranges to host. + * @brief Reads the given byte ranges into a caller-provided host buffer. * - * Holds a mutex while scheduling so each thread's host reads are submitted contiguously. + * Holds a mutex while scheduling so each thread's host reads are submitted contiguously. Each range + * is read into `dst` at its destination offset; zero-size ranges are skipped. * * @param datasources Input datasources * @param source_indices Datasource index for each byte range - * @param offsets Offset for each byte range + * @param offsets Datasource offset for each byte range * @param sizes Size for each byte range - * @return Host buffers for the read ranges, in input order + * @param dst_offsets Destination offset into `dst` for each byte range + * @param dst Host buffer that receives the read data */ -std::vector read_ranges_to_host( +void read_ranges_to_host( cudf::host_span const> datasources, cudf::host_span source_indices, cudf::host_span offsets, - cudf::host_span sizes) + cudf::host_span sizes, + cudf::host_span dst_offsets, + cudf::host_span dst) { - std::vector> host_read_tasks; + std::vector> host_read_tasks; + std::vector expected_sizes; host_read_tasks.reserve(source_indices.size()); + expected_sizes.reserve(source_indices.size()); - auto iter = cuda::make_zip_iterator(source_indices.begin(), offsets.begin(), sizes.begin()); + auto iter = cuda::make_zip_iterator( + source_indices.begin(), offsets.begin(), sizes.begin(), dst_offsets.begin()); // Schedule host reads holding the `host_read_mutex` so that all reads for a caller thread // are scheduled without interleaving with reads from other threads yielding better pipelining @@ -233,25 +249,30 @@ std::vector read_ranges_to_host( std::scoped_lock lock(host_read_mutex()); std::for_each(iter, iter + source_indices.size(), [&](auto const& tuple) { - auto const src_idx = cuda::std::get<0>(tuple); - auto const io_offset = cuda::std::get<1>(tuple); - auto const io_size = cuda::std::get<2>(tuple); + auto const src_idx = cuda::std::get<0>(tuple); + auto const io_offset = cuda::std::get<1>(tuple); + auto const io_size = cuda::std::get<2>(tuple); + auto const dst_offset = cuda::std::get<3>(tuple); - auto& datasource = datasources[src_idx].get(); + if (io_size == 0) { return; } + + auto& datasource = datasources[src_idx].get(); + auto* const dst_ptr = dst.data() + dst_offset; + expected_sizes.push_back(io_size); host_read_tasks.emplace_back(cudf::detail::host_worker_pool().submit_task( - [&datasource, io_offset, io_size]() -> host_read_buffer { - return datasource.host_read(io_offset, io_size); + [&datasource, io_offset, io_size, dst_ptr]() -> std::size_t { + return datasource.host_read(io_offset, io_size, dst_ptr); })); }); } - std::vector host_buffers; - host_buffers.reserve(host_read_tasks.size()); - std::transform(host_read_tasks.begin(), - host_read_tasks.end(), - std::back_inserter(host_buffers), - [](auto& task) { return task.get(); }); - return host_buffers; + // Complete the reads; every range must be read in full + std::for_each(cuda::counting_iterator(0), + cuda::counting_iterator(host_read_tasks.size()), + [&](std::size_t i) { + CUDF_EXPECTS(host_read_tasks[i].get() == expected_sizes[i], + "Failed to read complete byte range to host"); + }); } std::tuple, @@ -347,8 +368,6 @@ fetch_byte_ranges_to_device_async_impl( io_source_indices.size() == io_offsets.size(), "Unexpected number of IO source indices, offsets, sizes, or destinations"); - using host_read_buffer = std::unique_ptr; - // Vectors to hold futures from datasource std::vector> device_read_tasks{}; std::vector> host_read_tasks{}; @@ -464,53 +483,60 @@ fetch_bloom_filters_to_device_impl( std::size_t{0}, [](auto acc, auto const& bloom_ranges) { return acc + bloom_ranges.size(); }); - // Phase 1: Speculative reads; absent filters keep a zero-length slot - std::vector spec_source_indices; - std::vector spec_offsets; - std::vector spec_sizes; - spec_source_indices.reserve(total_filters); - spec_offsets.reserve(total_filters); - spec_sizes.reserve(total_filters); - - // Queue speculative reads - std::for_each(cuda::counting_iterator(0), - cuda::counting_iterator(num_sources), - [&](std::size_t source_idx) { - auto const& bloom_ranges = bloom_filter_byte_ranges_per_source[source_idx]; - std::for_each( - bloom_ranges.begin(), bloom_ranges.end(), [&](auto const& bloom_range) { - spec_source_indices.push_back(source_idx); - spec_offsets.push_back(static_cast(bloom_range.offset())); - spec_sizes.push_back(static_cast(bloom_range.size())); + // Phase 1: Initial read. Cover the complete bloom filter or enough bytes to parse the header + std::vector initial_source_indices(total_filters); + std::vector initial_offsets(total_filters); + std::vector initial_sizes(total_filters); + std::vector initial_dst_offsets(total_filters); + std::size_t total_initial_read_size = 0; + { + std::size_t filter_idx = 0; + std::for_each(cuda::counting_iterator(0), + cuda::counting_iterator(num_sources), + [&](auto const source_idx) { + auto const& bloom_ranges = bloom_filter_byte_ranges_per_source[source_idx]; + std::for_each(bloom_ranges.begin(), bloom_ranges.end(), [&](auto const& range) { + initial_source_indices[filter_idx] = source_idx; + initial_offsets[filter_idx] = static_cast(range.offset()); + initial_sizes[filter_idx] = static_cast(range.size()); + initial_dst_offsets[filter_idx] = total_initial_read_size; + total_initial_read_size += initial_sizes[filter_idx]; + ++filter_idx; }); - }); - - auto const spec_buffers = - read_ranges_to_host(datasources, spec_source_indices, spec_offsets, spec_sizes); - - // Phase 2: Follow-up reads; absent filters keep a zero-length slot - std::vector bloom_header_sizes(total_filters, 0); - std::vector bloom_bitset_sizes(total_filters, 0); - std::vector bloom_total_sizes(total_filters, 0); - std::vector followup_source_indices; - std::vector followup_offsets; - std::vector followup_sizes; + }); + } - // Accumulated bitset sizes for each source, rounded up to 32-byte boundaries - std::vector bitset_buffer_size_per_source(num_sources, 0); + // Read every initial bloom filter bytes into one host buffer + auto initial_buffer = cudf::detail::make_host_vector(total_initial_read_size, stream); + read_ranges_to_host(datasources, + initial_source_indices, + initial_offsets, + initial_sizes, + initial_dst_offsets, + initial_buffer); + + // Phase 2: Parse headers, organize bitset slots, and record deferred bitset reads + std::vector bitset_sizes(total_filters, 0); + std::vector bitset_src_addrs(total_filters, nullptr); + std::size_t total_device_size = 0; + + std::vector deferred_filter_indices; + std::vector deferred_source_indices; + std::vector deferred_offsets; + std::vector deferred_sizes; + std::vector deferred_dst_offsets; + std::size_t total_deferred_size = 0; - // Parse filters' headers and bitsets, and queue follow-up reads if needed std::for_each( cuda::counting_iterator(0), cuda::counting_iterator(total_filters), [&](std::size_t filter_idx) { - // Absent chunk: no bloom filter to read - if (spec_sizes[filter_idx] == 0) { return; } + // Absent filter: no bloom filter to read + if (initial_sizes[filter_idx] == 0) { return; } - // Parse the filter's header - auto const bytes_read = spec_buffers[filter_idx]->size(); + auto const* const filter_addr = initial_buffer.data() + initial_dst_offsets[filter_idx]; auto const header_info = - detail::parse_bloom_filter_header({spec_buffers[filter_idx]->data(), bytes_read}); + detail::parse_bloom_filter_header({filter_addr, initial_sizes[filter_idx]}); if (not header_info.has_value()) { CUDF_LOG_WARN("Encountered an invalid bloom filter header. Skipping"); return; @@ -521,32 +547,46 @@ fetch_bloom_filters_to_device_impl( "Encountered a bloom filter bitset size that is not a multiple of 32 bytes. Skipping"); return; } - bloom_header_sizes[filter_idx] = static_cast(header_bytes); - bloom_bitset_sizes[filter_idx] = static_cast(bitset_bytes); - bloom_total_sizes[filter_idx] = - bloom_header_sizes[filter_idx] + bloom_bitset_sizes[filter_idx]; - bitset_buffer_size_per_source[spec_source_indices[filter_idx]] += - cudf::util::round_up_safe(bloom_bitset_sizes[filter_idx], bloom_filter_block_bytes); - - // Speculative read did not reach the whole bitset: queue a follow-up read for it - if (bytes_read < bloom_total_sizes[filter_idx]) { - followup_source_indices.push_back(spec_source_indices[filter_idx]); - followup_offsets.push_back(spec_offsets[filter_idx] + bloom_header_sizes[filter_idx]); - followup_sizes.push_back(bloom_bitset_sizes[filter_idx]); + + bitset_sizes[filter_idx] = static_cast(bitset_bytes); + total_device_size += bitset_sizes[filter_idx]; + + auto const header_size = static_cast(header_bytes); + auto const total_bytes = header_size + bitset_bytes; + if (initial_sizes[filter_idx] >= total_bytes) { + // Whole bitset already in the host buffer: point at it, stripping the header + bitset_src_addrs[filter_idx] = filter_addr + header_size; + } else { + // Defer a read of only the bitset bytes + deferred_filter_indices.push_back(filter_idx); + deferred_source_indices.push_back(initial_source_indices[filter_idx]); + deferred_offsets.push_back(initial_offsets[filter_idx] + header_size); + deferred_sizes.push_back(bitset_sizes[filter_idx]); + deferred_dst_offsets.push_back(total_deferred_size); + total_deferred_size += bitset_sizes[filter_idx]; } }); - auto const followup_buffers = - read_ranges_to_host(datasources, followup_source_indices, followup_offsets, followup_sizes); - - // Phase 3: Batch copy all bitsets to their device slots - std::vector bitset_buffers_per_source; - bitset_buffers_per_source.reserve(num_sources); - std::transform(bitset_buffer_size_per_source.begin(), - bitset_buffer_size_per_source.end(), - std::back_inserter(bitset_buffers_per_source), - [&](std::size_t size) { return rmm::device_buffer(size, stream, aligned_mr); }); + // Phase 3: one device buffer holds every bitset + rmm::device_buffer bitset_buffer(total_device_size, stream, aligned_mr); + auto* const device_base = static_cast(bitset_buffer.data()); + + // Resolve deferred bitset reads + auto deferred_buffer = cudf::detail::make_host_vector(total_deferred_size, stream); + read_ranges_to_host(datasources, + deferred_source_indices, + deferred_offsets, + deferred_sizes, + deferred_dst_offsets, + deferred_buffer); + std::for_each(cuda::counting_iterator(0), + cuda::counting_iterator(deferred_filter_indices.size()), + [&](std::size_t i) { + bitset_src_addrs[deferred_filter_indices[i]] = + deferred_buffer.data() + deferred_dst_offsets[i]; + }); + // Build the per-source device spans and the batched copy list std::vector bitset_spans_per_source(num_sources); std::vector copy_dsts; std::vector copy_srcs; @@ -555,42 +595,40 @@ fetch_bloom_filters_to_device_impl( copy_srcs.reserve(total_filters); copy_sizes.reserve(total_filters); - // Queue bitset copies - std::size_t filter_idx = 0; - std::size_t followup_index = 0; - for (std::size_t source_idx = 0; source_idx < num_sources; ++source_idx) { - auto const& bloom_ranges = bloom_filter_byte_ranges_per_source[source_idx]; - bitset_spans_per_source[source_idx] = device_spans_per_source_type(bloom_ranges.size()); - auto* const device_base = static_cast(bitset_buffers_per_source[source_idx].data()); - std::size_t device_offset = 0; - for (std::size_t inner_idx = 0; inner_idx < bloom_ranges.size(); ++inner_idx, ++filter_idx) { - auto const bitset_size = bloom_bitset_sizes[filter_idx]; - - // Absent filter: leave an empty span - if (bitset_size == 0) { continue; } - - auto* const device_dst = device_base + device_offset; - CUDF_EXPECTS(reinterpret_cast(device_dst) % bloom_filter_block_bytes == 0, - "Encountered a misaligned bloom filter bitset"); - bitset_spans_per_source[source_idx][inner_idx] = - cudf::device_span{device_dst, bitset_size}; - - if (spec_buffers[filter_idx]->size() >= bloom_total_sizes[filter_idx]) { - // Whole bitset already read speculatively: copy it directly, stripping the header - copy_srcs.push_back(spec_buffers[filter_idx]->data() + bloom_header_sizes[filter_idx]); - } else { - // Bitset came from the follow-up read (queued in the same order) - auto const& buffer = followup_buffers[followup_index++]; - CUDF_EXPECTS(buffer->size() == bitset_size, "Failed to read complete bloom filter bitset"); - copy_srcs.push_back(buffer->data()); - } - copy_dsts.push_back(device_dst); - copy_sizes.push_back(bitset_size); - device_offset += cudf::util::round_up_safe(bitset_size, bloom_filter_block_bytes); - } - } + std::size_t filter_idx = 0; + std::size_t device_offset = 0; + std::for_each( + cuda::counting_iterator(0), + cuda::counting_iterator(num_sources), + [&](auto const source_idx) { + auto const& bloom_ranges = bloom_filter_byte_ranges_per_source[source_idx]; + bitset_spans_per_source[source_idx] = device_spans_per_source_type(bloom_ranges.size()); + std::for_each( + cuda::counting_iterator(0), + cuda::counting_iterator(bloom_ranges.size()), + [&](auto const inner_idx) { + auto const bitset_size = bitset_sizes[filter_idx]; + + // Absent or skipped filter: leave an empty span + if (bitset_size == 0) { + ++filter_idx; + return; + } + + auto* const device_dst = device_base + device_offset; + CUDF_EXPECTS(reinterpret_cast(device_dst) % bloom_filter_block_bytes == 0, + "Encountered a misaligned bloom filter bitset"); + bitset_spans_per_source[source_idx][inner_idx] = + cudf::device_span{device_dst, bitset_size}; + copy_dsts.push_back(device_dst); + copy_srcs.push_back(bitset_src_addrs[filter_idx]); + copy_sizes.push_back(bitset_size); + device_offset += bitset_size; + ++filter_idx; + }); + }); - // One batched host-to-device copy for every bitset (bloom filters never use direct device reads). + // One batched copy if (not copy_dsts.empty()) { { std::scoped_lock lock(device_read_mutex()); @@ -599,7 +637,10 @@ fetch_bloom_filters_to_device_impl( } stream.synchronize(); } - return {std::move(bitset_buffers_per_source), std::move(bitset_spans_per_source)}; + + std::vector bitset_buffers; + bitset_buffers.push_back(std::move(bitset_buffer)); + return {std::move(bitset_buffers), std::move(bitset_spans_per_source)}; } } // namespace From d36cddc9f176ecd48d1e97727e19119a05a1bef9 Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Fri, 17 Jul 2026 16:54:17 +0200 Subject: [PATCH 29/39] Refactor bloom filter handling in Parquet I/O. Updated function signatures to use a generic device memory resource instead of requiring aligned memory allocation. Removed unnecessary alignment checks and comments, simplifying the codebase. This change enhances clarity and maintains compatibility with existing memory management practices. --- cpp/include/cudf/io/parquet_io_utils.hpp | 14 ++++---------- cpp/src/io/parquet/bloom_filter_reader.cu | 19 +++---------------- .../io/parquet/io_utils/parquet_io_utils.cpp | 13 +++++++------ cpp/src/io/parquet/predicate_pushdown.cpp | 8 +------- cpp/src/io/parquet/reader_impl_helpers.hpp | 9 ++------- 5 files changed, 17 insertions(+), 46 deletions(-) diff --git a/cpp/include/cudf/io/parquet_io_utils.hpp b/cpp/include/cudf/io/parquet_io_utils.hpp index a7b13588162f..0ecfcfe4cfb3 100644 --- a/cpp/include/cudf/io/parquet_io_utils.hpp +++ b/cpp/include/cudf/io/parquet_io_utils.hpp @@ -157,14 +157,11 @@ fetch_byte_ranges_to_device_async( * * @ingroup io_utils * - * @note Device buffers for bloom filter byte ranges must be allocated using a 32 byte aligned - * memory resource - * * @param datasource Input datasource * @param bloom_filter_byte_ranges Byte ranges of complete bloom filters to fetch, must span a * complete bloom filter * @param stream CUDA stream - * @param aligned_mr Device memory resource to allocate aligned memory for bloom filters + * @param mr Device memory resource used to allocate the returned device buffers * * @return A pair containing the device buffers and the device spans of the bitset data */ @@ -172,21 +169,18 @@ std::pair, std::vector bloom_filter_byte_ranges, rmm::cuda_stream_view stream, - rmm::device_async_resource_ref aligned_mr); + rmm::device_async_resource_ref mr); /** * @brief Fetches Parquet bloom filter bitsets from multiple datasources into device buffers * * @ingroup io_utils * - * @note Device buffers for bloom filter byte ranges must be allocated using a 32 byte aligned - * memory resource - * * @param datasources Input datasources * @param bloom_filter_byte_ranges_per_source Byte ranges of complete bloom filters to fetch, one * vector per datasource. Each byte range must span a complete bloom filter. * @param stream CUDA stream - * @param aligned_mr Device memory resource to allocate aligned memory for bloom filters + * @param mr Device memory resource used to allocate the returned device buffers * * @return A pair containing a vector of device buffers and a vector of vectors of device spans */ @@ -196,7 +190,7 @@ fetch_bloom_filters_to_device( cudf::host_span const> datasources, cudf::host_span const> bloom_filter_byte_ranges_per_source, rmm::cuda_stream_view stream, - rmm::device_async_resource_ref aligned_mr); + rmm::device_async_resource_ref mr); /** @} */ // end of group } // namespace io::parquet diff --git a/cpp/src/io/parquet/bloom_filter_reader.cu b/cpp/src/io/parquet/bloom_filter_reader.cu index a8a3fe5a5eb8..7ed77446d933 100644 --- a/cpp/src/io/parquet/bloom_filter_reader.cu +++ b/cpp/src/io/parquet/bloom_filter_reader.cu @@ -327,19 +327,6 @@ std::optional> parse_bloom_filter_header( static_cast(header.num_bytes)}; } -std::size_t aggregate_reader_metadata::get_bloom_filter_alignment() const -{ - // Required alignment: - // https://github.com/NVIDIA/cuCollections/blob/deab5799f3e4226cb8a49acf2199c03b14941ee4/include/cuco/detail/bloom_filter/bloom_filter_impl.cuh#L55-L67 - using policy_type = arrow_filter_policy; - auto constexpr alignment = alignof(cuco::bloom_filter_ref, - cuco::thread_scope_thread, - policy_type>::filter_block_type); - static_assert((alignment & (alignment - 1)) == 0, "Alignment must be a power of 2"); - return std::max(alignment, rmm::CUDA_ALLOCATION_ALIGNMENT); -} - std::pair, std::vector>> aggregate_reader_metadata::read_bloom_filters( host_span const> sources, @@ -347,7 +334,7 @@ aggregate_reader_metadata::read_bloom_filters( host_span column_schemas, size_type total_row_groups, rmm::cuda_stream_view stream, - rmm::device_async_resource_ref aligned_mr) const + rmm::device_async_resource_ref mr) const { // Descriptors for all the chunks that make up the selected columns auto const num_input_columns = column_schemas.size(); @@ -405,8 +392,8 @@ aggregate_reader_metadata::read_bloom_filters( return std::ref(*source); }); - auto [bloom_filter_buffers, bitset_spans_per_source] = fetch_bloom_filters_to_device( - datasource_refs, bloom_filter_byte_ranges_per_source, stream, aligned_mr); + auto [bloom_filter_buffers, bitset_spans_per_source] = + fetch_bloom_filters_to_device(datasource_refs, bloom_filter_byte_ranges_per_source, stream, mr); // Flatten the per-source bitset spans into per-chunk order std::vector> bloom_filter_data; diff --git a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp index 6bdb29ec2cdd..26243a9b55a2 100644 --- a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp +++ b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp @@ -468,13 +468,14 @@ fetch_bloom_filters_to_device_impl( cudf::host_span const> bloom_filter_byte_ranges_per_source, rmm::cuda_stream_view stream, - rmm::device_async_resource_ref aligned_mr) + rmm::device_async_resource_ref mr) { auto const num_sources = datasources.size(); CUDF_EXPECTS(num_sources == bloom_filter_byte_ranges_per_source.size(), "Encountered mismatch in number of datasources and bloom filter byte range spans"); // Each bitset must align to 32-byte boundaries, as required by cuco's Arrow bloom filter bitsets. + // TODO(NVIDIA/cuCollections#829): replace with a cuco-provided block-size / alignment accessor. auto constexpr bloom_filter_block_bytes = std::size_t{32}; auto const total_filters = @@ -568,7 +569,7 @@ fetch_bloom_filters_to_device_impl( }); // Phase 3: one device buffer holds every bitset - rmm::device_buffer bitset_buffer(total_device_size, stream, aligned_mr); + rmm::device_buffer bitset_buffer(total_device_size, bloom_filter_block_bytes, stream, mr); auto* const device_base = static_cast(bitset_buffer.data()); // Resolve deferred bitset reads @@ -743,7 +744,7 @@ fetch_bloom_filters_to_device( cudf::io::datasource& datasource, cudf::host_span bloom_filter_byte_ranges, rmm::cuda_stream_view stream, - rmm::device_async_resource_ref aligned_mr) + rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -756,7 +757,7 @@ fetch_bloom_filters_to_device( {datasources.data(), datasources.size()}, {bloom_filter_byte_ranges_per_source.data(), bloom_filter_byte_ranges_per_source.size()}, stream, - aligned_mr); + mr); return {std::move(buffers), std::move(fetched_byte_ranges.front())}; } @@ -768,7 +769,7 @@ fetch_bloom_filters_to_device( cudf::host_span const> bloom_filter_byte_ranges_per_source, rmm::cuda_stream_view stream, - rmm::device_async_resource_ref aligned_mr) + rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -783,7 +784,7 @@ fetch_bloom_filters_to_device( {bloom_filter_byte_range_spans_per_source.data(), bloom_filter_byte_range_spans_per_source.size()}, stream, - aligned_mr); + mr); } } // namespace cudf::io::parquet diff --git a/cpp/src/io/parquet/predicate_pushdown.cpp b/cpp/src/io/parquet/predicate_pushdown.cpp index b12c77285232..a762d16d575c 100644 --- a/cpp/src/io/parquet/predicate_pushdown.cpp +++ b/cpp/src/io/parquet/predicate_pushdown.cpp @@ -18,8 +18,6 @@ #include #include -#include - #include #include @@ -319,10 +317,6 @@ aggregate_reader_metadata::filter_row_groups( {std::make_optional(num_stats_filtered_row_groups), std::nullopt}}; } - // Aligned resource adaptor to allocate bloom filter buffers with - auto aligned_mr = rmm::mr::aligned_resource_adaptor(cudf::get_current_device_resource_ref(), - get_bloom_filter_alignment()); - // Read a vector of bloom filter bitset device buffers for all columns with equality // predicate(s) across all row groups auto const [bloom_filter_buffers, bloom_filter_data] = @@ -331,7 +325,7 @@ aggregate_reader_metadata::filter_row_groups( equality_col_schemas, num_stats_filtered_row_groups, stream, - aligned_mr); + cudf::get_current_device_resource_ref()); // No bloom filters, return early if (bloom_filter_data.empty()) { diff --git a/cpp/src/io/parquet/reader_impl_helpers.hpp b/cpp/src/io/parquet/reader_impl_helpers.hpp index 7da6589978a1..aed18995cebf 100644 --- a/cpp/src/io/parquet/reader_impl_helpers.hpp +++ b/cpp/src/io/parquet/reader_impl_helpers.hpp @@ -269,11 +269,6 @@ class aggregate_reader_metadata { */ void column_info_for_row_group(row_group_info& rg_info, size_t chunk_start_row) const; - /** - * @brief Returns the required alignment for bloom filter buffers - */ - [[nodiscard]] size_t get_bloom_filter_alignment() const; - /** * @brief Reads bloom filter bitsets for the specified columns from the given lists of row * groups. @@ -283,7 +278,7 @@ class aggregate_reader_metadata { * @param column_schemas Schema indices of columns whose bloom filters will be read * @param num_row_groups Number of row groups in the file * @param stream CUDA stream used for device memory operations and kernel launches - * @param aligned_mr Aligned device memory resource to allocate bloom filter buffers + * @param mr Device memory resource used to allocate bloom filter buffers * * @return A pair of the device buffers backing the bloom filter bitsets and a flattened, * per-chunk list of bitset device spans (empty spans for chunks without a bloom filter) @@ -295,7 +290,7 @@ class aggregate_reader_metadata { host_span column_schemas, size_type num_row_groups, rmm::cuda_stream_view stream, - rmm::device_async_resource_ref aligned_mr) const; + rmm::device_async_resource_ref mr) const; /** * @brief Collects Parquet types for the columns with the specified schema indices From d0af7d288f4c5478d62a5864cd6eab7c3d03faca Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Fri, 17 Jul 2026 17:19:26 +0200 Subject: [PATCH 30/39] Optimize mutex commit --- cpp/src/io/parquet/io_utils/parquet_io_utils.cpp | 9 ++++----- 1 file changed, 4 insertions(+), 5 deletions(-) diff --git a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp index 26243a9b55a2..253b589472dd 100644 --- a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp +++ b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp @@ -58,10 +58,8 @@ using device_spans_per_source_type = std::vector; /** - * @brief Mutex serializing host reads across the byte-range and bloom-filter fetch paths - * - * Function-local static (shared, not a global) so each caller thread's host reads are scheduled - * contiguously without interleaving with other threads. + * @brief Serializes host-read submission batches to avoid cross-thread request interleaving and + * completion stalls */ std::mutex& host_read_mutex() { @@ -70,7 +68,8 @@ std::mutex& host_read_mutex() } /** - * @brief Mutex serializing device reads and host-to-device copies across those same fetch paths + * @brief Serializes device reads and host-to-device copies to avoid cross-thread request + * interleaving and completion stalls */ std::mutex& device_read_mutex() { From 59e7df39e9c097f61302210992ed23353b85d5e8 Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Mon, 20 Jul 2026 00:33:44 +0200 Subject: [PATCH 31/39] Nested --- .../io/parquet/io_utils/parquet_io_utils.cpp | 162 +++++++++--------- 1 file changed, 82 insertions(+), 80 deletions(-) diff --git a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp index 253b589472dd..276e7195f045 100644 --- a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp +++ b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp @@ -483,27 +483,37 @@ fetch_bloom_filters_to_device_impl( std::size_t{0}, [](auto acc, auto const& bloom_ranges) { return acc + bloom_ranges.size(); }); + std::vector bitset_spans_per_source(num_sources); + std::vector> output_span_indices(total_filters); + // Phase 1: Initial read. Cover the complete bloom filter or enough bytes to parse the header std::vector initial_source_indices(total_filters); std::vector initial_offsets(total_filters); std::vector initial_sizes(total_filters); std::vector initial_dst_offsets(total_filters); std::size_t total_initial_read_size = 0; + { std::size_t filter_idx = 0; - std::for_each(cuda::counting_iterator(0), - cuda::counting_iterator(num_sources), - [&](auto const source_idx) { - auto const& bloom_ranges = bloom_filter_byte_ranges_per_source[source_idx]; - std::for_each(bloom_ranges.begin(), bloom_ranges.end(), [&](auto const& range) { - initial_source_indices[filter_idx] = source_idx; - initial_offsets[filter_idx] = static_cast(range.offset()); - initial_sizes[filter_idx] = static_cast(range.size()); - initial_dst_offsets[filter_idx] = total_initial_read_size; - total_initial_read_size += initial_sizes[filter_idx]; - ++filter_idx; - }); - }); + std::for_each( + cuda::counting_iterator(0), + cuda::counting_iterator(num_sources), + [&](auto const source_idx) { + auto const& bloom_ranges = bloom_filter_byte_ranges_per_source[source_idx]; + bitset_spans_per_source[source_idx] = device_spans_per_source_type(bloom_ranges.size()); + std::for_each(cuda::counting_iterator(0), + cuda::counting_iterator(bloom_ranges.size()), + [&](auto const span_idx) { + auto const& range = bloom_ranges[span_idx]; + output_span_indices[filter_idx] = {source_idx, span_idx}; + initial_source_indices[filter_idx] = source_idx; + initial_offsets[filter_idx] = static_cast(range.offset()); + initial_sizes[filter_idx] = static_cast(range.size()); + initial_dst_offsets[filter_idx] = total_initial_read_size; + total_initial_read_size += initial_sizes[filter_idx]; + ++filter_idx; + }); + }); } // Read every initial bloom filter bytes into one host buffer @@ -516,8 +526,14 @@ fetch_bloom_filters_to_device_impl( initial_buffer); // Phase 2: Parse headers, organize bitset slots, and record deferred bitset reads - std::vector bitset_sizes(total_filters, 0); - std::vector bitset_src_addrs(total_filters, nullptr); + std::vector copy_dsts; + std::vector copy_dst_offsets; + std::vector copy_srcs; + std::vector copy_sizes; + copy_dsts.reserve(total_filters); + copy_dst_offsets.reserve(total_filters); + copy_srcs.reserve(total_filters); + copy_sizes.reserve(total_filters); std::size_t total_device_size = 0; std::vector deferred_filter_indices; @@ -531,47 +547,59 @@ fetch_bloom_filters_to_device_impl( cuda::counting_iterator(0), cuda::counting_iterator(total_filters), [&](std::size_t filter_idx) { + auto const push_empty_filter = [&]() { + copy_dsts.push_back(nullptr); + copy_dst_offsets.push_back(0); + copy_srcs.push_back(nullptr); + copy_sizes.push_back(0); + }; + // Absent filter: no bloom filter to read - if (initial_sizes[filter_idx] == 0) { return; } + if (initial_sizes[filter_idx] == 0) { + push_empty_filter(); + return; + } auto const* const filter_addr = initial_buffer.data() + initial_dst_offsets[filter_idx]; auto const header_info = detail::parse_bloom_filter_header({filter_addr, initial_sizes[filter_idx]}); if (not header_info.has_value()) { CUDF_LOG_WARN("Encountered an invalid bloom filter header. Skipping"); + push_empty_filter(); return; } auto const [header_bytes, bitset_bytes] = header_info.value(); if (bitset_bytes % bloom_filter_block_bytes != 0) { CUDF_LOG_WARN( "Encountered a bloom filter bitset size that is not a multiple of 32 bytes. Skipping"); + push_empty_filter(); return; } - bitset_sizes[filter_idx] = static_cast(bitset_bytes); - total_device_size += bitset_sizes[filter_idx]; - auto const header_size = static_cast(header_bytes); - auto const total_bytes = header_size + bitset_bytes; - if (initial_sizes[filter_idx] >= total_bytes) { + auto const bitset_size = static_cast(bitset_bytes); + copy_dsts.push_back(nullptr); + copy_dst_offsets.push_back(total_device_size); + copy_sizes.push_back(bitset_size); + total_device_size += bitset_size; + + if (initial_sizes[filter_idx] >= header_size + bitset_size) { // Whole bitset already in the host buffer: point at it, stripping the header - bitset_src_addrs[filter_idx] = filter_addr + header_size; + copy_srcs.push_back(filter_addr + header_size); } else { - // Defer a read of only the bitset bytes - deferred_filter_indices.push_back(filter_idx); + // Whole bitset not in the host buffer: defer the read + copy_srcs.push_back(nullptr); + + deferred_filter_indices.push_back(copy_srcs.size() - 1); deferred_source_indices.push_back(initial_source_indices[filter_idx]); deferred_offsets.push_back(initial_offsets[filter_idx] + header_size); - deferred_sizes.push_back(bitset_sizes[filter_idx]); + deferred_sizes.push_back(bitset_size); deferred_dst_offsets.push_back(total_deferred_size); - total_deferred_size += bitset_sizes[filter_idx]; + total_deferred_size += bitset_size; } }); - // Phase 3: one device buffer holds every bitset - rmm::device_buffer bitset_buffer(total_device_size, bloom_filter_block_bytes, stream, mr); - auto* const device_base = static_cast(bitset_buffer.data()); - - // Resolve deferred bitset reads + // Phase 3: Resolve deferred reads, then batch copy all bitsets to the device auto deferred_buffer = cudf::detail::make_host_vector(total_deferred_size, stream); read_ranges_to_host(datasources, deferred_source_indices, @@ -582,59 +610,33 @@ fetch_bloom_filters_to_device_impl( std::for_each(cuda::counting_iterator(0), cuda::counting_iterator(deferred_filter_indices.size()), [&](std::size_t i) { - bitset_src_addrs[deferred_filter_indices[i]] = + copy_srcs[deferred_filter_indices[i]] = deferred_buffer.data() + deferred_dst_offsets[i]; }); - // Build the per-source device spans and the batched copy list - std::vector bitset_spans_per_source(num_sources); - std::vector copy_dsts; - std::vector copy_srcs; - std::vector copy_sizes; - copy_dsts.reserve(total_filters); - copy_srcs.reserve(total_filters); - copy_sizes.reserve(total_filters); - - std::size_t filter_idx = 0; - std::size_t device_offset = 0; - std::for_each( - cuda::counting_iterator(0), - cuda::counting_iterator(num_sources), - [&](auto const source_idx) { - auto const& bloom_ranges = bloom_filter_byte_ranges_per_source[source_idx]; - bitset_spans_per_source[source_idx] = device_spans_per_source_type(bloom_ranges.size()); - std::for_each( - cuda::counting_iterator(0), - cuda::counting_iterator(bloom_ranges.size()), - [&](auto const inner_idx) { - auto const bitset_size = bitset_sizes[filter_idx]; - - // Absent or skipped filter: leave an empty span - if (bitset_size == 0) { - ++filter_idx; - return; - } - - auto* const device_dst = device_base + device_offset; - CUDF_EXPECTS(reinterpret_cast(device_dst) % bloom_filter_block_bytes == 0, - "Encountered a misaligned bloom filter bitset"); - bitset_spans_per_source[source_idx][inner_idx] = - cudf::device_span{device_dst, bitset_size}; - copy_dsts.push_back(device_dst); - copy_srcs.push_back(bitset_src_addrs[filter_idx]); - copy_sizes.push_back(bitset_size); - device_offset += bitset_size; - ++filter_idx; - }); - }); + // Add the buffer base to every non-empty output span and copy destination. + rmm::device_buffer bitset_buffer(total_device_size, bloom_filter_block_bytes, stream, mr); + auto* const device_base = static_cast(bitset_buffer.data()); + std::for_each(cuda::counting_iterator(0), + cuda::counting_iterator(total_filters), + [&](std::size_t filter_idx) { + auto const bitset_size = copy_sizes[filter_idx]; + auto* const device_dst = + bitset_size == 0 ? nullptr : device_base + copy_dst_offsets[filter_idx]; + if (bitset_size != 0) { + CUDF_EXPECTS(rmm::is_pointer_aligned(device_dst, bloom_filter_block_bytes), + "Encountered a misaligned bloom filter bitset"); + } + auto const [source_idx, span_idx] = output_span_indices[filter_idx]; + bitset_spans_per_source[source_idx][span_idx] = {device_dst, bitset_size}; + copy_dsts[filter_idx] = device_dst; + }); - // One batched copy - if (not copy_dsts.empty()) { - { - std::scoped_lock lock(device_read_mutex()); - CUDF_CUDA_TRY(cudf::detail::memcpy_batch_async( - copy_dsts.data(), copy_srcs.data(), copy_sizes.data(), copy_dsts.size(), stream)); - } + // One batched copy (entries with a null source or zero size are ignored by the batch API) + if (total_device_size != 0) { + std::scoped_lock lock(device_read_mutex()); + CUDF_CUDA_TRY(cudf::detail::memcpy_batch_async( + copy_dsts.data(), copy_srcs.data(), copy_sizes.data(), total_filters, stream)); stream.synchronize(); } From 4f458b180bb22f04645adc310dbfab16b8c04d88 Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Mon, 20 Jul 2026 12:00:14 +0200 Subject: [PATCH 32/39] Refactor Parquet I/O functions to enhance clarity and efficiency. Updated return types and documentation for `fetch_bloom_filters_to_device` and `read_ranges_to_host` to better reflect their functionality. Adjusted speculative read size in `aggregate_reader_metadata::read_bloom_filters` for improved performance. Improved iterator usage in `fetch_bloom_filters_to_device_impl` for better readability. --- cpp/include/cudf/io/parquet_io_utils.hpp | 6 +- cpp/src/io/parquet/bloom_filter_reader.cu | 2 +- .../io/parquet/io_utils/parquet_io_utils.cpp | 135 +++++++++--------- 3 files changed, 75 insertions(+), 68 deletions(-) diff --git a/cpp/include/cudf/io/parquet_io_utils.hpp b/cpp/include/cudf/io/parquet_io_utils.hpp index 0ecfcfe4cfb3..ca10c215dd5b 100644 --- a/cpp/include/cudf/io/parquet_io_utils.hpp +++ b/cpp/include/cudf/io/parquet_io_utils.hpp @@ -163,7 +163,8 @@ fetch_byte_ranges_to_device_async( * @param stream CUDA stream * @param mr Device memory resource used to allocate the returned device buffers * - * @return A pair containing the device buffers and the device spans of the bitset data + * @return A pair containing buffers that own the fetched bitsets and one device span per input byte + * range */ std::pair, std::vector>> fetch_bloom_filters_to_device(cudf::io::datasource& datasource, @@ -182,7 +183,8 @@ fetch_bloom_filters_to_device(cudf::io::datasource& datasource, * @param stream CUDA stream * @param mr Device memory resource used to allocate the returned device buffers * - * @return A pair containing a vector of device buffers and a vector of vectors of device spans + * @return A pair containing buffers that own the fetched bitsets and per-source device spans, with + * one inner vector per datasource */ std::pair, std::vector>>> diff --git a/cpp/src/io/parquet/bloom_filter_reader.cu b/cpp/src/io/parquet/bloom_filter_reader.cu index 7ed77446d933..0163f15d0473 100644 --- a/cpp/src/io/parquet/bloom_filter_reader.cu +++ b/cpp/src/io/parquet/bloom_filter_reader.cu @@ -344,7 +344,7 @@ aggregate_reader_metadata::read_bloom_filters( auto have_bloom_filters = false; // Speculatively read when a bloom filter's length is absent, enough to cover the header (and // often the whole bitset). - auto constexpr speculative_read_size = int64_t{512}; + auto constexpr speculative_read_size = int64_t{256}; // Build complete bloom filter byte ranges (header + bitset) for every column chunk std::vector> bloom_filter_byte_ranges_per_source( row_group_indices.size()); diff --git a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp index 276e7195f045..83d2012ecbbb 100644 --- a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp +++ b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp @@ -37,6 +37,7 @@ #include #include #include +#include #include #include #include @@ -216,42 +217,44 @@ std::vector> fetch_page_indexes_to /** * @brief Reads the given byte ranges into a caller-provided host buffer. * - * Holds a mutex while scheduling so each thread's host reads are submitted contiguously. Each range - * is read into `dst` at its destination offset; zero-size ranges are skipped. + * Holds a mutex while scheduling so each thread's host reads are submitted contiguously. The ranges + * are packed consecutively into `dst` in iterator order; zero-size ranges are skipped. * * @param datasources Input datasources - * @param source_indices Datasource index for each byte range - * @param offsets Datasource offset for each byte range - * @param sizes Size for each byte range - * @param dst_offsets Destination offset into `dst` for each byte range + * @param source_indices Iterator over the datasource index for each byte range + * @param offsets Iterator over the datasource offset for each byte range + * @param sizes Iterator over the size of each byte range + * @param count Number of byte ranges * @param dst Host buffer that receives the read data */ +template void read_ranges_to_host( cudf::host_span const> datasources, - cudf::host_span source_indices, - cudf::host_span offsets, - cudf::host_span sizes, - cudf::host_span dst_offsets, + SourceIndexIterator source_indices, + OffsetIterator offsets, + SizeIterator sizes, + std::size_t count, cudf::host_span dst) { std::vector> host_read_tasks; std::vector expected_sizes; - host_read_tasks.reserve(source_indices.size()); - expected_sizes.reserve(source_indices.size()); + host_read_tasks.reserve(count); + expected_sizes.reserve(count); - auto iter = cuda::make_zip_iterator( - source_indices.begin(), offsets.begin(), sizes.begin(), dst_offsets.begin()); + auto iter = cuda::make_zip_iterator(source_indices, offsets, sizes); + std::size_t dst_byte_offset = 0; // Schedule host reads holding the `host_read_mutex` so that all reads for a caller thread // are scheduled without interleaving with reads from other threads yielding better pipelining { std::scoped_lock lock(host_read_mutex()); - std::for_each(iter, iter + source_indices.size(), [&](auto const& tuple) { + std::for_each(iter, iter + count, [&](auto const& tuple) { auto const src_idx = cuda::std::get<0>(tuple); auto const io_offset = cuda::std::get<1>(tuple); auto const io_size = cuda::std::get<2>(tuple); - auto const dst_offset = cuda::std::get<3>(tuple); + auto const dst_offset = dst_byte_offset; + dst_byte_offset += io_size; if (io_size == 0) { return; } @@ -264,6 +267,7 @@ void read_ranges_to_host( })); }); } + CUDF_EXPECTS(dst_byte_offset == dst.size(), "Unexpected destination host buffer size"); // Complete the reads; every range must be read in full std::for_each(cuda::counting_iterator(0), @@ -484,7 +488,6 @@ fetch_bloom_filters_to_device_impl( [](auto acc, auto const& bloom_ranges) { return acc + bloom_ranges.size(); }); std::vector bitset_spans_per_source(num_sources); - std::vector> output_span_indices(total_filters); // Phase 1: Initial read. Cover the complete bloom filter or enough bytes to parse the header std::vector initial_source_indices(total_filters); @@ -495,52 +498,43 @@ fetch_bloom_filters_to_device_impl( { std::size_t filter_idx = 0; - std::for_each( - cuda::counting_iterator(0), - cuda::counting_iterator(num_sources), - [&](auto const source_idx) { - auto const& bloom_ranges = bloom_filter_byte_ranges_per_source[source_idx]; - bitset_spans_per_source[source_idx] = device_spans_per_source_type(bloom_ranges.size()); - std::for_each(cuda::counting_iterator(0), - cuda::counting_iterator(bloom_ranges.size()), - [&](auto const span_idx) { - auto const& range = bloom_ranges[span_idx]; - output_span_indices[filter_idx] = {source_idx, span_idx}; - initial_source_indices[filter_idx] = source_idx; - initial_offsets[filter_idx] = static_cast(range.offset()); - initial_sizes[filter_idx] = static_cast(range.size()); - initial_dst_offsets[filter_idx] = total_initial_read_size; - total_initial_read_size += initial_sizes[filter_idx]; - ++filter_idx; - }); - }); + std::for_each(cuda::counting_iterator(0), + cuda::counting_iterator(num_sources), + [&](auto const source_idx) { + auto const& bloom_ranges = bloom_filter_byte_ranges_per_source[source_idx]; + bitset_spans_per_source[source_idx].resize(bloom_ranges.size()); + std::for_each(bloom_ranges.begin(), bloom_ranges.end(), [&](auto const& range) { + initial_source_indices[filter_idx] = source_idx; + initial_offsets[filter_idx] = static_cast(range.offset()); + initial_sizes[filter_idx] = static_cast(range.size()); + initial_dst_offsets[filter_idx] = total_initial_read_size; + total_initial_read_size += initial_sizes[filter_idx]; + ++filter_idx; + }); + }); + CUDF_EXPECTS(filter_idx == total_filters, "Unexpected number of bloom filter byte ranges"); } // Read every initial bloom filter bytes into one host buffer auto initial_buffer = cudf::detail::make_host_vector(total_initial_read_size, stream); read_ranges_to_host(datasources, - initial_source_indices, - initial_offsets, - initial_sizes, - initial_dst_offsets, + initial_source_indices.cbegin(), + initial_offsets.cbegin(), + initial_sizes.cbegin(), + total_filters, initial_buffer); // Phase 2: Parse headers, organize bitset slots, and record deferred bitset reads std::vector copy_dsts; - std::vector copy_dst_offsets; std::vector copy_srcs; std::vector copy_sizes; copy_dsts.reserve(total_filters); - copy_dst_offsets.reserve(total_filters); copy_srcs.reserve(total_filters); copy_sizes.reserve(total_filters); std::size_t total_device_size = 0; std::vector deferred_filter_indices; - std::vector deferred_source_indices; std::vector deferred_offsets; - std::vector deferred_sizes; - std::vector deferred_dst_offsets; std::size_t total_deferred_size = 0; std::for_each( @@ -549,7 +543,6 @@ fetch_bloom_filters_to_device_impl( [&](std::size_t filter_idx) { auto const push_empty_filter = [&]() { copy_dsts.push_back(nullptr); - copy_dst_offsets.push_back(0); copy_srcs.push_back(nullptr); copy_sizes.push_back(0); }; @@ -579,7 +572,6 @@ fetch_bloom_filters_to_device_impl( auto const header_size = static_cast(header_bytes); auto const bitset_size = static_cast(bitset_bytes); copy_dsts.push_back(nullptr); - copy_dst_offsets.push_back(total_device_size); copy_sizes.push_back(bitset_size); total_device_size += bitset_size; @@ -590,47 +582,60 @@ fetch_bloom_filters_to_device_impl( // Whole bitset not in the host buffer: defer the read copy_srcs.push_back(nullptr); - deferred_filter_indices.push_back(copy_srcs.size() - 1); - deferred_source_indices.push_back(initial_source_indices[filter_idx]); + deferred_filter_indices.push_back(filter_idx); deferred_offsets.push_back(initial_offsets[filter_idx] + header_size); - deferred_sizes.push_back(bitset_size); - deferred_dst_offsets.push_back(total_deferred_size); total_deferred_size += bitset_size; } }); // Phase 3: Resolve deferred reads, then batch copy all bitsets to the device auto deferred_buffer = cudf::detail::make_host_vector(total_deferred_size, stream); + auto deferred_source_indices = + cuda::permutation_iterator{initial_source_indices.cbegin(), deferred_filter_indices.cbegin()}; + auto deferred_sizes = + cuda::permutation_iterator{copy_sizes.cbegin(), deferred_filter_indices.cbegin()}; read_ranges_to_host(datasources, deferred_source_indices, - deferred_offsets, + deferred_offsets.cbegin(), deferred_sizes, - deferred_dst_offsets, + deferred_filter_indices.size(), deferred_buffer); - std::for_each(cuda::counting_iterator(0), - cuda::counting_iterator(deferred_filter_indices.size()), - [&](std::size_t i) { - copy_srcs[deferred_filter_indices[i]] = - deferred_buffer.data() + deferred_dst_offsets[i]; - }); + std::size_t deferred_dst_offset = 0; + std::for_each( + deferred_filter_indices.begin(), deferred_filter_indices.end(), [&](auto const filter_idx) { + copy_srcs[filter_idx] = deferred_buffer.data() + deferred_dst_offset; + deferred_dst_offset += copy_sizes[filter_idx]; + }); + CUDF_EXPECTS(deferred_dst_offset == total_deferred_size, + "Unexpected deferred bloom filter buffer size"); // Add the buffer base to every non-empty output span and copy destination. rmm::device_buffer bitset_buffer(total_device_size, bloom_filter_block_bytes, stream, mr); - auto* const device_base = static_cast(bitset_buffer.data()); + auto* const device_base = static_cast(bitset_buffer.data()); + std::size_t device_offset = 0; std::for_each(cuda::counting_iterator(0), cuda::counting_iterator(total_filters), [&](std::size_t filter_idx) { auto const bitset_size = copy_sizes[filter_idx]; - auto* const device_dst = - bitset_size == 0 ? nullptr : device_base + copy_dst_offsets[filter_idx]; + auto* const device_dst = bitset_size == 0 ? nullptr : device_base + device_offset; if (bitset_size != 0) { CUDF_EXPECTS(rmm::is_pointer_aligned(device_dst, bloom_filter_block_bytes), "Encountered a misaligned bloom filter bitset"); } - auto const [source_idx, span_idx] = output_span_indices[filter_idx]; - bitset_spans_per_source[source_idx][span_idx] = {device_dst, bitset_size}; - copy_dsts[filter_idx] = device_dst; + copy_dsts[filter_idx] = device_dst; + device_offset += bitset_size; }); + CUDF_EXPECTS(device_offset == total_device_size, "Unexpected bloom filter device buffer size"); + + // Populate the nested per-source spans through a flattened view + auto flat_output_spans = bitset_spans_per_source | std::views::join; + std::transform(copy_dsts.begin(), + copy_dsts.end(), + copy_sizes.begin(), + flat_output_spans.begin(), + [](auto const dst, auto const size) { + return cudf::device_span{static_cast(dst), size}; + }); // One batched copy (entries with a null source or zero size are ignored by the batch API) if (total_device_size != 0) { From 3851e9f8a83e7c9eb668a04fea522080136ab9ec Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Mon, 20 Jul 2026 12:59:34 +0200 Subject: [PATCH 33/39] Refactor bloom filter data processing in `aggregate_reader_metadata::read_bloom_filters`. Simplified the flattening of bitset spans using ranges and improved readability by replacing `std::for_each` with a more concise transformation approach. This change enhances code clarity and efficiency. --- cpp/src/io/parquet/bloom_filter_reader.cu | 18 +++++++++--------- 1 file changed, 9 insertions(+), 9 deletions(-) diff --git a/cpp/src/io/parquet/bloom_filter_reader.cu b/cpp/src/io/parquet/bloom_filter_reader.cu index 0163f15d0473..3772bc195422 100644 --- a/cpp/src/io/parquet/bloom_filter_reader.cu +++ b/cpp/src/io/parquet/bloom_filter_reader.cu @@ -34,6 +34,7 @@ #include #include #include +#include #include namespace cudf::io::parquet::detail { @@ -398,15 +399,14 @@ aggregate_reader_metadata::read_bloom_filters( // Flatten the per-source bitset spans into per-chunk order std::vector> bloom_filter_data; bloom_filter_data.reserve(num_chunks); - std::for_each( - bitset_spans_per_source.begin(), bitset_spans_per_source.end(), [&](auto const& source_spans) { - std::transform(source_spans.begin(), - source_spans.end(), - std::back_inserter(bloom_filter_data), - [](auto const& span) { - return cudf::device_span{ - reinterpret_cast(span.data()), span.size()}; - }); + auto flat_bitset_spans = bitset_spans_per_source | std::views::join; + std::transform( + flat_bitset_spans.begin(), + flat_bitset_spans.end(), + std::back_inserter(bloom_filter_data), + [](auto const& span) { + return cudf::device_span{ + reinterpret_cast(span.data()), span.size()}; }); return {std::move(bloom_filter_buffers), std::move(bloom_filter_data)}; From 1159a8b9c6022c48318ca782b92bb2646fd31ea5 Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Mon, 20 Jul 2026 14:04:00 +0200 Subject: [PATCH 34/39] Remove unnecessary alignment checks in `fetch_bloom_filters_to_device_impl`. Simplified the bitset copying logic by directly assigning device pointers, enhancing code clarity and reducing complexity. --- cpp/src/io/parquet/io_utils/parquet_io_utils.cpp | 9 +-------- 1 file changed, 1 insertion(+), 8 deletions(-) diff --git a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp index 83d2012ecbbb..32b9f96e143e 100644 --- a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp +++ b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp @@ -20,7 +20,6 @@ #include #include -#include #include #include @@ -606,8 +605,6 @@ fetch_bloom_filters_to_device_impl( copy_srcs[filter_idx] = deferred_buffer.data() + deferred_dst_offset; deferred_dst_offset += copy_sizes[filter_idx]; }); - CUDF_EXPECTS(deferred_dst_offset == total_deferred_size, - "Unexpected deferred bloom filter buffer size"); // Add the buffer base to every non-empty output span and copy destination. rmm::device_buffer bitset_buffer(total_device_size, bloom_filter_block_bytes, stream, mr); @@ -618,11 +615,7 @@ fetch_bloom_filters_to_device_impl( [&](std::size_t filter_idx) { auto const bitset_size = copy_sizes[filter_idx]; auto* const device_dst = bitset_size == 0 ? nullptr : device_base + device_offset; - if (bitset_size != 0) { - CUDF_EXPECTS(rmm::is_pointer_aligned(device_dst, bloom_filter_block_bytes), - "Encountered a misaligned bloom filter bitset"); - } - copy_dsts[filter_idx] = device_dst; + copy_dsts[filter_idx] = device_dst; device_offset += bitset_size; }); CUDF_EXPECTS(device_offset == total_device_size, "Unexpected bloom filter device buffer size"); From 37d10c29218961915d98e0c83a917d9ece57b3ee Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Mon, 20 Jul 2026 19:18:29 +0200 Subject: [PATCH 35/39] Formatting --- cpp/src/io/parquet/bloom_filter_reader.cu | 15 +++++++-------- 1 file changed, 7 insertions(+), 8 deletions(-) diff --git a/cpp/src/io/parquet/bloom_filter_reader.cu b/cpp/src/io/parquet/bloom_filter_reader.cu index 3772bc195422..9f1747034d16 100644 --- a/cpp/src/io/parquet/bloom_filter_reader.cu +++ b/cpp/src/io/parquet/bloom_filter_reader.cu @@ -400,14 +400,13 @@ aggregate_reader_metadata::read_bloom_filters( std::vector> bloom_filter_data; bloom_filter_data.reserve(num_chunks); auto flat_bitset_spans = bitset_spans_per_source | std::views::join; - std::transform( - flat_bitset_spans.begin(), - flat_bitset_spans.end(), - std::back_inserter(bloom_filter_data), - [](auto const& span) { - return cudf::device_span{ - reinterpret_cast(span.data()), span.size()}; - }); + std::transform(flat_bitset_spans.begin(), + flat_bitset_spans.end(), + std::back_inserter(bloom_filter_data), + [](auto const& span) { + return cudf::device_span{ + reinterpret_cast(span.data()), span.size()}; + }); return {std::move(bloom_filter_buffers), std::move(bloom_filter_data)}; } From 4de4bcf803a76d1c1dba218cf72875d497713ef0 Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Tue, 21 Jul 2026 11:16:12 +0200 Subject: [PATCH 36/39] Refactor fetch_bloom_filters_to_device_impl to streamline memory management - Removed unnecessary `copy_dsts` vector and related null pointer handling. - Simplified the logic for populating `copy_srcs` and `copy_sizes`. - Ensured that device memory allocation is only performed when `device_base` is not null. - Improved clarity and maintainability of the code by restructuring deferred read handling. This change enhances the efficiency of the bloom filter fetching process in the Parquet I/O utilities. --- .../io/parquet/io_utils/parquet_io_utils.cpp | 57 +++++++++---------- 1 file changed, 28 insertions(+), 29 deletions(-) diff --git a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp index 32b9f96e143e..7bfd1975f05d 100644 --- a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp +++ b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp @@ -524,10 +524,8 @@ fetch_bloom_filters_to_device_impl( initial_buffer); // Phase 2: Parse headers, organize bitset slots, and record deferred bitset reads - std::vector copy_dsts; std::vector copy_srcs; std::vector copy_sizes; - copy_dsts.reserve(total_filters); copy_srcs.reserve(total_filters); copy_sizes.reserve(total_filters); std::size_t total_device_size = 0; @@ -541,7 +539,6 @@ fetch_bloom_filters_to_device_impl( cuda::counting_iterator(total_filters), [&](std::size_t filter_idx) { auto const push_empty_filter = [&]() { - copy_dsts.push_back(nullptr); copy_srcs.push_back(nullptr); copy_sizes.push_back(0); }; @@ -570,7 +567,6 @@ fetch_bloom_filters_to_device_impl( auto const header_size = static_cast(header_bytes); auto const bitset_size = static_cast(bitset_bytes); - copy_dsts.push_back(nullptr); copy_sizes.push_back(bitset_size); total_device_size += bitset_size; @@ -589,35 +585,38 @@ fetch_bloom_filters_to_device_impl( // Phase 3: Resolve deferred reads, then batch copy all bitsets to the device auto deferred_buffer = cudf::detail::make_host_vector(total_deferred_size, stream); - auto deferred_source_indices = - cuda::permutation_iterator{initial_source_indices.cbegin(), deferred_filter_indices.cbegin()}; - auto deferred_sizes = - cuda::permutation_iterator{copy_sizes.cbegin(), deferred_filter_indices.cbegin()}; - read_ranges_to_host(datasources, - deferred_source_indices, - deferred_offsets.cbegin(), - deferred_sizes, - deferred_filter_indices.size(), - deferred_buffer); - std::size_t deferred_dst_offset = 0; - std::for_each( - deferred_filter_indices.begin(), deferred_filter_indices.end(), [&](auto const filter_idx) { - copy_srcs[filter_idx] = deferred_buffer.data() + deferred_dst_offset; - deferred_dst_offset += copy_sizes[filter_idx]; - }); + { + auto deferred_source_indices = + cuda::permutation_iterator{initial_source_indices.cbegin(), deferred_filter_indices.cbegin()}; + auto deferred_sizes = + cuda::permutation_iterator{copy_sizes.cbegin(), deferred_filter_indices.cbegin()}; + read_ranges_to_host(datasources, + deferred_source_indices, + deferred_offsets.cbegin(), + deferred_sizes, + deferred_filter_indices.size(), + deferred_buffer); + std::size_t deferred_dst_offset = 0; + std::for_each( + deferred_filter_indices.begin(), deferred_filter_indices.end(), [&](auto const filter_idx) { + copy_srcs[filter_idx] = deferred_buffer.data() + deferred_dst_offset; + deferred_dst_offset += copy_sizes[filter_idx]; + }); + } - // Add the buffer base to every non-empty output span and copy destination. + // Add the buffer base to every output span and copy destination. rmm::device_buffer bitset_buffer(total_device_size, bloom_filter_block_bytes, stream, mr); + std::vector copy_dsts(total_filters); auto* const device_base = static_cast(bitset_buffer.data()); std::size_t device_offset = 0; - std::for_each(cuda::counting_iterator(0), - cuda::counting_iterator(total_filters), - [&](std::size_t filter_idx) { - auto const bitset_size = copy_sizes[filter_idx]; - auto* const device_dst = bitset_size == 0 ? nullptr : device_base + device_offset; - copy_dsts[filter_idx] = device_dst; - device_offset += bitset_size; - }); + if (device_base != nullptr) { + std::for_each(cuda::counting_iterator(0), + cuda::counting_iterator(total_filters), + [&](std::size_t filter_idx) { + copy_dsts[filter_idx] = device_base + device_offset; + device_offset += copy_sizes[filter_idx]; + }); + } CUDF_EXPECTS(device_offset == total_device_size, "Unexpected bloom filter device buffer size"); // Populate the nested per-source spans through a flattened view From 083e63d33445bd27b6a838cdb3ba4144e089b1ca Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Tue, 21 Jul 2026 16:32:34 +0200 Subject: [PATCH 37/39] use `cuda::std::as_bytes(span)` --- cpp/src/io/parquet/bloom_filter_reader.cu | 5 +---- 1 file changed, 1 insertion(+), 4 deletions(-) diff --git a/cpp/src/io/parquet/bloom_filter_reader.cu b/cpp/src/io/parquet/bloom_filter_reader.cu index 9f1747034d16..291a86b456e1 100644 --- a/cpp/src/io/parquet/bloom_filter_reader.cu +++ b/cpp/src/io/parquet/bloom_filter_reader.cu @@ -403,10 +403,7 @@ aggregate_reader_metadata::read_bloom_filters( std::transform(flat_bitset_spans.begin(), flat_bitset_spans.end(), std::back_inserter(bloom_filter_data), - [](auto const& span) { - return cudf::device_span{ - reinterpret_cast(span.data()), span.size()}; - }); + [](auto const& span) { return cuda::std::as_bytes(span); }); return {std::move(bloom_filter_buffers), std::move(bloom_filter_data)}; } From 66a58d348eee8ce4e97deb49819b1b81189f79bc Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Wed, 22 Jul 2026 09:57:21 +0200 Subject: [PATCH 38/39] Update bloom filter header parsing to use bytes per block for validation --- cpp/src/io/parquet/bloom_filter_reader.cu | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/cpp/src/io/parquet/bloom_filter_reader.cu b/cpp/src/io/parquet/bloom_filter_reader.cu index 291a86b456e1..eb7dfdc90b6b 100644 --- a/cpp/src/io/parquet/bloom_filter_reader.cu +++ b/cpp/src/io/parquet/bloom_filter_reader.cu @@ -309,7 +309,8 @@ std::optional> parse_bloom_filter_header( host_span bytes) { using policy_type = arrow_filter_policy; - auto constexpr words_per_block = policy_type::words_per_block; + using word_type = typename policy_type::word_type; + auto constexpr bytes_per_block = sizeof(word_type) * policy_type::words_per_block; // Deserialize the bloom filter header from the front of the buffer BloomFilterHeader header; @@ -318,7 +319,7 @@ std::optional> parse_bloom_filter_header( // Check if the bloom filter header is valid auto const is_header_valid = - (header.num_bytes % words_per_block) == 0 and + (header.num_bytes % bytes_per_block) == 0 and header.compression.compression == BloomFilterCompression::UNCOMPRESSED and header.algorithm.algorithm == BloomFilterAlgorithm::SPLIT_BLOCK and header.hash.hash == BloomFilterHash::XXHASH; From 1145ad557a164830b1840d0126230cefb8e21c43 Mon Sep 17 00:00:00 2001 From: Qi Chen Date: Fri, 24 Jul 2026 23:11:03 +0200 Subject: [PATCH 39/39] Enhance bloom filter logging and improve mutex scope in `fetch_bloom_filters_to_device_impl`. Updated warning message to include dynamic block size and refined mutex lock scope for better readability. --- cpp/src/io/parquet/io_utils/parquet_io_utils.cpp | 13 ++++++++----- 1 file changed, 8 insertions(+), 5 deletions(-) diff --git a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp index 7bfd1975f05d..b6f0413c57c1 100644 --- a/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp +++ b/cpp/src/io/parquet/io_utils/parquet_io_utils.cpp @@ -559,8 +559,9 @@ fetch_bloom_filters_to_device_impl( } auto const [header_bytes, bitset_bytes] = header_info.value(); if (bitset_bytes % bloom_filter_block_bytes != 0) { - CUDF_LOG_WARN( - "Encountered a bloom filter bitset size that is not a multiple of 32 bytes. Skipping"); + CUDF_LOG_WARN(std::format( + "Encountered a bloom filter bitset size that is not a multiple of {} bytes. Skipping", + bloom_filter_block_bytes)); push_empty_filter(); return; } @@ -631,9 +632,11 @@ fetch_bloom_filters_to_device_impl( // One batched copy (entries with a null source or zero size are ignored by the batch API) if (total_device_size != 0) { - std::scoped_lock lock(device_read_mutex()); - CUDF_CUDA_TRY(cudf::detail::memcpy_batch_async( - copy_dsts.data(), copy_srcs.data(), copy_sizes.data(), total_filters, stream)); + { + std::scoped_lock lock(device_read_mutex()); + CUDF_CUDA_TRY(cudf::detail::memcpy_batch_async( + copy_dsts.data(), copy_srcs.data(), copy_sizes.data(), total_filters, stream)); + } stream.synchronize(); }