Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
17 commits
Select commit Hold shift + click to select a range
45174ea
Support Parquet DELTA_BINARY_PACKED mini-blocks over 64 values
pramodsatya Jul 16, 2026
3c90717
Support DELTA_BYTE_ARRAY/DELTA_LENGTH_BYTE_ARRAY mini-blocks over 64 …
pramodsatya Jul 16, 2026
549ded7
Drop the DELTA mini-block size cap and fix nested decode races
pramodsatya Jul 16, 2026
aa8b101
Add DELTA encoding axes to the Parquet reader benchmarks
pramodsatya Jul 16, 2026
e732a18
Reserve the DELTA_BYTE_ARRAY prefix seed ahead of the skip scratch
pramodsatya Jul 21, 2026
71abede
Use canonical NVIDIA copyright notice in Parquet DELTA sources
pramodsatya Jul 21, 2026
995238a
Merge branch 'main' into parquet-delta-large-mini-blocks
pramodsatya Jul 21, 2026
25b2f69
Address review comments
pramodsatya Jul 22, 2026
5432caf
Reuse the compact protocol writer in the DELTA test fixtures and addr…
pramodsatya Jul 23, 2026
18a5946
Use cooperative_groups and a bits_per_byte constant in Parquet DELTA …
pramodsatya Jul 27, 2026
7dc4095
Add CompactProtocolWriter page-header overloads and use them in DELTA…
pramodsatya Jul 27, 2026
fb5c135
Address review comments on Parquet DELTA large mini-block support
pramodsatya Jul 29, 2026
31f9891
Merge branch 'main' into parquet-delta-large-mini-blocks
pramodsatya Jul 29, 2026
074c27a
Merge branch 'main' into pr-23314
vyasr Aug 2, 2026
6822b1b
Merge branch 'main' into parquet-delta-large-mini-blocks
vyasr Aug 2, 2026
f364b65
Merge remote-tracking branch 'upstream/main' into parquet-delta-large…
vyasr Aug 10, 2026
9ae76ed
Merge remote-tracking branch 'upstream/main' into parquet-delta-large…
vyasr Aug 11, 2026
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 2 additions & 2 deletions cpp/benchmarks/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -296,8 +296,8 @@ ConfigureNVBench(
# ##################################################################################################
# * parquet reader benchmark ----------------------------------------------------------------------
ConfigureNVBench(
PARQUET_READER_NVBENCH io/parquet/parquet_reader_input.cpp io/parquet/parquet_reader_options.cpp
io/parquet/reader_common.cpp
PARQUET_READER_NVBENCH io/parquet/parquet_reader_input.cpp io/parquet/parquet_reader_encoding.cpp
io/parquet/parquet_reader_options.cpp io/parquet/reader_common.cpp
)

# ##################################################################################################
Expand Down
103 changes: 103 additions & 0 deletions cpp/benchmarks/io/parquet/parquet_reader_encoding.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,103 @@
/*
* SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
* SPDX-License-Identifier: Apache-2.0
*/

#include "reader_common.hpp"

#include <benchmarks/common/generate_input.hpp>
#include <benchmarks/io/cuio_common.hpp>
#include <benchmarks/io/nvbench_helpers.hpp>

#include <cudf/io/parquet.hpp>
#include <cudf/utilities/error.hpp>

#include <nvbench/nvbench.cuh>

#include <string>
#include <string_view>
#include <utility>
#include <vector>

// Benchmarks decoding pages written with an explicitly requested column encoding. The writer's
// defaults never choose the DELTA_* encodings, so `parquet_read_decode` does not exercise their
// decode kernels; this benchmark covers them (with PLAIN as the baseline encoding).

namespace {

cudf::io::column_encoding retrieve_column_encoding_enum(std::string_view encoding_string)
{
if (encoding_string == "PLAIN") { return cudf::io::column_encoding::PLAIN; }
if (encoding_string == "DELTA_BINARY_PACKED") {
return cudf::io::column_encoding::DELTA_BINARY_PACKED;
}
if (encoding_string == "DELTA_LENGTH_BYTE_ARRAY") {
return cudf::io::column_encoding::DELTA_LENGTH_BYTE_ARRAY;
}
if (encoding_string == "DELTA_BYTE_ARRAY") { return cudf::io::column_encoding::DELTA_BYTE_ARRAY; }
CUDF_FAIL("Unsupported column encoding: " + std::string(encoding_string));
}

void bench_read_encoding(nvbench::state& state, std::vector<cudf::type_id> const& d_types)
{
auto const encoding = retrieve_column_encoding_enum(state.get_string("encoding"));
auto const source_type = retrieve_io_type_enum(state.get_string("io_type"));
auto const data_size = static_cast<size_t>(state.get_int64("data_size"));
auto const cardinality = static_cast<cudf::size_type>(state.get_int64("cardinality"));
auto const run_length = static_cast<cudf::size_type>(state.get_int64("run_length"));
cuio_source_sink_pair source_sink(source_type);

auto const num_rows_written = [&]() {
auto const tbl = create_random_table(
cycle_dtypes(d_types, num_cols),
table_size_bytes{data_size},
data_profile_builder().cardinality(cardinality).avg_run_length(run_length));
auto const view = tbl->view();

cudf::io::table_input_metadata metadata(view);
for (auto& col_meta : metadata.column_metadata) {
col_meta.set_encoding(encoding);
}

cudf::io::parquet_writer_options write_opts =
cudf::io::parquet_writer_options::builder(source_sink.make_sink_info(), view)
.metadata(std::move(metadata))
.compression(cudf::io::compression_type::NONE)
.dictionary_policy(cudf::io::dictionary_policy::NEVER)
.write_v2_headers(true);
cudf::io::write_parquet(write_opts);
return view.num_rows();
}();

parquet_read_common(num_rows_written, num_cols, source_sink, state);
}

} // namespace

void BM_parquet_read_delta_binary(nvbench::state& state)
{
bench_read_encoding(state, {cudf::type_id::INT32, cudf::type_id::INT64});
}

void BM_parquet_read_delta_string(nvbench::state& state)
{
bench_read_encoding(state, {cudf::type_id::STRING});

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

post benchmark results in the pr conversation page for before/after adding the decoding changes.

}

NVBENCH_BENCH(BM_parquet_read_delta_binary)
.set_name("parquet_read_delta_binary")
.add_string_axis("encoding", {"PLAIN", "DELTA_BINARY_PACKED"})
.add_string_axis("io_type", {"DEVICE_BUFFER"})
.set_min_samples(4)
.add_int64_axis("cardinality", {0, 1000})
.add_int64_axis("run_length", {1, 32})
.add_int64_axis("data_size", {512 << 20});

NVBENCH_BENCH(BM_parquet_read_delta_string)
.set_name("parquet_read_delta_string")
.add_string_axis("encoding", {"PLAIN", "DELTA_LENGTH_BYTE_ARRAY", "DELTA_BYTE_ARRAY"})
.add_string_axis("io_type", {"DEVICE_BUFFER"})
.set_min_samples(4)
.add_int64_axis("cardinality", {0, 1000})
.add_int64_axis("run_length", {1, 32})
.add_int64_axis("data_size", {512 << 20});
52 changes: 51 additions & 1 deletion cpp/src/io/parquet/compact_protocol_writer.cpp
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* SPDX-FileCopyrightText: Copyright (c) 2018-2025, NVIDIA CORPORATION.
* SPDX-FileCopyrightText: Copyright (c) 2018-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
* SPDX-License-Identifier: Apache-2.0
*/

Expand Down Expand Up @@ -263,6 +263,56 @@ size_t CompactProtocolWriter::write(SortingColumn const& sc)
return c.value();
}

size_t CompactProtocolWriter::write(DataPageHeader const& pg_hdr)
{
CompactProtocolFieldWriter c(*this);
c.field_int(1, pg_hdr.num_values);
c.field_int(2, static_cast<int32_t>(pg_hdr.encoding));
c.field_int(3, static_cast<int32_t>(pg_hdr.definition_level_encoding));
c.field_int(4, static_cast<int32_t>(pg_hdr.repetition_level_encoding));
return c.value();
}

size_t CompactProtocolWriter::write(DictionaryPageHeader const& dict_pg_hdr)
{
CompactProtocolFieldWriter c(*this);
c.field_int(1, dict_pg_hdr.num_values);
c.field_int(2, static_cast<int32_t>(dict_pg_hdr.encoding));
return c.value();
}

size_t CompactProtocolWriter::write(DataPageHeaderV2 const& pg_hdr_v2)
{
CompactProtocolFieldWriter c(*this);
c.field_int(1, pg_hdr_v2.num_values);
c.field_int(2, pg_hdr_v2.num_nulls);
c.field_int(3, pg_hdr_v2.num_rows);
c.field_int(4, static_cast<int32_t>(pg_hdr_v2.encoding));
c.field_int(5, pg_hdr_v2.definition_levels_byte_length);
c.field_int(6, pg_hdr_v2.repetition_levels_byte_length);
c.field_bool(7, pg_hdr_v2.is_compressed);
return c.value();
}

size_t CompactProtocolWriter::write(PageHeader const& page_hdr)
{
CompactProtocolFieldWriter c(*this);
c.field_int(1, static_cast<int32_t>(page_hdr.type));
c.field_int(2, page_hdr.uncompressed_page_size);
c.field_int(3, page_hdr.compressed_page_size);
// Exactly one page-specific header is set, selected by `type`; field ids (5/7/8) match
// CompactProtocolReader::read(PageHeader) and the GPU encoder in gpuEncodePageHeaders.
switch (page_hdr.type) {
case PageType::DATA_PAGE: c.field_struct(5, page_hdr.data_page_header); break;
case PageType::DICTIONARY_PAGE: c.field_struct(7, page_hdr.dictionary_page_header); break;
case PageType::DATA_PAGE_V2: c.field_struct(8, page_hdr.data_page_header_v2); break;
default:
CUDF_FAIL("Trying to write an invalid PageType " +
std::to_string(static_cast<int>(page_hdr.type)));
}
return c.value();
}

void CompactProtocolFieldWriter::put_byte(uint8_t v) { writer.m_buf.push_back(v); }

void CompactProtocolFieldWriter::put_byte(uint8_t const* raw, uint32_t len)
Expand Down
13 changes: 10 additions & 3 deletions cpp/src/io/parquet/compact_protocol_writer.hpp
Original file line number Diff line number Diff line change
@@ -1,19 +1,21 @@
/*
* SPDX-FileCopyrightText: Copyright (c) 2018-2025, NVIDIA CORPORATION.
* SPDX-FileCopyrightText: Copyright (c) 2018-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
* SPDX-License-Identifier: Apache-2.0
*/

#pragma once

#include <cudf/io/parquet_schema.hpp>
#include <cudf/utilities/export.hpp>

#include <algorithm>
#include <cstddef>
#include <cstdint>
#include <string>
#include <vector>

namespace cudf::io::parquet::detail {
namespace CUDF_EXPORT cudf {
namespace io::parquet::detail {

/**
* @brief Class for parsing Parquet's Thrift Compact Protocol encoded metadata
Expand Down Expand Up @@ -43,6 +45,10 @@ class CompactProtocolWriter {
size_t write(ColumnOrder const&);
size_t write(PageEncodingStats const&);
size_t write(SortingColumn const&);
size_t write(DataPageHeader const&);
size_t write(DictionaryPageHeader const&);
size_t write(DataPageHeaderV2 const&);
size_t write(PageHeader const&);

protected:
std::vector<uint8_t>& m_buf;
Expand Down Expand Up @@ -117,4 +123,5 @@ template <>
inline void CompactProtocolFieldWriter::field_int_list<int64_t>(int field,
std::vector<int64_t> const& val);

} // namespace cudf::io::parquet::detail
} // namespace io::parquet::detail
} // namespace CUDF_EXPORT cudf
Loading
Loading